intel · ai-fw-intg · Aug 19, 2026 · Aug 19, 2026 · Aug 19, 2026 · Aug 19, 2026
diff --git a/cmake/deps.txt b/cmake/deps.txt
@@ -54,10 +54,10 @@ pytorch_cpuinfo;https://github.com/pytorch/cpuinfo/archive/4628dc060ce4e82345dc1
 re2;https://github.com/google/re2/archive/refs/tags/2024-07-02.zip;646e1728269cde7fcef990bf4a8e87b047882e88
 safeint;https://github.com/dcleblanc/SafeInt/archive/refs/tags/3.0.28.zip;23f252040ff6cb9f1fd18575b32fa8fb5928daac
 tensorboard;https://github.com/tensorflow/tensorboard/archive/373eb09e4c5d2b3cc2493f0949dc4be6b6a45e81.zip;67b833913605a4f3f499894ab11528a702c2b381
-cutlass;https://github.com/NVIDIA/cutlass/archive/refs/tags/v4.4.2.zip;4b0bae4428b84370407c0a71778b13dc2eee5be1
+cutlass;https://github.com/NVIDIA/cutlass/archive/refs/tags/v4.7.0.zip;51d4f1ba4b9977b94606e23bbc307463a6741383
 extensions;https://github.com/microsoft/onnxruntime-extensions/archive/c24b7bab0c12f53da76d0c31b03b9f0f8ec8f3b4.zip;239063aee4946a9af147b473a4c3da78ba7413b4
 directx_headers;https://github.com/microsoft/DirectX-Headers/archive/refs/tags/v1.613.1.zip;47653509a3371eabb156360f42faf582f314bf2e
-cudnn_frontend;https://github.com/NVIDIA/cudnn-frontend/archive/refs/tags/v1.24.0.zip;a55a1980bf5c57692d66ae7bc3b39798f5535e1f
+cudnn_frontend;https://github.com/NVIDIA/cudnn-frontend/archive/refs/tags/v1.27.0.zip;1e4c9a464d3437e388ab0163f3be068dba783c08
 dawn;https://github.com/google/dawn/archive/refs/tags/v20260714.215939.zip;3056ed22d1606258ab43221b8c85b55b88614137
 kleidiai;https://github.com/ARM-software/kleidiai/archive/refs/tags/v1.20.0.tar.gz;6895e72b3d5cf1173358164cb3d64c9d7d33cc84
 # kleidiai-qmx is pinned to a specific commit as there are no tagged releases. When an appropriate tagged release becomes available,

diff --git a/cmake/external/cutlass.cmake b/cmake/external/cutlass.cmake
@@ -4,7 +4,7 @@ onnxruntime_fetchcontent_declare(
   URL ${DEP_URL_cutlass}
   URL_HASH SHA1=${DEP_SHA1_cutlass}
   EXCLUDE_FROM_ALL
-  PATCH_COMMAND ${Patch_EXECUTABLE} --ignore-whitespace -p1 < ${PROJECT_SOURCE_DIR}/patches/cutlass/cutlass_4.4.2.patch
+  PATCH_COMMAND ${Patch_EXECUTABLE} --ignore-whitespace -p1 < ${PROJECT_SOURCE_DIR}/patches/cutlass/cutlass_4.7.0.patch
 )
 
 # We only consume CUTLASS as a header-only dependency. Avoid FetchContent_MakeAvailable here

diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake
@@ -1066,6 +1066,7 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_P
       "${ONNXRUNTIME_ROOT}/core/providers/cuda/cudnn_common.cc"
       "${ONNXRUNTIME_ROOT}/core/providers/cuda/cudnn_loader.cc"
       "${ONNXRUNTIME_ROOT}/core/providers/cuda/cufft_loader.cc"
+      "${ONNXRUNTIME_ROOT}/core/providers/cuda/fpgeneric.cu"
       "${ONNXRUNTIME_ROOT}/core/providers/cuda/reduction/reduction_functions.cc"
       "${ONNXRUNTIME_ROOT}/core/providers/cuda/reduction/reduction_functions.cu"
       "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_plugin_test_shims.cc"

diff --git a/cmake/patches/cudnn_frontend/cudnn_frontend_win_dynamic_loading.patch b/cmake/patches/cudnn_frontend/cudnn_frontend_win_dynamic_loading.patch
@@ -15,7 +15,16 @@
      const char *error = reinterpret_cast<const char *>(dlerror());
      if (!handle || error) {
          // If opening the library fails, throw an exception with the error message
-@@ -85,13 +92,22 @@
+@@ -82,5 +89,8 @@
+      if (user_lib) {
+           if (user_lib[0] != '\0') {
++#ifdef _WIN32
++            SetLastError(0);
++#endif
+                HMODULE handle    = dlopen(user_lib, RTLD_NOW);
+                const char *error = reinterpret_cast<const char *>(dlerror());
+                if (!handle || error) {
+@@ -85,13 +95,22 @@
      dlerror();
 
      // List of potential libcudart libraries (Adding major version to support python package)

diff --git a/cmake/patches/cutlass/cutlass_4.4.2.patch → cmake/patches/cutlass/cutlass_4.7.0.patch b/cmake/patches/cutlass/cutlass_4.4.2.patch → cmake/patches/cutlass/cutlass_4.7.0.patch
@@ -26,11 +26,17 @@ index a8af62be..22e7332d 100644
    virtual Status memsetDeviceImpl(
      void* destination, ///< Device memory pointer to be filled
      void const* fill_value, ///< Value to be filled in the buffer
+diff --git a/include/cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp b/include/cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp
+index ae31620a..fd559cbb 100644
+--- a/include/cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp
++++ b/include/cutlass/gemm/kernel/sm100_static_tile_scheduler.hpp
+@@ -53 +52,0 @@ public:
+-  CUTLASS_HOST_DEVICE
 diff --git a/include/cutlass/exmy_base.h b/include/cutlass/exmy_base.h
 index be207a49..6028e01d 100644
 --- a/include/cutlass/exmy_base.h
 +++ b/include/cutlass/exmy_base.h
-@@ -1021,18 +1021,18 @@ struct float_exmy_base
+@@ -1021,17 +1021,17 @@ struct float_exmy_base
 
    /// Floating point conversion
    CUTLASS_HOST_DEVICE
@@ -51,4 +57,15 @@ index be207a49..6028e01d 100644
 +  explicit float_exmy_base(unsigned x) {
      storage = static_cast<Derived*>(this)->convert_from_float(float(x)).storage;
    }
-
+diff --git a/include/cutlass/subbyte_reference.h b/include/cutlass/subbyte_reference.h
+--- a/include/cutlass/subbyte_reference.h
++++ b/include/cutlass/subbyte_reference.h
+@@ -456,7 +456,7 @@ public:
+    //
+    Storage assumed;
+ #if (__CUDACC_VER_MAJOR__ > 12) || (__CUDACC_VER_MAJOR__ == 12 && __CUDACC_VER_MINOR__ >= 8)
+-    Storage original = __nv_atomic_load_n(ptr_, __NV_ATOMIC_RELAXED);
++    Storage original = __nv_atomic_load_n(ptr_, __NV_ATOMIC_RELAXED, __NV_THREAD_SCOPE_DEVICE);
+ #else
+    Storage original = *const_cast<Storage const volatile *>(ptr_);
+ #endif
diff --git a/docs/ContribOperators.md b/docs/ContribOperators.md
@@ -7214,6 +7214,8 @@ No versioning maintained for experimental ops.
   The embedding layer takes input_ids (word IDs) and segment_ids (sentence IDs) to look up word_embedding, position_embedding,
   and segment_emedding; the embeddings are added then applied layer normalization using gamma and beta tensors. The input_ids
   and segment_ids remain int32. All embeddings, gamma, and beta tensors are converted to int8/uint8. The last input mask is optional.
+  segment_ids, segment_embedding, segment_embedding_scale, and segment_embedding_zero_point must either all be provided or all
+  be omitted.
   If mask is provided, mask index (that is position of first 0 in mask, or number of words will be calculated.
 
 #### Version
@@ -7232,13 +7234,13 @@ No versioning maintained for experimental ops.
 <dt><tt>input_ids</tt> : T1</dt>
 <dd>2D words IDs with shape (batch_size, sequence_length)</dd>
 <dt><tt>segment_ids</tt> (optional) : T1</dt>
-<dd>2D segment IDs with shape (batch_size, sequence_length)</dd>
+<dd>2D segment IDs with shape (batch_size, sequence_length). Part of the all-or-none segment input group.</dd>
 <dt><tt>word_embedding_quant</tt> : T2</dt>
 <dd>2D with shape (,hidden_size)</dd>
 <dt><tt>position_embedding_quant</tt> : T2</dt>
 <dd>2D with shape (, hidden_size)</dd>
 <dt><tt>segment_embedding</tt> (optional) : T2</dt>
-<dd>2D with shape (, hidden_size)</dd>
+<dd>2D with shape (, hidden_size). Part of the all-or-none segment input group.</dd>
 <dt><tt>gamma_quant</tt> : T2</dt>
 <dd>1D gamma tensor for layer normalization with shape (hidden_size)</dd>
 <dt><tt>beta_quant</tt> : T2</dt>
@@ -7250,7 +7252,7 @@ No versioning maintained for experimental ops.
 <dt><tt>position_embedding_scale</tt> : T</dt>
 <dd>Scale for position embeddings</dd>
 <dt><tt>segment_embedding_scale</tt> (optional) : T</dt>
-<dd>Scale for segment embeddings</dd>
+<dd>Scale for segment embeddings. Part of the all-or-none segment input group.</dd>
 <dt><tt>gamma_scale</tt> : T</dt>
 <dd>Scale for 1D gamma tensor</dd>
 <dt><tt>beta_scale</tt> : T</dt>
@@ -7260,7 +7262,7 @@ No versioning maintained for experimental ops.
 <dt><tt>position_embedding_zero_point</tt> : T2</dt>
 <dd>Zero point for position embeddings</dd>
 <dt><tt>segment_embedding_zero_point</tt> (optional) : T2</dt>
-<dd>Zero Point for segment embeddings</dd>
+<dd>Zero Point for segment embeddings. Part of the all-or-none segment input group.</dd>
 <dt><tt>gamma_zero_point</tt> : T2</dt>
 <dd>Zero Point for 1D gamma tensor</dd>
 <dt><tt>beta_zero_point</tt> : T2</dt>

diff --git a/onnxruntime/contrib_ops/cpu/quantization/qembed_layer_norm.cc b/onnxruntime/contrib_ops/cpu/quantization/qembed_layer_norm.cc
@@ -187,6 +187,8 @@ Status ComputeInternal(OpKernelContext* context, float epsilon) {
 }
 
 Status CheckQuantizedInputs(OpKernelContext* context, bool* is_signed_inputs) {
+  const Tensor* segment_ids_tensor = context->Input<Tensor>(1);
+  const Tensor* segment_embedding_tensor = context->Input<Tensor>(4);
   const Tensor* word_embedding_scale_tensor = context->Input<Tensor>(8);
   const Tensor* position_embedding_scale_tensor = context->Input<Tensor>(9);
   const Tensor* segment_embedding_scale_tensor = context->Input<Tensor>(10);
@@ -198,8 +200,14 @@ Status CheckQuantizedInputs(OpKernelContext* context, bool* is_signed_inputs) {
   const Tensor* gamma_zero_point_tensor = context->Input<Tensor>(16);
   const Tensor* beta_zero_point_tensor = context->Input<Tensor>(17);
 
+  const bool has_segment_embedding = segment_ids_tensor != nullptr;
+  ORT_RETURN_IF(has_segment_embedding != (segment_embedding_tensor != nullptr) ||
+                    has_segment_embedding != (segment_embedding_scale_tensor != nullptr) ||
+                    has_segment_embedding != (segment_embedding_zero_point_tensor != nullptr),
+                "segment_ids, segment_embedding, segment_embedding_scale, and segment_embedding_zero_point "
+                "must either all be provided or all be omitted");
+
   bool word_embedding_is_signed_inputs = word_embedding_zero_point_tensor->IsDataType<int8_t>();
-  bool has_segment_embedding = context->Input<Tensor>(1) != nullptr;
 
   if (!IsScalarOr1ElementVector(word_embedding_scale_tensor)) {
     return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT,

diff --git a/...ime/contrib_ops/cuda/llm/cutlass_extensions/epilogue/collective/epilogue_moe_finalize.hpp b/...ime/contrib_ops/cuda/llm/cutlass_extensions/epilogue/collective/epilogue_moe_finalize.hpp
@@ -449,20 +449,20 @@ struct EpilogueMoeFusedFinalizeBuilder {
 
     // Dummy methods to perform different parts of TMA/Tensormap modifications
 
-    template <bool IsLoad, class ProblemShapeMNKL>
+    template <bool IsLoad, bool WaitForInflightTmaRequests = true, class ProblemShapeMNKL, class TensorMaps>
     CUTLASS_DEVICE void tensormaps_perform_update([[maybe_unused]] TensorMapStorage& shared_tensormaps,
                                                   [[maybe_unused]] typename EpilogueOp::Params const& params,
-                                                  [[maybe_unused]] cute::TmaDescriptor const* tensormap, [[maybe_unused]] ProblemShapeMNKL problem_shape,
+                                                  [[maybe_unused]] TensorMaps const& tensormaps, [[maybe_unused]] ProblemShapeMNKL problem_shape,
                                                   [[maybe_unused]] int32_t next_batch, [[maybe_unused]] int32_t warp_group_idx) {
     }
 
-    template <bool IsLoad>
+    template <bool IsLoad, bool WaitForInflightTmaRequests = true, class TensorMaps>
     CUTLASS_DEVICE void tensormaps_cp_fence_release([[maybe_unused]] TensorMapStorage& shared_tensormaps,
-                                                    [[maybe_unused]] cute::TmaDescriptor const* tensormap, [[maybe_unused]] int32_t warp_group_idx) {
+                                                    [[maybe_unused]] TensorMaps const& tensormaps, [[maybe_unused]] int32_t warp_group_idx) {
     }
 
-    template <bool IsLoad>
-    CUTLASS_DEVICE void tensormaps_fence_acquire([[maybe_unused]] cute::TmaDescriptor const* tensormap) {
+    template <bool IsLoad, class TensorMaps>
+    CUTLASS_DEVICE void tensormaps_fence_acquire([[maybe_unused]] TensorMaps const& tensormaps) {
     }
   };
 

diff --git a/onnxruntime/core/framework/tensor_external_data_info.cc b/onnxruntime/core/framework/tensor_external_data_info.cc
@@ -29,6 +29,7 @@ Status ExternalDataInfo::Create(const RepeatedPtrField<StringStringEntryProto>&
                                 std::unique_ptr<ExternalDataInfo>& external_data_info_result) {
   auto external_data_info = std::make_unique<ExternalDataInfo>();
   PrepackedInfos prepacked_infos;
+  bool has_location = false;
 
   const int input_size = input.size();
 
@@ -39,7 +40,11 @@ Status ExternalDataInfo::Create(const RepeatedPtrField<StringStringEntryProto>&
     if (!stringmap.has_value())
       return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "model format error! Need a value for the external data info");
 
-    if (stringmap.key() == "location" && !stringmap.value().empty()) {
+    if (stringmap.key() == "location") {
+      ORT_RETURN_IF(has_location,
+                    "model format error! TensorProto external data has duplicate 'location' entries");
+      has_location = true;
+      ORT_RETURN_IF(stringmap.value().empty(), "model format error! External data location cannot be empty");
       external_data_info->rel_path_ = ToWideString(stringmap.value());
     } else if (stringmap.key() == "offset" && !stringmap.value().empty()) {
       ORT_RETURN_IF_ERROR(ParseStringWithClassicLocale(stringmap.value(), external_data_info->offset_));

diff --git a/onnxruntime/core/framework/tensorprotoutils.cc b/onnxruntime/core/framework/tensorprotoutils.cc
@@ -312,11 +312,13 @@ namespace utils {
 
 bool HasExternalDataInMemory(const ONNX_NAMESPACE::TensorProto& ten_proto) {
   if (HasExternalData(ten_proto)) {
-    // Retrieve the external data info
     for (const auto& entry : ten_proto.external_data()) {
       if (entry.key() == "location") {
-        PathString location = ToWideString(entry.value());
-        return ((location == kTensorProtoLittleEndianMemoryAddressTag) || (location == kTensorProtoNativeEndianMemoryAddressTag));
+        const PathString location = ToWideString(entry.value());
+        if (location == kTensorProtoLittleEndianMemoryAddressTag ||
+            location == kTensorProtoNativeEndianMemoryAddressTag) {
+          return true;
+        }
       }
     }
   }
@@ -407,11 +409,12 @@ static bool HasPathComponentPrefix(const std::filesystem::path& prefix, const st
 ///
 /// Validation steps:
 ///   1. Reject empty paths
-///   2. Reject absolute paths (including Unix-style '/...' on Windows)
-///   3. Skip remaining checks on WASM if no filesystem is available
-///   4. Resolve `model_dir / external_data_path` to a canonical path (resolving symlinks for existing segments)
-///   5. Verify the canonical path is a prefix-child of the canonical model_dir (containment check)
-///   6. Verify the resolved file exists on disk
+///   2. Reject internal in-memory reference tags
+///   3. Reject absolute paths (including Unix-style '/...' on Windows)
+///   4. Skip remaining checks on WASM if no filesystem is available
+///   5. Resolve `model_dir / external_data_path` to a canonical path (resolving symlinks for existing segments)
+///   6. Verify the canonical path is a prefix-child of the canonical model_dir (containment check)
+///   7. Verify the resolved file exists on disk
 ///
 /// This function does NOT handle the symlinked-model fallback — that is the responsibility of
 /// ValidateExternalDataPath(), which calls this function as a first pass.
@@ -420,13 +423,18 @@ Status ValidateExternalDataPathFromDir(const std::filesystem::path& model_dir,
   // Step 1: Reject empty external data paths.
   ORT_RETURN_IF(external_data_path.empty(), "Empty external data path not allowed");
 
-  // Step 2: Reject absolute paths.
+  // Step 2: Reject internal in-memory reference tags.
+  ORT_RETURN_IF(external_data_path.native() == kTensorProtoLittleEndianMemoryAddressTag ||
+                    external_data_path.native() == kTensorProtoNativeEndianMemoryAddressTag,
+                "In-memory external data reference tag is not a valid file path");
+
+  // Step 3: Reject absolute paths.
   // Use !root_path().empty() to reject paths like '/some/path' even on Windows (where is_absolute()
   // requires a drive letter).
   ORT_RETURN_IF(!external_data_path.root_path().empty(), "Absolute path not allowed for external data location");
 
 #if defined(__wasm__)
-  // Step 3 (WASM only): If we can't access the current working directory, assume the WASM environment
+  // Step 4 (WASM only): If we can't access the current working directory, assume the WASM environment
   // does not have a virtual filesystem and defer validation to an ExternalDataLoader for the WASM EP.
   std::error_code error_code;
   std::filesystem::current_path(error_code);
@@ -435,7 +443,7 @@ Status ValidateExternalDataPathFromDir(const std::filesystem::path& model_dir,
   }
 #endif
 
-  // Step 4: Resolve both the model directory and the combined path to canonical forms.
+  // Step 5: Resolve both the model directory and the combined path to canonical forms.
   // WeaklyCanonicalPath resolves symlinks for existing path segments while lexically normalizing
   // non-existent trailing segments.
   std::filesystem::path resolved_dir = model_dir.empty() ? std::filesystem::path{"."} : model_dir;
@@ -445,8 +453,8 @@ Status ValidateExternalDataPathFromDir(const std::filesystem::path& model_dir,
   ORT_RETURN_IF_ERROR(WeaklyCanonicalPath(resolved_dir, model_dir_canonical));
   ORT_RETURN_IF_ERROR(WeaklyCanonicalPath(model_dir_canonical / external_data_path, external_data_path_canonical));
 
-  // Step 5: Containment check — verify the resolved external data path starts with the model directory.
-  // Step 6: Existence check — verify the file actually exists on disk.
+  // Step 6: Containment check — verify the resolved external data path starts with the model directory.
+  // Step 7: Existence check — verify the file actually exists on disk.
   if (HasPathComponentPrefix(model_dir_canonical, external_data_path_canonical)) {
     bool path_exists = false;
     ORT_RETURN_IF_ERROR(PathExists(external_data_path_canonical, path_exists));
@@ -466,7 +474,7 @@ Status ValidateExternalDataPathFromDir(const std::filesystem::path& model_dir,
 /// Validation flow:
 ///   1. Try ValidateExternalDataPathFromDir against the model file's parent directory.
 ///      If it passes, return success.
-///   2. If it fails due to empty/absolute external_data_path, return the error immediately
+///   2. If it fails due to an empty/absolute path or an in-memory reference tag, return the error immediately
 ///      (these are input errors unrelated to the model location).
 ///   3. If model_path is empty (model loaded from bytes), wrap the error with context.
 ///   4. If model_path is a symlink, try the symlink fallback:
@@ -494,9 +502,11 @@ Status ValidateExternalDataPath(const std::filesystem::path& model_path,
   }
 
   // --- Guard: Don't retry for input-validation errors ---
-  // Empty and absolute paths are always invalid regardless of model directory or symlinks.
+  // Empty paths, absolute paths, and in-memory reference tags are always invalid regardless of model directory.
   // Return the error directly without misleading "escapes directory" context.
-  if (external_data_path.empty() || !external_data_path.root_path().empty()) {
+  if (external_data_path.empty() || !external_data_path.root_path().empty() ||
+      external_data_path.native() == kTensorProtoLittleEndianMemoryAddressTag ||
+      external_data_path.native() == kTensorProtoNativeEndianMemoryAddressTag) {
     return status;
   }