Fix libcudart.so.13 hard dependency in pybind module breaking import on CPU-only Linux

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,69 @@ 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. The module is imported
+# lazily by onnxruntime.python.tools.quantization.cuda_quantizer only when CUDA
+# weight prepacking is requested.
+#
+# 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)
+  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/python/onnxruntime_pybind_cuda_quant.cc

@@ -0,0 +1,166 @@
+// Copyright (c) Microsoft Corporation. All rights reserved.
+// Licensed under the MIT License.
+
+// Standalone CUDA weight-preprocessing extension module.
+//
+// This module is intentionally kept SEPARATE from onnxruntime_pybind11_state so
+// that `import onnxruntime` never triggers a load-time dependency on the CUDA
+// runtime (libcudart). The CUDA weight packing kernels below link CUDA::cudart,
+// so this module has a hard libcudart dependency -- but it is imported lazily by
+// onnxruntime.python.tools.quantization.cuda_quantizer only when CUDA weight
+// prepacking is actually requested.
+//
+// This approach works for both the legacy in-tree CUDA EP build and the
+// CUDA-EP-as-plugin build (onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON), because it
+// does not rely on the provider bridge / ProviderInfo_CUDA interface (which is
+// not available in plugin builds).
+
+#include <pybind11/pybind11.h>
+#include <pybind11/numpy.h>
+
+#include <cuda_runtime.h>
+
+#include <memory>
+#include <sstream>
+#include <stdexcept>
+
+#include "contrib_ops/cuda/llm/fpA_intB_gemm_adaptor.h"
+#include "contrib_ops/cuda/llm/fpA_intB_gemm_preprocessors.h"
+
+namespace py = pybind11;
+
+namespace {
+
+void ThrowIfCudaError(cudaError_t status, const char* expression) {
+  if (status != cudaSuccess) {
+    std::ostringstream oss;
+    oss << expression << " failed: " << cudaGetErrorString(status);
+    throw std::runtime_error(oss.str());
+  }
+}
+
+struct CudaDeleter {
+  void operator()(void* p) const {
+    if (p) cudaFree(p);
+  }
+};
+
+using CudaPtr = std::unique_ptr<void, CudaDeleter>;
+
+// Preprocess quantized weights for CUDA mixed-precision GEMM kernels (FpA_IntB format).
+//
+// MatMulNBits/QMoE stores quantized weights in (N, K) layout:
+//   - N = number of output channels (columns in weight matrix W)
+//   - K = number of input features (rows in weight matrix W)
+//   - For 4-bit: shape is (N, K/2) bytes where each byte packs 2 elements
+//   - For 8-bit: shape is (N, K) bytes
+//
+// FpA_IntB GEMM kernels expect weights in (K, N) layout (transposed) for efficient
+// memory access during matrix multiplication. This function:
+//   1. Transposes from (N, K) to (K, N) layout
+//   2. Converts unsigned quantized values to signed int8 with zero-point adjustment
+//      - 4-bit: uint4 -> int8 with zero_point=8 (range [0,15] -> [-8,7])
+//      - 8-bit: uint8 -> int8 with zero_point=128 (range [0,255] -> [-128,127])
+//   3. Applies architecture-specific row permutation for optimized tensor core access
+//
+// Input:  q_weights - Quantized weights from MatMulNBits in (N, K) layout
+// Output: Preprocessed weights in (K, N) layout ready for fpA_intB GEMM kernels
+py::array_t<int8_t> PackWeightsForMixedGemm(
+    py::array_t<uint8_t> q_weights,
+    int32_t N,
+    int32_t K,
+    int32_t bits,
+    int32_t force_arch = -1) {
+  py::buffer_info q_weights_buf = q_weights.request();
+
+  if (bits != 4 && bits != 8) {
+    throw std::invalid_argument("bits must be 4 or 8");
+  }
+  if (N <= 0 || K <= 0) {
+    throw std::invalid_argument("N and K must be positive");
+  }
+  if (bits == 4 && K % 2 != 0) {
+    throw std::invalid_argument("K must be even for 4-bit packed weights");
+  }
+  if (q_weights_buf.ndim != 2 || q_weights_buf.shape[0] != N || q_weights_buf.shape[1] != K / (8 / bits)) {
+    throw std::invalid_argument("q_weights must have shape (N, K / (8 / bits))");
+  }
+
+  int n = static_cast<int>(N);
+  int k = static_cast<int>(K);
+
+  size_t packed_weight_bytes = static_cast<size_t>(n) * static_cast<size_t>(k) / (8 / bits);
+  py::array_t<int8_t> processed_weights({static_cast<pybind11::ssize_t>(packed_weight_bytes)});
+  py::buffer_info processed_weights_buf = processed_weights.request();
+
+  auto make_cuda_ptr = [](size_t bytes) -> CudaPtr {
+    void* p = nullptr;
+    ThrowIfCudaError(cudaMalloc(&p, bytes), "cudaMalloc");
+    return CudaPtr(p);
+  };
+
+  auto packed_transposed_weight_space = make_cuda_ptr(packed_weight_bytes);
+  int8_t* packed_transposed_weight = reinterpret_cast<int8_t*>(packed_transposed_weight_space.get());
+
+  auto fpA_intB_weight_buffer_ = make_cuda_ptr(packed_weight_bytes);
+  int8_t* preprocessed_weight = reinterpret_cast<int8_t*>(fpA_intB_weight_buffer_.get());
+
+  const uint8_t* blob_data_cpu = static_cast<const uint8_t*>(q_weights_buf.ptr);
+
+  auto blob_data_gpu_buf = make_cuda_ptr(packed_weight_bytes);
+  uint8_t* blob_data_gpu = reinterpret_cast<uint8_t*>(blob_data_gpu_buf.get());
+
+  cudaStream_t stream = cudaStreamLegacy;
+  ThrowIfCudaError(cudaMemcpyAsync(blob_data_gpu, blob_data_cpu, packed_weight_bytes, cudaMemcpyHostToDevice, stream),
+                   "cudaMemcpyAsync host-to-device");
+
+  if (bits == 4) {
+    ::onnxruntime::llm::kernels::fpA_intB_gemv::unpack_uint4_transposed_to_int8_direct_cuda(
+        stream, packed_transposed_weight, blob_data_gpu, n, k);
+  } else {
+    // 8 bits
+    ::onnxruntime::llm::kernels::fpA_intB_gemv::transpose_uint8_matrix_and_convert_to_int8(
+        stream, packed_transposed_weight, blob_data_gpu, n, k);
+  }
+
+  using ::onnxruntime::llm::kernels::weight_only::QuantType;
+  QuantType quant_type = bits == 4 ? QuantType::W4_A16 : QuantType::W8_A16;
+
+  int sm = force_arch;
+  if (sm < 0) {
+    int device_id = 0;
+    ThrowIfCudaError(cudaGetDevice(&device_id), "cudaGetDevice");
+    cudaDeviceProp device_prop;
+    ThrowIfCudaError(cudaGetDeviceProperties(&device_prop, device_id), "cudaGetDeviceProperties");
+    sm = device_prop.major * 10 + device_prop.minor;
+  }
+  sm = ::onnxruntime::llm::kernels::weight_only::get_arch_for_mixed_gemm_weight_preprocess(sm);
+
+  auto permutation_map_buffer = make_cuda_ptr(32 * sizeof(int32_t));
+
+  ::onnxruntime::llm::kernels::weight_only::preprocess_weights_for_mixed_gemm_cuda(
+      stream,
+      sm,
+      preprocessed_weight,
+      packed_transposed_weight,
+      reinterpret_cast<int32_t*>(permutation_map_buffer.get()),
+      {static_cast<size_t>(k), static_cast<size_t>(n)},
+      quant_type);
+
+  ThrowIfCudaError(cudaGetLastError(), "preprocess CUDA kernel launch");
+  ThrowIfCudaError(cudaMemcpyAsync(processed_weights_buf.ptr, preprocessed_weight, packed_weight_bytes,
+                                   cudaMemcpyDeviceToHost, stream),
+                   "cudaMemcpyAsync device-to-host");
+  ThrowIfCudaError(cudaStreamSynchronize(stream), "cudaStreamSynchronize");
+
+  return processed_weights;
+}
+
+}  // namespace
+
+PYBIND11_MODULE(onnxruntime_cuda_quant_preprocess, m) {
+  m.doc() = "CUDA weight-only quantization preprocessing helpers (loaded on demand).";
+  m.def("pack_weights_for_cuda_mixed_gemm", &PackWeightsForMixedGemm,
+        "Pack quantized weights for CUDA mixed-precision GEMM (FpA_IntB format)",
+        py::arg("q_weights"), py::arg("N"), py::arg("K"), py::arg("bits"), py::arg("force_arch") = -1);
+}

onnxruntime/python/onnxruntime_pybind_mlvalue.cc

@@ -23,10 +23,6 @@
 #include "core/framework/kernel_registry.h"
 #include "core/framework/provider_options_utils.h"
 
-#if defined(USE_CUDA) && !defined(ORT_NO_CUDA_IN_PYBIND)
-#include <cuda_runtime_api.h>
-#endif
-
 #ifdef USE_DML
 using Microsoft::WRL::ComPtr;
 
@@ -184,35 +180,14 @@ int32_t GetTensorProtoType(const OrtValue& ort_value) {
 }
 
 #ifdef USE_CUDA
-namespace {
-
-#if !defined(ORT_NO_CUDA_IN_PYBIND)
-void CudaRuntimeMemCpy(void* dst, const void* src, size_t num_bytes, cudaMemcpyKind kind) {
-  const auto copy_result = cudaMemcpy(dst, src, num_bytes, kind);
-  ORT_ENFORCE(copy_result == cudaSuccess, "cudaMemcpy failed: ", cudaGetErrorString(copy_result));
-
-  if (kind == cudaMemcpyHostToDevice) {
-    // Match ProviderInfo_CUDA::cudaMemcpy_HostToDevice: cudaMemcpy() uses the default
-    // stream, and pageable host-to-device copies can return before DMA to device is done.
-    const auto sync_result = cudaStreamSynchronize(0);
-    ORT_ENFORCE(sync_result == cudaSuccess, "cudaStreamSynchronize failed: ", cudaGetErrorString(sync_result));
-  }
-}
-#endif
-
-}  // namespace
 
 void CpuToCudaMemCpy(void* dst, const void* src, size_t num_bytes) {
   if (TryGetProviderInfo_CUDA() != nullptr) {
     GetProviderInfo_CUDA().cudaMemcpy_HostToDevice(dst, src, num_bytes);
     return;
   }
 
-#if !defined(ORT_NO_CUDA_IN_PYBIND)
-  CudaRuntimeMemCpy(dst, src, num_bytes, cudaMemcpyHostToDevice);
-#else
   ORT_THROW("CUDA provider interface is not available for host-to-device copy.");
-#endif
 }
 
 void CudaToCpuMemCpy(void* dst, const void* src, size_t num_bytes) {
@@ -221,11 +196,7 @@ void CudaToCpuMemCpy(void* dst, const void* src, size_t num_bytes) {
     return;
   }
 
-#if !defined(ORT_NO_CUDA_IN_PYBIND)
-  CudaRuntimeMemCpy(dst, src, num_bytes, cudaMemcpyDeviceToHost);
-#else
   ORT_THROW("CUDA provider interface is not available for device-to-host copy.");
-#endif
 }
 
 const std::unordered_map<OrtDevice, MemCpyFunc>* GetCudaToHostMemCpyFunction(const OrtDevice& device) {