From 830f96a242c6d06e8453960ac14030bae0be450b Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Mon, 5 Oct 2026 14:06:08 -0700 Subject: [PATCH 01/10] Prepare ONNX 1.23.1 reintegration Restore the reverted ONNX 1.23 integration and Float6 adapter mappings on current main. Align source, archive, vcpkg, CI and materialization pins with 1.23.1 and retain the released-opset 28 metadata workaround and regressions. Preserve subsequent free-threaded Python packaging support, repair vcpkg protobuf patch context, and derive equivalence-test version parity from the archive pin. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- .github/workflows/linux_ci.yml | 2 +- cmake/deps.txt | 7 +- cmake/external/onnx | 2 +- cmake/onnxruntime_unittests.cmake | 36 +-- cmake/patches/onnx/onnx.patch | 99 ++++++-- cmake/vcpkg-ports/onnx/binskim.patch | 99 ++++++-- .../onnx/fix-dependency-protobuf.patch | 6 +- cmake/vcpkg-ports/onnx/portfile.cmake | 4 +- cmake/vcpkg-ports/onnx/vcpkg.json | 2 +- .../InferenceTest.netcore.cs | 3 - docs/OperatorKernels.md | 11 +- docs/python/examples/plot_pipeline.py | 33 ++- .../onnxruntime/core/graph/schema_registry.h | 2 +- .../core/session/onnxruntime_c_api.h | 5 + .../onnxruntime/ep/adapter/tensor_helper.h | 4 + js/web/docs/webgl-operators.md | 7 +- .../framework/onnxruntime_map_type_info.h | 4 +- .../core/framework/tensor_type_and_shape.cc | 6 + .../core/framework/tensorprotoutils.cc | 2 + onnxruntime/core/framework/tensorprotoutils.h | 2 + .../core/optimizer/attention_fusion_helper.h | 2 +- .../core/optimizer/embed_layer_norm_fusion.cc | 12 +- .../core/optimizer/fast_gelu_fusion.cc | 4 +- .../core/optimizer/isinf_reducesum_fusion.cc | 4 +- .../core/optimizer/layer_norm_fusion.cc | 14 +- onnxruntime/core/optimizer/pad_fusion.cc | 2 +- .../transpose_optimization/optimizer_api.h | 2 +- .../coreml/builders/impl/base_op_builder.h | 2 +- .../providers/cpu/cpu_execution_provider.cc | 55 ++++- .../providers/cpu/math/element_wise_ops.cc | 141 ++++++++++-- .../core/providers/cuda/cu_inc/common.cuh | 34 +++ .../providers/cuda/cuda_execution_provider.cc | 6 +- .../cuda/math/binary_elementwise_ops.cc | 17 +- .../cuda/math/binary_elementwise_ops_impl.cu | 2 +- .../src/AbiCustomRegistry.cpp | 10 + .../builders/impl/base_op_builder.h | 2 +- .../impl/dequantizelinear_op_builder.cc | 2 + .../impl/quantizelinear_op_builder.cc | 2 + .../vsinpu/builders/impl/base_op_builder.h | 2 +- .../core/providers/webgpu/compute_context.h | 6 +- .../webnn/builders/impl/base_op_builder.h | 2 +- .../webnn/builders/impl/qdq_op_builder.cc | 2 + onnxruntime/core/session/inference_session.cc | 17 +- onnxruntime/python/backend/backend.py | 3 + onnxruntime/python/backend/backend_rep.py | 9 +- .../test/framework/tensorutils_test.cc | 2 + onnxruntime/test/framework/type_info_test.cc | 10 + onnxruntime/test/onnx/TestCase.cc | 4 +- onnxruntime/test/onnx/main.cc | 3 + onnxruntime/test/onnx/tensorprotoutils.cc | 2 + ...ansformation_potentially_added_ops_test.cc | 12 +- .../test/optimizer/qdq_transformer_test.cc | 5 +- .../providers/cpu/controlflow/scan_test.cc | 5 +- .../cpu/math/element_wise_ops_test.cc | 132 +++++++++++ .../providers/cpu/tensor/compress_op.test.cc | 12 +- .../providers/dml_onnx_tensor_wrapper_test.cc | 33 +-- .../test/python/onnx_backend_test_series.py | 6 + .../python/onnx_node_test_equivalence_test.py | 23 +- .../python/onnxruntime_test_python_backend.py | 118 ++++++++++ onnxruntime/test/python/requirements.txt | 2 +- .../transformers/test_cuda_plugin_ep.py | 4 +- .../test/python/transformers/test_moe_cuda.py | 8 +- .../onnx_backend_test_series_filters.jsonc | 215 ++++++++++++++++++ .../onnx_backend_test_series_overrides.jsonc | 5 +- .../graph_transform_test_builder.cc | 7 +- onnxruntime/test/unittest_util/op_tester.cc | 6 + requirements-materialize-onnx-node-tests.txt | 2 +- tools/ci_build/build.py | 15 +- .../linux/build_linux_python_package.sh | 2 +- .../python/cpu/scripts/requirements.txt | 2 +- .../docker/scripts/lort/requirements.txt | 2 +- .../docker/scripts/manylinux/requirements.txt | 2 +- .../linux/docker/scripts/requirements.txt | 2 +- .../github/linux/python/requirements.txt | 2 +- .../github/windows/python/requirements.txt | 2 +- tools/ci_build/test_build_args.py | 29 ++- tools/python/materialize_onnx_node_tests.py | 36 +-- 77 files changed, 1153 insertions(+), 247 deletions(-) diff --git a/.github/workflows/linux_ci.yml b/.github/workflows/linux_ci.yml index ad2b35e5e96a6..b15a0a92b9e82 100644 --- a/.github/workflows/linux_ci.yml +++ b/.github/workflows/linux_ci.yml @@ -79,7 +79,7 @@ jobs: dockerfile_path: tools/ci_build/github/linux/docker/Dockerfile.manylinux2_28_cpu docker_image_repo: onnxruntimecpubuildpythonx64 # MATERIALIZE=ON needs no explicit pre-build pip install on this leg: the manylinux cp314 - # image already ships onnx==1.22.0 + numpy==2.4.2 (the exact gate pins) via its + # image already ships the pinned onnx + numpy versions via its # requirements.txt (install_deps.sh installs it into cp311-cp314). Intentional, not an # omission -- unlike the QNN/android legs which run on bare agents and pip-install the pins. extra_build_flags: '--use_binskim_compliant_compile_flags --build_wheel --build_nuget --enable_transformers_tool_test --cmake_extra_defines onnxruntime_BUILD_BENCHMARKS=ON --cmake_extra_defines onnxruntime_MATERIALIZE_ONNX_NODE_TESTS=ON' diff --git a/cmake/deps.txt b/cmake/deps.txt index 2bec4eafdc4d1..17f39cef00f54 100644 --- a/cmake/deps.txt +++ b/cmake/deps.txt @@ -34,10 +34,9 @@ microsoft_gsl;https://github.com/microsoft/GSL/archive/refs/tags/v4.2.1.zip;1094 microsoft_wil;https://github.com/microsoft/wil/archive/refs/tags/v1.0.250325.1.zip;826c8bd47c2258ec61b8b218e031e5b33d27f761 mimalloc;https://github.com/microsoft/mimalloc/archive/refs/tags/v2.1.1.zip;d5ee7d34223d0567892db5179849939c8769dc41 mp11;https://github.com/boostorg/mp11/archive/refs/tags/boost-1.82.0.zip;9bc9e01dffb64d9e0773b2e44d2f22c51aace063 -# ONNX pinned to the 1.22.0 final release (tag v1.22.0). rc2->final changed only CI -# workflow yml + VERSION_NUMBER (no operator-schema, opset, or testdata change), so it is -# pure version plumbing. The 7 CI requirements.txt files are pinned to onnx==1.22.0 (now on PyPI). -onnx;https://github.com/onnx/onnx/archive/refs/tags/v1.22.0.zip;2b2cd58ac7a26df5371266149e0c76776330cdf1 +# ONNX 1.23.1 final release (opset 28). +# The 7 CI requirements.txt files are pinned to onnx==1.23.1. +onnx;https://github.com/onnx/onnx/archive/refs/tags/v1.23.1.zip;85060f3c52267c6185a5c1fa298083a514f5d201 # Use the latest commit of 10.9-GA onnx_tensorrt;https://github.com/onnx/onnx-tensorrt/archive/d5dce67db7c2e64b07e055571f5ec06f7f254de2.zip;01114d3b67650857281fa50faa2e412130a63b69 protobuf;https://github.com/protocolbuffers/protobuf/releases/download/v33.6/protobuf-33.6.tar.gz;051fee744fcca04dab37eb009d49aa3e9546acda diff --git a/cmake/external/onnx b/cmake/external/onnx index 2bb50465112fe..ecfb5f08b88b2 160000 --- a/cmake/external/onnx +++ b/cmake/external/onnx @@ -1 +1 @@ -Subproject commit 2bb50465112feca9003e1ed654d77f01ff1415ca +Subproject commit ecfb5f08b88b2c60733323bb8a691e99ba11e469 diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index dbdb0030cc933..86b11d46f28c3 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1813,13 +1813,13 @@ if (NOT onnxruntime_ENABLE_TRAINING_TORCH_INTEROP) # coverage gap this feature exists to close. # --------------------------------------------------------------------------- if (onnxruntime_MATERIALIZE_ONNX_NODE_TESTS AND NOT CMAKE_SYSTEM_NAME STREQUAL "Emscripten") - # onnx version pin: derive from the archive URL in cmake/deps.txt (single source of truth). - string(REGEX MATCH "v([0-9]+\\.[0-9]+\\.[0-9]+)" _onnx_url_ver "${DEP_URL_onnx}") + # Derive the expected wheel version from the archive URL so this also works with installed ONNX packages. + string(REGEX MATCH "v([0-9]+\\.[0-9]+\\.[0-9]+)" _onnx_url_version "${DEP_URL_onnx}") if(CMAKE_MATCH_1) set(_onnx_pinned_version ${CMAKE_MATCH_1}) else() - message(FATAL_ERROR "Could not parse the pinned onnx version from DEP_URL_onnx='${DEP_URL_onnx}' " - "(expected a vX.Y.Z tag). Fix cmake/deps.txt or this regex.") + message(FATAL_ERROR "Could not parse the pinned ONNX version from DEP_URL_onnx='${DEP_URL_onnx}' " + "(expected a vX.Y.Z tag). Fix cmake/deps.txt or this parser.") endif() # Python interpreter is not guaranteed for static test-only builds (the top-level @@ -1847,27 +1847,10 @@ if (NOT onnxruntime_ENABLE_TRAINING_TORCH_INTEROP) " OR reconfigure with -Donnxruntime_MATERIALIZE_ONNX_NODE_TESTS=OFF (node-test coverage will be dropped).\n" " Details: ${_onnx_err}") endif() - # onnx version gate: HARD FAIL on a genuine mismatch, but RC / pre-release AWARE. - # ONNX's opset-bump workflow ships wheels like 1.23.0rc1 or 1.23.0.dev20240101 whose - # COMPILED opset registry already matches the formal 1.23.0 tag, so we compare on the - # RELEASE BASE (major.minor.micro) rather than the raw string. This mirrors - # materialize_onnx_node_tests.py::_release_base EXACTLY (regex ^(\d+)\.(\d+)\.(\d+), with a - # raw-string fallback when there is no leading X.Y.Z) so the cmake and Python layers agree: - # an rcN/.devN wheel of the pinned tag passes, while a real major/minor/micro mismatch - # (e.g. 1.21.x, or 1.23.0 when pinned at 1.22.0) still FATALs. Both sides are normalized; - # _onnx_pinned_version is already a clean X.Y.Z (parsed from the deps.txt vX.Y.Z tag), so - # normalizing it is a no-op kept only for symmetry with the Python two-sided compare. - string(REGEX MATCH "^[0-9]+\\.[0-9]+\\.[0-9]+" _onnx_ver_base "${_onnx_ver}") - if(_onnx_ver_base STREQUAL "") - set(_onnx_ver_base "${_onnx_ver}") - endif() - string(REGEX MATCH "^[0-9]+\\.[0-9]+\\.[0-9]+" _onnx_pin_base "${_onnx_pinned_version}") - if(_onnx_pin_base STREQUAL "") - set(_onnx_pin_base "${_onnx_pinned_version}") - endif() - if(NOT _onnx_ver_base STREQUAL _onnx_pin_base) + # The source and wheel are both final releases, so require exact version parity. + if(NOT _onnx_ver STREQUAL _onnx_pinned_version) message(FATAL_ERROR - "onnx ${_onnx_ver} (release base ${_onnx_ver_base}) != pinned ${_onnx_pinned_version} " + "onnx ${_onnx_ver} != pinned ${_onnx_pinned_version} " "(cmake/deps.txt). A mismatched wheel bakes the wrong opset/IR into the materialized " "corpus (silent drift).\n" " Fix: pip install onnx==${_onnx_pinned_version}") @@ -1945,6 +1928,11 @@ if (NOT onnxruntime_ENABLE_TRAINING_TORCH_INTEROP) VERBATIM) add_custom_target(onnx_node_tests_materialized ALL DEPENDS ${_materialized_node_root}/.stamp) + # Keep the corpus available for targeted onnx_test_runner builds as well as + # normal ALL builds. The runner is the direct C++ consumer of this artifact. + if(TARGET onnx_test_runner) + add_dependencies(onnx_test_runner onnx_node_tests_materialized) + endif() if (NOT onnxruntime_REDUCED_OPS_BUILD) # First-class ctest over the materialized node corpus (the durable replacement for the diff --git a/cmake/patches/onnx/onnx.patch b/cmake/patches/onnx/onnx.patch index 8b0068b7b0c80..d4037eee8a0af 100644 --- a/cmake/patches/onnx/onnx.patch +++ b/cmake/patches/onnx/onnx.patch @@ -1,8 +1,8 @@ diff --git a/CMakeLists.txt b/CMakeLists.txt -index 791e24b..5f861a1 100644 +index 6f1d56949..7e6331dec 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt -@@ -46,6 +46,7 @@ option(ONNX_DISABLE_EXCEPTIONS "Disable exception handling." OFF) +@@ -53,6 +53,7 @@ option(ONNX_DISABLE_EXCEPTIONS "Disable exception handling." OFF) option(ONNX_DISABLE_STATIC_REGISTRATION "Disable static registration for ONNX operator schemas." OFF) option(ONNX_USE_UNITY_BUILD "Enable Unity (Jumbo) build for" OFF) option(ONNX_PYODIDE_BUILD "Build ONNX for Pyodide/Emscripten (wasm32)." OFF) @@ -10,7 +10,7 @@ index 791e24b..5f861a1 100644 option(ONNX_INSTALL "Install ONNX targets, headers, and CMake config files" ON) if(WIN32) option(ONNX_USE_MSVC_STATIC_RUNTIME "Build with MSVC static runtime. Ignored if CMAKE_MSVC_RUNTIME_LIBRARY is defined." OFF) -@@ -515,7 +516,22 @@ add_dependencies(onnx_proto onnx_proto_gen) +@@ -562,7 +563,23 @@ add_dependencies(onnx_proto onnx_proto_headers) # Core library. Hand-written sources come from target_sources(onnx ...) in the # subdirectories; the protobuf objects are pulled in from onnx_proto. add_library(onnx $) @@ -22,6 +22,7 @@ index 791e24b..5f861a1 100644 + # the full operator/schema sources via add_subdirectory(onnx). + target_sources(onnx PRIVATE + "${ONNX_ROOT}/onnx/defs/data_type_utils.cc" ++ "${ONNX_ROOT}/onnx/defs/doc_strings.cc" + ) + # Ensure ONNX_ML is treated as ON for minimal build consistency with ORT's file + set(ONNX_ML ON CACHE BOOL "Enable traditional ML API." FORCE) @@ -34,11 +35,37 @@ index 791e24b..5f861a1 100644 add_dependencies(onnx onnx_proto) set_target_properties(onnx PROPERTIES CXX_VISIBILITY_PRESET hidden) add_onnx_global_defines(onnx) +diff --git a/onnx/defs/doc_strings.cc b/onnx/defs/doc_strings.cc +index 976280392..8c22e55ff 100644 +--- a/onnx/defs/doc_strings.cc ++++ b/onnx/defs/doc_strings.cc +@@ -1705,6 +1705,8 @@ const char kDoc_DequantizeLinear_ver24[] = ""; + const char kDoc_RandomNormal_ver1[] = ""; + const char kDoc_Round_ver11[] = ""; + const char kDoc_SpaceToDepth_ver1[] = ""; ++const char kDoc_SpaceToDepth_ver28[] = ""; ++const char kDoc_DepthToSpace_ver28[] = ""; + const char kDoc_InstanceNormalization_ver6[] = ""; + const char kDoc_ThresholdedRelu_ver10[] = ""; + const char kDoc_Acosh_ver9[] = ""; +diff --git a/onnx/defs/nn/defs.cc b/onnx/defs/nn/defs.cc +index 4b58256e6..1a8103e7c 100644 +--- a/onnx/defs/nn/defs.cc ++++ b/onnx/defs/nn/defs.cc +@@ -2692,7 +2692,7 @@ ONNX_OPERATOR_SET_SCHEMA( + // positive value. + axis += input_ndim; + } +- if (axis < 0) { ++ if (axis < 0 || axis >= input_ndim) { + fail_shape_inference( + "Unexpected axis value (", + axis, diff --git a/onnx/defs/nn/old.cc b/onnx/defs/nn/old.cc -index 3887169..8432613 100644 +index cbb1532c0..6597a2ef9 100644 --- a/onnx/defs/nn/old.cc +++ b/onnx/defs/nn/old.cc -@@ -4078,7 +4078,6 @@ ONNX_OPERATOR_SET_SCHEMA( +@@ -4118,7 +4118,6 @@ ONNX_OPERATOR_SET_SCHEMA( GroupNormalization, 18, OpSchema() @@ -46,16 +73,52 @@ index 3887169..8432613 100644 .SetDoc(GroupNormalization_ver18_doc) .Attr("epsilon", "The epsilon value to use to avoid division by zero.", AttributeProto::FLOAT, 1e-5f) .Attr( -diff --git a/onnx/defs/nn/defs.cc b/onnx/defs/nn/defs.cc -index 0c271d5..6d64da7 100644 ---- a/onnx/defs/nn/defs.cc -+++ b/onnx/defs/nn/defs.cc -@@ -2670,7 +2670,7 @@ ONNX_OPERATOR_SET_SCHEMA( - // positive value. - axis += input_ndim; - } -- if (axis < 0) { -+ if (axis < 0 || axis >= input_ndim) { - fail_shape_inference( - "Unexpected axis value (", - axis, +diff --git a/onnx/defs/schema.h b/onnx/defs/schema.h +--- a/onnx/defs/schema.h ++++ b/onnx/defs/schema.h +@@ -937,7 +937,7 @@ + // Version corresponding last release of ONNX. Update this to match with + // the max version above in a *release* version of ONNX. But in other + // versions, the max version may be ahead of the last-release-version. +- last_release_version_map_[ONNX_DOMAIN] = 27; ++ last_release_version_map_[ONNX_DOMAIN] = 28; + last_release_version_map_[AI_ONNX_ML_DOMAIN] = 5; + last_release_version_map_[AI_ONNX_TRAINING_DOMAIN] = 1; + last_release_version_map_[AI_ONNX_PREVIEW_DOMAIN] = 1; +diff --git a/onnx/defs/printer.cc b/onnx/defs/printer.cc +index 1d43bccb0..b6beed6f3 100644 +--- a/onnx/defs/printer.cc ++++ b/onnx/defs/printer.cc +@@ -9,6 +9,9 @@ + #include + #include + #include ++#include ++#include ++#include + #include + #include + #include +@@ -31,6 +34,14 @@ std::vector ParseRawIntData(const TensorProto& tensor) { + // point, the shortest form that round-trips through the parser. + template + void PrintNumber(std::ostream& os, T value) { ++#if defined(__APPLE__) && defined(__ENVIRONMENT_IPHONE_OS_VERSION_MIN_REQUIRED__) && \ ++ __ENVIRONMENT_IPHONE_OS_VERSION_MIN_REQUIRED__ < 160300 ++ // std::to_chars is unavailable when targeting iOS releases before 16.3. ++ std::ostringstream value_stream; ++ value_stream.imbue(std::locale::classic()); ++ value_stream << std::setprecision(std::numeric_limits::max_digits10) << value; ++ os << value_stream.str(); ++#else + if constexpr (std::is_floating_point_v) { + // Normalize NaN: MSVC's to_chars emits payload forms like "nan(ind)" + // that the parser does not accept. +@@ -42,6 +53,7 @@ void PrintNumber(std::ostream& os, T value) { + std::array buf{}; + const auto res = std::to_chars(buf.data(), buf.data() + buf.size(), value); + os.write(buf.data(), res.ptr - buf.data()); ++#endif + } + + bool IsValidIdentifier(const std::string& str) { diff --git a/cmake/vcpkg-ports/onnx/binskim.patch b/cmake/vcpkg-ports/onnx/binskim.patch index 8b0068b7b0c80..d4037eee8a0af 100644 --- a/cmake/vcpkg-ports/onnx/binskim.patch +++ b/cmake/vcpkg-ports/onnx/binskim.patch @@ -1,8 +1,8 @@ diff --git a/CMakeLists.txt b/CMakeLists.txt -index 791e24b..5f861a1 100644 +index 6f1d56949..7e6331dec 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt -@@ -46,6 +46,7 @@ option(ONNX_DISABLE_EXCEPTIONS "Disable exception handling." OFF) +@@ -53,6 +53,7 @@ option(ONNX_DISABLE_EXCEPTIONS "Disable exception handling." OFF) option(ONNX_DISABLE_STATIC_REGISTRATION "Disable static registration for ONNX operator schemas." OFF) option(ONNX_USE_UNITY_BUILD "Enable Unity (Jumbo) build for" OFF) option(ONNX_PYODIDE_BUILD "Build ONNX for Pyodide/Emscripten (wasm32)." OFF) @@ -10,7 +10,7 @@ index 791e24b..5f861a1 100644 option(ONNX_INSTALL "Install ONNX targets, headers, and CMake config files" ON) if(WIN32) option(ONNX_USE_MSVC_STATIC_RUNTIME "Build with MSVC static runtime. Ignored if CMAKE_MSVC_RUNTIME_LIBRARY is defined." OFF) -@@ -515,7 +516,22 @@ add_dependencies(onnx_proto onnx_proto_gen) +@@ -562,7 +563,23 @@ add_dependencies(onnx_proto onnx_proto_headers) # Core library. Hand-written sources come from target_sources(onnx ...) in the # subdirectories; the protobuf objects are pulled in from onnx_proto. add_library(onnx $) @@ -22,6 +22,7 @@ index 791e24b..5f861a1 100644 + # the full operator/schema sources via add_subdirectory(onnx). + target_sources(onnx PRIVATE + "${ONNX_ROOT}/onnx/defs/data_type_utils.cc" ++ "${ONNX_ROOT}/onnx/defs/doc_strings.cc" + ) + # Ensure ONNX_ML is treated as ON for minimal build consistency with ORT's file + set(ONNX_ML ON CACHE BOOL "Enable traditional ML API." FORCE) @@ -34,11 +35,37 @@ index 791e24b..5f861a1 100644 add_dependencies(onnx onnx_proto) set_target_properties(onnx PROPERTIES CXX_VISIBILITY_PRESET hidden) add_onnx_global_defines(onnx) +diff --git a/onnx/defs/doc_strings.cc b/onnx/defs/doc_strings.cc +index 976280392..8c22e55ff 100644 +--- a/onnx/defs/doc_strings.cc ++++ b/onnx/defs/doc_strings.cc +@@ -1705,6 +1705,8 @@ const char kDoc_DequantizeLinear_ver24[] = ""; + const char kDoc_RandomNormal_ver1[] = ""; + const char kDoc_Round_ver11[] = ""; + const char kDoc_SpaceToDepth_ver1[] = ""; ++const char kDoc_SpaceToDepth_ver28[] = ""; ++const char kDoc_DepthToSpace_ver28[] = ""; + const char kDoc_InstanceNormalization_ver6[] = ""; + const char kDoc_ThresholdedRelu_ver10[] = ""; + const char kDoc_Acosh_ver9[] = ""; +diff --git a/onnx/defs/nn/defs.cc b/onnx/defs/nn/defs.cc +index 4b58256e6..1a8103e7c 100644 +--- a/onnx/defs/nn/defs.cc ++++ b/onnx/defs/nn/defs.cc +@@ -2692,7 +2692,7 @@ ONNX_OPERATOR_SET_SCHEMA( + // positive value. + axis += input_ndim; + } +- if (axis < 0) { ++ if (axis < 0 || axis >= input_ndim) { + fail_shape_inference( + "Unexpected axis value (", + axis, diff --git a/onnx/defs/nn/old.cc b/onnx/defs/nn/old.cc -index 3887169..8432613 100644 +index cbb1532c0..6597a2ef9 100644 --- a/onnx/defs/nn/old.cc +++ b/onnx/defs/nn/old.cc -@@ -4078,7 +4078,6 @@ ONNX_OPERATOR_SET_SCHEMA( +@@ -4118,7 +4118,6 @@ ONNX_OPERATOR_SET_SCHEMA( GroupNormalization, 18, OpSchema() @@ -46,16 +73,52 @@ index 3887169..8432613 100644 .SetDoc(GroupNormalization_ver18_doc) .Attr("epsilon", "The epsilon value to use to avoid division by zero.", AttributeProto::FLOAT, 1e-5f) .Attr( -diff --git a/onnx/defs/nn/defs.cc b/onnx/defs/nn/defs.cc -index 0c271d5..6d64da7 100644 ---- a/onnx/defs/nn/defs.cc -+++ b/onnx/defs/nn/defs.cc -@@ -2670,7 +2670,7 @@ ONNX_OPERATOR_SET_SCHEMA( - // positive value. - axis += input_ndim; - } -- if (axis < 0) { -+ if (axis < 0 || axis >= input_ndim) { - fail_shape_inference( - "Unexpected axis value (", - axis, +diff --git a/onnx/defs/schema.h b/onnx/defs/schema.h +--- a/onnx/defs/schema.h ++++ b/onnx/defs/schema.h +@@ -937,7 +937,7 @@ + // Version corresponding last release of ONNX. Update this to match with + // the max version above in a *release* version of ONNX. But in other + // versions, the max version may be ahead of the last-release-version. +- last_release_version_map_[ONNX_DOMAIN] = 27; ++ last_release_version_map_[ONNX_DOMAIN] = 28; + last_release_version_map_[AI_ONNX_ML_DOMAIN] = 5; + last_release_version_map_[AI_ONNX_TRAINING_DOMAIN] = 1; + last_release_version_map_[AI_ONNX_PREVIEW_DOMAIN] = 1; +diff --git a/onnx/defs/printer.cc b/onnx/defs/printer.cc +index 1d43bccb0..b6beed6f3 100644 +--- a/onnx/defs/printer.cc ++++ b/onnx/defs/printer.cc +@@ -9,6 +9,9 @@ + #include + #include + #include ++#include ++#include ++#include + #include + #include + #include +@@ -31,6 +34,14 @@ std::vector ParseRawIntData(const TensorProto& tensor) { + // point, the shortest form that round-trips through the parser. + template + void PrintNumber(std::ostream& os, T value) { ++#if defined(__APPLE__) && defined(__ENVIRONMENT_IPHONE_OS_VERSION_MIN_REQUIRED__) && \ ++ __ENVIRONMENT_IPHONE_OS_VERSION_MIN_REQUIRED__ < 160300 ++ // std::to_chars is unavailable when targeting iOS releases before 16.3. ++ std::ostringstream value_stream; ++ value_stream.imbue(std::locale::classic()); ++ value_stream << std::setprecision(std::numeric_limits::max_digits10) << value; ++ os << value_stream.str(); ++#else + if constexpr (std::is_floating_point_v) { + // Normalize NaN: MSVC's to_chars emits payload forms like "nan(ind)" + // that the parser does not accept. +@@ -42,6 +53,7 @@ void PrintNumber(std::ostream& os, T value) { + std::array buf{}; + const auto res = std::to_chars(buf.data(), buf.data() + buf.size(), value); + os.write(buf.data(), res.ptr - buf.data()); ++#endif + } + + bool IsValidIdentifier(const std::string& str) { diff --git a/cmake/vcpkg-ports/onnx/fix-dependency-protobuf.patch b/cmake/vcpkg-ports/onnx/fix-dependency-protobuf.patch index 4a438def6cf1b..12e1a7a662a25 100644 --- a/cmake/vcpkg-ports/onnx/fix-dependency-protobuf.patch +++ b/cmake/vcpkg-ports/onnx/fix-dependency-protobuf.patch @@ -2,14 +2,14 @@ diff --git a/CMakeLists.txt b/CMakeLists.txt index 4a7df7d71..257d77e61 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt -@@ -159,6 +159,7 @@ if(NOT DEFINED CMAKE_POSITION_INDEPENDENT_CODE) +@@ -198,6 +198,7 @@ if(NOT DEFINED CMAKE_POSITION_INDEPENDENT_CODE) set(CMAKE_POSITION_INDEPENDENT_CODE ON) endif() +find_package(protobuf CONFIG REQUIRED) list(APPEND CMAKE_MODULE_PATH ${ONNX_ROOT}/cmake/external) - if(NOT ONNX_BUILD_CUSTOM_PROTOBUF) - if((ONNX_USE_LITE_PROTO AND TARGET protobuf::libprotobuf-lite) OR ((NOT ONNX_USE_LITE_PROTO) AND TARGET protobuf::libprotobuf)) + if(ONNX_USE_ASAN OR ONNX_USE_TSAN OR ONNX_USE_TYPESAN) + find_package(Sanitizer REQUIRED) diff --git a/cmake/ONNXConfig.cmake.in b/cmake/ONNXConfig.cmake.in index aa1ef47d9..d450b51e5 100644 --- a/cmake/ONNXConfig.cmake.in diff --git a/cmake/vcpkg-ports/onnx/portfile.cmake b/cmake/vcpkg-ports/onnx/portfile.cmake index 7a266e1a07663..df43d91ef0040 100644 --- a/cmake/vcpkg-ports/onnx/portfile.cmake +++ b/cmake/vcpkg-ports/onnx/portfile.cmake @@ -3,8 +3,8 @@ vcpkg_check_linkage(ONLY_STATIC_LIBRARY) vcpkg_from_github( OUT_SOURCE_PATH SOURCE_PATH REPO onnx/onnx - REF "v1.22.0" - SHA512 13fafff073a8e0bcf67fd06195da979c3c6273b3f5fdd25f5e5c39ad6af20eefe107f13dc73d52b6021e986d87401e46121e00bba8fbd09f098334e34f78462f + REF "v1.23.1" + SHA512 85e377b73c412d25ebdc0e8ed5951adfb3643b20d1a75ea7e7a4186ff77dc2bcee9a6dfdc844a1f43eb935325e525b816dd4ba540fe2c045e3390acc8b16fadb PATCHES fix-cmakelists.patch fix-dependency-protobuf.patch diff --git a/cmake/vcpkg-ports/onnx/vcpkg.json b/cmake/vcpkg-ports/onnx/vcpkg.json index c399d4c1ab087..c1cb50f9bec8d 100644 --- a/cmake/vcpkg-ports/onnx/vcpkg.json +++ b/cmake/vcpkg-ports/onnx/vcpkg.json @@ -1,6 +1,6 @@ { "name": "onnx", - "version-semver": "1.22.0", + "version-semver": "1.23.1", "port-version": 0, "description": "Open standard for machine learning interoperability", "homepage": "https://onnx.ai", diff --git a/csharp/test/Microsoft.ML.OnnxRuntime.Tests.NetCoreApp/InferenceTest.netcore.cs b/csharp/test/Microsoft.ML.OnnxRuntime.Tests.NetCoreApp/InferenceTest.netcore.cs index c0475bb6102c1..14d4fffd55f56 100644 --- a/csharp/test/Microsoft.ML.OnnxRuntime.Tests.NetCoreApp/InferenceTest.netcore.cs +++ b/csharp/test/Microsoft.ML.OnnxRuntime.Tests.NetCoreApp/InferenceTest.netcore.cs @@ -491,9 +491,6 @@ private static Dictionary GetSkippedModels(DirectoryInfo modelsD { "test_mul_uint8", "Could not find an implementation for Mul(14) node with name" }, - { "test_bitshift_right_uint16", "Could not find an implementation for BitShift(11) nodeMeta with name ''"}, - { "test_bitshift_left_uint16", "Could not find an implementation for BitShift(11)"}, - { "test_pow_types_float32_uint64", "Could not find an implementation for Pow(15) node with name ''"}, { "test_pow_types_float32_uint32", "Could not find an implementation for Pow(15) node with name ''"}, diff --git a/docs/OperatorKernels.md b/docs/OperatorKernels.md index 10a1b5b96072b..fda11c1dc7fe9 100644 --- a/docs/OperatorKernels.md +++ b/docs/OperatorKernels.md @@ -63,7 +63,8 @@ The **OpSet Version** column uses the following notation: |||[9, 13]|**T** = tensor(double), tensor(float)| |||[7, 8]|**T** = tensor(double), tensor(float)| |BitCast|*in* input:**T1**
*out* output:**T2**|26+|**T1** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| -|BitShift|*in* X:**T**
*in* Y:**T**
*out* Z:**T**|11+|**T** = tensor(uint32), tensor(uint64), tensor(uint8)| +|BitShift|*in* X:**T**
*in* Y:**T**
*out* Z:**T**|28+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|||[11, 27]|**T** = tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseAnd|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseNot|*in* X:**T**
*out* Y:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseOr|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| @@ -287,7 +288,8 @@ The **OpSet Version** column uses the following notation: |||12|**T** = tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(int8), tensor(uint32), tensor(uint64), tensor(uint8)| |||[8, 11]|**T** = tensor(double), tensor(float)| |||[6, 7]|**T** = tensor(float)| -|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|13+|**T** = tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|28+|**T** = tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|||[13, 27]|**T** = tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||[10, 12]|**T** = tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |Mul|*in* A:**T**
*in* B:**T**
*out* C:**T**|14+|**T** = tensor(double), tensor(float), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||13|**T** = tensor(double), tensor(float), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| @@ -865,7 +867,8 @@ The **OpSet Version** column uses the following notation: |Min|*in* data_0:**T**
*out* min:**T**|13+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| |||12|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| |||[6, 11]|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16)| -|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|13+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| +|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|28+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| +|||[13, 27]|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| |||[10, 12]|**T** = tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| |Mul|*in* A:**T**
*in* B:**T**
*out* C:**T**|14+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||13|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| @@ -1239,7 +1242,7 @@ The **OpSet Version** column uses the following notation: |||14+|**T** = tensor(float), tensor(float16)| |||9+|**T** = tensor(float), tensor(float16)| |||7+|**T** = tensor(float), tensor(float16)| -|BitShift|*in* X:**T**
*in* Y:**T**
*out* Z:**T**|11+|**T** = tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|BitShift|*in* X:**T**
*in* Y:**T**
*out* Z:**T**|[11, 27]|**T** = tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseAnd|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseNot|*in* X:**T**
*out* Y:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseOr|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| diff --git a/docs/python/examples/plot_pipeline.py b/docs/python/examples/plot_pipeline.py index e0bae57b7e2c8..c73d3ba30c211 100644 --- a/docs/python/examples/plot_pipeline.py +++ b/docs/python/examples/plot_pipeline.py @@ -29,12 +29,10 @@ ################################# -# Draw a model with ONNX +# Draw a model with pydot # ++++++++++++++++++++++ -# We use `net_drawer.py `_ -# included in *onnx* package. -# We use *onnx* to load the model -# in a different way than before. +# We use *onnx* to load the model in a different way than before, then +# construct a pydot graph from its nodes and values. from onnx import ModelProto # noqa: E402 @@ -46,18 +44,29 @@ ################################### # We convert it into a graph. -from onnx.tools.net_drawer import GetOpNodeProducer, GetPydotGraph # noqa: E402 +import pydot # noqa: E402 + +pydot_graph = pydot.Dot(model.graph.name, rankdir="LR") +value_nodes = {} +for index, node in enumerate(model.graph.node): + node_name = node.name or f"{node.op_type}_{index}" + op_node = pydot.Node(node_name, label=node.op_type, shape="box") + pydot_graph.add_node(op_node) + for value_name in node.input: + value_node = value_nodes.setdefault(value_name, pydot.Node(value_name, shape="ellipse")) + pydot_graph.add_node(value_node) + pydot_graph.add_edge(pydot.Edge(value_node, op_node)) + for value_name in node.output: + value_node = pydot.Node(value_name, shape="ellipse") + value_nodes[value_name] = value_node + pydot_graph.add_node(value_node) + pydot_graph.add_edge(pydot.Edge(op_node, value_node)) -pydot_graph = GetPydotGraph( - model.graph, name=model.graph.name, rankdir="LR", node_producer=GetOpNodeProducer("docstring") -) pydot_graph.write_dot("graph.dot") ####################################### # Then into an image -import os # noqa: E402 - -os.system("dot -O -Tpng graph.dot") +pydot_graph.write_png("graph.dot.png") ################################ # Which we display... diff --git a/include/onnxruntime/core/graph/schema_registry.h b/include/onnxruntime/core/graph/schema_registry.h index ca51e3621b2c6..f60589a6d82e0 100644 --- a/include/onnxruntime/core/graph/schema_registry.h +++ b/include/onnxruntime/core/graph/schema_registry.h @@ -37,7 +37,7 @@ class IOnnxRuntimeOpSchemaCollection : public ONNX_NAMESPACE::ISchemaRegistry { using ISchemaRegistry::GetSchema; - const ONNX_NAMESPACE::OpSchema* GetSchema(const std::string& key, const int maxInclusiveVersion, + const ONNX_NAMESPACE::OpSchema* GetSchema(const std::string& key, int maxInclusiveVersion, const std::string& domain) const final { const ONNX_NAMESPACE::OpSchema* latest_schema = nullptr; int earliest_opset_where_unchanged = std::numeric_limits::max(); diff --git a/include/onnxruntime/core/session/onnxruntime_c_api.h b/include/onnxruntime/core/session/onnxruntime_c_api.h index 8c7f6e4389363..88ee8942e1695 100644 --- a/include/onnxruntime/core/session/onnxruntime_c_api.h +++ b/include/onnxruntime/core/session/onnxruntime_c_api.h @@ -223,6 +223,11 @@ typedef enum ONNXTensorElementDataType { ONNX_TENSOR_ELEMENT_DATA_TYPE_INT2, // maps to 4 packed int2 values (size == 1 byte) // Float8E8M0 type introduced in ONNX 1.21. 8-bit float with 8 exponent bits, 0 mantissa bits, no sign bit. ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT8E8M0, // Non-IEEE floating-point format, all values are powers of two + // Float6 types introduced in ONNX 1.23. + // ORT tensor storage uses one byte per element. TensorProto typed serialization uses int32_data; + // TensorProto raw_data uses bit-packed 6-bit values. + ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E2M3, + ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E3M2, } ONNXTensorElementDataType; // Synced with onnx TypeProto oneof diff --git a/include/onnxruntime/ep/adapter/tensor_helper.h b/include/onnxruntime/ep/adapter/tensor_helper.h index 949751d4cb444..b1cedbd89f63b 100644 --- a/include/onnxruntime/ep/adapter/tensor_helper.h +++ b/include/onnxruntime/ep/adapter/tensor_helper.h @@ -28,6 +28,10 @@ inline int OnnxTensorProtoTypeFromApiElementType(ONNXTensorElementDataType eleme return ONNX_NAMESPACE::TensorProto_DataType_INT2; case ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT8E8M0: return ONNX_NAMESPACE::TensorProto_DataType_FLOAT8E8M0; + case ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E2M3: + return ONNX_NAMESPACE::TensorProto_DataType_FLOAT6E2M3; + case ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E3M2: + return ONNX_NAMESPACE::TensorProto_DataType_FLOAT6E3M2; default: return static_cast(element_type); } diff --git a/js/web/docs/webgl-operators.md b/js/web/docs/webgl-operators.md index 1a8a99b88b2e0..2fbd03963809b 100644 --- a/js/web/docs/webgl-operators.md +++ b/js/web/docs/webgl-operators.md @@ -31,7 +31,7 @@ See [Compatibility](../README.md#Compatibility) for a list of the supported plat | [BitwiseOr](https://github.com/onnx/onnx/blob/main/docs/Operators.md#BitwiseOr) | | | [BitwiseXor](https://github.com/onnx/onnx/blob/main/docs/Operators.md#BitwiseXor) | | | [BlackmanWindow](https://github.com/onnx/onnx/blob/main/docs/Operators.md#BlackmanWindow) | | -| [Cast](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Cast) | [6-8](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-6), [9-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-9), [13-18](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-13), [19-20](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-19), [21-22](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-21), [23](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-23), [24](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-24), [25+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-25) | +| [Cast](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Cast) | [6-8](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-6), [9-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-9), [13-18](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-13), [19-20](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-19), [21-22](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-21), [23](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-23), [24](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-24), [25-27](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-25), [28+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Cast-28) | | [CastLike](https://github.com/onnx/onnx/blob/main/docs/Operators.md#CastLike) | | | [CausalConvWithState](https://github.com/onnx/onnx/blob/main/docs/Operators.md#CausalConvWithState) | | | [Ceil](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Ceil) | [6-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Ceil-6), [13+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Ceil-13) | @@ -53,7 +53,7 @@ See [Compatibility](../README.md#Compatibility) for a list of the supported plat | [CumSum](https://github.com/onnx/onnx/blob/main/docs/Operators.md#CumSum) | | | [DFT](https://github.com/onnx/onnx/blob/main/docs/Operators.md#DFT) | | | [DeformConv](https://github.com/onnx/onnx/blob/main/docs/Operators.md#DeformConv) | | -| [DepthToSpace](https://github.com/onnx/onnx/blob/main/docs/Operators.md#DepthToSpace) | [1-10](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#DepthToSpace-1), [11-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#DepthToSpace-11), [13+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#DepthToSpace-13) | +| [DepthToSpace](https://github.com/onnx/onnx/blob/main/docs/Operators.md#DepthToSpace) | [1-10](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#DepthToSpace-1), [11-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#DepthToSpace-11), [13-27](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#DepthToSpace-13), [28+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#DepthToSpace-28) | | [DequantizeLinear](https://github.com/onnx/onnx/blob/main/docs/Operators.md#DequantizeLinear) | | | [Det](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Det) | | | [Div](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Div) | [7-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Div-7), [13](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Div-13), [14+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Div-14) | @@ -144,7 +144,7 @@ See [Compatibility](../README.md#Compatibility) for a list of the supported plat | [Reciprocal](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Reciprocal) | | | [ReduceL1](https://github.com/onnx/onnx/blob/main/docs/Operators.md#ReduceL1) | | | [ReduceL2](https://github.com/onnx/onnx/blob/main/docs/Operators.md#ReduceL2) | | -| [ReduceLogSum](https://github.com/onnx/onnx/blob/main/docs/Operators.md#ReduceLogSum) | [1-10](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-1), [11-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-11), [13-17](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-13), [18+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-18) | +| [ReduceLogSum](https://github.com/onnx/onnx/blob/main/docs/Operators.md#ReduceLogSum) | [1-10](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-1), [11-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-11), [13-17](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-13), [18-27](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-18), [28+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceLogSum-28) | | [ReduceLogSumExp](https://github.com/onnx/onnx/blob/main/docs/Operators.md#ReduceLogSumExp) | | | [ReduceMax](https://github.com/onnx/onnx/blob/main/docs/Operators.md#ReduceMax) | [1-10](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMax-1), [11](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMax-11), [12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMax-12), [13-17](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMax-13), [18-19](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMax-18), [20+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMax-20) | | [ReduceMean](https://github.com/onnx/onnx/blob/main/docs/Operators.md#ReduceMean) | [1-10](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMean-1), [11-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMean-11), [13-17](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMean-13), [18+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#ReduceMean-18) | @@ -195,6 +195,7 @@ See [Compatibility](../README.md#Compatibility) for a list of the supported plat | [StringSplit](https://github.com/onnx/onnx/blob/main/docs/Operators.md#StringSplit) | | | [Sub](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Sub) | [7-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Sub-7), [13](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Sub-13), [14+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Sub-14) | | [Sum](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Sum) | [6-7](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Sum-6), [8-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Sum-8), [13+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Sum-13) | +| [SwiGLU](https://github.com/onnx/onnx/blob/main/docs/Operators.md#SwiGLU) | | | [Swish](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Swish) | | | [Tan](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Tan) | [7-21](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Tan-7), [22+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Tan-22) | | [Tanh](https://github.com/onnx/onnx/blob/main/docs/Operators.md#Tanh) | [6-12](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Tanh-6), [13+](https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Tanh-13) | diff --git a/onnxruntime/core/framework/onnxruntime_map_type_info.h b/onnxruntime/core/framework/onnxruntime_map_type_info.h index 71c680a2b1531..4e51aa00d3b3e 100644 --- a/onnxruntime/core/framework/onnxruntime_map_type_info.h +++ b/onnxruntime/core/framework/onnxruntime_map_type_info.h @@ -26,7 +26,7 @@ constexpr ONNXTensorElementDataType ToONNXTensorElementDataType( return ONNX_TENSOR_ELEMENT_DATA_TYPE_UNDEFINED; } - // The enums have the same values through FLOAT4E2M1. Their final three entries + // The enums have the same values through FLOAT4E2M1. The following three entries // differ because ONNX inserted FLOAT8E8M0 before UINT2 and INT2. switch (data_type) { case ONNX_NAMESPACE::TensorProto_DataType_FLOAT8E8M0: @@ -42,7 +42,7 @@ constexpr ONNXTensorElementDataType ToONNXTensorElementDataType( consteval bool IsTensorProtoToOrtElementTypeMapBijective() { constexpr size_t ort_element_type_count = - static_cast(ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT8E8M0) + 1; + static_cast(ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E3M2) + 1; constexpr size_t onnx_element_type_count = ONNX_NAMESPACE::TensorProto_DataType_DataType_ARRAYSIZE; if constexpr (onnx_element_type_count != ort_element_type_count) { return false; diff --git a/onnxruntime/core/framework/tensor_type_and_shape.cc b/onnxruntime/core/framework/tensor_type_and_shape.cc index e91d0396968e3..8f9ec22e49a3e 100644 --- a/onnxruntime/core/framework/tensor_type_and_shape.cc +++ b/onnxruntime/core/framework/tensor_type_and_shape.cc @@ -187,6 +187,12 @@ constexpr ONNXTensorElementDataType TensorDataTypeToOnnxRuntimeTensorElementData case o::TensorProto_DataType_FLOAT8E8M0: type = ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT8E8M0; break; + case o::TensorProto_DataType_FLOAT6E2M3: + type = ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E2M3; + break; + case o::TensorProto_DataType_FLOAT6E3M2: + type = ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E3M2; + break; case o::TensorProto_DataType_FLOAT16: type = ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT16; break; diff --git a/onnxruntime/core/framework/tensorprotoutils.cc b/onnxruntime/core/framework/tensorprotoutils.cc index d2586ddf97a9f..ebb57773bf8bd 100644 --- a/onnxruntime/core/framework/tensorprotoutils.cc +++ b/onnxruntime/core/framework/tensorprotoutils.cc @@ -2119,6 +2119,8 @@ ONNXTensorElementDataType CApiElementTypeFromProtoType(int type) { #if !defined(DISABLE_FLOAT4_TYPES) CASE_TYPE(FLOAT4E2M1) #endif + CASE_TYPE(FLOAT6E2M3) + CASE_TYPE(FLOAT6E3M2) default: return ONNX_TENSOR_ELEMENT_DATA_TYPE_UNDEFINED; diff --git a/onnxruntime/core/framework/tensorprotoutils.h b/onnxruntime/core/framework/tensorprotoutils.h index 384136951ec7d..b8728b2f6c9ef 100644 --- a/onnxruntime/core/framework/tensorprotoutils.h +++ b/onnxruntime/core/framework/tensorprotoutils.h @@ -58,6 +58,8 @@ inline constexpr std::array path1{ - {0, 0, "Cast", {9, 13, 19, 21, 23, 24, 25}, kOnnxDomain}, + {0, 0, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}, kOnnxDomain}, {0, 0, "Slice", {10, 11, 13}, kOnnxDomain}, // Last Slice {0, 0, "Slice", {10, 11, 13}, kOnnxDomain}, // Mask Slice {0, 1, "Unsqueeze", {1, 11, 13, 21, 23, 24, 25}, kOnnxDomain}, diff --git a/onnxruntime/core/optimizer/embed_layer_norm_fusion.cc b/onnxruntime/core/optimizer/embed_layer_norm_fusion.cc index 0101a4f8b9da0..045ac89970b31 100644 --- a/onnxruntime/core/optimizer/embed_layer_norm_fusion.cc +++ b/onnxruntime/core/optimizer/embed_layer_norm_fusion.cc @@ -248,8 +248,8 @@ static bool MatchPositionEmbeddingSubgraphsFromGather( std::vector parent_path_1{ {0, 1, "Expand", {8, 13}, kOnnxDomain}, {0, 0, "Unsqueeze", {1, 11, 13, 21, 23, 24, 25}, kOnnxDomain}, - {0, 0, "Cast", {9, 13, 19, 21, 23, 24, 25}, kOnnxDomain}, - {0, 0, "Squeeze", {1, 11, 13, 21, 23, 24, 25}, kOnnxDomain}, + {0, 0, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}, kOnnxDomain}, + {0, 0, "Squeeze", {1, 11, 13, 21, 23, 24, 25, 28}, kOnnxDomain}, {0, 0, "Transpose", {1, 13, 21, 23, 24, 25}, kOnnxDomain}, {0, 0, "NonZero", {9, 13}, kOnnxDomain}, {0, 0, "ConstantOfShape", {9, 20, 21, 23, 24, 25}, kOnnxDomain}, @@ -260,7 +260,7 @@ static bool MatchPositionEmbeddingSubgraphsFromGather( std::vector parent_path_2{ {0, 1, "Expand", {8, 13}, kOnnxDomain}, {0, 0, "Unsqueeze", {1, 11, 13, 21, 23, 24, 25}, kOnnxDomain}, - {0, 0, "Squeeze", {1, 11, 13, 21, 23, 24, 25}, kOnnxDomain}, + {0, 0, "Squeeze", {1, 11, 13, 21, 23, 24, 25, 28}, kOnnxDomain}, {0, 0, "Transpose", {1, 13, 21, 23, 24, 25}, kOnnxDomain}, {0, 0, "NonZero", {9, 13}, kOnnxDomain}, {0, 0, "ConstantOfShape", {9, 20, 21, 23, 24, 25}, kOnnxDomain}, @@ -272,15 +272,15 @@ static bool MatchPositionEmbeddingSubgraphsFromGather( std::vector parent_path_3{ {0, 1, "Expand", {8, 13}, kOnnxDomain}, {0, 0, "Unsqueeze", {1, 11, 13, 21, 23, 24, 25}, kOnnxDomain}, - {0, 0, "Range", {11, 27}, kOnnxDomain}, - {0, 1, "Cast", {9, 13, 19, 21, 23, 24, 25}, kOnnxDomain}, + {0, 0, "Range", {11, 27, 28}, kOnnxDomain}, + {0, 1, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}, kOnnxDomain}, {0, 0, "Gather", {1, 11, 13}, kOnnxDomain}, {0, 0, "Shape", {1, 13, 15, 19, 21, 23, 24, 25}, kOnnxDomain}}; // Path 4 pattern (Path 3 with no "Cast"): std::vector parent_path_4{ {0, 1, "Expand", {8, 13}, kOnnxDomain}, {0, 0, "Unsqueeze", {1, 11, 13, 21, 23, 24, 25}, kOnnxDomain}, - {0, 0, "Range", {11, 27}, kOnnxDomain}, + {0, 0, "Range", {11, 27, 28}, kOnnxDomain}, {0, 1, "Gather", {1, 11, 13}, kOnnxDomain}, {0, 0, "Shape", {1, 13, 15, 19, 21, 23, 24, 25}, kOnnxDomain}}; // Match one of the three path patterns. diff --git a/onnxruntime/core/optimizer/fast_gelu_fusion.cc b/onnxruntime/core/optimizer/fast_gelu_fusion.cc index cda93c9e09e31..8b911a6e63914 100644 --- a/onnxruntime/core/optimizer/fast_gelu_fusion.cc +++ b/onnxruntime/core/optimizer/fast_gelu_fusion.cc @@ -177,7 +177,7 @@ MatchResult FastGeluFusion::CheckSecondFormula(Graph& graph, Node& pow1_node, if (p_cast1_node != nullptr) { Node& cast1_node = *graph.GetNode(p_cast1_node->Index()); // this is fused Cast node, so expect 2 output edges - if (!(graph_utils::IsSupportedOptypeVersionAndDomain(cast1_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) && + if (!(graph_utils::IsSupportedOptypeVersionAndDomain(cast1_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) && cast1_node.InputDefs().size() == 1 && CheckNode(graph, cast1_node, pow1_node.GetExecutionProviderType(), false)) || cast1_node.GetOutputEdgesCount() != 2) { @@ -293,7 +293,7 @@ Status FastGeluFusion::ApplyImpl(Graph& graph, bool& modified, int graph_level, if (p_cast3_node == nullptr) continue; Node& cast3_node = *graph.GetNode(p_cast3_node->Index()); - if (!(graph_utils::IsSupportedOptypeVersionAndDomain(cast3_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) && + if (!(graph_utils::IsSupportedOptypeVersionAndDomain(cast3_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) && cast3_node.InputDefs().size() == 1 && CheckNode(graph, cast3_node, node.GetExecutionProviderType(), true))) { continue; diff --git a/onnxruntime/core/optimizer/isinf_reducesum_fusion.cc b/onnxruntime/core/optimizer/isinf_reducesum_fusion.cc index 3f963e86acb0d..2185261629e24 100644 --- a/onnxruntime/core/optimizer/isinf_reducesum_fusion.cc +++ b/onnxruntime/core/optimizer/isinf_reducesum_fusion.cc @@ -45,7 +45,7 @@ Status IsInfReduceSumFusion::ApplyImpl(Graph& graph, bool& modified, int graph_l // This Cast can be skipped as we are replacing the subgraph with IsAllFinite, which supports FP16 auto cast1_node_iter = isinf_node.InputNodesBegin(); if (cast1_node_iter != isinf_node.InputNodesEnd() && - graph_utils::IsSupportedOptypeVersionAndDomain(*cast1_node_iter, "Cast", {9, 13, 19, 21, 23, 24, 25}) && + graph_utils::IsSupportedOptypeVersionAndDomain(*cast1_node_iter, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) && cast1_node_iter->GetOutputEdgesCount() == 1) { // check input type of cast node Node& cast1_node = *graph.GetNode(cast1_node_iter->Index()); @@ -65,7 +65,7 @@ Status IsInfReduceSumFusion::ApplyImpl(Graph& graph, bool& modified, int graph_l } Node& cast2_node = *graph.GetNode(cast2_node_itr->Index()); - if (!graph_utils::IsSupportedOptypeVersionAndDomain(cast2_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) || + if (!graph_utils::IsSupportedOptypeVersionAndDomain(cast2_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) || cast2_node.GetOutputEdgesCount() != 1 || graph.NodeProducesGraphOutput(cast2_node)) { continue; diff --git a/onnxruntime/core/optimizer/layer_norm_fusion.cc b/onnxruntime/core/optimizer/layer_norm_fusion.cc index 753b23724d7eb..93d4ed851f8e6 100644 --- a/onnxruntime/core/optimizer/layer_norm_fusion.cc +++ b/onnxruntime/core/optimizer/layer_norm_fusion.cc @@ -145,7 +145,7 @@ static bool IsPowExponentTwo(const Graph& graph, const Node& pow_node) { const Node* exponent_input_node = graph_utils::GetInputNode(pow_node, 1); if (exponent_input_node == nullptr || - !graph_utils::IsSupportedOptypeVersionAndDomain(*exponent_input_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) || + !graph_utils::IsSupportedOptypeVersionAndDomain(*exponent_input_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) || exponent_input_node->InputDefs().empty() || exponent_input_node->InputDefs()[0] == nullptr) { return false; } @@ -334,7 +334,7 @@ Status LayerNormFusion::ApplyImpl(Graph& graph, bool& modified, int graph_level, if (p_reduce_mean_input_node) { Node& reduce_mean_input_node = *graph.GetNode(p_reduce_mean_input_node->Index()); // If input to the 1st ReduceMean is a Cast, and the Cast has same consumer count as subCnt + 1 - if (graph_utils::IsSupportedOptypeVersionAndDomain(reduce_mean_input_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) && + if (graph_utils::IsSupportedOptypeVersionAndDomain(reduce_mean_input_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) && reduce_mean_input_node.GetExecutionProviderType() == reduce_mean_node.GetExecutionProviderType() && optimizer_utils::CheckOutputEdges(graph, reduce_mean_input_node, static_cast(subCnt) + 1)) { nodes_to_remove.insert(nodes_to_remove.begin(), reduce_mean_input_node); @@ -347,7 +347,7 @@ Status LayerNormFusion::ApplyImpl(Graph& graph, bool& modified, int graph_level, const Node* p_cast1 = nullptr; if (!p_sub_node_dup && sub_node.GetOutputEdgesCount() == 1) { Node& cast_node = *graph.GetNode(sub_node.OutputNodesBegin()->Index()); - if (graph_utils::IsSupportedOptypeVersionAndDomain(cast_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) && + if (graph_utils::IsSupportedOptypeVersionAndDomain(cast_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) && cast_node.GetExecutionProviderType() == reduce_mean_node.GetExecutionProviderType() && optimizer_utils::CheckOutputEdges(graph, cast_node, 2u) && IsSupportedDataType(cast_node)) { p_cast1 = &cast_node; @@ -446,7 +446,7 @@ Status LayerNormFusion::ApplyImpl(Graph& graph, bool& modified, int graph_level, const Node* p_cast2 = graph_utils::FirstParentByType(pow_node, "Cast"); if (p_cast2 != nullptr && p_cast2 != p_cast1) { Node& cast_node = *graph.GetNode(p_cast2->Index()); - if (!graph_utils::IsSupportedOptypeVersionAndDomain(cast_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) || + if (!graph_utils::IsSupportedOptypeVersionAndDomain(cast_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) || cast_node.GetExecutionProviderType() != reduce_mean_node.GetExecutionProviderType() || !optimizer_utils::CheckOutputEdges(graph, cast_node, 1)) { continue; @@ -464,7 +464,7 @@ Status LayerNormFusion::ApplyImpl(Graph& graph, bool& modified, int graph_level, // can be removed. This is one possible place a Cast Op can exist, that is between Div and Mul nodes. // div --> mul or div --> cast --> mul Node* next_node = graph.GetNode(div_node.OutputNodesBegin()->Index()); - if (graph_utils::IsSupportedOptypeVersionAndDomain(*next_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) && + if (graph_utils::IsSupportedOptypeVersionAndDomain(*next_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) && optimizer_utils::CheckOutputEdges(graph, *next_node, 1)) { nodes_to_remove.push_back(*next_node); next_node = graph.GetNode(next_node->OutputNodesBegin()->Index()); @@ -745,7 +745,7 @@ Status SimplifiedLayerNormFusion::ApplyImpl(Graph& graph, bool& modified, int gr if (is_gpu_ep && p_pow_input_node) { Node& pow_input_node = *graph.GetNode(p_pow_input_node->Index()); // If input to Pow is a Cast, and the Cast has 2 consumers only (Pow, Div) - if (graph_utils::IsSupportedOptypeVersionAndDomain(pow_input_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) && + if (graph_utils::IsSupportedOptypeVersionAndDomain(pow_input_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) && pow_input_node.GetExecutionProviderType() == pow_node.GetExecutionProviderType() && optimizer_utils::CheckOutputEdges(graph, pow_input_node, 2)) { nodes_to_remove.insert(nodes_to_remove.begin(), pow_input_node); @@ -755,7 +755,7 @@ Status SimplifiedLayerNormFusion::ApplyImpl(Graph& graph, bool& modified, int gr // div --> mul or div --> cast --> mul Node* next_node = graph.GetNode(div_node.OutputNodesBegin()->Index()); - if (graph_utils::IsSupportedOptypeVersionAndDomain(*next_node, "Cast", {9, 13, 19, 21, 23, 24, 25}) && + if (graph_utils::IsSupportedOptypeVersionAndDomain(*next_node, "Cast", {9, 13, 19, 21, 23, 24, 25, 28}) && optimizer_utils::CheckOutputEdges(graph, *next_node, 1)) { if (!is_gpu_ep) continue; nodes_to_remove.push_back(*next_node); diff --git a/onnxruntime/core/optimizer/pad_fusion.cc b/onnxruntime/core/optimizer/pad_fusion.cc index 179a08e0c6f9e..9a54f06c57f09 100644 --- a/onnxruntime/core/optimizer/pad_fusion.cc +++ b/onnxruntime/core/optimizer/pad_fusion.cc @@ -132,7 +132,7 @@ bool PadFusion::SatisfyCondition(const Graph& graph, const Node& node, const log } const Node& child_node = *node.OutputNodesBegin(); - if (graph_utils::IsSupportedOptypeVersionAndDomain(child_node, "Cast", {1, 6, 9, 13, 19, 21, 23, 24, 25})) { + if (graph_utils::IsSupportedOptypeVersionAndDomain(child_node, "Cast", {1, 6, 9, 13, 19, 21, 23, 24, 25, 28})) { if (child_node.GetOutputEdgesCount() != 1) { return false; } diff --git a/onnxruntime/core/optimizer/transpose_optimization/optimizer_api.h b/onnxruntime/core/optimizer/transpose_optimization/optimizer_api.h index cd0094877657d..ea3af9abd3bba 100644 --- a/onnxruntime/core/optimizer/transpose_optimization/optimizer_api.h +++ b/onnxruntime/core/optimizer/transpose_optimization/optimizer_api.h @@ -478,7 +478,7 @@ class GraphRef { } // namespace api constexpr int64_t kMinSupportedOpset = 7; -constexpr int64_t kMaxSupportedOpset = 27; +constexpr int64_t kMaxSupportedOpset = 28; // enum of results that a CostCheckFn can return. enum class CostCheckResult { diff --git a/onnxruntime/core/providers/coreml/builders/impl/base_op_builder.h b/onnxruntime/core/providers/coreml/builders/impl/base_op_builder.h index 609a965fa5198..0e49f488f8684 100644 --- a/onnxruntime/core/providers/coreml/builders/impl/base_op_builder.h +++ b/onnxruntime/core/providers/coreml/builders/impl/base_op_builder.h @@ -46,7 +46,7 @@ class BaseOpBuilder : public IOpBuilder { const logging::Logger& logger) const; virtual int GetMinSupportedOpSet(const Node& /*node*/) const { return 1; } - virtual int GetMaxSupportedOpSet(const Node& /*node*/) const { return 27; } + virtual int GetMaxSupportedOpSet(const Node& /*node*/) const { return 28; } bool HasSupportedOpSet(const Node& node, const logging::Logger& logger) const; bool HasSupportedInputs(const Node& node, const OpBuilderInputParams& input_params, diff --git a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc index bdf0de496e8ef..add722a2e8082 100644 --- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc +++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc @@ -598,9 +598,10 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOn class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 12, MLFloat16, Gemm); #endif class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 12, GatherElements); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint8_t, BitShift); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint32_t, BitShift); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint64_t, BitShift); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 27, uint8_t, BitShift); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 27, uint16_t, BitShift); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 27, uint32_t, BitShift); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 27, uint64_t, BitShift); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 12, Pad); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, GatherND); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 26, Range); @@ -836,7 +837,8 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int16_t, Neg); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int32_t, Neg); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int64_t, Neg); -class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, Mod); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 27, Mod); +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, Mod); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, float, Abs); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, double, Abs); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int8_t, Abs); @@ -1554,6 +1556,16 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, // Opset 27 class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 27, Range); +// Opset 28 +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int8_t, BitShift); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int16_t, BitShift); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int32_t, BitShift); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int64_t, BitShift); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, uint8_t, BitShift); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, uint16_t, BitShift); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, uint32_t, BitShift); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, uint64_t, BitShift); + // !!PLEASE READ BELOW!! Following that, add new entries above this comment /* *** IMPORTANT! *** @@ -2320,12 +2332,14 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { double, Gemm)>, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -2769,7 +2783,8 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -3757,6 +3772,24 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { // opset 27 BuildKernelCreateInfo, + + // opset 28 + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, }; for (auto& function_table_entry : function_table) { KernelCreateInfo info = function_table_entry(); diff --git a/onnxruntime/core/providers/cpu/math/element_wise_ops.cc b/onnxruntime/core/providers/cpu/math/element_wise_ops.cc index 4a9c4646a2c34..f4d8b9b6211b7 100644 --- a/onnxruntime/core/providers/cpu/math/element_wise_ops.cc +++ b/onnxruntime/core/providers/cpu/math/element_wise_ops.cc @@ -12,6 +12,9 @@ #include "core/mlas/inc/mlas.h" #include +#include +#include +#include namespace onnxruntime { // Supported types for operators that have type reduction enabled @@ -489,10 +492,18 @@ REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Mean, 8, 12, float, Mean_8); // Supposed to add BFloat16 but we are not supporting now, however, separate registration REG_ELEMENTWISE_TYPED_KERNEL(Mean, 13, float, Mean_8); -REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 11, uint8_t, BitShift); -// REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 11, uint16_t, BitShift); -REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 11, uint32_t, BitShift); -REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 11, uint64_t, BitShift); +REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(BitShift, 11, 27, uint8_t, BitShift); +REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(BitShift, 11, 27, uint16_t, BitShift); +REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(BitShift, 11, 27, uint32_t, BitShift); +REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(BitShift, 11, 27, uint64_t, BitShift); +REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 28, int8_t, BitShift); +REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 28, int16_t, BitShift); +REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 28, int32_t, BitShift); +REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 28, int64_t, BitShift); +REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 28, uint8_t, BitShift); +REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 28, uint16_t, BitShift); +REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 28, uint32_t, BitShift); +REG_ELEMENTWISE_TYPED_KERNEL(BitShift, 28, uint64_t, BitShift); REG_ELEMENTWISE_TYPED_KERNEL(BitwiseAnd, 18, int8_t, BitwiseAnd); REG_ELEMENTWISE_TYPED_KERNEL(BitwiseAnd, 18, int16_t, BitwiseAnd); @@ -1334,17 +1345,54 @@ BitShift::BitShift(const OpKernelInfo& info) : OpKernel(info) { ORT_THROW("Invalid direction value of '", direction, "'. Valid values are 'LEFT' or 'RIGHT'."); } -// Shifting by >= the bit width of an unsigned type is undefined behavior in C++. -// On x86, 64-bit shifts mask the shift amount to 6 bits, so shift by 64 acts like shift by 0. -// Guard against this by returning 0 when the shift amount >= the bit width. +template +T FromUnsignedBits(std::make_unsigned_t bits) { + static_assert(std::is_integral_v); + T value; + std::memcpy(&value, &bits, sizeof(value)); + return value; +} + +template +inline bool IsInvalidShift(T shift) { + using UnsignedT = std::make_unsigned_t; + constexpr UnsignedT bit_width = std::numeric_limits::digits; + return (std::is_signed_v && shift < 0) || static_cast(shift) >= bit_width; +} + template inline T SafeShiftLeft(T value, T shift) { - return shift >= sizeof(T) * 8 ? T{0} : value << shift; + if (IsInvalidShift(shift)) { + return T{0}; + } + + using UnsignedT = std::make_unsigned_t; + const auto result = static_cast( + static_cast(value) << static_cast(shift)); + return FromUnsignedBits(result); } template inline T SafeShiftRight(T value, T shift) { - return shift >= sizeof(T) * 8 ? T{0} : value >> shift; + if (IsInvalidShift(shift)) { + if constexpr (std::is_signed_v) { + return value < 0 ? T{-1} : T{0}; + } else { + return T{0}; + } + } + + using UnsignedT = std::make_unsigned_t; + constexpr UnsignedT bit_width = std::numeric_limits::digits; + const auto shift_amount = static_cast(shift); + auto result = static_cast(static_cast(value) >> shift_amount); + if constexpr (std::is_signed_v) { + if (value < 0 && shift_amount != 0) { + result |= std::numeric_limits::max() << (bit_width - shift_amount); + } + } + + return FromUnsignedBits(result); } template @@ -2110,9 +2158,19 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL( BuildKernelDefConstraintsFromTypeList()), Mod); -ONNX_CPU_OPERATOR_KERNEL( +ONNX_CPU_OPERATOR_VERSIONED_KERNEL( Mod, 13, + 27, + KernelDefBuilder() + .TypeConstraint( + "T", + BuildKernelDefConstraintsFromTypeList()), + Mod); + +ONNX_CPU_OPERATOR_KERNEL( + Mod, + 28, KernelDefBuilder() .TypeConstraint( "T", @@ -2160,11 +2218,20 @@ void BroadCastFMod(OpKernelContext* context) { template inline T Modulus(T x, T y) { - auto res = x % y; + T res; + if constexpr (std::is_floating_point_v) { + res = std::fmod(x, y); + if (res == T{0}) { + return std::copysign(T{0}, y); + } + } else { + res = x % y; + } + if ((res < 0 && y > 0) || (res > 0 && y < 0)) { res += y; } - return static_cast(res); + return res; } template @@ -2242,6 +2309,42 @@ void BroadCastMLFloat16FMod(OpKernelContext* context) { UntypedBroadcastTwo(*context, funcs); } +void BroadCastMLFloat16Mod(OpKernelContext* context) { + ProcessBroadcastSpanFuncs funcs{ + [](BroadcastHelper& per_iter_bh) { + const auto X = per_iter_bh.ScalarInput0(); + auto Y = per_iter_bh.SpanInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(Y.begin(), Y.end(), output.begin(), + [X_fl = math::halfToFloat(X.val)](const MLFloat16& y) { + return MLFloat16(Modulus(X_fl, y.ToFloat())); + }); + }, + [](BroadcastHelper& per_iter_bh) { + auto X = per_iter_bh.SpanInput0(); + const MLFloat16 Y = per_iter_bh.ScalarInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(X.begin(), X.end(), output.begin(), + [Y_fl = math::halfToFloat(Y.val)](const MLFloat16& x) { + return MLFloat16(Modulus(x.ToFloat(), Y_fl)); + }); + }, + [](BroadcastHelper& per_iter_bh) { + auto X = per_iter_bh.SpanInput0(); + auto Y = per_iter_bh.SpanInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(X.begin(), X.end(), Y.begin(), output.begin(), + [](const MLFloat16& x, const MLFloat16& y) { + return MLFloat16(Modulus(x.ToFloat(), y.ToFloat())); + }); + }}; + + UntypedBroadcastTwo(*context, funcs); +} + template struct CallModImpl; @@ -2276,8 +2379,11 @@ struct CallModImpl::value> template struct CallModImpl::value, void>::type> { void operator()(bool fmod, OpKernelContext* ctx) const { - ORT_ENFORCE(fmod, "fmod attribute must be true for floating point types"); - BroadCastFMod(ctx); + if (fmod) { + BroadCastFMod(ctx); + } else { + BroadCastMod(ctx); + } } }; @@ -2285,8 +2391,11 @@ struct CallModImpl::value, template <> struct CallModImpl { void operator()(bool fmod, OpKernelContext* ctx) const { - ORT_ENFORCE(fmod, "fmod attribute must be true for floating point types"); - BroadCastMLFloat16FMod(ctx); + if (fmod) { + BroadCastMLFloat16FMod(ctx); + } else { + BroadCastMLFloat16Mod(ctx); + } } }; diff --git a/onnxruntime/core/providers/cuda/cu_inc/common.cuh b/onnxruntime/core/providers/cuda/cu_inc/common.cuh index f53da982538bd..a46385c9d2314 100644 --- a/onnxruntime/core/providers/cuda/cu_inc/common.cuh +++ b/onnxruntime/core/providers/cuda/cu_inc/common.cuh @@ -509,6 +509,40 @@ __device__ __inline__ T _Mod(T a, T b) { return r; } +template <> +__device__ __inline__ float _Mod(float a, float b) { + float r = fmodf(a, b); + if (r == 0.0f) { + return copysignf(0.0f, b); + } + if ((r > 0.0f && b < 0.0f) || (r < 0.0f && b > 0.0f)) { + r += b; + } + return r; +} + +template <> +__device__ __inline__ double _Mod(double a, double b) { + double r = fmod(a, b); + if (r == 0.0) { + return copysign(0.0, b); + } + if ((r > 0.0 && b < 0.0) || (r < 0.0 && b > 0.0)) { + r += b; + } + return r; +} + +template <> +__device__ __inline__ half _Mod(half a, half b) { + return half(_Mod(static_cast(a), static_cast(b))); +} + +template <> +__device__ __inline__ BFloat16 _Mod(BFloat16 a, BFloat16 b) { + return BFloat16(_Mod(static_cast(a), static_cast(b))); +} + template __device__ __inline__ T _Fmod(T a, T b) { return a % b; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index e6645a0d4cb77..80dbd322fa534 100755 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -1354,7 +1354,8 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Tanh); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Gemm); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, ReduceSum); -class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, Mod); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 27, Mod); +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Mod); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, QuantizeLinear); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, uint8_t, QuantizeLinear); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, DequantizeLinear); @@ -2647,7 +2648,8 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc index babbb4b3ba672..ce4d089fa7b67 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -467,7 +467,14 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX( BuildKernelDefConstraints()), Mod); -ONNX_OPERATOR_KERNEL_EX(Mod, kOnnxDomain, 13, kCudaExecutionProvider, +ONNX_OPERATOR_VERSIONED_KERNEL_EX( + Mod, kOnnxDomain, 13, 27, kCudaExecutionProvider, + (*KernelDefBuilder::Create()) + .TypeConstraint("T", BuildKernelDefConstraints()), + Mod); + +ONNX_OPERATOR_KERNEL_EX(Mod, kOnnxDomain, 28, kCudaExecutionProvider, (*KernelDefBuilder::Create()) .TypeConstraint("T", BuildKernelDefConstraints()), @@ -478,10 +485,6 @@ Status Mod::ComputeInternal(OpKernelContext* context) const { BinaryElementwisePreparation prepare; ORT_RETURN_IF_ERROR(Prepare(context, &prepare)); auto element_type = prepare.lhs_tensor->GetElementType(); - ORT_ENFORCE(fmod_ || element_type == on::TensorProto_DataType_INT32 || - element_type == on::TensorProto_DataType_INT64 || element_type == on::TensorProto_DataType_UINT32 || - element_type == on::TensorProto_DataType_UINT64, - "Non-fmod can support integer types only."); #define CASE_MOD_ELEMENT_TYPE(name, onnx_type, data_type) \ case onnx_type: { \ Impl_##name::MappedType>( \ @@ -514,6 +517,10 @@ Status Mod::ComputeInternal(OpKernelContext* context) const { CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_INT64, int64_t); CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_UINT32, uint32_t); CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_UINT64, uint64_t); + CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_FLOAT, float); + CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_DOUBLE, double); + CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_FLOAT16, MLFloat16); + CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_BFLOAT16, BFloat16); default: return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Unsupported element type: ", DataTypeImpl::ToString(prepare.lhs_tensor->DataType())); diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu index 4e71a9f7c5090..2b08614b1e9c4 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu @@ -162,7 +162,7 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Xor, bool) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(PRelu) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Max) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Min) -SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZIL(Mod) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Mod) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Fmod) // create declarations for impl for Pow diff --git a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/AbiCustomRegistry.cpp b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/AbiCustomRegistry.cpp index 076027dd3672f..3edc8d216f7c1 100644 --- a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/AbiCustomRegistry.cpp +++ b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/AbiCustomRegistry.cpp @@ -391,6 +391,16 @@ HRESULT STDMETHODCALLTYPE AbiCustomRegistry::RegisterOperatorKernel( builder.SetDomain(opKernel->domain) .SinceVersion(opKernel->minimumOperatorSetVersion) .Provider(providerType); + if (isInternalOperator && + opKernel->executionType == MLOperatorExecutionType::D3D12 && + opKernel->domain != nullptr && + std::string_view(opKernel->domain).empty() && + std::string_view(opKernel->name) == "BitShift" && + opKernel->minimumOperatorSetVersion == 11) + { + // DirectML masks out-of-range shift counts, which does not satisfy the opset-28 semantics. + builder.SinceVersion(11, 27); + } std::string_view name(opKernel->name); if (name == "MemcpyToHost") diff --git a/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/base_op_builder.h b/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/base_op_builder.h index a573b73d2bbb5..c52458b7046c5 100644 --- a/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/base_op_builder.h +++ b/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/base_op_builder.h @@ -72,7 +72,7 @@ class BaseOpBuilder : public IOpBuilder { const OpSupportCheckParams& params) const; virtual int GetMinSupportedOpSet(const NodeUnit& /* node_unit */) const { return 1; } - virtual int GetMaxSupportedOpSet(const NodeUnit& /* node_unit */) const { return 27; } + virtual int GetMaxSupportedOpSet(const NodeUnit& /* node_unit */) const { return 28; } // Check if this node_unit's type is supported // SingleNode type NodeUnit is supported diff --git a/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/dequantizelinear_op_builder.cc b/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/dequantizelinear_op_builder.cc index ed9062f894660..72544d87413cd 100644 --- a/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/dequantizelinear_op_builder.cc +++ b/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/dequantizelinear_op_builder.cc @@ -37,6 +37,8 @@ class DequantizeLinearOpBuilder : public BaseOpBuilder { return ANEURALNETWORKS_FEATURE_LEVEL_1; } + int GetMaxSupportedOpSet(const NodeUnit& /* node_unit */) const override { return 27; } + bool HasSupportedInputOutputsImpl( const GraphViewer& graph_viewer, const NodeUnit& node_unit, const OpSupportCheckParams& params) const override { diff --git a/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/quantizelinear_op_builder.cc b/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/quantizelinear_op_builder.cc index eb81f5e3f59ee..9a94f37b68ae6 100644 --- a/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/quantizelinear_op_builder.cc +++ b/onnxruntime/core/providers/nnapi/nnapi_builtin/builders/impl/quantizelinear_op_builder.cc @@ -37,6 +37,8 @@ class QuantizeLinearOpBuilder : public BaseOpBuilder { return ANEURALNETWORKS_FEATURE_LEVEL_3; } + int GetMaxSupportedOpSet(const NodeUnit& /* node_unit */) const override { return 27; } + bool HasSupportedInputOutputsImpl( const GraphViewer& graph_viewer, const NodeUnit& node_unit, const OpSupportCheckParams& params) const override { diff --git a/onnxruntime/core/providers/vsinpu/builders/impl/base_op_builder.h b/onnxruntime/core/providers/vsinpu/builders/impl/base_op_builder.h index 488581dbd0e97..8a9f342871f3d 100644 --- a/onnxruntime/core/providers/vsinpu/builders/impl/base_op_builder.h +++ b/onnxruntime/core/providers/vsinpu/builders/impl/base_op_builder.h @@ -49,7 +49,7 @@ class BaseOpBuilder : public IOpBuilder { virtual bool IsQuantizedOp(const NodeUnit& /* node_unit */) const { return false; } virtual int GetMinSupportedOpSet(const NodeUnit& /* node_unit */) const { return 1; } - virtual int GetMaxSupportedOpSet(const NodeUnit& /* node_unit */) const { return 27; } + virtual int GetMaxSupportedOpSet(const NodeUnit& /* node_unit */) const { return 28; } virtual bool HasSupportedInputOutputsImpl( const InitializedTensorSet& initializers, const NodeUnit& node_unit) const; diff --git a/onnxruntime/core/providers/webgpu/compute_context.h b/onnxruntime/core/providers/webgpu/compute_context.h index 03f5ff9ac5869..0d8b63e89b3b3 100644 --- a/onnxruntime/core/providers/webgpu/compute_context.h +++ b/onnxruntime/core/providers/webgpu/compute_context.h @@ -126,15 +126,13 @@ class ComputeContextBase { // // Get the logger. // + inline const auto& Logger() const { #if defined(ORT_USE_EP_API_ADAPTERS) - inline const onnxruntime::ep::adapter::Logger& Logger() const { return ep_.GetEpLogger(); - } #else - inline const logging::Logger& Logger() const { return *ep_.GetLogger(); - } #endif + } // // Run a compute shader program. diff --git a/onnxruntime/core/providers/webnn/builders/impl/base_op_builder.h b/onnxruntime/core/providers/webnn/builders/impl/base_op_builder.h index f4429fbb06457..d1334ba9a5420 100644 --- a/onnxruntime/core/providers/webnn/builders/impl/base_op_builder.h +++ b/onnxruntime/core/providers/webnn/builders/impl/base_op_builder.h @@ -52,7 +52,7 @@ class BaseOpBuilder : public IOpBuilder { // We still set the minimal supported opset to 1 as we couldn't // get the model opset version at this stage. virtual int GetMinSupportedOpSet(const Node& /* node */) const { return 1; } - virtual int GetMaxSupportedOpSet(const Node& /* node */) const { return 27; } + virtual int GetMaxSupportedOpSet(const Node& /* node */) const { return 28; } private: bool HasSupportedOpSet(const Node& node, const logging::Logger& logger) const; diff --git a/onnxruntime/core/providers/webnn/builders/impl/qdq_op_builder.cc b/onnxruntime/core/providers/webnn/builders/impl/qdq_op_builder.cc index e4e3a09391c3b..5b6ab6420ca45 100644 --- a/onnxruntime/core/providers/webnn/builders/impl/qdq_op_builder.cc +++ b/onnxruntime/core/providers/webnn/builders/impl/qdq_op_builder.cc @@ -25,6 +25,8 @@ class QDQOpBuilder : public BaseOpBuilder { const WebnnDeviceType /* device_type */, const logging::Logger& logger) const override; bool HasSupportedInputsImpl(const GraphViewer&, const Node& node, const emscripten::val& wnn_limits, const logging::Logger& logger) const override; + + int GetMaxSupportedOpSet(const Node& /* node */) const override { return 27; } }; Status QDQOpBuilder::AddToModelBuilderImpl(ModelBuilder& model_builder, diff --git a/onnxruntime/core/session/inference_session.cc b/onnxruntime/core/session/inference_session.cc index 5fb5b91fafe59..ce77c36661421 100644 --- a/onnxruntime/core/session/inference_session.cc +++ b/onnxruntime/core/session/inference_session.cc @@ -1241,7 +1241,9 @@ common::Status InferenceSession::LoadOnnxModel(const PathString& model_uri) { const bool strict_shape_type_inference = session_options_.config_options.GetConfigOrDefault( kOrtSessionOptionsConfigStrictShapeTypeInference, "0") == "1"; - ModelOptions model_opts(true, strict_shape_type_inference, check_load_cancellation_fn_); + const bool allow_released_opsets_only = session_options_.config_options.GetConfigOrDefault( + kOrtSessionOptionsConfigStrictAllowReleasedOpsetsOnly, "1") == "1"; + ModelOptions model_opts(allow_released_opsets_only, strict_shape_type_inference, check_load_cancellation_fn_); // When set, the external initializers folder overrides the model's own directory as the // base for resolving external data. The model bytes are still read from model_uri. @@ -1337,6 +1339,8 @@ common::Status InferenceSession::Load(const void* model_data, int model_data_len const bool strict_shape_type_inference = session_options_.config_options.GetConfigOrDefault( kOrtSessionOptionsConfigStrictShapeTypeInference, "0") == "1"; + const bool allow_released_opsets_only = session_options_.config_options.GetConfigOrDefault( + kOrtSessionOptionsConfigStrictAllowReleasedOpsetsOnly, "1") == "1"; PathString external_data_model_path = GetExternalInitializersFolderModelPath(session_options_.config_options); if (!external_data_model_path.empty()) { @@ -1345,7 +1349,7 @@ common::Status InferenceSession::Load(const void* model_data, int model_data_len return onnxruntime::Model::Load(std::move(model_proto), model_location_, model, HasLocalSchema() ? &custom_schema_registries_ : nullptr, *session_logger_, - ModelOptions(true, strict_shape_type_inference, + ModelOptions(allow_released_opsets_only, strict_shape_type_inference, check_load_cancellation_fn_)); }; @@ -1402,6 +1406,8 @@ common::Status InferenceSession::LoadOnnxModel(ModelProto model_proto) { #endif const bool strict_shape_type_inference = session_options_.config_options.GetConfigOrDefault( kOrtSessionOptionsConfigStrictShapeTypeInference, "0") == "1"; + const bool allow_released_opsets_only = session_options_.config_options.GetConfigOrDefault( + kOrtSessionOptionsConfigStrictAllowReleasedOpsetsOnly, "1") == "1"; PathString external_data_model_path = GetExternalInitializersFolderModelPath(session_options_.config_options); if (!external_data_model_path.empty()) { @@ -1411,7 +1417,7 @@ common::Status InferenceSession::LoadOnnxModel(ModelProto model_proto) { // This call will move model_proto to the constructed model instance return onnxruntime::Model::Load(std::move(model_proto), model_location_, model, HasLocalSchema() ? &custom_schema_registries_ : nullptr, *session_logger_, - ModelOptions(true, strict_shape_type_inference, + ModelOptions(allow_released_opsets_only, strict_shape_type_inference, check_load_cancellation_fn_)); }; @@ -1513,12 +1519,15 @@ common::Status InferenceSession::Load(const OrtModel& model_editor_api_model) { const bool strict_shape_type_inference = session_options_.config_options.GetConfigOrDefault( kOrtSessionOptionsConfigStrictShapeTypeInference, "0") == "1"; + const bool allow_released_opsets_only = session_options_.config_options.GetConfigOrDefault( + kOrtSessionOptionsConfigStrictAllowReleasedOpsetsOnly, "1") == "1"; // need to go from unique_ptr to shared_ptr when moving into model_ std::unique_ptr tmp_model; ORT_RETURN_IF_ERROR(Model::LoadFromModelEditorApiModel(model_editor_api_model, HasLocalSchema() ? &custom_schema_registries_ : nullptr, - ModelOptions(true, strict_shape_type_inference, + ModelOptions(allow_released_opsets_only, + strict_shape_type_inference, check_load_cancellation_fn_), *session_logger_, tmp_model)); diff --git a/onnxruntime/python/backend/backend.py b/onnxruntime/python/backend/backend.py index 69be7a7657adf..5edac6df88fc0 100644 --- a/onnxruntime/python/backend/backend.py +++ b/onnxruntime/python/backend/backend.py @@ -135,6 +135,9 @@ def prepare(cls, model, device=None, **kwargs): return OnnxRuntimeBackendRep(model) elif isinstance(model, (str, bytes)): options = SessionOptions() + options.add_session_config_entry( + "session.allow_released_opsets_only", "1" if cls.allowReleasedOpsetsOnly else "0" + ) for k, v in kwargs.items(): if k in _ALLOWED_SESSION_OPTIONS: setattr(options, k, v) diff --git a/onnxruntime/python/backend/backend_rep.py b/onnxruntime/python/backend/backend_rep.py index 950ce417c6c2d..75b682fc9002e 100644 --- a/onnxruntime/python/backend/backend_rep.py +++ b/onnxruntime/python/backend/backend_rep.py @@ -6,6 +6,7 @@ Implements ONNX's backend API. """ +import numpy as np from onnx.backend.base import BackendRep from onnxruntime import RunOptions @@ -23,6 +24,10 @@ ) +def _normalize_numpy_scalar(value): + return np.asarray(value) if isinstance(value, np.generic) else value + + class OnnxRuntimeBackendRep(BackendRep): """ Wraps an :class:`onnxruntime.InferenceSession` to implement ONNX's @@ -61,7 +66,7 @@ def run(self, inputs, **kwargs): # type: (Any, **Any) -> Tuple[Any, ...] if isinstance(inputs, list): inps = {} for i, inp in enumerate(self._session.get_inputs()): - inps[inp.name] = inputs[i] + inps[inp.name] = _normalize_numpy_scalar(inputs[i]) outs = self._session.run(None, inps, options) if isinstance(outs, list): return outs @@ -72,5 +77,5 @@ def run(self, inputs, **kwargs): # type: (Any, **Any) -> Tuple[Any, ...] inp = self._session.get_inputs() if len(inp) != 1: raise RuntimeError(f"Model expect {len(inp)} inputs") - inps = {inp[0].name: inputs} + inps = {inp[0].name: _normalize_numpy_scalar(inputs)} return self._session.run(None, inps, options) diff --git a/onnxruntime/test/framework/tensorutils_test.cc b/onnxruntime/test/framework/tensorutils_test.cc index 473952bfe7bb8..81845c362d937 100644 --- a/onnxruntime/test/framework/tensorutils_test.cc +++ b/onnxruntime/test/framework/tensorutils_test.cc @@ -61,6 +61,8 @@ constexpr bool TensorProtoElementSizesAreConstexpr() { sizeof(uint8_t), // FLOAT8E8M0 sizeof(uint8_t), // UINT2 sizeof(uint8_t), // INT2 + sizeof(uint8_t), // FLOAT6E2M3 + sizeof(uint8_t), // FLOAT6E3M2 }; for (size_t index = 0; index < expected_sizes.size(); ++index) { diff --git a/onnxruntime/test/framework/type_info_test.cc b/onnxruntime/test/framework/type_info_test.cc index c6119778162e0..afc28d83d9c80 100644 --- a/onnxruntime/test/framework/type_info_test.cc +++ b/onnxruntime/test/framework/type_info_test.cc @@ -10,6 +10,7 @@ #include "core/framework/onnxruntime_map_type_info.h" #include "core/framework/onnxruntime_sequence_type_info.h" #include "core/framework/tensor_type_and_shape.h" +#include "core/framework/tensorprotoutils.h" #include "core/framework/onnxruntime_typeinfo.h" namespace onnxruntime { @@ -47,6 +48,8 @@ constexpr bool TensorElementTypeConversionIsConstexpr() { ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT8E8M0, ONNX_TENSOR_ELEMENT_DATA_TYPE_UINT2, ONNX_TENSOR_ELEMENT_DATA_TYPE_INT2, + ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E2M3, + ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E3M2, }; for (size_t index = 0; index < expected_types.size(); ++index) { @@ -65,6 +68,13 @@ constexpr bool TensorElementTypeConversionIsConstexpr() { static_assert(TensorElementTypeConversionIsConstexpr()); +TEST(TypeInfoTests, CApiElementTypeFromProtoTypeFloat6) { + EXPECT_EQ(utils::CApiElementTypeFromProtoType(ONNX_NAMESPACE::TensorProto_DataType_FLOAT6E2M3), + ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E2M3); + EXPECT_EQ(utils::CApiElementTypeFromProtoType(ONNX_NAMESPACE::TensorProto_DataType_FLOAT6E3M2), + ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E3M2); +} + TEST(TypeInfoTests, TensorProto) { mb::Type tensor_type = {1, 2, 3, 4}; diff --git a/onnxruntime/test/onnx/TestCase.cc b/onnxruntime/test/onnx/TestCase.cc index 86eb9623b4905..a2c54a6904ea8 100644 --- a/onnxruntime/test/onnx/TestCase.cc +++ b/onnxruntime/test/onnx/TestCase.cc @@ -958,6 +958,8 @@ std::unique_ptr> GetBrokenTests(const std::string& provider {"cast_UINT4_to_FLOAT", "Skipped until onnxruntime/cmake/external/onnx points to onnx 1.19 which should include @onnx/onnx/pull/7074"}, {"cast_UINT4_to_FLOAT16", "Skipped until onnxruntime/cmake/external/onnx points to onnx 1.19 which should include @onnx/onnx/pull/7074"}, {"cast_UINT4_to_UINT8", "Skipped until onnxruntime/cmake/external/onnx points to onnx 1.19 which should include @onnx/onnx/pull/7074"}, + {"reduce_max_empty_set_bool", + "ReduceMax does not support bool; mirrors onnx_backend_test_series_filters.jsonc global exclusion"}, // onnx#7959: the materialized node corpus now runs on the CPU EP via the // onnx_test_node_materialized ctest (previously the C++ node runner was QNN-only). ORT has // no per-channel-quantization ConvInteger support on any EP -- the CPU kernel throws and @@ -1091,8 +1093,6 @@ std::unique_ptr> GetBrokenTests(const std::string& provider {"resize_upsample_sizes_nearest_ceil_half_pixel", "Bad onnx test output. Needs test fix."}, {"resize_upsample_sizes_nearest_floor_align_corners", "Bad onnx test output. Needs test fix."}, {"resize_upsample_sizes_nearest_round_prefer_ceil_asymmetric", "Bad onnx test output. Needs test fix."}, - {"bitshift_right_uint16", "BitShift(11) uint16 support not enabled currently"}, - {"bitshift_left_uint16", "BitShift(11) uint16 support not enabled currently"}, {"maxunpool_export_with_output_shape", "Invalid output in ONNX test. See https://github.com/onnx/onnx/issues/2398"}, {"cntk_simple_seg", "Bad onnx test output caused by wrong SAME_UPPER/SAME_LOWER for ConvTranspose"}, diff --git a/onnxruntime/test/onnx/main.cc b/onnxruntime/test/onnx/main.cc index 5740a407c7ef6..1817a92c91e7f 100644 --- a/onnxruntime/test/onnx/main.cc +++ b/onnxruntime/test/onnx/main.cc @@ -908,6 +908,9 @@ select from 'TF8', 'TF16', 'UINT8', 'FLOAT', 'ITENSOR'. \n)"); ORT_TSTR("rotary_embedding_3d_input_expanded"), ORT_TSTR("rotary_embedding_expanded"), ORT_TSTR("rotary_embedding_interleaved_expanded"), + // QNN CPU returns uniform Softmax probabilities for the expanded function's all-negative-infinity rows, + // so its zero-row guard is not preserved. Keep the native Attention case enabled. + ORT_TSTR("attention_4d_causal_nonpad_negative_offset_structural_empty_expanded"), // QNN don't support fmod = 1 ORT_TSTR("mod_mixed_sign_float64"), ORT_TSTR("mod_mixed_sign_float32"), diff --git a/onnxruntime/test/onnx/tensorprotoutils.cc b/onnxruntime/test/onnx/tensorprotoutils.cc index 5c8d8f9fd8f23..c9cbc05b2f391 100644 --- a/onnxruntime/test/onnx/tensorprotoutils.cc +++ b/onnxruntime/test/onnx/tensorprotoutils.cc @@ -652,6 +652,8 @@ ONNXTensorElementDataType CApiElementTypeFromProtoType(int type) { CASE_TYPE(INT4) CASE_TYPE(UINT2) CASE_TYPE(INT2) + CASE_TYPE(FLOAT6E2M3) + CASE_TYPE(FLOAT6E3M2) default: return ONNX_TENSOR_ELEMENT_DATA_TYPE_UNDEFINED; } diff --git a/onnxruntime/test/optimizer/layout_transformation_potentially_added_ops_test.cc b/onnxruntime/test/optimizer/layout_transformation_potentially_added_ops_test.cc index c092c660633b1..c9fa32d569cf7 100644 --- a/onnxruntime/test/optimizer/layout_transformation_potentially_added_ops_test.cc +++ b/onnxruntime/test/optimizer/layout_transformation_potentially_added_ops_test.cc @@ -27,9 +27,15 @@ TEST(LayoutTransformationPotentiallyAddedOpsTests, OpsHaveLatestVersions) { if (prev_domain != it->domain || prev_op_type != it->op_type) { const auto* schema = schema_registry->GetSchema(std::string{it->op_type}, INT_MAX, std::string{it->domain}); ASSERT_NE(schema, nullptr); - EXPECT_EQ(schema->SinceVersion(), it->since_version) - << "A new version for op " << it->op_type << " (" << schema->SinceVersion() - << ") is available. Please update kLayoutTransformationPotentiallyAddedOps to include it."; + const bool is_staged_opset_28_qdq = + it->domain == kOnnxDomain && schema->SinceVersion() == 28 && + (it->op_type == "DequantizeLinear" || it->op_type == "QuantizeLinear"); + if (!is_staged_opset_28_qdq) { + EXPECT_EQ(schema->SinceVersion(), it->since_version) + << "A new version for op " << it->op_type << " (" << schema->SinceVersion() + << ") is available. Please update kLayoutTransformationPotentiallyAddedOps to include it."; + } + // TODO(#32372): Remove this exception when the opset-28 Q/DQ kernels land. prev_domain = it->domain; prev_op_type = it->op_type; } diff --git a/onnxruntime/test/optimizer/qdq_transformer_test.cc b/onnxruntime/test/optimizer/qdq_transformer_test.cc index 64f521d5c6825..aeb9c4eebd3bc 100644 --- a/onnxruntime/test/optimizer/qdq_transformer_test.cc +++ b/onnxruntime/test/optimizer/qdq_transformer_test.cc @@ -5636,7 +5636,10 @@ TEST(QDQTransformerTests, QDQPropagation_GH11605_Opset13) { // test removal of Q->DQ pairs by QDQFinalCleanupTransformer TEST(QDQTransformerTests, QDQFinalCleanupTransformerReportsIntentionalRemoval) { auto& logger = DefaultLoggingManager().DefaultLogger(); - Model model("QDQFinalCleanupRemovalTester", false, logger); + std::unordered_map domain_to_version; + domain_to_version[kOnnxDomain] = 25; + Model model("QDQFinalCleanupRemovalTester", false, ModelMetaData(), PathString(), + IOnnxRuntimeOpSchemaRegistryList(), domain_to_version, {}, logger); Graph& graph = model.MainGraph(); ModelTestBuilder builder(graph); diff --git a/onnxruntime/test/providers/cpu/controlflow/scan_test.cc b/onnxruntime/test/providers/cpu/controlflow/scan_test.cc index 4583b772bb30c..c35a45d590eda 100644 --- a/onnxruntime/test/providers/cpu/controlflow/scan_test.cc +++ b/onnxruntime/test/providers/cpu/controlflow/scan_test.cc @@ -1061,7 +1061,8 @@ TEST(Scan8, NumScanInputsExceedsVariadicInputs) { test.AddOutput("scan_output_2", {1, 2, 1}, {0.f, 0.f}); test.AddOutput("scan_output_3", {1, 2, 1}, {0.f, 0.f}); - test.Run(OpTester::ExpectResult::kExpectFailure, "Invalid 'num_scan_inputs' of 10. Value must be between 1 and 3", + test.Run(OpTester::ExpectResult::kExpectFailure, + "num_scan_inputs (10) plus the sequence_lens input cannot exceed the number of Scan inputs (4).", options.excluded_provider_types); } @@ -1165,7 +1166,7 @@ TEST(Scan9, NumScanInputsExceedsVariadicInputs) { // inferencing error wrapping a shape-inference error, rather than the kernel-level message produced // by the opset 8 test above. test.Run(OpTester::ExpectResult::kExpectFailure, - "[ShapeInferenceError] Mismatch between number of inferred and declared dimensions", + "num_scan_inputs (10) cannot exceed the number of Scan inputs (3).", options.excluded_provider_types); } #endif // !defined(ORT_NO_EXCEPTIONS) diff --git a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc index f3cc72249b764..4348942186d17 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -15,6 +15,7 @@ #include #include #include +#include namespace onnxruntime { namespace test { @@ -4351,6 +4352,50 @@ TEST(ModOpTest, Fmod_float16_mixed_sign) { test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kQnnExecutionProvider}); } +template +void TestFloorModFloatingPointOpset28(const std::vector& x, + const std::vector& y, + const std::vector& expected) { + OpTester test("Mod", 28); + test.AddAttribute("fmod", 0); + const std::vector shape{static_cast(x.size())}; + test.AddInput("X", shape, x); + test.AddInput("Y", shape, y); + test.AddOutput("Z", shape, expected); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kQnnExecutionProvider, kOpenVINOExecutionProvider}); +} + +TEST(ModOpTest, FloorMod_float_mixed_sign_opset28) { + TestFloorModFloatingPointOpset28( + {-4.3f, 7.2f, 5.0f, 4.3f, -7.2f, 8.0f}, + {2.1f, -3.4f, 8.0f, -2.1f, 3.4f, 5.0f}, + {2.0f, -3.0f, 5.0f, -2.0f, 3.0f, 3.0f}); +} + +TEST(ModOpTest, FloorMod_double_mixed_sign_opset28) { + TestFloorModFloatingPointOpset28( + {-4.3, 7.2, 5.0, 4.3, -7.2, 8.0}, + {2.1, -3.4, 8.0, -2.1, 3.4, 5.0}, + {2.0, -3.0, 5.0, -2.0, 3.0, 3.0}); +} + +TEST(ModOpTest, FloorMod_float16_mixed_sign_opset28) { + TestFloorModFloatingPointOpset28( + MakeMLFloat16({-4.3f, 7.2f, 5.0f, 4.3f, -7.2f, 8.0f}), + MakeMLFloat16({2.1f, -3.4f, 8.0f, -2.1f, 3.4f, 5.0f}), + MakeMLFloat16({2.0f, -3.0f, 5.0f, -2.0f, 3.0f, 3.0f})); +} + +TEST(ModOpTest, FloorMod_float_special_values_opset28) { + constexpr auto infinity = std::numeric_limits::infinity(); + constexpr auto nan = std::numeric_limits::quiet_NaN(); + TestFloorModFloatingPointOpset28( + {6.0f, -6.0f, 5.0f, -5.0f, infinity, 1.0f, nan, 1.0f}, + {-3.0f, 3.0f, infinity, -infinity, 2.0f, 0.0f, 2.0f, nan}, + {-0.0f, 0.0f, 5.0f, -5.0f, nan, nan, nan, nan}); +} + #if defined(USE_CUDA) TEST(ModOpTest, Fmod_bfloat16_mixed_sign) { OpTester test("Mod", 13); @@ -4652,6 +4697,93 @@ TEST(BitShiftOpTest, BroadcastXRight_Uint8) { test.Run(); } +template +void TestSignedBitShiftOpset28() { + constexpr T bit_width = static_cast(std::numeric_limits>::digits); + + { + OpTester test("BitShift", 28); + test.AddAttribute("direction", "LEFT"); + test.AddInput("X", {6}, {T{-8}, T{-3}, T{1}, std::numeric_limits::max(), T{-2}, T{2}}); + test.AddInput("Y", {6}, {T{1}, T{2}, T{3}, T{1}, T{-1}, bit_width}); + test.AddOutput("Z", {6}, {T{-16}, T{-12}, T{8}, T{-2}, T{0}, T{0}}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kOpenVINOExecutionProvider}); + } + + { + OpTester test("BitShift", 28); + test.AddAttribute("direction", "RIGHT"); + test.AddInput("X", {6}, {T{-8}, T{-3}, T{1}, T{-2}, T{-2}, T{2}}); + test.AddInput("Y", {6}, {T{1}, T{2}, T{3}, T{-1}, bit_width, bit_width}); + test.AddOutput("Z", {6}, {T{-4}, T{-1}, T{0}, T{-1}, T{-1}, T{0}}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kOpenVINOExecutionProvider}); + } + + { + OpTester test("BitShift", 28); + test.AddAttribute("direction", "RIGHT"); + test.AddInput("X", {1}, {T{-8}}); + test.AddInput("Y", {4}, {T{1}, T{2}, bit_width, T{-1}}); + test.AddOutput("Z", {4}, {T{-4}, T{-2}, T{-1}, T{-1}}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kOpenVINOExecutionProvider}); + } + + { + OpTester test("BitShift", 28); + test.AddAttribute("direction", "LEFT"); + test.AddInput("X", {2}, {std::numeric_limits::max(), T{-8}}); + test.AddInput("Y", {1}, {T{1}}); + test.AddOutput("Z", {2}, {T{-2}, T{-16}}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kOpenVINOExecutionProvider}); + } +} + +TEST(BitShiftOpTest, SignedInt8Opset28) { + TestSignedBitShiftOpset28(); +} + +TEST(BitShiftOpTest, SignedInt16Opset28) { + TestSignedBitShiftOpset28(); +} + +TEST(BitShiftOpTest, SignedInt32Opset28) { + TestSignedBitShiftOpset28(); +} + +TEST(BitShiftOpTest, SignedInt64Opset28) { + TestSignedBitShiftOpset28(); +} + +template +void TestUnsignedBitShiftRegistration(int opset) { + OpTester test("BitShift", opset); + test.AddAttribute("direction", "LEFT"); + test.AddInput("X", {2}, {T{1}, T{3}}); + test.AddInput("Y", {2}, {T{1}, T{2}}); + test.AddOutput("Z", {2}, {T{2}, T{12}}); + test.Run(); +} + +TEST(BitShiftOpTest, Uint16Opset11) { + TestUnsignedBitShiftRegistration(11); +} + +TEST(BitShiftOpTest, UnsignedTypesOpset28) { + TestUnsignedBitShiftRegistration(28); + TestUnsignedBitShiftRegistration(28); + TestUnsignedBitShiftRegistration(28); + TestUnsignedBitShiftRegistration(28); +} + +TEST(BitShiftOpTest, RightShiftByBitWidth_Uint64Opset28) { + OpTester test("BitShift", 28); + test.AddAttribute("direction", "RIGHT"); + test.AddInput("X", {4}, {1000, 255, 1, 42}); + test.AddInput("Y", {4}, {64, 64, 64, 64}); + test.AddOutput("Z", {4}, {0, 0, 0, 0}); + test.Run(); +} + // Test that shift amounts >= bit width produce 0 (not undefined behavior). // DirectML EP has the same hardware-level shift masking behavior, so skip these tests for DML. TEST(BitShiftOpTest, RightShiftByBitWidth_Uint64) { diff --git a/onnxruntime/test/providers/cpu/tensor/compress_op.test.cc b/onnxruntime/test/providers/cpu/tensor/compress_op.test.cc index 495fea5735b32..36c3f61a59630 100644 --- a/onnxruntime/test/providers/cpu/tensor/compress_op.test.cc +++ b/onnxruntime/test/providers/cpu/tensor/compress_op.test.cc @@ -1,6 +1,9 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include +#include + #include "gtest/gtest.h" #include "test/providers/provider_test_utils.h" @@ -166,8 +169,13 @@ TEST(CompressTest, Compress_3dims_neg_axis) { // inputs to {0, 1}) and asserts the output is sized by truthiness. TEST(CompressTest, Compress_cuda_non_canonical_bool_condition) { // Build: output = Compress(input, condition, axis=0) - auto model = std::make_unique("compress_non_canonical_bool", false, - DefaultLoggingManager().DefaultLogger()); + const std::unordered_map domain_to_version{{kOnnxDomain, 27}}; + auto model = std::make_unique( + "compress_non_canonical_bool", false, ModelMetaData(), PathString(), IOnnxRuntimeOpSchemaRegistryList(), + domain_to_version, std::vector{}, DefaultLoggingManager().DefaultLogger()); + ASSERT_EQ(model->ToProto().opset_import_size(), 1); + EXPECT_EQ(model->ToProto().opset_import(0).domain(), kOnnxDomain); + EXPECT_EQ(model->ToProto().opset_import(0).version(), 27); onnxruntime::Graph& graph = model->MainGraph(); ONNX_NAMESPACE::TypeProto tensor_float; diff --git a/onnxruntime/test/providers/dml_onnx_tensor_wrapper_test.cc b/onnxruntime/test/providers/dml_onnx_tensor_wrapper_test.cc index c9d0d190e6ff9..6855b309f3e30 100644 --- a/onnxruntime/test/providers/dml_onnx_tensor_wrapper_test.cc +++ b/onnxruntime/test/providers/dml_onnx_tensor_wrapper_test.cc @@ -59,10 +59,9 @@ enum class ValueAttributePayload { // IMLOperatorAttributes1::GetTensorAttribute, which wraps the raw AttributeProto tensor in an // OnnxTensorWrapper. // -// Attribute tensors are not graph initializers, so none of the framework-level guards apply to them: -// Graph::ConvertInitializersIntoOrtValues never sees them, and ONNX shape inference for -// ConstantOfShape only reads the attribute's data_type (to type the output), never its data. That -// leaves OnnxTensorWrapper as the only place the payload can be validated. +// Attribute tensors are not graph initializers, so Graph::ConvertInitializersIntoOrtValues never +// sees them. Older ONNX versions also did not validate byte-aligned raw_data sizes in check_tensor, +// leaving OnnxTensorWrapper as the only place the payload could be validated. std::string BuildConstantOfShapeModel(ValueAttributePayload payload) { ONNX_NAMESPACE::ModelProto model; model.set_ir_version(ONNX_NAMESPACE::Version::IR_VERSION); @@ -113,8 +112,8 @@ std::string BuildConstantOfShapeModel(ValueAttributePayload payload) { return serialized; } -// Loads |serializedModel| into a session that has only the DML EP registered, and returns the status -// of Initialize(). |skipped| is set when no DirectML device is available. +// Loads |serializedModel| into a session that has only the DML EP registered, and returns the first +// failure from Load() or Initialize(). |skipped| is set when no DirectML device is available. // // Constant folding is disabled so the ConstantOfShape node survives graph optimization and is handed // to the DML EP. Without this, the CPU EP would evaluate the node during Level1 optimization and the @@ -166,15 +165,19 @@ TEST(DmlOnnxTensorWrapperTest, TruncatedRawDataAttributeTensorIsRejected) { GTEST_SKIP() << "DirectML execution provider is not available on this machine."; } - ASSERT_FALSE(status.IsOK()) << "Session initialization should reject a tensor whose declared shape " - "exceeds the data it carries."; - - // The detailed message raised inside the constructor does not survive: OnnxTensorWrapper is - // constructed behind the IMLOperatorAttributes1 COM boundary, whose ORT_CATCH_RETURN reduces the - // exception to a bare HRESULT before it is rethrown. 0x80070057 is E_INVALIDARG, so this asserts the - // model was rejected as malformed rather than failing for an unrelated reason. - EXPECT_NE(status.ErrorMessage().find("80070057"), std::string::npos) - << "Unexpected failure reason: " << status.ErrorMessage(); + ASSERT_FALSE(status.IsOK()) << "Model loading or session initialization should reject a tensor whose declared " + "shape exceeds the data it carries."; + + // ONNX 1.23 rejects the malformed attribute during Graph::Resolve. With older ONNX versions it + // reaches OnnxTensorWrapper, where the detailed constructor exception is reduced to E_INVALIDARG + // (0x80070057) at the IMLOperatorAttributes1 COM boundary. + const std::string& errorMessage = status.ErrorMessage(); + EXPECT_TRUE( + errorMessage.find("80070057") != std::string::npos || + errorMessage.find( + "raw_data size (1 bytes) is too small for the declared shape and type (8 bytes required)") != + std::string::npos) + << "Unexpected failure reason: " << errorMessage; } // The constructor's typed-field branch has the same defect in principle: UnpackTensor sizes the buffer diff --git a/onnxruntime/test/python/onnx_backend_test_series.py b/onnxruntime/test/python/onnx_backend_test_series.py index 62159721a2108..f12af65a6946f 100644 --- a/onnxruntime/test/python/onnx_backend_test_series.py +++ b/onnxruntime/test/python/onnx_backend_test_series.py @@ -17,6 +17,12 @@ import onnx.defs import onnxruntime.backend as backend # pylint: disable=consider-using-from-import +from onnxruntime.backend.backend import OnnxRuntimeBackend + +if os.getenv("ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS") == "1" and OnnxRuntimeBackend.allowReleasedOpsetsOnly: + raise RuntimeError( + "onnx_backend_test_series requires ALLOW_RELEASED_ONNX_OPSET_ONLY=0 before importing onnxruntime.backend." + ) pytest_plugins = ("onnx.backend.test.report",) diff --git a/onnxruntime/test/python/onnx_node_test_equivalence_test.py b/onnxruntime/test/python/onnx_node_test_equivalence_test.py index 1daa2f4d1eac7..0516f2c2c28a8 100644 --- a/onnxruntime/test/python/onnx_node_test_equivalence_test.py +++ b/onnxruntime/test/python/onnx_node_test_equivalence_test.py @@ -30,9 +30,8 @@ backstop. Version parity: the materialized model.onnx opset/IR is stamped from the -compiled onnx schema registry, so this test asserts the installed onnx equals -the ``cmake/deps.txt`` pin (compared on the release base, so a pre-release rcN/.dev -wheel of the pinned tag is accepted). numpy parity is NOT asserted -- the core +compiled onnx schema registry, so this test asserts the installed onnx exactly +equals the ``cmake/deps.txt`` pin. numpy parity is NOT asserted -- the core treats numpy as an informational soft-warning, and recomputed float outputs that differ by within a small ULP band are accepted as Class A numpy-gen skew. """ @@ -44,6 +43,7 @@ import sys import tempfile import unittest +from unittest import mock # Make the sibling materializer importable regardless of CWD. _TOOLS_PYTHON = os.path.normpath(os.path.join(os.path.dirname(__file__), "..", "..", "..", "tools", "python")) @@ -62,14 +62,14 @@ def _repo_root() -> str: def _expected_onnx_version() -> str | None: - """Parse the onnx pin (archive tag) from cmake/deps.txt.""" + """Derive the expected wheel version from the archive pin in cmake/deps.txt.""" deps = os.path.join(_repo_root(), "cmake", "deps.txt") if not os.path.isfile(deps): return None with open(deps) as f: for line in f: if line.startswith("onnx;"): - m = re.search(r"/tags/v([0-9]+\.[0-9]+\.[0-9]+)\.zip", line) + m = re.search(r"/v([0-9]+\.[0-9]+\.[0-9]+)\.zip;", line) if m: return m.group(1) return None @@ -123,6 +123,14 @@ def _find_oracle_node_dir() -> str | None: class NodeTestEquivalence(unittest.TestCase): + def test_expected_version_comes_from_archive_pin(self) -> None: + deps = "# ONNX_VERSION=1.23.0\nonnx;https://github.com/onnx/onnx/archive/refs/tags/v1.23.1.zip;archive_hash\n" + with ( + mock.patch("os.path.isfile", return_value=True), + mock.patch("builtins.open", mock.mock_open(read_data=deps)), + ): + self.assertEqual(_expected_onnx_version(), "1.23.1") + def test_materialized_equals_on_disk_oracle(self) -> None: oracle = _find_oracle_node_dir() if oracle is None: @@ -138,12 +146,11 @@ def test_materialized_equals_on_disk_oracle(self) -> None: except ImportError as exc: # pragma: no cover self.skipTest(f"onnx not importable: {exc}") - if expected_onnx and mat._release_base(onnx.__version__) != mat._release_base(expected_onnx): + if expected_onnx and onnx.__version__ != expected_onnx: self.skipTest( f"installed onnx=={onnx.__version__} != cmake/deps.txt pin " f"{expected_onnx}; cannot certify equivalence under version skew. " - f"Install the pinned onnx to run this check. (A pre-release rcN/.dev " - f"of the same release base is accepted.)" + "Install the pinned onnx to run this check." ) with tempfile.TemporaryDirectory(prefix="ort_node_mat_") as tmp: diff --git a/onnxruntime/test/python/onnxruntime_test_python_backend.py b/onnxruntime/test/python/onnxruntime_test_python_backend.py index bb83f6d36011f..e06d4f997514c 100644 --- a/onnxruntime/test/python/onnxruntime_test_python_backend.py +++ b/onnxruntime/test/python/onnxruntime_test_python_backend.py @@ -5,13 +5,17 @@ import os import tempfile import unittest +from types import SimpleNamespace import numpy as np from helper import get_name from numpy.testing import assert_allclose +from onnx import TensorProto, helper import onnxruntime as onnxrt import onnxruntime.backend as backend +from onnxruntime.backend.backend import OnnxRuntimeBackend +from onnxruntime.backend.backend_rep import OnnxRuntimeBackendRep class TestBackend(unittest.TestCase): @@ -23,6 +27,72 @@ def test_run_model(self): output_expected = np.array([[1.0, 4.0], [9.0, 16.0], [25.0, 36.0]], dtype=np.float32) np.testing.assert_allclose(res[0], output_expected, rtol=1e-05, atol=1e-08) + def test_prepare_released_opset28_model_when_policy_enabled(self): + model = self._make_opset28_model() + original_policy = OnnxRuntimeBackend.allowReleasedOpsetsOnly + try: + OnnxRuntimeBackend.allowReleasedOpsetsOnly = True + rep = backend.prepare(model) + finally: + OnnxRuntimeBackend.allowReleasedOpsetsOnly = original_policy + + result = rep.run(np.array([1.0], dtype=np.float32)) + np.testing.assert_array_equal(result[0], np.array([1.0], dtype=np.float32)) + + def test_prepare_model_bytes_rejects_unreleased_opset29_when_policy_enabled(self): + model = self._make_opset28_model() + model.opset_import[0].version = 29 + original_policy = OnnxRuntimeBackend.allowReleasedOpsetsOnly + try: + OnnxRuntimeBackend.allowReleasedOpsetsOnly = True + with self.assertRaisesRegex(onnxrt.capi.onnxruntime_pybind11_state.Fail, "Opset 29 is under development"): + backend.prepare(model.SerializeToString()) + finally: + OnnxRuntimeBackend.allowReleasedOpsetsOnly = original_policy + + def test_prepare_model_bytes_allows_opset28_when_policy_disabled(self): + model = self._make_opset28_model() + + original_policy = OnnxRuntimeBackend.allowReleasedOpsetsOnly + try: + OnnxRuntimeBackend.allowReleasedOpsetsOnly = False + rep = backend.prepare(model.SerializeToString()) + finally: + OnnxRuntimeBackend.allowReleasedOpsetsOnly = original_policy + + result = rep.run(np.array([1.0], dtype=np.float32)) + np.testing.assert_array_equal(result[0], np.array([1.0], dtype=np.float32)) + + def test_prepare_model_path_allows_opset28_when_policy_disabled(self): + model = self._make_opset28_model() + with tempfile.TemporaryDirectory() as tmpdir: + model_path = os.path.join(tmpdir, "opset28.onnx") + with open(model_path, "wb") as model_file: + model_file.write(model.SerializeToString()) + + original_policy = OnnxRuntimeBackend.allowReleasedOpsetsOnly + try: + OnnxRuntimeBackend.allowReleasedOpsetsOnly = False + rep = backend.prepare(model_path) + finally: + OnnxRuntimeBackend.allowReleasedOpsetsOnly = original_policy + + result = rep.run(np.array([1.0], dtype=np.float32)) + np.testing.assert_array_equal(result[0], np.array([1.0], dtype=np.float32)) + + @staticmethod + def _make_opset28_model(): + model = helper.make_model( + helper.make_graph( + [helper.make_node("Identity", ["X"], ["Y"])], + "opset28", + [helper.make_tensor_value_info("X", TensorProto.FLOAT, [1])], + [helper.make_tensor_value_info("Y", TensorProto.FLOAT, [1])], + ), + opset_imports=[helper.make_opsetid("", 28)], + ) + return model + def test_allocation_plan_works_with_only_execute_path_to_fetches_option(self): """ (inp0) (inp1) @@ -66,6 +136,54 @@ def test_allocation_plan_works_with_only_execute_path_to_fetches_option(self): assert_allclose(session_run_results[0], -(inp0 - inp1)) +class TestBackendRepInputNormalization(unittest.TestCase): + class _Session: + def __init__(self, input_names): + self._inputs = [SimpleNamespace(name=name) for name in input_names] + self.received_inputs = None + + def get_inputs(self): + return self._inputs + + def run(self, _, inputs, __): + self.received_inputs = inputs + return [inputs] + + def test_run_normalizes_numpy_scalar_inputs_to_zero_dimensional_arrays(self): + session = self._Session(["first", "second"]) + rep = OnnxRuntimeBackendRep(session) + + rep.run([np.int32(3), np.float32(4.5)]) + + self.assertEqual(session.received_inputs["first"].shape, ()) + self.assertEqual(session.received_inputs["second"].shape, ()) + self.assertEqual(session.received_inputs["first"].dtype, np.dtype(np.int32)) + self.assertEqual(session.received_inputs["second"].dtype, np.dtype(np.float32)) + + def test_run_normalizes_single_numpy_scalar_to_zero_dimensional_array(self): + session = self._Session(["input"]) + rep = OnnxRuntimeBackendRep(session) + + rep.run(np.int64(7)) + + self.assertEqual(session.received_inputs["input"].shape, ()) + self.assertEqual(session.received_inputs["input"].dtype, np.dtype(np.int64)) + + def test_run_preserves_non_scalar_inputs(self): + values = ( + (np.array([1], dtype=np.int32),), + {"key": np.array([2], dtype=np.int32)}, + ) + for value in values: + with self.subTest(value=value): + session = self._Session(["input"]) + rep = OnnxRuntimeBackendRep(session) + + rep.run(value) + + self.assertIs(session.received_inputs["input"], value) + + class TestBackendKwargsAllowlist(unittest.TestCase): """Tests that the SessionOptions/RunOptions kwargs allowlist correctly blocks dangerous attributes and allows safe ones, preventing arbitrary file writes diff --git a/onnxruntime/test/python/requirements.txt b/onnxruntime/test/python/requirements.txt index 296ac83f3b6b5..7d5b9bae47341 100644 --- a/onnxruntime/test/python/requirements.txt +++ b/onnxruntime/test/python/requirements.txt @@ -1,3 +1,3 @@ -onnx==1.22.0 +onnx==1.23.1 pytest onnx-ir diff --git a/onnxruntime/test/python/transformers/test_cuda_plugin_ep.py b/onnxruntime/test/python/transformers/test_cuda_plugin_ep.py index 95df52c3f6e61..670766f1c8fb0 100644 --- a/onnxruntime/test/python/transformers/test_cuda_plugin_ep.py +++ b/onnxruntime/test/python/transformers/test_cuda_plugin_ep.py @@ -30,8 +30,8 @@ EP_GRAPH_ASSIGNMENT_CONFIG_KEY = "session.record_ep_graph_assignment_info" NO_CUDNN_PLUGIN_TEST = os.getenv("ORT_TEST_CUDA_PLUGIN_NO_CUDNN", "").upper() in {"1", "ON", "TRUE", "YES"} requires_cudnn = unittest.skipIf(NO_CUDNN_PLUGIN_TEST, "test requires cuDNN-backed CUDA plugin kernels") -# Use the latest released ai.onnx opset so the model builders stay current as ONNX releases new opsets. -DEFAULT_ONNX_OPSET = max(v for (d, v) in helper.OP_SET_ID_VERSION_MAP if d == "ai.onnx") +# CUDA plugin runtime currently supports released ai.onnx opsets through 27. +DEFAULT_ONNX_OPSET = 27 def _make_released_opset_model(graph, producer_name="onnx-example"): diff --git a/onnxruntime/test/python/transformers/test_moe_cuda.py b/onnxruntime/test/python/transformers/test_moe_cuda.py index 5c92b5e197e4f..11731f295fad6 100644 --- a/onnxruntime/test/python/transformers/test_moe_cuda.py +++ b/onnxruntime/test/python/transformers/test_moe_cuda.py @@ -1734,7 +1734,11 @@ def create_sparse_mixer_onnx_graph( ], ) - return helper.make_model(graph, producer_name="MoE_Model") + return helper.make_model( + graph, + producer_name="MoE_Model", + opset_imports=[helper.make_opsetid("", 27), helper.make_opsetid("com.microsoft", 1)], + ) class TestSparseMixer(unittest.TestCase): @@ -1777,6 +1781,8 @@ def test_sparse_mixer_functional(self, onnx_dtype): fc2_bias, onnx_dtype, ) + opset_imports = {opset_import.domain: opset_import.version for opset_import in onnx_model.opset_import} + self.assertEqual(opset_imports, {"": 27, "com.microsoft": 1}) sess_options = onnxruntime.SessionOptions() sess = onnxruntime.InferenceSession(onnx_model.SerializeToString(), sess_options, providers=get_ort_provider()) diff --git a/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc b/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc index a10eb97496186..cf40b161dd008 100644 --- a/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc +++ b/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc @@ -55,6 +55,18 @@ "^test_attention_(23|24)_.*qk_matmul_output_mode3", // #28994 CUDA NOT_IMPLEMENTED mode>kQK (attention.cc:1477); #27712 to support // TODO: support qk_matmul_output modes beyond kQK in Attention-cuda (see issue #27712) // Tests combining qk_matmul with softcap need unfused-path qk_matmul support (deferred). + "^test_attention_3d_local_window_cuda$", // exits with 0xDEAD in the CUDA backend test + "^test_attention_bidirectional_window_cuda$", // exits with 0xDEAD in the CUDA backend test + "^test_attention_local_window_cuda$", // exits with 0xDEAD in the CUDA backend test + "^test_attention_local_window_ext_cache_float16_mask_cuda$", // exits with 0xDEAD in the CUDA backend test + "^test_attention_local_window_ext_cache_rank2_mask_cuda$", // exits with 0xDEAD in the CUDA backend test + "^test_attention_local_window_ext_cache_rank3_head_mask_cuda$", // exits with 0xDEAD in the CUDA backend test + "^test_attention_local_window_ext_cache_rank4_batch_mask_cuda$", // exits with 0xDEAD in the CUDA backend test + "^test_attention_local_window_gqa_rank4_mask_cuda$", // exits with an access violation in the CUDA backend test + "^test_attention_local_window_gqa_rank4_mask_expanded_cuda$", // expanded Attention output mismatch + "^test_attention_local_window_rank1_boolean_mask_cuda$", // exits with an access violation in the CUDA backend test + "^test_attention_local_window_with_past_cuda$", // exits with an access violation in the CUDA backend test + "^test_attention_4d_causal_nonpad_negative_offset_structural_empty_expanded_cuda$", // expanded Attention output mismatch "^test_attention_3d_with_past_and_present_qk_matmul_softcap_cuda", // qk_matmul modes beyond kQK not supported "^test_attention_4d_with_qk_matmul_softcap_cuda", // qk_matmul modes beyond kQK not supported "^test_attention_3d_with_past_and_present_qk_matmul_bias_cuda", // QK matmul + bias not supported in Attention-cuda @@ -84,6 +96,209 @@ "^test_castlike_UINT2_to*", // ORT does not support uint2 "^test_cast_INT2_to*", // ORT does not support int2 "^test_cast_UINT2_to*", // ORT does not support uint2 + // TODO(#32378): Remove the exact BF16 Attention failures when the function dependencies support BF16. + "^test_attention_3d_causal_bf16_cpu$", + "^test_attention_3d_causal_bf16_expanded_cpu$", + "^test_attention_4d_attn_mask_causal_bf16_cpu$", + "^test_attention_4d_attn_mask_causal_bf16_expanded_cpu$", + "^test_attention_4d_causal_bf16_cpu$", + "^test_attention_4d_causal_bf16_expanded_cpu$", + "^test_attention_4d_causal_padded_kv_bf16_cpu$", + "^test_attention_4d_causal_padded_kv_bf16_expanded_cpu$", + "^test_attention_4d_padded_kv_bf16_cpu$", + "^test_attention_4d_padded_kv_bf16_expanded_cpu$", + "^test_attention_3d_causal_bf16_cuda$", + "^test_attention_3d_causal_bf16_expanded_cuda$", + "^test_attention_4d_attn_mask_causal_bf16_cuda$", + "^test_attention_4d_attn_mask_causal_bf16_expanded_cuda$", + "^test_attention_4d_causal_bf16_cuda$", + "^test_attention_4d_causal_bf16_expanded_cuda$", + "^test_attention_4d_causal_padded_kv_bf16_cuda$", + "^test_attention_4d_causal_padded_kv_bf16_expanded_cuda$", + "^test_attention_4d_padded_kv_bf16_cuda$", + "^test_attention_4d_padded_kv_bf16_expanded_cuda$", + "^test_attention_local_window_ext_cache_float16_mask_cpu$", + // TODO: Remove the exact Cast(28) failures when CPU Cast(28) is available in all builds. + "^test_cast_DOUBLE_to_FLOAT_cpu$", + "^test_celu_bfloat16_cpu$", + "^test_celu_bfloat16_expanded_cpu$", + "^test_celu_cpu$", + "^test_celu_expanded_cpu$", + "^test_celu_float16_cpu$", + "^test_celu_float16_expanded_cpu$", + // TODO(#32381): Remove the exact log-reduction failures when their Cast(28) dependency is available. + "^test_reduce_log_sum_exp_default_axes_keepdims_example_cpu$", + "^test_reduce_log_sum_exp_default_axes_keepdims_example_expanded_cpu$", + "^test_reduce_log_sum_exp_default_axes_keepdims_random_cpu$", + "^test_reduce_log_sum_exp_default_axes_keepdims_random_expanded_cpu$", + "^test_reduce_log_sum_exp_do_not_keepdims_example_cpu$", + "^test_reduce_log_sum_exp_do_not_keepdims_example_expanded_cpu$", + "^test_reduce_log_sum_exp_do_not_keepdims_random_cpu$", + "^test_reduce_log_sum_exp_do_not_keepdims_random_expanded_cpu$", + "^test_reduce_log_sum_exp_empty_set_cpu$", + "^test_reduce_log_sum_exp_empty_set_expanded_cpu$", + "^test_reduce_log_sum_exp_keepdims_example_cpu$", + "^test_reduce_log_sum_exp_keepdims_example_expanded_cpu$", + "^test_reduce_log_sum_exp_keepdims_random_cpu$", + "^test_reduce_log_sum_exp_keepdims_random_expanded_cpu$", + "^test_reduce_log_sum_exp_negative_axes_keepdims_example_cpu$", + "^test_reduce_log_sum_exp_negative_axes_keepdims_example_expanded_cpu$", + "^test_reduce_log_sum_exp_negative_axes_keepdims_random_cpu$", + "^test_reduce_log_sum_exp_negative_axes_keepdims_random_expanded_cpu$", + // TODO: Remove the exact Compress(28) failures when the existing kernel supports opset 28. + "^test_compress_0_cpu$", + "^test_compress_1_cpu$", + "^test_compress_bfloat16_cpu$", + "^test_compress_default_axis_cpu$", + "^test_compress_negative_axis_cpu$", + // TODO: Remove the exact Einsum(28) failures when the existing kernel supports opset 28. + "^test_einsum_batch_diagonal_cpu$", + "^test_einsum_batch_matmul_bfloat16_cpu$", + "^test_einsum_batch_matmul_cpu$", + "^test_einsum_inner_prod_cpu$", + "^test_einsum_scalar_cpu$", + "^test_einsum_sum_bfloat16_cpu$", + "^test_einsum_sum_cpu$", + "^test_einsum_transpose_bfloat16_cpu$", + "^test_einsum_transpose_cpu$", + // TODO: Remove the exact OneHot(28) failures when the existing kernel supports opset 28. + "^test_onehot_negative_indices_cpu$", + "^test_onehot_out_of_range_indices_cpu$", + "^test_onehot_with_axis_cpu$", + "^test_onehot_with_bfloat16_values_cpu$", + "^test_onehot_with_negative_axis_cpu$", + "^test_onehot_without_axis_cpu$", + // TODO: Remove this test when ReduceMax(28) supports an empty bool input. + "^test_reduce_max_empty_set_bool_cpu$", + "^test_reduce_max_empty_set_bool_cuda$", + // TODO(#32381): Remove the exact Swish(24) function-dependency failures when SwiGLU is supported. + "^test_swiglu_alpha_cpu$", + "^test_swiglu_alpha_expanded_cpu$", + "^test_swiglu_cpu$", + "^test_swiglu_expanded_cpu$", + "^test_swiglu_float16_cpu$", + "^test_swiglu_float16_expanded_cpu$", + // TODO: Remove the exact Unique(28) failures when the existing kernel supports opset 28. + "^test_unique_bfloat16_sorted_without_axis_cpu$", + "^test_unique_length_1_cpu$", + "^test_unique_not_sorted_without_axis_cpu$", + "^test_unique_sorted_with_axis_3d_cpu$", + "^test_unique_sorted_with_axis_cpu$", + "^test_unique_sorted_with_negative_axis_cpu$", + "^test_unique_sorted_without_axis_cpu$", + // TODO: Remove these exact CUDA skips when fallback nodes have providers and Swish(24) has a CPU kernel. + "^test_cast_DOUBLE_to_FLOAT_cuda$", + "^test_celu_bfloat16_cuda$", + "^test_celu_bfloat16_expanded_cuda$", + "^test_celu_cuda$", + "^test_celu_expanded_cuda$", + "^test_celu_float16_cuda$", + "^test_celu_float16_expanded_cuda$", + "^test_compress_0_cuda$", + "^test_compress_1_cuda$", + "^test_compress_bfloat16_cuda$", + "^test_compress_default_axis_cuda$", + "^test_compress_negative_axis_cuda$", + "^test_einsum_batch_diagonal_cuda$", + "^test_einsum_batch_matmul_bfloat16_cuda$", + "^test_einsum_batch_matmul_cuda$", + "^test_einsum_inner_prod_cuda$", + "^test_einsum_scalar_cuda$", + "^test_einsum_sum_bfloat16_cuda$", + "^test_einsum_sum_cuda$", + "^test_einsum_transpose_bfloat16_cuda$", + "^test_einsum_transpose_cuda$", + "^test_onehot_negative_indices_cuda$", + "^test_onehot_out_of_range_indices_cuda$", + "^test_onehot_with_axis_cuda$", + "^test_onehot_with_bfloat16_values_cuda$", + "^test_onehot_with_negative_axis_cuda$", + "^test_onehot_without_axis_cuda$", + "^test_reduce_log_sum_exp_default_axes_keepdims_example_cuda$", + "^test_reduce_log_sum_exp_default_axes_keepdims_example_expanded_cuda$", + "^test_reduce_log_sum_exp_default_axes_keepdims_random_cuda$", + "^test_reduce_log_sum_exp_default_axes_keepdims_random_expanded_cuda$", + "^test_reduce_log_sum_exp_do_not_keepdims_example_cuda$", + "^test_reduce_log_sum_exp_do_not_keepdims_example_expanded_cuda$", + "^test_reduce_log_sum_exp_do_not_keepdims_random_cuda$", + "^test_reduce_log_sum_exp_do_not_keepdims_random_expanded_cuda$", + "^test_reduce_log_sum_exp_keepdims_example_cuda$", + "^test_reduce_log_sum_exp_keepdims_example_expanded_cuda$", + "^test_reduce_log_sum_exp_keepdims_random_cuda$", + "^test_reduce_log_sum_exp_keepdims_random_expanded_cuda$", + "^test_reduce_log_sum_exp_negative_axes_keepdims_example_cuda$", + "^test_reduce_log_sum_exp_negative_axes_keepdims_example_expanded_cuda$", + "^test_reduce_log_sum_exp_negative_axes_keepdims_random_cuda$", + "^test_reduce_log_sum_exp_negative_axes_keepdims_random_expanded_cuda$", + "^test_swiglu_alpha_cuda$", + "^test_swiglu_alpha_expanded_cuda$", + "^test_swiglu_cuda$", + "^test_swiglu_expanded_cuda$", + "^test_swiglu_float16_cuda$", + "^test_swiglu_float16_expanded_cuda$", + "^test_unique_bfloat16_sorted_without_axis_cuda$", + "^test_unique_length_1_cuda$", + "^test_unique_not_sorted_without_axis_cuda$", + "^test_unique_sorted_with_axis_3d_cuda$", + "^test_unique_sorted_with_axis_cuda$", + "^test_unique_sorted_with_negative_axis_cuda$", + "^test_unique_sorted_without_axis_cuda$", + // TODO(#32361): Remove these exact skips when the opset-28 BitShift kernels land. + ["opset28", "^test_bitshift_left_int16(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int32(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int32_negative_shift(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int32_overflow(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int32_shift_ge_width(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int64(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int8(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int8_negative_shift(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int8_overflow(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_int8_shift_ge_width(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_uint32(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_uint64(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_left_uint8(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int16(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int32(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int32_negative_input(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int32_negative_shift(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int32_shift_ge_width(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int64(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int8(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int8_negative_input(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int8_negative_shift(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_int8_shift_ge_width(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_uint32(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_uint64(_(cpu|cuda))?$"], + ["opset28", "^test_bitshift_right_uint8(_(cpu|cuda))?$"], + "^test_mod_mixed_sign_int16_cuda$", + "^test_mod_mixed_sign_int8_cuda$", + "^test_mod_uint16_cuda$", + "^test_mod_uint8_cuda$", + // TODO(#32372): Remove these exact skips when the opset-28 Q/DQ kernels land. + "^test_dequantizelinear(_(cpu|cuda))?$", + "^test_dequantizelinear_axis(_(cpu|cuda))?$", + "^test_dequantizelinear_e4m3fn(_(cpu|cuda))?$", + "^test_dequantizelinear_e4m3fn_float16(_(cpu|cuda))?$", + "^test_dequantizelinear_e4m3fn_zero_point(_(cpu|cuda))?$", + "^test_dequantizelinear_e5m2(_(cpu|cuda))?$", + "^test_dequantizelinear_float4e2m1(_(cpu|cuda))?$", + "^test_dequantizelinear_int16(_(cpu|cuda))?$", + "^test_dequantizelinear_int2(_(cpu|cuda))?$", + "^test_dequantizelinear_uint16(_(cpu|cuda))?$", + "^test_dequantizelinear_uint2(_(cpu|cuda))?$", + "^test_quantizelinear(_(cpu|cuda))?$", + "^test_quantizelinear_axis(_(cpu|cuda))?$", + "^test_quantizelinear_e4m3fn(_(cpu|cuda))?$", + "^test_quantizelinear_e5m2(_(cpu|cuda))?$", + "^test_quantizelinear_float4e2m1(_(cpu|cuda))?$", + "^test_quantizelinear_int16(_(cpu|cuda))?$", + "^test_quantizelinear_int2(_(cpu|cuda))?$", + "^test_quantizelinear_uint16(_(cpu|cuda))?$", + "^test_quantizelinear_uint2(_(cpu|cuda))?$", + // TODO(#32378): Remove these exact skips when opset-28 ReverseSequence support lands. + ["opset28", "^test_reversesequence_batch(_(cpu|cuda))?$"], + ["opset28", "^test_reversesequence_bfloat16(_(cpu|cuda))?$"], + ["opset28", "^test_reversesequence_time(_(cpu|cuda))?$"], "^test_convinteger_with_padding*", // ORT does not support per-channel quantization for ConvInteger "^test_cast_e8m0_*", // ORT does not support float8e8m0 "^test_batchnorm_epsilon_training_mode", diff --git a/onnxruntime/test/testdata/onnx_backend_test_series_overrides.jsonc b/onnxruntime/test/testdata/onnx_backend_test_series_overrides.jsonc index f4da6b0d37a47..b56b0e5a97720 100644 --- a/onnxruntime/test/testdata/onnx_backend_test_series_overrides.jsonc +++ b/onnxruntime/test/testdata/onnx_backend_test_series_overrides.jsonc @@ -21,5 +21,8 @@ }, // Key: str, the name of the test as defined by ONNX without any device suffix. // Val: float, max relative difference between expected and actual. - "rtol_overrides": {} + "rtol_overrides": { + "test_attention_4d_causal_fp16": 0.002, + "test_attention_4d_causal_fp16_expanded": 0.002 + } } diff --git a/onnxruntime/test/unittest_util/graph_transform_test_builder.cc b/onnxruntime/test/unittest_util/graph_transform_test_builder.cc index 325f820343091..4e67113a4d76a 100644 --- a/onnxruntime/test/unittest_util/graph_transform_test_builder.cc +++ b/onnxruntime/test/unittest_util/graph_transform_test_builder.cc @@ -169,9 +169,7 @@ void TransformerTester(const std::function& buil SessionOptions session_options; session_options.graph_optimization_level = transformer ? baseline_level : level; // Mirror the model build's shape/type-inference strictness onto the session so that - // loading the serialized model applies the same options. (allow_released_opsets_only is - // passed explicitly to Load below, since the byte/proto Load overloads don't read it - // from the config options.) + // loading the serialized model applies the same options. ASSERT_STATUS_OK(session_options.config_options.AddConfigEntry( kOrtSessionOptionsConfigStrictShapeTypeInference, model_options.strict_shape_type_inference ? "1" : "0")); @@ -187,8 +185,7 @@ void TransformerTester(const std::function& buil ASSERT_STATUS_OK(session.RegisterExecutionProvider(ep_shared)); } - // Load via an istream so the under-development opset flag can be honored: the istream Load - // overload takes allow_released_opsets_only explicitly (the byte/proto overloads hardcode it). + // Load via an istream so the under-development opset flag can be passed explicitly. // This lets models built at an unreleased ONNX opset (e.g. opset 27 in ONNX 1.22) load on // strict (ALLOW_RELEASED_ONNX_OPSET_ONLY default) CI legs. std::istringstream model_istream(model_data); diff --git a/onnxruntime/test/unittest_util/op_tester.cc b/onnxruntime/test/unittest_util/op_tester.cc index b633989866277..2b1f4d87df2e1 100644 --- a/onnxruntime/test/unittest_util/op_tester.cc +++ b/onnxruntime/test/unittest_util/op_tester.cc @@ -40,6 +40,12 @@ onnxruntime::Model& OpTester::BuildModel(const std::unordered_map domain_to_version(extra_domain_to_version); const auto& domain = Domain(); + if (domain != kOnnxDomain && domain_to_version.count(kOnnxDomain) == 0) { + const auto& last_released_versions = + ONNX_NAMESPACE::OpSchemaRegistry::DomainToVersionRange::Instance().LastReleaseVersionMap(); + domain_to_version.insert({kOnnxDomain, last_released_versions.at(kOnnxDomain)}); + } + if (domain_to_version.count(domain) == 0) { domain_to_version.insert({domain, Opset()}); } else { diff --git a/requirements-materialize-onnx-node-tests.txt b/requirements-materialize-onnx-node-tests.txt index ee46a7f7ea7ed..adc62024df3a9 100644 --- a/requirements-materialize-onnx-node-tests.txt +++ b/requirements-materialize-onnx-node-tests.txt @@ -10,6 +10,6 @@ # PIN-LOCK: these values MUST stay in lockstep with the CMake gate (cmake/onnxruntime_unittests.cmake, # which HARD-FATALs when the installed versions differ from its expected onnx/numpy pins) and with # the Python MIN pins. Bump here and in cmake/deps.txt + the numpy MIN pins together. -onnx==1.22.0 +onnx==1.23.1 numpy==2.2.6; python_version < "3.11" numpy==2.4.2; python_version >= "3.11" diff --git a/tools/ci_build/build.py b/tools/ci_build/build.py index 25d08695fcfde..32711a96b3334 100644 --- a/tools/ci_build/build.py +++ b/tools/ci_build/build.py @@ -149,6 +149,13 @@ def run_subprocess( return run(*args, cwd=cwd, capture_stdout=capture_stdout, shell=shell, env=my_env) +def get_onnx_backend_test_environment(_use_cuda): + return { + "ALLOW_RELEASED_ONNX_OPSET_ONLY": "0", + "ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS": "1", + } + + def update_submodules(source_dir): run_subprocess(["git", "submodule", "sync", "--recursive"], cwd=source_dir) run_subprocess(["git", "submodule", "update", "--init", "--recursive"], cwd=source_dir) @@ -2069,7 +2076,13 @@ def run_onnxruntime_tests(args, source_dir, ctest_path, build_dir, configs): if not args.skip_onnx_tests: run_subprocess([os.path.join(cwd, "onnx_test_runner"), "test_models"], cwd=cwd) if config != "Debug": - run_subprocess([sys.executable, "onnx_backend_test_series.py"], cwd=cwd, dll_path=dll_path) + # Set the opset policy explicitly so the child process does not inherit the CI default. + run_subprocess( + [sys.executable, "onnx_backend_test_series.py"], + cwd=cwd, + dll_path=dll_path, + env=get_onnx_backend_test_environment(args.use_cuda), + ) if not args.skip_keras_test: try: diff --git a/tools/ci_build/github/linux/build_linux_python_package.sh b/tools/ci_build/github/linux/build_linux_python_package.sh index 6dfbc25ca26fd..d8264e1449739 100755 --- a/tools/ci_build/github/linux/build_linux_python_package.sh +++ b/tools/ci_build/github/linux/build_linux_python_package.sh @@ -142,7 +142,7 @@ do ls "$python3_dir" PIP_REQUIREMENTS=(-r /onnxruntime_src/tools/ci_build/github/linux/python/requirements.txt) if [[ "${PYTHON_EXE}" == */cp3??-cp3??t/* ]]; then - # ONNX does not publish free-threaded wheels, and ONNX-dependent tooling would resolve its unsupported source build. + # ONNX-dependent tooling lacks wheels for some free-threaded versions; avoid unsupported source builds. PIP_REQUIREMENTS=(-r "${FREE_THREADED_REQUIREMENTS}") fi if [[ "${PYTHON_EXE}" == */cp313-cp313t/* ]]; then diff --git a/tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt b/tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt index e5b4b9b2dbd94..fa93a018f0481 100644 --- a/tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt +++ b/tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt @@ -7,4 +7,4 @@ wheel protobuf==6.33.6 sympy==1.14 flatbuffers -onnx==1.22.0 +onnx==1.23.1 diff --git a/tools/ci_build/github/linux/docker/scripts/lort/requirements.txt b/tools/ci_build/github/linux/docker/scripts/lort/requirements.txt index 1f99fc4c9a25b..c9c9038d2a3a4 100644 --- a/tools/ci_build/github/linux/docker/scripts/lort/requirements.txt +++ b/tools/ci_build/github/linux/docker/scripts/lort/requirements.txt @@ -3,7 +3,7 @@ beartype==0.15.0 flatbuffers cerberus h5py -onnx==1.22.0 +onnx==1.23.1 # Python dependencies required for pytorch development astunparse expecttest!=0.2.0 diff --git a/tools/ci_build/github/linux/docker/scripts/manylinux/requirements.txt b/tools/ci_build/github/linux/docker/scripts/manylinux/requirements.txt index 5b0a682d9d419..818763d672e58 100644 --- a/tools/ci_build/github/linux/docker/scripts/manylinux/requirements.txt +++ b/tools/ci_build/github/linux/docker/scripts/manylinux/requirements.txt @@ -9,4 +9,4 @@ sympy==1.14 flatbuffers neural-compressor>=2.2.1 triton==3.5.0 -onnx==1.22.0 +onnx==1.23.1 diff --git a/tools/ci_build/github/linux/docker/scripts/requirements.txt b/tools/ci_build/github/linux/docker/scripts/requirements.txt index 462de90e67c3f..ba56c68dd5ea5 100644 --- a/tools/ci_build/github/linux/docker/scripts/requirements.txt +++ b/tools/ci_build/github/linux/docker/scripts/requirements.txt @@ -12,4 +12,4 @@ protobuf==6.33.6 packaging onnxscript==0.6.2 onnx-ir==0.1.16 -onnx==1.22.0 +onnx==1.23.1 diff --git a/tools/ci_build/github/linux/python/requirements.txt b/tools/ci_build/github/linux/python/requirements.txt index 4730f767c9b50..ada77b03d35a1 100644 --- a/tools/ci_build/github/linux/python/requirements.txt +++ b/tools/ci_build/github/linux/python/requirements.txt @@ -12,4 +12,4 @@ onnxscript==0.6.2 onnx-ir==0.1.16 jinja2 markupsafe -onnx==1.22.0 +onnx==1.23.1 diff --git a/tools/ci_build/github/windows/python/requirements.txt b/tools/ci_build/github/windows/python/requirements.txt index ef80b8af11b05..c0f1de5dfbb7d 100644 --- a/tools/ci_build/github/windows/python/requirements.txt +++ b/tools/ci_build/github/windows/python/requirements.txt @@ -14,4 +14,4 @@ jinja2 markupsafe semver packaging -onnx==1.22.0 +onnx==1.23.1 diff --git a/tools/ci_build/test_build_args.py b/tools/ci_build/test_build_args.py index a98eaaaed1974..b2c93ed54248f 100644 --- a/tools/ci_build/test_build_args.py +++ b/tools/ci_build/test_build_args.py @@ -2,6 +2,7 @@ # Copyright (c) Microsoft Corporation. All rights reserved. # Licensed under the MIT License. +import importlib import sys import unittest from pathlib import Path @@ -10,7 +11,13 @@ sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "python")) sys.path.insert(0, str(Path(__file__).resolve().parent)) -import build_args +build = importlib.import_module("build") +build_args = importlib.import_module("build_args") + +_UNRELEASED_OPSET_ENVIRONMENT = { + "ALLOW_RELEASED_ONNX_OPSET_ONLY": "0", + "ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS": "1", +} class BuildArgsTest(unittest.TestCase): @@ -65,5 +72,25 @@ def test_acl_deprecation_warning_not_emitted_without_use_acl(self): warn.assert_not_called() +class OnnxBackendTestEnvironmentTest(unittest.TestCase): + def test_cpu_and_cuda_enable_unreleased_opsets_by_default(self): + with mock.patch.dict(build.os.environ, {}, clear=True): + for use_cuda in (False, True): + with self.subTest(use_cuda=use_cuda): + self.assertEqual( + build.get_onnx_backend_test_environment(use_cuda), + _UNRELEASED_OPSET_ENVIRONMENT, + ) + + def test_cpu_and_cuda_override_explicit_parent_strict_opset_mode(self): + with mock.patch.dict(build.os.environ, {"ALLOW_RELEASED_ONNX_OPSET_ONLY": "1"}, clear=True): + for use_cuda in (False, True): + with self.subTest(use_cuda=use_cuda): + self.assertEqual( + build.get_onnx_backend_test_environment(use_cuda), + _UNRELEASED_OPSET_ENVIRONMENT, + ) + + if __name__ == "__main__": unittest.main() diff --git a/tools/python/materialize_onnx_node_tests.py b/tools/python/materialize_onnx_node_tests.py index 5071dfca75a70..0e4b6de46c924 100644 --- a/tools/python/materialize_onnx_node_tests.py +++ b/tools/python/materialize_onnx_node_tests.py @@ -28,8 +28,7 @@ * ``onnx.__version__ == --expected-onnx-version`` -- HARD FAIL. The model opset/IR is stamped from the *compiled* schema registry (``get_schema().since_version``), so a mismatched wheel silently bakes the - wrong opset into ``model.onnx``. Compared on the release base so a pre-release - (``rcN`` / ``.dev``) wheel of the pinned tag does not false-FATAL. + wrong opset into ``model.onnx``. The source and wheel pins must match exactly. * ``numpy.__version__`` vs ``--expected-numpy-version`` -- SOFT WARNING only, never a hard gate. ONNX's own corpus-gen numpy is uncontrollable/unknowable, so a fixed pin cannot guarantee byte-identity and a FATAL would false-red on @@ -52,7 +51,6 @@ import argparse import os -import re import shutil import sys import warnings @@ -93,19 +91,6 @@ class MaterializeError(RuntimeError): # --------------------------------------------------------------------------- # # Version / count guards # --------------------------------------------------------------------------- # -def _release_base(version: str) -> str: - """Return the release base (major.minor.micro) of a possibly-prerelease version. - - The ONNX opset-bump workflow ships release-candidate / dev wheels like - ``1.23.0rc1`` or ``1.23.0.dev20240101`` whose COMPILED opset registry already - matches the formal ``1.23.0`` tag. Comparing on the release base lets an RC - wheel of the pinned tag pass, while a genuine major/minor/micro mismatch - (e.g. ``1.22.0`` vs ``1.23.0``) still fails loud. - """ - m = re.match(r"^\s*(\d+)\.(\d+)\.(\d+)", version) - return ".".join(m.groups()) if m else version.strip() - - def _check_versions(expected_onnx: str, expected_numpy: str | None) -> None: """Assert wheel parity BEFORE importing/collecting cases. @@ -124,17 +109,14 @@ def _check_versions(expected_onnx: str, expected_numpy: str | None) -> None: import numpy # noqa: PLC0415 import onnx # noqa: PLC0415 - # onnx: HARD FAIL, RC/pre-release aware (compare release base, not the raw - # string) so an rcN/dev wheel of the pinned tag isn't a false FATAL. - if expected_onnx and _release_base(onnx.__version__) != _release_base(expected_onnx): + # onnx: HARD FAIL on any source/wheel version mismatch. + if expected_onnx and onnx.__version__ != expected_onnx: raise MaterializeError( f"onnx version mismatch: installed onnx=={onnx.__version__} but " - f"--expected-onnx-version=={expected_onnx} (release base " - f"{_release_base(onnx.__version__)} != {_release_base(expected_onnx)}). " + f"--expected-onnx-version=={expected_onnx}. " f"The node-test corpus opset/IR is baked from the COMPILED onnx schema " f"registry, so a mismatched wheel produces a silently-drifted corpus. " - f"Install onnx=={expected_onnx} (the cmake/deps.txt pin) before materializing. " - f"A pre-release (rcN/.dev) of the SAME release base is accepted." + f"Install onnx=={expected_onnx} (the cmake/deps.txt pin) before materializing." ) # numpy: SOFT WARNING, never a hard gate. ONNX's OWN corpus-gen numpy is # uncontrollable and unknowable, so a fixed pin can neither guarantee @@ -197,11 +179,9 @@ def _write_case(case, output_root: str) -> None: n_graph_in = len(case.model.graph.input) n_graph_out = len(case.model.graph.output) for i, (inputs, outputs) in enumerate(case.data_sets): - # Fail loud BEFORE writing any .pb for this dataset: the writer indexes - # graph.input[j]/graph.output[j] per collected value, so more collected - # values than graph arity would raise a bare IndexError mid-write (not - # caught by main's MaterializeError handler) and leave a partial tree. - if len(inputs) > n_graph_in or len(outputs) > n_graph_out: + # Fail loud before writing a partial dataset if positional values do not + # exactly match their graph-value declarations. + if len(inputs) != n_graph_in or len(outputs) != n_graph_out: raise MaterializeError( f"node case {case.name!r} data_set {i}: arity mismatch -- collected " f"{len(inputs)} input(s)/{len(outputs)} output(s) but the model graph " From f5277349b262914008e6d6f678dc2d63320cd6c4 Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Tue, 6 Oct 2026 10:10:45 -0700 Subject: [PATCH 02/10] Restore existing Cast types at opset 28 across execution providers Split prior Cast registrations and preserve existing conversion type sets so CPU fallback casts and current-opset provider tests bind valid kernels. Keep FP6 conversion support deferred, add version-boundary regressions, remove resolved Cast backend exclusions, and update operator documentation. Isolate strict Python backend policy tests from CI's global permissive opset environment. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- docs/OperatorKernels.md | 9 +- .../providers/cpu/cpu_execution_provider.cc | 6 +- .../core/providers/cpu/tensor/cast_op.cc | 13 ++- .../providers/cuda/cuda_execution_provider.cc | 107 ++++++++++++------ .../core/providers/cuda/tensor/cast_op.cc | 36 ++++-- .../src/Operators/OperatorRegistration.cpp | 1 + .../providers/js/js_execution_provider.cc | 6 +- .../core/providers/js/operators/cast.cc | 11 +- .../webgpu/webgpu_execution_provider.cc | 3 +- .../test/providers/cpu/tensor/cast_op_test.cc | 25 +++- .../python/onnxruntime_test_python_backend.py | 3 + .../onnx_backend_test_series_filters.jsonc | 3 - 12 files changed, 169 insertions(+), 54 deletions(-) diff --git a/docs/OperatorKernels.md b/docs/OperatorKernels.md index 2351e770abee4..6d8cfdd113d77 100644 --- a/docs/OperatorKernels.md +++ b/docs/OperatorKernels.md @@ -70,7 +70,8 @@ The **OpSet Version** column uses the following notation: |BitwiseOr|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseXor|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BlackmanWindow|*in* size:**T1**
*out* output:**T2**|17+|**T1** = tensor(int32), tensor(int64)
**T2** = tensor(double), tensor(float), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| -|Cast|*in* input:**T1**
*out* output:**T2**|25+|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(float8e8m0), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(float8e8m0), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)| +|Cast|*in* input:**T1**
*out* output:**T2**|28+|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(float8e8m0), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(float8e8m0), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)| +|||[25, 27]|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(float8e8m0), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(float8e8m0), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)| |||24|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(float8e8m0), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(float8e8m0), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)| |||23|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)| |||[21, 22]|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e4m3fnuz), tensor(float8e5m2), tensor(float8e5m2fnuz), tensor(int16), tensor(int2), tensor(int32), tensor(int4), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint2), tensor(uint32), tensor(uint4), tensor(uint64), tensor(uint8)| @@ -694,7 +695,8 @@ The **OpSet Version** column uses the following notation: |||14|**T** = tensor(double), tensor(float), tensor(float16)
**U** = tensor(double), tensor(float), tensor(float16)| |||[9, 13]|**T** = tensor(double), tensor(float), tensor(float16)| |||[7, 8]|**T** = tensor(double), tensor(float), tensor(float16)| -|Cast|*in* input:**T1**
*out* output:**T2**|25+|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|Cast|*in* input:**T1**
*out* output:**T2**|28+|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|||[25, 27]|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||[23, 24]|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||[21, 22]|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||[19, 20]|**T1** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(float4e2m1), tensor(float8e4m3fn), tensor(float8e5m2), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| @@ -1247,7 +1249,8 @@ The **OpSet Version** column uses the following notation: |BitwiseNot|*in* X:**T**
*out* Y:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseOr|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |BitwiseXor|*in* A:**T**
*in* B:**T**
*out* C:**T**|18+|**T** = tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| -|Cast|*in* input:**T1**
*out* output:**T2**|21+|**T1** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|Cast|*in* input:**T1**
*out* output:**T2**|28+|**T1** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|||21+|**T1** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||19+|**T1** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||13+|**T1** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||9+|**T1** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)
**T2** = tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| diff --git a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc index add722a2e8082..c05ba89a11496 100644 --- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc +++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc @@ -1500,7 +1500,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 24, uint8_t, TopK); // Opset 25 -class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, Cast); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, 27, Cast); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, ConstantOfShape); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, int32_t, DequantizeLinear); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, uint8_t, DequantizeLinear); @@ -1557,6 +1557,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 27, Range); // Opset 28 +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, Cast); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int8_t, BitShift); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int16_t, BitShift); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int32_t, BitShift); @@ -3692,7 +3693,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, // opset 25 - BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -3774,6 +3775,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, // opset 28 + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo()) + .TypeConstraint("T2", BuildKernelDefConstraintsFromTypeList()) + .MayInplace(0, 0), + Cast); + +// FP6 conversions are deferred; preserve the existing types for CPU fallback casts. +ONNX_CPU_OPERATOR_KERNEL( + Cast, + 28, KernelDefBuilder() .TypeConstraint("T1", BuildKernelDefConstraintsFromTypeList()) .TypeConstraint("T2", BuildKernelDefConstraintsFromTypeList()) diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 80dbd322fa534..b1098d97cdad0 100755 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -1843,25 +1843,25 @@ class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 24, T class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 24, 24, Unsqueeze); // Opset 25. -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, float, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, double, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, MLFloat16, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, BFloat16, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, int8_t, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, int16_t, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, int32_t, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, int64_t, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint8_t, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint16_t, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint32_t, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint64_t, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, bool, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, float, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, double, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, MLFloat16, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, BFloat16, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, int8_t, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, int16_t, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, int32_t, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, int64_t, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, uint8_t, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, uint16_t, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, uint32_t, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, uint64_t, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, bool, Cast); #if !defined(DISABLE_FLOAT8_TYPES) -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, Float8E4M3FN, Cast); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, Float8E5M2, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, Float8E4M3FN, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, Float8E5M2, Cast); #endif #if !defined(DISABLE_FLOAT4_TYPES) -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, Float4E2M1x2, Cast); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, Float4E2M1x2, Cast); #endif class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, ConstantOfShape); class ONNX_OPERATOR_TWO_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint8_t, float, DequantizeLinear); @@ -1910,6 +1910,28 @@ class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, U // Opset 27. class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 27, Range); + +// Opset 28. +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, float, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, double, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, MLFloat16, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, BFloat16, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int8_t, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int16_t, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int32_t, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int64_t, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint8_t, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint16_t, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint32_t, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint64_t, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, bool, Cast); +#if !defined(DISABLE_FLOAT8_TYPES) +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Float8E4M3FN, Cast); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Float8E5M2, Cast); +#endif +#if !defined(DISABLE_FLOAT4_TYPES) +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Float4E2M1x2, Cast); +#endif #endif static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { @@ -3138,25 +3160,25 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, // Opset 25 - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, #if !defined(DISABLE_FLOAT8_TYPES) - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, #endif #if !defined(DISABLE_FLOAT4_TYPES) - BuildKernelCreateInfo, + BuildKernelCreateInfo, #endif BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -3204,6 +3226,27 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, // Opset 27 BuildKernelCreateInfo, + // Opset 28 + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, +#if !defined(DISABLE_FLOAT8_TYPES) + BuildKernelCreateInfo, + BuildKernelCreateInfo, +#endif +#if !defined(DISABLE_FLOAT4_TYPES) + BuildKernelCreateInfo, +#endif #endif }; diff --git a/onnxruntime/core/providers/cuda/tensor/cast_op.cc b/onnxruntime/core/providers/cuda/tensor/cast_op.cc index 2ed08e25d02d2..6146b2ca934d1 100644 --- a/onnxruntime/core/providers/cuda/tensor/cast_op.cc +++ b/onnxruntime/core/providers/cuda/tensor/cast_op.cc @@ -100,10 +100,20 @@ const std::vector& CastOpTypeConstraints() { .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ .TypeConstraint("T2", CastOpTypeConstraints()), \ Cast); \ + ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \ + Cast, \ + kOnnxDomain, \ + 25, 27, \ + T, \ + kCudaExecutionProvider, \ + (*KernelDefBuilder::Create()) \ + .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ + .TypeConstraint("T2", CastOpTypeConstraints()), \ + Cast); \ ONNX_OPERATOR_TYPED_KERNEL_EX( \ Cast, \ kOnnxDomain, \ - 25, \ + 28, \ T, \ kCudaExecutionProvider, \ (*KernelDefBuilder::Create()) \ @@ -411,11 +421,21 @@ SPECIALIZE_IMPL(BFloat16) .TypeConstraint("T2", OutputTypeConstraints), \ Cast); -#define REGISTER_KERNEL_TYPED_25(T, OutputTypeConstraints) \ +#define REGISTER_KERNEL_TYPED_25_TO_28(T, OutputTypeConstraints) \ + ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \ + Cast, \ + kOnnxDomain, \ + 25, 27, \ + T, \ + kCudaExecutionProvider, \ + (*KernelDefBuilder::Create()) \ + .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ + .TypeConstraint("T2", OutputTypeConstraints), \ + Cast); \ ONNX_OPERATOR_TYPED_KERNEL_EX( \ Cast, \ kOnnxDomain, \ - 25, \ + 28, \ T, \ kCudaExecutionProvider, \ (*KernelDefBuilder::Create()) \ @@ -425,20 +445,20 @@ SPECIALIZE_IMPL(BFloat16) #if !defined(DISABLE_FLOAT8_TYPES) -#define SPECIALIZE_IMPL_19_TO_25(T) \ +#define SPECIALIZE_IMPL_19_TO_28(T) \ REGISTER_KERNEL_TYPED_19_TO_22(T) \ REGISTER_KERNEL_TYPED_23_TO_24(T, CastOpTypeConstraints()) \ - REGISTER_KERNEL_TYPED_25(T, CastOpTypeConstraints()) \ + REGISTER_KERNEL_TYPED_25_TO_28(T, CastOpTypeConstraints()) \ template Status Cast::ComputeInternal(OpKernelContext* context) const; -SPECIALIZE_IMPL_19_TO_25(Float8E4M3FN) -SPECIALIZE_IMPL_19_TO_25(Float8E5M2) +SPECIALIZE_IMPL_19_TO_28(Float8E4M3FN) +SPECIALIZE_IMPL_19_TO_28(Float8E5M2) #endif #if !defined(DISABLE_FLOAT4_TYPES) REGISTER_KERNEL_TYPED_23_TO_24(Float4E2M1x2, {DataTypeImpl::GetTensorType()}) -REGISTER_KERNEL_TYPED_25(Float4E2M1x2, {DataTypeImpl::GetTensorType()}) +REGISTER_KERNEL_TYPED_25_TO_28(Float4E2M1x2, {DataTypeImpl::GetTensorType()}) template Status Cast::ComputeInternal(OpKernelContext* context) const; #endif diff --git a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp index 26f998c7521a2..b4c05dcfaa51a 100644 --- a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp +++ b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp @@ -1092,6 +1092,7 @@ constexpr static OperatorRegistrationInformation operatorRegistrationInformation {REG_INFO( 13, Cast, typeNameListTwo, supportedTypeListCast, DmlGraphSupport::Supported)}, {REG_INFO( 19, Cast, typeNameListTwo, supportedTypeListCast, DmlGraphSupport::Supported)}, {REG_INFO( 21, Cast, typeNameListTwo, supportedTypeListCast, DmlGraphSupport::Supported)}, + {REG_INFO( 28, Cast, typeNameListTwo, supportedTypeListCast, DmlGraphSupport::Supported)}, {REG_INFO_VER( 15, CastLike, typeNameListTwo, supportedTypeListCast, DmlGraphSupport::Supported)}, {REG_INFO_VER( 19, CastLike, typeNameListTwo, supportedTypeListCast, DmlGraphSupport::Supported)}, {REG_INFO_VER( 21, CastLike, typeNameListTwo, supportedTypeListCast, DmlGraphSupport::Supported)}, diff --git a/onnxruntime/core/providers/js/js_execution_provider.cc b/onnxruntime/core/providers/js/js_execution_provider.cc index 135aba2e32b59..c3e0072fd0dab 100644 --- a/onnxruntime/core/providers/js/js_execution_provider.cc +++ b/onnxruntime/core/providers/js/js_execution_provider.cc @@ -123,7 +123,8 @@ class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomai class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomain, 9, 12, Cast); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomain, 13, 18, Cast); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomain, 19, 20, Cast); -class ONNX_OPERATOR_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomain, 21, Cast); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomain, 21, 27, Cast); +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomain, 28, Cast); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomain, 6, 10, Clip); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kJsExecutionProvider, kOnnxDomain, 11, 11, Clip); @@ -466,7 +467,8 @@ std::unique_ptr RegisterKernels() { KERNEL_CREATE_INFO_VERSIONED(9, 12, Cast), KERNEL_CREATE_INFO_VERSIONED(13, 18, Cast), KERNEL_CREATE_INFO_VERSIONED(19, 20, Cast), - KERNEL_CREATE_INFO(21, Cast), + KERNEL_CREATE_INFO_VERSIONED(21, 27, Cast), + KERNEL_CREATE_INFO(28, Cast), // activations KERNEL_CREATE_INFO_VERSIONED(6, 10, Clip), diff --git a/onnxruntime/core/providers/js/operators/cast.cc b/onnxruntime/core/providers/js/operators/cast.cc index f499d0627e032..f03a59f0bd922 100644 --- a/onnxruntime/core/providers/js/operators/cast.cc +++ b/onnxruntime/core/providers/js/operators/cast.cc @@ -58,10 +58,19 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX( .TypeConstraint("T1", CastOpTypeConstraints()) .TypeConstraint("T2", CastOpTypeConstraints()), Cast); +ONNX_OPERATOR_VERSIONED_KERNEL_EX( + Cast, + kOnnxDomain, + 21, 27, + kJsExecutionProvider, + (*KernelDefBuilder::Create()) + .TypeConstraint("T1", CastOpTypeConstraints()) + .TypeConstraint("T2", CastOpTypeConstraints()), + Cast); ONNX_OPERATOR_KERNEL_EX( Cast, kOnnxDomain, - 21, + 28, kJsExecutionProvider, (*KernelDefBuilder::Create()) .TypeConstraint("T1", CastOpTypeConstraints()) diff --git a/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc b/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc index ea7cf7a71d9fd..f6dddba3a0686 100644 --- a/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc +++ b/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc @@ -465,7 +465,8 @@ std::unique_ptr RegisterKernels(bool enable_graph_capture, bool ORT_THROW_IF_ERROR(kernel_registry->Register(CreateCastVersionedKernelInfo(21, 22, enable_int64))); ORT_THROW_IF_ERROR(kernel_registry->Register(CreateCastVersionedKernelInfo(23, 23, enable_int64))); ORT_THROW_IF_ERROR(kernel_registry->Register(CreateCastVersionedKernelInfo(24, 24, enable_int64))); - ORT_THROW_IF_ERROR(kernel_registry->Register(CreateCastKernelInfo(25, enable_int64))); + ORT_THROW_IF_ERROR(kernel_registry->Register(CreateCastVersionedKernelInfo(25, 27, enable_int64))); + ORT_THROW_IF_ERROR(kernel_registry->Register(CreateCastKernelInfo(28, enable_int64))); // Register int64 Clip kernels with conditional int64 support RegisterClipInt64Kernels(*kernel_registry, enable_int64); diff --git a/onnxruntime/test/providers/cpu/tensor/cast_op_test.cc b/onnxruntime/test/providers/cpu/tensor/cast_op_test.cc index 05bd7aef1e440..07eca67baed4a 100644 --- a/onnxruntime/test/providers/cpu/tensor/cast_op_test.cc +++ b/onnxruntime/test/providers/cpu/tensor/cast_op_test.cc @@ -155,6 +155,9 @@ std::vector CastedValues(gsl::span src) { } struct CastNonStringTester { + int opset{21}; + bool cuda_only{false}; + template void operator()(const std::pair&) { SCOPED_TRACE( @@ -177,7 +180,8 @@ struct CastNonStringTester { auto output_span = gsl::make_span(output_buffer.get(), size); CastSpan(input_span, output_span); - TestCastOp(input_span, output_span, shape.AsShapeVector()); + TestCastOp(input_span, output_span, shape.AsShapeVector(), + OpTester::ExpectResult::kExpectSuccess, "", opset, Saturate::None, cuda_only); } }; @@ -194,6 +198,25 @@ TEST(CastOpTest, NonStringTypes) { CastNonStringTester{}); } +TEST(CastOpTest, NonStringTypes_Opset25To28) { + for (int opset : {25, 27, 28}) { + SCOPED_TRACE(opset); + boost::mp11::mp_for_each>( + CastNonStringTester{opset}); + } +} + +#if defined(USE_CUDA) +TEST(CastOpTest, NonStringTypes_Opset28_Cuda) { + if (!HasCudaEnvironment(530)) { + GTEST_SKIP() << "CUDA Cast tests require a CUDA device with compute capability >= 5.3."; + } + + boost::mp11::mp_for_each>( + CastNonStringTester{28, true}); +} +#endif + TEST(CastOpTest, FromString) { const std::vector shape{2, 2, 2}; const std::vector string_data = {"-inf", "+INF", "0.9767611", "0.28280696", diff --git a/onnxruntime/test/python/onnxruntime_test_python_backend.py b/onnxruntime/test/python/onnxruntime_test_python_backend.py index e06d4f997514c..946d164f8f7e4 100644 --- a/onnxruntime/test/python/onnxruntime_test_python_backend.py +++ b/onnxruntime/test/python/onnxruntime_test_python_backend.py @@ -6,6 +6,7 @@ import tempfile import unittest from types import SimpleNamespace +from unittest.mock import patch import numpy as np from helper import get_name @@ -27,6 +28,7 @@ def test_run_model(self): output_expected = np.array([[1.0, 4.0], [9.0, 16.0], [25.0, 36.0]], dtype=np.float32) np.testing.assert_allclose(res[0], output_expected, rtol=1e-05, atol=1e-08) + @patch.dict(os.environ, {"ALLOW_RELEASED_ONNX_OPSET_ONLY": "1"}) def test_prepare_released_opset28_model_when_policy_enabled(self): model = self._make_opset28_model() original_policy = OnnxRuntimeBackend.allowReleasedOpsetsOnly @@ -39,6 +41,7 @@ def test_prepare_released_opset28_model_when_policy_enabled(self): result = rep.run(np.array([1.0], dtype=np.float32)) np.testing.assert_array_equal(result[0], np.array([1.0], dtype=np.float32)) + @patch.dict(os.environ, {"ALLOW_RELEASED_ONNX_OPSET_ONLY": "1"}) def test_prepare_model_bytes_rejects_unreleased_opset29_when_policy_enabled(self): model = self._make_opset28_model() model.opset_import[0].version = 29 diff --git a/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc b/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc index cf40b161dd008..d5d3feb855c8d 100644 --- a/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc +++ b/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc @@ -118,8 +118,6 @@ "^test_attention_4d_padded_kv_bf16_cuda$", "^test_attention_4d_padded_kv_bf16_expanded_cuda$", "^test_attention_local_window_ext_cache_float16_mask_cpu$", - // TODO: Remove the exact Cast(28) failures when CPU Cast(28) is available in all builds. - "^test_cast_DOUBLE_to_FLOAT_cpu$", "^test_celu_bfloat16_cpu$", "^test_celu_bfloat16_expanded_cpu$", "^test_celu_cpu$", @@ -187,7 +185,6 @@ "^test_unique_sorted_with_negative_axis_cpu$", "^test_unique_sorted_without_axis_cpu$", // TODO: Remove these exact CUDA skips when fallback nodes have providers and Swish(24) has a CPU kernel. - "^test_cast_DOUBLE_to_FLOAT_cuda$", "^test_celu_bfloat16_cuda$", "^test_celu_bfloat16_expanded_cuda$", "^test_celu_cuda$", From 35aacac653343d2473b7aa63e49a7cbdf6a75098 Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Tue, 6 Oct 2026 13:12:26 -0700 Subject: [PATCH 03/10] Fix remaining ONNX reintegration CI regressions Define the DirectML Cast-28 version constant and regenerate the web operator table. Reject unsupported floating floor-Mod and CRD SpaceToDepth in QNN capability checks, with focused regression coverage preserving integer Mod and DCR support. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- js/web/docs/webgpu-operators.md | 2 +- .../OperatorAuthorHelper/OperatorVersions.h | 5 +++ .../qnn/builder/opbuilder/mod_op_builder.cc | 11 ++++- .../builder/opbuilder/simple_op_builder.cc | 6 +++ onnxruntime/test/providers/qnn/mod_op_test.cc | 20 +++++++++ .../providers/qnn/qnn_model_wrapper_test.cc | 42 +++++++++++++++++++ 6 files changed, 84 insertions(+), 2 deletions(-) diff --git a/js/web/docs/webgpu-operators.md b/js/web/docs/webgpu-operators.md index 05e930c6a658c..ebc69778302e7 100644 --- a/js/web/docs/webgpu-operators.md +++ b/js/web/docs/webgpu-operators.md @@ -25,7 +25,7 @@ Do not modify directly.* | BatchNormalization | ai.onnx(7-8,9-13,14,15+); com.ms.internal.nhwc(7-8,9-13,14,15+) | | | BiasAdd | com.microsoft(1+) | | | BiasSplitGelu | com.microsoft(1+) | | -| Cast | ai.onnx(6-8,9-12,13-18,19-20,21+) | | +| Cast | ai.onnx(6-8,9-12,13-18,19-20,21-27,28+) | | | Ceil | ai.onnx(6-12,13+) | | | Clip | ai.onnx(6-10,11,12,13+) | | | Concat | ai.onnx(1-3,4-10,11-12,13+) | | diff --git a/onnxruntime/core/providers/dml/OperatorAuthorHelper/OperatorVersions.h b/onnxruntime/core/providers/dml/OperatorAuthorHelper/OperatorVersions.h index b4d402a1d9e77..3dc41cbd01b12 100644 --- a/onnxruntime/core/providers/dml/OperatorAuthorHelper/OperatorVersions.h +++ b/onnxruntime/core/providers/dml/OperatorAuthorHelper/OperatorVersions.h @@ -457,6 +457,11 @@ namespace OperatorHelper static const int sc_sinceVer_GroupNorm = 21; } + namespace OnnxOperatorSet28 + { + static const int sc_sinceVer_Cast = 28; + } + namespace MsftOperatorSet1 { static const int sc_sinceVer_DmlFusedConv = 1; diff --git a/onnxruntime/core/providers/qnn/builder/opbuilder/mod_op_builder.cc b/onnxruntime/core/providers/qnn/builder/opbuilder/mod_op_builder.cc index 1d61ac9b9599d..5db8e638e83c1 100644 --- a/onnxruntime/core/providers/qnn/builder/opbuilder/mod_op_builder.cc +++ b/onnxruntime/core/providers/qnn/builder/opbuilder/mod_op_builder.cc @@ -47,12 +47,21 @@ Status ModOpBuilder::ProcessInputs(QnnModelWrapper& qnn_model_wrapper, bool do_op_validation) const { ORT_UNUSED_PARAMETER(do_op_validation); NodeAttrHelper node_helper(node_unit); - int64_t fmod = node_helper.Get("fmod", static_cast(0)); // 0=integer mod. 1=float mod. + int64_t fmod = node_helper.Get("fmod", static_cast(0)); if (1 == fmod) { return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "QNN Mod Op only support fmod == 0 for now."); } const auto& inputs = node_unit.Inputs(); + int32_t input_type = 0; + ORT_RETURN_IF_ERROR(utils::GetOnnxTensorElemDataType(inputs[0].node_arg, input_type)); + // a - b * floor(a / b) does not preserve the ONNX floating-point edge-case semantics. + ORT_RETURN_IF(input_type == ONNX_NAMESPACE::TensorProto_DataType_FLOAT || + input_type == ONNX_NAMESPACE::TensorProto_DataType_FLOAT16 || + input_type == ONNX_NAMESPACE::TensorProto_DataType_DOUBLE || + input_type == ONNX_NAMESPACE::TensorProto_DataType_BFLOAT16, + "QNN EP does not support floating-point Mod with fmod == 0."); + const auto input_count = GetInputCountQnnRequired(node_unit); for (size_t input_i = 0; input_i < input_count; ++input_i) { ORT_RETURN_IF_ERROR(ProcessInput(qnn_model_wrapper, inputs[input_i], logger, input_names)); diff --git a/onnxruntime/core/providers/qnn/builder/opbuilder/simple_op_builder.cc b/onnxruntime/core/providers/qnn/builder/opbuilder/simple_op_builder.cc index 5537a388f598c..9d59ac1ea5417 100644 --- a/onnxruntime/core/providers/qnn/builder/opbuilder/simple_op_builder.cc +++ b/onnxruntime/core/providers/qnn/builder/opbuilder/simple_op_builder.cc @@ -48,6 +48,12 @@ Status SimpleOpBuilder::ExplicitOpCheck(QnnModelWrapper& qnn_model_wrapper, const NodeUnit& node_unit) const { const std::string& op_type = node_unit.OpType(); + if (op_type == "SpaceToDepth") { + NodeAttrHelper node_helper(node_unit); + ORT_RETURN_IF_NOT(node_helper.Get("mode", "DCR") == "DCR", + "QNN EP only supports SpaceToDepth with DCR mode."); + } + if (op_type == "GridSample") { NodeAttrHelper node_helper(node_unit); std::string mode = node_helper.Get("mode", "linear"); diff --git a/onnxruntime/test/providers/qnn/mod_op_test.cc b/onnxruntime/test/providers/qnn/mod_op_test.cc index 271fa76a6d639..06793f72cda51 100644 --- a/onnxruntime/test/providers/qnn/mod_op_test.cc +++ b/onnxruntime/test/providers/qnn/mod_op_test.cc @@ -4,6 +4,7 @@ #if !defined(ORT_MINIMAL_BUILD) #include +#include #include #include "test/providers/qnn/qnn_test_utils.h" @@ -38,6 +39,25 @@ static void RunModTest(const std::vector>& input_defs, // CPU tests: // +TEST_F(QnnCPUBackendTests, Mod_float_Floor_Opset28_Fallback) { + const float inf = std::numeric_limits::infinity(); + RunModTest({TestInputDef({8}, false, {-5.0f, 5.0f, -5.0f, 5.0f, -5.0f, 5.0f, -5.0f, 5.0f}), + TestInputDef({8}, false, {3.0f, 3.0f, -3.0f, -3.0f, inf, inf, -inf, -inf})}, + {utils::MakeAttribute("fmod", static_cast(0))}, + ExpectedEPNodeAssignment::None, + "cpu", + 28); +} + +TEST_F(QnnCPUBackendTests, Mod_integer_Floor_Opset28) { + RunModTest({TestInputDef({4}, false, {-5, 5, -5, 5}), + TestInputDef({4}, false, {3, 3, -3, -3})}, + {utils::MakeAttribute("fmod", static_cast(0))}, + ExpectedEPNodeAssignment::All, + "cpu", + 28); +} + // Test that Mod with dynamic divisor. TEST_F(QnnCPUBackendTests, Mod_dynamic_Divisor) { RandomValueGenerator rand_gen{optional{2345}}; diff --git a/onnxruntime/test/providers/qnn/qnn_model_wrapper_test.cc b/onnxruntime/test/providers/qnn/qnn_model_wrapper_test.cc index 54f966928e577..505d7018cccc6 100644 --- a/onnxruntime/test/providers/qnn/qnn_model_wrapper_test.cc +++ b/onnxruntime/test/providers/qnn/qnn_model_wrapper_test.cc @@ -10,9 +10,11 @@ #if !defined(ORT_MINIMAL_BUILD) && BUILD_QNN_EP_STATIC_LIB #include "core/graph/model.h" +#include "core/providers/qnn/builder/op_builder_factory.h" #include "core/providers/qnn/builder/qnn_model_wrapper.h" #include "core/providers/qnn/builder/qnn_def.h" #include "test/util/include/default_providers.h" +#include "test/util/include/asserts.h" #include "test/util/include/test_environment.h" using namespace onnxruntime; @@ -56,6 +58,46 @@ struct QnnModelWrapperTestContext { } // namespace +TEST(QnnModelWrapperTest, SpaceToDepth_RejectsCRDBeforeLayoutTransformation) { + for (const std::string mode : {"DCR", "CRD"}) { + QnnModelWrapperTestContext ctx; + Graph& graph = ctx.model->MainGraph(); + ONNX_NAMESPACE::TypeProto input_type; + auto* input_tensor = input_type.mutable_tensor_type(); + input_tensor->set_elem_type(ONNX_NAMESPACE::TensorProto::FLOAT); + for (int64_t dim : {1, 2, 4, 4}) { + input_tensor->mutable_shape()->add_dim()->set_dim_value(dim); + } + ONNX_NAMESPACE::TypeProto output_type; + auto* output_tensor = output_type.mutable_tensor_type(); + output_tensor->set_elem_type(ONNX_NAMESPACE::TensorProto::FLOAT); + for (int64_t dim : {1, 8, 2, 2}) { + output_tensor->mutable_shape()->add_dim()->set_dim_value(dim); + } + auto& input = graph.GetOrCreateNodeArg("input", &input_type); + auto& output = graph.GetOrCreateNodeArg("output", &output_type); + Node& node = graph.AddNode("SpaceToDepth", "SpaceToDepth", "", {&input}, {&output}); + node.AddAttribute("blocksize", static_cast(2)); + node.AddAttribute("mode", mode); + ASSERT_STATUS_OK(graph.Resolve()); + ctx.graph_viewer = std::make_unique(graph); + ctx.input_index_map = {{"input", 0}}; + ctx.output_index_map = {{"output", 0}}; + + auto wrapper = ctx.CreateWrapper(ModelSettings{}); + const auto* builder = GetOpBuilder("SpaceToDepth"); + ASSERT_NE(builder, nullptr); + const NodeUnit node_unit(node); + auto status = builder->IsOpSupported(*wrapper, node_unit, DefaultLoggingManager().DefaultLogger()); + if (mode == "DCR") { + ASSERT_STATUS_OK(status); + } else { + EXPECT_FALSE(status.IsOK()); + EXPECT_NE(status.ErrorMessage().find("only supports SpaceToDepth with DCR mode"), std::string::npos); + } + } +} + // Verifies that when htp_shared_memory is disabled (default), the mem type of a // graph input tensor remains QNN_TENSORMEMTYPE_RAW. TEST(QnnModelWrapperTest, AddTensorWrapper_SharedMemoryDisabled_GraphInput_MemTypeIsRaw) { From 2a61404f59b50030cb3f6f0efe57794a9a0853cb Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Tue, 6 Oct 2026 16:31:32 -0700 Subject: [PATCH 04/10] Address ONNX reintegration review feedback Cap unverified Mod-28 graph-provider support, add an assignment-aware fallback regression, restore strict backend-corpus policy, exercise the per-session opt-out with unreleased opset29 passthrough graphs, and enable the implemented BitShift corpus cases. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- onnxruntime/core/providers/cann/cann_graph.cc | 3 +- .../migraphx/migraphx_execution_provider.cc | 5 +++ .../openvino/ov_versions/data_ops.cc | 4 ++ .../cpu/math/element_wise_ops_test.cc | 41 +++++++++++++++++++ .../python/onnxruntime_test_python_backend.py | 24 +++++++---- .../onnx_backend_test_series_filters.jsonc | 27 ------------ tools/ci_build/build.py | 4 +- tools/ci_build/test_build_args.py | 27 ++++++------ 8 files changed, 85 insertions(+), 50 deletions(-) diff --git a/onnxruntime/core/providers/cann/cann_graph.cc b/onnxruntime/core/providers/cann/cann_graph.cc index fc2c6fcbf361b..064bcfd706336 100644 --- a/onnxruntime/core/providers/cann/cann_graph.cc +++ b/onnxruntime/core/providers/cann/cann_graph.cc @@ -70,7 +70,8 @@ std::vector SupportONNXModel(const GraphViewer& graph_viewer) { const auto& node = graph_viewer.GetNode(index); if (node->Domain() != kOnnxDomain || domain_version < lower_bound || - !cann_supported_ops.count(node->OpType())) { + !cann_supported_ops.count(node->OpType()) || + (node->OpType() == "Mod" && node->SinceVersion() >= 28)) { unsupported_nodes.push_back(index); continue; } diff --git a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc index 8908013a52291..b2e8efa883275 100644 --- a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc +++ b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc @@ -377,6 +377,11 @@ std::vector toVector(const ONNX_NAMESPACE::int64s& nums) { static bool IsUnsupportedOpMode(const onnxruntime::GraphViewer& graph_viewer, const Node* node) { std::vector input_nodes; const auto& optype = node->OpType(); + // Mod-28 floating floor/edge-case semantics require the CPU fallback. + if (optype == "Mod" && node->SinceVersion() >= 28) { + return true; + } + if (optype == "ArgMax" || optype == "ArgMin") { const auto& attributes = node->GetAttributes(); // we do not support select_last_index = 1 for now diff --git a/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc index 7aae81b987c0d..49797978471f8 100644 --- a/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc +++ b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc @@ -727,6 +727,10 @@ bool DataOps::dimension_unsupported(const Node* node) { bool DataOps::node_is_supported(const NodeIndex node_idx, bool& has_external_weights_) { const auto& node = graph_viewer_.GetNode(node_idx); const auto& optype = node->OpType(); + // Mod-28 floating floor/edge-case semantics require the CPU fallback. + if (optype == "Mod" && node->SinceVersion() >= 28) { + return false; + } #ifndef NDEBUG if (openvino_ep::backend_utils::IsDebugEnabled()) { diff --git a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc index 4348942186d17..5048de8c1e47b 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -4396,6 +4396,47 @@ TEST(ModOpTest, FloorMod_float_special_values_opset28) { {-0.0f, 0.0f, 5.0f, -5.0f, nan, nan, nan, nan}); } +TEST(ModOpTest, FloorMod_opset28_provider_fallback) { + InlinedVector> providers; +#ifdef USE_CANN + OrtCANNProviderOptions cann_options{}; + cann_options.enable_cann_graph = 1; + auto cann_factory = CannProviderFactoryCreator::Create(&cann_options); + ASSERT_NE(cann_factory, nullptr); + providers.push_back(cann_factory->CreateProvider()); +#endif + providers.push_back(DefaultMIGraphXExecutionProvider()); + providers.push_back(DefaultOpenVINOExecutionProvider()); + bool tested_provider = false; + for (auto& provider : providers) { + if (!provider) { + continue; + } + tested_provider = true; + SCOPED_TRACE(provider->Type()); + const std::vector shape{4}; + const float infinity = std::numeric_limits::infinity(); + const std::vector x{-5.0f, 5.0f, -5.0f, 5.0f}; + const std::vector y{3.0f, -3.0f, infinity, -infinity}; + OpTester test("Mod", 28); + test.AddAttribute("fmod", 0); + test.AddInput("X", shape, x); + test.AddInput("Y", shape, y); + test.AddOutput("Z", shape, {1.0f, -1.0f, infinity, -infinity}); + std::string model_data; + test.BuildModel().ToProto().SerializeToString(&model_data); + EPVerificationParams params; + params.ep_node_assignment = ExpectedEPNodeAssignment::None; + RunAndVerifyOutputsWithEP( + AsByteSpan(model_data.data(), model_data.size()), CurrentTestName(), std::move(provider), + {{"X", CreateInputOrtValueOnCPU(shape, x)}, {"Y", CreateInputOrtValueOnCPU(shape, y)}}, + params); + } + if (!tested_provider) { + GTEST_SKIP() << "Requires CANN, MIGraphX, or OpenVINO."; + } +} + #if defined(USE_CUDA) TEST(ModOpTest, Fmod_bfloat16_mixed_sign) { OpTester test("Mod", 13); diff --git a/onnxruntime/test/python/onnxruntime_test_python_backend.py b/onnxruntime/test/python/onnxruntime_test_python_backend.py index 946d164f8f7e4..2cfa152b94a17 100644 --- a/onnxruntime/test/python/onnxruntime_test_python_backend.py +++ b/onnxruntime/test/python/onnxruntime_test_python_backend.py @@ -43,8 +43,7 @@ def test_prepare_released_opset28_model_when_policy_enabled(self): @patch.dict(os.environ, {"ALLOW_RELEASED_ONNX_OPSET_ONLY": "1"}) def test_prepare_model_bytes_rejects_unreleased_opset29_when_policy_enabled(self): - model = self._make_opset28_model() - model.opset_import[0].version = 29 + model = self._make_opset29_model() original_policy = OnnxRuntimeBackend.allowReleasedOpsetsOnly try: OnnxRuntimeBackend.allowReleasedOpsetsOnly = True @@ -53,8 +52,9 @@ def test_prepare_model_bytes_rejects_unreleased_opset29_when_policy_enabled(self finally: OnnxRuntimeBackend.allowReleasedOpsetsOnly = original_policy - def test_prepare_model_bytes_allows_opset28_when_policy_disabled(self): - model = self._make_opset28_model() + @patch.dict(os.environ, {"ALLOW_RELEASED_ONNX_OPSET_ONLY": "1"}) + def test_prepare_model_bytes_allows_unreleased_opset29_when_policy_disabled(self): + model = self._make_opset29_model() original_policy = OnnxRuntimeBackend.allowReleasedOpsetsOnly try: @@ -66,10 +66,11 @@ def test_prepare_model_bytes_allows_opset28_when_policy_disabled(self): result = rep.run(np.array([1.0], dtype=np.float32)) np.testing.assert_array_equal(result[0], np.array([1.0], dtype=np.float32)) - def test_prepare_model_path_allows_opset28_when_policy_disabled(self): - model = self._make_opset28_model() + @patch.dict(os.environ, {"ALLOW_RELEASED_ONNX_OPSET_ONLY": "1"}) + def test_prepare_model_path_allows_unreleased_opset29_when_policy_disabled(self): + model = self._make_opset29_model() with tempfile.TemporaryDirectory() as tmpdir: - model_path = os.path.join(tmpdir, "opset28.onnx") + model_path = os.path.join(tmpdir, "opset29.onnx") with open(model_path, "wb") as model_file: model_file.write(model.SerializeToString()) @@ -83,6 +84,15 @@ def test_prepare_model_path_allows_opset28_when_policy_disabled(self): result = rep.run(np.array([1.0], dtype=np.float32)) np.testing.assert_array_equal(result[0], np.array([1.0], dtype=np.float32)) + @staticmethod + def _make_opset29_model(): + # A passthrough graph tests load policy without requiring unregistered opset-29 schemas. + value_info = helper.make_tensor_value_info("X", TensorProto.FLOAT, [1]) + return helper.make_model( + helper.make_graph([], "opset29", [value_info], [value_info]), + opset_imports=[helper.make_opsetid("", 29)], + ) + @staticmethod def _make_opset28_model(): model = helper.make_model( diff --git a/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc b/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc index d5d3feb855c8d..325fa58772f76 100644 --- a/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc +++ b/onnxruntime/test/testdata/onnx_backend_test_series_filters.jsonc @@ -240,33 +240,6 @@ "^test_unique_sorted_with_axis_cuda$", "^test_unique_sorted_with_negative_axis_cuda$", "^test_unique_sorted_without_axis_cuda$", - // TODO(#32361): Remove these exact skips when the opset-28 BitShift kernels land. - ["opset28", "^test_bitshift_left_int16(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int32(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int32_negative_shift(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int32_overflow(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int32_shift_ge_width(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int64(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int8(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int8_negative_shift(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int8_overflow(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_int8_shift_ge_width(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_uint32(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_uint64(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_left_uint8(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int16(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int32(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int32_negative_input(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int32_negative_shift(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int32_shift_ge_width(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int64(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int8(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int8_negative_input(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int8_negative_shift(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_int8_shift_ge_width(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_uint32(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_uint64(_(cpu|cuda))?$"], - ["opset28", "^test_bitshift_right_uint8(_(cpu|cuda))?$"], "^test_mod_mixed_sign_int16_cuda$", "^test_mod_mixed_sign_int8_cuda$", "^test_mod_uint16_cuda$", diff --git a/tools/ci_build/build.py b/tools/ci_build/build.py index 32711a96b3334..249fcfee6754c 100644 --- a/tools/ci_build/build.py +++ b/tools/ci_build/build.py @@ -151,8 +151,8 @@ def run_subprocess( def get_onnx_backend_test_environment(_use_cuda): return { - "ALLOW_RELEASED_ONNX_OPSET_ONLY": "0", - "ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS": "1", + "ALLOW_RELEASED_ONNX_OPSET_ONLY": "1", + "ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS": "0", } diff --git a/tools/ci_build/test_build_args.py b/tools/ci_build/test_build_args.py index b2c93ed54248f..f6814fcc06562 100644 --- a/tools/ci_build/test_build_args.py +++ b/tools/ci_build/test_build_args.py @@ -14,9 +14,9 @@ build = importlib.import_module("build") build_args = importlib.import_module("build_args") -_UNRELEASED_OPSET_ENVIRONMENT = { - "ALLOW_RELEASED_ONNX_OPSET_ONLY": "0", - "ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS": "1", +_RELEASED_OPSET_ENVIRONMENT = { + "ALLOW_RELEASED_ONNX_OPSET_ONLY": "1", + "ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS": "0", } @@ -73,23 +73,24 @@ def test_acl_deprecation_warning_not_emitted_without_use_acl(self): class OnnxBackendTestEnvironmentTest(unittest.TestCase): - def test_cpu_and_cuda_enable_unreleased_opsets_by_default(self): + def test_cpu_and_cuda_require_released_opsets_by_default(self): with mock.patch.dict(build.os.environ, {}, clear=True): for use_cuda in (False, True): with self.subTest(use_cuda=use_cuda): self.assertEqual( build.get_onnx_backend_test_environment(use_cuda), - _UNRELEASED_OPSET_ENVIRONMENT, + _RELEASED_OPSET_ENVIRONMENT, ) - def test_cpu_and_cuda_override_explicit_parent_strict_opset_mode(self): - with mock.patch.dict(build.os.environ, {"ALLOW_RELEASED_ONNX_OPSET_ONLY": "1"}, clear=True): - for use_cuda in (False, True): - with self.subTest(use_cuda=use_cuda): - self.assertEqual( - build.get_onnx_backend_test_environment(use_cuda), - _UNRELEASED_OPSET_ENVIRONMENT, - ) + def test_cpu_and_cuda_require_released_opsets_regardless_of_parent_policy(self): + for parent_policy in ("0", "1"): + with mock.patch.dict(build.os.environ, {"ALLOW_RELEASED_ONNX_OPSET_ONLY": parent_policy}, clear=True): + for use_cuda in (False, True): + with self.subTest(use_cuda=use_cuda, parent_policy=parent_policy): + self.assertEqual( + build.get_onnx_backend_test_environment(use_cuda), + _RELEASED_OPSET_ENVIRONMENT, + ) if __name__ == "__main__": From a8baf7707f125d1458efc15bb5b67ed911aff27d Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Tue, 6 Oct 2026 22:02:05 -0700 Subject: [PATCH 05/10] Resolve Mod fallback test model before serialization Populate graph inputs and outputs before passing the serialized OpTester model to the session loader. Exercise the same serialization path with CPU in every build so missing graph resolution cannot be masked by unavailable optional providers. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- .../providers/cpu/math/element_wise_ops_test.cc | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc index 5048de8c1e47b..f7cd660c7eeda 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -4398,6 +4398,7 @@ TEST(ModOpTest, FloorMod_float_special_values_opset28) { TEST(ModOpTest, FloorMod_opset28_provider_fallback) { InlinedVector> providers; + providers.push_back(DefaultCpuExecutionProvider()); #ifdef USE_CANN OrtCANNProviderOptions cann_options{}; cann_options.enable_cann_graph = 1; @@ -4407,12 +4408,10 @@ TEST(ModOpTest, FloorMod_opset28_provider_fallback) { #endif providers.push_back(DefaultMIGraphXExecutionProvider()); providers.push_back(DefaultOpenVINOExecutionProvider()); - bool tested_provider = false; for (auto& provider : providers) { if (!provider) { continue; } - tested_provider = true; SCOPED_TRACE(provider->Type()); const std::vector shape{4}; const float infinity = std::numeric_limits::infinity(); @@ -4424,17 +4423,18 @@ TEST(ModOpTest, FloorMod_opset28_provider_fallback) { test.AddInput("Y", shape, y); test.AddOutput("Z", shape, {1.0f, -1.0f, infinity, -infinity}); std::string model_data; - test.BuildModel().ToProto().SerializeToString(&model_data); + auto& model = test.BuildModel(); + ASSERT_STATUS_OK(model.MainGraph().Resolve()); + model.ToProto().SerializeToString(&model_data); EPVerificationParams params; - params.ep_node_assignment = ExpectedEPNodeAssignment::None; + params.ep_node_assignment = provider->Type() == kCpuExecutionProvider + ? ExpectedEPNodeAssignment::All + : ExpectedEPNodeAssignment::None; RunAndVerifyOutputsWithEP( AsByteSpan(model_data.data(), model_data.size()), CurrentTestName(), std::move(provider), {{"X", CreateInputOrtValueOnCPU(shape, x)}, {"Y", CreateInputOrtValueOnCPU(shape, y)}}, params); } - if (!tested_provider) { - GTEST_SKIP() << "Requires CANN, MIGraphX, or OpenVINO."; - } } #if defined(USE_CUDA) From 6acf1b143d825967266cbfcfd925d34f9c5199ea Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Wed, 7 Oct 2026 03:47:56 -0700 Subject: [PATCH 06/10] Execute Mod fallback tester before external session verification Run the CPU reference through OpTester so its Debug-only destructor guard recognizes that the test executed. Preserve the serialized-model checks for provider assignment and CPU fallback. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc | 1 + 1 file changed, 1 insertion(+) diff --git a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc index f7cd660c7eeda..a8b47051d385c 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -4422,6 +4422,7 @@ TEST(ModOpTest, FloorMod_opset28_provider_fallback) { test.AddInput("X", shape, x); test.AddInput("Y", shape, y); test.AddOutput("Z", shape, {1.0f, -1.0f, infinity, -infinity}); + test.ConfigEp(DefaultCpuExecutionProvider()).RunWithConfig(); std::string model_data; auto& model = test.BuildModel(); ASSERT_STATUS_OK(model.MainGraph().Resolve()); From b780a3438c715e1683c5764fe85ab5d7b3f51565 Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Wed, 7 Oct 2026 18:33:54 -0700 Subject: [PATCH 07/10] Update ONNX integration to 1.23.2 Use the released opset metadata fix upstream and remove the redundant local schema patch. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- cmake/deps.txt | 6 +++--- cmake/external/onnx | 2 +- cmake/patches/onnx/onnx.patch | 12 ------------ cmake/vcpkg-ports/onnx/binskim.patch | 12 ------------ cmake/vcpkg-ports/onnx/portfile.cmake | 4 ++-- cmake/vcpkg-ports/onnx/vcpkg.json | 2 +- onnxruntime/test/python/requirements.txt | 2 +- requirements-materialize-onnx-node-tests.txt | 2 +- .../aarch64/python/cpu/scripts/requirements.txt | 2 +- .../linux/docker/scripts/lort/requirements.txt | 2 +- .../linux/docker/scripts/manylinux/requirements.txt | 2 +- .../github/linux/docker/scripts/requirements.txt | 2 +- tools/ci_build/github/linux/python/requirements.txt | 2 +- .../ci_build/github/windows/python/requirements.txt | 2 +- 14 files changed, 15 insertions(+), 39 deletions(-) diff --git a/cmake/deps.txt b/cmake/deps.txt index 17f39cef00f54..3f70238cf7a72 100644 --- a/cmake/deps.txt +++ b/cmake/deps.txt @@ -34,9 +34,9 @@ microsoft_gsl;https://github.com/microsoft/GSL/archive/refs/tags/v4.2.1.zip;1094 microsoft_wil;https://github.com/microsoft/wil/archive/refs/tags/v1.0.250325.1.zip;826c8bd47c2258ec61b8b218e031e5b33d27f761 mimalloc;https://github.com/microsoft/mimalloc/archive/refs/tags/v2.1.1.zip;d5ee7d34223d0567892db5179849939c8769dc41 mp11;https://github.com/boostorg/mp11/archive/refs/tags/boost-1.82.0.zip;9bc9e01dffb64d9e0773b2e44d2f22c51aace063 -# ONNX 1.23.1 final release (opset 28). -# The 7 CI requirements.txt files are pinned to onnx==1.23.1. -onnx;https://github.com/onnx/onnx/archive/refs/tags/v1.23.1.zip;85060f3c52267c6185a5c1fa298083a514f5d201 +# ONNX 1.23.2 final release (opset 28). +# The 7 CI requirements.txt files are pinned to onnx==1.23.2. +onnx;https://github.com/onnx/onnx/archive/refs/tags/v1.23.2.zip;8846de32ab45e653c2cd651d8d6c4c299fb1583d # Use the latest commit of 10.9-GA onnx_tensorrt;https://github.com/onnx/onnx-tensorrt/archive/d5dce67db7c2e64b07e055571f5ec06f7f254de2.zip;01114d3b67650857281fa50faa2e412130a63b69 protobuf;https://github.com/protocolbuffers/protobuf/releases/download/v33.6/protobuf-33.6.tar.gz;051fee744fcca04dab37eb009d49aa3e9546acda diff --git a/cmake/external/onnx b/cmake/external/onnx index ecfb5f08b88b2..e1ac8e74b620b 160000 --- a/cmake/external/onnx +++ b/cmake/external/onnx @@ -1 +1 @@ -Subproject commit ecfb5f08b88b2c60733323bb8a691e99ba11e469 +Subproject commit e1ac8e74b620bbdc7ba07f0440806bd93f651a28 diff --git a/cmake/patches/onnx/onnx.patch b/cmake/patches/onnx/onnx.patch index d4037eee8a0af..e86a723d96a76 100644 --- a/cmake/patches/onnx/onnx.patch +++ b/cmake/patches/onnx/onnx.patch @@ -73,18 +73,6 @@ index cbb1532c0..6597a2ef9 100644 .SetDoc(GroupNormalization_ver18_doc) .Attr("epsilon", "The epsilon value to use to avoid division by zero.", AttributeProto::FLOAT, 1e-5f) .Attr( -diff --git a/onnx/defs/schema.h b/onnx/defs/schema.h ---- a/onnx/defs/schema.h -+++ b/onnx/defs/schema.h -@@ -937,7 +937,7 @@ - // Version corresponding last release of ONNX. Update this to match with - // the max version above in a *release* version of ONNX. But in other - // versions, the max version may be ahead of the last-release-version. -- last_release_version_map_[ONNX_DOMAIN] = 27; -+ last_release_version_map_[ONNX_DOMAIN] = 28; - last_release_version_map_[AI_ONNX_ML_DOMAIN] = 5; - last_release_version_map_[AI_ONNX_TRAINING_DOMAIN] = 1; - last_release_version_map_[AI_ONNX_PREVIEW_DOMAIN] = 1; diff --git a/onnx/defs/printer.cc b/onnx/defs/printer.cc index 1d43bccb0..b6beed6f3 100644 --- a/onnx/defs/printer.cc diff --git a/cmake/vcpkg-ports/onnx/binskim.patch b/cmake/vcpkg-ports/onnx/binskim.patch index d4037eee8a0af..e86a723d96a76 100644 --- a/cmake/vcpkg-ports/onnx/binskim.patch +++ b/cmake/vcpkg-ports/onnx/binskim.patch @@ -73,18 +73,6 @@ index cbb1532c0..6597a2ef9 100644 .SetDoc(GroupNormalization_ver18_doc) .Attr("epsilon", "The epsilon value to use to avoid division by zero.", AttributeProto::FLOAT, 1e-5f) .Attr( -diff --git a/onnx/defs/schema.h b/onnx/defs/schema.h ---- a/onnx/defs/schema.h -+++ b/onnx/defs/schema.h -@@ -937,7 +937,7 @@ - // Version corresponding last release of ONNX. Update this to match with - // the max version above in a *release* version of ONNX. But in other - // versions, the max version may be ahead of the last-release-version. -- last_release_version_map_[ONNX_DOMAIN] = 27; -+ last_release_version_map_[ONNX_DOMAIN] = 28; - last_release_version_map_[AI_ONNX_ML_DOMAIN] = 5; - last_release_version_map_[AI_ONNX_TRAINING_DOMAIN] = 1; - last_release_version_map_[AI_ONNX_PREVIEW_DOMAIN] = 1; diff --git a/onnx/defs/printer.cc b/onnx/defs/printer.cc index 1d43bccb0..b6beed6f3 100644 --- a/onnx/defs/printer.cc diff --git a/cmake/vcpkg-ports/onnx/portfile.cmake b/cmake/vcpkg-ports/onnx/portfile.cmake index df43d91ef0040..c1391ed9a5aed 100644 --- a/cmake/vcpkg-ports/onnx/portfile.cmake +++ b/cmake/vcpkg-ports/onnx/portfile.cmake @@ -3,8 +3,8 @@ vcpkg_check_linkage(ONLY_STATIC_LIBRARY) vcpkg_from_github( OUT_SOURCE_PATH SOURCE_PATH REPO onnx/onnx - REF "v1.23.1" - SHA512 85e377b73c412d25ebdc0e8ed5951adfb3643b20d1a75ea7e7a4186ff77dc2bcee9a6dfdc844a1f43eb935325e525b816dd4ba540fe2c045e3390acc8b16fadb + REF "v1.23.2" + SHA512 fc93f19acc87fa0f4b75897c4f7c7e3e14d29f078f6e2b1a2806c5e2454789abdb44067437bc073d612fb48a564d3e5616fbe5e254d021ccec3e301170b1781d PATCHES fix-cmakelists.patch fix-dependency-protobuf.patch diff --git a/cmake/vcpkg-ports/onnx/vcpkg.json b/cmake/vcpkg-ports/onnx/vcpkg.json index c1cb50f9bec8d..888a94bfc7f40 100644 --- a/cmake/vcpkg-ports/onnx/vcpkg.json +++ b/cmake/vcpkg-ports/onnx/vcpkg.json @@ -1,6 +1,6 @@ { "name": "onnx", - "version-semver": "1.23.1", + "version-semver": "1.23.2", "port-version": 0, "description": "Open standard for machine learning interoperability", "homepage": "https://onnx.ai", diff --git a/onnxruntime/test/python/requirements.txt b/onnxruntime/test/python/requirements.txt index 7d5b9bae47341..fc11cff5648c0 100644 --- a/onnxruntime/test/python/requirements.txt +++ b/onnxruntime/test/python/requirements.txt @@ -1,3 +1,3 @@ -onnx==1.23.1 +onnx==1.23.2 pytest onnx-ir diff --git a/requirements-materialize-onnx-node-tests.txt b/requirements-materialize-onnx-node-tests.txt index adc62024df3a9..a9c2ff3193b22 100644 --- a/requirements-materialize-onnx-node-tests.txt +++ b/requirements-materialize-onnx-node-tests.txt @@ -10,6 +10,6 @@ # PIN-LOCK: these values MUST stay in lockstep with the CMake gate (cmake/onnxruntime_unittests.cmake, # which HARD-FATALs when the installed versions differ from its expected onnx/numpy pins) and with # the Python MIN pins. Bump here and in cmake/deps.txt + the numpy MIN pins together. -onnx==1.23.1 +onnx==1.23.2 numpy==2.2.6; python_version < "3.11" numpy==2.4.2; python_version >= "3.11" diff --git a/tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt b/tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt index fa93a018f0481..6970a1a5651c1 100644 --- a/tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt +++ b/tools/ci_build/github/linux/docker/inference/aarch64/python/cpu/scripts/requirements.txt @@ -7,4 +7,4 @@ wheel protobuf==6.33.6 sympy==1.14 flatbuffers -onnx==1.23.1 +onnx==1.23.2 diff --git a/tools/ci_build/github/linux/docker/scripts/lort/requirements.txt b/tools/ci_build/github/linux/docker/scripts/lort/requirements.txt index c9c9038d2a3a4..314aaa3f6a541 100644 --- a/tools/ci_build/github/linux/docker/scripts/lort/requirements.txt +++ b/tools/ci_build/github/linux/docker/scripts/lort/requirements.txt @@ -3,7 +3,7 @@ beartype==0.15.0 flatbuffers cerberus h5py -onnx==1.23.1 +onnx==1.23.2 # Python dependencies required for pytorch development astunparse expecttest!=0.2.0 diff --git a/tools/ci_build/github/linux/docker/scripts/manylinux/requirements.txt b/tools/ci_build/github/linux/docker/scripts/manylinux/requirements.txt index 818763d672e58..ddbecaf6cb32e 100644 --- a/tools/ci_build/github/linux/docker/scripts/manylinux/requirements.txt +++ b/tools/ci_build/github/linux/docker/scripts/manylinux/requirements.txt @@ -9,4 +9,4 @@ sympy==1.14 flatbuffers neural-compressor>=2.2.1 triton==3.5.0 -onnx==1.23.1 +onnx==1.23.2 diff --git a/tools/ci_build/github/linux/docker/scripts/requirements.txt b/tools/ci_build/github/linux/docker/scripts/requirements.txt index ba56c68dd5ea5..5007bb45d1dfb 100644 --- a/tools/ci_build/github/linux/docker/scripts/requirements.txt +++ b/tools/ci_build/github/linux/docker/scripts/requirements.txt @@ -12,4 +12,4 @@ protobuf==6.33.6 packaging onnxscript==0.6.2 onnx-ir==0.1.16 -onnx==1.23.1 +onnx==1.23.2 diff --git a/tools/ci_build/github/linux/python/requirements.txt b/tools/ci_build/github/linux/python/requirements.txt index ada77b03d35a1..6f9ba1eaf4ba2 100644 --- a/tools/ci_build/github/linux/python/requirements.txt +++ b/tools/ci_build/github/linux/python/requirements.txt @@ -12,4 +12,4 @@ onnxscript==0.6.2 onnx-ir==0.1.16 jinja2 markupsafe -onnx==1.23.1 +onnx==1.23.2 diff --git a/tools/ci_build/github/windows/python/requirements.txt b/tools/ci_build/github/windows/python/requirements.txt index c0f1de5dfbb7d..24ef0cb30c081 100644 --- a/tools/ci_build/github/windows/python/requirements.txt +++ b/tools/ci_build/github/windows/python/requirements.txt @@ -14,4 +14,4 @@ jinja2 markupsafe semver packaging -onnx==1.23.1 +onnx==1.23.2 From ae295a3db141c88dff00e1569d44ffbca4b9766c Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Wed, 7 Oct 2026 18:33:56 -0700 Subject: [PATCH 08/10] Keep opset-28 BitShift and SpaceToDepth routing safe Reject unverified provider semantics, expand SpaceToDepth-28 through its ONNX function, and preserve pre-28 fusion and native behavior. Check literal DCR/CRD outputs and provider routing. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- docs/OperatorKernels.md | 6 +- .../slice_concat_to_space_to_depth_fusion.cc | 5 ++ onnxruntime/core/providers/cann/cann_graph.cc | 2 +- .../providers/cpu/cpu_execution_provider.cc | 4 +- .../providers/cpu/tensor/space_depth_ops.cc | 3 +- .../providers/cuda/cuda_execution_provider.cc | 4 +- .../core/providers/cuda/cuda_nhwc_kernels.cc | 4 +- .../providers/cuda/tensor/space_depth_ops.cc | 6 +- .../src/Operators/DmlOperatorSpaceToDepth.cpp | 6 ++ .../src/Operators/OperatorRegistration.cpp | 5 +- .../migraphx/migraphx_execution_provider.cc | 4 +- .../openvino/ov_versions/data_ops.cc | 4 +- .../test/optimizer/graph_transform_test.cc | 25 ++++-- .../cpu/math/element_wise_ops_test.cc | 32 +++++++ .../cpu/tensor/space_depth_ops_test.cc | 85 +++++++++++++++++++ 15 files changed, 168 insertions(+), 27 deletions(-) diff --git a/docs/OperatorKernels.md b/docs/OperatorKernels.md index 6d8cfdd113d77..6d5f8fc2423de 100644 --- a/docs/OperatorKernels.md +++ b/docs/OperatorKernels.md @@ -480,7 +480,7 @@ The **OpSet Version** column uses the following notation: |||[1, 21]|**T** = tensor(float)| |Softsign|*in* input:**T**
*out* output:**T**|22+|**T** = tensor(float)| |||[1, 21]|**T** = tensor(float)| -|SpaceToDepth|*in* input:**T**
*out* output:**T**|13+|**T** = tensor(double), tensor(float), tensor(int8), tensor(uint8)| +|SpaceToDepth|*in* input:**T**
*out* output:**T**|[13, 27]|**T** = tensor(double), tensor(float), tensor(int8), tensor(uint8)| |||[1, 12]|**T** = tensor(double), tensor(float), tensor(int8), tensor(uint8)| |Split|*in* input:**T**
*in* split:**T**
*out* outputs...:**T**

or

*in* input:**T**
*in* split:**tensor(int64)**
*out* outputs:**T**

or

*in* input:**T**
*out* outputs:**T**|18+|**T** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||[13, 17]|**T** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(string), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| @@ -1021,7 +1021,7 @@ The **OpSet Version** column uses the following notation: |||[1, 21]|**T** = tensor(double), tensor(float), tensor(float16)| |Softsign|*in* input:**T**
*out* output:**T**|22+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16)| |||[1, 21]|**T** = tensor(double), tensor(float), tensor(float16)| -|SpaceToDepth|*in* input:**T**
*out* output:**T**|13+|**T** = tensor(double), tensor(float), tensor(float16)| +|SpaceToDepth|*in* input:**T**
*out* output:**T**|[13, 27]|**T** = tensor(double), tensor(float), tensor(float16)| |||[1, 12]|**T** = tensor(double), tensor(float), tensor(float16)| |Split|*in* input:**T**
*in* split:**T**
*out* outputs...:**T**

or

*in* input:**T**
*in* split:**tensor(int64)**
*out* outputs:**T**

or

*in* input:**T**
*out* outputs:**T**|18+|**T** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||[13, 17]|**T** = tensor(bfloat16), tensor(bool), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| @@ -1202,7 +1202,7 @@ The **OpSet Version** column uses the following notation: |||10|**I** = tensor(int64)
**T** = tensor(float), tensor(float16)| |||[8, 9]|**I** = tensor(int64)
**T** = tensor(float), tensor(float16)| |||[1, 7]|**T** = tensor(float), tensor(float16)| -|SpaceToDepth|*in* input:**T**
*out* output:**T**|13+|**T** = tensor(double), tensor(float), tensor(float16)| +|SpaceToDepth|*in* input:**T**
*out* output:**T**|[13, 27]|**T** = tensor(double), tensor(float), tensor(float16)| |||[1, 12]|**T** = tensor(double), tensor(float), tensor(float16)| | | | | diff --git a/onnxruntime/core/optimizer/slice_concat_to_space_to_depth_fusion.cc b/onnxruntime/core/optimizer/slice_concat_to_space_to_depth_fusion.cc index 8caea2c150990..3df3b985431f6 100644 --- a/onnxruntime/core/optimizer/slice_concat_to_space_to_depth_fusion.cc +++ b/onnxruntime/core/optimizer/slice_concat_to_space_to_depth_fusion.cc @@ -399,6 +399,11 @@ NodeArg* CreateInt64Initializer(Graph& graph, } bool FuseSliceConcatToSpaceToDepth(Node& concat, Graph& graph, const logging::Logger& logger) { + // Fusion runs after function expansion, and SpaceToDepth-28 has no native CPU kernel. + if (graph.DomainToVersionMap().at(kOnnxDomain) >= 28) { + return false; + } + if (!graph_utils::IsSupportedOptypeVersionAndDomain(concat, "Concat", {4, 11, 13}, kOnnxDomain) || concat.InputDefs().size() != 4) { return false; diff --git a/onnxruntime/core/providers/cann/cann_graph.cc b/onnxruntime/core/providers/cann/cann_graph.cc index 064bcfd706336..d431c89d4d219 100644 --- a/onnxruntime/core/providers/cann/cann_graph.cc +++ b/onnxruntime/core/providers/cann/cann_graph.cc @@ -71,7 +71,7 @@ std::vector SupportONNXModel(const GraphViewer& graph_viewer) { if (node->Domain() != kOnnxDomain || domain_version < lower_bound || !cann_supported_ops.count(node->OpType()) || - (node->OpType() == "Mod" && node->SinceVersion() >= 28)) { + ((node->OpType() == "Mod" || node->OpType() == "SpaceToDepth") && node->SinceVersion() >= 28)) { unsupported_nodes.push_back(index); continue; } diff --git a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc index c9d5fcdcd8695..701424a532d07 100644 --- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc +++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc @@ -873,7 +873,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, double, Log); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 14, Pow); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, DepthToSpace); -class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, SpaceToDepth); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 27, SpaceToDepth); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, Slice); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 17, Split); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 20, Unsqueeze); @@ -2830,7 +2830,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo()}), SpaceToDepth); -ONNX_CPU_OPERATOR_KERNEL( +ONNX_CPU_OPERATOR_VERSIONED_KERNEL( SpaceToDepth, 13, + 27, KernelDefBuilder() .TypeConstraint("T", {DataTypeImpl::GetTensorType(), DataTypeImpl::GetTensorType(), diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index b1098d97cdad0..7ce6bebfa1965 100755 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -1327,7 +1327,7 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kO class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 17, double, Pad); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 17, MLFloat16, Pad); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 17, bool, Pad); -class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, SpaceToDepth); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 27, SpaceToDepth); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, DepthToSpace); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, int8_t, Sign); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, int16_t, Sign); @@ -2643,7 +2643,7 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/onnxruntime/core/providers/cuda/cuda_nhwc_kernels.cc b/onnxruntime/core/providers/cuda/cuda_nhwc_kernels.cc index 388a30d4a289f..22f14b6123eb8 100755 --- a/onnxruntime/core/providers/cuda/cuda_nhwc_kernels.cc +++ b/onnxruntime/core/providers/cuda/cuda_nhwc_kernels.cc @@ -85,7 +85,7 @@ class CUDA_NHWC_OP_VERSIONED_CLASS_NAME(1, 10, DepthToSpace); class CUDA_NHWC_OP_VERSIONED_CLASS_NAME(11, 12, DepthToSpace); class CUDA_NHWC_OP_CLASS_NAME(13, DepthToSpace); class CUDA_NHWC_OP_VERSIONED_CLASS_NAME(1, 12, SpaceToDepth); -class CUDA_NHWC_OP_CLASS_NAME(13, SpaceToDepth); +class CUDA_NHWC_OP_VERSIONED_CLASS_NAME(13, 27, SpaceToDepth); class CUDA_NHWC_OP_VERSIONED_TYPED_CLASS_NAME(1, 12, float, LRN); class CUDA_NHWC_OP_VERSIONED_TYPED_CLASS_NAME(1, 12, double, LRN); class CUDA_NHWC_OP_VERSIONED_TYPED_CLASS_NAME(1, 12, MLFloat16, LRN); @@ -154,7 +154,7 @@ Status RegisterCudaNhwcKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/onnxruntime/core/providers/cuda/tensor/space_depth_ops.cc b/onnxruntime/core/providers/cuda/tensor/space_depth_ops.cc index cbfc62494fde4..3c5fe31f35309 100644 --- a/onnxruntime/core/providers/cuda/tensor/space_depth_ops.cc +++ b/onnxruntime/core/providers/cuda/tensor/space_depth_ops.cc @@ -37,10 +37,11 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX( SpaceToDepth); #endif -ONNX_OPERATOR_KERNEL_EX( +ONNX_OPERATOR_VERSIONED_KERNEL_EX( SpaceToDepth, kOnnxDomain, 13, + 27, kCudaExecutionProvider, (*KernelDefBuilder::Create()) .TypeConstraint("T", @@ -50,10 +51,11 @@ ONNX_OPERATOR_KERNEL_EX( SpaceToDepth); #ifdef ENABLE_CUDA_NHWC_OPS -ONNX_OPERATOR_KERNEL_EX( +ONNX_OPERATOR_VERSIONED_KERNEL_EX( SpaceToDepth, kMSInternalNHWCDomain, 13, + 27, kCudaExecutionProvider, (*KernelDefBuilder::Create()) .TypeConstraint("T", diff --git a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/DmlOperatorSpaceToDepth.cpp b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/DmlOperatorSpaceToDepth.cpp index 9d1faef3fd2a6..81eff170f1312 100644 --- a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/DmlOperatorSpaceToDepth.cpp +++ b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/DmlOperatorSpaceToDepth.cpp @@ -30,6 +30,12 @@ class DmlOperatorSpaceToDepth : public DmlOperator, public SpaceToDepthHelper } }; +void CALLBACK QuerySpaceToDepth(IMLOperatorSupportQueryContextPrivate* context, bool* isSupported) +{ + MLOperatorAttributes attributes(context); + *isSupported = attributes.GetOptionalAttribute(AttrName::Mode, "DCR") == "DCR"; +} + DML_OP_DEFINE_CREATION_FUNCTION(SpaceToDepth, DmlOperatorSpaceToDepth); } // namespace Dml diff --git a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp index b4c05dcfaa51a..c5e32af3bc90d 100644 --- a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp +++ b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp @@ -543,6 +543,7 @@ DML_OP_EXTERN_CREATION_FUNCTION(MatMulNBits); DML_OP_EXTERN_QUERY_FUNCTION(MaxPool); DML_OP_EXTERN_QUERY_FUNCTION(Slice); DML_OP_EXTERN_QUERY_FUNCTION(Resize); +DML_OP_EXTERN_QUERY_FUNCTION(SpaceToDepth); DML_OP_EXTERN_QUERY_FUNCTION(Scatter); DML_OP_EXTERN_QUERY_FUNCTION(EinSum); DML_OP_EXTERN_QUERY_FUNCTION(RecurrentNeuralNetwork); @@ -813,8 +814,8 @@ constexpr static OperatorRegistrationInformation operatorRegistrationInformation {REG_INFO_VER( 19, Pad, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported, requiredConstantCpuInputs(1, 2, 3) /*pads, value, axes*/)}, #endif - {REG_INFO( 7, SpaceToDepth, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported)}, - {REG_INFO( 13, SpaceToDepth, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported)}, + {REG_INFO( 7, SpaceToDepth, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported, requiredConstantCpuInputs(), std::nullopt, QuerySpaceToDepth)}, + {REG_INFO( 13, SpaceToDepth, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported, requiredConstantCpuInputs(), std::nullopt, QuerySpaceToDepth)}, {REG_INFO( 7, DepthToSpace, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported)}, {REG_INFO( 11, DepthToSpace, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported)}, {REG_INFO( 13, DepthToSpace, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported)}, diff --git a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc index b2e8efa883275..72c5676c20fa0 100644 --- a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc +++ b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc @@ -377,8 +377,8 @@ std::vector toVector(const ONNX_NAMESPACE::int64s& nums) { static bool IsUnsupportedOpMode(const onnxruntime::GraphViewer& graph_viewer, const Node* node) { std::vector input_nodes; const auto& optype = node->OpType(); - // Mod-28 floating floor/edge-case semantics require the CPU fallback. - if (optype == "Mod" && node->SinceVersion() >= 28) { + // These opset-28 semantics require CPU kernels or ONNX function expansion. + if ((optype == "Mod" || optype == "SpaceToDepth") && node->SinceVersion() >= 28) { return true; } diff --git a/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc index 49797978471f8..24c5388fe25c7 100644 --- a/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc +++ b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc @@ -727,8 +727,8 @@ bool DataOps::dimension_unsupported(const Node* node) { bool DataOps::node_is_supported(const NodeIndex node_idx, bool& has_external_weights_) { const auto& node = graph_viewer_.GetNode(node_idx); const auto& optype = node->OpType(); - // Mod-28 floating floor/edge-case semantics require the CPU fallback. - if (optype == "Mod" && node->SinceVersion() >= 28) { + // These opset-28 semantics require CPU kernels or ONNX function expansion. + if ((optype == "Mod" || optype == "BitShift" || optype == "SpaceToDepth") && node->SinceVersion() >= 28) { return false; } diff --git a/onnxruntime/test/optimizer/graph_transform_test.cc b/onnxruntime/test/optimizer/graph_transform_test.cc index c303dfc5e4954..b67c014d9fb4a 100644 --- a/onnxruntime/test/optimizer/graph_transform_test.cc +++ b/onnxruntime/test/optimizer/graph_transform_test.cc @@ -7470,6 +7470,12 @@ TEST_F(GraphTransformationTests, SliceConcatToSpaceToDepthFusionTest) { auto check_transformed_graph = [get_op_count](InferenceSessionWrapper& session) { const Graph& graph = session.GetGraph(); const auto op_to_count = CountOpsInGraph(graph); + if (graph.DomainToVersionMap().at(kOnnxDomain) >= 28) { + ASSERT_EQ(get_op_count(op_to_count, "Slice"), 4); + ASSERT_EQ(get_op_count(op_to_count, "Concat"), 1); + ASSERT_EQ(get_op_count(op_to_count, "SpaceToDepth"), 0); + return; + } ASSERT_TRUE(op_to_count.count("Slice") == 0 || op_to_count.at("Slice") == 0); ASSERT_TRUE(op_to_count.count("Concat") == 0 || op_to_count.at("Concat") == 0); ASSERT_EQ(get_op_count(op_to_count, "SpaceToDepth"), 1); @@ -7482,14 +7488,17 @@ TEST_F(GraphTransformationTests, SliceConcatToSpaceToDepthFusionTest) { } }; - TransformerTester(build_test_case, - check_transformed_graph, - TransformerLevel::Default, - TransformerLevel::Level1, - 13, - 0.0, - 0.0, - std::make_unique()); + for (int opset : {13, 27, 28}) { + SCOPED_TRACE(opset); + TransformerTester(build_test_case, + check_transformed_graph, + TransformerLevel::Default, + TransformerLevel::Level1, + opset, + 0.0, + 0.0, + std::make_unique()); + } } TEST_F(GraphTransformationTests, SliceConcatToSpaceToDepthFusionWithConstantNodesTest) { diff --git a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc index a8b47051d385c..c15df7ddc90cb 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -4796,6 +4796,38 @@ TEST(BitShiftOpTest, SignedInt64Opset28) { TestSignedBitShiftOpset28(); } +template +void TestBitShiftOpset28ProviderFallback(const std::vector& x, const std::vector& y, + const std::vector& expected) { + const std::vector shape{static_cast(x.size())}; + OpTester test("BitShift", 28); + test.AddAttribute("direction", "RIGHT"); + test.AddInput("X", shape, x); + test.AddInput("Y", shape, y); + test.AddOutput("Z", shape, expected); + test.ConfigEp(DefaultCpuExecutionProvider()).RunWithConfig(); + auto provider = DefaultOpenVINOExecutionProvider(); + if (!provider) { + return; + } + + auto& model = test.BuildModel(); + ASSERT_STATUS_OK(model.MainGraph().Resolve()); + std::string model_data; + ASSERT_TRUE(model.ToProto().SerializeToString(&model_data)); + EPVerificationParams params; + params.ep_node_assignment = ExpectedEPNodeAssignment::None; + RunAndVerifyOutputsWithEP( + AsByteSpan(model_data.data(), model_data.size()), CurrentTestName(), std::move(provider), + {{"X", CreateInputOrtValueOnCPU(shape, x)}, {"Y", CreateInputOrtValueOnCPU(shape, y)}}, + params); +} + +TEST(BitShiftOpTest, Opset28ProviderFallback) { + TestBitShiftOpset28ProviderFallback({-8, -3, -2, -2}, {1, 2, -1, 32}, {-4, -1, -1, -1}); + TestBitShiftOpset28ProviderFallback({8, 255, 1, 42}, {1, 32, 33, 64}, {4, 0, 0, 0}); +} + template void TestUnsignedBitShiftRegistration(int opset) { OpTester test("BitShift", opset); diff --git a/onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc b/onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc index ac7d896258603..7df530e0ff078 100644 --- a/onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc +++ b/onnxruntime/test/providers/cpu/tensor/space_depth_ops_test.cc @@ -9,6 +9,9 @@ #include "core/framework/allocator.h" #include "core/providers/cpu/tensor/space_depth_ops.h" #include "core/mlas/inc/mlas.h" +#include "test/unittest_util/graph_transform_test_builder.h" +#include "test/util/include/default_providers.h" +#include "test/util/include/inference_session_wrapper.h" namespace onnxruntime { namespace test { @@ -121,6 +124,88 @@ TEST(TensorOpTest, SpaceToDepthTest_1_double) { test.AddOutput("output", {N, C * blocksize * blocksize, H / blocksize, W / blocksize}, result); test.Run(); } + +TEST(TensorOpTest, SpaceToDepthOpset28ModesAndProviderRouting) { + const std::vector input{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15}; + const std::vector dcr{0, 2, 8, 10, 1, 3, 9, 11, 4, 6, 12, 14, 5, 7, 13, 15}; + const std::vector crd{0, 2, 1, 3, 4, 6, 5, 7, 8, 10, 9, 11, 12, 14, 13, 15}; + for (int opset : {13, 27, 28}) { + for (const std::string mode : {"", "DCR", "CRD"}) { + if (opset < 28 && !mode.empty()) { + continue; + } + SCOPED_TRACE(opset); + SCOPED_TRACE(mode); + Model model("SpaceToDepthRouting", false, ModelMetaData(), PathString(), + IOnnxRuntimeOpSchemaRegistryList(), {{kOnnxDomain, opset}}, {}, + DefaultLoggingManager().DefaultLogger()); + ModelTestBuilder builder(model.MainGraph()); + auto* x = builder.MakeInput({1, 2, 2, 4}, input); + auto* z = builder.MakeOutput(); + auto& node = builder.AddNode("SpaceToDepth", {x}, {z}); + node.AddAttribute("blocksize", int64_t{2}); + if (!mode.empty()) { + node.AddAttribute("mode", mode); + } + builder.SetGraphOutputs(); + ASSERT_STATUS_OK(model.MainGraph().Resolve()); + std::string model_data; + ASSERT_TRUE(model.ToProto().SerializeToString(&model_data)); + + InlinedVector> providers; + providers.push_back(DefaultCpuExecutionProvider()); + providers.push_back(DefaultCudaExecutionProvider()); +#ifdef ENABLE_CUDA_NHWC_OPS + providers.push_back(DefaultCudaNHWCExecutionProvider()); +#endif + providers.push_back(DefaultDmlExecutionProvider()); + providers.push_back(DefaultOpenVINOExecutionProvider()); + providers.push_back(DefaultMIGraphXExecutionProvider()); +#ifdef USE_CANN + OrtCANNProviderOptions cann_options{}; + cann_options.enable_cann_graph = 1; + auto cann_factory = CannProviderFactoryCreator::Create(&cann_options); + ASSERT_NE(cann_factory, nullptr); + providers.push_back(cann_factory->CreateProvider()); +#endif + for (auto& provider : providers) { + if (!provider) { + continue; + } + const std::string provider_type = provider->Type(); + SCOPED_TRACE(provider_type); + SessionOptions options; + options.graph_optimization_level = TransformerLevel::Default; + options.enable_mem_pattern = false; + options.execution_mode = ExecutionMode::ORT_SEQUENTIAL; + InferenceSessionWrapper session{options, GetEnvironment()}; + ASSERT_STATUS_OK(session.RegisterExecutionProvider(std::move(provider))); + ASSERT_STATUS_OK(session.Load(model_data.data(), static_cast(model_data.size()))); + ASSERT_STATUS_OK(session.Initialize()); + bool has_space_to_depth = false; + for (const auto& resolved_node : session.GetGraph().Nodes()) { + has_space_to_depth |= resolved_node.OpType() == "SpaceToDepth"; + } + if (opset == 28 && (mode == "CRD" || provider_type == kCpuExecutionProvider || + provider_type == kCudaExecutionProvider)) { + EXPECT_FALSE(has_space_to_depth); + } else if (opset < 28 && provider_type == kCpuExecutionProvider) { + EXPECT_TRUE(has_space_to_depth); + } + + std::vector fetches; + ASSERT_STATUS_OK(session.Run(RunOptions{}, builder.feeds_, builder.output_names_, &fetches)); + ASSERT_EQ(fetches.size(), 1u); + const auto& tensor = fetches[0].Get(); + EXPECT_EQ(tensor.Shape(), TensorShape({1, 8, 1, 2})); + const auto& expected = mode == "CRD" ? crd : dcr; + const auto* actual = tensor.Data(); + EXPECT_EQ(std::vector(actual, actual + tensor.Shape().Size()), expected); + } + } + } +} + TEST(TensorOpTest, SpaceToDepthTest_2) { OpTester test("SpaceToDepth"); constexpr int64_t blocksize = 3; From 9ec563988f8af815a4f364731307338abb71f118 Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Wed, 7 Oct 2026 18:33:57 -0700 Subject: [PATCH 09/10] Use ABI-safe tensor type conversion in ONNX node runner Reuse shared wire-number conversion for INT2, UINT2, and FLOAT8E8M0 after the main merge. This fixes 34 native corpus failures without adding skips. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- onnxruntime/test/onnx/tensorprotoutils.cc | 39 ++--------------------- 1 file changed, 3 insertions(+), 36 deletions(-) diff --git a/onnxruntime/test/onnx/tensorprotoutils.cc b/onnxruntime/test/onnx/tensorprotoutils.cc index c9cbc05b2f391..ac465b7db2ae7 100644 --- a/onnxruntime/test/onnx/tensorprotoutils.cc +++ b/onnxruntime/test/onnx/tensorprotoutils.cc @@ -19,6 +19,7 @@ #include "core/framework/endian_utils.h" #include "core/graph/onnx_protobuf.h" #include "core/session/onnxruntime_cxx_api.h" +#include "core/session/onnxruntime_type_conversion.h" #include "mem_buffer.h" struct OrtStatus { @@ -621,42 +622,8 @@ ORT_API(void, OrtUninitializeBuffer, _In_opt_ void* input, size_t input_len, enu (Y*)preallocated, static_cast(tensor_size)); \ break; -#define CASE_TYPE(X) \ - case onnx::TensorProto_DataType_##X: \ - return ONNX_TENSOR_ELEMENT_DATA_TYPE_##X; - ONNXTensorElementDataType CApiElementTypeFromProtoType(int type) { - switch (type) { - CASE_TYPE(FLOAT) - CASE_TYPE(UINT8) - CASE_TYPE(INT8) - CASE_TYPE(UINT16) - CASE_TYPE(INT16) - CASE_TYPE(INT32) - CASE_TYPE(INT64) - CASE_TYPE(STRING) - CASE_TYPE(BOOL) - CASE_TYPE(FLOAT16) - CASE_TYPE(DOUBLE) - CASE_TYPE(UINT32) - CASE_TYPE(UINT64) - CASE_TYPE(COMPLEX64) - CASE_TYPE(COMPLEX128) - CASE_TYPE(BFLOAT16) - CASE_TYPE(FLOAT8E4M3FN) - CASE_TYPE(FLOAT8E4M3FNUZ) - CASE_TYPE(FLOAT8E5M2) - CASE_TYPE(FLOAT8E5M2FNUZ) - CASE_TYPE(FLOAT4E2M1) - CASE_TYPE(UINT4) - CASE_TYPE(INT4) - CASE_TYPE(UINT2) - CASE_TYPE(INT2) - CASE_TYPE(FLOAT6E2M3) - CASE_TYPE(FLOAT6E3M2) - default: - return ONNX_TENSOR_ELEMENT_DATA_TYPE_UNDEFINED; - } + return utils::ToOrtTensorElementDataType(type); } ONNXTensorElementDataType GetTensorElementType(const onnx::TensorProto& tensor_proto) { @@ -750,7 +717,7 @@ Status TensorProtoToMLValue(const onnx::TensorProto& tensor_proto, const MemBuff } std::vector tensor_shape_vec = GetTensorShapeFromTensorProto(tensor_proto); // Note: We permit an empty tensor_shape_vec, and treat it as a scalar (a tensor of size 1). - value = Ort::Value::CreateTensor(&allocator, tensor_data, m.GetLen(), tensor_shape_vec.data(), tensor_shape_vec.size(), (ONNXTensorElementDataType)tensor_proto.data_type()); + value = Ort::Value::CreateTensor(&allocator, tensor_data, m.GetLen(), tensor_shape_vec.data(), tensor_shape_vec.size(), ele_type); return Status::OK(); } From 59f21887a70d8eb44a00fec7323bddb29f845d8b Mon Sep 17 00:00:00 2001 From: Ganesan Ramalingam Date: Fri, 9 Oct 2026 08:15:49 -0700 Subject: [PATCH 10/10] Fix registration validation and WASM external-initializer tests Accept the exact SpaceToDepth native-kernel boundary before opset-28 function expansion and cover it with validator regressions. Keep external-buffer compilation and in-memory subgraph coverage in WASM without attempting unsupported host-filesystem external-data loads. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- .../test/framework/inference_session_test.cc | 7 +++- tools/ci_build/op_registration_validator.py | 11 ++++- .../test_op_registration_validator.py | 42 +++++++++++++++++++ 3 files changed, 58 insertions(+), 2 deletions(-) create mode 100644 tools/ci_build/test_op_registration_validator.py diff --git a/onnxruntime/test/framework/inference_session_test.cc b/onnxruntime/test/framework/inference_session_test.cc index 7b1ceefbb499b..dd4a9c8965504 100644 --- a/onnxruntime/test/framework/inference_session_test.cc +++ b/onnxruntime/test/framework/inference_session_test.cc @@ -5243,6 +5243,10 @@ TEST(InferenceSessionTests, CompileApiExternalBufferKeepsSubgraphInitializersEmb } #if defined(DISABLE_EXTERNAL_INITIALIZERS) if (buffer_destination && has_main_initializer) continue; +#endif +#if defined(__EMSCRIPTEN__) + // File-backed subgraph tensors require a Module.MountedFiles mapping in WASM. + if (!buffer_destination) continue; #endif { std::ofstream output(input_path, std::ios::binary); @@ -5291,7 +5295,8 @@ TEST(InferenceSessionTests, CompileApiExternalBufferKeepsSubgraphInitializersEmb } } -#if !defined(DISABLE_EXTERNAL_INITIALIZERS) +// Subgraph and attribute external tensors in this test load from the host filesystem. +#if !defined(DISABLE_EXTERNAL_INITIALIZERS) && !defined(__EMSCRIPTEN__) TEST(InferenceSessionTests, ExternalInitializerInjectionIsMainGraphOnly) { const std::filesystem::path model_dir = ORT_TSTR("external_buffers_attribute_model_dir"); std::filesystem::remove_all(model_dir); diff --git a/tools/ci_build/op_registration_validator.py b/tools/ci_build/op_registration_validator.py index 6cc7f3bb5ec6d..30dcd73a8bb2d 100644 --- a/tools/ci_build/op_registration_validator.py +++ b/tools/ci_build/op_registration_validator.py @@ -33,6 +33,11 @@ "kOnnxDomain:ThresholdedRelu": 10, } +# Native kernels stop before these versions; newer schemas execute via function expansion. +function_only_ops = { + "kOnnxDomain:SpaceToDepth": 28, +} + @dataclasses.dataclass class RegistrationInfo: @@ -160,7 +165,7 @@ def _validate_last_registration(self, last_r: RegistrationInfo) -> bool: """ Validates the last registration in sorted order for a single domain and op and returns whether it is valid. """ - # make sure we have an unversioned last entry for each operator unless it's deprecated + # Require an unversioned last entry unless a known native-kernel boundary applies. # TODO If the operator is deprecated, validation is more lax. I.e., it doesn't require a versioned registration. # This could be tightened up but we would need to handle the deprecated contrib ops registered in the ONNX @@ -173,6 +178,10 @@ def _validate_last_registration(self, last_r: RegistrationInfo) -> bool: deprecation_version is not None and last_r.end_version == deprecation_version - 1 ) + function_version = function_only_ops.get(domain_and_op_str) + if function_version is not None and last_r.end_version == function_version - 1: + allow_missing_unversioned_registration = True + # special handling for ArgMin/ArgMax, which CUDA EP doesn't yet support for opset 12+ # TODO remove once CUDA EP supports ArgMin/ArgMax for opset 12+ ops_with_incomplete_support = ["kOnnxDomain:ArgMin", "kOnnxDomain:ArgMax"] diff --git a/tools/ci_build/test_op_registration_validator.py b/tools/ci_build/test_op_registration_validator.py new file mode 100644 index 0000000000000..572023d05f597 --- /dev/null +++ b/tools/ci_build/test_op_registration_validator.py @@ -0,0 +1,42 @@ +#!/usr/bin/env python3 +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. + +import sys +import unittest +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parents[1] / "python")) +sys.path.insert(0, str(Path(__file__).resolve().parent)) + +from op_registration_validator import RegistrationValidator + + +class RegistrationValidatorTest(unittest.TestCase): + def _validate(self, operator, versions): + validator = RegistrationValidator() + for start_version, end_version in versions: + validator.process_registration([], "kOnnxDomain", operator, start_version, end_version) + return validator.ok() + + def test_space_to_depth_native_kernels_end_before_function_version(self): + self.assertTrue(self._validate("SpaceToDepth", [(1, 12), (13, 27)])) + + def test_space_to_depth_function_exception_requires_exact_boundary(self): + for end_version in (26, 28): + with self.subTest(end_version=end_version): + self.assertFalse(self._validate("SpaceToDepth", [(1, 12), (13, end_version)])) + + def test_space_to_depth_function_exception_does_not_allow_registration_gaps(self): + self.assertFalse(self._validate("SpaceToDepth", [(1, 11), (13, 27)])) + + def test_space_to_depth_accepts_future_native_kernel(self): + self.assertTrue(self._validate("SpaceToDepth", [(1, 12), (13, 27), (28, None)])) + + def test_other_operators_still_require_unversioned_registration(self): + self.assertFalse(self._validate("DepthToSpace", [(1, 12), (13, 27)])) + self.assertTrue(self._validate("DepthToSpace", [(1, 12), (13, None)])) + + +if __name__ == "__main__": + unittest.main()