From d87c7ad1cbd1f9694a64312a869c30eeceb69192 Mon Sep 17 00:00:00 2001 From: Vineeth Chelur Date: Wed, 2 Sep 2026 13:52:10 -0700 Subject: [PATCH 01/13] Update cpuinfo and apply thread-safe deinitialization patch (#32300) ### Description - Update `pytorch/cpuinfo` from `4628dc060ce4e82345dc166bbac875609db4ff69` to `66ee79c038d70dad9f08705b2c9b3e58f6d8f512`, the latest commit on cpuinfo `main` as of August 27, 2026. - Carry the thread-safe, reference-counted initialization and deinitialization changes from [pytorch/cpuinfo#400](https://github.com/pytorch/cpuinfo/pull/400) as one shared ORT patch used by both FetchContent and vcpkg. - Reset Windows ARM64 cache-population state on each initialization so repeated DLL load/unload cycles cannot reuse stale cache indices. - Scope XNNPACK's cpuinfo references to hardware discovery so XNNPACK-enabled ORT builds do not retain unmatched references during DLL unload. The cpuinfo patch can be removed after pytorch/cpuinfo#400, including the ARM64 reinitialization fix, is merged and ORT updates to a revision containing it. The XNNPACK compatibility patch can be removed after the corresponding lifecycle fix is available upstream. ### Motivation and Context [microsoft/onnxruntime#28245](https://github.com/microsoft/onnxruntime/pull/28245) integrated `cpuinfo_deinitialize()` after [pytorch/cpuinfo#387](https://github.com/pytorch/cpuinfo/pull/387) added it upstream. That implementation was later reverted by [pytorch/cpuinfo#411](https://github.com/pytorch/cpuinfo/pull/411) because initialization and deinitialization were not safe for multiple consumers. Pinning the latest cpuinfo `main` without an ORT-side patch would therefore make `cpuinfo_deinitialize()` a no-op again. Carrying the corrected implementation keeps ORT independent of the pending upstream review while preserving safe cleanup during dynamic DLL unload. ### Testing - `onnxruntime_cpuinfo_refcount_test` covers sequential consumers, concurrent consumers, and reinitialization after final release. - Verified this test fails against ORT `main`'s pinned cpuinfo revision (`4628dc060ce4e82345dc166bbac875609db4ff69`) and passes against the patched revision (`66ee79c038d70dad9f08705b2c9b3e58f6d8f512`). - In XNNPACK-enabled builds, the refcount test initializes XNNPACK and verifies hardware discovery does not retain a cpuinfo reference. - `onnxruntime_shared_lib_cpuinfo_dlopen_test` loads a small DLL containing ORT's `CPUIDInfo` and cpuinfo, captures a cpuinfo process-heap allocation, unloads the DLL, and verifies that allocation was released. - Verified the FetchContent patch sequences for the default, Linux, and Windows ARM64/ARM64EC paths, plus the vcpkg patch sequence, apply with zero rejected hunks. --------- Co-authored-by: Vineeth Chelur --- cmake/deps.txt | 2 +- .../external/onnxruntime_external_deps.cmake | 78 ++- cmake/external/xnnpack.cmake | 4 +- cmake/onnxruntime_unittests.cmake | 69 ++ cmake/patches/.gitattributes | 3 + .../cpuinfo/enable_deinit_refcounting.patch | 650 ++++++++++++++++++ ...elease_cpuinfo_after_hardware_config.patch | 46 ++ cmake/vcpkg-ports/cpuinfo/portfile.cmake | 6 +- cmake/vcpkg-ports/xnnpack/portfile.cmake | 1 + .../test/shared_lib/cpuinfo_dlopen_test.cc | 48 ++ .../shared_lib/cpuinfo_dlopen_test_library.cc | 11 + .../cpuinfo_dlopen_test_library.def | 3 + .../test/shared_lib/cpuinfo_refcount_test.cc | 223 ++++++ 13 files changed, 1105 insertions(+), 39 deletions(-) create mode 100644 cmake/patches/cpuinfo/enable_deinit_refcounting.patch create mode 100644 cmake/patches/xnnpack/release_cpuinfo_after_hardware_config.patch create mode 100644 onnxruntime/test/shared_lib/cpuinfo_dlopen_test.cc create mode 100644 onnxruntime/test/shared_lib/cpuinfo_dlopen_test_library.cc create mode 100644 onnxruntime/test/shared_lib/cpuinfo_dlopen_test_library.def create mode 100644 onnxruntime/test/shared_lib/cpuinfo_refcount_test.cc diff --git a/cmake/deps.txt b/cmake/deps.txt index c1604e7a5112a..5928ac2c13140 100644 --- a/cmake/deps.txt +++ b/cmake/deps.txt @@ -50,7 +50,7 @@ protoc_mac_universal;https://github.com/protocolbuffers/protobuf/releases/downlo psimd;https://github.com/Maratyszcza/psimd/archive/072586a71b55b7f8c584153d223e95687148a900.zip;1f5454b01f06f9656b77e4a5e2e31d7422487013 pthreadpool;https://github.com/google/pthreadpool/archive/dcc9f28589066af0dbd4555579281230abbf74dd.zip;533a77943203ef15ca608bcd9dbe2c94da7451d2 pybind11;https://github.com/pybind/pybind11/archive/refs/tags/v3.0.2.zip;a064e663b4d7a337ac291d1bef7337ef4e60a1ae -pytorch_cpuinfo;https://github.com/pytorch/cpuinfo/archive/4628dc060ce4e82345dc166bbac875609db4ff69.zip;e58d4b47c16a982111c897e669ae4f1821a393d7 +pytorch_cpuinfo;https://github.com/pytorch/cpuinfo/archive/66ee79c038d70dad9f08705b2c9b3e58f6d8f512.zip;6e2be17c110e4f4184b059303c722b3c5bd080a6 re2;https://github.com/google/re2/archive/refs/tags/2024-07-02.zip;646e1728269cde7fcef990bf4a8e87b047882e88 safeint;https://github.com/dcleblanc/SafeInt/archive/refs/tags/3.0.28.zip;23f252040ff6cb9f1fd18575b32fa8fb5928daac tensorboard;https://github.com/tensorflow/tensorboard/archive/373eb09e4c5d2b3cc2493f0949dc4be6b6a45e81.zip;67b833913605a4f3f499894ab11528a702c2b381 diff --git a/cmake/external/onnxruntime_external_deps.cmake b/cmake/external/onnxruntime_external_deps.cmake index 905451c8df2dc..6cad792f2aebd 100644 --- a/cmake/external/onnxruntime_external_deps.cmake +++ b/cmake/external/onnxruntime_external_deps.cmake @@ -363,42 +363,52 @@ if (CPUINFO_SUPPORTED) set(CPUINFO_BUILD_UNIT_TESTS OFF CACHE INTERNAL "") set(CPUINFO_BUILD_MOCK_TESTS OFF CACHE INTERNAL "") set(CPUINFO_BUILD_BENCHMARKS OFF CACHE INTERNAL "") - if (onnxruntime_target_platform STREQUAL "ARM64EC" OR onnxruntime_target_platform STREQUAL "ARM64") - message(STATUS "Applying patches for Windows ARM64/ARM64EC in cpuinfo") - onnxruntime_fetchcontent_declare( - pytorch_cpuinfo - URL ${DEP_URL_pytorch_cpuinfo} - URL_HASH SHA1=${DEP_SHA1_pytorch_cpuinfo} - EXCLUDE_FROM_ALL - PATCH_COMMAND - ${Patch_EXECUTABLE} -p1 < ${PROJECT_SOURCE_DIR}/patches/cpuinfo/patch_cpuinfo_h_for_arm64ec.patch && - # https://github.com/pytorch/cpuinfo/pull/324 - ${Patch_EXECUTABLE} -p1 < ${PROJECT_SOURCE_DIR}/patches/cpuinfo/patch_vcpkg_arm64ec_support.patch - FIND_PACKAGE_ARGS NAMES cpuinfo - ) - elseif(CMAKE_SYSTEM_NAME STREQUAL "Linux") - message(STATUS "Applying sysfs fallback patch for cpuinfo on Linux") - onnxruntime_fetchcontent_declare( - pytorch_cpuinfo - URL ${DEP_URL_pytorch_cpuinfo} - URL_HASH SHA1=${DEP_SHA1_pytorch_cpuinfo} - EXCLUDE_FROM_ALL - PATCH_COMMAND - # https://github.com/microsoft/onnxruntime/issues/10038 - ${Patch_EXECUTABLE} -p1 < ${PROJECT_SOURCE_DIR}/patches/cpuinfo/fix_missing_sysfs_fallback.patch - FIND_PACKAGE_ARGS NAMES cpuinfo - ) + if(onnxruntime_USE_VCPKG AND NOT APPLE) + find_package(cpuinfo CONFIG REQUIRED) else() - onnxruntime_fetchcontent_declare( - pytorch_cpuinfo - URL ${DEP_URL_pytorch_cpuinfo} - URL_HASH SHA1=${DEP_SHA1_pytorch_cpuinfo} - EXCLUDE_FROM_ALL - FIND_PACKAGE_ARGS NAMES cpuinfo - ) + if (onnxruntime_target_platform STREQUAL "ARM64EC" OR onnxruntime_target_platform STREQUAL "ARM64") + message(STATUS "Applying patches for Windows ARM64/ARM64EC in cpuinfo") + onnxruntime_fetchcontent_declare( + pytorch_cpuinfo + URL ${DEP_URL_pytorch_cpuinfo} + URL_HASH SHA1=${DEP_SHA1_pytorch_cpuinfo} + EXCLUDE_FROM_ALL + PATCH_COMMAND + ${Patch_EXECUTABLE} -p1 < ${PROJECT_SOURCE_DIR}/patches/cpuinfo/patch_cpuinfo_h_for_arm64ec.patch && + # https://github.com/pytorch/cpuinfo/pull/324 + ${Patch_EXECUTABLE} -p1 < ${PROJECT_SOURCE_DIR}/patches/cpuinfo/patch_vcpkg_arm64ec_support.patch && + # https://github.com/pytorch/cpuinfo/pull/400 + ${Patch_EXECUTABLE} --binary --ignore-whitespace -p1 < + ${PROJECT_SOURCE_DIR}/patches/cpuinfo/enable_deinit_refcounting.patch + ) + elseif(CMAKE_SYSTEM_NAME STREQUAL "Linux") + message(STATUS "Applying sysfs fallback patch for cpuinfo on Linux") + onnxruntime_fetchcontent_declare( + pytorch_cpuinfo + URL ${DEP_URL_pytorch_cpuinfo} + URL_HASH SHA1=${DEP_SHA1_pytorch_cpuinfo} + EXCLUDE_FROM_ALL + PATCH_COMMAND + # https://github.com/microsoft/onnxruntime/issues/10038 + ${Patch_EXECUTABLE} -p1 < ${PROJECT_SOURCE_DIR}/patches/cpuinfo/fix_missing_sysfs_fallback.patch && + # https://github.com/pytorch/cpuinfo/pull/400 + ${Patch_EXECUTABLE} --binary --ignore-whitespace -p1 < + ${PROJECT_SOURCE_DIR}/patches/cpuinfo/enable_deinit_refcounting.patch + ) + else() + onnxruntime_fetchcontent_declare( + pytorch_cpuinfo + URL ${DEP_URL_pytorch_cpuinfo} + URL_HASH SHA1=${DEP_SHA1_pytorch_cpuinfo} + EXCLUDE_FROM_ALL + PATCH_COMMAND + # https://github.com/pytorch/cpuinfo/pull/400 + ${Patch_EXECUTABLE} --binary --ignore-whitespace -p1 < + ${PROJECT_SOURCE_DIR}/patches/cpuinfo/enable_deinit_refcounting.patch + ) + endif() + onnxruntime_fetchcontent_makeavailable(pytorch_cpuinfo) endif() - set(ONNXRUNTIME_CPUINFO_PROJ pytorch_cpuinfo) - onnxruntime_fetchcontent_makeavailable(${ONNXRUNTIME_CPUINFO_PROJ}) if(TARGET cpuinfo::cpuinfo AND NOT TARGET cpuinfo) message(STATUS "Aliasing cpuinfo::cpuinfo to cpuinfo") add_library(cpuinfo ALIAS cpuinfo::cpuinfo) diff --git a/cmake/external/xnnpack.cmake b/cmake/external/xnnpack.cmake index 571283c33c713..59cc26316a770 100644 --- a/cmake/external/xnnpack.cmake +++ b/cmake/external/xnnpack.cmake @@ -84,9 +84,9 @@ endif() onnxruntime_fetchcontent_declare(googlexnnpack URL ${DEP_URL_googlexnnpack} URL_HASH SHA1=${DEP_SHA1_googlexnnpack} - PATCH_COMMAND ${Patch_EXECUTABLE} --binary --ignore-whitespace -p1 < ${PROJECT_SOURCE_DIR}/patches/xnnpack/AddEmscriptenAndIosSupport.patch + PATCH_COMMAND ${Patch_EXECUTABLE} --binary --ignore-whitespace -p1 < ${PROJECT_SOURCE_DIR}/patches/xnnpack/AddEmscriptenAndIosSupport.patch && + ${Patch_EXECUTABLE} --ignore-whitespace -p1 < ${PROJECT_SOURCE_DIR}/patches/xnnpack/release_cpuinfo_after_hardware_config.patch EXCLUDE_FROM_ALL - FIND_PACKAGE_ARGS NAMES xnnpack ) onnxruntime_fetchcontent_makeavailable(googlexnnpack) set(XNNPACK_DIR ${googlexnnpack_SOURCE_DIR}) diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index ebcc2cc9458d0..715c0efd12a1d 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1987,6 +1987,43 @@ endif() set_target_properties(ep_weight_sharing_ctx_gen PROPERTIES FOLDER "ONNXRuntimeTest") endif() + if (CPUINFO_SUPPORTED AND NOT CMAKE_SYSTEM_NAME STREQUAL "Emscripten") + set(onnxruntime_cpuinfo_test_library cpuinfo) + get_target_property(onnxruntime_cpuinfo_aliased_target cpuinfo ALIASED_TARGET) + if(onnxruntime_cpuinfo_aliased_target) + set(onnxruntime_cpuinfo_test_library ${onnxruntime_cpuinfo_aliased_target}) + endif() + get_target_property(onnxruntime_cpuinfo_library_type ${onnxruntime_cpuinfo_test_library} TYPE) + if(onnxruntime_cpuinfo_library_type STREQUAL "STATIC_LIBRARY") + onnxruntime_add_executable( + onnxruntime_cpuinfo_refcount_test + ${ONNXRUNTIME_SHARED_LIB_TEST_SRC_DIR}/cpuinfo_refcount_test.cc) + target_compile_definitions( + onnxruntime_cpuinfo_refcount_test + PRIVATE ORT_CPUINFO_TEST_HAS_INTERNAL_STATE) + if(onnxruntime_USE_XNNPACK) + target_compile_definitions( + onnxruntime_cpuinfo_refcount_test + PRIVATE ORT_CPUINFO_TEST_USE_XNNPACK) + if(onnxruntime_USE_VCPKG) + target_include_directories(onnxruntime_cpuinfo_refcount_test PRIVATE ${XNNPACK_HDR}) + else() + target_include_directories(onnxruntime_cpuinfo_refcount_test PRIVATE ${XNNPACK_INCLUDE_DIR}) + endif() + target_link_libraries( + onnxruntime_cpuinfo_refcount_test + PRIVATE ${onnxruntime_EXTERNAL_LIBRARIES_XNNPACK}) + endif() + target_link_libraries(onnxruntime_cpuinfo_refcount_test PRIVATE cpuinfo Threads::Threads) + add_test( + NAME onnxruntime_cpuinfo_refcount_test + COMMAND onnxruntime_cpuinfo_refcount_test) + set_target_properties( + onnxruntime_cpuinfo_refcount_test + PROPERTIES FOLDER "ONNXRuntimeTest") + endif() + endif() + # shared lib if (onnxruntime_BUILD_SHARED_LIB) if(WIN32) @@ -1995,6 +2032,38 @@ endif() add_test(NAME onnxruntime_shared_lib_dlopen_test COMMAND onnxruntime_shared_lib_dlopen_test WORKING_DIRECTORY $) set_target_properties(onnxruntime_shared_lib_dlopen_test PROPERTIES FOLDER "ONNXRuntimeTest") + if(onnxruntime_cpuinfo_library_type STREQUAL "STATIC_LIBRARY") + onnxruntime_add_shared_library( + onnxruntime_cpuinfo_dlopen_test_library + ${ONNXRUNTIME_SHARED_LIB_TEST_SRC_DIR}/cpuinfo_dlopen_test_library.cc + ${ONNXRUNTIME_SHARED_LIB_TEST_SRC_DIR}/cpuinfo_dlopen_test_library.def) + target_include_directories( + onnxruntime_cpuinfo_dlopen_test_library + PRIVATE ${ONNXRUNTIME_ROOT}) + target_link_libraries( + onnxruntime_cpuinfo_dlopen_test_library + PRIVATE onnxruntime_common cpuinfo) + + onnxruntime_add_executable( + onnxruntime_shared_lib_cpuinfo_dlopen_test + ${ONNXRUNTIME_SHARED_LIB_TEST_SRC_DIR}/cpuinfo_dlopen_test.cc) + add_dependencies( + onnxruntime_shared_lib_cpuinfo_dlopen_test + onnxruntime_cpuinfo_dlopen_test_library) + target_compile_definitions( + onnxruntime_shared_lib_cpuinfo_dlopen_test + PRIVATE + ORT_CPUINFO_DLOPEN_TEST_LIBRARY=L"$") + add_test( + NAME onnxruntime_shared_lib_cpuinfo_dlopen_test + COMMAND onnxruntime_shared_lib_cpuinfo_dlopen_test + WORKING_DIRECTORY $) + set_target_properties( + onnxruntime_cpuinfo_dlopen_test_library + onnxruntime_shared_lib_cpuinfo_dlopen_test + PROPERTIES FOLDER "ONNXRuntimeTest") + endif() + if (MSVC) # set VS debugger working directory to the test program's directory set_target_properties(onnxruntime_shared_lib_dlopen_test PROPERTIES VS_DEBUGGER_WORKING_DIRECTORY $) diff --git a/cmake/patches/.gitattributes b/cmake/patches/.gitattributes index 9812ceb1ffd9b..24d5f860ef2a8 100644 --- a/cmake/patches/.gitattributes +++ b/cmake/patches/.gitattributes @@ -1 +1,4 @@ *.patch text eol=lf + +# cpuinfo contains both LF and CRLF source files, so this patch must preserve both. +cpuinfo/enable_deinit_refcounting.patch -text diff --git a/cmake/patches/cpuinfo/enable_deinit_refcounting.patch b/cmake/patches/cpuinfo/enable_deinit_refcounting.patch new file mode 100644 index 0000000000000..960cfc6df58fa --- /dev/null +++ b/cmake/patches/cpuinfo/enable_deinit_refcounting.patch @@ -0,0 +1,650 @@ +diff --git a/src/arm/linux/init.c b/src/arm/linux/init.c +index 48c9f4a050412220b5c817814b643fc4850c7626..cb3170f279474ee344dd61480eab2831f2d027b9 100644 +--- a/src/arm/linux/init.c ++++ b/src/arm/linux/init.c +@@ -1001,3 +1001,43 @@ cleanup: + free(linux_cpu_to_core_map); + free(linux_cpu_to_uarch_index_map); + } ++ ++void cpuinfo_arm_linux_deinit(void) { ++ free(cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ free(cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ free(cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ /* cpuinfo_packages points to static storage (&package) — do not free */ ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ free(cpuinfo_uarchs); ++ cpuinfo_uarchs = NULL; ++ cpuinfo_uarchs_count = 0; ++ ++ for (int lvl = 0; lvl < cpuinfo_cache_level_max; ++lvl) { ++ free(cpuinfo_cache[lvl]); ++ cpuinfo_cache[lvl] = NULL; ++ cpuinfo_cache_count[lvl] = 0; ++ } ++ cpuinfo_max_cache_size = 0; ++ ++ free(cpuinfo_linux_cpu_to_processor_map); ++ cpuinfo_linux_cpu_to_processor_map = NULL; ++ ++ free(cpuinfo_linux_cpu_to_core_map); ++ cpuinfo_linux_cpu_to_core_map = NULL; ++ ++ free((void*)cpuinfo_linux_cpu_to_uarch_index_map); ++ cpuinfo_linux_cpu_to_uarch_index_map = NULL; ++ ++ cpuinfo_linux_cpu_max = 0; ++} +diff --git a/src/arm/mach/init.c b/src/arm/mach/init.c +index 76224b9d40eb5cd047f1353e81b5999196d3b881..9a9fb4276c531687bf7984e2bd41152ed91e94a6 100644 +--- a/src/arm/mach/init.c ++++ b/src/arm/mach/init.c +@@ -873,3 +873,32 @@ cleanup: + free(l2); + free(l3); + } ++ ++void cpuinfo_arm_mach_deinit(void) { ++ free(cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ free(cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ free(cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ free(cpuinfo_packages); ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ free(cpuinfo_uarchs); ++ cpuinfo_uarchs = NULL; ++ cpuinfo_uarchs_count = 0; ++ ++ for (int lvl = 0; lvl < cpuinfo_cache_level_max; ++lvl) { ++ free(cpuinfo_cache[lvl]); ++ cpuinfo_cache[lvl] = NULL; ++ cpuinfo_cache_count[lvl] = 0; ++ } ++ cpuinfo_max_cache_size = 0; ++} +diff --git a/src/arm/windows/init-by-logical-sys-info.c b/src/arm/windows/init-by-logical-sys-info.c +index 815ecb770a5bc5c48c7dff049f50b111b427954a..f0d66223f0d189776659248f525200ebdcda0116 100644 +--- a/src/arm/windows/init-by-logical-sys-info.c ++++ b/src/arm/windows/init-by-logical-sys-info.c +@@ -14,6 +14,8 @@ + + #define MAX_NR_OF_CACHES (cpuinfo_cache_level_max - 1) + ++static uint32_t cache_counters[MAX_NR_OF_CACHES]; ++ + /* Call chain: + * cpu_info_init_by_logical_sys_info + * read_packages_for_processors +@@ -126,6 +128,8 @@ bool cpu_info_init_by_logical_sys_info(const struct woa_chip_info* chip_info, co + uint32_t nr_of_uarchs = 0; + bool result = false; + ++ memset(cache_counters, 0, sizeof(cache_counters)); ++ + HANDLE heap = GetProcessHeap(); + + /* 1. Count available logical processor groups and processors */ +@@ -372,6 +376,40 @@ clean_up: + return result; + } + ++BOOL CALLBACK cpuinfo_arm_windows_deinit(PINIT_ONCE init_once, PVOID parameter, PVOID* context) { ++ HANDLE heap = GetProcessHeap(); ++ ++ HeapFree(heap, 0, cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ HeapFree(heap, 0, cpuinfo_packages); ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ HeapFree(heap, 0, cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ HeapFree(heap, 0, cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ HeapFree(heap, 0, cpuinfo_uarchs); ++ cpuinfo_uarchs = NULL; ++ cpuinfo_uarchs_count = 0; ++ ++ /* Caches are allocated as a single contiguous block starting at l1i */ ++ HeapFree(heap, 0, cpuinfo_cache[cpuinfo_cache_level_1i]); ++ for (int lvl = 0; lvl < cpuinfo_cache_level_max; ++lvl) { ++ cpuinfo_cache[lvl] = NULL; ++ cpuinfo_cache_count[lvl] = 0; ++ } ++ cpuinfo_max_cache_size = 0; ++ ++ return TRUE; ++} ++ + static uint32_t count_logical_processors(const uint32_t max_group_count, uint32_t* global_proc_index_per_group) { + uint32_t nr_of_processors = 0; + +@@ -606,11 +644,6 @@ static bool parse_relation_cache_info( + uint32_t* numbers_of_caches, + const uint32_t* global_proc_index_per_group, + PSYSTEM_LOGICAL_PROCESSOR_INFORMATION_EX info) { +- static uint32_t l1i_counter = 0; +- static uint32_t l1d_counter = 0; +- static uint32_t l2_counter = 0; +- static uint32_t l3_counter = 0; +- + /* Count cache types for allocation at first. */ + if (caches == NULL) { + switch (info->Cache.Level) { +@@ -665,12 +698,12 @@ static bool parse_relation_cache_info( + case 1: + switch (info->Cache.Type) { + case CacheInstruction: +- current_cache = l1i_base + l1i_counter; +- l1i_counter++; ++ current_cache = l1i_base + cache_counters[cpuinfo_cache_level_1i]; ++ cache_counters[cpuinfo_cache_level_1i]++; + break; + case CacheData: +- current_cache = l1d_base + l1d_counter; +- l1d_counter++; ++ current_cache = l1d_base + cache_counters[cpuinfo_cache_level_1d]; ++ cache_counters[cpuinfo_cache_level_1d]++; + break; + case CacheUnified: + break; +@@ -681,12 +714,12 @@ static bool parse_relation_cache_info( + } + break; + case 2: +- current_cache = l2_base + l2_counter; +- l2_counter++; ++ current_cache = l2_base + cache_counters[cpuinfo_cache_level_2]; ++ cache_counters[cpuinfo_cache_level_2]++; + break; + case 3: +- current_cache = l3_base + l3_counter; +- l3_counter++; ++ current_cache = l3_base + cache_counters[cpuinfo_cache_level_3]; ++ cache_counters[cpuinfo_cache_level_3]++; + break; + } + current_cache->size = info->Cache.CacheSize; +diff --git a/src/arm/windows/init.c b/src/arm/windows/init.c +index 528ae3279f5e979409ac07de6cee833691e034cd..fcf219df153419249211be4fdf014f50fc5198da 100644 +--- a/src/arm/windows/init.c ++++ b/src/arm/windows/init.c +@@ -25,12 +25,18 @@ BOOL CALLBACK cpuinfo_arm_windows_init(PINIT_ONCE init_once, PVOID parameter, PV + set_cpuinfo_isa_fields(); + + chip_info = get_system_info_from_registry(); ++ const bool chip_info_from_registry = chip_info != NULL; + if (chip_info == NULL) { + chip_info = &woa_chip_unknown; + } + + cpuinfo_is_initialized = cpu_info_init_by_logical_sys_info(chip_info, chip_info->uarchs[0].vendor); + ++ if (chip_info_from_registry) { ++ HeapFree(GetProcessHeap(), 0, chip_info->chip_name_string); ++ HeapFree(GetProcessHeap(), 0, chip_info); ++ } ++ + return true; + } + +@@ -177,6 +183,7 @@ static struct woa_chip_info* get_system_info_from_registry(void) { + uint64_t midr_qword = (uint32_t)read_registry_qword(cpu0_subkey, chip_midr_value); + if (midr_qword == 0) { + cpuinfo_log_error("Registry read error for MIDR value"); ++ HeapFree(GetProcessHeap(), 0, text_buffer); + return NULL; + } + // MIDR is only 32 bits, so we need to cast it to uint32_t +@@ -187,6 +194,7 @@ static struct woa_chip_info* get_system_info_from_registry(void) { + uint64_t frequency_mhz = read_registry_dword(cpu0_subkey, chip_mhz_value); + if (frequency_mhz == 0) { + cpuinfo_log_error("Registry read error for frequency value"); ++ HeapFree(GetProcessHeap(), 0, text_buffer); + return NULL; + } + // Convert MHz to Hz +@@ -196,11 +204,13 @@ static struct woa_chip_info* get_system_info_from_registry(void) { + chip_info = HeapAlloc(GetProcessHeap(), HEAP_ZERO_MEMORY, sizeof(struct woa_chip_info)); + if (chip_info == NULL) { + cpuinfo_log_error("Heap allocation error for chip_info"); ++ HeapFree(GetProcessHeap(), 0, text_buffer); + return NULL; + } + + // set chip_info fields + chip_info->chip_name_string = wcsndup(text_buffer, CPUINFO_PACKAGE_NAME_MAX - 1); ++ HeapFree(GetProcessHeap(), 0, text_buffer); + chip_info->uarchs[0] = get_core_info_from_midr(midr_value, frequency_hz); + + cpuinfo_log_debug("detected chip model name: %ls", chip_info->chip_name_string); +diff --git a/src/cpuinfo/internal-api.h b/src/cpuinfo/internal-api.h +index d84b26a89cfc38e726216688e490f50799f31455..9c1588872da413535f2fb68fb96d6106b420ac49 100644 +--- a/src/cpuinfo/internal-api.h ++++ b/src/cpuinfo/internal-api.h +@@ -45,6 +45,9 @@ extern CPUINFO_INTERNAL struct cpuinfo_uarch_info cpuinfo_global_uarch; + extern CPUINFO_INTERNAL uint32_t cpuinfo_linux_cpu_max; + extern CPUINFO_INTERNAL const struct cpuinfo_processor** cpuinfo_linux_cpu_to_processor_map; + extern CPUINFO_INTERNAL const struct cpuinfo_core** cpuinfo_linux_cpu_to_core_map; ++#if CPUINFO_ARCH_ARM || CPUINFO_ARCH_ARM64 || CPUINFO_ARCH_RISCV32 || CPUINFO_ARCH_RISCV64 ++extern CPUINFO_INTERNAL const uint32_t* cpuinfo_linux_cpu_to_uarch_index_map; ++#endif + #endif + + CPUINFO_PRIVATE void cpuinfo_x86_mach_init(void); +@@ -62,6 +65,23 @@ CPUINFO_PRIVATE void cpuinfo_arm_linux_init(void); + CPUINFO_PRIVATE void cpuinfo_riscv_linux_init(void); + CPUINFO_PRIVATE void cpuinfo_emscripten_init(void); + ++/* Platform-specific deinitialization functions. ++ * Each frees the globals allocated by its corresponding init function. */ ++CPUINFO_PRIVATE void cpuinfo_x86_mach_deinit(void); ++CPUINFO_PRIVATE void cpuinfo_x86_linux_deinit(void); ++CPUINFO_PRIVATE void cpuinfo_x86_freebsd_deinit(void); ++#if defined(_WIN32) || defined(__CYGWIN__) ++#if CPUINFO_ARCH_ARM64 ++CPUINFO_PRIVATE BOOL CALLBACK cpuinfo_arm_windows_deinit(PINIT_ONCE init_once, PVOID parameter, PVOID* context); ++#else ++CPUINFO_PRIVATE BOOL CALLBACK cpuinfo_x86_windows_deinit(PINIT_ONCE init_once, PVOID parameter, PVOID* context); ++#endif ++#endif ++CPUINFO_PRIVATE void cpuinfo_arm_mach_deinit(void); ++CPUINFO_PRIVATE void cpuinfo_arm_linux_deinit(void); ++CPUINFO_PRIVATE void cpuinfo_riscv_linux_deinit(void); ++CPUINFO_PRIVATE void cpuinfo_emscripten_deinit(void); ++ + CPUINFO_PRIVATE uint32_t cpuinfo_compute_max_cache_size(const struct cpuinfo_processor* processor); + + typedef void (*cpuinfo_processor_callback)(uint32_t); +diff --git a/src/emscripten/init.c b/src/emscripten/init.c +index c2393243755816fb341f892b7b092a58891603c4..6c8210f2c23ddc0f0efde954d80fcf974aba74fb 100644 +--- a/src/emscripten/init.c ++++ b/src/emscripten/init.c +@@ -286,3 +286,35 @@ cleanup: + free(l1d); + free(l2); + } ++ ++void cpuinfo_emscripten_deinit(void) { ++ free(cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ free(cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ free(cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ /* cpuinfo_packages points to static storage (&static_package) — do not free */ ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ /* L1i, L1d, L2 are dynamically allocated; L3 may point to static ++ * storage (&static_x86_l3) so only free levels that were calloc'd */ ++ free(cpuinfo_cache[cpuinfo_cache_level_1i]); ++ free(cpuinfo_cache[cpuinfo_cache_level_1d]); ++ free(cpuinfo_cache[cpuinfo_cache_level_2]); ++ /* cpuinfo_cache[cpuinfo_cache_level_3] may be &static_x86_l3 — do not free */ ++ for (int lvl = 0; lvl < cpuinfo_cache_level_max; ++lvl) { ++ cpuinfo_cache[lvl] = NULL; ++ cpuinfo_cache_count[lvl] = 0; ++ } ++ cpuinfo_max_cache_size = 0; ++ ++ cpuinfo_global_uarch = (struct cpuinfo_uarch_info){0}; ++} +diff --git a/src/init.c b/src/init.c +index 81d5721c446f363cefcbb39d52507a2c07298c35..e3cd9217a9acfc9615f97a1aa6bb66d96f373910 100644 +--- a/src/init.c ++++ b/src/init.c +@@ -13,55 +13,98 @@ + #endif + + #if defined(_WIN32) || defined(__CYGWIN__) +-static INIT_ONCE init_guard = INIT_ONCE_STATIC_INIT; ++static SRWLOCK cpuinfo_lifecycle_lock = SRWLOCK_INIT; ++#define CPUINFO_LIFECYCLE_LOCK() AcquireSRWLockExclusive(&cpuinfo_lifecycle_lock) ++#define CPUINFO_LIFECYCLE_UNLOCK() ReleaseSRWLockExclusive(&cpuinfo_lifecycle_lock) + #elif !defined(__EMSCRIPTEN__) || defined(__EMSCRIPTEN_PTHREADS__) +-static pthread_once_t init_guard = PTHREAD_ONCE_INIT; ++static pthread_mutex_t cpuinfo_lifecycle_lock = PTHREAD_MUTEX_INITIALIZER; ++#define CPUINFO_LIFECYCLE_LOCK() pthread_mutex_lock(&cpuinfo_lifecycle_lock) ++#define CPUINFO_LIFECYCLE_UNLOCK() pthread_mutex_unlock(&cpuinfo_lifecycle_lock) + #else +-static bool init_guard = false; ++/* Single-threaded build (e.g. Emscripten without pthreads): no lock needed. */ ++#define CPUINFO_LIFECYCLE_LOCK() \ ++ do { \ ++ } while (0) ++#define CPUINFO_LIFECYCLE_UNLOCK() \ ++ do { \ ++ } while (0) + #endif ++static unsigned int cpuinfo_init_refcount = 0; + + bool CPUINFO_ABI cpuinfo_initialize(void) { ++ CPUINFO_LIFECYCLE_LOCK(); ++ if (cpuinfo_init_refcount == 0) { + #if CPUINFO_ARCH_X86 || CPUINFO_ARCH_X86_64 + #if defined(__MACH__) && defined(__APPLE__) +- pthread_once(&init_guard, &cpuinfo_x86_mach_init); ++ cpuinfo_x86_mach_init(); + #elif defined(__FreeBSD__) +- pthread_once(&init_guard, &cpuinfo_x86_freebsd_init); ++ cpuinfo_x86_freebsd_init(); + #elif defined(__linux__) +- pthread_once(&init_guard, &cpuinfo_x86_linux_init); ++ cpuinfo_x86_linux_init(); + #elif defined(_WIN32) || defined(__CYGWIN__) +- InitOnceExecuteOnce(&init_guard, &cpuinfo_x86_windows_init, NULL, NULL); ++ cpuinfo_x86_windows_init(NULL, NULL, NULL); + #else +- cpuinfo_log_error("operating system is not supported in cpuinfo"); ++ cpuinfo_log_error("operating system is not supported in cpuinfo"); + #endif + #elif CPUINFO_ARCH_ARM || CPUINFO_ARCH_ARM64 + #if defined(__linux__) +- pthread_once(&init_guard, &cpuinfo_arm_linux_init); ++ cpuinfo_arm_linux_init(); + #elif defined(__MACH__) && defined(__APPLE__) +- pthread_once(&init_guard, &cpuinfo_arm_mach_init); ++ cpuinfo_arm_mach_init(); + #elif defined(_WIN32) +- InitOnceExecuteOnce(&init_guard, &cpuinfo_arm_windows_init, NULL, NULL); ++ cpuinfo_arm_windows_init(NULL, NULL, NULL); + #else +- cpuinfo_log_error("operating system is not supported in cpuinfo"); ++ cpuinfo_log_error("operating system is not supported in cpuinfo"); + #endif + #elif CPUINFO_ARCH_RISCV32 || CPUINFO_ARCH_RISCV64 + #if defined(__linux__) +- pthread_once(&init_guard, &cpuinfo_riscv_linux_init); ++ cpuinfo_riscv_linux_init(); + #else +- cpuinfo_log_error("operating system is not supported in cpuinfo"); ++ cpuinfo_log_error("operating system is not supported in cpuinfo"); + #endif + #elif CPUINFO_ARCH_ASMJS || CPUINFO_ARCH_WASM || CPUINFO_ARCH_WASMSIMD +-#if defined(__EMSCRIPTEN_PTHREADS__) +- pthread_once(&init_guard, &cpuinfo_emscripten_init); +-#else +- if (!init_guard) { + cpuinfo_emscripten_init(); +- } +- init_guard = true; +-#endif + #else +- cpuinfo_log_error("processor architecture is not supported in cpuinfo"); ++ cpuinfo_log_error("processor architecture is not supported in cpuinfo"); + #endif +- return cpuinfo_is_initialized; ++ } ++ if (cpuinfo_is_initialized) { ++ cpuinfo_init_refcount++; ++ } ++ const bool initialized = cpuinfo_is_initialized; ++ CPUINFO_LIFECYCLE_UNLOCK(); ++ return initialized; + } + +-void CPUINFO_ABI cpuinfo_deinitialize(void) {} ++void CPUINFO_ABI cpuinfo_deinitialize(void) { ++ CPUINFO_LIFECYCLE_LOCK(); ++ if (cpuinfo_init_refcount != 0 && --cpuinfo_init_refcount == 0) { ++ cpuinfo_is_initialized = false; ++#if CPUINFO_ARCH_X86 || CPUINFO_ARCH_X86_64 ++#if defined(__MACH__) && defined(__APPLE__) ++ cpuinfo_x86_mach_deinit(); ++#elif defined(__FreeBSD__) ++ cpuinfo_x86_freebsd_deinit(); ++#elif defined(__linux__) ++ cpuinfo_x86_linux_deinit(); ++#elif defined(_WIN32) || defined(__CYGWIN__) ++ cpuinfo_x86_windows_deinit(NULL, NULL, NULL); ++#endif ++#elif CPUINFO_ARCH_ARM || CPUINFO_ARCH_ARM64 ++#if defined(__linux__) ++ cpuinfo_arm_linux_deinit(); ++#elif defined(__MACH__) && defined(__APPLE__) ++ cpuinfo_arm_mach_deinit(); ++#elif defined(_WIN32) ++ cpuinfo_arm_windows_deinit(NULL, NULL, NULL); ++#endif ++#elif CPUINFO_ARCH_RISCV32 || CPUINFO_ARCH_RISCV64 ++#if defined(__linux__) ++ cpuinfo_riscv_linux_deinit(); ++#endif ++#elif CPUINFO_ARCH_ASMJS || CPUINFO_ARCH_WASM || CPUINFO_ARCH_WASMSIMD ++ cpuinfo_emscripten_deinit(); ++#endif ++ } ++ CPUINFO_LIFECYCLE_UNLOCK(); ++} +diff --git a/src/riscv/linux/init.c b/src/riscv/linux/init.c +index 45168c7e66b80e7f8a712a6afe7bab379a07b95b..3da8c3ffd59f02bebb0193b832d67ddf0841cab3 100644 +--- a/src/riscv/linux/init.c ++++ b/src/riscv/linux/init.c +@@ -618,3 +618,36 @@ cleanup: + free(linux_cpu_to_core_map); + free(linux_cpu_to_uarch_index_map); + } ++ ++void cpuinfo_riscv_linux_deinit(void) { ++ free(cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ free(cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ free(cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ free(cpuinfo_packages); ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ free(cpuinfo_uarchs); ++ cpuinfo_uarchs = NULL; ++ cpuinfo_uarchs_count = 0; ++ ++ free(cpuinfo_linux_cpu_to_processor_map); ++ cpuinfo_linux_cpu_to_processor_map = NULL; ++ ++ free(cpuinfo_linux_cpu_to_core_map); ++ cpuinfo_linux_cpu_to_core_map = NULL; ++ ++ free((void*)cpuinfo_linux_cpu_to_uarch_index_map); ++ cpuinfo_linux_cpu_to_uarch_index_map = NULL; ++ ++ cpuinfo_linux_cpu_max = 0; ++} +diff --git a/src/x86/freebsd/init.c b/src/x86/freebsd/init.c +index 797fa24b95a55e0503459dd2fdaf8476f060a631..ba30a5d9c89c209115f6797d9b52263c7271e17f 100644 +--- a/src/x86/freebsd/init.c ++++ b/src/x86/freebsd/init.c +@@ -396,3 +396,30 @@ cleanup: + free(l3); + free(l4); + } ++ ++void cpuinfo_x86_freebsd_deinit(void) { ++ free(cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ free(cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ free(cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ free(cpuinfo_packages); ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ for (int lvl = 0; lvl < cpuinfo_cache_level_max; ++lvl) { ++ free(cpuinfo_cache[lvl]); ++ cpuinfo_cache[lvl] = NULL; ++ cpuinfo_cache_count[lvl] = 0; ++ } ++ cpuinfo_max_cache_size = 0; ++ ++ cpuinfo_global_uarch = (struct cpuinfo_uarch_info){0}; ++} +diff --git a/src/x86/linux/init.c b/src/x86/linux/init.c +index d2b2d47514c8cc75b502c78b1d1176ae49fe3304..62e1187e0631b791c378d23963d63e10095a12be 100644 +--- a/src/x86/linux/init.c ++++ b/src/x86/linux/init.c +@@ -676,3 +676,38 @@ cleanup: + free(linux_cpu_to_processor_map); + free(linux_cpu_to_core_map); + } ++ ++void cpuinfo_x86_linux_deinit(void) { ++ free(cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ free(cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ free(cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ free(cpuinfo_packages); ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ for (int lvl = 0; lvl < cpuinfo_cache_level_max; ++lvl) { ++ free(cpuinfo_cache[lvl]); ++ cpuinfo_cache[lvl] = NULL; ++ cpuinfo_cache_count[lvl] = 0; ++ } ++ cpuinfo_max_cache_size = 0; ++ ++ free(cpuinfo_linux_cpu_to_processor_map); ++ cpuinfo_linux_cpu_to_processor_map = NULL; ++ ++ free(cpuinfo_linux_cpu_to_core_map); ++ cpuinfo_linux_cpu_to_core_map = NULL; ++ ++ cpuinfo_linux_cpu_max = 0; ++ ++ cpuinfo_global_uarch = (struct cpuinfo_uarch_info){0}; ++} +diff --git a/src/x86/mach/init.c b/src/x86/mach/init.c +index b8ea6047ac043fc067190013e719188b50aca3f5..9009e830ab4ac15610af72e700266b800d88e0fc 100644 +--- a/src/x86/mach/init.c ++++ b/src/x86/mach/init.c +@@ -378,3 +378,30 @@ cleanup: + free(l3); + free(l4); + } ++ ++void cpuinfo_x86_mach_deinit(void) { ++ free(cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ free(cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ free(cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ free(cpuinfo_packages); ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ for (int lvl = 0; lvl < cpuinfo_cache_level_max; ++lvl) { ++ free(cpuinfo_cache[lvl]); ++ cpuinfo_cache[lvl] = NULL; ++ cpuinfo_cache_count[lvl] = 0; ++ } ++ cpuinfo_max_cache_size = 0; ++ ++ cpuinfo_global_uarch = (struct cpuinfo_uarch_info){0}; ++} +diff --git a/src/x86/windows/init.c b/src/x86/windows/init.c +index 7380bb46d5f2ee96e98b77fd378be9d70de813ee..f1401beb60de989503402be98fdbac4c20909c1b 100644 +--- a/src/x86/windows/init.c ++++ b/src/x86/windows/init.c +@@ -665,5 +665,39 @@ cleanup: + if (l4 != NULL) { + HeapFree(heap, 0, l4); + } ++ if (processor_infos != NULL) { ++ HeapFree(heap, 0, processor_infos); ++ } ++ return TRUE; ++} ++ ++BOOL CALLBACK cpuinfo_x86_windows_deinit(PINIT_ONCE init_once, PVOID parameter, PVOID* context) { ++ HANDLE heap = GetProcessHeap(); ++ ++ HeapFree(heap, 0, cpuinfo_processors); ++ cpuinfo_processors = NULL; ++ cpuinfo_processors_count = 0; ++ ++ HeapFree(heap, 0, cpuinfo_cores); ++ cpuinfo_cores = NULL; ++ cpuinfo_cores_count = 0; ++ ++ HeapFree(heap, 0, cpuinfo_clusters); ++ cpuinfo_clusters = NULL; ++ cpuinfo_clusters_count = 0; ++ ++ HeapFree(heap, 0, cpuinfo_packages); ++ cpuinfo_packages = NULL; ++ cpuinfo_packages_count = 0; ++ ++ for (int lvl = 0; lvl < cpuinfo_cache_level_max; ++lvl) { ++ HeapFree(heap, 0, cpuinfo_cache[lvl]); ++ cpuinfo_cache[lvl] = NULL; ++ cpuinfo_cache_count[lvl] = 0; ++ } ++ cpuinfo_max_cache_size = 0; ++ ++ cpuinfo_global_uarch = (struct cpuinfo_uarch_info){0}; ++ + return TRUE; + } diff --git a/cmake/patches/xnnpack/release_cpuinfo_after_hardware_config.patch b/cmake/patches/xnnpack/release_cpuinfo_after_hardware_config.patch new file mode 100644 index 0000000000000..8aea6bfc4d156 --- /dev/null +++ b/cmake/patches/xnnpack/release_cpuinfo_after_hardware_config.patch @@ -0,0 +1,46 @@ +diff --git a/src/configs/hardware-config.c b/src/configs/hardware-config.c +index b4e942aa6955360e3ea85619d549410e65ab7a9c..6cf6e9a213b6f72c94940ff815e11f70a658c419 100644 +--- a/src/configs/hardware-config.c ++++ b/src/configs/hardware-config.c +@@ -343,6 +343,7 @@ static void init_hardware_config(void) { + xnn_log_info("cpu_get_uarch(%i): 0x%x", i, cpuinfo_get_uarch(i)->uarch); + } + #endif // XNN_MAX_UARCH_TYPES > 1 ++ cpuinfo_deinitialize(); + #else + xnn_log_warning("Unable to determine L1/L2 data cache properties."); + #endif // XNN_ENABLE_CPUINFO +@@ -358,21 +359,33 @@ const struct xnn_hardware_config* xnn_init_hardware_config() { + #if XNN_ARCH_ARM + if (!cpuinfo_has_arm_v6()) { + xnn_log_debug("unsupported hardware: ARMv6 not detected"); ++ #if !XNN_PLATFORM_WEB && XNN_ENABLE_CPUINFO ++ cpuinfo_deinitialize(); ++ #endif + return NULL; + } + + if (!cpuinfo_has_arm_vfpv2() && !cpuinfo_has_arm_vfpv3()) { + xnn_log_debug("unsupported hardware: VFP FPU not detected"); ++ #if !XNN_PLATFORM_WEB && XNN_ENABLE_CPUINFO ++ cpuinfo_deinitialize(); ++ #endif + return NULL; + } + #endif // XNN_ARCH_ARM + #if XNN_ARCH_X86 + if (!cpuinfo_has_x86_sse2()) { + xnn_log_debug("unsupported hardware: SSE2 not detected"); ++ #if !XNN_PLATFORM_WEB && XNN_ENABLE_CPUINFO ++ cpuinfo_deinitialize(); ++ #endif + return NULL; + } + #endif // XNN_ARCH_X86 + + XNN_INIT_ONCE(hardware); ++ #if !XNN_PLATFORM_WEB && !XNN_ARCH_RISCV && !XNN_ARCH_PPC64 && XNN_ENABLE_CPUINFO ++ cpuinfo_deinitialize(); ++ #endif + return &hardware_config; + } diff --git a/cmake/vcpkg-ports/cpuinfo/portfile.cmake b/cmake/vcpkg-ports/cpuinfo/portfile.cmake index 9140a233e2ccd..aa2205b699e07 100644 --- a/cmake/vcpkg-ports/cpuinfo/portfile.cmake +++ b/cmake/vcpkg-ports/cpuinfo/portfile.cmake @@ -6,12 +6,14 @@ endif() vcpkg_from_github( OUT_SOURCE_PATH SOURCE_PATH REPO pytorch/cpuinfo - REF 4628dc060ce4e82345dc166bbac875609db4ff69 - SHA512 db7a93279f2f6daaf825fbd8552935d8ed671d276b65ad614e11f722b6a6848e663850d65180d33b554d67ef1a36aae842feb368699f90be8f21172a1af1924e + REF 66ee79c038d70dad9f08705b2c9b3e58f6d8f512 + SHA512 ec1df3a03e52f50f221fc5f25251b7181316b59b32be61f11dec440ef245bc2070d0e16579ae8fd2748711f120c914f4479619ce1bd6b0bc9dbe462a8526298f HEAD_REF main PATCHES patch_cpuinfo_h_for_arm64ec.patch patch_vcpkg_arm64ec_support.patch # https://github.com/pytorch/cpuinfo/pull/324 + # https://github.com/pytorch/cpuinfo/pull/400 + "${CMAKE_CURRENT_LIST_DIR}/../../patches/cpuinfo/enable_deinit_refcounting.patch" ) vcpkg_check_features(OUT_FEATURE_OPTIONS FEATURE_OPTIONS diff --git a/cmake/vcpkg-ports/xnnpack/portfile.cmake b/cmake/vcpkg-ports/xnnpack/portfile.cmake index 60b3566629e10..6ccfd2909454f 100644 --- a/cmake/vcpkg-ports/xnnpack/portfile.cmake +++ b/cmake/vcpkg-ports/xnnpack/portfile.cmake @@ -11,6 +11,7 @@ vcpkg_from_github( PATCHES fix-build.patch disable_gcc_warning.patch + "${CMAKE_CURRENT_LIST_DIR}/../../patches/xnnpack/release_cpuinfo_after_hardware_config.patch" ) vcpkg_find_acquire_program(PYTHON3) diff --git a/onnxruntime/test/shared_lib/cpuinfo_dlopen_test.cc b/onnxruntime/test/shared_lib/cpuinfo_dlopen_test.cc new file mode 100644 index 0000000000000..e2f97769f0b32 --- /dev/null +++ b/onnxruntime/test/shared_lib/cpuinfo_dlopen_test.cc @@ -0,0 +1,48 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include + +#include +#include + +#ifndef ORT_CPUINFO_DLOPEN_TEST_LIBRARY +#error ORT_CPUINFO_DLOPEN_TEST_LIBRARY must name the test DLL. +#endif + +int wmain() { + HMODULE library = LoadLibraryW(ORT_CPUINFO_DLOPEN_TEST_LIBRARY); + if (library == nullptr) { + std::cerr << "LoadLibraryW failed with error " << GetLastError() << std::endl; + return EXIT_FAILURE; + } + + using GetCpuinfoAllocation = const void* (*)(); + const auto get_cpuinfo_allocation = reinterpret_cast( + GetProcAddress(library, "OrtGetCpuinfoAllocationForTesting")); + if (get_cpuinfo_allocation == nullptr) { + std::cerr << "GetProcAddress failed with error " << GetLastError() << std::endl; + FreeLibrary(library); + return EXIT_FAILURE; + } + + const void* allocation = get_cpuinfo_allocation(); + HANDLE process_heap = GetProcessHeap(); + if (allocation == nullptr || !HeapValidate(process_heap, 0, allocation)) { + std::cerr << "cpuinfo did not return a valid process-heap allocation" << std::endl; + FreeLibrary(library); + return EXIT_FAILURE; + } + + if (!FreeLibrary(library)) { + std::cerr << "FreeLibrary failed with error " << GetLastError() << std::endl; + return EXIT_FAILURE; + } + + if (HeapValidate(process_heap, 0, allocation)) { + std::cerr << "cpuinfo allocation remained valid after the CPUIDInfo test DLL was unloaded" << std::endl; + return EXIT_FAILURE; + } + + return EXIT_SUCCESS; +} diff --git a/onnxruntime/test/shared_lib/cpuinfo_dlopen_test_library.cc b/onnxruntime/test/shared_lib/cpuinfo_dlopen_test_library.cc new file mode 100644 index 0000000000000..edad81c76388d --- /dev/null +++ b/onnxruntime/test/shared_lib/cpuinfo_dlopen_test_library.cc @@ -0,0 +1,11 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include + +#include "core/common/cpuid_info.h" + +extern "C" const void* OrtGetCpuinfoAllocationForTesting() { + static_cast(onnxruntime::CPUIDInfo::GetCPUIDInfo()); + return cpuinfo_get_processors(); +} diff --git a/onnxruntime/test/shared_lib/cpuinfo_dlopen_test_library.def b/onnxruntime/test/shared_lib/cpuinfo_dlopen_test_library.def new file mode 100644 index 0000000000000..ea4221dc1e5d3 --- /dev/null +++ b/onnxruntime/test/shared_lib/cpuinfo_dlopen_test_library.def @@ -0,0 +1,3 @@ +LIBRARY onnxruntime_cpuinfo_dlopen_test_library +EXPORTS + OrtGetCpuinfoAllocationForTesting diff --git a/onnxruntime/test/shared_lib/cpuinfo_refcount_test.cc b/onnxruntime/test/shared_lib/cpuinfo_refcount_test.cc new file mode 100644 index 0000000000000..7c91ce9f88afd --- /dev/null +++ b/onnxruntime/test/shared_lib/cpuinfo_refcount_test.cc @@ -0,0 +1,223 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include + +#include +#include +#include +#include +#include + +#if defined(ORT_CPUINFO_TEST_USE_XNNPACK) +#include +#endif + +#if defined(ORT_CPUINFO_TEST_HAS_INTERNAL_STATE) +constexpr int kCpuinfoCacheLevelCount = 5; + +extern "C" { +extern bool cpuinfo_is_initialized; +extern struct cpuinfo_processor* cpuinfo_processors; +extern struct cpuinfo_core* cpuinfo_cores; +extern struct cpuinfo_cluster* cpuinfo_clusters; +extern struct cpuinfo_package* cpuinfo_packages; +extern struct cpuinfo_cache* cpuinfo_cache[kCpuinfoCacheLevelCount]; +extern uint32_t cpuinfo_processors_count; +extern uint32_t cpuinfo_cores_count; +extern uint32_t cpuinfo_clusters_count; +extern uint32_t cpuinfo_packages_count; +extern uint32_t cpuinfo_cache_count[kCpuinfoCacheLevelCount]; +extern uint32_t cpuinfo_max_cache_size; +} +#endif + +namespace { + +bool HasValidCpuinfoState() { + return cpuinfo_get_processors_count() != 0 && + cpuinfo_get_processors() != nullptr && + cpuinfo_get_processor(0) != nullptr; +} + +bool IsCpuinfoDeinitialized() { +#if defined(ORT_CPUINFO_TEST_HAS_INTERNAL_STATE) + if (cpuinfo_is_initialized || + cpuinfo_processors != nullptr || + cpuinfo_cores != nullptr || + cpuinfo_clusters != nullptr || + cpuinfo_packages != nullptr || + cpuinfo_processors_count != 0 || + cpuinfo_cores_count != 0 || + cpuinfo_clusters_count != 0 || + cpuinfo_packages_count != 0 || + cpuinfo_max_cache_size != 0) { + return false; + } + + for (int level = 0; level < kCpuinfoCacheLevelCount; ++level) { + if (cpuinfo_cache[level] != nullptr || cpuinfo_cache_count[level] != 0) { + return false; + } + } + + return true; +#else + return true; +#endif +} + +bool TestSequentialConsumers() { + if (!cpuinfo_initialize()) { + std::cerr << "The first cpuinfo_initialize call failed" << std::endl; + return false; + } + + if (!cpuinfo_initialize()) { + std::cerr << "The second cpuinfo_initialize call failed" << std::endl; + cpuinfo_deinitialize(); + return false; + } + + const uint32_t processor_count = cpuinfo_get_processors_count(); + const cpuinfo_processor* processors = cpuinfo_get_processors(); + if (processor_count == 0 || processors == nullptr) { + std::cerr << "cpuinfo did not expose valid processor data" << std::endl; + cpuinfo_deinitialize(); + cpuinfo_deinitialize(); + return false; + } + + cpuinfo_deinitialize(); + + const bool first_consumer_release_preserved_state = + cpuinfo_get_processors_count() == processor_count && + cpuinfo_get_processors() == processors; + + cpuinfo_deinitialize(); + + if (!IsCpuinfoDeinitialized()) { + std::cerr << "cpuinfo remained initialized after the final consumer released it" << std::endl; + return false; + } + + if (!first_consumer_release_preserved_state) { + std::cerr << "cpuinfo released shared state while another consumer was still active" << std::endl; + return false; + } + + if (!cpuinfo_initialize()) { + std::cerr << "cpuinfo failed to reinitialize after the final consumer released it" << std::endl; + return false; + } + cpuinfo_deinitialize(); + + if (!IsCpuinfoDeinitialized()) { + std::cerr << "cpuinfo remained initialized after the reinitialized consumer released it" << std::endl; + return false; + } + + return true; +} + +bool TestConcurrentConsumers() { + constexpr size_t kThreadCount = 8; + constexpr size_t kIterations = 100; + std::atomic initialized_count{0}; + std::atomic validated_count{0}; + std::atomic deinitialized_count{0}; + std::atomic failed{false}; + + const auto consumer = [&]() { + for (size_t iteration = 0; iteration < kIterations; ++iteration) { + const bool initialized = cpuinfo_initialize(); + if (!initialized) { + failed = true; + } + + const size_t expected_count = (iteration + 1) * kThreadCount; + ++initialized_count; + while (initialized_count < expected_count) { + std::this_thread::yield(); + } + + if (initialized && !HasValidCpuinfoState()) { + failed = true; + } + + ++validated_count; + while (validated_count < expected_count) { + std::this_thread::yield(); + } + + if (initialized) { + cpuinfo_deinitialize(); + } + + ++deinitialized_count; + while (deinitialized_count < expected_count) { + std::this_thread::yield(); + } + } + }; + + std::vector consumers; + consumers.reserve(kThreadCount); + for (size_t thread = 0; thread < kThreadCount; ++thread) { + consumers.emplace_back(consumer); + } + for (std::thread& consumer_thread : consumers) { + consumer_thread.join(); + } + + if (failed) { + std::cerr << "cpuinfo failed during concurrent initialization and deinitialization" << std::endl; + return false; + } + + if (!cpuinfo_initialize() || !HasValidCpuinfoState()) { + std::cerr << "cpuinfo failed to reinitialize after concurrent use" << std::endl; + return false; + } + cpuinfo_deinitialize(); + + if (!IsCpuinfoDeinitialized()) { + std::cerr << "cpuinfo remained initialized after concurrent use" << std::endl; + return false; + } + + return true; +} + +#if defined(ORT_CPUINFO_TEST_USE_XNNPACK) +bool TestXnnpackReleasesCpuinfo() { + if (xnn_initialize(nullptr) != xnn_status_success) { + std::cerr << "XNNPACK initialization failed" << std::endl; + return false; + } + + if (!IsCpuinfoDeinitialized()) { + std::cerr << "XNNPACK retained a cpuinfo reference after hardware discovery" << std::endl; + return false; + } + + xnn_deinitialize(); + return true; +} +#endif + +} // namespace + +int main() { + if (!TestSequentialConsumers() || !TestConcurrentConsumers()) { + return EXIT_FAILURE; + } + +#if defined(ORT_CPUINFO_TEST_USE_XNNPACK) + if (!TestXnnpackReleasesCpuinfo()) { + return EXIT_FAILURE; + } +#endif + + return EXIT_SUCCESS; +} From 89e9f623cb2f9922e39a6cfa4c7845917d72e142 Mon Sep 17 00:00:00 2001 From: Ananya Anand <148831356+4n4ny4@users.noreply.github.com> Date: Wed, 2 Sep 2026 14:15:13 -0700 Subject: [PATCH 02/13] [WebGPU] Select the pooling path by occupancy rather than output size (#32313) ### Description Pool picks between a serial path, where one invocation loops over the whole kernel window, and a workgroup-cooperative path that splits the window across invocations and reduces in shared memory. The choice was made on output size, which does not describe how much of the GPU either path will fill. This selects on occupancy instead: the serial path is only competitive when the output is large enough to keep the device busy on its own, and the cooperative path is what recovers a large kernel window over a small output. ### Motivation and Context EfficientNet-B0 has a global average pool that reduces the entire feature map to one pixel per channel, which is the shape the old heuristic handled worst: a small output that the serial path cannot parallelise over, with a long reduction behind it. Measured on an NVIDIA TITAN V, driver 560.94: - EfficientNet-B0 native end to end: 3.73 ms -> 1.78 ms - Pool as a share of GPU time: 54.6% -> 11.1% Numerically identical output; this only changes how the reduction is scheduled. --------- Co-authored-by: Ananya Anand --- onnxruntime/core/providers/webgpu/nn/pool.cc | 39 ++++++++---- onnxruntime/core/providers/webgpu/nn/pool.h | 6 +- .../test/providers/cpu/nn/pool_op_test.cc | 63 +++++++++++++++++++ 3 files changed, 92 insertions(+), 16 deletions(-) diff --git a/onnxruntime/core/providers/webgpu/nn/pool.cc b/onnxruntime/core/providers/webgpu/nn/pool.cc index 07213805380c0..e1e09681d575b 100644 --- a/onnxruntime/core/providers/webgpu/nn/pool.cc +++ b/onnxruntime/core/providers/webgpu/nn/pool.cc @@ -95,7 +95,7 @@ Status PoolProgram::GenerateShaderCode(ShaderHelper& shader) const { var_decl_code = SS_GET(var_decl_ss); sampling_code = " value = max(value, x_val);\n"; - if (are_small_output_big_kernel_) { + if (use_parallel_reduction_) { downsampling_code = " sum_or_max_shared[local_idx] = value;\n"; } } else { @@ -118,7 +118,7 @@ Status PoolProgram::GenerateShaderCode(ShaderHelper& shader) const { : " if (!is_pad) {\n count++;\n }\n"; SS(downsampling_ss, kStringInitialSize); - if (are_small_output_big_kernel_) { + if (use_parallel_reduction_) { downsampling_ss << " sum_or_max_shared[local_idx] = value;\n" << " count_shared[local_idx] = count;\n"; } else { @@ -165,7 +165,7 @@ Status PoolProgram::GenerateShaderCode(ShaderHelper& shader) const { std::string pad_break_code = overflow_check_code.empty() ? " break;\n" : ""; std::string sum_or_max_shared; - if (are_small_output_big_kernel_) { + if (use_parallel_reduction_) { shader.AdditionalImplementation() << "var sum_or_max_shared : array<" << (is_float16_ ? "f16" : "f32") << ",workgroup_size_x >;\n" << (!is_max_pool_ ? "var count_shared : array;\n" : ""); @@ -192,10 +192,10 @@ Status PoolProgram::GenerateShaderCode(ShaderHelper& shader) const { << " }\n"; sum_or_max_shared = SS_GET(shared_ss); } - std::string kernel_loop_decl_code = are_small_output_big_kernel_ ? " for (var i: u32 = local_idx; i < uniforms.kernel_size; i += workgroup_size_x) {\n" : " for (var i: u32 = 0; i < uniforms.kernel_size; i++) {\n"; + std::string kernel_loop_decl_code = use_parallel_reduction_ ? " for (var i: u32 = local_idx; i < uniforms.kernel_size; i += workgroup_size_x) {\n" : " for (var i: u32 = 0; i < uniforms.kernel_size; i++) {\n"; SS(output_ss, kStringInitialSize); - if (are_small_output_big_kernel_) { + if (use_parallel_reduction_) { output_ss << " if (local_idx == 0) {\n" << " value = sum_or_max_shared[0];\n"; if (!is_max_pool_) { @@ -212,8 +212,8 @@ Status PoolProgram::GenerateShaderCode(ShaderHelper& shader) const { std::string output_code = SS_GET(output_ss); auto& body = shader.MainFunctionBody(); - body << (are_small_output_big_kernel_ ? "" : shader.GuardAgainstOutOfBoundsWorkgroupSizes("uniforms.output_size")) - << " let y_indices = " << output.OffsetToIndices((are_small_output_big_kernel_ ? "workgroup_idx" : "global_idx")) << ";\n" + body << (use_parallel_reduction_ ? "" : shader.GuardAgainstOutOfBoundsWorkgroupSizes("uniforms.output_size")) + << " let y_indices = " << output.OffsetToIndices((use_parallel_reduction_ ? "workgroup_idx" : "global_idx")) << ";\n" << " var x_indices = y_indices;\n" << " var k_indices: array;\n" << var_decl_code @@ -260,6 +260,13 @@ Status PoolProgram::GenerateShaderCode(ShaderHelper& shader) const { return Status::OK(); } +// Threads a parallel-reduction workgroup uses to co-operate on one output element's kernel window. +constexpr uint32_t kParallelReductionWorkgroupSize = 128; + +// Below this many workgroups a dispatch leaves compute units idle. WebGPU reports no compute-unit +// count, so this is a fixed value chosen to sit under discrete GPUs and over integrated ones. +constexpr uint32_t kMinDispatchGroupsToFillDevice = 64; + template Status Pool::ComputeInternal(ComputeContext& context) const { // TODO: support 'column major' storage_order. @@ -327,10 +334,16 @@ Status Pool::ComputeInternal(ComputeContext& context) const { const auto strides_u32 = NarrowToU32(strides); const auto dilations_u32 = NarrowToU32(dilations); - bool are_small_output_big_kernel = output_size <= 128 && kernel_size >= 128; - PoolProgram program{is_max_pool, is_nhwc, kernel_shape, is_float16, count_include_pad, are_small_output_big_kernel}; + const uint32_t serial_dispatch_groups = static_cast((output_size + WORKGROUP_SIZE - 1) / WORKGROUP_SIZE); + // The serial path gives each output element one thread that loops the whole window, so its + // parallelism is output_size alone. The parallel path spends a workgroup per output element and + // reduces the window across its threads. Select on whether the serial path fills the device, not + // on output size: a small output with a large window holds plenty of work either way. + const bool use_parallel_reduction = + serial_dispatch_groups < kMinDispatchGroupsToFillDevice && kernel_size >= kParallelReductionWorkgroupSize; + PoolProgram program{is_max_pool, is_nhwc, kernel_shape, is_float16, count_include_pad, use_parallel_reduction}; - program.CacheHint(kernel_shape.size(), is_max_pool, is_nhwc, is_float16, count_include_pad, are_small_output_big_kernel) + program.CacheHint(kernel_shape.size(), is_max_pool, is_nhwc, is_float16, count_include_pad, use_parallel_reduction) .AddInputs({{X, ProgramTensorMetadataDependency::TypeAndRank}}) .AddOutputs({{Y}}) .AddUniformVariables({output_size, kernel_size, @@ -339,11 +352,11 @@ Status Pool::ComputeInternal(ComputeContext& context) const { gsl::span(strides_u32.data(), strides_u32.size()), gsl::span(dilations_u32.data(), dilations_u32.size())}); - if (are_small_output_big_kernel) { - program.SetWorkgroupSize(128) + if (use_parallel_reduction) { + program.SetWorkgroupSize(kParallelReductionWorkgroupSize) .SetDispatchGroupSize(output_size); } else { - program.SetDispatchGroupSize((output_size + WORKGROUP_SIZE - 1) / WORKGROUP_SIZE); + program.SetDispatchGroupSize(serial_dispatch_groups); } return context.RunProgram(program); diff --git a/onnxruntime/core/providers/webgpu/nn/pool.h b/onnxruntime/core/providers/webgpu/nn/pool.h index 57bdc64954acd..fbcabf0615256 100644 --- a/onnxruntime/core/providers/webgpu/nn/pool.h +++ b/onnxruntime/core/providers/webgpu/nn/pool.h @@ -14,14 +14,14 @@ namespace webgpu { class PoolProgram final : public Program { public: PoolProgram(bool is_max_pool, bool is_nhwc, const TensorShapeVector& kernel_shape, bool is_float16, - bool count_include_pad, bool are_small_output_big_kernel) + bool count_include_pad, bool use_parallel_reduction) : Program{"Pool"}, is_max_pool_{is_max_pool}, is_nhwc_{is_nhwc}, kernel_shape_{kernel_shape}, is_float16_{is_float16}, count_include_pad_{count_include_pad}, - are_small_output_big_kernel_{are_small_output_big_kernel} {} + use_parallel_reduction_{use_parallel_reduction} {} Status GenerateShaderCode(ShaderHelper& sh) const override; @@ -40,7 +40,7 @@ class PoolProgram final : public Program { const TensorShapeVector kernel_shape_; const bool is_float16_; const bool count_include_pad_; - const bool are_small_output_big_kernel_; + const bool use_parallel_reduction_; }; template diff --git a/onnxruntime/test/providers/cpu/nn/pool_op_test.cc b/onnxruntime/test/providers/cpu/nn/pool_op_test.cc index f89cf8c19c161..2288d5fbdb1b5 100644 --- a/onnxruntime/test/providers/cpu/nn/pool_op_test.cc +++ b/onnxruntime/test/providers/cpu/nn/pool_op_test.cc @@ -1847,6 +1847,69 @@ TEST(PoolTest, GlobalAveragePool_Large_256) { test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}); } +// The WebGPU EP runs a pool as a workgroup-cooperative reduction when the serial path would not +// fill the device and the kernel window is large. Global pooling only reaches that path with a +// single output element per channel, which leaves two things unexercised: an output index taken +// from workgroup_idx, and a divisor that has to be reduced across the workgroup. +static void RunWebGpuLargeKernelPoolTest(const char* op_type, int opset, bool is_max_pool) { + auto webgpu_ep = DefaultWebGpuExecutionProvider(); + if (webgpu_ep == nullptr) { + GTEST_SKIP() << "WebGPU EP is not available in this build."; + } + + // A 12x12 window with pad 2 over a 24x24 input: kernel_size 144, output_size 578. + constexpr int64_t kChannels = 2, kSpatial = 24, kKernel = 12, kPad = 2; + constexpr int64_t kOutSpatial = kSpatial + 2 * kPad - kKernel + 1; + + std::vector x_vals(kChannels * kSpatial * kSpatial); + for (size_t i = 0; i < x_vals.size(); ++i) { + x_vals[i] = static_cast(i) * 0.01f; + } + + std::vector expected_vals(kChannels * kOutSpatial * kOutSpatial); + for (int64_t c = 0; c < kChannels; ++c) { + for (int64_t oh = 0; oh < kOutSpatial; ++oh) { + for (int64_t ow = 0; ow < kOutSpatial; ++ow) { + float acc = is_max_pool ? std::numeric_limits::lowest() : 0.0f; + int64_t count = 0; + for (int64_t kh = 0; kh < kKernel; ++kh) { + for (int64_t kw = 0; kw < kKernel; ++kw) { + const int64_t ih = oh + kh - kPad; + const int64_t iw = ow + kw - kPad; + if (ih < 0 || ih >= kSpatial || iw < 0 || iw >= kSpatial) { + continue; + } + const float v = x_vals[static_cast((c * kSpatial + ih) * kSpatial + iw)]; + acc = is_max_pool ? (v > acc ? v : acc) : acc + v; + ++count; + } + } + expected_vals[static_cast((c * kOutSpatial + oh) * kOutSpatial + ow)] = + is_max_pool ? acc : acc / static_cast(count); + } + } + } + + OpTester test(op_type, opset); + test.AddAttribute("kernel_shape", std::vector{kKernel, kKernel}); + test.AddAttribute("pads", std::vector{kPad, kPad, kPad, kPad}); + test.AddInput("X", {1, kChannels, kSpatial, kSpatial}, x_vals); + test.AddOutput("Y", {1, kChannels, kOutSpatial, kOutSpatial}, expected_vals, + /*sort_output=*/false, /*rel_error=*/1e-3f, /*abs_error=*/1e-2f); + + std::vector> execution_providers; + execution_providers.push_back(std::move(webgpu_ep)); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +} + +TEST(PoolTest, AveragePool_LargeKernelMultiElementOutput_WebGpu) { + RunWebGpuLargeKernelPoolTest("AveragePool", 11, /*is_max_pool=*/false); +} + +TEST(PoolTest, MaxPool_LargeKernelMultiElementOutput_WebGpu) { + RunWebGpuLargeKernelPoolTest("MaxPool", 12, /*is_max_pool=*/true); +} + TEST(PoolTest, LpPool) { OpTester test("LpPool"); From c4231eecc4508248369f1f742a7614c914886fcb Mon Sep 17 00:00:00 2001 From: Ananya Anand <148831356+4n4ny4@users.noreply.github.com> Date: Wed, 2 Sep 2026 14:21:07 -0700 Subject: [PATCH 03/13] [WebGPU] Report copy_tensors misuse instead of terminating the process (#32315) ### Description `BufferManager::MemCpy` reports misuse with `ORT_ENFORCE`, which throws: an aliased src/dst pair, an undersized destination and a still-mapped buffer all throw rather than returning a Status. The shared data transfer reaches it through `WebGpuDataTransferImpl::CopyTensorsImpl`, which is a `noexcept` C ABI callback and only handled a returned non-OK Status. The exception escaped a `noexcept` frame, so `std::terminate` took the process down with no message and no stack. This wraps the callback body so any throw becomes an `OrtStatus`, which is what the surrounding C API already expects. The two sibling callbacks are deliberately left alone: `CanCopyImpl` only calls C ABI function pointers, and `ReleaseImpl` returns void, so converting a throw there would mean swallowing it rather than reporting it, which is a separate decision from this one. ### Motivation and Context Not currently reachable from Python, because `copy_tensors` rejects every WebGPU OrtValue outright. It becomes reachable as soon as the supported copies are allowed, which is what #32074 does, so it is worth fixing at the boundary on its own rather than landing inside a larger change. The regression test for it uses the session-scoped allocation APIs under discussion in #32074 and is held there rather than duplicated here. Co-authored-by: Ananya Anand <4n4ny4@users.noreply.github.com> --- .../webgpu/webgpu_provider_factory.cc | 16 ++++++- .../test/autoep/test_webgpu_allocators.cc | 43 +++++++++++++++++++ 2 files changed, 58 insertions(+), 1 deletion(-) diff --git a/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc b/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc index 801307f185bd4..c2f89680eecf1 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc +++ b/onnxruntime/core/providers/webgpu/webgpu_provider_factory.cc @@ -412,12 +412,26 @@ struct WebGpuDataTransferImpl : OrtDataTransferImpl { (src_type == OrtMemoryInfoDeviceType_CPU && dst_type == OrtMemoryInfoDeviceType_GPU); } + // BufferManager::MemCpy reports misuse by throwing (ORT_ENFORCE), not by returning a Status. + // This callback crosses the C ABI and is noexcept, so an escaping exception would terminate the + // process. Convert throws to an OrtStatus. static OrtStatus* ORT_API_CALL CopyTensorsImpl( OrtDataTransferImpl* this_ptr, const OrtValue** src_tensors, OrtValue** dst_tensors, - OrtSyncStream** /*streams*/, + OrtSyncStream** streams, size_t num_tensors) noexcept { + API_IMPL_BEGIN + return CopyTensorsOrThrow(this_ptr, src_tensors, dst_tensors, streams, num_tensors); + API_IMPL_END + } + + static OrtStatus* CopyTensorsOrThrow( + OrtDataTransferImpl* this_ptr, + const OrtValue** src_tensors, + OrtValue** dst_tensors, + OrtSyncStream** /*streams*/, + size_t num_tensors) { auto& impl = *static_cast(this_ptr); if (num_tensors == 0) { diff --git a/onnxruntime/test/autoep/test_webgpu_allocators.cc b/onnxruntime/test/autoep/test_webgpu_allocators.cc index f7f77df598d87..13d6b551ca270 100644 --- a/onnxruntime/test/autoep/test_webgpu_allocators.cc +++ b/onnxruntime/test/autoep/test_webgpu_allocators.cc @@ -316,6 +316,49 @@ TEST_F(WebGpuPluginSharedAllocatorTest, DeviceTensorDataRoundTripsWithSharedAndS } } +// BufferManager::MemCpy rejects a self-copy with ORT_ENFORCE, which throws rather than returning a +// Status. CopyTensorsImpl is a noexcept C ABI callback, so before the throw was converted to an +// OrtStatus this ended the process instead of reporting the misuse. +TEST_F(WebGpuPluginSharedAllocatorTest, DeviceTensorSelfCopyIsReportedInsteadOfTerminating) { + constexpr std::array shape{8}; + std::array input_data{1.0f, -2.0f, 3.5f, 4.0f, -5.25f, 6.0f, 7.75f, -8.0f}; + auto cpu_memory_info = Ort::MemoryInfo::CreateCpu(OrtArenaAllocator, OrtMemTypeDefault); + + for (const auto& ep_device : EpDevices()) { + auto device_memory_info = ep_device.GetMemoryInfo(OrtDeviceMemoryType_DEFAULT); + ASSERT_NE(device_memory_info, nullptr); + SCOPED_TRACE(::testing::Message() << "device " << device_memory_info.GetDeviceId()); + + Ort::KeyValuePairs allocator_options; + auto shared_allocator = Env().CreateSharedAllocator( + ep_device, OrtDeviceMemoryType_DEFAULT, OrtDeviceAllocator, allocator_options); + ASSERT_NE(shared_allocator, nullptr); + + Ort::SessionOptions session_options; + Ort::KeyValuePairs ep_options; + session_options.AppendExecutionProvider_V2(Env(), {ep_device}, ep_options); + Ort::Session session(Env(), ORT_TSTR("testdata/mul_1.onnx"), session_options); + + auto device_tensor = Ort::Value::CreateTensor(shared_allocator, shape.data(), shape.size()); + + // Source and destination resolve to the same WGPUBuffer, which MemCpy refuses. + Ort::Status self_copy = Env().CopyTensor(device_tensor, device_tensor, nullptr); + ASSERT_FALSE(self_copy.IsOK()); + EXPECT_THAT(self_copy.GetErrorMessage(), ::testing::HasSubstr("must be different")); + + // The misuse was reported, not fatal: the same tensor still round trips. + std::array output_data{}; + auto cpu_input = Ort::Value::CreateTensor( + cpu_memory_info, input_data.data(), input_data.size(), shape.data(), shape.size()); + auto cpu_output = Ort::Value::CreateTensor( + cpu_memory_info, output_data.data(), output_data.size(), shape.data(), shape.size()); + + ASSERT_ORTSTATUS_OK(Env().CopyTensor(cpu_input, device_tensor, nullptr)); + ASSERT_ORTSTATUS_OK(Env().CopyTensor(device_tensor, cpu_output, nullptr)); + EXPECT_EQ(output_data, input_data); + } +} + // Manual Task Manager test. This intentionally reserves substantial GPU memory and runs for an extended period. // Enable explicitly with --gtest_also_run_disabled_tests and this test's full name. TEST_F(WebGpuPluginSharedAllocatorTest, DISABLED_ManualPerDeviceMemoryAndComputeLoad) { From 44100ea41f2ab72076c07a22462e9755d75db251 Mon Sep 17 00:00:00 2001 From: Arne H Juul Date: Wed, 2 Sep 2026 23:58:06 +0200 Subject: [PATCH 04/13] Give the CUDA plugin build the CUDA 13 CCCL include path (#32392) ### Description CUDA 13.0 moved libcu++, CUB and Thrust from /include to /include/cccl. onnxruntime_providers_cuda.cmake compensates for that inside config_cuda_provider_shared_module, but the plugin EP's cmake never did. The plugin target globs the same host .cc files that include CUTLASS headers (contrib_ops/cuda/llm/cutlass_heuristic.cc and others), and those reach , so on a CUDA 13 toolkit they failed with cutlass/cutlass.h:40:33: fatal error: cuda/std/utility: No such file or directory Move the include handling - and the CUDA 13.3 cudafe++ header workaround it depends on - out of onnxruntime_providers_cuda.cmake and into the new shared cmake/onnxruntime_cuda_cccl.cmake, then call it from both provider cmakes. The include flags of the in-tree provider build are unchanged. The only adjustments made while extracting the code are: * The 13.0 version check becomes an early return instead of wrapping the function body. * The generated-header directory is pinned to ${CMAKE_BINARY_DIR} rather than ${CMAKE_CURRENT_BINARY_DIR}, which in a function expands in the caller's directory scope. Both are the same directory for the two call sites (cmake/CMakeLists.txt is the top-level list file), but a shared module should not generate files in a caller-dependent location. This matches op_reduction_root in onnxruntime_providers.cmake, the closest analogue: a configure-time generated header tree injected per target. * The entry point is named ort_configure_cuda_cccl, since it generates files in the build tree as well as adding include directories. Comments record the two things that are easy to break later: the SM-specific OBJECT libraries of the plugin target inherit these include directories through a generator expression, which is why calling the function after they are created still covers them; and the CUDA 13.3 workaround is UNIX-only because it is untested on Windows, not because MSVC is known to be unaffected. ### Motivation and Context Trying a plugin build failed with: cutlass/cutlass.h:40:33: fatal error: cuda/std/utility: No such file or directory --------- Co-authored-by: Claude Opus 5 (1M context) --- .github/workflows/linux_cuda_no_cudnn.yml | 1 + .github/workflows/windows_cuda_no_cudnn.yml | 1 + cmake/onnxruntime_cuda_cccl.cmake | 95 +++++++++++++++++++ cmake/onnxruntime_providers_cuda.cmake | 66 +------------ cmake/onnxruntime_providers_cuda_plugin.cmake | 13 +++ plugin-ep-cuda/paths.txt | 1 + 6 files changed, 113 insertions(+), 64 deletions(-) create mode 100644 cmake/onnxruntime_cuda_cccl.cmake diff --git a/.github/workflows/linux_cuda_no_cudnn.yml b/.github/workflows/linux_cuda_no_cudnn.yml index 005d39147a83a..ce0a7d701dc17 100644 --- a/.github/workflows/linux_cuda_no_cudnn.yml +++ b/.github/workflows/linux_cuda_no_cudnn.yml @@ -5,6 +5,7 @@ on: branches: [main, 'rel-*'] paths: - '.github/workflows/linux_cuda_no_cudnn.yml' + - 'cmake/onnxruntime_cuda_cccl.cmake' - 'cmake/onnxruntime_providers_cuda.cmake' - 'cmake/onnxruntime_python.cmake' - 'onnxruntime/__init__.py' diff --git a/.github/workflows/windows_cuda_no_cudnn.yml b/.github/workflows/windows_cuda_no_cudnn.yml index 4c37320f74254..3ef1db73a0e1c 100644 --- a/.github/workflows/windows_cuda_no_cudnn.yml +++ b/.github/workflows/windows_cuda_no_cudnn.yml @@ -5,6 +5,7 @@ on: branches: [main, 'rel-*'] paths: - '.github/workflows/windows_cuda_no_cudnn.yml' + - 'cmake/onnxruntime_cuda_cccl.cmake' - 'cmake/onnxruntime_providers_cuda.cmake' - 'cmake/onnxruntime_providers_cuda_plugin.cmake' - 'cmake/onnxruntime_python.cmake' diff --git a/cmake/onnxruntime_cuda_cccl.cmake b/cmake/onnxruntime_cuda_cccl.cmake new file mode 100644 index 0000000000000..f0588b763e0c9 --- /dev/null +++ b/cmake/onnxruntime_cuda_cccl.cmake @@ -0,0 +1,95 @@ +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. + +# CCCL (libcu++/CUB/Thrust) header handling shared by the in-tree CUDA provider +# (onnxruntime_providers_cuda.cmake) and the CUDA plugin EP (onnxruntime_providers_cuda_plugin.cmake). +# Both compile host C++ translation units that include CUTLASS headers, which pull in , +# so both need the same include path handling. + +include_guard(GLOBAL) + +# Work around a CUDA 13.3 cudafe++ (EDG front-end) regression that mis-parses CCCL's +# global-qualified partial specializations, e.g. in : +# template +# struct ::cuda::proclaims_copyable_arguments<...> : ::cuda::std::true_type {}; +# nvcc fails with "global qualification of class name is invalid before ':' token". +# The fix is to write the specialization with the namespace reopened instead of using a +# global-qualified name. We cannot edit the (often read-only) toolkit headers, so generate +# corrected copies of the affected headers into the build tree and place that directory +# ahead of the toolkit cccl include path. This is a no-op on toolkits whose headers do not +# contain the offending pattern (e.g. once NVIDIA fixes it), so it is safe to keep enabled. +function(ort_cuda133_patch_cccl_header src dst) + if (NOT EXISTS "${src}") + return() + endif() + file(READ "${src}" _content) + set(_orig "${_content}") + # + string(REPLACE + "template \nstruct ::cuda::proclaims_copyable_arguments> : ::cuda::std::true_type\n{};" + "_CCCL_BEGIN_NAMESPACE_CUDA\ntemplate \nstruct proclaims_copyable_arguments> : ::cuda::std::true_type\n{};\n_CCCL_END_NAMESPACE_CUDA" + _content "${_content}") + # + string(REPLACE + "template <>\nstruct ::cuda::proclaims_copyable_arguments\n : ::cuda::std::true_type\n{};" + "_CCCL_BEGIN_NAMESPACE_CUDA\ntemplate <>\nstruct proclaims_copyable_arguments\n : ::cuda::std::true_type\n{};\n_CCCL_END_NAMESPACE_CUDA" + _content "${_content}") + if (NOT _content STREQUAL _orig) + get_filename_component(_dst_dir "${dst}" DIRECTORY) + file(MAKE_DIRECTORY "${_dst_dir}") + file(WRITE "${dst}" "${_content}") + elseif (EXISTS "${dst}") + # The toolkit header no longer matches the offending pattern (e.g. after a CUDA + # upgrade in an existing build tree). Remove any previously generated copy so a + # stale patched header does not keep shadowing the toolkit header. + file(REMOVE "${dst}") + endif() +endfunction() + +# Give ${target} everything it needs to compile against the CCCL headers of a CUDA 13 toolkit: +# +# * Handle the CUDA 13.0 CCCL header directory move: libcu++, CUB and Thrust moved from +# /include to /include/cccl, so - reached from the CUTLASS +# headers that host .cc files include - is no longer on the default include path of the host +# compiler. nvcc adds it by itself, so this only matters for targets that compile host C++. +# * On CUDA 13.3, generate the patched CCCL headers described above into ${CMAKE_BINARY_DIR} +# (and remove stale ones), then put that directory first on the include path. +# +# Note that this generates files in the build tree as a side effect, not just include flags. +# It must be called for every target that compiles CUDA or CUTLASS-including host sources; +# targets that inherit $ from such a target are +# covered by the parent call, but only because that generator expression is evaluated after +# configuration - see the call in onnxruntime_providers_cuda_plugin.cmake. +function(ort_configure_cuda_cccl target) + if (CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.0) + return() + endif() + + foreach(inc_dir ${CUDAToolkit_INCLUDE_DIRS}) + if (EXISTS "${inc_dir}/cccl") + # The UNIX guard is not a statement about MSVC being unaffected: the cudafe++ regression + # is simply untested on Windows, where no CUDA 13.3 build has been run. If a Windows + # CUDA 13.3 build hits the same "global qualification of class name is invalid" error, + # dropping UNIX from this condition is expected to be all that is needed. + if (UNIX AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.3 AND CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.4) + # Generate cudafe++-parseable copies of the CCCL headers that contain global-qualified + # partial specializations (see ort_cuda133_patch_cccl_header above) and put the fixed + # directory ahead of the toolkit cccl include so the corrected headers win. + set(_ort_cccl_fix_dir "${CMAKE_BINARY_DIR}/cccl_cuda13_fix") + ort_cuda133_patch_cccl_header( + "${inc_dir}/cccl/cub/device/device_transform.cuh" + "${_ort_cccl_fix_dir}/cub/device/device_transform.cuh") + ort_cuda133_patch_cccl_header( + "${inc_dir}/cccl/cub/device/dispatch/tuning/tuning_transform.cuh" + "${_ort_cccl_fix_dir}/cub/device/dispatch/tuning/tuning_transform.cuh") + if (EXISTS "${_ort_cccl_fix_dir}/cub/device/device_transform.cuh" OR + EXISTS "${_ort_cccl_fix_dir}/cub/device/dispatch/tuning/tuning_transform.cuh") + target_include_directories(${target} BEFORE PRIVATE "${_ort_cccl_fix_dir}") + endif() + endif() + + # Add the cccl subdirectory to the include path so can be found + target_include_directories(${target} PRIVATE "${inc_dir}/cccl") + endif() + endforeach() +endfunction() diff --git a/cmake/onnxruntime_providers_cuda.cmake b/cmake/onnxruntime_providers_cuda.cmake index 5deb723303c63..6cab6bf1c2844 100644 --- a/cmake/onnxruntime_providers_cuda.cmake +++ b/cmake/onnxruntime_providers_cuda.cmake @@ -66,6 +66,7 @@ ) include(onnxruntime_cuda_source_filters.cmake) + include(onnxruntime_cuda_cccl.cmake) onnxruntime_filter_cuda_cu_sources(onnxruntime_cuda_contrib_ops_cu_srcs) if (NOT onnxruntime_USE_TRT_FUSED_ATTENTION) @@ -217,44 +218,6 @@ target_compile_definitions(onnxruntime_providers_cuda PRIVATE FILE_NAME=\"onnxruntime_providers_cuda.dll\") endif() - # Work around a CUDA 13.3 cudafe++ (EDG front-end) regression that mis-parses CCCL's - # global-qualified partial specializations, e.g. in : - # template - # struct ::cuda::proclaims_copyable_arguments<...> : ::cuda::std::true_type {}; - # nvcc fails with "global qualification of class name is invalid before ':' token". - # The fix is to write the specialization with the namespace reopened instead of using a - # global-qualified name. We cannot edit the (often read-only) toolkit headers, so generate - # corrected copies of the affected headers into the build tree and place that directory - # ahead of the toolkit cccl include path. This is a no-op on toolkits whose headers do not - # contain the offending pattern (e.g. once NVIDIA fixes it), so it is safe to keep enabled. - function(ort_cuda133_patch_cccl_header src dst) - if (NOT EXISTS "${src}") - return() - endif() - file(READ "${src}" _content) - set(_orig "${_content}") - # - string(REPLACE - "template \nstruct ::cuda::proclaims_copyable_arguments> : ::cuda::std::true_type\n{};" - "_CCCL_BEGIN_NAMESPACE_CUDA\ntemplate \nstruct proclaims_copyable_arguments> : ::cuda::std::true_type\n{};\n_CCCL_END_NAMESPACE_CUDA" - _content "${_content}") - # - string(REPLACE - "template <>\nstruct ::cuda::proclaims_copyable_arguments\n : ::cuda::std::true_type\n{};" - "_CCCL_BEGIN_NAMESPACE_CUDA\ntemplate <>\nstruct proclaims_copyable_arguments\n : ::cuda::std::true_type\n{};\n_CCCL_END_NAMESPACE_CUDA" - _content "${_content}") - if (NOT _content STREQUAL _orig) - get_filename_component(_dst_dir "${dst}" DIRECTORY) - file(MAKE_DIRECTORY "${_dst_dir}") - file(WRITE "${dst}" "${_content}") - elseif (EXISTS "${dst}") - # The toolkit header no longer matches the offending pattern (e.g. after a CUDA - # upgrade in an existing build tree). Remove any previously generated copy so a - # stale patched header does not keep shadowing the toolkit header. - file(REMOVE "${dst}") - endif() - endfunction() - # config_cuda_provider_shared_module can be used to config onnxruntime_providers_cuda_obj, onnxruntime_providers_cuda & onnxruntime_providers_cuda_ut. # This function guarantees that all 3 targets have the same configurations. function(config_cuda_provider_shared_module target) @@ -422,32 +385,7 @@ target_link_libraries(${target} PRIVATE Eigen3::Eigen) target_include_directories(${target} PRIVATE ${ONNXRUNTIME_ROOT} ${CMAKE_CURRENT_BINARY_DIR} PUBLIC ${CUDAToolkit_INCLUDE_DIRS}) - # Handle CUDA 13.0 CCCL header directory move - if (CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0) - foreach(inc_dir ${CUDAToolkit_INCLUDE_DIRS}) - if (EXISTS "${inc_dir}/cccl") - if (UNIX AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.3 AND CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.4) - # Generate cudafe++-parseable copies of the CCCL headers that contain global-qualified - # partial specializations (see ort_cuda133_patch_cccl_header above) and put the fixed - # directory ahead of the toolkit cccl include so the corrected headers win. - set(_ort_cccl_fix_dir "${CMAKE_CURRENT_BINARY_DIR}/cccl_cuda13_fix") - ort_cuda133_patch_cccl_header( - "${inc_dir}/cccl/cub/device/device_transform.cuh" - "${_ort_cccl_fix_dir}/cub/device/device_transform.cuh") - ort_cuda133_patch_cccl_header( - "${inc_dir}/cccl/cub/device/dispatch/tuning/tuning_transform.cuh" - "${_ort_cccl_fix_dir}/cub/device/dispatch/tuning/tuning_transform.cuh") - if (EXISTS "${_ort_cccl_fix_dir}/cub/device/device_transform.cuh" OR - EXISTS "${_ort_cccl_fix_dir}/cub/device/dispatch/tuning/tuning_transform.cuh") - target_include_directories(${target} BEFORE PRIVATE "${_ort_cccl_fix_dir}") - endif() - endif() - - # Add the cccl subdirectory to the include path so can be found - target_include_directories(${target} PRIVATE "${inc_dir}/cccl") - endif() - endforeach() - endif() + ort_configure_cuda_cccl(${target}) # ${CMAKE_CURRENT_BINARY_DIR} is so that #include "onnxruntime_config.h" inside tensor_shape.h is found set_target_properties(${target} PROPERTIES LINKER_LANGUAGE CUDA) diff --git a/cmake/onnxruntime_providers_cuda_plugin.cmake b/cmake/onnxruntime_providers_cuda_plugin.cmake index f7032d8bcc4b7..dd5d25639b70f 100644 --- a/cmake/onnxruntime_providers_cuda_plugin.cmake +++ b/cmake/onnxruntime_providers_cuda_plugin.cmake @@ -111,6 +111,7 @@ list(FILTER CUDA_PLUGIN_EP_CU_SRCS EXCLUDE REGEX ".*/contrib_ops/cuda/transforme # Apply shared CUDA .cu source filtering (flash attention quick build, MoE GEMM FP4/FP8). include(onnxruntime_cuda_source_filters.cmake) +include(onnxruntime_cuda_cccl.cmake) onnxruntime_filter_cuda_cu_sources(CUDA_PLUGIN_EP_CU_SRCS) onnxruntime_extract_sm_specific_cuda_sources(CUDA_PLUGIN_EP_CU_SRCS SM90_SOURCES _cuda_plugin_sm90_tma_srcs @@ -447,6 +448,18 @@ target_include_directories(onnxruntime_providers_cuda_plugin PRIVATE ${cutlass_SOURCE_DIR}/tools/util/include ) +# The host .cc files globbed into this target (contrib_ops/cuda/llm/*.cc and friends) include +# CUTLASS headers, which reach . In the non-plugin build the same files are part +# of onnxruntime_providers_cuda, which gets this from config_cuda_provider_shared_module. +# +# The SM-specific OBJECT libraries created above are covered by this call even though they +# already exist: onnxruntime_add_cuda_plugin_object_library gives them this target's includes +# as $, which is +# evaluated after configuration and so picks up whatever is added here - order included, so +# the CUDA 13.3 patched-header directory keeps shadowing the toolkit CCCL headers for their +# .cu sources. Keep that indirection in mind before making the inheritance eager. +ort_configure_cuda_cccl(onnxruntime_providers_cuda_plugin) + onnxruntime_add_include_to_target( onnxruntime_providers_cuda_plugin onnxruntime_common diff --git a/plugin-ep-cuda/paths.txt b/plugin-ep-cuda/paths.txt index b39da54676e30..4d9cf8c2b108e 100644 --- a/plugin-ep-cuda/paths.txt +++ b/plugin-ep-cuda/paths.txt @@ -11,6 +11,7 @@ :(top)cmake/external/cutlass.cmake :(top)cmake/patches/cudnn_frontend :(top)cmake/patches/cutlass +:(top)cmake/onnxruntime_cuda_cccl.cmake :(top)cmake/onnxruntime_cuda_source_filters.cmake :(top)cmake/onnxruntime_providers_cuda.cmake :(top)cmake/onnxruntime_providers_cuda_plugin.cmake From a7f80f7ea734e7cfce210db2c0136dc1efc1d3a8 Mon Sep 17 00:00:00 2001 From: Ananya Anand <148831356+4n4ny4@users.noreply.github.com> Date: Wed, 2 Sep 2026 16:45:05 -0700 Subject: [PATCH 05/13] [WebGPU] Forward -i provider options to the WebGPU EP in perf_test (#32316) ### Description `perf_test -i` did not forward provider options to the WebGPU EP, so options passed that way were accepted and then silently ignored. This parses them the way every other EP branch in the file already does and hands them to `AppendExecutionProvider`. The comment also records the short key form, because the fully-qualified `ep.webgpuexecutionprovider.` form gets double-prefixed and dropped without a warning. ### Motivation and Context Both were found while trying to sweep WebGPU EP settings from `perf_test`. The failure mode is quiet in both cases: the run completes, reports a plausible number, and uses the default value. A dispatch-window sweep had to be discarded and re-run after hitting the second one. Co-authored-by: Ananya Anand <4n4ny4@users.noreply.github.com> --- onnxruntime/test/perftest/ort_test_session.cc | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) diff --git a/onnxruntime/test/perftest/ort_test_session.cc b/onnxruntime/test/perftest/ort_test_session.cc index ae9ed2a504320..1beb4e4d0c92b 100644 --- a/onnxruntime/test/perftest/ort_test_session.cc +++ b/onnxruntime/test/perftest/ort_test_session.cc @@ -607,7 +607,17 @@ select from 'TF8', 'TF16', 'UINT8', 'FLOAT', 'ITENSOR'. \n)"); #endif } else if (provider_name_ == onnxruntime::kWebGpuExecutionProvider) { #ifdef USE_WEBGPU - session_options.AppendExecutionProvider("WebGPU", {}); + // Use the short key form here: -i "enableGraphCapture|1". AppendExecutionProvider prefixes + // provider-option keys with "ep.webgpuexecutionprovider." itself, so a fully qualified key is + // double-prefixed and silently ignored. Session config entries (-C) do not reach the EP + // factory, which reads these at append time. +#ifdef _MSC_VER + std::string option_string = ToUTF8String(performance_test_config.run_config.ep_runtime_config_string); +#else + std::string option_string = performance_test_config.run_config.ep_runtime_config_string; +#endif + ParseSessionConfigs(option_string, provider_options); + session_options.AppendExecutionProvider("WebGPU", provider_options); #else ORT_THROW("WebGPU is not supported in this build\n"); #endif From c41c291836246043d4dc17b9e6a21904841fbd10 Mon Sep 17 00:00:00 2001 From: eserscor Date: Wed, 2 Sep 2026 21:21:00 -0400 Subject: [PATCH 06/13] Use core available for CUDA plugin ARM64 build. Build time 4 hours -> 1.5 hours. (#32399) ### Description ### Motivation and Context --- tools/ci_build/github/linux/build_cuda_plugin_package.sh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tools/ci_build/github/linux/build_cuda_plugin_package.sh b/tools/ci_build/github/linux/build_cuda_plugin_package.sh index a8bc55df45aca..8d014cba26377 100755 --- a/tools/ci_build/github/linux/build_cuda_plugin_package.sh +++ b/tools/ci_build/github/linux/build_cuda_plugin_package.sh @@ -27,7 +27,7 @@ done # Update parallel to the minimum of the number of processors and 8 for aarch64 to avoid OOM errors during build. if [ "$arch" = "aarch64" ]; then - PARALLEL=$(( $(nproc) < 8 ? $(nproc) : 8 )) + PARALLEL=$(( $(nproc) < 8 ? $(nproc) : $(nproc) - 2 )) else PARALLEL="" fi From 3ad12b032b5fc25a446de6c559c58276709ea712 Mon Sep 17 00:00:00 2001 From: eserscor Date: Wed, 2 Sep 2026 23:13:13 -0400 Subject: [PATCH 07/13] Primary nuget pipeline runs with CFS clean. (#32400) ### Description ### Motivation and Context --- ...pi-noopenmp-packaging-pipelines-cuda13.yml | 2 +- .../c-api-noopenmp-packaging-pipelines.yml | 3 +++ .../c-api-noopenmp-test-pipelines.yml | 13 ------------- .../custom-nuget-packaging-pipeline.yml | 2 ++ .../templates/android-java-api-aar.yml | 10 ++++++++++ .../templates/c-api-linux-cpu.yml | 1 + .../azure-pipelines/templates/esrp_nuget.yml | 2 +- .../jobs/checkout-and-git-redirect.yml | 13 ++++++++++--- .../templates/mac-esrp-dylib.yml | 2 +- .../setup-feeds-and-python-steps.yml | 19 ++++++++++++++++++- .../templates/win-esrp-dll.yml | 2 +- 11 files changed, 48 insertions(+), 21 deletions(-) diff --git a/tools/ci_build/github/azure-pipelines/c-api-noopenmp-packaging-pipelines-cuda13.yml b/tools/ci_build/github/azure-pipelines/c-api-noopenmp-packaging-pipelines-cuda13.yml index 7be878df9fde4..3eb7e9869b4aa 100644 --- a/tools/ci_build/github/azure-pipelines/c-api-noopenmp-packaging-pipelines-cuda13.yml +++ b/tools/ci_build/github/azure-pipelines/c-api-noopenmp-packaging-pipelines-cuda13.yml @@ -87,7 +87,7 @@ extends: template: v1/1ES.Official.PipelineTemplate.yml@1esPipelines parameters: settings: - networkIsolationPolicy: Permissive + networkIsolationPolicy: CFSClean2,CFSClean3,DefaultDeny featureFlags: binskimScanAllExtensions: true sdl: diff --git a/tools/ci_build/github/azure-pipelines/c-api-noopenmp-packaging-pipelines.yml b/tools/ci_build/github/azure-pipelines/c-api-noopenmp-packaging-pipelines.yml index 59f5e872a3d38..20a303add5349 100644 --- a/tools/ci_build/github/azure-pipelines/c-api-noopenmp-packaging-pipelines.yml +++ b/tools/ci_build/github/azure-pipelines/c-api-noopenmp-packaging-pipelines.yml @@ -120,6 +120,9 @@ extends: # For productions pipelines, use "Official". template: v1/1ES.Official.PipelineTemplate.yml@1esPipelines parameters: + # Waiting on the Dawn changes to land + settings: + networkIsolationPolicy: CFSClean,CFSClean2,CFSClean3,DefaultDeny featureFlags: binskimScanAllExtensions: true sdl: diff --git a/tools/ci_build/github/azure-pipelines/c-api-noopenmp-test-pipelines.yml b/tools/ci_build/github/azure-pipelines/c-api-noopenmp-test-pipelines.yml index c3a9fd05c510f..6d41f59674b95 100644 --- a/tools/ci_build/github/azure-pipelines/c-api-noopenmp-test-pipelines.yml +++ b/tools/ci_build/github/azure-pipelines/c-api-noopenmp-test-pipelines.yml @@ -77,19 +77,6 @@ stages: artifactName: 'onnxruntime-android-full-aar' ReleaseVersionSuffix: $(ReleaseVersionSuffix) -- stage: Final_AAR_Testing_Android_QNN - dependsOn: Setup - variables: - ReleaseVersionSuffix: $[ stageDependencies.Setup.Restore_And_Use_Variables.outputs['Set_Release_Version_Suffix.ReleaseVersionSuffix'] ] - jobs: - - template: templates/android-java-api-aar-test.yml - parameters: - artifactName: 'onnxruntime-android-qnn-aar' - packageName: 'onnxruntime-android-qnn' - #TODO: get this information from the setup stage - QnnSDKVersion: '2.42.0.251225' - ReleaseVersionSuffix: $(ReleaseVersionSuffix) - - template: nuget/templates/test_win.yml parameters: AgentPool: 'onnxruntime-Win-CPU-VS2022-Latest' diff --git a/tools/ci_build/github/azure-pipelines/custom-nuget-packaging-pipeline.yml b/tools/ci_build/github/azure-pipelines/custom-nuget-packaging-pipeline.yml index 24070a8f1da06..89531d9a414ac 100644 --- a/tools/ci_build/github/azure-pipelines/custom-nuget-packaging-pipeline.yml +++ b/tools/ci_build/github/azure-pipelines/custom-nuget-packaging-pipeline.yml @@ -50,6 +50,8 @@ extends: # For productions pipelines, use "Official". template: v1/1ES.Official.PipelineTemplate.yml@1esPipelines parameters: + settings: + networkIsolationPolicy: CFSClean,CFSClean2,CFSClean3,DefaultDeny featureFlags: binskimScanAllExtensions: true sdl: diff --git a/tools/ci_build/github/azure-pipelines/templates/android-java-api-aar.yml b/tools/ci_build/github/azure-pipelines/templates/android-java-api-aar.yml index 275273dd77adb..62c690c11ca60 100644 --- a/tools/ci_build/github/azure-pipelines/templates/android-java-api-aar.yml +++ b/tools/ci_build/github/azure-pipelines/templates/android-java-api-aar.yml @@ -82,6 +82,10 @@ jobs: steps: - template: jobs/checkout-and-git-redirect.yml + - template: setup-feeds-and-python-steps.yml + parameters: + architecture: x64 + - task: CmdLine@2 displayName: Create artifacts directory inputs: @@ -127,6 +131,10 @@ jobs: --volume $NDK_HOME:/ndk_home \ --volume $(Build.BinariesDirectory)/.artifacts:/home/onnxruntimedev/.artifacts \ --volume $(Build.BinariesDirectory)/.build_settings:/home/onnxruntimedev/.build_settings \ + --volume "${NPM_CONFIG_USERCONFIG}:/tmp/.npmrc:ro" \ + --volume "$HOME/.m2:/home/onnxruntimedev/.m2:ro" \ + --volume "$HOME/.gradle:/home/onnxruntimedev/.gradle" \ + --volume "${HOME}/.gitconfig:/home/onnxruntimedev/.gitconfig:ro" \ $QNN_VOLUME \ -e NIGHTLY_BUILD \ -e BUILD_BUILDNUMBER \ @@ -135,6 +143,8 @@ jobs: -e PUBLISH_EXECUTABLES=${{parameters.publish_executables}} \ -e PACKAGE_NAME=${{parameters.packageName}} \ -e RELEASE_VERSION_SUFFIX=${{parameters.ReleaseVersionSuffix}} \ + -e NPM_CONFIG_USERCONFIG=/tmp/.npmrc \ + -e PIP_INDEX_URL \ onnxruntimecpubuildcentos8x64_packaging \ /bin/bash /onnxruntime_src/tools/ci_build/github/android/build_aar_and_copy_artifacts.sh $USE_QNN workingDirectory: $(Build.SourcesDirectory) diff --git a/tools/ci_build/github/azure-pipelines/templates/c-api-linux-cpu.yml b/tools/ci_build/github/azure-pipelines/templates/c-api-linux-cpu.yml index d8efcf11be121..db436c44c5025 100644 --- a/tools/ci_build/github/azure-pipelines/templates/c-api-linux-cpu.yml +++ b/tools/ci_build/github/azure-pipelines/templates/c-api-linux-cpu.yml @@ -75,6 +75,7 @@ jobs: -e NPM_CONFIG_USERCONFIG=/tmp/.npmrc \ -e PIP_INDEX_URL \ --volume "${NPM_CONFIG_USERCONFIG}:/tmp/.npmrc:ro" \ + --volume "${HOME}/.gitconfig:/home/onnxruntimedev/.gitconfig:ro" \ --volume $HOME/.m2:/home/onnxruntimedev/.m2:ro \ --volume $HOME/.gradle:/home/onnxruntimedev/.gradle \ --volume $HOME/.onnx:/home/onnxruntimedev/.onnx -e NIGHTLY_BUILD onnxruntimecpubuildcentos8${{parameters.OnnxruntimeArch}}_packaging /bin/bash -c "python3 \ diff --git a/tools/ci_build/github/azure-pipelines/templates/esrp_nuget.yml b/tools/ci_build/github/azure-pipelines/templates/esrp_nuget.yml index ffec479474721..e79d880eb9f3f 100644 --- a/tools/ci_build/github/azure-pipelines/templates/esrp_nuget.yml +++ b/tools/ci_build/github/azure-pipelines/templates/esrp_nuget.yml @@ -5,7 +5,7 @@ parameters: steps: - ${{ if eq(parameters['DoEsrp'], 'true') }}: - - task: SFP.build-tasks.custom-build-task-1.EsrpCodeSigning@5 + - task: EsrpCodeSigning@6 displayName: 'ESRP CodeSigning' inputs: ConnectedServiceName: 'OnnxrunTimeCodeSign_20240611' diff --git a/tools/ci_build/github/azure-pipelines/templates/jobs/checkout-and-git-redirect.yml b/tools/ci_build/github/azure-pipelines/templates/jobs/checkout-and-git-redirect.yml index dee8c604dd4c6..dad80d8628b38 100644 --- a/tools/ci_build/github/azure-pipelines/templates/jobs/checkout-and-git-redirect.yml +++ b/tools/ci_build/github/azure-pipelines/templates/jobs/checkout-and-git-redirect.yml @@ -4,6 +4,12 @@ parameters: default: false steps: +- checkout: self + clean: true + fetchDepth: 1 + persistCredentials: ${{ ne(parameters.submodules, 'false') }} + submodules: false + - pwsh: | git config --global --add url."https://PAT:${env:TOKEN}@aiinfra.visualstudio.com/Lotus-Dependencies/_git/".insteadOf "https://chromium.googlesource.com/chromium/src/third_party/" git config --global --add url."https://github.com/".insteadOf "https://chromium.googlesource.com/external/github.com/" @@ -11,6 +17,7 @@ steps: env: TOKEN: $(System.AccessToken) -- checkout: self - clean: true - submodules: ${{ parameters.submodules }} +- ${{ if ne(parameters.submodules, 'false') }}: + - pwsh: | + git submodule update --init --recursive + displayName: Update submodules diff --git a/tools/ci_build/github/azure-pipelines/templates/mac-esrp-dylib.yml b/tools/ci_build/github/azure-pipelines/templates/mac-esrp-dylib.yml index 5e6cd2240feba..3e8991c3e5e3a 100644 --- a/tools/ci_build/github/azure-pipelines/templates/mac-esrp-dylib.yml +++ b/tools/ci_build/github/azure-pipelines/templates/mac-esrp-dylib.yml @@ -8,7 +8,7 @@ parameters: default: '*.zip' steps: -- task: SFP.build-tasks.custom-build-task-1.EsrpCodeSigning@5 +- task: EsrpCodeSigning@6 displayName: 'ESRP CodeSigning' inputs: ConnectedServiceName: 'OnnxrunTimeCodeSign_20240611' diff --git a/tools/ci_build/github/azure-pipelines/templates/setup-feeds-and-python-steps.yml b/tools/ci_build/github/azure-pipelines/templates/setup-feeds-and-python-steps.yml index 04eb2793d14d4..3979d9dd41b4f 100644 --- a/tools/ci_build/github/azure-pipelines/templates/setup-feeds-and-python-steps.yml +++ b/tools/ci_build/github/azure-pipelines/templates/setup-feeds-and-python-steps.yml @@ -163,7 +163,8 @@ steps: Set-Content $m2SettingsXmlPath - # gradle - inject 1ES feed repo into all projects/buildscripts. Auth using creds from maven settings.xml. + # gradle - inject 1ES feed repo into plugin management, all projects, and buildscripts. + # Auth using creds from maven settings.xml. # nb: see HACK note above. $gradleInitDir = if ($env:GRADLE_HOME_DIR) { $env:GRADLE_HOME_DIR } else { "~/.gradle" } if (!(Test-Path -Path "${gradleInitDir}/init.d")) { @@ -185,6 +186,22 @@ steps: } } + settingsEvaluated { settings -> + settings.pluginManagement { + repositories { + clear() + maven { + name "1ES Feed" + url uri("https://aiinfra.pkgs.visualstudio.com/_packaging/${env:central_maven_feed_name}/maven/v1") + credentials(PasswordCredentials) { + username oneEsFeedCredentials().username + password oneEsFeedCredentials().password + } + } + } + } + } + allprojects { buildscript { repositories { diff --git a/tools/ci_build/github/azure-pipelines/templates/win-esrp-dll.yml b/tools/ci_build/github/azure-pipelines/templates/win-esrp-dll.yml index 0476bc74349bf..39d8824813844 100644 --- a/tools/ci_build/github/azure-pipelines/templates/win-esrp-dll.yml +++ b/tools/ci_build/github/azure-pipelines/templates/win-esrp-dll.yml @@ -16,7 +16,7 @@ parameters: default: '*.dll' steps: -- task: SFP.build-tasks.custom-build-task-1.EsrpCodeSigning@5 +- task: EsrpCodeSigning@6 displayName: ${{ parameters.DisplayName }} condition: and(succeeded(), eq('${{ parameters.DoEsrp }}', true)) inputs: From c5ead598c8ec2483eea127a5709c20cec57ac33a Mon Sep 17 00:00:00 2001 From: Wanming Lin Date: Thu, 3 Sep 2026 11:36:56 +0800 Subject: [PATCH 08/13] [WebNN EP] Support SkipLayerNormalization op (#32377) `SkipLayerNormalization` can be decomposed by `LayerNormalization` + `Add`. > output=LayerNorm(X+skip+bias) --- js/web/docs/webnn-operators.md | 1 + .../builders/impl/normalization_op_builder.cc | 77 +++++++++++++++---- .../core/providers/webnn/builders/map_info.h | 1 + .../webnn/builders/op_builder_factory.cc | 1 + 4 files changed, 66 insertions(+), 14 deletions(-) diff --git a/js/web/docs/webnn-operators.md b/js/web/docs/webnn-operators.md index dbf2ebb01ad4f..f9871680c3260 100644 --- a/js/web/docs/webnn-operators.md +++ b/js/web/docs/webnn-operators.md @@ -107,6 +107,7 @@ platforms. Check the [WebNN status](https://webmachinelearning.github.io/webnn-s | SimplifiedLayerNormalization | ai.onnx(1+) | pow, reduceMean, add, sqrt, div, mul | | | Sigmoid | ai.onnx(7-12, 13+) | sigmoid | | | Sign | ai.onnx(9-12, 13+) | sign | | +| SkipLayerNormalization | com.microsoft(1+) | add, layerNormalization | | | SkipSimplifiedLayerNormalization | com.microsoft(1+) | pow, reduceMean, add, sqrt, div, mul | | | Softplus | ai.onnx(7+) | softplus | | | Softsign | ai.onnx(7+) | softsign | | diff --git a/onnxruntime/core/providers/webnn/builders/impl/normalization_op_builder.cc b/onnxruntime/core/providers/webnn/builders/impl/normalization_op_builder.cc index 8a42a6e3dd74e..ace35020a1b6b 100644 --- a/onnxruntime/core/providers/webnn/builders/impl/normalization_op_builder.cc +++ b/onnxruntime/core/providers/webnn/builders/impl/normalization_op_builder.cc @@ -46,11 +46,16 @@ Status NormalizationOpBuilder::AddToModelBuilderImpl(ModelBuilder& model_builder emscripten::val options = emscripten::val::object(); options.set("label", node.Name()); - const size_t scale_input_index = op_type == "SkipSimplifiedLayerNormalization" ? 2 : 1; + // Skip normalization variants (Skip[Simplified]LayerNormalization) take input and skip as + // inputs 0 and 1, so their scale/bias inputs are shifted by one compared to the non-skip ops. + const bool is_skip_norm = + op_type == "SkipLayerNormalization" || op_type == "SkipSimplifiedLayerNormalization"; + + const size_t scale_input_index = is_skip_norm ? 2 : 1; emscripten::val scale = model_builder.GetOperand(input_defs[scale_input_index]->Name()); options.set("scale", scale); - const size_t bias_input_index = op_type == "SkipSimplifiedLayerNormalization" ? 3 : 2; + const size_t bias_input_index = is_skip_norm ? 3 : 2; emscripten::val bias = emscripten::val::undefined(); if (TensorExists(input_defs, bias_input_index)) { // Bias input exists. @@ -59,7 +64,10 @@ Status NormalizationOpBuilder::AddToModelBuilderImpl(ModelBuilder& model_builder } NodeAttrHelper helper(node); - const auto epsilon = helper.Get("epsilon", 1e-05f); + // The com.microsoft Skip[Simplified]LayerNormalization schema defaults epsilon to 1e-12f + // (kDefaultSkipLayerNormEpsilon), whereas the other normalization ops default to 1e-5f. + const float default_epsilon = is_skip_norm ? 1e-12f : 1e-05f; + const auto epsilon = helper.Get("epsilon", default_epsilon); options.set("epsilon", epsilon); emscripten::val output = emscripten::val::undefined(); @@ -69,6 +77,7 @@ Status NormalizationOpBuilder::AddToModelBuilderImpl(ModelBuilder& model_builder emscripten::val variance = model_builder.GetOperand(input_defs[4]->Name()); output = model_builder.GetBuilder().call("batchNormalization", input, mean, variance, options); } else if (op_type == "LayerNormalization" || + op_type == "SkipLayerNormalization" || op_type == "SimplifiedLayerNormalization" || op_type == "SkipSimplifiedLayerNormalization") { int64_t axis = helper.Get("axis", -1); @@ -76,7 +85,34 @@ Status NormalizationOpBuilder::AddToModelBuilderImpl(ModelBuilder& model_builder std::vector axes(rank - SafeInt(axis)); std::iota(axes.begin(), axes.end(), axis); - if (op_type == "LayerNormalization") { + if (op_type == "LayerNormalization" || op_type == "SkipLayerNormalization") { + /** + WebNN's layerNormalization already implements the full (non-simplified) layer normalization. + SkipLayerNormalization has no native WebNN op, but it is just a layerNormalization applied to + the residual sum, so we compute the residual first and reuse layerNormalization: + input_skip_bias_sum = X + skip + bias (bias is optional, input index 4) + output = layerNormalization(input_skip_bias_sum) with scale (gamma) and bias (beta). + */ + if (op_type == "SkipLayerNormalization") { + emscripten::val common_options = emscripten::val::object(); + emscripten::val skip = model_builder.GetOperand(input_defs[1]->Name()); + common_options.set("label", node.Name() + "_add_skip"); + input = model_builder.GetBuilder().call("add", input, skip, common_options); + + // Optional pre-normalization bias (input index 4), added to the residual sum. + if (TensorExists(input_defs, 4)) { + emscripten::val pre_norm_bias = model_builder.GetOperand(input_defs[4]->Name()); + common_options.set("label", node.Name() + "_add_bias"); + input = model_builder.GetBuilder().call("add", input, pre_norm_bias, common_options); + } + + // Optional output input_skip_bias_sum (output index 3). Now input equals to input_skip_bias_sum. + if (TensorExists(output_defs, 3)) { + emscripten::val input_skip_bias_sum = input; + model_builder.AddOperand(output_defs[3]->Name(), input_skip_bias_sum); + } + } + options.set("axes", emscripten::val::array(axes)); output = model_builder.GetBuilder().call("layerNormalization", input, options); } else { // SimplifiedLayerNormalization or SkipSimplifiedLayerNormalization @@ -145,8 +181,10 @@ Status NormalizationOpBuilder::AddToModelBuilderImpl(ModelBuilder& model_builder common_options.set("label", node.Name() + "_mul"); output = model_builder.GetBuilder().call("mul", scale, div, common_options); - // Add (if bias exists) - if (!bias.isUndefined()) { + // Add (if bias exists). + // For SkipSimplifiedLayerNormalization the bias is the residual bias and has already been + // folded into input_skip_bias_sum above, so it must not be added again post-normalization. + if (!bias.isUndefined() && op_type != "SkipSimplifiedLayerNormalization") { common_options.set("label", node.Name() + "_add_bias"); output = model_builder.GetBuilder().call("add", output, bias, common_options); } @@ -218,14 +256,14 @@ bool NormalizationOpBuilder::IsOpSupportedImpl(const GraphViewer&, } const auto& output_defs = node.OutputDefs(); - if (op_type == "SkipSimplifiedLayerNormalization") { + if (op_type == "SkipLayerNormalization" || op_type == "SkipSimplifiedLayerNormalization") { if (output_defs.size() > 4) { - LOGS(logger, VERBOSE) << "SkipSimplifiedLayerNormalization output count must not exceed 4."; + LOGS(logger, VERBOSE) << op_type << " output count must not exceed 4."; return false; } if (TensorExists(output_defs, 1) || TensorExists(output_defs, 2)) { // Output mean and inv_std_var are used for training mode, which is not supported. - LOGS(logger, VERBOSE) << "SkipSimplifiedLayerNormalization's output mean and inv_std_var are not supported."; + LOGS(logger, VERBOSE) << op_type << "'s output mean and inv_std_var are not supported."; return false; } } else { @@ -274,10 +312,13 @@ bool NormalizationOpBuilder::HasSupportedInputsImpl(const GraphViewer&, const No return false; } - if (op_type == "SimplifiedLayerNormalization" || op_type == "SkipSimplifiedLayerNormalization") { - // SkipSimplifiedLayerNormalization and SimplifiedLayerNormalization are supported by decomposed WebNN ops. + if (op_type == "SkipLayerNormalization" || + op_type == "SimplifiedLayerNormalization" || + op_type == "SkipSimplifiedLayerNormalization") { + // These ops are supported by decomposed WebNN ops (see decomposed_op_map). // Check if the input data type is supported by each decomposed WebNN op. - // Decomposed ops include: "Add", "Div", "Mul", "Pow", "ReduceMean" and "Sqrt". + // SkipLayerNormalization decomposes into: "Add", "LayerNormalization". + // Simplified/SkipSimplifiedLayerNormalization: "Add", "Div", "Mul", "Pow", "ReduceMean", "Sqrt". for (const std::string_view decomposed_op_type : decomposed_op_map.at(op_type)) { const std::string_view webnn_op_type = GetWebNNOpType(decomposed_op_type); const std::string_view webnn_input_name = GetWebNNOpFirstInputName(decomposed_op_type); @@ -292,7 +333,12 @@ bool NormalizationOpBuilder::HasSupportedInputsImpl(const GraphViewer&, const No return false; } // It's complicated to check all the decomposed ops' input rank support. - // Ensure at least the first input rank is supported by the decomposed ops (pow and div accept the first input). + // Ensure at least the first input rank is supported by the ops that consume the (residual) input. + if (op_type == "SkipLayerNormalization") { + return IsRankSupportedByWebNNOp(wnn_limits, "add", "a", input_shape.size(), node.Name(), logger) && + IsRankSupportedByWebNNOp(wnn_limits, "layerNormalization", "input", input_shape.size(), + node.Name(), logger); + } return IsRankSupportedByWebNNOp(wnn_limits, "pow", "a", input_shape.size(), node.Name(), logger) && IsRankSupportedByWebNNOp(wnn_limits, "div", "a", input_shape.size(), node.Name(), logger); } else { @@ -327,7 +373,9 @@ bool NormalizationOpBuilder::HasSupportedOutputsImpl(const Node& node, return false; } - if (op_type == "SimplifiedLayerNormalization" || op_type == "SkipSimplifiedLayerNormalization") { + if (op_type == "SkipLayerNormalization" || + op_type == "SimplifiedLayerNormalization" || + op_type == "SkipSimplifiedLayerNormalization") { // Check if the output data type is supported by every decomposed WebNN op. for (const std::string_view decomposed_op_type : decomposed_op_map.at(op_type)) { const std::string_view webnn_op_type = GetWebNNOpType(decomposed_op_type); @@ -365,6 +413,7 @@ void CreateNormalizationOpBuilder(const std::string& op_type, OpBuilderRegistrat "BatchNormalization", "InstanceNormalization", "LayerNormalization", + "SkipLayerNormalization", "SimplifiedLayerNormalization", "SkipSimplifiedLayerNormalization", }; diff --git a/onnxruntime/core/providers/webnn/builders/map_info.h b/onnxruntime/core/providers/webnn/builders/map_info.h index 77af87410fa14..481da4fabf2fb 100644 --- a/onnxruntime/core/providers/webnn/builders/map_info.h +++ b/onnxruntime/core/providers/webnn/builders/map_info.h @@ -62,6 +62,7 @@ const std::map> decomposed_op_ma {"MultiHeadAttention", {"Add", "Cast", "Concat", "Div", "MatMul", "Reshape", "Softmax", "Transpose"}}, {"RotaryEmbedding", {"Add", "Concat", "Gather", "Mul", "Reshape", "Slice", "Split"}}, {"SimplifiedLayerNormalization", {"Add", "Div", "Mul", "Pow", "ReduceMean", "Sqrt"}}, + {"SkipLayerNormalization", {"Add", "LayerNormalization"}}, {"SkipSimplifiedLayerNormalization", {"Add", "Div", "Mul", "Pow", "ReduceMean", "Sqrt"}}, }; diff --git a/onnxruntime/core/providers/webnn/builders/op_builder_factory.cc b/onnxruntime/core/providers/webnn/builders/op_builder_factory.cc index 9e03f6f9aaf89..b1fc870b3d249 100644 --- a/onnxruntime/core/providers/webnn/builders/op_builder_factory.cc +++ b/onnxruntime/core/providers/webnn/builders/op_builder_factory.cc @@ -185,6 +185,7 @@ static OpBuilderRegistrations CreateOpBuilderRegistrations() { CreateNormalizationOpBuilder("BatchNormalization", op_registrations); CreateNormalizationOpBuilder("InstanceNormalization", op_registrations); CreateNormalizationOpBuilder("LayerNormalization", op_registrations); + CreateNormalizationOpBuilder("SkipLayerNormalization", op_registrations); CreateNormalizationOpBuilder("SimplifiedLayerNormalization", op_registrations); CreateNormalizationOpBuilder("SkipSimplifiedLayerNormalization", op_registrations); } From db3bc5a06b0ba92c318e460ab6f6e9c738670dda Mon Sep 17 00:00:00 2001 From: Jianhui Dai Date: Thu, 3 Sep 2026 15:09:59 +0800 Subject: [PATCH 09/13] [WebGPU] Fix 1D-dispatch shader fast path (#32343) ### Description ShaderHelper::Init() special-cased shaders whose dispatch grid has a single row of workgroups (dispatch_group_size_y_ == 1 && dispatch_group_size_z_ == 1) by computing global_idx directly from `@builtin(global_invocation_id).x` and workgroup_idx from `@builtin(workgroup_id).x`. This is only correct when the workgroup itself is also 1D (workgroup_size_y == 1 && workgroup_size_z == 1). For a program dispatched as a single row of 2D/3D workgroups (e.g. workgroup_size = (8, 8, 1)), global_id.x omits the contribution of local_invocation_id.y/z that local_invocation_index folds in, so global_idx (and workgroup_idx-derived offsets) come out wrong and the shader reads/writes incorrect elements. Remove the special-cased fast path so every non-indirect dispatch uses the general num_workgroups-based formula, which is correct regardless of workgroup or dispatch shape. Also drop the now-unnecessary is_1d_dispatch distinction from the program cache key (program_cache_key.h/.cc, webgpu_context.cc), since generated shader source no longer varies with dispatch dimensionality. ### Motivation and Context See above. --- onnxruntime/core/providers/webgpu/program_cache_key.cc | 7 ++----- onnxruntime/core/providers/webgpu/program_cache_key.h | 3 +-- onnxruntime/core/providers/webgpu/shader_helper.cc | 5 ----- onnxruntime/core/providers/webgpu/webgpu_context.cc | 4 +--- 4 files changed, 4 insertions(+), 15 deletions(-) diff --git a/onnxruntime/core/providers/webgpu/program_cache_key.cc b/onnxruntime/core/providers/webgpu/program_cache_key.cc index ac40e8f4de07c..ecb49dc7485c1 100644 --- a/onnxruntime/core/providers/webgpu/program_cache_key.cc +++ b/onnxruntime/core/providers/webgpu/program_cache_key.cc @@ -52,17 +52,15 @@ void AppendTensorInfo(OStringStream& ss, std::string CalculateProgramCacheKey(const ProgramBase& program, std::span inputs_segments, - std::span outputs_segments, - bool is_1d_dispatch) { + std::span outputs_segments) { SS(ss, kStringInitialSizeCacheKey); // final key format: - // =[]::::: + // =[]:::: // // = ||... // = ,, // = - // = // = ||... // = // = ||... @@ -97,7 +95,6 @@ std::string CalculateProgramCacheKey(const ProgramBase& program, ss << ":" D("SubgroupSize=") << subgroup_size; } - ss << ":" D("DispatchDim=") << (is_1d_dispatch ? "1" : "3"); ss << ":" D("UniformSizes="); bool first = true; for (const auto& uniform : program.UniformVariables()) { diff --git a/onnxruntime/core/providers/webgpu/program_cache_key.h b/onnxruntime/core/providers/webgpu/program_cache_key.h index e34a7ab8efe55..b56000c8e5c00 100644 --- a/onnxruntime/core/providers/webgpu/program_cache_key.h +++ b/onnxruntime/core/providers/webgpu/program_cache_key.h @@ -13,8 +13,7 @@ namespace webgpu { std::string CalculateProgramCacheKey(const ProgramBase& program, std::span inputs_segments, - std::span outputs_segments, - bool is_1d_dispatch); + std::span outputs_segments); } // namespace webgpu } // namespace onnxruntime diff --git a/onnxruntime/core/providers/webgpu/shader_helper.cc b/onnxruntime/core/providers/webgpu/shader_helper.cc index 137d02d86a0c8..c69994c8dded5 100644 --- a/onnxruntime/core/providers/webgpu/shader_helper.cc +++ b/onnxruntime/core/providers/webgpu/shader_helper.cc @@ -73,7 +73,6 @@ Status ShaderHelper::Init() { } // init body string stream - bool is_1d_dispatch = dispatch_group_size_y_ == 1 && dispatch_group_size_z_ == 1; bool use_indirect_dispatch = program_.IndirectDispatchTensor() != nullptr; // append header for main function so it is ready for user to append main function body @@ -103,10 +102,6 @@ Status ShaderHelper::Init() { " let num_workgroups_y = indirect_buffer[1];\n" " let workgroup_idx = workgroup_id.z * num_workgroups_x * num_workgroups_y + workgroup_id.y * num_workgroups_x + workgroup_id.x;\n" " let global_idx = workgroup_idx * (workgroup_size_x * workgroup_size_y * workgroup_size_z) + local_idx;\n"; - } else if (is_1d_dispatch) { - body_ss_ << ") {\n"; - body_ss_ << " let global_idx = global_id.x;\n" - " let workgroup_idx = workgroup_id.x;\n"; } else { body_ss_ << ",\n" " @builtin(num_workgroups) num_workgroups : vec3) {\n"; diff --git a/onnxruntime/core/providers/webgpu/webgpu_context.cc b/onnxruntime/core/providers/webgpu/webgpu_context.cc index 37ccccd2ca3b6..ff85ff95e3866 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_context.cc +++ b/onnxruntime/core/providers/webgpu/webgpu_context.cc @@ -507,9 +507,7 @@ Status WebGpuContext::Run(ComputeContextBase& context, const ProgramBase& progra "Only one of SetIndirectDispatchTensor and SetDispatchGroupSize should be called for program", program.Name()); } - bool is_1d_dispatch = (y == 1 && z == 1); - - auto key = CalculateProgramCacheKey(program, inputs_segments, outputs_segments, is_1d_dispatch); + auto key = CalculateProgramCacheKey(program, inputs_segments, outputs_segments); LOGS(context.Logger(), INFO) << "Starting program \"" << key << "\" (" << x << ", " << y << ", " << z << ")"; // The program cache prevents duplicate builds across encoded windows. From ff8d83b82d87a0b019e40a8c0ff11621a7caf39b Mon Sep 17 00:00:00 2001 From: NIKHIL Date: Thu, 3 Sep 2026 13:28:15 +0530 Subject: [PATCH 10/13] Java: release provider option values with the correct strings (#31828) ## Summary - release CUDA provider option values against the `jstring` objects they came from - apply the same correction to the TensorRT provider option binding Fixes #31706 ## Testing - `git diff --check` - A JNI/Gradle build could not be run locally because this Windows environment has no JDK or C compiler. The change is limited to the two provider-option release loops and is ready for repository CI. ## AI assistance disclosure I identified and reproduced the issue, then used OpenAI Codex to assist with the implementation and regression-test work. I personally reviewed the resulting diff and ran the validation commands listed above to verify the fix. Any full-suite or local-environment limitations are documented in the validation section. --------- Co-authored-by: Scott McKay --- .../native/ai_onnxruntime_providers_OrtCUDAProviderOptions.c | 2 +- .../ai_onnxruntime_providers_OrtTensorRTProviderOptions.c | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/java/src/main/native/ai_onnxruntime_providers_OrtCUDAProviderOptions.c b/java/src/main/native/ai_onnxruntime_providers_OrtCUDAProviderOptions.c index 46df515c2e235..6f07ac8f1c9fc 100644 --- a/java/src/main/native/ai_onnxruntime_providers_OrtCUDAProviderOptions.c +++ b/java/src/main/native/ai_onnxruntime_providers_OrtCUDAProviderOptions.c @@ -58,7 +58,7 @@ JNIEXPORT void JNICALL Java_ai_onnxruntime_providers_OrtCUDAProviderOptions_appl for (jsize i = 0; i < keyLength; i++) { jobject key = (*jniEnv)->GetObjectArrayElement(jniEnv, jKeyArr, i); (*jniEnv)->ReleaseStringUTFChars(jniEnv,key,keys[i]); - jobject value = (*jniEnv)->GetObjectArrayElement(jniEnv, jKeyArr, i); + jobject value = (*jniEnv)->GetObjectArrayElement(jniEnv, jValueArr, i); (*jniEnv)->ReleaseStringUTFChars(jniEnv,value,values[i]); } free((void*)keys); diff --git a/java/src/main/native/ai_onnxruntime_providers_OrtTensorRTProviderOptions.c b/java/src/main/native/ai_onnxruntime_providers_OrtTensorRTProviderOptions.c index 404a80f118306..b355867988527 100644 --- a/java/src/main/native/ai_onnxruntime_providers_OrtTensorRTProviderOptions.c +++ b/java/src/main/native/ai_onnxruntime_providers_OrtTensorRTProviderOptions.c @@ -57,7 +57,7 @@ JNIEXPORT void JNICALL Java_ai_onnxruntime_providers_OrtTensorRTProviderOptions_ for (jsize i = 0; i < keyLength; i++) { jobject key = (*jniEnv)->GetObjectArrayElement(jniEnv, jKeyArr, i); (*jniEnv)->ReleaseStringUTFChars(jniEnv,key,keys[i]); - jobject value = (*jniEnv)->GetObjectArrayElement(jniEnv, jKeyArr, i); + jobject value = (*jniEnv)->GetObjectArrayElement(jniEnv, jValueArr, i); (*jniEnv)->ReleaseStringUTFChars(jniEnv,value,values[i]); } free((void*)keys); From d46218a8ca85da83542cbf2625bfcc84cddca5b8 Mon Sep 17 00:00:00 2001 From: Yateng Hong Date: Thu, 3 Sep 2026 20:26:29 +0800 Subject: [PATCH 11/13] [TRT] Fix trt shape values (#32415) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ### Description Fix the TensorRT 10.11+ shape-tensor optimization profile setup by passing the configured optimum and maximum values to their matching  kOPT  and  kMAX  profile selectors. ### Motivation and Context Swapping OPT and MAX can make an valid profile invalid, causing TensorRT engine construction to fail when the configured optimum and maximum values differ. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .../providers/tensorrt/tensorrt_execution_provider.cc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/onnxruntime/core/providers/tensorrt/tensorrt_execution_provider.cc b/onnxruntime/core/providers/tensorrt/tensorrt_execution_provider.cc index ea7fc2b7bd275..2d206a1251f84 100644 --- a/onnxruntime/core/providers/tensorrt/tensorrt_execution_provider.cc +++ b/onnxruntime/core/providers/tensorrt/tensorrt_execution_provider.cc @@ -545,8 +545,8 @@ bool ApplyProfileShapesFromProviderOptions(std::vector shapes_opt_64(shapes_opt.begin(), shapes_opt.end()); std::vector shapes_max_64(shapes_max.begin(), shapes_max.end()); trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kMIN, &shapes_min_64[0], shape_size); - trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kMAX, &shapes_opt_64[0], shape_size); - trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kOPT, &shapes_max_64[0], shape_size); + trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kMAX, &shapes_max_64[0], shape_size); + trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kOPT, &shapes_opt_64[0], shape_size); #else trt_profile->setShapeValues(input_name.c_str(), nvinfer1::OptProfileSelector::kMIN, &shapes_min[0], shape_size); trt_profile->setShapeValues(input_name.c_str(), nvinfer1::OptProfileSelector::kMAX, &shapes_max[0], shape_size); @@ -761,8 +761,8 @@ Status ApplyProfileShapesFromInputTensorValue(std::vector shapes_opt_64(shapes_opt.begin(), shapes_opt.end()); std::vector shapes_max_64(shapes_max.begin(), shapes_max.end()); trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kMIN, &shapes_min_64[0], shape_size); - trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kMAX, &shapes_opt_64[0], shape_size); - trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kOPT, &shapes_max_64[0], shape_size); + trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kMAX, &shapes_max_64[0], shape_size); + trt_profile->setShapeValuesV2(input_name.c_str(), nvinfer1::OptProfileSelector::kOPT, &shapes_opt_64[0], shape_size); #else trt_profile->setShapeValues(input_name.c_str(), nvinfer1::OptProfileSelector::kMIN, &shapes_min[0], shape_size); trt_profile->setShapeValues(input_name.c_str(), nvinfer1::OptProfileSelector::kMAX, &shapes_max[0], shape_size); From dcc70f5cfbdc4f5ee0f8e75580b8e26de66518a9 Mon Sep 17 00:00:00 2001 From: Copilot <198982749+Copilot@users.noreply.github.com> Date: Thu, 3 Sep 2026 11:36:06 -0700 Subject: [PATCH 12/13] [CUDA EP] Slice: single-memcpy fast path for contiguous subregions (#28902) ### Description CUDA EP's `Slice` always allocated an output buffer and launched the per-element `_SliceKernel`, even when the output is a contiguous subregion of the input (e.g. `tensor[i:j, :]`). This adds a fast path that copies such slices with a single `cudaMemcpyAsync` (DeviceToDevice) from `input_ptr + offset`, skipping the per-element index decomposition. - **`TryComputeContiguousSliceOffset`** (`core/providers/cuda/tensor/slice.cc`): detects the contiguous case and returns the start element offset. Conditions: all steps == 1; scanning from the rightmost axis, axes are fully included up to the first trimmed (pivot) axis; every axis left of the pivot selects exactly one element. Offset is computed from row-major strides. - **`Slice::CallSliceImp`**: takes the memcpy path when the slice is contiguous, otherwise falls back to the existing kernel. Placed in `CallSliceImp` rather than the shared `ComputeInternal`, so training `SliceGrad` (which overrides `CallSliceImp` with scatter semantics) is unaffected. - **Tests** (`slice_op.test.cc`): added `Slice3D_ContiguousLeadingAxis` and `Slice3D_ContiguousSingleLeadingIndex`, exercising the path across CPU/CUDA EPs. Covered patterns: leading-axis slicing, batch splitting on axis 0, and single-index leading selection. Strided, middle-axis, and inner-axis slices remain on the kernel path. ### Motivation and Context Models with repeated leading-axis slicing (e.g. per-layer embedding extraction) incur one kernel launch per `Slice`. Replacing the trivially-contiguous cases with a single memcpy removes the per-element index computation and launch overhead, matching how `Reshape`/`Squeeze`/`Flatten` avoid copies. Full zero-copy aliasing would require allocation-planner support for offset aliases; this is the lower-risk memcpy alternative noted in the issue. --------- Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: Justin Chu Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- .../core/providers/cuda/tensor/slice.cc | 132 +++++++++++++++++- .../core/providers/cuda/tensor/slice.h | 5 +- .../providers/cpu/tensor/slice_op.test.cc | 48 +++++++ .../training_ops/cuda/tensor/slice_grad.cc | 5 +- .../training_ops/cuda/tensor/slice_grad.h | 5 +- 5 files changed, 190 insertions(+), 5 deletions(-) diff --git a/onnxruntime/core/providers/cuda/tensor/slice.cc b/onnxruntime/core/providers/cuda/tensor/slice.cc index 34de6eeac3ea2..37b83ad9c6da0 100644 --- a/onnxruntime/core/providers/cuda/tensor/slice.cc +++ b/onnxruntime/core/providers/cuda/tensor/slice.cc @@ -6,6 +6,9 @@ #include "core/providers/cpu/tensor/slice_helper.h" #include "core/providers/cuda/tensor/slice_impl.h" +#include +#include + namespace onnxruntime { namespace cuda { // this really doesn't need to be a typed registration as the indices come from attributes and can only be int64. @@ -102,6 +105,90 @@ static Status SliceImpCore(cudaStream_t stream, output_shape.Size()); } +static bool TryConvertInt64ToSizeT(int64_t value, size_t& converted_value) { + if (value < 0 || static_cast(value) > std::numeric_limits::max()) { + return false; + } + + converted_value = static_cast(value); + return true; +} + +static bool TryMultiply(size_t lhs, size_t rhs, size_t& product) { + if (lhs != 0 && rhs > std::numeric_limits::max() / lhs) { + return false; + } + + product = lhs * rhs; + return true; +} + +// Detect whether a step-1 slice selects a single contiguous block of the input tensor. +// This is the case when the slice only trims leading dimensions: scanning from the right, +// every axis is fully included until the first trimmed ("pivot") axis, and every axis to the +// left of the pivot selects exactly one element. When true, the output is the contiguous +// sub-region input_ptr + offset_in_elements and can be produced with a single device-to-device +// memcpy instead of the per-element slice kernel. +static bool TryComputeContiguousSliceOffset(gsl::span input_dims, + gsl::span output_dims, + const TArray& starts_buffer, + const TArray& steps_buffer, + size_t& offset_in_elements) { + const int32_t rank = static_cast(input_dims.size()); + if (rank == 0 || static_cast(output_dims.size()) != rank || + starts_buffer.Size() != rank || steps_buffer.Size() != rank) { + return false; + } + + // Only step-1 slices can be contiguous. + for (int32_t i = 0; i < rank; ++i) { + if (steps_buffer[i] != 1) { + return false; + } + } + + // Find the first trimmed axis scanning from the right (the pivot). + int32_t pivot = -1; + for (int32_t i = rank - 1; i >= 0; --i) { + if (output_dims[i] != input_dims[i]) { + pivot = i; + break; + } + } + + // Every axis to the left of the pivot must select exactly one element, otherwise the + // selected region is split into multiple non-adjacent blocks. + for (int32_t i = 0; i < pivot; ++i) { + if (output_dims[i] != 1) { + return false; + } + } + + // Compute the offset of the first selected element (in elements). + size_t offset = 0; + size_t stride = 1; + for (int32_t i = rank - 1; i >= 0; --i) { + size_t start = 0; + size_t input_dim = 0; + if (!TryConvertInt64ToSizeT(starts_buffer[i], start) || + !TryConvertInt64ToSizeT(input_dims[i], input_dim) || + start > input_dim) { + return false; + } + + size_t offset_increment = 0; + if (!TryMultiply(start, stride, offset_increment) || + offset > std::numeric_limits::max() - offset_increment || + !TryMultiply(stride, input_dim, stride)) { + return false; + } + + offset += offset_increment; + } + offset_in_elements = offset; + return true; +} + namespace SliceCuda { static Status ComputeSliceStrides(const TensorShape& input_shape, TArray& input_strides, @@ -192,10 +279,18 @@ Status Slice::ComputeInternal(OpKernelContext* ctx) const { ORT_RETURN_IF_ERROR(SliceCuda::ComputeSliceStrides(input_shape, input_strides, output_strides, compute_metadata)); + gsl::span sliced_input_dims = input_dimensions; + gsl::span sliced_output_dims = compute_metadata.output_dims_; + if (compute_metadata.p_flattened_input_dims_) { + sliced_input_dims = compute_metadata.flattened_input_dims_; + sliced_output_dims = compute_metadata.flattened_output_dims_; + } + // It may seem that we may use `SliceImpCore()` directly, but we need to go through `CallSliceImp()` because // `ComputeInternal()` is shared between the inferencing and training kernels and the training kernel overrides // `CallSliceImp()` - ORT_RETURN_IF_ERROR(CallSliceImp(input_tensor->DataType()->Size(), input_dimensions.size(), starts_buffer, + ORT_RETURN_IF_ERROR(CallSliceImp(input_tensor->DataType()->Size(), input_dimensions.size(), + sliced_input_dims, sliced_output_dims, starts_buffer, steps_buffer, input_strides, output_strides, ctx, output_shape)); @@ -217,13 +312,46 @@ Status Slice::FillInputVectors(OpKernelContext* ctx, TensorShapeVector& } template -Status Slice::CallSliceImp(size_t element_size, size_t dimension_count, const TArray& starts_buffer, +Status Slice::CallSliceImp(size_t element_size, size_t dimension_count, + gsl::span sliced_input_dims, + gsl::span sliced_output_dims, + const TArray& starts_buffer, const TArray& steps_buffer, const TArray& input_strides, const TArray& output_strides, OpKernelContext* ctx, const TensorShape& output_shape) const { const auto* input_tensor = ctx->Input(0); auto* output_tensor = ctx->Output(0, output_shape); + const int64_t output_size = output_shape.Size(); + if (output_size == 0) { + return Status::OK(); + } + + // Fast path: when the slice selects a single contiguous block of the input (only leading + // dimensions are trimmed and all steps are 1), we can copy the block directly with a single + // device-to-device memcpy and skip the per-element slice kernel entirely. + size_t offset_in_elements = 0; + size_t output_elements = 0; + size_t input_elements = 0; + if (TryComputeContiguousSliceOffset(sliced_input_dims, sliced_output_dims, + starts_buffer, steps_buffer, offset_in_elements)) { + if (TryConvertInt64ToSizeT(output_size, output_elements) && + TryConvertInt64ToSizeT(input_tensor->Shape().Size(), input_elements) && + offset_in_elements <= input_elements && + output_elements <= input_elements - offset_in_elements) { + size_t byte_offset = 0; + size_t copy_size = 0; + if (TryMultiply(offset_in_elements, element_size, byte_offset) && + TryMultiply(output_elements, element_size, copy_size)) { + const char* input_data = static_cast(input_tensor->DataRaw()) + byte_offset; + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(output_tensor->MutableDataRaw(), input_data, + copy_size, + cudaMemcpyDeviceToDevice, Stream(ctx))); + return Status::OK(); + } + } + } + return SliceImpCore(Stream(ctx), input_tensor->DataRaw(), output_tensor->MutableDataRaw(), diff --git a/onnxruntime/core/providers/cuda/tensor/slice.h b/onnxruntime/core/providers/cuda/tensor/slice.h index 050206f47217b..67eee43be182a 100644 --- a/onnxruntime/core/providers/cuda/tensor/slice.h +++ b/onnxruntime/core/providers/cuda/tensor/slice.h @@ -35,7 +35,10 @@ class Slice : public CudaKernel, public SliceBase { TensorShapeVector& input_ends, TensorShapeVector& input_axes, TensorShapeVector& input_steps) const; - virtual Status CallSliceImp(size_t element_size, size_t dimension_count, const TArray& starts_buffer, + virtual Status CallSliceImp(size_t element_size, size_t dimension_count, + gsl::span sliced_input_dims, + gsl::span sliced_output_dims, + const TArray& starts_buffer, const TArray& steps_buffer, const TArray& input_strides, const TArray& output_strides, OpKernelContext* ctx, const TensorShape& output_shape) const; diff --git a/onnxruntime/test/providers/cpu/tensor/slice_op.test.cc b/onnxruntime/test/providers/cpu/tensor/slice_op.test.cc index f34018d11dca4..e18190c842ce7 100644 --- a/onnxruntime/test/providers/cpu/tensor/slice_op.test.cc +++ b/onnxruntime/test/providers/cpu/tensor/slice_op.test.cc @@ -265,6 +265,54 @@ TEST(SliceTest, Slice3D) { 332.0f, 333.0f}); } +// Contiguous leading-axis slices exercise the CUDA EP single-memcpy fast path +// (only leading dimensions are trimmed, all trailing dimensions fully included). +TEST(SliceTest, Slice3D_ContiguousLeadingAxis) { + RunSliceTest({3, 2, 2}, + {111.0f, 112.0f, + 121.0f, 122.0f, + + 211.0f, 212.0f, + 221.0f, 222.0f, + + 311.0f, 312.0f, + 321.0f, 322.0f}, + {1}, + {3}, + {0}, + {}, + {2, 2, 2}, + {211.0f, 212.0f, + 221.0f, 222.0f, + + 311.0f, 312.0f, + 321.0f, 322.0f}); +} + +// Selecting a single leading index plus a partial second axis with full trailing +// dims is still contiguous (input_ptr + offset). +TEST(SliceTest, Slice3D_ContiguousSingleLeadingIndex) { + RunSliceTest({3, 3, 2}, + {111.0f, 112.0f, + 121.0f, 122.0f, + 131.0f, 132.0f, + + 211.0f, 212.0f, + 221.0f, 222.0f, + 231.0f, 232.0f, + + 311.0f, 312.0f, + 321.0f, 322.0f, + 331.0f, 332.0f}, + {1, 1}, + {2, 3}, + {0, 1}, + {}, + {1, 2, 2}, + {221.0f, 222.0f, + 231.0f, 232.0f}); +} + template static void TestSlice1DIntData() { // static_assert(std::is_integral_v); diff --git a/orttraining/orttraining/training_ops/cuda/tensor/slice_grad.cc b/orttraining/orttraining/training_ops/cuda/tensor/slice_grad.cc index 94ced5ee85b6d..5d1dceaff23d2 100644 --- a/orttraining/orttraining/training_ops/cuda/tensor/slice_grad.cc +++ b/orttraining/orttraining/training_ops/cuda/tensor/slice_grad.cc @@ -47,7 +47,10 @@ Status SliceGrad::FillInputVectors(OpKernelContext* ctx, TensorShapeVector& inpu ctx->Input(5), input_starts, input_ends, input_axes, input_steps); } -Status SliceGrad::CallSliceImp(size_t element_size, size_t dimension_count, const TArray& starts_buffer, +Status SliceGrad::CallSliceImp(size_t element_size, size_t dimension_count, + gsl::span /*sliced_input_dims*/, + gsl::span /*sliced_output_dims*/, + const TArray& starts_buffer, const TArray& steps_buffer, const TArray& input_strides, const TArray& output_strides, OpKernelContext* ctx, const TensorShape& output_shape) const { diff --git a/orttraining/orttraining/training_ops/cuda/tensor/slice_grad.h b/orttraining/orttraining/training_ops/cuda/tensor/slice_grad.h index cbf500ec23870..ead90d8ce492c 100644 --- a/orttraining/orttraining/training_ops/cuda/tensor/slice_grad.h +++ b/orttraining/orttraining/training_ops/cuda/tensor/slice_grad.h @@ -15,7 +15,10 @@ class SliceGrad final : public Slice { Status FillInputVectors(OpKernelContext* ctx, TensorShapeVector& input_starts, TensorShapeVector& input_ends, TensorShapeVector& input_axes, TensorShapeVector& input_steps) const override; - Status CallSliceImp(size_t element_size, size_t dimension_count, const TArray& starts_buffer, + Status CallSliceImp(size_t element_size, size_t dimension_count, + gsl::span sliced_input_dims, + gsl::span sliced_output_dims, + const TArray& starts_buffer, const TArray& steps_buffer, const TArray& input_strides, const TArray& output_strides, OpKernelContext* ctx, const TensorShape& output_shape) const override; From 2ea96f75dac830a1c1c89103e713f7b89362bd8f Mon Sep 17 00:00:00 2001 From: Tianlei Wu Date: Thu, 3 Sep 2026 13:12:49 -0700 Subject: [PATCH 13/13] [CUDA] Fix PagedAttention build without FlashAttention (#32327) ## Description Fix the CUDA PagedAttention build when FlashAttention is disabled. Two values used only by the FlashAttention path otherwise trigger `-Werror=unused-variable` and `-Werror=unused-but-set-variable`. ## Summary of Changes - Define `flash_block_size_ok` only when `USE_FLASH_ATTENTION` is enabled. - Mark `max_kv_len_lower_bound` unused in builds without FlashAttention. - Leave runtime behavior unchanged for both configurations. ## Testing - Built `onnxruntime_providers_cuda` with `onnxruntime_USE_FLASH_ATTENTION=OFF` and `-Werror`. - Compiled the rebased `paged_attention.cc` directly with the same CUDA Debug configuration. ## Checklist - [x] Tests not required for this build-only fix - [x] No breaking changes - [x] Documentation not required --- onnxruntime/contrib_ops/cuda/bert/paged_attention.cc | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/onnxruntime/contrib_ops/cuda/bert/paged_attention.cc b/onnxruntime/contrib_ops/cuda/bert/paged_attention.cc index 75d5a0a73e3c1..18ae17272fe14 100644 --- a/onnxruntime/contrib_ops/cuda/bert/paged_attention.cc +++ b/onnxruntime/contrib_ops/cuda/bert/paged_attention.cc @@ -299,7 +299,6 @@ Status PagedAttention::ComputeInternal(OpKernelContext* context) cons // which has no page-alignment requirement. const int flash_min_block_size = parameters.head_size <= 64 ? 256 : (parameters.head_size <= 128 ? 128 : 64); - const bool flash_block_size_ok = kIsQuantizedCache || (parameters.block_size % flash_min_block_size) == 0; // LATENT (absorbed MLA) has exactly one eligible backend: neither FlashAttention nor the CUTLASS // fMHA wrapper supports v_head_size != head_size or a head_size of 576, and the paged decode @@ -317,6 +316,7 @@ Status PagedAttention::ComputeInternal(OpKernelContext* context) cons } #if USE_FLASH_ATTENTION + const bool flash_block_size_ok = kIsQuantizedCache || (parameters.block_size % flash_min_block_size) == 0; const bool flash_eligible = !use_latent_attention && !disable_flash_attention_ && flash_block_size_ok && @@ -651,6 +651,8 @@ Status PagedAttention::ComputeInternal(OpKernelContext* context) cons GetScratchBuffer(out_accum_bytes, GetComputeStream(context)); } } +#else + ORT_UNUSED_PARAMETER(max_kv_len_lower_bound); #endif if (needs_dense_kv) {