Sync with Microsoft ONNX Runtime - 09072026

cmake/CMakeLists.txt

@@ -77,6 +77,7 @@ cmake_dependent_option(onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS "Build with CUD
 
 cmake_dependent_option(onnxruntime_USE_CUDA_NHWC_OPS "Build CUDA with NHWC op support" ON "onnxruntime_USE_CUDA" OFF)
 cmake_dependent_option(onnxruntime_BUILD_CUDA_EP_AS_PLUGIN "Build CUDA EP as a separate plugin shared library instead of the legacy in-tree provider" OFF "onnxruntime_USE_CUDA" OFF)
+option(onnxruntime_BUILD_CUDA_QUANT_PREPROCESS "Build CUDA weight-packing module onnxruntime_cuda_quant_preprocess.so" ON)
 option(onnxruntime_CUDA_MINIMAL "Build CUDA without any operations apart from memcpy ops. Usefuel for a very minial TRT build" OFF)
 option(onnxruntime_ENABLE_CUDA_LINE_NUMBER_INFO "When building with CUDA support, generate device code line number information." OFF)
 option(onnxruntime_USE_OPENVINO "Build with OpenVINO support" OFF)

cmake/onnxruntime_python.cmake

@@ -20,6 +20,13 @@ file(GLOB onnxruntime_pybind_srcs CONFIGURE_DEPENDS
   ${onnxruntime_pybind_srcs_pattern}
   )
 
+# onnxruntime_pybind_cuda_quant.cc is compiled into the standalone
+# onnxruntime_cuda_quant_preprocess extension module (see below), not into
+# onnxruntime_pybind11_state. It includes <cuda_runtime.h> and links CUDA::cudart,
+# so compiling it into the main pybind module would break CPU-only builds and
+# re-introduce the hard libcudart dependency this design avoids.
+list(REMOVE_ITEM onnxruntime_pybind_srcs ${ONNXRUNTIME_ROOT}/python/onnxruntime_pybind_cuda_quant.cc)
+
 if(onnxruntime_ENABLE_TRAINING)
   list(REMOVE_ITEM onnxruntime_pybind_srcs  ${ONNXRUNTIME_ROOT}/python/onnxruntime_pybind_module.cc)
 endif()
@@ -231,22 +238,11 @@ target_link_libraries(onnxruntime_pybind11_state PRIVATE
     Python::NumPy
 )
 
-# Starting with Python 3.8 on Windows, PATH environment variable are no longer used to resolve DLL dependencies
-# for extension modules or libraries loaded via ctypes.
-# To avoid package import issues, we do not link pybind module against the CUDA runtime on Windows, instead of
-# os.add_dll_directory() to deal with CUDA paths.
-if (onnxruntime_USE_CUDA AND NOT WIN32)
-  target_sources(onnxruntime_pybind11_state PRIVATE
-    "${ONNXRUNTIME_ROOT}/contrib_ops/cuda/llm/fpA_intB_gemm_adaptor.cu"
-    "${ONNXRUNTIME_ROOT}/contrib_ops/cuda/llm/fpA_intB_gemm_preprocessors_impl.cu"
-  )
-  include(cutlass)
-  target_include_directories(onnxruntime_pybind11_state PRIVATE ${cutlass_SOURCE_DIR}/include)
-  target_link_libraries(onnxruntime_pybind11_state PRIVATE CUDA::cudart)
-endif()
-if (onnxruntime_USE_CUDA AND WIN32)
-  target_compile_definitions(onnxruntime_pybind11_state PRIVATE ORT_NO_CUDA_IN_PYBIND)
-endif()
+# The CUDA quantization helpers (pack_weights_for_cuda_mixed_gemm) are built into a
+# separate extension module (onnxruntime_cuda_quant_preprocess) that is imported on
+# demand. Do NOT compile CUDA source files directly into onnxruntime_pybind11_state or
+# link CUDA::cudart from it: that would create a hard libcudart.so dependency that
+# prevents importing the Python module on CPU-only machines.
 
 set(onnxruntime_pybind11_state_dependencies
     ${onnxruntime_EXTERNAL_DEPENDENCIES}
@@ -315,6 +311,71 @@ else()
   set_target_properties(onnxruntime_pybind11_state PROPERTIES SUFFIX ".so")
 endif()
 
+# ---------------------------------------------------------------------------
+# Standalone CUDA weight-preprocessing extension module.
+#
+# The CUDA weight-packing kernels (pack_weights_for_cuda_mixed_gemm) are compiled
+# into their OWN Python extension module instead of onnxruntime_pybind11_state.
+# This keeps the hard libcudart dependency out of the main pybind module so that
+# `import onnxruntime` still works on CPU-only machines.
+#
+# Production weight packing is done in PyTorch (cuda_quantizer.py); this module is
+# retained only as a byte-parity oracle for that PyTorch packer. It is gated by
+# onnxruntime_BUILD_CUDA_QUANT_PREPROCESS.
+#
+# It does NOT go through the provider bridge / ProviderInfo_CUDA, so it works for
+# both the legacy in-tree CUDA EP build and the CUDA-EP-as-plugin build.
+#
+# Not built on Windows: matching the previous behavior where CUDA runtime was not
+# linked into Python extension modules (DLL search path constraints since
+# Python 3.8), so pack_weights_for_cuda_mixed_gemm was unavailable there.
+if (onnxruntime_USE_CUDA AND NOT WIN32 AND onnxruntime_BUILD_CUDA_QUANT_PREPROCESS)
+  onnxruntime_add_shared_library_module(onnxruntime_cuda_quant_preprocess
+    "${ONNXRUNTIME_ROOT}/python/onnxruntime_pybind_cuda_quant.cc"
+    "${ONNXRUNTIME_ROOT}/contrib_ops/cuda/llm/fpA_intB_gemm_adaptor.cu"
+    "${ONNXRUNTIME_ROOT}/contrib_ops/cuda/llm/fpA_intB_gemm_preprocessors_impl.cu"
+  )
+  include(cutlass)
+  onnxruntime_add_include_to_target(onnxruntime_cuda_quant_preprocess Python::Module onnxruntime_common)
+  target_include_directories(onnxruntime_cuda_quant_preprocess PRIVATE
+    ${ONNXRUNTIME_ROOT}
+    ${pybind11_INCLUDE_DIRS}
+    ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES}
+    ${cutlass_SOURCE_DIR}/include
+    ${cutlass_SOURCE_DIR}/tools/util/include
+  )
+  target_compile_definitions(onnxruntime_cuda_quant_preprocess PRIVATE USE_CUDA)
+  target_link_libraries(onnxruntime_cuda_quant_preprocess PRIVATE
+    onnxruntime_common
+    Boost::mp11
+    safeint_interface
+    ${ABSEIL_LIBS}
+    CUDA::cudart
+    ${pybind11_lib}
+    Python::NumPy
+  )
+  if (NOT MSVC)
+    target_compile_options(onnxruntime_cuda_quant_preprocess PRIVATE "-fvisibility=hidden")
+  endif()
+  set_target_properties(onnxruntime_cuda_quant_preprocess PROPERTIES PREFIX "" SUFFIX ".so" FOLDER "ONNXRuntime")
+  if (APPLE)
+    set_target_properties(onnxruntime_cuda_quant_preprocess PROPERTIES
+      INSTALL_RPATH "@loader_path"
+      BUILD_WITH_INSTALL_RPATH TRUE
+      INSTALL_RPATH_USE_LINK_PATH FALSE)
+  elseif (NOT CMAKE_SYSTEM_NAME MATCHES "AIX")
+    target_link_options(onnxruntime_cuda_quant_preprocess PRIVATE "LINKER:-rpath=\$ORIGIN")
+  endif()
+  # Place the module next to the main pybind module inside onnxruntime/capi.
+  add_custom_command(
+    TARGET onnxruntime_cuda_quant_preprocess POST_BUILD
+    COMMAND ${CMAKE_COMMAND} -E make_directory $<TARGET_FILE_DIR:onnxruntime_common>/onnxruntime/capi
+    COMMAND ${CMAKE_COMMAND} -E copy
+        $<TARGET_FILE:onnxruntime_cuda_quant_preprocess>
+        $<TARGET_FILE_DIR:onnxruntime_common>/onnxruntime/capi/
+  )
+endif()
+
 # Generate version_info.py in Windows build.
 # Has to be done before onnxruntime_python_srcs is set.
 if (WIN32)

docs/contrib_ops/cuda/matmul_nbits.md

@@ -87,9 +87,9 @@ step is **not** performed.
 The offline CUDA packer exposed through Python produces this layout:
 
 ```python
-from onnxruntime.capi import _pybind_state as _pybind
+from onnxruntime.capi import onnxruntime_cuda_quant_preprocess as _cuda_quant
 
-prepacked_flat = _pybind.pack_weights_for_cuda_mixed_gemm(
+prepacked_flat = _cuda_quant.pack_weights_for_cuda_mixed_gemm(
   q_weight.reshape(N, -1), N, K, bits, 80
 )
 prepacked_b = np.asarray(prepacked_flat, dtype=np.int8).view(np.uint8).reshape(q_weight.shape)
@@ -309,7 +309,7 @@ present. `ComputeInternal` then:
   GEMV profiling helpers, e.g. `profile_qmoe_gemv.sh`).
 - CUDA prepacked-weight parity tests:
   [onnxruntime/test/python/quantization/test_op_matmulnbits_prepacked_cuda.py](../../../onnxruntime/test/python/quantization/test_op_matmulnbits_prepacked_cuda.py).
-  These use `_pybind_state.pack_weights_for_cuda_mixed_gemm(..., 80)` to produce
+  These use `onnxruntime_cuda_quant_preprocess.pack_weights_for_cuda_mixed_gemm(..., 80)` to produce
   `weight_prepacked=1` initializers and compare their outputs against runtime
   fpA_intB prepacking for int4/int8 and GEMV/GEMM-shaped `M` values.
 - Constructor failure tests for unsupported prepacked configurations live in

docs/cuda_plugin_ep/QUICK_START.md

@@ -146,7 +146,7 @@ sess = ort.InferenceSession(
 
 **Python `OrtValue` host/device copies:**
 
-`OrtValue.update_inplace()` and `OrtValue.numpy()` work with CUDA plugin tensors after the plugin has been registered. On Linux, the ONNX Runtime Python binding links the CUDA runtime and can fall back to direct `cudaMemcpy` if the legacy CUDA provider bridge is unavailable. On Windows, the Python binding is built with `ORT_NO_CUDA_IN_PYBIND`, so it cannot call CUDA runtime APIs directly; host/device copies must use the data-transfer implementation registered by the CUDA plugin library. If `OrtValue.update_inplace()` fails with a message about the CUDA provider interface or an unsupported GPU device, verify that the plugin library is registered before creating or updating CUDA `OrtValue` objects.
+`OrtValue.update_inplace()` and `OrtValue.numpy()` work with CUDA plugin tensors after the plugin has been registered. The Python binding cannot call CUDA runtime APIs directly; host/device copies must use the data-transfer implementation registered by the CUDA plugin library. If `OrtValue.update_inplace()` fails with a message about the CUDA provider interface or an unsupported GPU device, verify that the plugin library is registered before creating or updating CUDA `OrtValue` objects.
 
 ### External GPU Allocator Options
 

docs/cuda_plugin_ep/cuda_plugin_ep_design.md

@@ -344,10 +344,12 @@ This is intentionally conservative and correct for the plugin EP's first sync in
 
 The Python `OrtValue` helpers (`update_inplace()` for host-to-device and `numpy()` for device-to-host) historically reached CUDA copies through the legacy provider bridge (`GetProviderInfo_CUDA()`). That bridge requires the provider shared library to export `GetProvider()`, which the CUDA plugin intentionally does not export.
 
-The fallback path is platform-specific:
+To keep working when the bridge is absent (as with the plugin EP), the pybind can reach CUDA copies two ways: the legacy provider bridge (`TryGetProviderInfo_CUDA()`) and a plugin-registered `OrtDataTransfer` copy function (`CreateDataTransferMemCpy()`, backed by the plugin EP's `IDataTransfer`). It tries whichever is available and throws if neither is, in which case a CUDA `OrtValue` copy cannot be performed.
 
-- On non-Windows CUDA builds, `onnxruntime_pybind11_state` links `CUDA::cudart`. If `TryGetProviderInfo_CUDA()` fails, pybind can copy directly with `cudaMemcpy`; host-to-device copies synchronize the default stream, matching `ProviderInfo_CUDA::cudaMemcpy_HostToDevice()`.
-- On Windows CUDA builds, pybind is compiled with `ORT_NO_CUDA_IN_PYBIND` and does not link CUDA runtime APIs. If `TryGetProviderInfo_CUDA()` fails, pybind must obtain an `OrtDataTransfer` copy function from the registered plugin EP. Without a registered plugin data-transfer implementation, CUDA `OrtValue.update_inplace()` cannot copy host data into the plugin-owned device tensor.
+The two code paths differ only in which mechanism they try first, and this does not change the outcome (exactly one applies in a given build):
+
+- `OrtValue.update_inplace(numpy_array)` / `OrtValue.numpy()` (in `onnxruntime_pybind_ortvalue.cc`) try the provider bridge first, then fall back to the plugin `OrtDataTransfer`.
+- `OrtValue.update_inplace(OrtValue)` (`UpdateOrtValueInplace` in `onnxruntime_pybind_mlvalue.cc`) tries the plugin `OrtDataTransfer` first, then falls back to the built-in CUDA provider copy functions.
 
 ### 5.2 Handle Access Path
 

onnxruntime/contrib_ops/cpu/crop_and_resize.cc

@@ -17,16 +17,12 @@ limitations under the License.
 #include "contrib_ops/cpu/crop_and_resize.h"
 
 #include <cmath>
+#include "core/common/safeint.h"
 #include "core/util/math_cpuonly.h"
 #include "core/common/common.h"
 #include "core/framework/tensor.h"
 #include "core/platform/threadpool.h"
 #include "core/providers/cpu/object_detection/roialign.h"
-// TODO: fix the warnings
-#if defined(_MSC_VER) && !defined(__clang__)
-// Chance of arithmetic overflow could be reduced
-#pragma warning(disable : 26451)
-#endif
 using namespace onnxruntime::concurrency;
 
 namespace onnxruntime {
@@ -97,7 +93,11 @@ void CropAndResizeForward(const TensorShape& output_shape,
                                       ? roi_start_h * (height - 1)
                                       : 0.5 * (roi_start_h + roi_end_h) * (height - 1));
           }
-          if (in_y < 0 || in_y > height - 1) {
+          // Route non-finite (NaN/inf) or out-of-image coordinates to the extrapolation
+          // branch. The negated-conjunction form is NaN-safe: every comparison with NaN is
+          // false, so !(in_y >= 0 && in_y <= height - 1) is true and NaN cannot reach the
+          // integer index math below. For finite values this is identical to the < / > form.
+          if (!(in_y >= 0 && in_y <= static_cast<T>(height - 1))) {
             for (int64_t pw = 0; pw < pooled_width; pw++) {
               for (int64_t c = 0; c < channels; c++) {
                 int64_t index_n_c = index_n + c * pooled_width * pooled_height;
@@ -126,7 +126,8 @@ void CropAndResizeForward(const TensorShape& output_shape,
                                         ? roi_start_w * (width - 1)
                                         : 0.5 * (roi_start_w + roi_end_w) * (width - 1));
             }
-            if (in_x < 0 || in_x > width - 1) {
+            // NaN-safe bounds guard (see the in_y guard above for rationale).
+            if (!(in_x >= 0 && in_x <= static_cast<T>(width - 1))) {
               for (int64_t c = 0; c < channels; c++) {
                 int64_t index_n_c = index_n + c * pooled_width * pooled_height;
                 int64_t index = index_n_c + ph * pooled_width + pw;
@@ -140,16 +141,17 @@ void CropAndResizeForward(const TensorShape& output_shape,
               const int left_x_index = static_cast<int>(floorf(static_cast<float>(in_x)));
               const int right_x_index = static_cast<int>(ceilf(static_cast<float>(in_x)));
               const float x_lerp = static_cast<float>(in_x - left_x_index);
-              auto top_left_index = top_y_index * width + left_x_index;
-              auto top_right_index = top_y_index * width + right_x_index;
-              auto bottom_left_index = bottom_y_index * width + left_x_index;
-              auto bottom_right_index = bottom_y_index * width + right_x_index;
+              auto top_left_index = SafeInt<int64_t>(top_y_index) * width + left_x_index;
+              auto top_right_index = SafeInt<int64_t>(top_y_index) * width + right_x_index;
+              auto bottom_left_index = SafeInt<int64_t>(bottom_y_index) * width + left_x_index;
+              auto bottom_right_index = SafeInt<int64_t>(bottom_y_index) * width + right_x_index;
 
               for (auto c = 0; c < channels; c++) {
                 int64_t index_n_c = index_n + c * pooled_width * pooled_height;
                 int64_t index = index_n_c + ph * pooled_width + pw;
                 const T* offset_bottom_data =
-                    bottom_data + static_cast<int64_t>((roi_batch_ind * channels + c) * height * width);
+                    bottom_data +
+                    static_cast<int64_t>((SafeInt<int64_t>(roi_batch_ind) * channels + c) * height * width);
                 const float top_left(static_cast<float>(offset_bottom_data[top_left_index]));
                 const float top_right(static_cast<float>(offset_bottom_data[top_right_index]));
                 const float bottom_left(static_cast<float>(offset_bottom_data[bottom_left_index]));
@@ -162,13 +164,14 @@ void CropAndResizeForward(const TensorShape& output_shape,
             } else {  // mode == "nearest"
               const int closest_x_index = static_cast<int>(roundf(static_cast<float>(in_x)));
               const int closest_y_index = static_cast<int>(roundf(static_cast<float>(in_y)));
-              auto closest_index = closest_y_index * width + closest_x_index;
+              auto closest_index = SafeInt<int64_t>(closest_y_index) * width + closest_x_index;
 
               for (auto c = 0; c < channels; c++) {
                 int64_t index_n_c = index_n + c * pooled_width * pooled_height;
                 int64_t index = index_n_c + ph * pooled_width + pw;
                 const T* offset_bottom_data =
-                    bottom_data + static_cast<int64_t>((roi_batch_ind * channels + c) * height * width);
+                    bottom_data +
+                    static_cast<int64_t>((SafeInt<int64_t>(roi_batch_ind) * channels + c) * height * width);
                 top_data[index] = static_cast<float>(offset_bottom_data[closest_index]);
               }
             }
@@ -217,6 +220,9 @@ Status CropAndResize<T>::Compute(OpKernelContext* context) const {
   auto crop_height = crop_size_data[0];
   auto crop_width = crop_size_data[1];
 
+  ORT_RETURN_IF_NOT(crop_height > 0 && crop_width > 0,
+                    "crop_size values must be positive; got [", crop_height, ", ", crop_width, "]");
+
   auto status = CheckROIAlignValidInput(X_ptr, rois_ptr, batch_indices_ptr);
   if (status != Status::OK()) {
     return status;

onnxruntime/contrib_ops/cpu/transformers/logits_processor.h

@@ -332,7 +332,10 @@ class LogitsProcessorList : public ILogitsProcessorList {
       processor_list_.push_back(prefix_vocab_mask_processor_.get());
     }
 
-    if (parameters.min_length > 0) {
+    // For a negative "no eos" sentinel there is no token to demote, so a MinLength processor here
+    // would be a guaranteed no-op (SetScore ignores negative token ids). Skip constructing it as a
+    // defensive, minor performance optimization, consistent with the other conditional-adds above.
+    if (parameters.min_length > 0 && parameters.eos_token_id >= 0) {
       min_length_processor_ = std::make_unique<MinLengthLogitsProcessor<float>>(parameters.min_length,
                                                                                 parameters.eos_token_id);
       processor_list_.push_back(min_length_processor_.get());

onnxruntime/contrib_ops/cuda/llm/generate_moe_kernels.py

@@ -6,6 +6,7 @@
 #   python generate_moe_kernels.py -a "80;90" -o ./moe_gemm/launchers
 
 import argparse
+import glob
 import os
 from itertools import product
 
@@ -333,19 +334,31 @@ def has_arch(sm):
     if has_arch(80) or has_arch(90):  # SM90 also uses SM80 kernels for non-TMA path
         operations = generate_sm80_moe_operations()
 
-        # Group by element type for separate files to reduce compile time
+        # Split by (element type, tile shape) into separate files. Each SM80 template
+        # instantiation is a full CUTLASS GEMM kernel that nvcc compiles serially within
+        # a translation unit, so packing all of them into one file per dtype creates a
+        # long single-file critical path. Emitting one file per tile shape lets the build
+        # system (Ninja) compile these kernels in parallel and shrinks the slowest object.
         groups = {}
         for op in operations:
-            key = op["element_type"]
-            if key not in groups:
-                groups[key] = []
-            groups[key].append(op)
-
-        for dtype, ops in groups.items():
-            output_file = os.path.join(output_dir, f"fused_moe_gemm_sm80_{dtype}.generated.cu")
+            key = (op["element_type"], op["tile_m"], op["tile_n"], op["tile_k"])
+            groups.setdefault(key, []).append(op)
+
+        current_files = set()
+        for (dtype, tile_m, tile_n, tile_k), ops in groups.items():
+            file_name = f"fused_moe_gemm_sm80_{dtype}_m{tile_m}_n{tile_n}_k{tile_k}.generated.cu"
+            current_files.add(file_name)
+            output_file = os.path.join(output_dir, file_name)
             content = get_sm80_file_content(ops, 80)
             write_file(content, output_file)
 
+        # Remove stale SM80 generated files (e.g. the old monolithic per-dtype files or
+        # split files from a previous tile configuration) so they are not compiled.
+        for stale in glob.glob(os.path.join(output_dir, "fused_moe_gemm_sm80_*.generated.cu")):
+            if os.path.basename(stale) not in current_files:
+                os.remove(stale)
+                print(f"Removed stale {stale}")
+
     # Generate SM90 TMA Warp Specialized Grouped GEMM kernels
     if has_arch(90):
         operations = generate_sm90_tma_ws_operations()