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 467f2c2d5d1a3..2be571e06fbe5 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.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 2bb50465112fe..e1ac8e74b620b 160000
--- a/cmake/external/onnx
+++ b/cmake/external/onnx
@@ -1 +1 @@
-Subproject commit 2bb50465112feca9003e1ed654d77f01ff1415ca
+Subproject commit e1ac8e74b620bbdc7ba07f0440806bd93f651a28
diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake
index d7f82262d3a5e..238c670a02a96 100644
--- a/cmake/onnxruntime_unittests.cmake
+++ b/cmake/onnxruntime_unittests.cmake
@@ -1942,13 +1942,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
@@ -1976,27 +1976,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}")
@@ -2074,6 +2057,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 3f3f0aad774e1..e86a723d96a76 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 $<TARGET_OBJECTS:onnx_proto>)
@@ -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,146 +73,40 @@ 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/common/assertions.h b/onnx/common/assertions.h
---- a/onnx/common/assertions.h
-+++ b/onnx/common/assertions.h
-@@ -27,9 +27,9 @@
- #endif
- std::string barf(const char* fmt, ...);
-
--[[noreturn]] void throw_assert_error(std::string& /*msg*/);
-+[[noreturn]] void throw_assert_error(const std::string& /*msg*/);
-
--[[noreturn]] void throw_tensor_error(std::string& /*msg*/);
-+[[noreturn]] void throw_tensor_error(const std::string& /*msg*/);
-
- } // namespace ONNX_NAMESPACE
-
-diff --git a/onnx/common/assertions.cc b/onnx/common/assertions.cc
---- a/onnx/common/assertions.cc
-+++ b/onnx/common/assertions.cc
-@@ -42,11 +42,11 @@
- }
-
--void throw_assert_error(std::string& msg) {
-+void throw_assert_error(const std::string& msg) {
-     ONNX_THROW_EX(assert_error(msg));
- }
-
--void throw_tensor_error(std::string& msg) {
-+void throw_tensor_error(const std::string& msg) {
-     ONNX_THROW_EX(tensor_error(msg));
+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 <cmath>
+ #include <cstdint>
+ #include <iomanip>
++#include <limits>
++#include <locale>
++#include <sstream>
+ #include <string>
+ #include <type_traits>
+ #include <vector>
+@@ -31,6 +34,14 @@ std::vector<Widened> ParseRawIntData(const TensorProto& tensor) {
+ // point, the shortest form that round-trips through the parser.
+ template <typename T>
+ 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<T>::max_digits10) << value;
++  os << value_stream.str();
++#else
+   if constexpr (std::is_floating_point_v<T>) {
+     // 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<char, 32> buf{};
+   const auto res = std::to_chars(buf.data(), buf.data() + buf.size(), value);
+   os.write(buf.data(), res.ptr - buf.data());
++#endif
  }
-
- } // namespace ONNX_NAMESPACE
-diff --git a/onnx/common/tensor.h b/onnx/common/tensor.h
---- a/onnx/common/tensor.h
-+++ b/onnx/common/tensor.h
-@@ -51,13 +51,13 @@
-     /// size_from_dim() cannot handle this case, while elem_num() handles it correctly
-     int64_t elem_num() const {
--    return safe_dim_product(sizes_, [](const char* msg) { throw tensor_error(msg); });
-+    return safe_dim_product(sizes_, [](const char* msg) { throw_tensor_error(msg); });
-     }
-     int64_t size_from_dim(int dim) const {
-       if (dim < 0) {
-           dim += static_cast<int>(sizes_.size());
-       }
-       ONNX_ASSERT(dim >= 0 && (size_t)dim < sizes_.size())
--    return safe_dim_product(sizes_.begin() + dim, sizes_.end(), [](const char* msg) { throw tensor_error(msg); });
-+    return safe_dim_product(sizes_.begin() + dim, sizes_.end(), [](const char* msg) { throw_tensor_error(msg); });
-     }
-
-     int32_t elem_type() const {
-diff --git a/onnx/common/scoped_resource.h b/onnx/common/scoped_resource.h
---- a/onnx/common/scoped_resource.h
-+++ b/onnx/common/scoped_resource.h
-@@ -25,18 +25,18 @@
- // Usage:
--//   using ScopedFd = ScopedResource<-1, close_fd>;
--//   ScopedFd guard(fd);          // destructor calls close_fd(fd)
-+//   using ScopedFd = ScopedResource<FdTraits>;
-+//   ScopedFd guard(fd);          // destructor calls FdTraits::close(fd)
- //   int raw = guard.release();   // relinquish ownership
--template <auto Invalid, void (*Close)(decltype(Invalid))>
-+template <typename Traits>
- class ScopedResource {
--  using T = decltype(Invalid);
-+  using T = typename Traits::type;
-     T val_;
-
-  public:
-     explicit ScopedResource(T v) : val_(v) {}
-     ~ScopedResource() {
--    if (val_ != Invalid) {
--      Close(val_);
-+    if (val_ != Traits::invalid()) {
-+      Traits::close(val_);
-       }
-     }
-     T get() const {
-       return val_;
-@@ -44,7 +44,7 @@
-     }
-     T release() {
-       T tmp = val_;
--    val_ = Invalid;
-+    val_ = Traits::invalid();
-       return tmp;
-     }
-     ScopedResource(const ScopedResource&) = delete;
-@@ -55,20 +55,32 @@
- #ifdef _WIN32
--inline void close_handle(HANDLE h) {
--  CloseHandle(h);
--}
--using ScopedHandle = ScopedResource<INVALID_HANDLE_VALUE, close_handle>;
-+struct HandleTraits {
-+  using type = HANDLE;
-+  static HANDLE invalid() {
-+    return INVALID_HANDLE_VALUE;
-+  }
-+  static void close(HANDLE h) {
-+    CloseHandle(h);
-+  }
-+};
-+using ScopedHandle = ScopedResource<HandleTraits>;
- #endif
-
--inline void close_fd(int fd) {
-+struct FdTraits {
-+  using type = int;
-+  static constexpr int invalid() {
-+    return -1;
-+  }
-+  static void close(int fd) {
- #ifdef _WIN32
--  _close(fd);
-+    _close(fd);
- #else
--  close(fd);
-+    ::close(fd);
- #endif
--}
--using ScopedFd = ScopedResource<-1, close_fd>;
-+  }
-+};
-+using ScopedFd = ScopedResource<FdTraits>;
-
- // RAII guard that invokes a callable on destruction (scope exit).
- template <typename F>
- class ScopeExit {
-     F fn_;
+ 
+ bool IsValidIdentifier(const std::string& str) {
diff --git a/cmake/vcpkg-ports/onnx/binskim.patch b/cmake/vcpkg-ports/onnx/binskim.patch
index 3f3f0aad774e1..e86a723d96a76 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 $<TARGET_OBJECTS:onnx_proto>)
@@ -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,146 +73,40 @@ 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/common/assertions.h b/onnx/common/assertions.h
---- a/onnx/common/assertions.h
-+++ b/onnx/common/assertions.h
-@@ -27,9 +27,9 @@
- #endif
- std::string barf(const char* fmt, ...);
-
--[[noreturn]] void throw_assert_error(std::string& /*msg*/);
-+[[noreturn]] void throw_assert_error(const std::string& /*msg*/);
-
--[[noreturn]] void throw_tensor_error(std::string& /*msg*/);
-+[[noreturn]] void throw_tensor_error(const std::string& /*msg*/);
-
- } // namespace ONNX_NAMESPACE
-
-diff --git a/onnx/common/assertions.cc b/onnx/common/assertions.cc
---- a/onnx/common/assertions.cc
-+++ b/onnx/common/assertions.cc
-@@ -42,11 +42,11 @@
- }
-
--void throw_assert_error(std::string& msg) {
-+void throw_assert_error(const std::string& msg) {
-     ONNX_THROW_EX(assert_error(msg));
- }
-
--void throw_tensor_error(std::string& msg) {
-+void throw_tensor_error(const std::string& msg) {
-     ONNX_THROW_EX(tensor_error(msg));
+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 <cmath>
+ #include <cstdint>
+ #include <iomanip>
++#include <limits>
++#include <locale>
++#include <sstream>
+ #include <string>
+ #include <type_traits>
+ #include <vector>
+@@ -31,6 +34,14 @@ std::vector<Widened> ParseRawIntData(const TensorProto& tensor) {
+ // point, the shortest form that round-trips through the parser.
+ template <typename T>
+ 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<T>::max_digits10) << value;
++  os << value_stream.str();
++#else
+   if constexpr (std::is_floating_point_v<T>) {
+     // 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<char, 32> buf{};
+   const auto res = std::to_chars(buf.data(), buf.data() + buf.size(), value);
+   os.write(buf.data(), res.ptr - buf.data());
++#endif
  }
-
- } // namespace ONNX_NAMESPACE
-diff --git a/onnx/common/tensor.h b/onnx/common/tensor.h
---- a/onnx/common/tensor.h
-+++ b/onnx/common/tensor.h
-@@ -51,13 +51,13 @@
-     /// size_from_dim() cannot handle this case, while elem_num() handles it correctly
-     int64_t elem_num() const {
--    return safe_dim_product(sizes_, [](const char* msg) { throw tensor_error(msg); });
-+    return safe_dim_product(sizes_, [](const char* msg) { throw_tensor_error(msg); });
-     }
-     int64_t size_from_dim(int dim) const {
-       if (dim < 0) {
-           dim += static_cast<int>(sizes_.size());
-       }
-       ONNX_ASSERT(dim >= 0 && (size_t)dim < sizes_.size())
--    return safe_dim_product(sizes_.begin() + dim, sizes_.end(), [](const char* msg) { throw tensor_error(msg); });
-+    return safe_dim_product(sizes_.begin() + dim, sizes_.end(), [](const char* msg) { throw_tensor_error(msg); });
-     }
-
-     int32_t elem_type() const {
-diff --git a/onnx/common/scoped_resource.h b/onnx/common/scoped_resource.h
---- a/onnx/common/scoped_resource.h
-+++ b/onnx/common/scoped_resource.h
-@@ -25,18 +25,18 @@
- // Usage:
--//   using ScopedFd = ScopedResource<-1, close_fd>;
--//   ScopedFd guard(fd);          // destructor calls close_fd(fd)
-+//   using ScopedFd = ScopedResource<FdTraits>;
-+//   ScopedFd guard(fd);          // destructor calls FdTraits::close(fd)
- //   int raw = guard.release();   // relinquish ownership
--template <auto Invalid, void (*Close)(decltype(Invalid))>
-+template <typename Traits>
- class ScopedResource {
--  using T = decltype(Invalid);
-+  using T = typename Traits::type;
-     T val_;
-
-  public:
-     explicit ScopedResource(T v) : val_(v) {}
-     ~ScopedResource() {
--    if (val_ != Invalid) {
--      Close(val_);
-+    if (val_ != Traits::invalid()) {
-+      Traits::close(val_);
-       }
-     }
-     T get() const {
-       return val_;
-@@ -44,7 +44,7 @@
-     }
-     T release() {
-       T tmp = val_;
--    val_ = Invalid;
-+    val_ = Traits::invalid();
-       return tmp;
-     }
-     ScopedResource(const ScopedResource&) = delete;
-@@ -55,20 +55,32 @@
- #ifdef _WIN32
--inline void close_handle(HANDLE h) {
--  CloseHandle(h);
--}
--using ScopedHandle = ScopedResource<INVALID_HANDLE_VALUE, close_handle>;
-+struct HandleTraits {
-+  using type = HANDLE;
-+  static HANDLE invalid() {
-+    return INVALID_HANDLE_VALUE;
-+  }
-+  static void close(HANDLE h) {
-+    CloseHandle(h);
-+  }
-+};
-+using ScopedHandle = ScopedResource<HandleTraits>;
- #endif
-
--inline void close_fd(int fd) {
-+struct FdTraits {
-+  using type = int;
-+  static constexpr int invalid() {
-+    return -1;
-+  }
-+  static void close(int fd) {
- #ifdef _WIN32
--  _close(fd);
-+    _close(fd);
- #else
--  close(fd);
-+    ::close(fd);
- #endif
--}
--using ScopedFd = ScopedResource<-1, close_fd>;
-+  }
-+};
-+using ScopedFd = ScopedResource<FdTraits>;
-
- // RAII guard that invokes a callable on destruction (scope exit).
- template <typename F>
- class ScopeExit {
-     F fn_;
+ 
+ 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..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.22.0"
-    SHA512 13fafff073a8e0bcf67fd06195da979c3c6273b3f5fdd25f5e5c39ad6af20eefe107f13dc73d52b6021e986d87401e46121e00bba8fbd09f098334e34f78462f
+    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 c399d4c1ab087..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.22.0",
+  "version-semver": "1.23.2",
   "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<string, string> 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 1e9c7dbb2a45c..6d5f8fc2423de 100644
--- a/docs/OperatorKernels.md
+++ b/docs/OperatorKernels.md
@@ -63,13 +63,15 @@ 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**<br> *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)<br/> **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**<br> *in* Y:**T**<br> *out* Z:**T**|11+|**T** = tensor(uint32), tensor(uint64), tensor(uint8)|
+|BitShift|*in* X:**T**<br> *in* Y:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *out* output:**T2**|17+|**T1** = tensor(int32), tensor(int64)<br/> **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**<br> *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)<br/> **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**<br> *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)<br/> **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)<br/> **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)<br/> **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)<br/> **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)<br/> **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)|
@@ -287,7 +289,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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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)|
@@ -477,7 +480,7 @@ The **OpSet Version** column uses the following notation:
 |||[1, 21]|**T** = tensor(float)|
 |Softsign|*in* input:**T**<br> *out* output:**T**|22+|**T** = tensor(float)|
 |||[1, 21]|**T** = tensor(float)|
-|SpaceToDepth|*in* input:**T**<br> *out* output:**T**|13+|**T** = tensor(double), tensor(float), tensor(int8), tensor(uint8)|
+|SpaceToDepth|*in* input:**T**<br> *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**<br> *in* split:**T**<br> *out* outputs...:**T**<br><br>or<br><br>*in* input:**T**<br> *in* split:**tensor(int64)**<br> *out* outputs:**T**<br><br>or<br><br>*in* input:**T**<br> *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)|
@@ -692,7 +695,8 @@ The **OpSet Version** column uses the following notation:
 |||14|**T** = tensor(double), tensor(float), tensor(float16)<br/> **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**<br> *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)<br/> **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**<br> *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)<br/> **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)<br/> **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)<br/> **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)<br/> **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)<br/> **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)|
@@ -865,7 +869,8 @@ The **OpSet Version** column uses the following notation:
 |Min|*in* data_0:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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)|
@@ -1016,7 +1021,7 @@ The **OpSet Version** column uses the following notation:
 |||[1, 21]|**T** = tensor(double), tensor(float), tensor(float16)|
 |Softsign|*in* input:**T**<br> *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**<br> *out* output:**T**|13+|**T** = tensor(double), tensor(float), tensor(float16)|
+|SpaceToDepth|*in* input:**T**<br> *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**<br> *in* split:**T**<br> *out* outputs...:**T**<br><br>or<br><br>*in* input:**T**<br> *in* split:**tensor(int64)**<br> *out* outputs:**T**<br><br>or<br><br>*in* input:**T**<br> *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)|
@@ -1197,7 +1202,7 @@ The **OpSet Version** column uses the following notation:
 |||10|**I** = tensor(int64)<br/> **T** = tensor(float), tensor(float16)|
 |||[8, 9]|**I** = tensor(int64)<br/> **T** = tensor(float), tensor(float16)|
 |||[1, 7]|**T** = tensor(float), tensor(float16)|
-|SpaceToDepth|*in* input:**T**<br> *out* output:**T**|13+|**T** = tensor(double), tensor(float), tensor(float16)|
+|SpaceToDepth|*in* input:**T**<br> *out* output:**T**|[13, 27]|**T** = tensor(double), tensor(float), tensor(float16)|
 |||[1, 12]|**T** = tensor(double), tensor(float), tensor(float16)|
 | |
 | |
@@ -1239,12 +1244,13 @@ 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**<br> *in* Y:**T**<br> *out* Z:**T**|11+|**T** = tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
+|BitShift|*in* X:**T**<br> *in* Y:**T**<br> *out* Z:**T**|[11, 27]|**T** = tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)|
 |BitwiseAnd|*in* A:**T**<br> *in* B:**T**<br> *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**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *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)<br/> **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**<br> *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)<br/> **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)<br/> **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)<br/> **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)<br/> **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)<br/> **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/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 <https://github.com/onnx/onnx/blob/main/onnx/tools/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<int>::max();
diff --git a/include/onnxruntime/core/session/onnxruntime_c_api.h b/include/onnxruntime/core/session/onnxruntime_c_api.h
index 4bec5ed2669c7..572dbe2aa0a22 100644
--- a/include/onnxruntime/core/session/onnxruntime_c_api.h
+++ b/include/onnxruntime/core/session/onnxruntime_c_api.h
@@ -224,6 +224,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/core/session/onnxruntime_type_conversion.h b/include/onnxruntime/core/session/onnxruntime_type_conversion.h
index 0941b83b3edc9..d5c818a916a54 100644
--- a/include/onnxruntime/core/session/onnxruntime_type_conversion.h
+++ b/include/onnxruntime/core/session/onnxruntime_type_conversion.h
@@ -33,7 +33,7 @@ constexpr ONNXTensorElementDataType ToOrtTensorElementDataType(int tensor_proto_
       return ONNX_TENSOR_ELEMENT_DATA_TYPE_INT2;
     default:
       if (tensor_proto_element_type < 0 ||
-          tensor_proto_element_type > static_cast<int>(ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT8E8M0)) {
+          tensor_proto_element_type > static_cast<int>(ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E3M2)) {
         return ONNX_TENSOR_ELEMENT_DATA_TYPE_UNDEFINED;
       }
       return static_cast<ONNXTensorElementDataType>(tensor_proto_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/js/web/docs/webgpu-operators.md b/js/web/docs/webgpu-operators.md
index ee471fc657ce5..6307b1fa24a03 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/framework/onnxruntime_map_type_info.h b/onnxruntime/core/framework/onnxruntime_map_type_info.h
index 09674df403fe3..21e8f3e7a2503 100644
--- a/onnxruntime/core/framework/onnxruntime_map_type_info.h
+++ b/onnxruntime/core/framework/onnxruntime_map_type_info.h
@@ -32,7 +32,7 @@ constexpr ONNXTensorElementDataType ToONNXTensorElementDataType(
 
 consteval bool IsTensorProtoToOrtElementTypeMapBijective() {
   constexpr size_t ort_element_type_count =
-      static_cast<size_t>(ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT8E8M0) + 1;
+      static_cast<size_t>(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 cd00b02bfac2f..49e2d12b93d45 100644
--- a/onnxruntime/core/framework/tensorprotoutils.cc
+++ b/onnxruntime/core/framework/tensorprotoutils.cc
@@ -2237,6 +2237,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 59bd955c6519e..d1ff89c12120c 100644
--- a/onnxruntime/core/framework/tensorprotoutils.h
+++ b/onnxruntime/core/framework/tensorprotoutils.h
@@ -58,6 +58,8 @@ inline constexpr std::array<uint8_t, ONNX_NAMESPACE::TensorProto_DataType_DataTy
         sizeof(uint8_t),   // FLOAT8E8M0
         sizeof(uint8_t),   // UINT2
         sizeof(uint8_t),   // INT2
+        sizeof(uint8_t),   // FLOAT6E2M3: packed 6-bit elements
+        sizeof(uint8_t),   // FLOAT6E3M2: packed 6-bit elements
     };
 
 consteval bool IsTensorProtoDataTypeElementSizeMapComplete() {
diff --git a/onnxruntime/core/optimizer/attention_fusion_helper.h b/onnxruntime/core/optimizer/attention_fusion_helper.h
index 814c2c2ea3259..4e1850bf4a267 100644
--- a/onnxruntime/core/optimizer/attention_fusion_helper.h
+++ b/onnxruntime/core/optimizer/attention_fusion_helper.h
@@ -400,7 +400,7 @@ bool MatchUnidirMaskSubgraph(const Graph& graph, const Node& add_node, MatchUnid
   }
 
   std::vector<graph_utils::EdgeEndToMatch> 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<graph_utils::EdgeEndToMatch> 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<graph_utils::EdgeEndToMatch> 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<graph_utils::EdgeEndToMatch> 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<graph_utils::EdgeEndToMatch> 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<size_t>(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 05310af8cdcad..6310904b8464a 100644
--- a/onnxruntime/core/optimizer/pad_fusion.cc
+++ b/onnxruntime/core/optimizer/pad_fusion.cc
@@ -195,7 +195,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/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/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/cann/cann_graph.cc b/onnxruntime/core/providers/cann/cann_graph.cc
index fc2c6fcbf361b..d431c89d4d219 100644
--- a/onnxruntime/core/providers/cann/cann_graph.cc
+++ b/onnxruntime/core/providers/cann/cann_graph.cc
@@ -70,7 +70,8 @@ std::vector<NodeIndex> 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->OpType() == "SpaceToDepth") && node->SinceVersion() >= 28)) {
       unsupported_nodes.push_back(index);
       continue;
     }
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 c3e062611d5cf..701424a532d07 100644
--- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc
+++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc
@@ -599,9 +599,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);
@@ -837,7 +838,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);
@@ -871,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);
@@ -1499,7 +1501,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);
@@ -1555,6 +1557,17 @@ 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_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);
+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! ***
@@ -2321,12 +2334,14 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
                                                                             double, Gemm)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 12,
                                                                       GatherElements)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint8_t,
-                                                                  BitShift)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint32_t,
-                                                                  BitShift)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint64_t,
-                                                                  BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 27,
+                                                                            uint8_t, BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 27,
+                                                                            uint16_t, BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 27,
+                                                                            uint32_t, BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 27,
+                                                                            uint64_t, BitShift)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 12, Pad)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11,
                                                                       GatherND)>,
@@ -2770,7 +2785,8 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int16_t, Neg)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int32_t, Neg)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int64_t, Neg)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, Mod)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 27, Mod)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, Mod)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, float, Abs)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, double, Abs)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int8_t, Abs)>,
@@ -2814,7 +2830,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, Gather)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, GatherElements)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, DepthToSpace)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, SpaceToDepth)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 27, SpaceToDepth)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 15,
                                                                       ScatterElements)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 15,
@@ -3678,7 +3694,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 24, SplitToSequence)>,
 
       // opset 25
-      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, 27, Cast)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, ConstantOfShape)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 25, uint8_t,
                                                                   DequantizeLinear)>,
@@ -3758,6 +3774,25 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
 
       // opset 27
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 27, Range)>,
+
+      // opset 28
+      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int8_t,
+                                                                  BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int16_t,
+                                                                  BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int32_t,
+                                                                  BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, int64_t,
+                                                                  BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, uint8_t,
+                                                                  BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, uint16_t,
+                                                                  BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, uint32_t,
+                                                                  BitShift)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, uint64_t,
+                                                                  BitShift)>,
   };
   static_assert(registration_internal::IsKernelRegistrationTableValid(
       function_table, BuildKernelCreateInfo<void>));
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 <cmath>
+#include <cstring>
+#include <limits>
+#include <type_traits>
 
 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<T>::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 <typename T>
+T FromUnsignedBits(std::make_unsigned_t<T> bits) {
+  static_assert(std::is_integral_v<T>);
+  T value;
+  std::memcpy(&value, &bits, sizeof(value));
+  return value;
+}
+
+template <typename T>
+inline bool IsInvalidShift(T shift) {
+  using UnsignedT = std::make_unsigned_t<T>;
+  constexpr UnsignedT bit_width = std::numeric_limits<UnsignedT>::digits;
+  return (std::is_signed_v<T> && shift < 0) || static_cast<UnsignedT>(shift) >= bit_width;
+}
+
 template <typename T>
 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<T>;
+  const auto result = static_cast<UnsignedT>(
+      static_cast<UnsignedT>(value) << static_cast<UnsignedT>(shift));
+  return FromUnsignedBits<T>(result);
 }
 
 template <typename T>
 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<T>) {
+      return value < 0 ? T{-1} : T{0};
+    } else {
+      return T{0};
+    }
+  }
+
+  using UnsignedT = std::make_unsigned_t<T>;
+  constexpr UnsignedT bit_width = std::numeric_limits<UnsignedT>::digits;
+  const auto shift_amount = static_cast<UnsignedT>(shift);
+  auto result = static_cast<UnsignedT>(static_cast<UnsignedT>(value) >> shift_amount);
+  if constexpr (std::is_signed_v<T>) {
+    if (value < 0 && shift_amount != 0) {
+      result |= std::numeric_limits<UnsignedT>::max() << (bit_width - shift_amount);
+    }
+  }
+
+  return FromUnsignedBits<T>(result);
 }
 
 template <typename T>
@@ -2110,9 +2158,19 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
             BuildKernelDefConstraintsFromTypeList<EnabledModTypes>()),
     Mod);
 
-ONNX_CPU_OPERATOR_KERNEL(
+ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
     Mod,
     13,
+    27,
+    KernelDefBuilder()
+        .TypeConstraint(
+            "T",
+            BuildKernelDefConstraintsFromTypeList<EnabledModTypes>()),
+    Mod);
+
+ONNX_CPU_OPERATOR_KERNEL(
+    Mod,
+    28,
     KernelDefBuilder()
         .TypeConstraint(
             "T",
@@ -2160,11 +2218,20 @@ void BroadCastFMod(OpKernelContext* context) {
 
 template <class T>
 inline T Modulus(T x, T y) {
-  auto res = x % y;
+  T res;
+  if constexpr (std::is_floating_point_v<T>) {
+    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<T>(res);
+  return res;
 }
 
 template <class T>
@@ -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<MLFloat16>();
+        auto Y = per_iter_bh.SpanInput1<MLFloat16>();
+        auto output = per_iter_bh.OutputSpan<MLFloat16>();
+
+        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<MLFloat16>();
+        const MLFloat16 Y = per_iter_bh.ScalarInput1<MLFloat16>();
+        auto output = per_iter_bh.OutputSpan<MLFloat16>();
+
+        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<MLFloat16>();
+        auto Y = per_iter_bh.SpanInput1<MLFloat16>();
+        auto output = per_iter_bh.OutputSpan<MLFloat16>();
+
+        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 <class T, typename Enable = void>
 struct CallModImpl;
 
@@ -2276,8 +2379,11 @@ struct CallModImpl<T, typename std::enable_if<!std::is_floating_point<T>::value>
 template <class T>
 struct CallModImpl<T, typename std::enable_if<std::is_floating_point<T>::value, void>::type> {
   void operator()(bool fmod, OpKernelContext* ctx) const {
-    ORT_ENFORCE(fmod, "fmod attribute must be true for floating point types");
-    BroadCastFMod<T>(ctx);
+    if (fmod) {
+      BroadCastFMod<T>(ctx);
+    } else {
+      BroadCastMod<T>(ctx);
+    }
   }
 };
 
@@ -2285,8 +2391,11 @@ struct CallModImpl<T, typename std::enable_if<std::is_floating_point<T>::value,
 template <>
 struct CallModImpl<MLFloat16> {
   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/cpu/tensor/cast_op.cc b/onnxruntime/core/providers/cpu/tensor/cast_op.cc
index 1d5a7c63228b3..22a3e7b62c458 100644
--- a/onnxruntime/core/providers/cpu/tensor/cast_op.cc
+++ b/onnxruntime/core/providers/cpu/tensor/cast_op.cc
@@ -1212,9 +1212,20 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
         .MayInplace(0, 0),  // allocation planner will check input and output sizes match before inplacing
     Cast);
 
-ONNX_CPU_OPERATOR_KERNEL(
+ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
     Cast,
     25,
+    27,
+    KernelDefBuilder()
+        .TypeConstraint("T1", BuildKernelDefConstraintsFromTypeList<EnabledSrcTypes>())
+        .TypeConstraint("T2", BuildKernelDefConstraintsFromTypeList<EnabledDstTypes>())
+        .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<EnabledSrcTypes>())
         .TypeConstraint("T2", BuildKernelDefConstraintsFromTypeList<EnabledDstTypes>())
diff --git a/onnxruntime/core/providers/cpu/tensor/space_depth_ops.cc b/onnxruntime/core/providers/cpu/tensor/space_depth_ops.cc
index f8b78e80e7cac..f59cd88508b71 100644
--- a/onnxruntime/core/providers/cpu/tensor/space_depth_ops.cc
+++ b/onnxruntime/core/providers/cpu/tensor/space_depth_ops.cc
@@ -18,9 +18,10 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
                               DataTypeImpl::GetTensorType<int8_t>()}),
     SpaceToDepth);
 
-ONNX_CPU_OPERATOR_KERNEL(
+ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
     SpaceToDepth,
     13,
+    27,
     KernelDefBuilder()
         .TypeConstraint("T", {DataTypeImpl::GetTensorType<float>(),
                               DataTypeImpl::GetTensorType<double>(),
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<float>(a), static_cast<float>(b)));
+}
+
+template <>
+__device__ __inline__ BFloat16 _Mod(BFloat16 a, BFloat16 b) {
+  return BFloat16(_Mod(static_cast<float>(a), static_cast<float>(b)));
+}
+
 template <typename T>
 __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 5fa508ee6448d..a2ba138e6b8f2 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);
@@ -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);
@@ -1842,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);
@@ -1909,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) {
@@ -2620,7 +2643,7 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 17, double, Pad)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 17, MLFloat16, Pad)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 17, bool, Pad)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, SpaceToDepth)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 27, SpaceToDepth)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, DepthToSpace)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, int8_t, Sign)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, int16_t, Sign)>,
@@ -2647,7 +2670,8 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Tanh)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Gemm)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, ReduceSum)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, Mod)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 27, Mod)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Mod)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, QuantizeLinear)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, uint8_t, QuantizeLinear)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, DequantizeLinear)>,
@@ -3136,25 +3160,25 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 24, 24, Unsqueeze)>,
 
       // Opset 25
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, float, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, double, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, MLFloat16, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, int8_t, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, int16_t, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, int32_t, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, int64_t, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint8_t, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint16_t, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint32_t, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint64_t, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, bool, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, BFloat16, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, float, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, double, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, MLFloat16, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, int8_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, int16_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, int32_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, int64_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, uint8_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, uint16_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, uint32_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, uint64_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, bool, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, BFloat16, Cast)>,
 #if !defined(DISABLE_FLOAT8_TYPES)
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, Float8E4M3FN, Cast)>,
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, Float8E5M2, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, Float8E4M3FN, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, Float8E5M2, Cast)>,
 #endif
 #if !defined(DISABLE_FLOAT4_TYPES)
-      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, Float4E2M1x2, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, 27, Float4E2M1x2, Cast)>,
 #endif
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, ConstantOfShape)>,
       BuildKernelCreateInfo<ONNX_OPERATOR_TWO_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, uint8_t, float, DequantizeLinear)>,
@@ -3202,6 +3226,27 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, Unsqueeze)>,
       // Opset 27
       BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 27, Range)>,
+      // Opset 28
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, float, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, double, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, MLFloat16, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, BFloat16, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int8_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int16_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int32_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, int64_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint8_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint16_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint32_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, uint64_t, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, bool, Cast)>,
+#if !defined(DISABLE_FLOAT8_TYPES)
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Float8E4M3FN, Cast)>,
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Float8E5M2, Cast)>,
+#endif
+#if !defined(DISABLE_FLOAT4_TYPES)
+      BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Float4E2M1x2, Cast)>,
+#endif
 #endif
   };
 
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<CUDA_NHWC_OP_VERSIONED_CLASS_NAME(11, 12, DepthToSpace)>,
       BuildKernelCreateInfo<CUDA_NHWC_OP_CLASS_NAME(13, DepthToSpace)>,
       BuildKernelCreateInfo<CUDA_NHWC_OP_VERSIONED_CLASS_NAME(1, 12, SpaceToDepth)>,
-      BuildKernelCreateInfo<CUDA_NHWC_OP_CLASS_NAME(13, SpaceToDepth)>,
+      BuildKernelCreateInfo<CUDA_NHWC_OP_VERSIONED_CLASS_NAME(13, 27, SpaceToDepth)>,
       BuildKernelCreateInfo<CUDA_NHWC_OP_VERSIONED_TYPED_CLASS_NAME(1, 12, float, LRN)>,
       BuildKernelCreateInfo<CUDA_NHWC_OP_VERSIONED_TYPED_CLASS_NAME(1, 12, double, LRN)>,
       BuildKernelCreateInfo<CUDA_NHWC_OP_VERSIONED_TYPED_CLASS_NAME(1, 12, MLFloat16, LRN)>,
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<int32_t, int64_t, uint32_t, uint64_t, float, double, MLFloat16>()),
     Mod);
 
-ONNX_OPERATOR_KERNEL_EX(Mod, kOnnxDomain, 13, kCudaExecutionProvider,
+ONNX_OPERATOR_VERSIONED_KERNEL_EX(
+    Mod, kOnnxDomain, 13, 27, kCudaExecutionProvider,
+    (*KernelDefBuilder::Create())
+        .TypeConstraint("T", BuildKernelDefConstraints<int32_t, int64_t, uint32_t, uint64_t, float, double, MLFloat16,
+                                                       BFloat16>()),
+    Mod);
+
+ONNX_OPERATOR_KERNEL_EX(Mod, kOnnxDomain, 28, kCudaExecutionProvider,
                         (*KernelDefBuilder::Create())
                             .TypeConstraint("T", BuildKernelDefConstraints<int32_t, int64_t, uint32_t, uint64_t, float,
                                                                            double, MLFloat16, BFloat16>()),
@@ -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<typename ToCudaType<data_type>::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/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<MLDataType>& CastOpTypeConstraints() {
           .TypeConstraint("T1", DataTypeImpl::GetTensorType<T>()) \
           .TypeConstraint("T2", CastOpTypeConstraints()),         \
       Cast<T>);                                                   \
+  ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX(                        \
+      Cast,                                                       \
+      kOnnxDomain,                                                \
+      25, 27,                                                     \
+      T,                                                          \
+      kCudaExecutionProvider,                                     \
+      (*KernelDefBuilder::Create())                               \
+          .TypeConstraint("T1", DataTypeImpl::GetTensorType<T>()) \
+          .TypeConstraint("T2", CastOpTypeConstraints()),         \
+      Cast<T>);                                                   \
   ONNX_OPERATOR_TYPED_KERNEL_EX(                                  \
       Cast,                                                       \
       kOnnxDomain,                                                \
-      25,                                                         \
+      28,                                                         \
       T,                                                          \
       kCudaExecutionProvider,                                     \
       (*KernelDefBuilder::Create())                               \
@@ -411,11 +421,21 @@ SPECIALIZE_IMPL(BFloat16)
           .TypeConstraint("T2", OutputTypeConstraints),           \
       Cast<T>);
 
-#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<T>()) \
+          .TypeConstraint("T2", OutputTypeConstraints),           \
+      Cast<T>);                                                   \
   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<T>::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<float>()})
-REGISTER_KERNEL_TYPED_25(Float4E2M1x2, {DataTypeImpl::GetTensorType<float>()})
+REGISTER_KERNEL_TYPED_25_TO_28(Float4E2M1x2, {DataTypeImpl::GetTensorType<float>()})
 template Status Cast<Float4E2M1x2>::ComputeInternal(OpKernelContext* context) const;
 #endif
 
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<LAYOUT_NHWC>);
 #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<LAYOUT_NCHW>);
 
 #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/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/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<std::string>(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 26f998c7521a2..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)},
@@ -1092,6 +1093,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/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/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<KernelRegistry> 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/migraphx/migraphx_execution_provider.cc b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc
index 8908013a52291..72c5676c20fa0 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<int> toVector(const ONNX_NAMESPACE::int64s& nums) {
 static bool IsUnsupportedOpMode(const onnxruntime::GraphViewer& graph_viewer, const Node* node) {
   std::vector<NodeIndex> input_nodes;
   const auto& optype = node->OpType();
+  // These opset-28 semantics require CPU kernels or ONNX function expansion.
+  if ((optype == "Mod" || optype == "SpaceToDepth") && 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/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/openvino/ov_versions/data_ops.cc b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc
index 7aae81b987c0d..24c5388fe25c7 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();
+  // These opset-28 semantics require CPU kernels or ONNX function expansion.
+  if ((optype == "Mod" || optype == "BitShift" || optype == "SpaceToDepth") && node->SinceVersion() >= 28) {
+    return false;
+  }
 
 #ifndef NDEBUG
   if (openvino_ep::backend_utils::IsDebugEnabled()) {
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<int64_t>(0));  // 0=integer mod. 1=float mod.
+  int64_t fmod = node_helper.Get("fmod", static_cast<int64_t>(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/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/webgpu/webgpu_execution_provider.cc b/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc
index ec2ed34439ab6..dba4f109db6d7 100644
--- a/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc
+++ b/onnxruntime/core/providers/webgpu/webgpu_execution_provider.cc
@@ -468,7 +468,8 @@ std::unique_ptr<KernelRegistry> 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/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 9227930f312dd..e9015b3e40b09 100644
--- a/onnxruntime/core/session/inference_session.cc
+++ b/onnxruntime/core/session/inference_session.cc
@@ -1249,7 +1249,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.
@@ -1345,6 +1347,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()) {
@@ -1353,7 +1357,8 @@ 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, check_load_cancellation_fn_});
+                                    ModelOptions{allow_released_opsets_only, strict_shape_type_inference,
+                                                 check_load_cancellation_fn_});
   };
 
   return LoadWithLoader(loader, "model_loading_array");
@@ -1409,6 +1414,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()) {
@@ -1418,7 +1425,8 @@ 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, check_load_cancellation_fn_});
+                                    ModelOptions{allow_released_opsets_only, strict_shape_type_inference,
+                                                 check_load_cancellation_fn_});
   };
 
   return LoadWithLoader(loader, "model_loading_proto");
@@ -1516,12 +1524,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<Model> 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/inference_session_test.cc b/onnxruntime/test/framework/inference_session_test.cc
index f7fab9b320c4f..dd4a9c8965504 100644
--- a/onnxruntime/test/framework/inference_session_test.cc
+++ b/onnxruntime/test/framework/inference_session_test.cc
@@ -5245,7 +5245,7 @@ TEST(InferenceSessionTests, CompileApiExternalBufferKeepsSubgraphInitializersEmb
       if (buffer_destination && has_main_initializer) continue;
 #endif
 #if defined(__EMSCRIPTEN__)
-      // The WASM test module does not provide Module.MountedFiles for file-backed external data.
+      // File-backed subgraph tensors require a Module.MountedFiles mapping in WASM.
       if (!buffer_destination) continue;
 #endif
       {
@@ -5295,7 +5295,7 @@ TEST(InferenceSessionTests, CompileApiExternalBufferKeepsSubgraphInitializersEmb
   }
 }
 
-// The WASM test module cannot load host filesystem external data without Module.MountedFiles.
+// 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");
diff --git a/onnxruntime/test/framework/tensorutils_test.cc b/onnxruntime/test/framework/tensorutils_test.cc
index 6f941cd95d741..779c6cfff5366 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 9368f5da889b0..7631b5eaaf58b 100644
--- a/onnxruntime/test/framework/type_info_test.cc
+++ b/onnxruntime/test/framework/type_info_test.cc
@@ -11,6 +11,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"
 #include "core/session/onnxruntime_type_conversion.h"
 
@@ -49,6 +50,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) {
@@ -72,6 +75,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, TensorElementTypeConversions) {
   struct TypePair {
     ONNXTensorElementDataType api_type;
@@ -105,6 +115,8 @@ TEST(TypeInfoTests, TensorElementTypeConversions) {
       {ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT8E8M0, ONNX_NAMESPACE::TensorProto_DataType_FLOAT8E8M0},
       {ONNX_TENSOR_ELEMENT_DATA_TYPE_UINT2, ONNX_NAMESPACE::TensorProto_DataType_UINT2},
       {ONNX_TENSOR_ELEMENT_DATA_TYPE_INT2, ONNX_NAMESPACE::TensorProto_DataType_INT2},
+      {ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E2M3, ONNX_NAMESPACE::TensorProto_DataType_FLOAT6E2M3},
+      {ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT6E3M2, ONNX_NAMESPACE::TensorProto_DataType_FLOAT6E3M2},
   };
   static_assert(std::size(types) == ONNX_NAMESPACE::TensorProto_DataType_DataType_ARRAYSIZE);
   for (const auto& [api_type, proto_type] : types) {
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<std::set<BrokenTest>> 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<std::set<BrokenTest>> 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 2af930ff75bcf..4ec48f558741b 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,41 +622,8 @@ ORT_API(void, OrtUninitializeBuffer, _In_opt_ void* input, size_t input_len, enu
                                          (Y*)preallocated, static_cast<size_t>(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(FLOAT8E8M0)
-    CASE_TYPE(FLOAT4E2M1)
-    CASE_TYPE(UINT4)
-    CASE_TYPE(INT4)
-    CASE_TYPE(UINT2)
-    CASE_TYPE(INT2)
-    default:
-      return ONNX_TENSOR_ELEMENT_DATA_TYPE_UNDEFINED;
-  }
+  return utils::ToOrtTensorElementDataType(type);
 }
 
 ONNXTensorElementDataType GetTensorElementType(const onnx::TensorProto& tensor_proto) {
diff --git a/onnxruntime/test/optimizer/graph_transform_test.cc b/onnxruntime/test/optimizer/graph_transform_test.cc
index 933167791554b..3333fc2b8121e 100644
--- a/onnxruntime/test/optimizer/graph_transform_test.cc
+++ b/onnxruntime/test/optimizer/graph_transform_test.cc
@@ -7802,6 +7802,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);
@@ -7814,14 +7820,17 @@ TEST_F(GraphTransformationTests, SliceConcatToSpaceToDepthFusionTest) {
     }
   };
 
-  TransformerTester(build_test_case,
-                    check_transformed_graph,
-                    TransformerLevel::Default,
-                    TransformerLevel::Level1,
-                    13,
-                    0.0,
-                    0.0,
-                    std::make_unique<SliceConcatToSpaceToDepthFusion>());
+  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<SliceConcatToSpaceToDepthFusion>());
+  }
 }
 
 TEST_F(GraphTransformationTests, SliceConcatToSpaceToDepthFusionWithConstantNodesTest) {
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 de9090d905ca8..04e7ed5cc11de 100644
--- a/onnxruntime/test/optimizer/qdq_transformer_test.cc
+++ b/onnxruntime/test/optimizer/qdq_transformer_test.cc
@@ -5690,7 +5690,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<std::string, int> 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<float>("scan_output_2", {1, 2, 1}, {0.f, 0.f});
   test.AddOutput<float>("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..c15df7ddc90cb 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 <algorithm>
 #include <limits>
 #include <math.h>
+#include <type_traits>
 
 namespace onnxruntime {
 namespace test {
@@ -4351,6 +4352,92 @@ TEST(ModOpTest, Fmod_float16_mixed_sign) {
   test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kQnnExecutionProvider});
 }
 
+template <typename T>
+void TestFloorModFloatingPointOpset28(const std::vector<T>& x,
+                                      const std::vector<T>& y,
+                                      const std::vector<T>& expected) {
+  OpTester test("Mod", 28);
+  test.AddAttribute<int64_t>("fmod", 0);
+  const std::vector<int64_t> shape{static_cast<int64_t>(x.size())};
+  test.AddInput<T>("X", shape, x);
+  test.AddInput<T>("Y", shape, y);
+  test.AddOutput<T>("Z", shape, expected);
+  test.Run(OpTester::ExpectResult::kExpectSuccess, "",
+           {kQnnExecutionProvider, kOpenVINOExecutionProvider});
+}
+
+TEST(ModOpTest, FloorMod_float_mixed_sign_opset28) {
+  TestFloorModFloatingPointOpset28<float>(
+      {-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<double>(
+      {-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<MLFloat16>(
+      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<float>::infinity();
+  constexpr auto nan = std::numeric_limits<float>::quiet_NaN();
+  TestFloorModFloatingPointOpset28<float>(
+      {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});
+}
+
+TEST(ModOpTest, FloorMod_opset28_provider_fallback) {
+  InlinedVector<std::unique_ptr<IExecutionProvider>> providers;
+  providers.push_back(DefaultCpuExecutionProvider());
+#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());
+  for (auto& provider : providers) {
+    if (!provider) {
+      continue;
+    }
+    SCOPED_TRACE(provider->Type());
+    const std::vector<int64_t> shape{4};
+    const float infinity = std::numeric_limits<float>::infinity();
+    const std::vector<float> x{-5.0f, 5.0f, -5.0f, 5.0f};
+    const std::vector<float> y{3.0f, -3.0f, infinity, -infinity};
+    OpTester test("Mod", 28);
+    test.AddAttribute<int64_t>("fmod", 0);
+    test.AddInput<float>("X", shape, x);
+    test.AddInput<float>("Y", shape, y);
+    test.AddOutput<float>("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());
+    model.ToProto().SerializeToString(&model_data);
+    EPVerificationParams params;
+    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<float>(shape, x)}, {"Y", CreateInputOrtValueOnCPU<float>(shape, y)}},
+        params);
+  }
+}
+
 #if defined(USE_CUDA)
 TEST(ModOpTest, Fmod_bfloat16_mixed_sign) {
   OpTester test("Mod", 13);
@@ -4652,6 +4739,125 @@ TEST(BitShiftOpTest, BroadcastXRight_Uint8) {
   test.Run();
 }
 
+template <typename T>
+void TestSignedBitShiftOpset28() {
+  constexpr T bit_width = static_cast<T>(std::numeric_limits<std::make_unsigned_t<T>>::digits);
+
+  {
+    OpTester test("BitShift", 28);
+    test.AddAttribute("direction", "LEFT");
+    test.AddInput<T>("X", {6}, {T{-8}, T{-3}, T{1}, std::numeric_limits<T>::max(), T{-2}, T{2}});
+    test.AddInput<T>("Y", {6}, {T{1}, T{2}, T{3}, T{1}, T{-1}, bit_width});
+    test.AddOutput<T>("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<T>("X", {6}, {T{-8}, T{-3}, T{1}, T{-2}, T{-2}, T{2}});
+    test.AddInput<T>("Y", {6}, {T{1}, T{2}, T{3}, T{-1}, bit_width, bit_width});
+    test.AddOutput<T>("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<T>("X", {1}, {T{-8}});
+    test.AddInput<T>("Y", {4}, {T{1}, T{2}, bit_width, T{-1}});
+    test.AddOutput<T>("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<T>("X", {2}, {std::numeric_limits<T>::max(), T{-8}});
+    test.AddInput<T>("Y", {1}, {T{1}});
+    test.AddOutput<T>("Z", {2}, {T{-2}, T{-16}});
+    test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kOpenVINOExecutionProvider});
+  }
+}
+
+TEST(BitShiftOpTest, SignedInt8Opset28) {
+  TestSignedBitShiftOpset28<int8_t>();
+}
+
+TEST(BitShiftOpTest, SignedInt16Opset28) {
+  TestSignedBitShiftOpset28<int16_t>();
+}
+
+TEST(BitShiftOpTest, SignedInt32Opset28) {
+  TestSignedBitShiftOpset28<int32_t>();
+}
+
+TEST(BitShiftOpTest, SignedInt64Opset28) {
+  TestSignedBitShiftOpset28<int64_t>();
+}
+
+template <typename T>
+void TestBitShiftOpset28ProviderFallback(const std::vector<T>& x, const std::vector<T>& y,
+                                         const std::vector<T>& expected) {
+  const std::vector<int64_t> shape{static_cast<int64_t>(x.size())};
+  OpTester test("BitShift", 28);
+  test.AddAttribute("direction", "RIGHT");
+  test.AddInput<T>("X", shape, x);
+  test.AddInput<T>("Y", shape, y);
+  test.AddOutput<T>("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<T>(shape, x)}, {"Y", CreateInputOrtValueOnCPU<T>(shape, y)}},
+      params);
+}
+
+TEST(BitShiftOpTest, Opset28ProviderFallback) {
+  TestBitShiftOpset28ProviderFallback<int32_t>({-8, -3, -2, -2}, {1, 2, -1, 32}, {-4, -1, -1, -1});
+  TestBitShiftOpset28ProviderFallback<uint32_t>({8, 255, 1, 42}, {1, 32, 33, 64}, {4, 0, 0, 0});
+}
+
+template <typename T>
+void TestUnsignedBitShiftRegistration(int opset) {
+  OpTester test("BitShift", opset);
+  test.AddAttribute("direction", "LEFT");
+  test.AddInput<T>("X", {2}, {T{1}, T{3}});
+  test.AddInput<T>("Y", {2}, {T{1}, T{2}});
+  test.AddOutput<T>("Z", {2}, {T{2}, T{12}});
+  test.Run();
+}
+
+TEST(BitShiftOpTest, Uint16Opset11) {
+  TestUnsignedBitShiftRegistration<uint16_t>(11);
+}
+
+TEST(BitShiftOpTest, UnsignedTypesOpset28) {
+  TestUnsignedBitShiftRegistration<uint8_t>(28);
+  TestUnsignedBitShiftRegistration<uint16_t>(28);
+  TestUnsignedBitShiftRegistration<uint32_t>(28);
+  TestUnsignedBitShiftRegistration<uint64_t>(28);
+}
+
+TEST(BitShiftOpTest, RightShiftByBitWidth_Uint64Opset28) {
+  OpTester test("BitShift", 28);
+  test.AddAttribute("direction", "RIGHT");
+  test.AddInput<uint64_t>("X", {4}, {1000, 255, 1, 42});
+  test.AddInput<uint64_t>("Y", {4}, {64, 64, 64, 64});
+  test.AddOutput<uint64_t>("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/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<DstType> CastedValues(gsl::span<const SrcType> src) {
 }
 
 struct CastNonStringTester {
+  int opset{21};
+  bool cuda_only{false};
+
   template <typename SrcType, typename DstType>
   void operator()(const std::pair<SrcType, DstType>&) {
     SCOPED_TRACE(
@@ -177,7 +180,8 @@ struct CastNonStringTester {
     auto output_span = gsl::make_span<DstType>(output_buffer.get(), size);
     CastSpan<SrcType, DstType>(input_span, output_span);
 
-    TestCastOp<SrcType, DstType>(input_span, output_span, shape.AsShapeVector());
+    TestCastOp<SrcType, DstType>(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<boost::mp11::mp_product<std::pair, CastNonStringTypes, CastNonStringTypes>>(
+        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<boost::mp11::mp_product<std::pair, CastNonStringTypes, CastNonStringTypes>>(
+      CastNonStringTester{28, true});
+}
+#endif
+
 TEST(CastOpTest, FromString) {
   const std::vector<int64_t> shape{2, 2, 2};
   const std::vector<std::string> string_data = {"-inf", "+INF", "0.9767611", "0.28280696",
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 <unordered_map>
+#include <vector>
+
 #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<onnxruntime::Model>("compress_non_canonical_bool", false,
-                                                    DefaultLoggingManager().DefaultLogger());
+  const std::unordered_map<std::string, int> domain_to_version{{kOnnxDomain, 27}};
+  auto model = std::make_unique<onnxruntime::Model>(
+      "compress_non_canonical_bool", false, ModelMetaData(), PathString(), IOnnxRuntimeOpSchemaRegistryList(),
+      domain_to_version, std::vector<ONNX_NAMESPACE::FunctionProto>{}, 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/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<double>("output", {N, C * blocksize * blocksize, H / blocksize, W / blocksize}, result);
   test.Run();
 }
+
+TEST(TensorOpTest, SpaceToDepthOpset28ModesAndProviderRouting) {
+  const std::vector<float> input{0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15};
+  const std::vector<float> dcr{0, 2, 8, 10, 1, 3, 9, 11, 4, 6, 12, 14, 5, 7, 13, 15};
+  const std::vector<float> 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<float>({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<std::unique_ptr<IExecutionProvider>> 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<int>(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<OrtValue> fetches;
+        ASSERT_STATUS_OK(session.Run(RunOptions{}, builder.feeds_, builder.output_names_, &fetches));
+        ASSERT_EQ(fetches.size(), 1u);
+        const auto& tensor = fetches[0].Get<Tensor>();
+        EXPECT_EQ(tensor.Shape(), TensorShape({1, 8, 1, 2}));
+        const auto& expected = mode == "CRD" ? crd : dcr;
+        const auto* actual = tensor.Data<float>();
+        EXPECT_EQ(std::vector<float>(actual, actual + tensor.Shape().Size()), expected);
+      }
+    }
+  }
+}
+
 TEST(TensorOpTest, SpaceToDepthTest_2) {
   OpTester test("SpaceToDepth");
   constexpr int64_t blocksize = 3;
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/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 <cassert>
+#include <limits>
 #include <string>
 
 #include "test/providers/qnn/qnn_test_utils.h"
@@ -38,6 +39,25 @@ static void RunModTest(const std::vector<TestInputDef<DataType>>& input_defs,
 // CPU tests:
 //
 
+TEST_F(QnnCPUBackendTests, Mod_float_Floor_Opset28_Fallback) {
+  const float inf = std::numeric_limits<float>::infinity();
+  RunModTest<float>({TestInputDef<float>({8}, false, {-5.0f, 5.0f, -5.0f, 5.0f, -5.0f, 5.0f, -5.0f, 5.0f}),
+                     TestInputDef<float>({8}, false, {3.0f, 3.0f, -3.0f, -3.0f, inf, inf, -inf, -inf})},
+                    {utils::MakeAttribute("fmod", static_cast<int64_t>(0))},
+                    ExpectedEPNodeAssignment::None,
+                    "cpu",
+                    28);
+}
+
+TEST_F(QnnCPUBackendTests, Mod_integer_Floor_Opset28) {
+  RunModTest<int32_t>({TestInputDef<int32_t>({4}, false, {-5, 5, -5, 5}),
+                       TestInputDef<int32_t>({4}, false, {3, 3, -3, -3})},
+                      {utils::MakeAttribute("fmod", static_cast<int64_t>(0))},
+                      ExpectedEPNodeAssignment::All,
+                      "cpu",
+                      28);
+}
+
 // Test that Mod with dynamic divisor.
 TEST_F(QnnCPUBackendTests, Mod_dynamic_Divisor) {
   RandomValueGenerator rand_gen{optional<RandomValueGenerator::RandomSeedType>{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<int64_t>(2));
+    node.AddAttribute("mode", mode);
+    ASSERT_STATUS_OK(graph.Resolve());
+    ctx.graph_viewer = std::make_unique<GraphViewer>(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) {
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..2cfa152b94a17 100644
--- a/onnxruntime/test/python/onnxruntime_test_python_backend.py
+++ b/onnxruntime/test/python/onnxruntime_test_python_backend.py
@@ -5,13 +5,18 @@
 import os
 import tempfile
 import unittest
+from types import SimpleNamespace
+from unittest.mock import patch
 
 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 +28,84 @@ 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
+        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))
+
+    @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_opset29_model()
+        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
+
+    @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:
+            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))
+
+    @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, "opset29.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_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(
+            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 +149,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..fc11cff5648c0 100644
--- a/onnxruntime/test/python/requirements.txt
+++ b/onnxruntime/test/python/requirements.txt
@@ -1,3 +1,3 @@
-onnx==1.22.0
+onnx==1.23.2
 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..325fa58772f76 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,179 @@
         "^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$",
+        "^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_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$",
+        "^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<void(ModelTestBuilder& helper)>& 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<void(ModelTestBuilder& helper)>& 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<std::string, i
   // Create a simple model
   std::unordered_map<std::string, int> 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..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.22.0
+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/build.py b/tools/ci_build/build.py
index 1bcdae465d400..c8e031a640275 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": "1",
+        "ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS": "0",
+    }
+
+
 def update_submodules(source_dir):
     run_subprocess(["git", "submodule", "sync", "--recursive"], cwd=source_dir)
     run_subprocess(["git", "submodule", "update", "--init", "--recursive"], cwd=source_dir)
@@ -2071,7 +2078,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..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.22.0
+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 1f99fc4c9a25b..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.22.0
+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 5b0a682d9d419..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.22.0
+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 462de90e67c3f..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.22.0
+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 4730f767c9b50..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.22.0
+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 ef80b8af11b05..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.22.0
+onnx==1.23.2
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_build_args.py b/tools/ci_build/test_build_args.py
index ffcd3680b65dc..a44de459a569a 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")
+
+_RELEASED_OPSET_ENVIRONMENT = {
+    "ALLOW_RELEASED_ONNX_OPSET_ONLY": "1",
+    "ORT_BACKEND_TEST_ALLOW_UNRELEASED_OPSETS": "0",
+}
 
 
 class BuildArgsTest(unittest.TestCase):
@@ -114,5 +121,26 @@ def test_acl_deprecation_warning_not_emitted_without_use_acl(self):
         warn.assert_not_called()
 
 
+class OnnxBackendTestEnvironmentTest(unittest.TestCase):
+    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),
+                        _RELEASED_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__":
     unittest.main()
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()
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 "