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/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_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/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/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); 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/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) { 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/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); 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/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. 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/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); } 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) { 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 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"); 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/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; +} 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; 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 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: 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