From 123044bde9ea7838930eed0db184893ca041174b Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 21 Sep 2026 12:08:58 +0000 Subject: [PATCH 01/17] Add GPUDirect Storage model loading Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- docs/CUDA_GPU_Direct_Storage.md | 79 +++++++ docs/FAQ.md | 5 + .../providers/cuda/cuda_provider_options.h | 3 +- .../providers/cuda/cuda_execution_provider.cc | 8 +- .../cuda/cuda_execution_provider_info.cc | 8 + .../cuda/cuda_execution_provider_info.h | 4 +- .../cuda/cuda_external_data_loader.cc | 37 +++- .../cuda/cuda_external_data_loader.h | 9 +- .../cuda/cuda_external_data_loader_gds.cc | 193 ++++++++++++++++++ .../cuda/cuda_external_data_loader_gds.h | 32 +++ .../providers/cuda/cuda_provider_factory.cc | 6 + .../core/session/provider_bridge_ort.cc | 6 + .../cuda_external_data_loader_test.cc | 146 +++++++++++++ 13 files changed, 529 insertions(+), 7 deletions(-) create mode 100644 docs/CUDA_GPU_Direct_Storage.md create mode 100644 onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc create mode 100644 onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md new file mode 100644 index 0000000000000..723e4a9910434 --- /dev/null +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -0,0 +1,79 @@ +# CUDA external-data loading with GPUDirect Storage + +The CUDA execution provider can use NVIDIA GPUDirect Storage (GDS) to load ONNX external initializers without +staging file data in CPU memory. GDS is opt-in. If it cannot be initialized or cannot read an external-data file, +ONNX Runtime logs a warning and uses the existing pinned-host-buffer loader for the rest of the session. + +## Data path + +With GDS enabled, ONNX Runtime opens each external-data file with `O_DIRECT` and uses `libcufile` to read 64 MiB +blocks into a reusable, cuFile-registered CUDA buffer. Each block is then copied device-to-device into the +initializer allocation owned by the CUDA arena: + +```text +external-data file -> registered CUDA staging buffer -> CUDA arena initializer + cuFileRead device-to-device copy +``` + +The reusable staging buffer bounds additional GPU memory to 64 MiB per CUDA external-data loader. Each +device-to-device copy completes before that buffer is reused. String and Boolean initializers retain the existing +loading path because they require host-side conversion. + +GDS requires: + +- Linux and a CUDA toolkit that provides `cufile.h`; +- `libcufile.so` at runtime; +- a working GDS driver and supported storage/filesystem configuration; and +- external weights stored in a file that can be opened with `O_DIRECT`. + +ONNX Runtime loads `libcufile` dynamically, so enabling the option does not add a mandatory runtime dependency for +users who keep GDS disabled. + +## Configuration + +The following CUDA execution provider options control the primary and fallback paths: + +| Option | Values | Default | Purpose | +|---|---|---:|---| +| `external_data_loader_use_gds` | `0` or `1` | `0` | Try GDS before another external-data loading path | +| `external_data_loader_reading_threads` | `0` to `64` | `4` | Configure the pinned-buffer fallback; `0` disables it | + +Keep `external_data_loader_reading_threads` greater than zero when enabling GDS to retain the pinned-buffer backup. +`1` uses synchronous reads into pinned memory. Values from `2` through `64` use that many parallel read tasks per +64 MiB pinned buffer. If the value is `0` and GDS is unavailable, ONNX Runtime falls back to pageable host memory. + +Python: + +```python +import onnxruntime as ort + +providers = [ + ( + "CUDAExecutionProvider", + { + "external_data_loader_use_gds": "1", + "external_data_loader_reading_threads": "4", + }, + ), + "CPUExecutionProvider", +] + +session = ort.InferenceSession("model.onnx", providers=providers) +``` + +C++: + +```cpp +Ort::SessionOptions session_options; +Ort::CUDAProviderOptions cuda_options; +cuda_options.Update({ + {"external_data_loader_use_gds", "1"}, + {"external_data_loader_reading_threads", "4"}, +}); +session_options.AppendExecutionProvider_CUDA_V2(cuda_options); + +Ort::Session session(env, ORT_TSTR("model.onnx"), session_options); +``` + +The option only affects initializers stored as ONNX external data and assigned to CUDA memory. Embedded initializers +and initializers assigned to other execution providers retain their existing paths. diff --git a/docs/FAQ.md b/docs/FAQ.md index 70bedbd02e944..97511d3410eb7 100644 --- a/docs/FAQ.md +++ b/docs/FAQ.md @@ -4,6 +4,11 @@ Here are some commonly raised questions from users of ONNX Runtime and brought u ## Do the GPU builds support quantized models? The default CUDA build supports 3 standard quantization operators: QuantizeLinear, DequantizeLinear, and MatMulInteger. The TensorRT EP has limited support for INT8 quantized ops. In general, support of quantized models through ORT is continuing to expand on a model-driven basis. For performance improvements, quantization is not always required, and we suggest trying alternative strategies to [performance tune](https://onnxruntime.ai/docs/performance/tune-performance/) before determining that quantization is necessary. +## Can CUDA load model weights with GPUDirect Storage? + +Yes. See [CUDA external-data loading with GPUDirect Storage](CUDA_GPU_Direct_Storage.md) for prerequisites, +configuration, and the pinned-buffer fallback. + ## How do I change the severity level of the default logger to something other than the default (WARNING)? Setting the severity level to VERBOSE is most useful when debugging errors. diff --git a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h index 0afbb13739205..00a40399faace 100644 --- a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h +++ b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h @@ -42,5 +42,6 @@ struct OrtCUDAProviderOptionsV2 { int use_tf32 = 1; // use TF32 int fuse_conv_bias = 0; // Enable CUDNN Frontend kernel fusing, results in JIT compiles int sdpa_kernel = 0; // Scaled Dot Product Attention kernel option - size_t external_data_loader_reading_threads = 4; // Number of CPU read tasks per external-data staging buffer. 0 disables the loader; 1 disables parallel reads. + size_t external_data_loader_reading_threads = 4; // Number of CPU read tasks per external-data staging buffer. 0 disables pinned-buffer fallback; 1 disables parallel reads. + int external_data_loader_use_gds = 0; // Try GPUDirect Storage before falling back to the pinned-buffer loader. }; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index e6645a0d4cb77..454f319e41554 100755 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -3433,12 +3433,16 @@ std::unique_ptr CUDAExecutionProvider::GetDataTransf } std::unique_ptr CUDAExecutionProvider::GetExternalDataLoader() const { - if (info_.external_data_loader_reading_threads == 0) { + if (info_.external_data_loader_reading_threads == 0 && + !info_.external_data_loader_use_gds) { return nullptr; } return std::make_unique( - info_.device_id, info_.external_data_loader_reading_threads); + info_.device_id, info_.external_data_loader_reading_threads, + static_cast(cudaMallocHost), + static_cast(cudaStreamCreateWithFlags), + info_.external_data_loader_use_gds); } std::vector> diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc index 14899aee419ec..5dedcad1b711e 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc @@ -39,6 +39,7 @@ constexpr const char* kUseTF32 = "use_tf32"; constexpr const char* kFuseConvBias = "fuse_conv_bias"; constexpr const char* kSdpaKernel = "sdpa_kernel"; constexpr const char* kExternalDataLoaderReadingThreads = "external_data_loader_reading_threads"; +constexpr const char* kExternalDataLoaderUseGds = "external_data_loader_use_gds"; } // namespace provider_option_names } // namespace cuda @@ -146,6 +147,9 @@ CUDAExecutionProviderInfo CUDAExecutionProviderInfo::FromProviderOptions(const P OrtCUDAProviderOptionsV2::kMaxExternalDataLoaderReadingThreadCount, "."); return Status::OK(); }) + .AddAssignmentToReference( + cuda::provider_option_names::kExternalDataLoaderUseGds, + info.external_data_loader_use_gds) .AddValueParser( cuda::provider_option_names::kTunableOpEnable, [&info](const std::string& value_str) -> Status { @@ -203,6 +207,8 @@ ProviderOptions CUDAExecutionProviderInfo::ToProviderOptions(const CUDAExecution {cuda::provider_option_names::kFuseConvBias, MakeStringWithClassicLocale(info.fuse_conv_bias)}, {cuda::provider_option_names::kExternalDataLoaderReadingThreads, MakeStringWithClassicLocale(info.external_data_loader_reading_threads)}, + {cuda::provider_option_names::kExternalDataLoaderUseGds, + MakeStringWithClassicLocale(info.external_data_loader_use_gds)}, }; return options; @@ -230,6 +236,8 @@ ProviderOptions CUDAExecutionProviderInfo::ToProviderOptions(const OrtCUDAProvid {cuda::provider_option_names::kSdpaKernel, MakeStringWithClassicLocale(info.sdpa_kernel)}, {cuda::provider_option_names::kExternalDataLoaderReadingThreads, MakeStringWithClassicLocale(info.external_data_loader_reading_threads)}, + {cuda::provider_option_names::kExternalDataLoaderUseGds, + MakeStringWithClassicLocale(info.external_data_loader_use_gds)}, }; return options; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h index b4d4cb758d2f9..796217a2e7b36 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h @@ -84,9 +84,10 @@ struct CUDAExecutionProviderInfo { int sdpa_kernel{0}; - // 0 disables the custom external-data loader and retains the framework's existing path. + // 0 disables the pinned-buffer loader and uses the pageable fallback if GDS is unavailable. // 1 uses the pinned-buffer loader with synchronous reads. 2..64 use that many parallel read tasks per block. size_t external_data_loader_reading_threads{4}; + bool external_data_loader_use_gds{false}; static CUDAExecutionProviderInfo FromProviderOptions(const ProviderOptions& options); static ProviderOptions ToProviderOptions(const CUDAExecutionProviderInfo& info); @@ -122,6 +123,7 @@ struct std::hash<::onnxruntime::CUDAExecutionProviderInfo> { onnxruntime::HashCombine(info.sdpa_kernel, value); onnxruntime::HashCombine(info.enable_cudnn, value); onnxruntime::HashCombine(info.external_data_loader_reading_threads, value); + onnxruntime::HashCombine(info.external_data_loader_use_gds, value); // Memory pointers onnxruntime::HashCombine(reinterpret_cast(info.user_compute_stream), value); diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc index a587c252c3b34..2d27d07788a8c 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc @@ -118,11 +118,15 @@ common::Status LoadWithPageableBuffer(const RandomAccessFile& file, FileOffsetTy ExternalDataLoader::ExternalDataLoader(int device_id, size_t reading_thread_count, AllocatePinnedBufferFn allocate_pinned_buffer, - CreateStreamFn create_stream) + CreateStreamFn create_stream, + bool use_gds, + GdsLoader::CreateFn create_gds_loader) : device_id_(device_id), reading_thread_count_(reading_thread_count), allocate_pinned_buffer_(allocate_pinned_buffer), - create_stream_(create_stream) {} + create_stream_(create_stream), + use_gds_(use_gds), + create_gds_loader_(create_gds_loader) {} ExternalDataLoader::~ExternalDataLoader() { reader_pool_.reset(); @@ -166,6 +170,8 @@ void ExternalDataLoader::ReleaseResources() const noexcept { previous_device != device_id_ && cudaSetDevice(device_id_) == cudaSuccess; + gds_loader_.reset(); + for (auto& stream : streams_) { if (stream != nullptr) { ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaStreamSynchronize(stream))); @@ -210,6 +216,33 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, std::lock_guard lock(mutex_); CudaDeviceGuard device_guard; ORT_RETURN_IF_ERROR(device_guard.SetDevice(device_id_)); + + if (use_gds_ && !gds_disabled_ && + std::endian::native == std::endian::little && + !tensor.IsDataType()) { + Status gds_status = Status::OK(); + if (!gds_loader_) { + gds_status = create_gds_loader_(device_id_, gds_loader_); + } + if (gds_status.IsOK()) { + gds_status = gds_loader_->Load(data_file_path, data_offset, length, tensor); + } + if (gds_status.IsOK()) { + return Status::OK(); + } + + gds_disabled_ = true; + gds_loader_.reset(); + LOGS_DEFAULT(WARNING) << "GPUDirect Storage could not load external data; falling back to the CUDA " + << (reading_thread_count_ == 0 ? "pageable-buffer" : "pinned-buffer") + << " loader. " + << gds_status.ErrorMessage(); + } + + if (reading_thread_count_ == 0) { + return LoadWithPageableBuffer(*file, data_offset, length, tensor, 1, reader_pool_); + } + const auto resource_status = EnsureResources(); if (!resource_status.IsOK()) { // TODO: Remember setup failures during initialization and report the first CUDA error diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader.h index bcd795021a71a..31dbf9a5f4a93 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.h +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.h @@ -8,6 +8,7 @@ #include #include "core/framework/external_data_loader.h" +#include "core/providers/cuda/cuda_external_data_loader_gds.h" #include "cuda_pch.h" namespace onnxruntime { @@ -69,7 +70,9 @@ class ExternalDataLoader final : public IExternalDataLoader { ExternalDataLoader(int device_id, size_t reading_thread_count, AllocatePinnedBufferFn allocate_pinned_buffer = cudaMallocHost, - CreateStreamFn create_stream = cudaStreamCreateWithFlags); + CreateStreamFn create_stream = cudaStreamCreateWithFlags, + bool use_gds = false, + GdsLoader::CreateFn create_gds_loader = GdsLoader::Create); ~ExternalDataLoader() override; bool CanLoad(const OrtMemoryInfo& target_memory_info) const override; @@ -91,6 +94,10 @@ class ExternalDataLoader final : public IExternalDataLoader { const size_t reading_thread_count_; const AllocatePinnedBufferFn allocate_pinned_buffer_; const CreateStreamFn create_stream_; + const bool use_gds_; + const GdsLoader::CreateFn create_gds_loader_; + mutable bool gds_disabled_{false}; + mutable std::unique_ptr gds_loader_; mutable std::unique_ptr reader_pool_; }; diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc new file mode 100644 index 0000000000000..76a8ee7421c1a --- /dev/null +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -0,0 +1,193 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +// provider_api.h must be first to set SHARED_PROVIDER +#include "core/providers/shared_library/provider_api.h" + +#include "core/providers/cuda/cuda_external_data_loader_gds.h" + +#include +#include +#include +#include +#include +#include + +#include "core/common/common.h" +#include "core/common/safeint.h" +#include "core/providers/cuda/cuda_common.h" + +#if defined(__linux__) && __has_include() +#define ORT_CUDA_GDS_AVAILABLE 1 +#include +#include +#include +#include +#endif + +namespace onnxruntime { +namespace cuda { +namespace { + +constexpr size_t kGdsBufferSize = 64 * 1024 * 1024; + +#if defined(ORT_CUDA_GDS_AVAILABLE) + +template +common::Status LoadSymbol(void* library, const char* name, T& function) { + dlerror(); + function = reinterpret_cast(dlsym(library, name)); + const char* error = dlerror(); + ORT_RETURN_IF(function == nullptr || error != nullptr, + "Failed to load ", name, " from libcufile: ", + error == nullptr ? "symbol not found" : error); + return Status::OK(); +} + +common::Status CheckCuFileStatus(CUfileError_t status, std::string_view operation) { + ORT_RETURN_IF(status.err != CU_FILE_SUCCESS, operation, " failed: ", + cufileop_status_error(status.err), " (", static_cast(status.err), ")"); + return Status::OK(); +} + +class LinuxGdsLoader final : public GdsLoader { + public: + using DriverOpenFn = decltype(&cuFileDriverOpen); + using DriverCloseFn = CUfileError_t (*)(); + using HandleRegisterFn = decltype(&cuFileHandleRegister); + using HandleDeregisterFn = decltype(&cuFileHandleDeregister); + using BufferRegisterFn = decltype(&cuFileBufRegister); + using BufferDeregisterFn = decltype(&cuFileBufDeregister); + using ReadFn = decltype(&cuFileRead); + + ~LinuxGdsLoader() override { + if (gds_buffer_registered_) { + ORT_IGNORE_RETURN_VALUE(buffer_deregister_(gds_buffer_)); + } + if (gds_buffer_ != nullptr) { + ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaFree(gds_buffer_))); + } + if (driver_initialized_) { + ORT_IGNORE_RETURN_VALUE(driver_close_()); + } + if (library_ != nullptr) { + ORT_IGNORE_RETURN_VALUE(dlclose(library_)); + } + } + + static common::Status Create(int device_id, std::unique_ptr& loader) { + auto candidate = std::unique_ptr(new LinuxGdsLoader()); + ORT_RETURN_IF_ERROR(candidate->Initialize(device_id)); + loader = std::move(candidate); + return Status::OK(); + } + + common::Status Load(const std::filesystem::path& data_file_path, + int64_t data_offset, + size_t data_length, + Tensor& tensor) const override { + const int file_descriptor = open(data_file_path.c_str(), O_RDONLY | O_DIRECT); + ORT_RETURN_IF(file_descriptor < 0, "Failed to open external data for GPUDirect Storage: ", + std::strerror(errno)); + auto close_file = gsl::finally([file_descriptor]() { ORT_IGNORE_RETURN_VALUE(close(file_descriptor)); }); + + CUfileDescr_t descriptor{}; + descriptor.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; + descriptor.handle.fd = file_descriptor; + CUfileHandle_t file_handle = nullptr; + ORT_RETURN_IF_ERROR(CheckCuFileStatus(handle_register_(&file_handle, &descriptor), + "cuFileHandleRegister")); + auto deregister_file = gsl::finally([&]() { handle_deregister_(file_handle); }); + + auto* destination = static_cast(tensor.MutableDataRaw()); + for (size_t offset = 0; offset < data_length;) { + const size_t chunk_size = std::min(kGdsBufferSize, data_length - offset); + const auto file_offset = SafeInt(data_offset) + offset; + const ssize_t bytes_read = read_(file_handle, gds_buffer_, chunk_size, file_offset, 0); + if (bytes_read != static_cast(chunk_size)) { + if (bytes_read == -1) { + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", std::strerror(errno)); + } + if (bytes_read < 0) { + const auto cu_file_error = static_cast(-bytes_read); + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", + cufileop_status_error(cu_file_error), + " (", static_cast(cu_file_error), ")"); + } + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead returned ", bytes_read, + " bytes; expected ", chunk_size, "."); + } + + CUDA_RETURN_IF_ERROR( + cudaMemcpy(destination + offset, gds_buffer_, chunk_size, cudaMemcpyDeviceToDevice)); + CUDA_RETURN_IF_ERROR(cudaStreamSynchronize(nullptr)); + offset += chunk_size; + } + + return Status::OK(); + } + + private: + common::Status Initialize(int device_id) { + library_ = dlopen("libcufile.so.0", RTLD_NOW | RTLD_LOCAL); + if (library_ == nullptr) { + library_ = dlopen("libcufile.so", RTLD_NOW | RTLD_LOCAL); + } + const char* library_error = dlerror(); + ORT_RETURN_IF(library_ == nullptr, "GPUDirect Storage is unavailable: ", + library_error == nullptr ? "libcufile could not be loaded" : library_error); + + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileDriverOpen", driver_open_)); + auto close_status = LoadSymbol(library_, "cuFileDriverClose_v2", driver_close_); + if (!close_status.IsOK()) { + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileDriverClose", driver_close_)); + } + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileHandleRegister", handle_register_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileHandleDeregister", handle_deregister_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileBufRegister", buffer_register_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileBufDeregister", buffer_deregister_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileRead", read_)); + + ORT_RETURN_IF_ERROR(CheckCuFileStatus(driver_open_(), "cuFileDriverOpen")); + driver_initialized_ = true; + + CUDA_RETURN_IF_ERROR(cudaSetDevice(device_id)); + CUDA_RETURN_IF_ERROR(cudaMalloc(&gds_buffer_, kGdsBufferSize)); + ORT_RETURN_IF_ERROR(CheckCuFileStatus( + buffer_register_(gds_buffer_, kGdsBufferSize, 0), "cuFileBufRegister")); + gds_buffer_registered_ = true; + return Status::OK(); + } + + LinuxGdsLoader() = default; + + void* library_{nullptr}; + void* gds_buffer_{nullptr}; + bool driver_initialized_{false}; + bool gds_buffer_registered_{false}; + DriverOpenFn driver_open_{nullptr}; + DriverCloseFn driver_close_{nullptr}; + HandleRegisterFn handle_register_{nullptr}; + HandleDeregisterFn handle_deregister_{nullptr}; + BufferRegisterFn buffer_register_{nullptr}; + BufferDeregisterFn buffer_deregister_{nullptr}; + ReadFn read_{nullptr}; +}; + +#endif + +} // namespace + +common::Status GdsLoader::Create(int device_id, std::unique_ptr& loader) { +#if defined(ORT_CUDA_GDS_AVAILABLE) + return LinuxGdsLoader::Create(device_id, loader); +#else + ORT_UNUSED_PARAMETER(device_id); + ORT_UNUSED_PARAMETER(loader); + return ORT_MAKE_STATUS(ONNXRUNTIME, NOT_IMPLEMENTED, + "GPUDirect Storage requires Linux and a CUDA toolkit with cuFile headers."); +#endif +} + +} // namespace cuda +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h new file mode 100644 index 0000000000000..4b18552232b33 --- /dev/null +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h @@ -0,0 +1,32 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include +#include + +#include "core/common/status.h" + +namespace onnxruntime { +class Tensor; + +namespace cuda { + +class GdsLoader { + public: + using CreateFn = common::Status (*)(int device_id, std::unique_ptr& loader); + + virtual ~GdsLoader() = default; + + virtual common::Status Load(const std::filesystem::path& data_file_path, + int64_t data_offset, + size_t data_length, + Tensor& tensor) const = 0; + + static common::Status Create(int device_id, std::unique_ptr& loader); +}; + +} // namespace cuda +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc index 00519584950b8..bc216a792a481 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc @@ -210,6 +210,9 @@ struct CUDA_Provider : Provider { OrtCUDAProviderOptionsV2::kMaxExternalDataLoaderReadingThreadCount, "external_data_loader_reading_threads must be between 0 and ", OrtCUDAProviderOptionsV2::kMaxExternalDataLoaderReadingThreadCount, "."); + ORT_ENFORCE(params->external_data_loader_use_gds == 0 || + params->external_data_loader_use_gds == 1, + "external_data_loader_use_gds must be 0 or 1."); // Calling a function like ::cudaDeviceSynchronize will cause CUDA to ensure there is binary code for the current GPU architecture // Ideally this will be already part of the binary, but if not, CUDA will JIT it during this call. This can take a very long time @@ -251,6 +254,7 @@ struct CUDA_Provider : Provider { info.use_tf32 = params->use_tf32 != 0; info.sdpa_kernel = params->sdpa_kernel; info.external_data_loader_reading_threads = params->external_data_loader_reading_threads; + info.external_data_loader_use_gds = params->external_data_loader_use_gds != 0; return std::make_shared(info); } @@ -288,6 +292,8 @@ struct CUDA_Provider : Provider { cuda_options.fuse_conv_bias = internal_options.fuse_conv_bias; cuda_options.external_data_loader_reading_threads = internal_options.external_data_loader_reading_threads; + cuda_options.external_data_loader_use_gds = + internal_options.external_data_loader_use_gds; } ProviderOptions GetProviderOptions(const void* provider_options) override { diff --git a/onnxruntime/core/session/provider_bridge_ort.cc b/onnxruntime/core/session/provider_bridge_ort.cc index 8ece3b74db06a..7db658a83699a 100644 --- a/onnxruntime/core/session/provider_bridge_ort.cc +++ b/onnxruntime/core/session/provider_bridge_ort.cc @@ -2958,6 +2958,12 @@ ORT_API_STATUS_IMPL(OrtApis::SessionOptionsAppendExecutionProvider_CUDA_V2, _In_ ORT_INVALID_ARGUMENT, message.c_str()); } + if (cuda_options->external_data_loader_use_gds != 0 && + cuda_options->external_data_loader_use_gds != 1) { + return OrtApis::CreateStatus( + ORT_INVALID_ARGUMENT, + "external_data_loader_use_gds must be 0 or 1."); + } auto factory = onnxruntime::CudaProviderFactoryCreator::Create(cuda_options); if (!factory) { diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc index 9c688e2788561..e8b096476c02d 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc @@ -193,6 +193,17 @@ TEST(CudaExternalDataLoaderTest, DisablesLoaderWhenConfiguredWithZeroReadingThre EXPECT_EQ(execution_provider->GetExternalDataLoader(), nullptr); } +TEST(CudaExternalDataLoaderTest, EnablesLoaderForGdsWithZeroFallbackReadingThreads) { + OrtCUDAProviderOptionsV2 provider_options{}; + provider_options.do_copy_in_default_stream = true; + provider_options.use_tf32 = false; + provider_options.external_data_loader_reading_threads = 0; + provider_options.external_data_loader_use_gds = 1; + auto execution_provider = CudaExecutionProviderWithOptions(&provider_options); + ASSERT_NE(execution_provider, nullptr); + EXPECT_NE(execution_provider->GetExternalDataLoader(), nullptr); +} + TEST(CudaExternalDataLoaderTest, RejectsTooManyReadingThreadsFromStructOptions) { OrtCUDAProviderOptionsV2 provider_options{}; provider_options.external_data_loader_reading_threads = @@ -213,6 +224,20 @@ TEST(CudaExternalDataLoaderTest, RejectsTooManyReadingThreadsFromStructOptions) EXPECT_EQ(CudaExecutionProviderWithOptions(&provider_options), nullptr); } +TEST(CudaExternalDataLoaderTest, RejectsInvalidGdsOptionFromStructOptions) { + OrtCUDAProviderOptionsV2 provider_options{}; + provider_options.external_data_loader_use_gds = 2; + Ort::SessionOptions session_options; + try { + session_options.AppendExecutionProvider_CUDA_V2(provider_options); + FAIL() << "Expected an invalid external_data_loader_use_gds value to be rejected."; + } catch (const Ort::Exception& ex) { + EXPECT_THAT(ex.what(), testing::HasSubstr("external_data_loader_use_gds")); + } + + EXPECT_EQ(CudaExecutionProviderWithOptions(&provider_options), nullptr); +} + TEST(CudaExternalDataLoaderTest, ReusesAlternatingBuffersAcrossRepeatedLoads) { VerifyLoad(2 * cuda::kExternalDataLoaderBufferSize + 1, 2); } @@ -225,6 +250,127 @@ cudaError_t FailStreamCreation(cudaStream_t*, unsigned int) { return cudaErrorInitializationError; } +class TestGdsLoader final : public cuda::GdsLoader { + public: + explicit TestGdsLoader(uint8_t value) : value_(value) {} + + Status Load(const std::filesystem::path&, int64_t, size_t data_length, + Tensor& tensor) const override { + const auto result = cudaMemset(tensor.MutableDataRaw(), value_, data_length); + ORT_RETURN_IF(result != cudaSuccess, "cudaMemset failed: ", cudaGetErrorString(result)); + return Status::OK(); + } + + private: + uint8_t value_; +}; + +Status CreateTestGdsLoader(int, std::unique_ptr& loader) { + loader = std::make_unique(0x5a); + return Status::OK(); +} + +Status FailTestGdsLoaderCreation(int, std::unique_ptr&) { + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "GDS unavailable for test"); +} + +size_t gds_create_attempt_count = 0; + +Status CountAndFailTestGdsLoaderCreation(int, std::unique_ptr&) { + ++gds_create_attempt_count; + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "GDS unavailable for test"); +} + +TEST(CudaExternalDataLoaderTest, UsesGdsWithoutAllocatingPinnedBuffers) { + constexpr size_t kLength = 1024; + PathString path; + CreateExternalDataFile(kLength, path); + ScopedFileDeleter file_deleter{path}; + + auto execution_provider = DefaultCudaExecutionProvider(); + ASSERT_NE(execution_provider, nullptr); + auto allocators = execution_provider->CreatePreferredAllocators(); + const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { + return candidate->Info().device.Type() == OrtDevice::GPU && + candidate->Info().mem_type == OrtMemTypeDefault; + }); + ASSERT_NE(allocator, allocators.end()); + Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); + + cuda::ExternalDataLoader loader( + 0, 4, FailPinnedBufferAllocation, + static_cast(cudaStreamCreateWithFlags), + true, CreateTestGdsLoader); + ASSERT_STATUS_OK(loader.LoadTensor( + Env::Default(), path, kFilePrefixSize, kLength, tensor)); + + std::array output{}; + ASSERT_EQ(cudaSuccess, cudaMemcpy( + output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); + EXPECT_TRUE(std::all_of(output.begin(), output.end(), [](uint8_t value) { return value == 0x5a; })); +} + +TEST(CudaExternalDataLoaderTest, FallsBackToPinnedBuffersWhenGdsIsUnavailable) { + constexpr size_t kLength = 1024; + PathString path; + CreateExternalDataFile(kLength, path); + ScopedFileDeleter file_deleter{path}; + + auto execution_provider = DefaultCudaExecutionProvider(); + ASSERT_NE(execution_provider, nullptr); + auto allocators = execution_provider->CreatePreferredAllocators(); + const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { + return candidate->Info().device.Type() == OrtDevice::GPU && + candidate->Info().mem_type == OrtMemTypeDefault; + }); + ASSERT_NE(allocator, allocators.end()); + Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); + + cuda::ExternalDataLoader loader( + 0, 4, + static_cast(cudaMallocHost), + static_cast(cudaStreamCreateWithFlags), + true, FailTestGdsLoaderCreation); + ASSERT_STATUS_OK(loader.LoadTensor( + Env::Default(), path, kFilePrefixSize, kLength, tensor)); + + std::array output{}; + ASSERT_EQ(cudaSuccess, cudaMemcpy( + output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); ++i) { + ASSERT_EQ(TestValue(i), output[i]) << "Mismatch at byte " << i; + } +} + +TEST(CudaExternalDataLoaderTest, DoesNotRetryGdsAfterFailure) { + constexpr size_t kLength = 1024; + PathString path; + CreateExternalDataFile(kLength, path); + ScopedFileDeleter file_deleter{path}; + + auto execution_provider = DefaultCudaExecutionProvider(); + ASSERT_NE(execution_provider, nullptr); + auto allocators = execution_provider->CreatePreferredAllocators(); + const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { + return candidate->Info().device.Type() == OrtDevice::GPU && + candidate->Info().mem_type == OrtMemTypeDefault; + }); + ASSERT_NE(allocator, allocators.end()); + Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); + + gds_create_attempt_count = 0; + cuda::ExternalDataLoader loader( + 0, 4, + static_cast(cudaMallocHost), + static_cast(cudaStreamCreateWithFlags), + true, CountAndFailTestGdsLoaderCreation); + ASSERT_STATUS_OK(loader.LoadTensor( + Env::Default(), path, kFilePrefixSize, kLength, tensor)); + ASSERT_STATUS_OK(loader.LoadTensor( + Env::Default(), path, kFilePrefixSize, kLength, tensor)); + EXPECT_EQ(gds_create_attempt_count, 1U); +} + TEST(CudaExternalDataLoaderTest, NormalizesBoolWithPinnedAndPageableFallback) { const std::array input{0, 1, 2, 255}; const std::array expected{0, 1, 1, 1}; From 606d6a20ceff84a8a19db47b527c81363e3855cb Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 21 Sep 2026 13:43:50 +0000 Subject: [PATCH 02/17] Require native GPUDirect Storage support Disable cuFile compatibility mode so unsupported systems use the configured ONNX Runtime pinned-buffer fallback instead of cuFile's internal POSIX path. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- docs/CUDA_GPU_Direct_Storage.md | 3 ++- .../core/providers/cuda/cuda_external_data_loader_gds.cc | 6 ++++++ 2 files changed, 8 insertions(+), 1 deletion(-) diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index 723e4a9910434..512a945fc5a8d 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -27,7 +27,8 @@ GDS requires: - external weights stored in a file that can be opened with `O_DIRECT`. ONNX Runtime loads `libcufile` dynamically, so enabling the option does not add a mandatory runtime dependency for -users who keep GDS disabled. +users who keep GDS disabled. It also disables cuFile compatibility mode: if the storage stack cannot provide the +direct GDS path, ONNX Runtime uses its configured pinned-buffer fallback instead of cuFile's internal POSIX fallback. ## Configuration diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index 76a8ee7421c1a..2686a6d35d0a5 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -58,6 +58,7 @@ class LinuxGdsLoader final : public GdsLoader { using HandleDeregisterFn = decltype(&cuFileHandleDeregister); using BufferRegisterFn = decltype(&cuFileBufRegister); using BufferDeregisterFn = decltype(&cuFileBufDeregister); + using SetBoolParameterFn = decltype(&cuFileSetParameterBool); using ReadFn = decltype(&cuFileRead); ~LinuxGdsLoader() override { @@ -146,8 +147,12 @@ class LinuxGdsLoader final : public GdsLoader { ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileHandleDeregister", handle_deregister_)); ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileBufRegister", buffer_register_)); ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileBufDeregister", buffer_deregister_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileSetParameterBool", set_bool_parameter_)); ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileRead", read_)); + ORT_RETURN_IF_ERROR(CheckCuFileStatus( + set_bool_parameter_(CUFILE_PARAM_PROPERTIES_ALLOW_COMPAT_MODE, false), + "Disabling cuFile compatibility mode")); ORT_RETURN_IF_ERROR(CheckCuFileStatus(driver_open_(), "cuFileDriverOpen")); driver_initialized_ = true; @@ -171,6 +176,7 @@ class LinuxGdsLoader final : public GdsLoader { HandleDeregisterFn handle_deregister_{nullptr}; BufferRegisterFn buffer_register_{nullptr}; BufferDeregisterFn buffer_deregister_{nullptr}; + SetBoolParameterFn set_bool_parameter_{nullptr}; ReadFn read_{nullptr}; }; From de484960a9917a635401e80c6cb12bdc801f496f Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 21 Sep 2026 14:09:04 +0000 Subject: [PATCH 03/17] Try PCI P2PDMA for native GDS Request the driverless PCI P2PDMA path before opening cuFile while retaining pinned-buffer fallback when the host kernel driver or PCIe topology is unsupported. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- docs/CUDA_GPU_Direct_Storage.md | 8 +++++--- .../core/providers/cuda/cuda_external_data_loader_gds.cc | 3 +++ 2 files changed, 8 insertions(+), 3 deletions(-) diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index 512a945fc5a8d..56ca1f197f586 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -23,12 +23,14 @@ GDS requires: - Linux and a CUDA toolkit that provides `cufile.h`; - `libcufile.so` at runtime; -- a working GDS driver and supported storage/filesystem configuration; and +- either `nvidia-fs` or a recent open NVIDIA kernel module that supports PCI P2PDMA; +- a supported storage/filesystem and PCIe topology; and - external weights stored in a file that can be opened with `O_DIRECT`. ONNX Runtime loads `libcufile` dynamically, so enabling the option does not add a mandatory runtime dependency for -users who keep GDS disabled. It also disables cuFile compatibility mode: if the storage stack cannot provide the -direct GDS path, ONNX Runtime uses its configured pinned-buffer fallback instead of cuFile's internal POSIX fallback. +users who keep GDS disabled. It requests PCI P2PDMA, which can provide GDS without `nvidia-fs` on supported recent +kernels, GPUs, and storage devices. It also disables cuFile compatibility mode: if the storage stack cannot provide +a native GDS path, ONNX Runtime uses its configured pinned-buffer fallback instead of cuFile's internal POSIX fallback. ## Configuration diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index 2686a6d35d0a5..8929b036fe837 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -150,6 +150,9 @@ class LinuxGdsLoader final : public GdsLoader { ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileSetParameterBool", set_bool_parameter_)); ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileRead", read_)); + ORT_RETURN_IF_ERROR(CheckCuFileStatus( + set_bool_parameter_(CUFILE_PARAM_USE_PCIP2PDMA, true), + "Enabling cuFile PCI P2PDMA")); ORT_RETURN_IF_ERROR(CheckCuFileStatus( set_bool_parameter_(CUFILE_PARAM_PROPERTIES_ALLOW_COMPAT_MODE, false), "Disabling cuFile compatibility mode")); From 72ed26c6405a3d286803243dec53425bc0232b8a Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 21 Sep 2026 15:34:13 +0000 Subject: [PATCH 04/17] Address GDS loader review feedback Reuse validated file descriptors, share the process-wide cuFile driver, route unaligned ranges to fallback, and cover partial GDS failures. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- docs/CUDA_GPU_Direct_Storage.md | 2 + onnxruntime/core/platform/env.h | 4 + onnxruntime/core/platform/posix/env.cc | 4 + .../cuda/cuda_external_data_loader.cc | 7 +- .../cuda/cuda_external_data_loader_gds.cc | 187 +++++++++++++----- .../cuda/cuda_external_data_loader_gds.h | 4 +- .../cuda_external_data_loader_test.cc | 100 ++++++++-- 7 files changed, 235 insertions(+), 73 deletions(-) diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index 56ca1f197f586..0c40a8eaa4006 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -31,6 +31,8 @@ ONNX Runtime loads `libcufile` dynamically, so enabling the option does not add users who keep GDS disabled. It requests PCI P2PDMA, which can provide GDS without `nvidia-fs` on supported recent kernels, GPUs, and storage devices. It also disables cuFile compatibility mode: if the storage stack cannot provide a native GDS path, ONNX Runtime uses its configured pinned-buffer fallback instead of cuFile's internal POSIX fallback. +GDS is attempted only for external-data ranges whose offset and length are both 4 KiB aligned. An unaligned +initializer uses the configured fallback without disabling GDS for later aligned initializers. ## Configuration diff --git a/onnxruntime/core/platform/env.h b/onnxruntime/core/platform/env.h index 8e0f6669a9dbc..5e3be8c9974bf 100644 --- a/onnxruntime/core/platform/env.h +++ b/onnxruntime/core/platform/env.h @@ -130,6 +130,10 @@ class RandomAccessFile { */ virtual common::Status Read(FileOffsetType offset, gsl::span buffer) const = 0; + // Returns the POSIX file descriptor for the owned open file, or -1 when unavailable. + // The descriptor remains owned by this object and is valid only for its lifetime. + virtual int GetFileDescriptor() const { return -1; } + protected: RandomAccessFile() = default; diff --git a/onnxruntime/core/platform/posix/env.cc b/onnxruntime/core/platform/posix/env.cc index 43b2c4b9a73ae..d144abc60aa07 100644 --- a/onnxruntime/core/platform/posix/env.cc +++ b/onnxruntime/core/platform/posix/env.cc @@ -163,6 +163,10 @@ class PosixRandomAccessFile final : public RandomAccessFile { return common::Status::OK(); } + int GetFileDescriptor() const override { + return descriptor_.Get(); + } + private: ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(PosixRandomAccessFile); ScopedFileDescriptor descriptor_; diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc index 2d27d07788a8c..d92353f7c4414 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc @@ -217,7 +217,10 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, CudaDeviceGuard device_guard; ORT_RETURN_IF_ERROR(device_guard.SetDevice(device_id_)); - if (use_gds_ && !gds_disabled_ && + const bool gds_range_is_aligned = + data_offset % static_cast(kGdsIoAlignment) == 0 && + length % kGdsIoAlignment == 0; + if (use_gds_ && !gds_disabled_ && gds_range_is_aligned && std::endian::native == std::endian::little && !tensor.IsDataType()) { Status gds_status = Status::OK(); @@ -225,7 +228,7 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, gds_status = create_gds_loader_(device_id_, gds_loader_); } if (gds_status.IsOK()) { - gds_status = gds_loader_->Load(data_file_path, data_offset, length, tensor); + gds_status = gds_loader_->Load(file->GetFileDescriptor(), data_offset, length, tensor); } if (gds_status.IsOK()) { return Status::OK(); diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index 8929b036fe837..9a6490f06d48c 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -10,6 +10,7 @@ #include #include #include +#include #include #include @@ -50,7 +51,7 @@ common::Status CheckCuFileStatus(CUfileError_t status, std::string_view operatio return Status::OK(); } -class LinuxGdsLoader final : public GdsLoader { +class CuFileDriver { public: using DriverOpenFn = decltype(&cuFileDriverOpen); using DriverCloseFn = CUfileError_t (*)(); @@ -61,12 +62,10 @@ class LinuxGdsLoader final : public GdsLoader { using SetBoolParameterFn = decltype(&cuFileSetParameterBool); using ReadFn = decltype(&cuFileRead); - ~LinuxGdsLoader() override { - if (gds_buffer_registered_) { - ORT_IGNORE_RETURN_VALUE(buffer_deregister_(gds_buffer_)); - } - if (gds_buffer_ != nullptr) { - ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaFree(gds_buffer_))); + ~CuFileDriver() { + std::unique_lock lock(GlobalMutex(), std::defer_lock); + if (registered_as_active_) { + lock.lock(); } if (driver_initialized_) { ORT_IGNORE_RETURN_VALUE(driver_close_()); @@ -76,6 +75,108 @@ class LinuxGdsLoader final : public GdsLoader { } } + static common::Status Acquire(std::shared_ptr& driver) { + static std::weak_ptr active_driver; + std::lock_guard lock(GlobalMutex()); + + driver = active_driver.lock(); + if (driver) { + return Status::OK(); + } + + auto candidate = std::shared_ptr(new CuFileDriver()); + ORT_RETURN_IF_ERROR(candidate->Initialize()); + candidate->registered_as_active_ = true; + active_driver = candidate; + driver = std::move(candidate); + return Status::OK(); + } + + CUfileError_t RegisterHandle(CUfileHandle_t* handle, CUfileDescr_t* descriptor) const { + return handle_register_(handle, descriptor); + } + + void DeregisterHandle(CUfileHandle_t handle) const { + handle_deregister_(handle); + } + + CUfileError_t RegisterBuffer(const void* buffer, size_t length) const { + return buffer_register_(buffer, length, 0); + } + + CUfileError_t DeregisterBuffer(const void* buffer) const { + return buffer_deregister_(buffer); + } + + ssize_t Read(CUfileHandle_t handle, void* buffer, size_t length, + off_t file_offset, off_t buffer_offset) const { + return read_(handle, buffer, length, file_offset, buffer_offset); + } + + private: + static std::mutex& GlobalMutex() { + static std::mutex mutex; + return mutex; + } + + common::Status Initialize() { + library_ = dlopen("libcufile.so.0", RTLD_NOW | RTLD_LOCAL); + if (library_ == nullptr) { + library_ = dlopen("libcufile.so", RTLD_NOW | RTLD_LOCAL); + } + const char* library_error = dlerror(); + ORT_RETURN_IF(library_ == nullptr, "GPUDirect Storage is unavailable: ", + library_error == nullptr ? "libcufile could not be loaded" : library_error); + + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileDriverOpen", driver_open_)); + auto close_status = LoadSymbol(library_, "cuFileDriverClose_v2", driver_close_); + if (!close_status.IsOK()) { + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileDriverClose", driver_close_)); + } + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileHandleRegister", handle_register_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileHandleDeregister", handle_deregister_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileBufRegister", buffer_register_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileBufDeregister", buffer_deregister_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileSetParameterBool", set_bool_parameter_)); + ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileRead", read_)); + + ORT_RETURN_IF_ERROR(CheckCuFileStatus( + set_bool_parameter_(CUFILE_PARAM_USE_PCIP2PDMA, true), + "Enabling cuFile PCI P2PDMA")); + ORT_RETURN_IF_ERROR(CheckCuFileStatus( + set_bool_parameter_(CUFILE_PARAM_PROPERTIES_ALLOW_COMPAT_MODE, false), + "Disabling cuFile compatibility mode")); + ORT_RETURN_IF_ERROR(CheckCuFileStatus(driver_open_(), "cuFileDriverOpen")); + driver_initialized_ = true; + return Status::OK(); + } + + CuFileDriver() = default; + + void* library_{nullptr}; + bool driver_initialized_{false}; + bool registered_as_active_{false}; + DriverOpenFn driver_open_{nullptr}; + DriverCloseFn driver_close_{nullptr}; + HandleRegisterFn handle_register_{nullptr}; + HandleDeregisterFn handle_deregister_{nullptr}; + BufferRegisterFn buffer_register_{nullptr}; + BufferDeregisterFn buffer_deregister_{nullptr}; + SetBoolParameterFn set_bool_parameter_{nullptr}; + ReadFn read_{nullptr}; +}; + +class LinuxGdsLoader final : public GdsLoader { + public: + ~LinuxGdsLoader() override { + if (gds_buffer_registered_) { + ORT_IGNORE_RETURN_VALUE(driver_->DeregisterBuffer(gds_buffer_)); + } + if (gds_buffer_ != nullptr) { + ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaFree(gds_buffer_))); + } + } + static common::Status Create(int device_id, std::unique_ptr& loader) { auto candidate = std::unique_ptr(new LinuxGdsLoader()); ORT_RETURN_IF_ERROR(candidate->Initialize(device_id)); @@ -83,28 +184,42 @@ class LinuxGdsLoader final : public GdsLoader { return Status::OK(); } - common::Status Load(const std::filesystem::path& data_file_path, + common::Status Load(int file_descriptor, int64_t data_offset, size_t data_length, Tensor& tensor) const override { - const int file_descriptor = open(data_file_path.c_str(), O_RDONLY | O_DIRECT); - ORT_RETURN_IF(file_descriptor < 0, "Failed to open external data for GPUDirect Storage: ", + ORT_RETURN_IF(file_descriptor < 0, + "GPUDirect Storage requires an open POSIX file descriptor."); + + const int direct_descriptor = dup(file_descriptor); + ORT_RETURN_IF(direct_descriptor < 0, "Failed to duplicate external-data file descriptor: ", + std::strerror(errno)); + auto close_file = gsl::finally([direct_descriptor]() { + ORT_IGNORE_RETURN_VALUE(close(direct_descriptor)); + }); + + const int original_flags = fcntl(direct_descriptor, F_GETFL); + ORT_RETURN_IF(original_flags < 0, "Failed to query external-data file flags: ", std::strerror(errno)); - auto close_file = gsl::finally([file_descriptor]() { ORT_IGNORE_RETURN_VALUE(close(file_descriptor)); }); + ORT_RETURN_IF(fcntl(direct_descriptor, F_SETFL, original_flags | O_DIRECT) < 0, + "Failed to enable O_DIRECT for GPUDirect Storage: ", std::strerror(errno)); + auto restore_flags = gsl::finally([direct_descriptor, original_flags]() { + ORT_IGNORE_RETURN_VALUE(fcntl(direct_descriptor, F_SETFL, original_flags)); + }); CUfileDescr_t descriptor{}; descriptor.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; - descriptor.handle.fd = file_descriptor; + descriptor.handle.fd = direct_descriptor; CUfileHandle_t file_handle = nullptr; - ORT_RETURN_IF_ERROR(CheckCuFileStatus(handle_register_(&file_handle, &descriptor), + ORT_RETURN_IF_ERROR(CheckCuFileStatus(driver_->RegisterHandle(&file_handle, &descriptor), "cuFileHandleRegister")); - auto deregister_file = gsl::finally([&]() { handle_deregister_(file_handle); }); + auto deregister_file = gsl::finally([&]() { driver_->DeregisterHandle(file_handle); }); auto* destination = static_cast(tensor.MutableDataRaw()); for (size_t offset = 0; offset < data_length;) { const size_t chunk_size = std::min(kGdsBufferSize, data_length - offset); const auto file_offset = SafeInt(data_offset) + offset; - const ssize_t bytes_read = read_(file_handle, gds_buffer_, chunk_size, file_offset, 0); + const ssize_t bytes_read = driver_->Read(file_handle, gds_buffer_, chunk_size, file_offset, 0); if (bytes_read != static_cast(chunk_size)) { if (bytes_read == -1) { return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", std::strerror(errno)); @@ -130,57 +245,21 @@ class LinuxGdsLoader final : public GdsLoader { private: common::Status Initialize(int device_id) { - library_ = dlopen("libcufile.so.0", RTLD_NOW | RTLD_LOCAL); - if (library_ == nullptr) { - library_ = dlopen("libcufile.so", RTLD_NOW | RTLD_LOCAL); - } - const char* library_error = dlerror(); - ORT_RETURN_IF(library_ == nullptr, "GPUDirect Storage is unavailable: ", - library_error == nullptr ? "libcufile could not be loaded" : library_error); - - ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileDriverOpen", driver_open_)); - auto close_status = LoadSymbol(library_, "cuFileDriverClose_v2", driver_close_); - if (!close_status.IsOK()) { - ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileDriverClose", driver_close_)); - } - ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileHandleRegister", handle_register_)); - ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileHandleDeregister", handle_deregister_)); - ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileBufRegister", buffer_register_)); - ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileBufDeregister", buffer_deregister_)); - ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileSetParameterBool", set_bool_parameter_)); - ORT_RETURN_IF_ERROR(LoadSymbol(library_, "cuFileRead", read_)); - - ORT_RETURN_IF_ERROR(CheckCuFileStatus( - set_bool_parameter_(CUFILE_PARAM_USE_PCIP2PDMA, true), - "Enabling cuFile PCI P2PDMA")); - ORT_RETURN_IF_ERROR(CheckCuFileStatus( - set_bool_parameter_(CUFILE_PARAM_PROPERTIES_ALLOW_COMPAT_MODE, false), - "Disabling cuFile compatibility mode")); - ORT_RETURN_IF_ERROR(CheckCuFileStatus(driver_open_(), "cuFileDriverOpen")); - driver_initialized_ = true; + ORT_RETURN_IF_ERROR(CuFileDriver::Acquire(driver_)); CUDA_RETURN_IF_ERROR(cudaSetDevice(device_id)); CUDA_RETURN_IF_ERROR(cudaMalloc(&gds_buffer_, kGdsBufferSize)); ORT_RETURN_IF_ERROR(CheckCuFileStatus( - buffer_register_(gds_buffer_, kGdsBufferSize, 0), "cuFileBufRegister")); + driver_->RegisterBuffer(gds_buffer_, kGdsBufferSize), "cuFileBufRegister")); gds_buffer_registered_ = true; return Status::OK(); } LinuxGdsLoader() = default; - void* library_{nullptr}; + std::shared_ptr driver_; void* gds_buffer_{nullptr}; - bool driver_initialized_{false}; bool gds_buffer_registered_{false}; - DriverOpenFn driver_open_{nullptr}; - DriverCloseFn driver_close_{nullptr}; - HandleRegisterFn handle_register_{nullptr}; - HandleDeregisterFn handle_deregister_{nullptr}; - BufferRegisterFn buffer_register_{nullptr}; - BufferDeregisterFn buffer_deregister_{nullptr}; - SetBoolParameterFn set_bool_parameter_{nullptr}; - ReadFn read_{nullptr}; }; #endif diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h index 4b18552232b33..5af5197186f31 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h @@ -14,13 +14,15 @@ class Tensor; namespace cuda { +inline constexpr size_t kGdsIoAlignment = 4096; + class GdsLoader { public: using CreateFn = common::Status (*)(int device_id, std::unique_ptr& loader); virtual ~GdsLoader() = default; - virtual common::Status Load(const std::filesystem::path& data_file_path, + virtual common::Status Load(int file_descriptor, int64_t data_offset, size_t data_length, Tensor& tensor) const = 0; diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc index e8b096476c02d..7feebe3c584cc 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc @@ -42,13 +42,15 @@ uint8_t TestValue(size_t index, size_t load = 0) { void CreateExternalDataFile(size_t length, PathString& path, gsl::span suffix = {}, - size_t load = 0) { + size_t load = 0, + size_t prefix_size = kFilePrefixSize) { FILE* file = nullptr; path = ORT_TSTR("cuda_external_data_loader_XXXXXX"); CreateTestFile(file, path); std::vector chunk(1024 * 1024); - ASSERT_EQ(kFilePrefixSize, fwrite(chunk.data(), 1, kFilePrefixSize, file)); + ASSERT_LE(prefix_size, chunk.size()); + ASSERT_EQ(prefix_size, fwrite(chunk.data(), 1, prefix_size, file)); for (size_t offset = 0; offset < length;) { const size_t chunk_size = std::min(chunk.size(), length - offset); for (size_t i = 0; i < chunk_size; ++i) { @@ -254,7 +256,7 @@ class TestGdsLoader final : public cuda::GdsLoader { public: explicit TestGdsLoader(uint8_t value) : value_(value) {} - Status Load(const std::filesystem::path&, int64_t, size_t data_length, + Status Load(int, int64_t, size_t data_length, Tensor& tensor) const override { const auto result = cudaMemset(tensor.MutableDataRaw(), value_, data_length); ORT_RETURN_IF(result != cudaSuccess, "cudaMemset failed: ", cudaGetErrorString(result)); @@ -275,16 +277,35 @@ Status FailTestGdsLoaderCreation(int, std::unique_ptr&) { } size_t gds_create_attempt_count = 0; +size_t gds_load_attempt_count = 0; -Status CountAndFailTestGdsLoaderCreation(int, std::unique_ptr&) { +Status CreateCountingTestGdsLoader(int, std::unique_ptr& loader) { ++gds_create_attempt_count; - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "GDS unavailable for test"); + loader = std::make_unique(0x5a); + return Status::OK(); +} + +class PartialWriteFailingGdsLoader final : public cuda::GdsLoader { + public: + Status Load(int, int64_t, size_t data_length, Tensor& tensor) const override { + ++gds_load_attempt_count; + const auto result = cudaMemset(tensor.MutableDataRaw(), 0xee, data_length / 2); + ORT_RETURN_IF(result != cudaSuccess, "cudaMemset failed: ", cudaGetErrorString(result)); + ORT_RETURN_IF(cudaStreamSynchronize(nullptr) != cudaSuccess, "cudaStreamSynchronize failed"); + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "GDS read failed after a partial write"); + } +}; + +Status CreatePartialWriteFailingGdsLoader(int, std::unique_ptr& loader) { + ++gds_create_attempt_count; + loader = std::make_unique(); + return Status::OK(); } TEST(CudaExternalDataLoaderTest, UsesGdsWithoutAllocatingPinnedBuffers) { - constexpr size_t kLength = 1024; + constexpr size_t kLength = cuda::kGdsIoAlignment; PathString path; - CreateExternalDataFile(kLength, path); + CreateExternalDataFile(kLength, path, {}, 0, cuda::kGdsIoAlignment); ScopedFileDeleter file_deleter{path}; auto execution_provider = DefaultCudaExecutionProvider(); @@ -302,7 +323,7 @@ TEST(CudaExternalDataLoaderTest, UsesGdsWithoutAllocatingPinnedBuffers) { static_cast(cudaStreamCreateWithFlags), true, CreateTestGdsLoader); ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), path, kFilePrefixSize, kLength, tensor)); + Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); std::array output{}; ASSERT_EQ(cudaSuccess, cudaMemcpy( @@ -311,9 +332,9 @@ TEST(CudaExternalDataLoaderTest, UsesGdsWithoutAllocatingPinnedBuffers) { } TEST(CudaExternalDataLoaderTest, FallsBackToPinnedBuffersWhenGdsIsUnavailable) { - constexpr size_t kLength = 1024; + constexpr size_t kLength = cuda::kGdsIoAlignment; PathString path; - CreateExternalDataFile(kLength, path); + CreateExternalDataFile(kLength, path, {}, 0, cuda::kGdsIoAlignment); ScopedFileDeleter file_deleter{path}; auto execution_provider = DefaultCudaExecutionProvider(); @@ -332,7 +353,7 @@ TEST(CudaExternalDataLoaderTest, FallsBackToPinnedBuffersWhenGdsIsUnavailable) { static_cast(cudaStreamCreateWithFlags), true, FailTestGdsLoaderCreation); ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), path, kFilePrefixSize, kLength, tensor)); + Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); std::array output{}; ASSERT_EQ(cudaSuccess, cudaMemcpy( @@ -343,9 +364,9 @@ TEST(CudaExternalDataLoaderTest, FallsBackToPinnedBuffersWhenGdsIsUnavailable) { } TEST(CudaExternalDataLoaderTest, DoesNotRetryGdsAfterFailure) { - constexpr size_t kLength = 1024; + constexpr size_t kLength = cuda::kGdsIoAlignment; PathString path; - CreateExternalDataFile(kLength, path); + CreateExternalDataFile(kLength, path, {}, 0, cuda::kGdsIoAlignment); ScopedFileDeleter file_deleter{path}; auto execution_provider = DefaultCudaExecutionProvider(); @@ -359,16 +380,63 @@ TEST(CudaExternalDataLoaderTest, DoesNotRetryGdsAfterFailure) { Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); gds_create_attempt_count = 0; + gds_load_attempt_count = 0; cuda::ExternalDataLoader loader( 0, 4, static_cast(cudaMallocHost), static_cast(cudaStreamCreateWithFlags), - true, CountAndFailTestGdsLoaderCreation); + true, CreatePartialWriteFailingGdsLoader); ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), path, kFilePrefixSize, kLength, tensor)); + Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), path, kFilePrefixSize, kLength, tensor)); + Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); EXPECT_EQ(gds_create_attempt_count, 1U); + EXPECT_EQ(gds_load_attempt_count, 1U); + + std::array output{}; + ASSERT_EQ(cudaSuccess, cudaMemcpy( + output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); + for (size_t i = 0; i < output.size(); ++i) { + ASSERT_EQ(TestValue(i), output[i]) << "Mismatch at byte " << i; + } +} + +TEST(CudaExternalDataLoaderTest, UnalignedRangeDoesNotDisableGds) { + constexpr size_t kLength = cuda::kGdsIoAlignment; + PathString unaligned_path; + CreateExternalDataFile(kLength, unaligned_path); + ScopedFileDeleter unaligned_file_deleter{unaligned_path}; + PathString aligned_path; + CreateExternalDataFile(kLength, aligned_path, {}, 0, cuda::kGdsIoAlignment); + ScopedFileDeleter aligned_file_deleter{aligned_path}; + + auto execution_provider = DefaultCudaExecutionProvider(); + ASSERT_NE(execution_provider, nullptr); + auto allocators = execution_provider->CreatePreferredAllocators(); + const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { + return candidate->Info().device.Type() == OrtDevice::GPU && + candidate->Info().mem_type == OrtMemTypeDefault; + }); + ASSERT_NE(allocator, allocators.end()); + Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); + + gds_create_attempt_count = 0; + cuda::ExternalDataLoader loader( + 0, 4, + static_cast(cudaMallocHost), + static_cast(cudaStreamCreateWithFlags), + true, CreateCountingTestGdsLoader); + ASSERT_STATUS_OK(loader.LoadTensor( + Env::Default(), unaligned_path, kFilePrefixSize, kLength, tensor)); + EXPECT_EQ(gds_create_attempt_count, 0U); + ASSERT_STATUS_OK(loader.LoadTensor( + Env::Default(), aligned_path, cuda::kGdsIoAlignment, kLength, tensor)); + EXPECT_EQ(gds_create_attempt_count, 1U); + + std::array output{}; + ASSERT_EQ(cudaSuccess, cudaMemcpy( + output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); + EXPECT_TRUE(std::all_of(output.begin(), output.end(), [](uint8_t value) { return value == 0x5a; })); } TEST(CudaExternalDataLoaderTest, NormalizesBoolWithPinnedAndPageableFallback) { From df468f795316129a90dfc83280331ef36067ab41 Mon Sep 17 00:00:00 2001 From: xadupre Date: Mon, 21 Sep 2026 17:14:30 +0000 Subject: [PATCH 05/17] Address GDS follow-up review comments Avoid changing the RandomAccessFile vtable, guard optional cuFile headers, make GDS test counters atomic, and clarify diagnostics and documentation. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- docs/CUDA_GPU_Direct_Storage.md | 2 +- onnxruntime/core/platform/env.h | 12 +++++++---- onnxruntime/core/platform/posix/env.cc | 2 +- .../cuda/cuda_external_data_loader.cc | 10 ++++++++- .../cuda/cuda_external_data_loader_gds.cc | 4 +++- .../core/session/provider_bridge_ort.cc | 6 +++++- .../cuda_external_data_loader_test.cc | 21 ++++++++++++------- 7 files changed, 41 insertions(+), 16 deletions(-) diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index 0c40a8eaa4006..fddfc797d2818 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -15,7 +15,7 @@ external-data file -> registered CUDA staging buffer -> CUDA arena initializer cuFileRead device-to-device copy ``` -The reusable staging buffer bounds additional GPU memory to 64 MiB per CUDA external-data loader. Each +The reusable staging buffer bounds additional GPU memory usage to 64 MiB per CUDA external-data loader. Each device-to-device copy completes before that buffer is reused. String and Boolean initializers retain the existing loading path because they require host-side conversion. diff --git a/onnxruntime/core/platform/env.h b/onnxruntime/core/platform/env.h index 5e3be8c9974bf..de6eb1d083bca 100644 --- a/onnxruntime/core/platform/env.h +++ b/onnxruntime/core/platform/env.h @@ -130,10 +130,6 @@ class RandomAccessFile { */ virtual common::Status Read(FileOffsetType offset, gsl::span buffer) const = 0; - // Returns the POSIX file descriptor for the owned open file, or -1 when unavailable. - // The descriptor remains owned by this object and is valid only for its lifetime. - virtual int GetFileDescriptor() const { return -1; } - protected: RandomAccessFile() = default; @@ -141,6 +137,14 @@ class RandomAccessFile { ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(RandomAccessFile); }; +class PosixFileDescriptorProvider { + public: + virtual ~PosixFileDescriptorProvider() = default; + + // The descriptor remains owned by the provider and is valid only for its lifetime. + virtual int GetFileDescriptor() const = 0; +}; + /// \brief An interface used by the onnxruntime implementation to /// access operating system functionality like the filesystem etc. /// diff --git a/onnxruntime/core/platform/posix/env.cc b/onnxruntime/core/platform/posix/env.cc index d144abc60aa07..8c20cb5ce112c 100644 --- a/onnxruntime/core/platform/posix/env.cc +++ b/onnxruntime/core/platform/posix/env.cc @@ -130,7 +130,7 @@ common::Status GetFileLength(int fd, size_t& file_size) { return common::Status::OK(); } -class PosixRandomAccessFile final : public RandomAccessFile { +class PosixRandomAccessFile final : public RandomAccessFile, public PosixFileDescriptorProvider { public: PosixRandomAccessFile(ScopedFileDescriptor descriptor, std::string path) : descriptor_(std::move(descriptor)), path_(std::move(path)) {} diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc index d92353f7c4414..7a2979abccda9 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc @@ -228,7 +228,15 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, gds_status = create_gds_loader_(device_id_, gds_loader_); } if (gds_status.IsOK()) { - gds_status = gds_loader_->Load(file->GetFileDescriptor(), data_offset, length, tensor); +#if defined(ORT_NO_RTTI) + constexpr int file_descriptor = -1; +#else + const auto* descriptor_provider = dynamic_cast(file.get()); + const int file_descriptor = + descriptor_provider == nullptr ? -1 : descriptor_provider->GetFileDescriptor(); +#endif + gds_status = gds_loader_->Load( + file_descriptor, data_offset, length, tensor); } if (gds_status.IsOK()) { return Status::OK(); diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index 9a6490f06d48c..dc0cb3d017b47 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -18,13 +18,15 @@ #include "core/common/safeint.h" #include "core/providers/cuda/cuda_common.h" -#if defined(__linux__) && __has_include() +#if defined(__linux__) && defined(__has_include) +#if __has_include() #define ORT_CUDA_GDS_AVAILABLE 1 #include #include #include #include #endif +#endif namespace onnxruntime { namespace cuda { diff --git a/onnxruntime/core/session/provider_bridge_ort.cc b/onnxruntime/core/session/provider_bridge_ort.cc index 7db658a83699a..c686485dcfafa 100644 --- a/onnxruntime/core/session/provider_bridge_ort.cc +++ b/onnxruntime/core/session/provider_bridge_ort.cc @@ -2960,9 +2960,13 @@ ORT_API_STATUS_IMPL(OrtApis::SessionOptionsAppendExecutionProvider_CUDA_V2, _In_ } if (cuda_options->external_data_loader_use_gds != 0 && cuda_options->external_data_loader_use_gds != 1) { + const auto message = onnxruntime::MakeString( + "external_data_loader_use_gds got ", + cuda_options->external_data_loader_use_gds, + "; must be 0 or 1."); return OrtApis::CreateStatus( ORT_INVALID_ARGUMENT, - "external_data_loader_use_gds must be 0 or 1."); + message.c_str()); } auto factory = onnxruntime::CudaProviderFactoryCreator::Create(cuda_options); diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc index 7feebe3c584cc..095200eedc854 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc @@ -3,6 +3,7 @@ #include #include +#include #include #include #include @@ -235,6 +236,7 @@ TEST(CudaExternalDataLoaderTest, RejectsInvalidGdsOptionFromStructOptions) { FAIL() << "Expected an invalid external_data_loader_use_gds value to be rejected."; } catch (const Ort::Exception& ex) { EXPECT_THAT(ex.what(), testing::HasSubstr("external_data_loader_use_gds")); + EXPECT_THAT(ex.what(), testing::HasSubstr("got 2")); } EXPECT_EQ(CudaExecutionProviderWithOptions(&provider_options), nullptr); @@ -256,8 +258,13 @@ class TestGdsLoader final : public cuda::GdsLoader { public: explicit TestGdsLoader(uint8_t value) : value_(value) {} - Status Load(int, int64_t, size_t data_length, + Status Load(int file_descriptor, int64_t, size_t data_length, Tensor& tensor) const override { +#ifdef __linux__ + ORT_RETURN_IF(file_descriptor < 0, "Expected a POSIX file descriptor."); +#else + ORT_UNUSED_PARAMETER(file_descriptor); +#endif const auto result = cudaMemset(tensor.MutableDataRaw(), value_, data_length); ORT_RETURN_IF(result != cudaSuccess, "cudaMemset failed: ", cudaGetErrorString(result)); return Status::OK(); @@ -276,8 +283,8 @@ Status FailTestGdsLoaderCreation(int, std::unique_ptr&) { return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "GDS unavailable for test"); } -size_t gds_create_attempt_count = 0; -size_t gds_load_attempt_count = 0; +std::atomic gds_create_attempt_count{0}; +std::atomic gds_load_attempt_count{0}; Status CreateCountingTestGdsLoader(int, std::unique_ptr& loader) { ++gds_create_attempt_count; @@ -390,8 +397,8 @@ TEST(CudaExternalDataLoaderTest, DoesNotRetryGdsAfterFailure) { Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); ASSERT_STATUS_OK(loader.LoadTensor( Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); - EXPECT_EQ(gds_create_attempt_count, 1U); - EXPECT_EQ(gds_load_attempt_count, 1U); + EXPECT_EQ(gds_create_attempt_count.load(), 1U); + EXPECT_EQ(gds_load_attempt_count.load(), 1U); std::array output{}; ASSERT_EQ(cudaSuccess, cudaMemcpy( @@ -428,10 +435,10 @@ TEST(CudaExternalDataLoaderTest, UnalignedRangeDoesNotDisableGds) { true, CreateCountingTestGdsLoader); ASSERT_STATUS_OK(loader.LoadTensor( Env::Default(), unaligned_path, kFilePrefixSize, kLength, tensor)); - EXPECT_EQ(gds_create_attempt_count, 0U); + EXPECT_EQ(gds_create_attempt_count.load(), 0U); ASSERT_STATUS_OK(loader.LoadTensor( Env::Default(), aligned_path, cuda::kGdsIoAlignment, kLength, tensor)); - EXPECT_EQ(gds_create_attempt_count, 1U); + EXPECT_EQ(gds_create_attempt_count.load(), 1U); std::array output{}; ASSERT_EQ(cudaSuccess, cudaMemcpy( From 857b793e543b02e65cea042e6c64231c9a8255dd Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 22 Sep 2026 10:38:36 +0000 Subject: [PATCH 06/17] Fix GDS builds with older cuFile headers and MSVC Probe the required cuFile configuration API before enabling native GDS, so older CUDA toolkits compile and retain the logged host-memory fallback. Guard the core Tensor forward declaration from SHARED_PROVIDER builds, where Tensor is a struct. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/CMakeLists.txt | 18 ++++++++++++++++++ docs/CUDA_GPU_Direct_Storage.md | 5 ++++- .../cuda/cuda_external_data_loader_gds.cc | 8 +++----- .../cuda/cuda_external_data_loader_gds.h | 2 ++ 4 files changed, 27 insertions(+), 6 deletions(-) diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index a0edf161eb334..eda31967a00a9 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -1545,6 +1545,24 @@ if (onnxruntime_USE_CUDA) endif() find_package(CUDAToolkit REQUIRED) + if(CMAKE_SYSTEM_NAME STREQUAL "Linux") + include(CheckCXXSourceCompiles) + include(CMakePushCheckState) + cmake_push_check_state(RESET) + set(CMAKE_REQUIRED_INCLUDES ${CUDAToolkit_INCLUDE_DIRS}) + check_cxx_source_compiles(" + #include + using SetBoolParameterFn = decltype(&cuFileSetParameterBool); + int main() { + return CUFILE_PARAM_USE_PCIP2PDMA == CUFILE_PARAM_PROPERTIES_ALLOW_COMPAT_MODE; + }" onnxruntime_CUFILE_CONFIG_API_SUPPORTED) + cmake_pop_check_state() + if(onnxruntime_CUFILE_CONFIG_API_SUPPORTED) + set_property(SOURCE "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_gds.cc" + APPEND PROPERTY COMPILE_DEFINITIONS ORT_CUDA_GDS_AVAILABLE) + endif() + endif() + if(MSVC AND CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.9 AND CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.0) foreach(_cuda_include_dir IN LISTS CUDAToolkit_INCLUDE_DIRS) diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index fddfc797d2818..952970e2aa2b5 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -21,7 +21,8 @@ loading path because they require host-side conversion. GDS requires: -- Linux and a CUDA toolkit that provides `cufile.h`; +- Linux and `cufile.h` with `cuFileSetParameterBool`, `CUFILE_PARAM_USE_PCIP2PDMA`, and + `CUFILE_PARAM_PROPERTIES_ALLOW_COMPAT_MODE`; - `libcufile.so` at runtime; - either `nvidia-fs` or a recent open NVIDIA kernel module that supports PCI P2PDMA; - a supported storage/filesystem and PCIe topology; and @@ -33,6 +34,8 @@ kernels, GPUs, and storage devices. It also disables cuFile compatibility mode: a native GDS path, ONNX Runtime uses its configured pinned-buffer fallback instead of cuFile's internal POSIX fallback. GDS is attempted only for external-data ranges whose offset and length are both 4 KiB aligned. An unaligned initializer uses the configured fallback without disabling GDS for later aligned initializers. +The build checks for the required cuFile configuration API. Older CUDA toolkits without it remain supported, +but enabling GDS in those builds logs a warning and uses the configured host-memory fallback. ## Configuration diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index dc0cb3d017b47..6a585e22bf07c 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -18,15 +18,12 @@ #include "core/common/safeint.h" #include "core/providers/cuda/cuda_common.h" -#if defined(__linux__) && defined(__has_include) -#if __has_include() -#define ORT_CUDA_GDS_AVAILABLE 1 +#if defined(ORT_CUDA_GDS_AVAILABLE) #include #include #include #include #endif -#endif namespace onnxruntime { namespace cuda { @@ -275,7 +272,8 @@ common::Status GdsLoader::Create(int device_id, std::unique_ptr& load ORT_UNUSED_PARAMETER(device_id); ORT_UNUSED_PARAMETER(loader); return ORT_MAKE_STATUS(ONNXRUNTIME, NOT_IMPLEMENTED, - "GPUDirect Storage requires Linux and a CUDA toolkit with cuFile headers."); + "GPUDirect Storage requires Linux and cuFile headers with cuFileSetParameterBool, " + "CUFILE_PARAM_USE_PCIP2PDMA, and CUFILE_PARAM_PROPERTIES_ALLOW_COMPAT_MODE."); #endif } diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h index 5af5197186f31..591d4f2f6d33e 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h @@ -10,7 +10,9 @@ #include "core/common/status.h" namespace onnxruntime { +#ifndef SHARED_PROVIDER class Tensor; +#endif namespace cuda { From 4c171d681783fe33e45dad97ac90c983e9fe890d Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 22 Sep 2026 11:59:24 +0000 Subject: [PATCH 07/17] Serialize final GDS driver release before reacquisition Encapsulate shared driver ownership in a handle that acquires the lifetime mutex before dropping its strong reference. Driver teardown therefore completes before another loader can configure a replacement. Keep the mutex state alive with each handle, preserve shared reuse and initialization-error cleanup, and cover final close versus reacquire without requiring native GDS. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/onnxruntime_unittests.cmake | 1 + docs/CUDA_GPU_Direct_Storage.md | 2 + .../cuda/cuda_external_data_loader_gds.cc | 39 +---- .../cuda/cuda_external_data_loader_gds.h | 50 +++++++ .../cuda_external_data_loader_gds_test.cc | 137 ++++++++++++++++++ 5 files changed, 196 insertions(+), 33 deletions(-) create mode 100644 onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_test.cc diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index dfede7b813e6e..7f76a65d5fb1c 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1059,6 +1059,7 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_P NOT onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_REDUCED_OPS_BUILD) set(onnxruntime_test_providers_cuda_plugin_internal_test_src "${TEST_SRC_DIR}/providers/cuda/test_cases/allocator_cuda_test.cc" + "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_external_data_loader_gds_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_utils_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/group_query_attention_workspace_header_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/packed_attention_workspace_header_test.cc" diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index 952970e2aa2b5..f45e913f329f9 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -18,6 +18,8 @@ external-data file -> registered CUDA staging buffer -> CUDA arena initializer The reusable staging buffer bounds additional GPU memory usage to 64 MiB per CUDA external-data loader. Each device-to-device copy completes before that buffer is reused. String and Boolean initializers retain the existing loading path because they require host-side conversion. +Loaders share a process-wide cuFile driver. Its final release and subsequent initialization are serialized, so a +new loader cannot configure or reopen the driver until the previous driver has finished closing. GDS requires: diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index 6a585e22bf07c..a2d86c9dbfbea 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -10,7 +10,6 @@ #include #include #include -#include #include #include @@ -61,11 +60,10 @@ class CuFileDriver { using SetBoolParameterFn = decltype(&cuFileSetParameterBool); using ReadFn = decltype(&cuFileRead); + CuFileDriver() = default; + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(CuFileDriver); + ~CuFileDriver() { - std::unique_lock lock(GlobalMutex(), std::defer_lock); - if (registered_as_active_) { - lock.lock(); - } if (driver_initialized_) { ORT_IGNORE_RETURN_VALUE(driver_close_()); } @@ -74,23 +72,6 @@ class CuFileDriver { } } - static common::Status Acquire(std::shared_ptr& driver) { - static std::weak_ptr active_driver; - std::lock_guard lock(GlobalMutex()); - - driver = active_driver.lock(); - if (driver) { - return Status::OK(); - } - - auto candidate = std::shared_ptr(new CuFileDriver()); - ORT_RETURN_IF_ERROR(candidate->Initialize()); - candidate->registered_as_active_ = true; - active_driver = candidate; - driver = std::move(candidate); - return Status::OK(); - } - CUfileError_t RegisterHandle(CUfileHandle_t* handle, CUfileDescr_t* descriptor) const { return handle_register_(handle, descriptor); } @@ -112,12 +93,6 @@ class CuFileDriver { return read_(handle, buffer, length, file_offset, buffer_offset); } - private: - static std::mutex& GlobalMutex() { - static std::mutex mutex; - return mutex; - } - common::Status Initialize() { library_ = dlopen("libcufile.so.0", RTLD_NOW | RTLD_LOCAL); if (library_ == nullptr) { @@ -150,11 +125,9 @@ class CuFileDriver { return Status::OK(); } - CuFileDriver() = default; - + private: void* library_{nullptr}; bool driver_initialized_{false}; - bool registered_as_active_{false}; DriverOpenFn driver_open_{nullptr}; DriverCloseFn driver_close_{nullptr}; HandleRegisterFn handle_register_{nullptr}; @@ -244,7 +217,7 @@ class LinuxGdsLoader final : public GdsLoader { private: common::Status Initialize(int device_id) { - ORT_RETURN_IF_ERROR(CuFileDriver::Acquire(driver_)); + ORT_RETURN_IF_ERROR(driver_.Acquire()); CUDA_RETURN_IF_ERROR(cudaSetDevice(device_id)); CUDA_RETURN_IF_ERROR(cudaMalloc(&gds_buffer_, kGdsBufferSize)); @@ -256,7 +229,7 @@ class LinuxGdsLoader final : public GdsLoader { LinuxGdsLoader() = default; - std::shared_ptr driver_; + GdsDriverHandle driver_; void* gds_buffer_{nullptr}; bool gds_buffer_registered_{false}; }; diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h index 591d4f2f6d33e..29e9bf625ad03 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h @@ -6,7 +6,10 @@ #include #include #include +#include +#include +#include "core/common/common.h" #include "core/common/status.h" namespace onnxruntime { @@ -18,6 +21,53 @@ namespace cuda { inline constexpr size_t kGdsIoAlignment = 4096; +template +class GdsDriverHandle { + public: + GdsDriverHandle() : state_(SharedState()) {} + + ~GdsDriverHandle() { + // Lock before the final strong reference expires, not inside Driver's destructor. + std::lock_guard lock(state_->mutex); + driver_.reset(); + } + + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(GdsDriverHandle); + + template + common::Status Acquire(Args&&... args) { + std::lock_guard lock(state_->mutex); + if (driver_) { + return common::Status::OK(); + } + + driver_ = state_->driver.lock(); + if (!driver_) { + auto candidate = std::make_shared(std::forward(args)...); + ORT_RETURN_IF_ERROR(candidate->Initialize()); + state_->driver = candidate; + driver_ = std::move(candidate); + } + return common::Status::OK(); + } + + Driver* operator->() const { return driver_.get(); } + + private: + struct State { + std::mutex mutex; + std::weak_ptr driver; + }; + + static std::shared_ptr SharedState() { + static auto state = std::make_shared(); + return state; + } + + std::shared_ptr state_; + std::shared_ptr driver_; +}; + class GdsLoader { public: using CreateFn = common::Status (*)(int device_id, std::unique_ptr& loader); diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_test.cc new file mode 100644 index 0000000000000..75c9978ec9892 --- /dev/null +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_test.cc @@ -0,0 +1,137 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include +#include +#include +#include +#include +#include + +#include "core/providers/cuda/cuda_external_data_loader_gds.h" +#include "gtest/gtest.h" + +namespace onnxruntime { +namespace test { +namespace { + +struct DriverState { + std::atomic initialize_attempts{0}; + std::atomic opens{0}; + std::atomic closes{0}; + std::atomic active{0}; + bool fail_initialization{false}; + bool block_first_close{false}; + std::latch close_started{1}; + std::latch allow_close{1}; +}; + +class TestDriver { + public: + explicit TestDriver(DriverState& state) : state_(state) {} + + ~TestDriver() { + if (initialized_) { + if (state_.block_first_close && state_.closes.load() == 0) { + state_.close_started.count_down(); + state_.allow_close.wait(); + } + --state_.active; + ++state_.closes; + } + } + + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(TestDriver); + + Status Initialize() { + ++state_.initialize_attempts; + ORT_RETURN_IF(state_.fail_initialization, "Driver initialization failed"); + ORT_RETURN_IF(state_.active.exchange(1) != 0, "Previous driver is still open"); + ++state_.opens; + initialized_ = true; + return Status::OK(); + } + + private: + DriverState& state_; + bool initialized_{false}; +}; + +using DriverHandle = cuda::GdsDriverHandle; + +TEST(CudaGdsDriverTest, SharesDriverUntilLastHandleIsReleased) { + DriverState state; + { + auto first = std::make_unique(); + auto second = std::make_unique(); + ASSERT_TRUE(first->Acquire(state).IsOK()); + ASSERT_TRUE(first->Acquire(state).IsOK()); + ASSERT_TRUE(second->Acquire(state).IsOK()); + EXPECT_EQ(first->operator->(), second->operator->()); + EXPECT_EQ(state.initialize_attempts.load(), 1U); + + first.reset(); + EXPECT_EQ(state.closes.load(), 0U); + second.reset(); + EXPECT_EQ(state.closes.load(), 1U); + EXPECT_EQ(state.active.load(), 0U); + + DriverHandle next; + ASSERT_TRUE(next.Acquire(state).IsOK()); + EXPECT_EQ(state.opens.load(), 2U); + } + EXPECT_EQ(state.closes.load(), 2U); +} + +TEST(CudaGdsDriverTest, FailedInitializationCanBeRetried) { + DriverState state; + state.fail_initialization = true; + { + DriverHandle handle; + EXPECT_FALSE(handle.Acquire(state).IsOK()); + EXPECT_EQ(state.closes.load(), 0U); + EXPECT_EQ(state.active.load(), 0U); + + state.fail_initialization = false; + ASSERT_TRUE(handle.Acquire(state).IsOK()); + EXPECT_EQ(state.initialize_attempts.load(), 2U); + EXPECT_EQ(state.opens.load(), 1U); + } + EXPECT_EQ(state.closes.load(), 1U); +} + +TEST(CudaGdsDriverTest, ReacquireWaitsForFinalClose) { + DriverState state; + state.block_first_close = true; + auto first = std::make_unique(); + ASSERT_TRUE(first->Acquire(state).IsOK()); + + std::jthread releasing([handle = std::move(first)]() mutable { handle.reset(); }); + state.close_started.wait(); + + std::latch acquire_started{1}; + std::promise acquired; + auto result = acquired.get_future(); + std::jthread acquiring([&]() { + DriverHandle next; + acquire_started.count_down(); + acquired.set_value(next.Acquire(state)); + }); + acquire_started.wait(); + + const auto waiting = result.wait_for(std::chrono::milliseconds(100)); + state.allow_close.count_down(); + releasing.join(); + acquiring.join(); + + EXPECT_EQ(waiting, std::future_status::timeout); + EXPECT_TRUE(result.get().IsOK()); + EXPECT_EQ(state.initialize_attempts.load(), 2U); + EXPECT_EQ(state.opens.load(), 2U); + EXPECT_EQ(state.closes.load(), 2U); + EXPECT_EQ(state.active.load(), 0U); +} + +} // namespace +} // namespace test +} // namespace onnxruntime From d78f7503fc9592fc3fa789031b366af2f0f47ef8 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 22 Sep 2026 12:40:48 +0000 Subject: [PATCH 08/17] Test native GDS reads with injectable I/O callbacks Extract the existing native descriptor and chunked-read implementation behind cuFile and device-operation callbacks. Exercise that same implementation with host-only tests for multi-chunk reads, error decoding, copy/synchronization failures, and descriptor/handle cleanup. Wire coverage into regular and plugin CUDA internal tests. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/CMakeLists.txt | 2 + cmake/onnxruntime_unittests.cmake | 2 + docs/CUDA_GPU_Direct_Storage.md | 7 + .../cuda/cuda_external_data_loader_gds.cc | 95 +------ .../cuda/cuda_external_data_loader_gds_api.h | 38 +++ .../cuda/cuda_external_data_loader_gds_io.cc | 88 ++++++ .../cuda_external_data_loader_gds_io_test.cc | 260 ++++++++++++++++++ 7 files changed, 409 insertions(+), 83 deletions(-) create mode 100644 onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_api.h create mode 100644 onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_io.cc create mode 100644 onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index eda31967a00a9..974ee16b32de4 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -1559,6 +1559,8 @@ if (onnxruntime_USE_CUDA) cmake_pop_check_state() if(onnxruntime_CUFILE_CONFIG_API_SUPPORTED) set_property(SOURCE "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_gds.cc" + "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_gds_io.cc" + "${ONNXRUNTIME_ROOT}/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc" APPEND PROPERTY COMPILE_DEFINITIONS ORT_CUDA_GDS_AVAILABLE) endif() endif() diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index 7f76a65d5fb1c..751ce41d33d67 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1059,6 +1059,7 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_P NOT onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_REDUCED_OPS_BUILD) set(onnxruntime_test_providers_cuda_plugin_internal_test_src "${TEST_SRC_DIR}/providers/cuda/test_cases/allocator_cuda_test.cc" + "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_external_data_loader_gds_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_utils_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/group_query_attention_workspace_header_test.cc" @@ -1085,6 +1086,7 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_P set(onnxruntime_providers_cuda_plugin_ut_impl_src "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_allocator.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_call.cc" + "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_gds_io.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_utils.cu" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cudnn_common.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cudnn_loader.cc" diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index f45e913f329f9..16b99fdd8a5b6 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -87,3 +87,10 @@ Ort::Session session(env, ORT_TSTR("model.onnx"), session_options); The option only affects initializers stored as ONNX external data and assigned to CUDA memory. Embedded initializers and initializers assigned to other execution providers retain their existing paths. + +## Native-path tests + +On Linux builds with the required cuFile headers, `CudaGdsIoTest.*` exercises the same read loop used by the native +loader, with injected cuFile and device-copy callbacks. These tests require neither a GPU nor a working GDS installation. +They cover 64 MiB chunk boundaries, short reads, cuFile/POSIX errors, copy/synchronization failures, handle cleanup, +and restoration of the original file descriptor's flags. They do not validate native GDS hardware support or performance. diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index a2d86c9dbfbea..38b5317223fd2 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -6,30 +6,19 @@ #include "core/providers/cuda/cuda_external_data_loader_gds.h" -#include -#include -#include -#include -#include -#include - #include "core/common/common.h" -#include "core/common/safeint.h" #include "core/providers/cuda/cuda_common.h" #if defined(ORT_CUDA_GDS_AVAILABLE) -#include +#include "core/providers/cuda/cuda_external_data_loader_gds_api.h" + #include -#include -#include #endif namespace onnxruntime { namespace cuda { namespace { -constexpr size_t kGdsBufferSize = 64 * 1024 * 1024; - #if defined(ORT_CUDA_GDS_AVAILABLE) template @@ -43,12 +32,6 @@ common::Status LoadSymbol(void* library, const char* name, T& function) { return Status::OK(); } -common::Status CheckCuFileStatus(CUfileError_t status, std::string_view operation) { - ORT_RETURN_IF(status.err != CU_FILE_SUCCESS, operation, " failed: ", - cufileop_status_error(status.err), " (", static_cast(status.err), ")"); - return Status::OK(); -} - class CuFileDriver { public: using DriverOpenFn = decltype(&cuFileDriverOpen); @@ -72,12 +55,14 @@ class CuFileDriver { } } - CUfileError_t RegisterHandle(CUfileHandle_t* handle, CUfileDescr_t* descriptor) const { - return handle_register_(handle, descriptor); - } - - void DeregisterHandle(CUfileHandle_t handle) const { - handle_deregister_(handle); + GdsReadApi GetReadApi() const { + return {handle_register_, handle_deregister_, read_, + [](void* destination, const void* source, size_t length, cudaMemcpyKind kind) { + return CUDA_CALL(cudaMemcpy(destination, source, length, kind)); + }, + [](cudaStream_t stream) { + return CUDA_CALL(cudaStreamSynchronize(stream)); + }}; } CUfileError_t RegisterBuffer(const void* buffer, size_t length) const { @@ -88,11 +73,6 @@ class CuFileDriver { return buffer_deregister_(buffer); } - ssize_t Read(CUfileHandle_t handle, void* buffer, size_t length, - off_t file_offset, off_t buffer_offset) const { - return read_(handle, buffer, length, file_offset, buffer_offset); - } - common::Status Initialize() { library_ = dlopen("libcufile.so.0", RTLD_NOW | RTLD_LOCAL); if (library_ == nullptr) { @@ -160,59 +140,8 @@ class LinuxGdsLoader final : public GdsLoader { int64_t data_offset, size_t data_length, Tensor& tensor) const override { - ORT_RETURN_IF(file_descriptor < 0, - "GPUDirect Storage requires an open POSIX file descriptor."); - - const int direct_descriptor = dup(file_descriptor); - ORT_RETURN_IF(direct_descriptor < 0, "Failed to duplicate external-data file descriptor: ", - std::strerror(errno)); - auto close_file = gsl::finally([direct_descriptor]() { - ORT_IGNORE_RETURN_VALUE(close(direct_descriptor)); - }); - - const int original_flags = fcntl(direct_descriptor, F_GETFL); - ORT_RETURN_IF(original_flags < 0, "Failed to query external-data file flags: ", - std::strerror(errno)); - ORT_RETURN_IF(fcntl(direct_descriptor, F_SETFL, original_flags | O_DIRECT) < 0, - "Failed to enable O_DIRECT for GPUDirect Storage: ", std::strerror(errno)); - auto restore_flags = gsl::finally([direct_descriptor, original_flags]() { - ORT_IGNORE_RETURN_VALUE(fcntl(direct_descriptor, F_SETFL, original_flags)); - }); - - CUfileDescr_t descriptor{}; - descriptor.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; - descriptor.handle.fd = direct_descriptor; - CUfileHandle_t file_handle = nullptr; - ORT_RETURN_IF_ERROR(CheckCuFileStatus(driver_->RegisterHandle(&file_handle, &descriptor), - "cuFileHandleRegister")); - auto deregister_file = gsl::finally([&]() { driver_->DeregisterHandle(file_handle); }); - - auto* destination = static_cast(tensor.MutableDataRaw()); - for (size_t offset = 0; offset < data_length;) { - const size_t chunk_size = std::min(kGdsBufferSize, data_length - offset); - const auto file_offset = SafeInt(data_offset) + offset; - const ssize_t bytes_read = driver_->Read(file_handle, gds_buffer_, chunk_size, file_offset, 0); - if (bytes_read != static_cast(chunk_size)) { - if (bytes_read == -1) { - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", std::strerror(errno)); - } - if (bytes_read < 0) { - const auto cu_file_error = static_cast(-bytes_read); - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", - cufileop_status_error(cu_file_error), - " (", static_cast(cu_file_error), ")"); - } - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead returned ", bytes_read, - " bytes; expected ", chunk_size, "."); - } - - CUDA_RETURN_IF_ERROR( - cudaMemcpy(destination + offset, gds_buffer_, chunk_size, cudaMemcpyDeviceToDevice)); - CUDA_RETURN_IF_ERROR(cudaStreamSynchronize(nullptr)); - offset += chunk_size; - } - - return Status::OK(); + return LoadGdsFile(driver_->GetReadApi(), gds_buffer_, + file_descriptor, data_offset, data_length, tensor.MutableDataRaw()); } private: diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_api.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_api.h new file mode 100644 index 0000000000000..c4b68168bb61b --- /dev/null +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_api.h @@ -0,0 +1,38 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include +#include +#include +#include + +#include +#include + +#include "core/common/status.h" + +namespace onnxruntime { +namespace cuda { + +inline constexpr size_t kGdsBufferSize = 64 * 1024 * 1024; + +common::Status CheckCuFileStatus(CUfileError_t status, std::string_view operation); + +struct GdsReadApi { + std::function> handle_register; + std::function> handle_deregister; + std::function> read; + std::function copy; + std::function synchronize; +}; + +// The caller owns the registered staging buffer, which must hold at least kGdsBufferSize bytes. +common::Status LoadGdsFile(const GdsReadApi& api, void* staging_buffer, + int file_descriptor, int64_t data_offset, size_t data_length, + void* destination); + +} // namespace cuda +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_io.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_io.cc new file mode 100644 index 0000000000000..6f90eaa2fc65f --- /dev/null +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_io.cc @@ -0,0 +1,88 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#if defined(ORT_CUDA_GDS_AVAILABLE) + +#include "core/providers/cuda/cuda_external_data_loader_gds_api.h" + +#include +#include +#include + +#include +#include +#include + +#include "core/common/common.h" +#include "core/common/safeint.h" + +namespace onnxruntime { +namespace cuda { + +common::Status CheckCuFileStatus(CUfileError_t status, std::string_view operation) { + ORT_RETURN_IF(status.err != CU_FILE_SUCCESS, operation, " failed: ", + cufileop_status_error(status.err), " (", static_cast(status.err), ")"); + return Status::OK(); +} + +common::Status LoadGdsFile(const GdsReadApi& api, void* staging_buffer, + int file_descriptor, int64_t data_offset, size_t data_length, + void* destination_buffer) { + ORT_RETURN_IF(file_descriptor < 0, + "GPUDirect Storage requires an open POSIX file descriptor."); + + const int direct_descriptor = dup(file_descriptor); + ORT_RETURN_IF(direct_descriptor < 0, "Failed to duplicate external-data file descriptor: ", + std::strerror(errno)); + auto close_file = gsl::finally([direct_descriptor]() { + ORT_IGNORE_RETURN_VALUE(close(direct_descriptor)); + }); + + const int original_flags = fcntl(direct_descriptor, F_GETFL); + ORT_RETURN_IF(original_flags < 0, "Failed to query external-data file flags: ", + std::strerror(errno)); + ORT_RETURN_IF(fcntl(direct_descriptor, F_SETFL, original_flags | O_DIRECT) < 0, + "Failed to enable O_DIRECT for GPUDirect Storage: ", std::strerror(errno)); + auto restore_flags = gsl::finally([direct_descriptor, original_flags]() { + ORT_IGNORE_RETURN_VALUE(fcntl(direct_descriptor, F_SETFL, original_flags)); + }); + + CUfileDescr_t descriptor{}; + descriptor.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; + descriptor.handle.fd = direct_descriptor; + CUfileHandle_t file_handle = nullptr; + ORT_RETURN_IF_ERROR(CheckCuFileStatus(api.handle_register(&file_handle, &descriptor), + "cuFileHandleRegister")); + auto deregister_file = gsl::finally([&]() { api.handle_deregister(file_handle); }); + + auto* destination = static_cast(destination_buffer); + for (size_t offset = 0; offset < data_length;) { + const size_t chunk_size = std::min(kGdsBufferSize, data_length - offset); + const auto file_offset = SafeInt(data_offset) + offset; + const ssize_t bytes_read = api.read(file_handle, staging_buffer, chunk_size, file_offset, 0); + if (bytes_read != static_cast(chunk_size)) { + if (bytes_read == -1) { + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", std::strerror(errno)); + } + if (bytes_read < 0) { + const auto cu_file_error = static_cast(-bytes_read); + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", + cufileop_status_error(cu_file_error), + " (", static_cast(cu_file_error), ")"); + } + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead returned ", bytes_read, + " bytes; expected ", chunk_size, "."); + } + + ORT_RETURN_IF_ERROR(api.copy(destination + offset, staging_buffer, chunk_size, cudaMemcpyDeviceToDevice)); + ORT_RETURN_IF_ERROR(api.synchronize(nullptr)); + offset += chunk_size; + } + + return Status::OK(); +} + +} // namespace cuda +} // namespace onnxruntime + +#endif diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc new file mode 100644 index 0000000000000..9fc0e06622e1f --- /dev/null +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc @@ -0,0 +1,260 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#if defined(ORT_CUDA_GDS_AVAILABLE) + +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include "core/common/inlined_containers.h" +#include "core/providers/cuda/cuda_external_data_loader_gds.h" +#include "core/providers/cuda/cuda_external_data_loader_gds_api.h" +#include "gtest/gtest.h" + +namespace onnxruntime { +namespace test { +namespace { + +class CudaGdsIoTest : public ::testing::Test { + protected: + static constexpr size_t kAlignment = cuda::kGdsIoAlignment; + static constexpr size_t kBufferSize = cuda::kGdsBufferSize; + static constexpr uint8_t kUntouched = 0xff; + + struct Read { + size_t length; + off_t file_offset; + }; + + void SetUp() override { + file_.reset(std::tmpfile()); + ASSERT_NE(file_, nullptr); + original_flags_ = fcntl(fileno(file_.get()), F_GETFL); + ASSERT_GE(original_flags_, 0); + + api_.handle_register = [this](CUfileHandle_t* handle, CUfileDescr_t* descriptor) { + calls_.push_back("register"); + EXPECT_EQ(descriptor->type, CU_FILE_HANDLE_TYPE_OPAQUE_FD); + registered_fd_ = descriptor->handle.fd; + EXPECT_NE(registered_fd_, fileno(file_.get())); + EXPECT_EQ(fcntl(registered_fd_, F_GETFL), original_flags_ | O_DIRECT); + if (register_error_ == CU_FILE_SUCCESS) { + *handle = this; + } + return CUfileError_t{register_error_, CUDA_SUCCESS}; + }; + api_.handle_deregister = [this](CUfileHandle_t handle) { + calls_.push_back("deregister"); + EXPECT_EQ(handle, this); + EXPECT_EQ(fcntl(registered_fd_, F_GETFL), original_flags_ | O_DIRECT); + }; + api_.read = [this](CUfileHandle_t handle, void* buffer, size_t length, + off_t file_offset, off_t buffer_offset) -> ssize_t { + calls_.push_back("read"); + EXPECT_EQ(handle, this); + EXPECT_EQ(buffer, staging_.data()); + EXPECT_EQ(buffer_offset, 0); + reads_.push_back({length, file_offset}); + if (reads_.size() == fail_read_) { + errno = EIO; + return read_result_; + } + std::memset(buffer, ReadValue(file_offset), length); + return static_cast(length); + }; + api_.copy = [this](void* destination, const void* source, size_t length, cudaMemcpyKind kind) { + calls_.push_back("copy"); + EXPECT_EQ(kind, cudaMemcpyDeviceToDevice); + EXPECT_EQ(source, staging_.data()); + EXPECT_EQ(destination, destination_.data() + kAlignment + copied_); + EXPECT_EQ(length, reads_.back().length); + if (fail_copy_) { + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "Injected device copy failure"); + } + std::memcpy(destination, source, length); + copied_ += length; + return Status::OK(); + }; + api_.synchronize = [this](cudaStream_t stream) { + calls_.push_back("synchronize"); + EXPECT_EQ(stream, nullptr); + if (fail_synchronize_) { + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "Injected device synchronization failure"); + } + return Status::OK(); + }; + } + + Status Load(size_t length = kAlignment) { + destination_.assign(length + 2 * kAlignment, kUntouched); + return cuda::LoadGdsFile(api_, staging_.data(), fileno(file_.get()), kAlignment, + length, destination_.data() + kAlignment); + } + + void ExpectCleanup(bool registered = true) { + ASSERT_GE(registered_fd_, 0); + EXPECT_EQ(std::count(calls_.begin(), calls_.end(), "deregister"), registered ? 1 : 0); + EXPECT_EQ(fcntl(registered_fd_, F_GETFD), -1); + EXPECT_EQ(errno, EBADF); + EXPECT_EQ(fcntl(fileno(file_.get()), F_GETFL), original_flags_); + ExpectBytes(0, kAlignment, kUntouched); + ExpectBytes(destination_.size() - kAlignment, kAlignment, kUntouched); + } + + void ExpectBytes(size_t offset, size_t length, uint8_t value) { + ASSERT_LE(offset + length, destination_.size()); + EXPECT_TRUE(std::all_of(destination_.begin() + offset, destination_.begin() + offset + length, + [value](uint8_t byte) { return byte == value; })) + << "offset=" << offset << ", length=" << length; + } + + void ExpectError(const Status& status, std::string_view message) { + ASSERT_FALSE(status.IsOK()); + EXPECT_NE(status.ErrorMessage().find(message), std::string::npos) << status; + } + + static uint8_t ReadValue(off_t file_offset) { + return static_cast((file_offset / kAlignment) % 251); + } + + std::unique_ptr file_{nullptr, std::fclose}; + int original_flags_{-1}; + int registered_fd_{-1}; + cuda::GdsReadApi api_; + InlinedVector staging_ = InlinedVector(kBufferSize); + InlinedVector destination_; + InlinedVector reads_; + InlinedVector calls_; + CUfileOpError register_error_{CU_FILE_SUCCESS}; + size_t fail_read_{0}; + ssize_t read_result_{0}; + size_t copied_{0}; + bool fail_copy_{false}; + bool fail_synchronize_{false}; +}; + +TEST_F(CudaGdsIoTest, ReadsMultipleChunksAndTailBeforeReusingStagingBuffer) { + const size_t length = 2 * kBufferSize + kAlignment; + const auto status = Load(length); + ASSERT_TRUE(status.IsOK()) << status; + + ASSERT_EQ(reads_.size(), 3U); + EXPECT_EQ(copied_, length); + for (size_t i = 0; i < reads_.size(); ++i) { + const size_t chunk_length = i == 2 ? kAlignment : kBufferSize; + const auto file_offset = static_cast(kAlignment + i * kBufferSize); + EXPECT_EQ(reads_[i].length, chunk_length); + EXPECT_EQ(reads_[i].file_offset, file_offset); + ExpectBytes(kAlignment + i * kBufferSize, chunk_length, ReadValue(file_offset)); + } + EXPECT_EQ(calls_, (InlinedVector{ + "register", "read", "copy", "synchronize", + "read", "copy", "synchronize", "read", "copy", "synchronize", "deregister"})); + ExpectCleanup(); +} + +TEST_F(CudaGdsIoTest, ShortReadDoesNotCopyIncompleteData) { + fail_read_ = 1; + read_result_ = kAlignment - 1; + ExpectError(Load(), "cuFileRead returned 4095 bytes; expected 4096."); + EXPECT_EQ(calls_, (InlinedVector{"register", "read", "deregister"})); + ExpectBytes(kAlignment, kAlignment, kUntouched); + ExpectCleanup(); +} + +TEST_F(CudaGdsIoTest, EndOfFileDoesNotCopyStaleData) { + fail_read_ = 1; + read_result_ = 0; + ExpectError(Load(), "cuFileRead returned 0 bytes; expected 4096."); + EXPECT_EQ(calls_, (InlinedVector{"register", "read", "deregister"})); + ExpectBytes(kAlignment, kAlignment, kUntouched); + ExpectCleanup(); +} + +TEST_F(CudaGdsIoTest, DecodesPosixReadError) { + fail_read_ = 1; + read_result_ = -1; + ExpectError(Load(), std::string("cuFileRead failed: ") + std::strerror(EIO)); + EXPECT_EQ(calls_, (InlinedVector{"register", "read", "deregister"})); + ExpectBytes(kAlignment, kAlignment, kUntouched); + ExpectCleanup(); +} + +TEST_F(CudaGdsIoTest, DecodesCuFileReadError) { + fail_read_ = 1; + read_result_ = -CU_FILE_IO_NOT_SUPPORTED; + const auto status = Load(); + ExpectError(status, std::string("cuFileRead failed: ") + cufileop_status_error(CU_FILE_IO_NOT_SUPPORTED)); + ExpectError(status, "(" + std::to_string(CU_FILE_IO_NOT_SUPPORTED) + ")"); + EXPECT_EQ(calls_, (InlinedVector{"register", "read", "deregister"})); + ExpectBytes(kAlignment, kAlignment, kUntouched); + ExpectCleanup(); +} + +TEST_F(CudaGdsIoTest, RegistrationFailureClosesDescriptorAndRestoresFlags) { + register_error_ = CU_FILE_INVALID_VALUE; + ExpectError(Load(), std::string("cuFileHandleRegister failed: ") + cufileop_status_error(register_error_)); + EXPECT_EQ(calls_, (InlinedVector{"register"})); + ExpectBytes(kAlignment, kAlignment, kUntouched); + ExpectCleanup(false); +} + +TEST_F(CudaGdsIoTest, CopyFailureStopsBeforeSynchronizationAndNextRead) { + fail_copy_ = true; + ExpectError(Load(kBufferSize + kAlignment), "Injected device copy failure"); + EXPECT_EQ(calls_, (InlinedVector{"register", "read", "copy", "deregister"})); + ExpectBytes(kAlignment, kBufferSize + kAlignment, kUntouched); + ExpectCleanup(); +} + +TEST_F(CudaGdsIoTest, SynchronizationFailureStopsBeforeStagingBufferReuse) { + fail_synchronize_ = true; + ExpectError(Load(kBufferSize + kAlignment), "Injected device synchronization failure"); + EXPECT_EQ(calls_, (InlinedVector{"register", "read", "copy", "synchronize", "deregister"})); + ExpectBytes(kAlignment, kBufferSize, ReadValue(kAlignment)); + ExpectBytes(kAlignment + kBufferSize, kAlignment, kUntouched); + ExpectCleanup(); +} + +TEST_F(CudaGdsIoTest, LaterReadFailureLeavesRemainingDestinationUntouched) { + fail_read_ = 2; + read_result_ = -1; + ExpectError(Load(2 * kBufferSize + kAlignment), std::strerror(EIO)); + ASSERT_EQ(reads_.size(), 2U); + EXPECT_EQ(reads_[1].file_offset, static_cast(kAlignment + kBufferSize)); + EXPECT_EQ(calls_, (InlinedVector{ + "register", "read", "copy", "synchronize", "read", "deregister"})); + ExpectBytes(kAlignment, kBufferSize, ReadValue(kAlignment)); + ExpectBytes(kAlignment + kBufferSize, kBufferSize + kAlignment, kUntouched); + ExpectCleanup(); +} + +TEST_F(CudaGdsIoTest, InvalidDescriptorDoesNotCallCuFile) { + ExpectError(cuda::LoadGdsFile(api_, staging_.data(), -1, 0, kAlignment, staging_.data()), + "requires an open POSIX file descriptor"); + EXPECT_TRUE(calls_.empty()); +} + +TEST_F(CudaGdsIoTest, ClosedDescriptorReportsDuplicationFailure) { + const int descriptor = dup(fileno(file_.get())); + ASSERT_GE(descriptor, 0); + ASSERT_EQ(close(descriptor), 0); + ExpectError(cuda::LoadGdsFile(api_, staging_.data(), descriptor, 0, kAlignment, staging_.data()), + "Failed to duplicate external-data file descriptor"); + EXPECT_TRUE(calls_.empty()); + EXPECT_EQ(fcntl(fileno(file_.get()), F_GETFL), original_flags_); +} + +} // namespace +} // namespace test +} // namespace onnxruntime + +#endif From 3704b43cb3f2edda214a6d12abca1d93ee238c41 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 22 Sep 2026 13:18:49 +0000 Subject: [PATCH 09/17] Remove test-only CUDA external loader injection APIs Call CUDA allocation, stream creation, and GdsLoader::Create directly. Remove the injectable native GDS read API and restore the private native implementation. Replace injected loader tests with real-provider integration coverage while retaining shared-driver lifetime safety. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/CMakeLists.txt | 2 - cmake/onnxruntime_unittests.cmake | 2 - docs/CUDA_GPU_Direct_Storage.md | 10 +- .../providers/cuda/cuda_execution_provider.cc | 2 - .../cuda/cuda_external_data_loader.cc | 17 +- .../cuda/cuda_external_data_loader.h | 12 +- .../cuda/cuda_external_data_loader_gds.cc | 93 ++++++- .../cuda/cuda_external_data_loader_gds.h | 2 - .../cuda/cuda_external_data_loader_gds_api.h | 38 --- .../cuda/cuda_external_data_loader_gds_io.cc | 88 ------ .../cuda_external_data_loader_gds_io_test.cc | 260 ------------------ .../cuda_external_data_loader_test.cc | 231 ++-------------- 12 files changed, 111 insertions(+), 646 deletions(-) delete mode 100644 onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_api.h delete mode 100644 onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_io.cc delete mode 100644 onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index 974ee16b32de4..eda31967a00a9 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -1559,8 +1559,6 @@ if (onnxruntime_USE_CUDA) cmake_pop_check_state() if(onnxruntime_CUFILE_CONFIG_API_SUPPORTED) set_property(SOURCE "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_gds.cc" - "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_gds_io.cc" - "${ONNXRUNTIME_ROOT}/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc" APPEND PROPERTY COMPILE_DEFINITIONS ORT_CUDA_GDS_AVAILABLE) endif() endif() diff --git a/cmake/onnxruntime_unittests.cmake b/cmake/onnxruntime_unittests.cmake index 751ce41d33d67..7f76a65d5fb1c 100644 --- a/cmake/onnxruntime_unittests.cmake +++ b/cmake/onnxruntime_unittests.cmake @@ -1059,7 +1059,6 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_P NOT onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_REDUCED_OPS_BUILD) set(onnxruntime_test_providers_cuda_plugin_internal_test_src "${TEST_SRC_DIR}/providers/cuda/test_cases/allocator_cuda_test.cc" - "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_external_data_loader_gds_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/cuda_utils_test.cc" "${TEST_SRC_DIR}/providers/cuda/test_cases/group_query_attention_workspace_header_test.cc" @@ -1086,7 +1085,6 @@ if (onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS AND onnxruntime_BUILD_CUDA_EP_AS_P set(onnxruntime_providers_cuda_plugin_ut_impl_src "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_allocator.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_call.cc" - "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_gds_io.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_utils.cu" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cudnn_common.cc" "${ONNXRUNTIME_ROOT}/core/providers/cuda/cudnn_loader.cc" diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index 16b99fdd8a5b6..842e548e9f2ca 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -88,9 +88,9 @@ Ort::Session session(env, ORT_TSTR("model.onnx"), session_options); The option only affects initializers stored as ONNX external data and assigned to CUDA memory. Embedded initializers and initializers assigned to other execution providers retain their existing paths. -## Native-path tests +## Tests -On Linux builds with the required cuFile headers, `CudaGdsIoTest.*` exercises the same read loop used by the native -loader, with injected cuFile and device-copy callbacks. These tests require neither a GPU nor a working GDS installation. -They cover 64 MiB chunk boundaries, short reads, cuFile/POSIX errors, copy/synchronization failures, handle cleanup, -and restoration of the original file descriptor's flags. They do not validate native GDS hardware support or performance. +`CudaExternalDataLoaderTest.*Gds*` exercises the real loader with GDS enabled, including aligned and unaligned ranges, +multiple buffers, repeated loads, and different host-memory fallback configurations. These tests require a CUDA GPU, +but use the configured fallback when native GDS is unavailable. A passing result alone does not prove native GDS usage +or performance. `CudaGdsDriverTest.*` separately checks shared-driver lifetime synchronization without GPU hardware. diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 454f319e41554..007d714fdbfa4 100755 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -3440,8 +3440,6 @@ std::unique_ptr CUDAExecutionProvider::GetExte return std::make_unique( info_.device_id, info_.external_data_loader_reading_threads, - static_cast(cudaMallocHost), - static_cast(cudaStreamCreateWithFlags), info_.external_data_loader_use_gds); } diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc index 7a2979abccda9..3976275637bd2 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc @@ -116,17 +116,10 @@ common::Status LoadWithPageableBuffer(const RandomAccessFile& file, FileOffsetTy } // namespace -ExternalDataLoader::ExternalDataLoader(int device_id, size_t reading_thread_count, - AllocatePinnedBufferFn allocate_pinned_buffer, - CreateStreamFn create_stream, - bool use_gds, - GdsLoader::CreateFn create_gds_loader) +ExternalDataLoader::ExternalDataLoader(int device_id, size_t reading_thread_count, bool use_gds) : device_id_(device_id), reading_thread_count_(reading_thread_count), - allocate_pinned_buffer_(allocate_pinned_buffer), - create_stream_(create_stream), - use_gds_(use_gds), - create_gds_loader_(create_gds_loader) {} + use_gds_(use_gds) {} ExternalDataLoader::~ExternalDataLoader() { reader_pool_.reset(); @@ -147,13 +140,13 @@ common::Status ExternalDataLoader::EnsureResources() const { } for (size_t i = 0; i < buffers_.size(); ++i) { - auto status = CUDA_CALL(allocate_pinned_buffer_(&buffers_[i], kExternalDataLoaderBufferSize)); + auto status = CUDA_CALL(cudaMallocHost(&buffers_[i], kExternalDataLoaderBufferSize)); if (!status.IsOK()) { ReleaseResources(); return status; } - status = CUDA_CALL(create_stream_(&streams_[i], cudaStreamNonBlocking)); + status = CUDA_CALL(cudaStreamCreateWithFlags(&streams_[i], cudaStreamNonBlocking)); if (!status.IsOK()) { ReleaseResources(); return status; @@ -225,7 +218,7 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, !tensor.IsDataType()) { Status gds_status = Status::OK(); if (!gds_loader_) { - gds_status = create_gds_loader_(device_id_, gds_loader_); + gds_status = GdsLoader::Create(device_id_, gds_loader_); } if (gds_status.IsOK()) { #if defined(ORT_NO_RTTI) diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader.h index 31dbf9a5f4a93..27f04ebd76882 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.h +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.h @@ -65,14 +65,7 @@ class ExternalDataLoaderThreadPool; */ class ExternalDataLoader final : public IExternalDataLoader { public: - using AllocatePinnedBufferFn = cudaError_t (*)(void**, size_t); - using CreateStreamFn = cudaError_t (*)(cudaStream_t*, unsigned int); - - ExternalDataLoader(int device_id, size_t reading_thread_count, - AllocatePinnedBufferFn allocate_pinned_buffer = cudaMallocHost, - CreateStreamFn create_stream = cudaStreamCreateWithFlags, - bool use_gds = false, - GdsLoader::CreateFn create_gds_loader = GdsLoader::Create); + ExternalDataLoader(int device_id, size_t reading_thread_count, bool use_gds = false); ~ExternalDataLoader() override; bool CanLoad(const OrtMemoryInfo& target_memory_info) const override; @@ -92,10 +85,7 @@ class ExternalDataLoader final : public IExternalDataLoader { mutable std::array buffers_{}; mutable std::array streams_{}; const size_t reading_thread_count_; - const AllocatePinnedBufferFn allocate_pinned_buffer_; - const CreateStreamFn create_stream_; const bool use_gds_; - const GdsLoader::CreateFn create_gds_loader_; mutable bool gds_disabled_{false}; mutable std::unique_ptr gds_loader_; mutable std::unique_ptr reader_pool_; diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index 38b5317223fd2..fc5244a065de2 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -6,13 +6,20 @@ #include "core/providers/cuda/cuda_external_data_loader_gds.h" +#include +#include +#include +#include + #include "core/common/common.h" +#include "core/common/safeint.h" #include "core/providers/cuda/cuda_common.h" #if defined(ORT_CUDA_GDS_AVAILABLE) -#include "core/providers/cuda/cuda_external_data_loader_gds_api.h" - +#include #include +#include +#include #endif namespace onnxruntime { @@ -21,6 +28,8 @@ namespace { #if defined(ORT_CUDA_GDS_AVAILABLE) +constexpr size_t kGdsBufferSize = 64 * 1024 * 1024; + template common::Status LoadSymbol(void* library, const char* name, T& function) { dlerror(); @@ -32,6 +41,12 @@ common::Status LoadSymbol(void* library, const char* name, T& function) { return Status::OK(); } +common::Status CheckCuFileStatus(CUfileError_t status, std::string_view operation) { + ORT_RETURN_IF(status.err != CU_FILE_SUCCESS, operation, " failed: ", + cufileop_status_error(status.err), " (", static_cast(status.err), ")"); + return Status::OK(); +} + class CuFileDriver { public: using DriverOpenFn = decltype(&cuFileDriverOpen); @@ -55,14 +70,12 @@ class CuFileDriver { } } - GdsReadApi GetReadApi() const { - return {handle_register_, handle_deregister_, read_, - [](void* destination, const void* source, size_t length, cudaMemcpyKind kind) { - return CUDA_CALL(cudaMemcpy(destination, source, length, kind)); - }, - [](cudaStream_t stream) { - return CUDA_CALL(cudaStreamSynchronize(stream)); - }}; + CUfileError_t RegisterHandle(CUfileHandle_t* handle, CUfileDescr_t* descriptor) const { + return handle_register_(handle, descriptor); + } + + void DeregisterHandle(CUfileHandle_t handle) const { + handle_deregister_(handle); } CUfileError_t RegisterBuffer(const void* buffer, size_t length) const { @@ -73,6 +86,11 @@ class CuFileDriver { return buffer_deregister_(buffer); } + ssize_t Read(CUfileHandle_t handle, void* buffer, size_t length, + off_t file_offset, off_t buffer_offset) const { + return read_(handle, buffer, length, file_offset, buffer_offset); + } + common::Status Initialize() { library_ = dlopen("libcufile.so.0", RTLD_NOW | RTLD_LOCAL); if (library_ == nullptr) { @@ -140,8 +158,59 @@ class LinuxGdsLoader final : public GdsLoader { int64_t data_offset, size_t data_length, Tensor& tensor) const override { - return LoadGdsFile(driver_->GetReadApi(), gds_buffer_, - file_descriptor, data_offset, data_length, tensor.MutableDataRaw()); + ORT_RETURN_IF(file_descriptor < 0, + "GPUDirect Storage requires an open POSIX file descriptor."); + + const int direct_descriptor = dup(file_descriptor); + ORT_RETURN_IF(direct_descriptor < 0, "Failed to duplicate external-data file descriptor: ", + std::strerror(errno)); + auto close_file = gsl::finally([direct_descriptor]() { + ORT_IGNORE_RETURN_VALUE(close(direct_descriptor)); + }); + + const int original_flags = fcntl(direct_descriptor, F_GETFL); + ORT_RETURN_IF(original_flags < 0, "Failed to query external-data file flags: ", + std::strerror(errno)); + ORT_RETURN_IF(fcntl(direct_descriptor, F_SETFL, original_flags | O_DIRECT) < 0, + "Failed to enable O_DIRECT for GPUDirect Storage: ", std::strerror(errno)); + auto restore_flags = gsl::finally([direct_descriptor, original_flags]() { + ORT_IGNORE_RETURN_VALUE(fcntl(direct_descriptor, F_SETFL, original_flags)); + }); + + CUfileDescr_t descriptor{}; + descriptor.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; + descriptor.handle.fd = direct_descriptor; + CUfileHandle_t file_handle = nullptr; + ORT_RETURN_IF_ERROR(CheckCuFileStatus(driver_->RegisterHandle(&file_handle, &descriptor), + "cuFileHandleRegister")); + auto deregister_file = gsl::finally([&]() { driver_->DeregisterHandle(file_handle); }); + + auto* destination = static_cast(tensor.MutableDataRaw()); + for (size_t offset = 0; offset < data_length;) { + const size_t chunk_size = std::min(kGdsBufferSize, data_length - offset); + const auto file_offset = SafeInt(data_offset) + offset; + const ssize_t bytes_read = driver_->Read(file_handle, gds_buffer_, chunk_size, file_offset, 0); + if (bytes_read != static_cast(chunk_size)) { + if (bytes_read == -1) { + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", std::strerror(errno)); + } + if (bytes_read < 0) { + const auto cu_file_error = static_cast(-bytes_read); + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", + cufileop_status_error(cu_file_error), + " (", static_cast(cu_file_error), ")"); + } + return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead returned ", bytes_read, + " bytes; expected ", chunk_size, "."); + } + + CUDA_RETURN_IF_ERROR( + cudaMemcpy(destination + offset, gds_buffer_, chunk_size, cudaMemcpyDeviceToDevice)); + CUDA_RETURN_IF_ERROR(cudaStreamSynchronize(nullptr)); + offset += chunk_size; + } + + return Status::OK(); } private: diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h index 29e9bf625ad03..2b15fdd3a0cd6 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.h @@ -70,8 +70,6 @@ class GdsDriverHandle { class GdsLoader { public: - using CreateFn = common::Status (*)(int device_id, std::unique_ptr& loader); - virtual ~GdsLoader() = default; virtual common::Status Load(int file_descriptor, diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_api.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_api.h deleted file mode 100644 index c4b68168bb61b..0000000000000 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_api.h +++ /dev/null @@ -1,38 +0,0 @@ -// Copyright (c) Microsoft Corporation. All rights reserved. -// Licensed under the MIT License. - -#pragma once - -#include -#include -#include -#include -#include - -#include -#include - -#include "core/common/status.h" - -namespace onnxruntime { -namespace cuda { - -inline constexpr size_t kGdsBufferSize = 64 * 1024 * 1024; - -common::Status CheckCuFileStatus(CUfileError_t status, std::string_view operation); - -struct GdsReadApi { - std::function> handle_register; - std::function> handle_deregister; - std::function> read; - std::function copy; - std::function synchronize; -}; - -// The caller owns the registered staging buffer, which must hold at least kGdsBufferSize bytes. -common::Status LoadGdsFile(const GdsReadApi& api, void* staging_buffer, - int file_descriptor, int64_t data_offset, size_t data_length, - void* destination); - -} // namespace cuda -} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_io.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_io.cc deleted file mode 100644 index 6f90eaa2fc65f..0000000000000 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds_io.cc +++ /dev/null @@ -1,88 +0,0 @@ -// Copyright (c) Microsoft Corporation. All rights reserved. -// Licensed under the MIT License. - -#if defined(ORT_CUDA_GDS_AVAILABLE) - -#include "core/providers/cuda/cuda_external_data_loader_gds_api.h" - -#include -#include -#include - -#include -#include -#include - -#include "core/common/common.h" -#include "core/common/safeint.h" - -namespace onnxruntime { -namespace cuda { - -common::Status CheckCuFileStatus(CUfileError_t status, std::string_view operation) { - ORT_RETURN_IF(status.err != CU_FILE_SUCCESS, operation, " failed: ", - cufileop_status_error(status.err), " (", static_cast(status.err), ")"); - return Status::OK(); -} - -common::Status LoadGdsFile(const GdsReadApi& api, void* staging_buffer, - int file_descriptor, int64_t data_offset, size_t data_length, - void* destination_buffer) { - ORT_RETURN_IF(file_descriptor < 0, - "GPUDirect Storage requires an open POSIX file descriptor."); - - const int direct_descriptor = dup(file_descriptor); - ORT_RETURN_IF(direct_descriptor < 0, "Failed to duplicate external-data file descriptor: ", - std::strerror(errno)); - auto close_file = gsl::finally([direct_descriptor]() { - ORT_IGNORE_RETURN_VALUE(close(direct_descriptor)); - }); - - const int original_flags = fcntl(direct_descriptor, F_GETFL); - ORT_RETURN_IF(original_flags < 0, "Failed to query external-data file flags: ", - std::strerror(errno)); - ORT_RETURN_IF(fcntl(direct_descriptor, F_SETFL, original_flags | O_DIRECT) < 0, - "Failed to enable O_DIRECT for GPUDirect Storage: ", std::strerror(errno)); - auto restore_flags = gsl::finally([direct_descriptor, original_flags]() { - ORT_IGNORE_RETURN_VALUE(fcntl(direct_descriptor, F_SETFL, original_flags)); - }); - - CUfileDescr_t descriptor{}; - descriptor.type = CU_FILE_HANDLE_TYPE_OPAQUE_FD; - descriptor.handle.fd = direct_descriptor; - CUfileHandle_t file_handle = nullptr; - ORT_RETURN_IF_ERROR(CheckCuFileStatus(api.handle_register(&file_handle, &descriptor), - "cuFileHandleRegister")); - auto deregister_file = gsl::finally([&]() { api.handle_deregister(file_handle); }); - - auto* destination = static_cast(destination_buffer); - for (size_t offset = 0; offset < data_length;) { - const size_t chunk_size = std::min(kGdsBufferSize, data_length - offset); - const auto file_offset = SafeInt(data_offset) + offset; - const ssize_t bytes_read = api.read(file_handle, staging_buffer, chunk_size, file_offset, 0); - if (bytes_read != static_cast(chunk_size)) { - if (bytes_read == -1) { - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", std::strerror(errno)); - } - if (bytes_read < 0) { - const auto cu_file_error = static_cast(-bytes_read); - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead failed: ", - cufileop_status_error(cu_file_error), - " (", static_cast(cu_file_error), ")"); - } - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "cuFileRead returned ", bytes_read, - " bytes; expected ", chunk_size, "."); - } - - ORT_RETURN_IF_ERROR(api.copy(destination + offset, staging_buffer, chunk_size, cudaMemcpyDeviceToDevice)); - ORT_RETURN_IF_ERROR(api.synchronize(nullptr)); - offset += chunk_size; - } - - return Status::OK(); -} - -} // namespace cuda -} // namespace onnxruntime - -#endif diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc deleted file mode 100644 index 9fc0e06622e1f..0000000000000 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_gds_io_test.cc +++ /dev/null @@ -1,260 +0,0 @@ -// Copyright (c) Microsoft Corporation. All rights reserved. -// Licensed under the MIT License. - -#if defined(ORT_CUDA_GDS_AVAILABLE) - -#include -#include -#include -#include -#include -#include -#include - -#include -#include - -#include "core/common/inlined_containers.h" -#include "core/providers/cuda/cuda_external_data_loader_gds.h" -#include "core/providers/cuda/cuda_external_data_loader_gds_api.h" -#include "gtest/gtest.h" - -namespace onnxruntime { -namespace test { -namespace { - -class CudaGdsIoTest : public ::testing::Test { - protected: - static constexpr size_t kAlignment = cuda::kGdsIoAlignment; - static constexpr size_t kBufferSize = cuda::kGdsBufferSize; - static constexpr uint8_t kUntouched = 0xff; - - struct Read { - size_t length; - off_t file_offset; - }; - - void SetUp() override { - file_.reset(std::tmpfile()); - ASSERT_NE(file_, nullptr); - original_flags_ = fcntl(fileno(file_.get()), F_GETFL); - ASSERT_GE(original_flags_, 0); - - api_.handle_register = [this](CUfileHandle_t* handle, CUfileDescr_t* descriptor) { - calls_.push_back("register"); - EXPECT_EQ(descriptor->type, CU_FILE_HANDLE_TYPE_OPAQUE_FD); - registered_fd_ = descriptor->handle.fd; - EXPECT_NE(registered_fd_, fileno(file_.get())); - EXPECT_EQ(fcntl(registered_fd_, F_GETFL), original_flags_ | O_DIRECT); - if (register_error_ == CU_FILE_SUCCESS) { - *handle = this; - } - return CUfileError_t{register_error_, CUDA_SUCCESS}; - }; - api_.handle_deregister = [this](CUfileHandle_t handle) { - calls_.push_back("deregister"); - EXPECT_EQ(handle, this); - EXPECT_EQ(fcntl(registered_fd_, F_GETFL), original_flags_ | O_DIRECT); - }; - api_.read = [this](CUfileHandle_t handle, void* buffer, size_t length, - off_t file_offset, off_t buffer_offset) -> ssize_t { - calls_.push_back("read"); - EXPECT_EQ(handle, this); - EXPECT_EQ(buffer, staging_.data()); - EXPECT_EQ(buffer_offset, 0); - reads_.push_back({length, file_offset}); - if (reads_.size() == fail_read_) { - errno = EIO; - return read_result_; - } - std::memset(buffer, ReadValue(file_offset), length); - return static_cast(length); - }; - api_.copy = [this](void* destination, const void* source, size_t length, cudaMemcpyKind kind) { - calls_.push_back("copy"); - EXPECT_EQ(kind, cudaMemcpyDeviceToDevice); - EXPECT_EQ(source, staging_.data()); - EXPECT_EQ(destination, destination_.data() + kAlignment + copied_); - EXPECT_EQ(length, reads_.back().length); - if (fail_copy_) { - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "Injected device copy failure"); - } - std::memcpy(destination, source, length); - copied_ += length; - return Status::OK(); - }; - api_.synchronize = [this](cudaStream_t stream) { - calls_.push_back("synchronize"); - EXPECT_EQ(stream, nullptr); - if (fail_synchronize_) { - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "Injected device synchronization failure"); - } - return Status::OK(); - }; - } - - Status Load(size_t length = kAlignment) { - destination_.assign(length + 2 * kAlignment, kUntouched); - return cuda::LoadGdsFile(api_, staging_.data(), fileno(file_.get()), kAlignment, - length, destination_.data() + kAlignment); - } - - void ExpectCleanup(bool registered = true) { - ASSERT_GE(registered_fd_, 0); - EXPECT_EQ(std::count(calls_.begin(), calls_.end(), "deregister"), registered ? 1 : 0); - EXPECT_EQ(fcntl(registered_fd_, F_GETFD), -1); - EXPECT_EQ(errno, EBADF); - EXPECT_EQ(fcntl(fileno(file_.get()), F_GETFL), original_flags_); - ExpectBytes(0, kAlignment, kUntouched); - ExpectBytes(destination_.size() - kAlignment, kAlignment, kUntouched); - } - - void ExpectBytes(size_t offset, size_t length, uint8_t value) { - ASSERT_LE(offset + length, destination_.size()); - EXPECT_TRUE(std::all_of(destination_.begin() + offset, destination_.begin() + offset + length, - [value](uint8_t byte) { return byte == value; })) - << "offset=" << offset << ", length=" << length; - } - - void ExpectError(const Status& status, std::string_view message) { - ASSERT_FALSE(status.IsOK()); - EXPECT_NE(status.ErrorMessage().find(message), std::string::npos) << status; - } - - static uint8_t ReadValue(off_t file_offset) { - return static_cast((file_offset / kAlignment) % 251); - } - - std::unique_ptr file_{nullptr, std::fclose}; - int original_flags_{-1}; - int registered_fd_{-1}; - cuda::GdsReadApi api_; - InlinedVector staging_ = InlinedVector(kBufferSize); - InlinedVector destination_; - InlinedVector reads_; - InlinedVector calls_; - CUfileOpError register_error_{CU_FILE_SUCCESS}; - size_t fail_read_{0}; - ssize_t read_result_{0}; - size_t copied_{0}; - bool fail_copy_{false}; - bool fail_synchronize_{false}; -}; - -TEST_F(CudaGdsIoTest, ReadsMultipleChunksAndTailBeforeReusingStagingBuffer) { - const size_t length = 2 * kBufferSize + kAlignment; - const auto status = Load(length); - ASSERT_TRUE(status.IsOK()) << status; - - ASSERT_EQ(reads_.size(), 3U); - EXPECT_EQ(copied_, length); - for (size_t i = 0; i < reads_.size(); ++i) { - const size_t chunk_length = i == 2 ? kAlignment : kBufferSize; - const auto file_offset = static_cast(kAlignment + i * kBufferSize); - EXPECT_EQ(reads_[i].length, chunk_length); - EXPECT_EQ(reads_[i].file_offset, file_offset); - ExpectBytes(kAlignment + i * kBufferSize, chunk_length, ReadValue(file_offset)); - } - EXPECT_EQ(calls_, (InlinedVector{ - "register", "read", "copy", "synchronize", - "read", "copy", "synchronize", "read", "copy", "synchronize", "deregister"})); - ExpectCleanup(); -} - -TEST_F(CudaGdsIoTest, ShortReadDoesNotCopyIncompleteData) { - fail_read_ = 1; - read_result_ = kAlignment - 1; - ExpectError(Load(), "cuFileRead returned 4095 bytes; expected 4096."); - EXPECT_EQ(calls_, (InlinedVector{"register", "read", "deregister"})); - ExpectBytes(kAlignment, kAlignment, kUntouched); - ExpectCleanup(); -} - -TEST_F(CudaGdsIoTest, EndOfFileDoesNotCopyStaleData) { - fail_read_ = 1; - read_result_ = 0; - ExpectError(Load(), "cuFileRead returned 0 bytes; expected 4096."); - EXPECT_EQ(calls_, (InlinedVector{"register", "read", "deregister"})); - ExpectBytes(kAlignment, kAlignment, kUntouched); - ExpectCleanup(); -} - -TEST_F(CudaGdsIoTest, DecodesPosixReadError) { - fail_read_ = 1; - read_result_ = -1; - ExpectError(Load(), std::string("cuFileRead failed: ") + std::strerror(EIO)); - EXPECT_EQ(calls_, (InlinedVector{"register", "read", "deregister"})); - ExpectBytes(kAlignment, kAlignment, kUntouched); - ExpectCleanup(); -} - -TEST_F(CudaGdsIoTest, DecodesCuFileReadError) { - fail_read_ = 1; - read_result_ = -CU_FILE_IO_NOT_SUPPORTED; - const auto status = Load(); - ExpectError(status, std::string("cuFileRead failed: ") + cufileop_status_error(CU_FILE_IO_NOT_SUPPORTED)); - ExpectError(status, "(" + std::to_string(CU_FILE_IO_NOT_SUPPORTED) + ")"); - EXPECT_EQ(calls_, (InlinedVector{"register", "read", "deregister"})); - ExpectBytes(kAlignment, kAlignment, kUntouched); - ExpectCleanup(); -} - -TEST_F(CudaGdsIoTest, RegistrationFailureClosesDescriptorAndRestoresFlags) { - register_error_ = CU_FILE_INVALID_VALUE; - ExpectError(Load(), std::string("cuFileHandleRegister failed: ") + cufileop_status_error(register_error_)); - EXPECT_EQ(calls_, (InlinedVector{"register"})); - ExpectBytes(kAlignment, kAlignment, kUntouched); - ExpectCleanup(false); -} - -TEST_F(CudaGdsIoTest, CopyFailureStopsBeforeSynchronizationAndNextRead) { - fail_copy_ = true; - ExpectError(Load(kBufferSize + kAlignment), "Injected device copy failure"); - EXPECT_EQ(calls_, (InlinedVector{"register", "read", "copy", "deregister"})); - ExpectBytes(kAlignment, kBufferSize + kAlignment, kUntouched); - ExpectCleanup(); -} - -TEST_F(CudaGdsIoTest, SynchronizationFailureStopsBeforeStagingBufferReuse) { - fail_synchronize_ = true; - ExpectError(Load(kBufferSize + kAlignment), "Injected device synchronization failure"); - EXPECT_EQ(calls_, (InlinedVector{"register", "read", "copy", "synchronize", "deregister"})); - ExpectBytes(kAlignment, kBufferSize, ReadValue(kAlignment)); - ExpectBytes(kAlignment + kBufferSize, kAlignment, kUntouched); - ExpectCleanup(); -} - -TEST_F(CudaGdsIoTest, LaterReadFailureLeavesRemainingDestinationUntouched) { - fail_read_ = 2; - read_result_ = -1; - ExpectError(Load(2 * kBufferSize + kAlignment), std::strerror(EIO)); - ASSERT_EQ(reads_.size(), 2U); - EXPECT_EQ(reads_[1].file_offset, static_cast(kAlignment + kBufferSize)); - EXPECT_EQ(calls_, (InlinedVector{ - "register", "read", "copy", "synchronize", "read", "deregister"})); - ExpectBytes(kAlignment, kBufferSize, ReadValue(kAlignment)); - ExpectBytes(kAlignment + kBufferSize, kBufferSize + kAlignment, kUntouched); - ExpectCleanup(); -} - -TEST_F(CudaGdsIoTest, InvalidDescriptorDoesNotCallCuFile) { - ExpectError(cuda::LoadGdsFile(api_, staging_.data(), -1, 0, kAlignment, staging_.data()), - "requires an open POSIX file descriptor"); - EXPECT_TRUE(calls_.empty()); -} - -TEST_F(CudaGdsIoTest, ClosedDescriptorReportsDuplicationFailure) { - const int descriptor = dup(fileno(file_.get())); - ASSERT_GE(descriptor, 0); - ASSERT_EQ(close(descriptor), 0); - ExpectError(cuda::LoadGdsFile(api_, staging_.data(), descriptor, 0, kAlignment, staging_.data()), - "Failed to duplicate external-data file descriptor"); - EXPECT_TRUE(calls_.empty()); - EXPECT_EQ(fcntl(fileno(file_.get()), F_GETFL), original_flags_); -} - -} // namespace -} // namespace test -} // namespace onnxruntime - -#endif diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc index 095200eedc854..c9593f6333ecc 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc @@ -3,7 +3,6 @@ #include #include -#include #include #include #include @@ -66,11 +65,13 @@ void CreateExternalDataFile(size_t length, PathString& path, EXPECT_EQ(0, fclose(file)); } -void VerifyLoad(size_t length, size_t load_count = 1, size_t reading_thread_count = 4) { +void VerifyLoad(size_t length, size_t load_count = 1, size_t reading_thread_count = 4, + bool use_gds = false, size_t prefix_size = kFilePrefixSize) { OrtCUDAProviderOptionsV2 provider_options{}; provider_options.do_copy_in_default_stream = true; provider_options.use_tf32 = false; provider_options.external_data_loader_reading_threads = reading_thread_count; + provider_options.external_data_loader_use_gds = use_gds; auto execution_provider = CudaExecutionProviderWithOptions(&provider_options); ASSERT_NE(execution_provider, nullptr); auto loader = execution_provider->GetExternalDataLoader(); @@ -85,11 +86,11 @@ void VerifyLoad(size_t length, size_t load_count = 1, size_t reading_thread_coun for (size_t load = 0; load < load_count; ++load) { PathString path; - CreateExternalDataFile(length, path, {}, load); + CreateExternalDataFile(length, path, {}, load, prefix_size); ScopedFileDeleter file_deleter{path}; ASSERT_EQ(cudaSuccess, cudaMemset(tensor.MutableDataRaw(), 0xa5, length)); ASSERT_EQ(cudaSuccess, cudaStreamSynchronize(nullptr)); - ASSERT_STATUS_OK(loader->LoadTensor(Env::Default(), path, kFilePrefixSize, length, tensor)); + ASSERT_STATUS_OK(loader->LoadTensor(Env::Default(), path, prefix_size, length, tensor)); std::vector output(length); ASSERT_EQ(cudaSuccess, cudaMemcpy(output.data(), tensor.DataRaw(), length, cudaMemcpyDeviceToHost)); @@ -246,207 +247,23 @@ TEST(CudaExternalDataLoaderTest, ReusesAlternatingBuffersAcrossRepeatedLoads) { VerifyLoad(2 * cuda::kExternalDataLoaderBufferSize + 1, 2); } -cudaError_t FailPinnedBufferAllocation(void**, size_t) { - return cudaErrorMemoryAllocation; -} - -cudaError_t FailStreamCreation(cudaStream_t*, unsigned int) { - return cudaErrorInitializationError; -} - -class TestGdsLoader final : public cuda::GdsLoader { - public: - explicit TestGdsLoader(uint8_t value) : value_(value) {} - - Status Load(int file_descriptor, int64_t, size_t data_length, - Tensor& tensor) const override { -#ifdef __linux__ - ORT_RETURN_IF(file_descriptor < 0, "Expected a POSIX file descriptor."); -#else - ORT_UNUSED_PARAMETER(file_descriptor); -#endif - const auto result = cudaMemset(tensor.MutableDataRaw(), value_, data_length); - ORT_RETURN_IF(result != cudaSuccess, "cudaMemset failed: ", cudaGetErrorString(result)); - return Status::OK(); - } - - private: - uint8_t value_; -}; - -Status CreateTestGdsLoader(int, std::unique_ptr& loader) { - loader = std::make_unique(0x5a); - return Status::OK(); -} - -Status FailTestGdsLoaderCreation(int, std::unique_ptr&) { - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "GDS unavailable for test"); -} - -std::atomic gds_create_attempt_count{0}; -std::atomic gds_load_attempt_count{0}; - -Status CreateCountingTestGdsLoader(int, std::unique_ptr& loader) { - ++gds_create_attempt_count; - loader = std::make_unique(0x5a); - return Status::OK(); -} - -class PartialWriteFailingGdsLoader final : public cuda::GdsLoader { - public: - Status Load(int, int64_t, size_t data_length, Tensor& tensor) const override { - ++gds_load_attempt_count; - const auto result = cudaMemset(tensor.MutableDataRaw(), 0xee, data_length / 2); - ORT_RETURN_IF(result != cudaSuccess, "cudaMemset failed: ", cudaGetErrorString(result)); - ORT_RETURN_IF(cudaStreamSynchronize(nullptr) != cudaSuccess, "cudaStreamSynchronize failed"); - return ORT_MAKE_STATUS(ONNXRUNTIME, FAIL, "GDS read failed after a partial write"); - } -}; - -Status CreatePartialWriteFailingGdsLoader(int, std::unique_ptr& loader) { - ++gds_create_attempt_count; - loader = std::make_unique(); - return Status::OK(); -} - -TEST(CudaExternalDataLoaderTest, UsesGdsWithoutAllocatingPinnedBuffers) { - constexpr size_t kLength = cuda::kGdsIoAlignment; - PathString path; - CreateExternalDataFile(kLength, path, {}, 0, cuda::kGdsIoAlignment); - ScopedFileDeleter file_deleter{path}; - - auto execution_provider = DefaultCudaExecutionProvider(); - ASSERT_NE(execution_provider, nullptr); - auto allocators = execution_provider->CreatePreferredAllocators(); - const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { - return candidate->Info().device.Type() == OrtDevice::GPU && - candidate->Info().mem_type == OrtMemTypeDefault; - }); - ASSERT_NE(allocator, allocators.end()); - Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); - - cuda::ExternalDataLoader loader( - 0, 4, FailPinnedBufferAllocation, - static_cast(cudaStreamCreateWithFlags), - true, CreateTestGdsLoader); - ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); - - std::array output{}; - ASSERT_EQ(cudaSuccess, cudaMemcpy( - output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); - EXPECT_TRUE(std::all_of(output.begin(), output.end(), [](uint8_t value) { return value == 0x5a; })); -} - -TEST(CudaExternalDataLoaderTest, FallsBackToPinnedBuffersWhenGdsIsUnavailable) { - constexpr size_t kLength = cuda::kGdsIoAlignment; - PathString path; - CreateExternalDataFile(kLength, path, {}, 0, cuda::kGdsIoAlignment); - ScopedFileDeleter file_deleter{path}; - - auto execution_provider = DefaultCudaExecutionProvider(); - ASSERT_NE(execution_provider, nullptr); - auto allocators = execution_provider->CreatePreferredAllocators(); - const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { - return candidate->Info().device.Type() == OrtDevice::GPU && - candidate->Info().mem_type == OrtMemTypeDefault; - }); - ASSERT_NE(allocator, allocators.end()); - Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); - - cuda::ExternalDataLoader loader( - 0, 4, - static_cast(cudaMallocHost), - static_cast(cudaStreamCreateWithFlags), - true, FailTestGdsLoaderCreation); - ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); - - std::array output{}; - ASSERT_EQ(cudaSuccess, cudaMemcpy( - output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); - for (size_t i = 0; i < output.size(); ++i) { - ASSERT_EQ(TestValue(i), output[i]) << "Mismatch at byte " << i; +TEST(CudaExternalDataLoaderTest, LoadsAlignedDataWithGdsEnabled) { + for (const size_t reading_thread_count : {0, 1, 4}) { + SCOPED_TRACE(reading_thread_count); + VerifyLoad(cuda::kGdsIoAlignment, 2, reading_thread_count, true, cuda::kGdsIoAlignment); } } -TEST(CudaExternalDataLoaderTest, DoesNotRetryGdsAfterFailure) { - constexpr size_t kLength = cuda::kGdsIoAlignment; - PathString path; - CreateExternalDataFile(kLength, path, {}, 0, cuda::kGdsIoAlignment); - ScopedFileDeleter file_deleter{path}; - - auto execution_provider = DefaultCudaExecutionProvider(); - ASSERT_NE(execution_provider, nullptr); - auto allocators = execution_provider->CreatePreferredAllocators(); - const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { - return candidate->Info().device.Type() == OrtDevice::GPU && - candidate->Info().mem_type == OrtMemTypeDefault; - }); - ASSERT_NE(allocator, allocators.end()); - Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); - - gds_create_attempt_count = 0; - gds_load_attempt_count = 0; - cuda::ExternalDataLoader loader( - 0, 4, - static_cast(cudaMallocHost), - static_cast(cudaStreamCreateWithFlags), - true, CreatePartialWriteFailingGdsLoader); - ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); - ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), path, cuda::kGdsIoAlignment, kLength, tensor)); - EXPECT_EQ(gds_create_attempt_count.load(), 1U); - EXPECT_EQ(gds_load_attempt_count.load(), 1U); - - std::array output{}; - ASSERT_EQ(cudaSuccess, cudaMemcpy( - output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); - for (size_t i = 0; i < output.size(); ++i) { - ASSERT_EQ(TestValue(i), output[i]) << "Mismatch at byte " << i; - } +TEST(CudaExternalDataLoaderTest, LoadsMultipleBuffersWithGdsEnabled) { + VerifyLoad(2 * cuda::kExternalDataLoaderBufferSize + cuda::kGdsIoAlignment, + 2, 4, true, cuda::kGdsIoAlignment); } -TEST(CudaExternalDataLoaderTest, UnalignedRangeDoesNotDisableGds) { - constexpr size_t kLength = cuda::kGdsIoAlignment; - PathString unaligned_path; - CreateExternalDataFile(kLength, unaligned_path); - ScopedFileDeleter unaligned_file_deleter{unaligned_path}; - PathString aligned_path; - CreateExternalDataFile(kLength, aligned_path, {}, 0, cuda::kGdsIoAlignment); - ScopedFileDeleter aligned_file_deleter{aligned_path}; - - auto execution_provider = DefaultCudaExecutionProvider(); - ASSERT_NE(execution_provider, nullptr); - auto allocators = execution_provider->CreatePreferredAllocators(); - const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { - return candidate->Info().device.Type() == OrtDevice::GPU && - candidate->Info().mem_type == OrtMemTypeDefault; - }); - ASSERT_NE(allocator, allocators.end()); - Tensor tensor(DataTypeImpl::GetType(), TensorShape({kLength}), *allocator); - - gds_create_attempt_count = 0; - cuda::ExternalDataLoader loader( - 0, 4, - static_cast(cudaMallocHost), - static_cast(cudaStreamCreateWithFlags), - true, CreateCountingTestGdsLoader); - ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), unaligned_path, kFilePrefixSize, kLength, tensor)); - EXPECT_EQ(gds_create_attempt_count.load(), 0U); - ASSERT_STATUS_OK(loader.LoadTensor( - Env::Default(), aligned_path, cuda::kGdsIoAlignment, kLength, tensor)); - EXPECT_EQ(gds_create_attempt_count.load(), 1U); - - std::array output{}; - ASSERT_EQ(cudaSuccess, cudaMemcpy( - output.data(), tensor.DataRaw(), output.size(), cudaMemcpyDeviceToHost)); - EXPECT_TRUE(std::all_of(output.begin(), output.end(), [](uint8_t value) { return value == 0x5a; })); +TEST(CudaExternalDataLoaderTest, LoadsUnalignedDataWithGdsEnabled) { + VerifyLoad(cuda::kExternalDataLoaderParallelReadThreshold + 1, 2, 4, true); } -TEST(CudaExternalDataLoaderTest, NormalizesBoolWithPinnedAndPageableFallback) { +TEST(CudaExternalDataLoaderTest, NormalizesBoolWithPinnedAndPageableLoading) { const std::array input{0, 1, 2, 255}; const std::array expected{0, 1, 1, 1}; PathString path; @@ -462,22 +279,12 @@ TEST(CudaExternalDataLoaderTest, NormalizesBoolWithPinnedAndPageableFallback) { }); ASSERT_NE(allocator, allocators.end()); - for (int failure_mode = 0; failure_mode < 3; ++failure_mode) { - SCOPED_TRACE(failure_mode); - std::unique_ptr loader; - if (failure_mode == 1) { - loader = std::make_unique( - 0, 4, FailPinnedBufferAllocation); - } else if (failure_mode == 2) { - loader = std::make_unique( - 0, 4, static_cast(cudaMallocHost), - FailStreamCreation); - } else { - loader = std::make_unique(0, 4); - } + for (const size_t reading_thread_count : {0, 1, 4}) { + SCOPED_TRACE(reading_thread_count); + cuda::ExternalDataLoader loader(0, reading_thread_count); Tensor tensor(DataTypeImpl::GetType(), TensorShape({static_cast(input.size())}), *allocator); - ASSERT_STATUS_OK(loader->LoadTensor( + ASSERT_STATUS_OK(loader.LoadTensor( Env::Default(), path, kFilePrefixSize, input.size(), tensor)); std::array output{}; ASSERT_EQ(cudaSuccess, cudaMemcpy( From fdea334f58b6d1bd509e46d9973272a3f3fdbd40 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 22 Sep 2026 13:48:57 +0000 Subject: [PATCH 10/17] Clarify configured GDS host-memory fallback Address the documentation note embedded in the Copilot review summary: fallback can use pinned or pageable memory, depending on the configured reading-thread count. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- docs/CUDA_GPU_Direct_Storage.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index 842e548e9f2ca..b1755468deb28 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -2,7 +2,7 @@ The CUDA execution provider can use NVIDIA GPUDirect Storage (GDS) to load ONNX external initializers without staging file data in CPU memory. GDS is opt-in. If it cannot be initialized or cannot read an external-data file, -ONNX Runtime logs a warning and uses the existing pinned-host-buffer loader for the rest of the session. +ONNX Runtime logs a warning and uses the configured pinned/pageable host-memory fallback for the rest of the session. ## Data path @@ -33,7 +33,7 @@ GDS requires: ONNX Runtime loads `libcufile` dynamically, so enabling the option does not add a mandatory runtime dependency for users who keep GDS disabled. It requests PCI P2PDMA, which can provide GDS without `nvidia-fs` on supported recent kernels, GPUs, and storage devices. It also disables cuFile compatibility mode: if the storage stack cannot provide -a native GDS path, ONNX Runtime uses its configured pinned-buffer fallback instead of cuFile's internal POSIX fallback. +a native GDS path, ONNX Runtime uses its configured host-memory fallback instead of cuFile's internal POSIX fallback. GDS is attempted only for external-data ranges whose offset and length are both 4 KiB aligned. An unaligned initializer uses the configured fallback without disabling GDS for later aligned initializers. The build checks for the required cuFile configuration API. Older CUDA toolkits without it remain supported, From 9efeabcea154358ac99978060ea6fb53b0fdc332 Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 22 Sep 2026 14:08:20 +0000 Subject: [PATCH 11/17] Preserve close-on-exec on duplicated GDS descriptors Use F_DUPFD_CLOEXEC so the already-open external-data descriptor is duplicated atomically without leaking it across a concurrent exec. Clarify that the configured GDS fallback can use pinned or pageable host memory. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- include/onnxruntime/core/providers/cuda/cuda_provider_options.h | 2 +- .../core/providers/cuda/cuda_external_data_loader_gds.cc | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h index 00a40399faace..5d19f97ee1f9a 100644 --- a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h +++ b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h @@ -43,5 +43,5 @@ struct OrtCUDAProviderOptionsV2 { int fuse_conv_bias = 0; // Enable CUDNN Frontend kernel fusing, results in JIT compiles int sdpa_kernel = 0; // Scaled Dot Product Attention kernel option size_t external_data_loader_reading_threads = 4; // Number of CPU read tasks per external-data staging buffer. 0 disables pinned-buffer fallback; 1 disables parallel reads. - int external_data_loader_use_gds = 0; // Try GPUDirect Storage before falling back to the pinned-buffer loader. + int external_data_loader_use_gds = 0; // Try GPUDirect Storage before the configured pinned/pageable host-memory fallback. }; diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc index fc5244a065de2..9ce52309a7be8 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_gds.cc @@ -161,7 +161,7 @@ class LinuxGdsLoader final : public GdsLoader { ORT_RETURN_IF(file_descriptor < 0, "GPUDirect Storage requires an open POSIX file descriptor."); - const int direct_descriptor = dup(file_descriptor); + const int direct_descriptor = fcntl(file_descriptor, F_DUPFD_CLOEXEC, 0); ORT_RETURN_IF(direct_descriptor < 0, "Failed to duplicate external-data file descriptor: ", std::strerror(errno)); auto close_file = gsl::finally([direct_descriptor]() { From 0a3a7f9e41ddad3430af4b677418ee4490d6c2cb Mon Sep 17 00:00:00 2001 From: xadupre Date: Tue, 22 Sep 2026 15:21:49 +0000 Subject: [PATCH 12/17] Cover GDS provider option parsing and serialization Exercise the public CUDA provider-options API for defaults, 0/1 parsing, serialized round trips, invalid values, and unchanged state after rejected updates. Use the actual parser and serializer without adding production test hooks. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- docs/CUDA_GPU_Direct_Storage.md | 2 + onnxruntime/test/shared_lib/test_inference.cc | 47 +++++++++++++++++++ 2 files changed, 49 insertions(+) diff --git a/docs/CUDA_GPU_Direct_Storage.md b/docs/CUDA_GPU_Direct_Storage.md index b1755468deb28..af4ce8b955063 100644 --- a/docs/CUDA_GPU_Direct_Storage.md +++ b/docs/CUDA_GPU_Direct_Storage.md @@ -94,3 +94,5 @@ and initializers assigned to other execution providers retain their existing pat multiple buffers, repeated loads, and different host-memory fallback configurations. These tests require a CUDA GPU, but use the configured fallback when native GDS is unavailable. A passing result alone does not prove native GDS usage or performance. `CudaGdsDriverTest.*` separately checks shared-driver lifetime synchronization without GPU hardware. + +`CApiTest.CUDAProviderOptions*Gds*` checks string-based configuration, invalid values, and serialization round trips. diff --git a/onnxruntime/test/shared_lib/test_inference.cc b/onnxruntime/test/shared_lib/test_inference.cc index 0113854bc81a8..3d4ee2af3d246 100644 --- a/onnxruntime/test/shared_lib/test_inference.cc +++ b/onnxruntime/test/shared_lib/test_inference.cc @@ -4174,6 +4174,53 @@ INSTANTIATE_TEST_SUITE_P(CApiTensorRTTest, CApiTensorRTTest, #ifdef USE_CUDA +TEST(CApiTest, CUDAProviderOptionsGdsRoundTrip) { + constexpr const char* key = "external_data_loader_use_gds"; + Ort::CUDAProviderOptions cuda_options; + cuda_options.Update({}); + EXPECT_EQ((*cuda_options).external_data_loader_use_gds, 0); + + for (const char* value : {"0", "1", "0"}) { + SCOPED_TRACE(value); + cuda_options.Update({{key, value}}); + EXPECT_EQ((*cuda_options).external_data_loader_use_gds, value[0] - '0'); + + const auto serialized = cuda_options.GetCUDAProviderOptionsAsString(); + std::istringstream stream(serialized); + std::unordered_map round_trip_options; + for (std::string entry; std::getline(stream, entry, ';');) { + const auto separator = entry.find('='); + ASSERT_NE(separator, std::string::npos) << entry; + ASSERT_TRUE(round_trip_options.emplace( + entry.substr(0, separator), entry.substr(separator + 1)) + .second); + } + ASSERT_EQ(round_trip_options.at(key), value); + + Ort::CUDAProviderOptions restored; + restored.Update(round_trip_options); + EXPECT_EQ((*restored).external_data_loader_use_gds, value[0] - '0'); + EXPECT_EQ(restored.GetCUDAProviderOptionsAsString(), serialized); + } +} + +#ifndef ORT_NO_EXCEPTIONS +TEST(CApiTest, CUDAProviderOptionsRejectInvalidGdsValue) { + const char* keys[] = {"external_data_loader_use_gds"}; + for (const char* value : {"2", "-1", "invalid", ""}) { + SCOPED_TRACE(value); + Ort::CUDAProviderOptions cuda_options; + cuda_options.Update({{keys[0], "1"}}); + const char* values[] = {value}; + Ort::Status status(Ort::GetApi().UpdateCUDAProviderOptions(cuda_options, keys, values, 1)); + ASSERT_FALSE(status.IsOK()); + const char* expected_error = value[0] == '\0' ? "key/value cannot be empty" : keys[0]; + EXPECT_THAT(status.GetErrorMessage(), testing::HasSubstr(expected_error)); + EXPECT_EQ((*cuda_options).external_data_loader_use_gds, 1); + } +} +#endif + // This test uses CreateCUDAProviderOptions/UpdateCUDAProviderOptions/UpdateCUDAProviderOptionsWithValue APIs to configure and create a CUDA Execution Provider instance TEST(CApiTest, TestConfigureCUDAProviderOptions) { Ort::CUDAProviderOptions cuda_options; From 990945ac50e23bd67a5b008962eef788973129a9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Xavier=20Dupr=C3=A9?= Date: Wed, 23 Sep 2026 18:06:34 +0200 Subject: [PATCH 13/17] Add Windows DirectStorage CUDA loading and three-path benchmarks Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 8af39836-82f0-478b-b1a0-338c518c1828 --- cmake/CMakeLists.txt | 1 + cmake/deps.txt | 1 + cmake/external/directstorage.cmake | 11 + cmake/onnxruntime_providers_cuda.cmake | 9 + docs/FAQ.md | 2 +- docs/Model_Loading_Performance.md | 128 ++++- .../providers/cuda/cuda_provider_options.h | 1 + .../core/framework/session_state_utils.cc | 13 +- onnxruntime/core/platform/env.h | 8 + onnxruntime/core/platform/windows/env.cc | 4 +- .../providers/cuda/cuda_execution_provider.cc | 5 +- .../cuda/cuda_execution_provider_info.cc | 8 + .../cuda/cuda_execution_provider_info.h | 4 +- .../cuda/cuda_external_data_loader.cc | 51 +- .../cuda/cuda_external_data_loader.h | 7 +- ...cuda_external_data_loader_directstorage.cc | 255 +++++++++ .../cuda_external_data_loader_directstorage.h | 36 ++ .../providers/cuda/cuda_provider_factory.cc | 6 + .../core/session/provider_bridge_ort.cc | 8 + .../cuda_external_data_loader_test.cc | 72 ++- .../benchmark_cuda_model_loading.py | 508 ++++++++++++++++-- .../test_benchmark_cuda_model_loading.py | 261 +++++++++ onnxruntime/test/shared_lib/test_inference.cc | 36 ++ 23 files changed, 1354 insertions(+), 81 deletions(-) create mode 100644 cmake/external/directstorage.cmake create mode 100644 onnxruntime/core/providers/cuda/cuda_external_data_loader_directstorage.cc create mode 100644 onnxruntime/core/providers/cuda/cuda_external_data_loader_directstorage.h create mode 100644 onnxruntime/test/python/transformers/test_benchmark_cuda_model_loading.py diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index fa1500904d77f..8832a9145dced 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -70,6 +70,7 @@ option(onnxruntime_ENABLE_PYTHON "Enable python bindings" OFF) option(onnxruntime_ENABLE_MEMLEAK_CHECKER "Experimental: Enable memory leak checker in Windows debug build" OFF) option(onnxruntime_ENABLE_CONVSYMKERNELAVX2_SAT_CHECKER "Experimental: Enable ConvSymKernelAvx2 assembly saturation checker in build" OFF) option(onnxruntime_USE_CUDA "Build with CUDA support" OFF) +cmake_dependent_option(onnxruntime_USE_CUDA_DIRECTSTORAGE "Build Microsoft DirectStorage CUDA loading support" OFF "onnxruntime_USE_CUDA;WIN32" OFF) # Enable ONNX Runtime CUDA EP's internal unit tests that directly access the EP's internal functions instead of through # OpKernels. When the option is ON, we will have two copies of GTest library in the same process. It is not a typical # use. If you hit any problem with that, please do not report it to GTest. Turn OFF the following build option instead. diff --git a/cmake/deps.txt b/cmake/deps.txt index c6ee3c572dc3a..d098e984185dc 100644 --- a/cmake/deps.txt +++ b/cmake/deps.txt @@ -58,6 +58,7 @@ cutlass;https://github.com/NVIDIA/cutlass/archive/refs/tags/v4.7.0.zip;51d4f1ba4 deep_gemm;https://github.com/deepseek-ai/DeepGEMM/archive/559d79fb6994a58b8a15b4b93bf13ccc16edf247.tar.gz;76a0076386991cac8e5d32c3e7e74d9bb8102115 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 +directstorage;https://www.nuget.org/api/v2/package/Microsoft.Direct3D.DirectStorage/1.2.3;be08f099a75c54997a753f444224af96ed0ee3d9 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/v20260828.013844.zip;83f269a6b0c2a5a085c96a4e360d337950740018 dawn_agility_sdk;https://www.nuget.org/api/v2/package/Microsoft.Direct3D.D3D12/1.721.3-preview;fa5f5fc8d0c8c209cfbb530be634960d20595841 diff --git a/cmake/external/directstorage.cmake b/cmake/external/directstorage.cmake new file mode 100644 index 0000000000000..448a4d1fd1b06 --- /dev/null +++ b/cmake/external/directstorage.cmake @@ -0,0 +1,11 @@ +# Copyright (c) Microsoft Corporation. All rights reserved. +# Licensed under the MIT License. + +include_guard(GLOBAL) +onnxruntime_fetchcontent_declare( + directstorage + URL ${DEP_URL_directstorage} + URL_HASH SHA1=${DEP_SHA1_directstorage} + DOWNLOAD_NAME directstorage.zip +) +onnxruntime_fetchcontent_makeavailable(directstorage) diff --git a/cmake/onnxruntime_providers_cuda.cmake b/cmake/onnxruntime_providers_cuda.cmake index f0e37e8aa4c50..923c12819aad7 100644 --- a/cmake/onnxruntime_providers_cuda.cmake +++ b/cmake/onnxruntime_providers_cuda.cmake @@ -236,7 +236,16 @@ # config_cuda_provider_shared_module can be used to config onnxruntime_providers_cuda_obj, onnxruntime_providers_cuda & onnxruntime_providers_cuda_ut. # This function guarantees that all 3 targets have the same configurations. + if(onnxruntime_USE_CUDA_DIRECTSTORAGE) + include(external/directstorage.cmake) + endif() + function(config_cuda_provider_shared_module target) + if(onnxruntime_USE_CUDA_DIRECTSTORAGE) + target_compile_definitions(${target} PRIVATE ORT_CUDA_DIRECTSTORAGE_AVAILABLE) + target_include_directories(${target} PRIVATE "${directstorage_SOURCE_DIR}/native/include") + target_link_libraries(${target} PRIVATE d3d12 dxgi) + endif() if (onnxruntime_REDUCED_OPS_BUILD) add_op_reduction_include_dirs(${target}) endif() diff --git a/docs/FAQ.md b/docs/FAQ.md index 0b90620483c48..a8387edebbb86 100644 --- a/docs/FAQ.md +++ b/docs/FAQ.md @@ -7,7 +7,7 @@ The default CUDA build supports 3 standard quantization operators: QuantizeLinea ## How can I reduce model loading time for large CPU or CUDA models? See [Accelerate model loading](Model_Loading_Performance.md) for parallel CPU weight prepacking and CUDA external-data -loading through GPUDirect Storage or pinned host buffers, including the session and execution provider options that +loading through GPUDirect Storage, Microsoft DirectStorage, or pinned host buffers, including the session and execution provider options that control them. ## How do I change the severity level of the default logger to something other than the default (WARNING)? diff --git a/docs/Model_Loading_Performance.md b/docs/Model_Loading_Performance.md index 90be96a4f03c9..8bb51bf58ad5d 100644 --- a/docs/Model_Loading_Performance.md +++ b/docs/Model_Loading_Performance.md @@ -5,7 +5,7 @@ ONNX Runtime provides two independent loading optimizations for models with larg | Target | Mechanism | Configuration | Default | |---|---|---|---| | CPU prepacking | Run eligible CPU kernel `PrePack()` calls concurrently | `session.prepack.enable_parallel` | Disabled (`"0"`) | -| CUDA external initializers | Read external data with GPUDirect Storage or reusable pinned buffers | CUDA EP options `external_data_loader_use_gds` and `external_data_loader_reading_threads` | GDS disabled; 4 readers | +| CUDA external initializers | Read external data with GPUDirect Storage, Microsoft DirectStorage, or reusable pinned buffers | CUDA EP options `external_data_loader_use_gds`, `external_data_loader_use_directstorage`, and `external_data_loader_reading_threads` | Direct storage disabled; 4 readers | The CPU option is a session configuration entry. The CUDA options are execution provider options passed when the CUDA EP is appended to `SessionOptions`. @@ -50,7 +50,7 @@ The best thread count depends on available CPU cores, memory bandwidth, storage, Models saved with [external data](https://onnx.ai/onnx/repo-docs/ExternalData.html) normally load weights through pageable CPU memory before copying them to the GPU. The CUDA execution provider can instead use NVIDIA GPUDirect -Storage (GDS) or reusable pinned host buffers. +Storage (GDS) on Linux, Microsoft DirectStorage on Windows, or reusable pinned host buffers. ### GPUDirect Storage @@ -92,6 +92,47 @@ initializer uses the configured fallback without disabling GDS for later aligned The build checks for the required cuFile configuration API. Older CUDA toolkits without it remain supported, but enabling GDS in those builds logs a warning and uses the configured host-memory fallback. +### Microsoft DirectStorage (Windows) + +The built-in CUDA execution provider also supports Microsoft's DirectStorage API through D3D12/CUDA +interoperability. Build with `--cmake_extra_defines onnxruntime_USE_CUDA_DIRECTSTORAGE=ON` in addition to the +usual CUDA build options. This opt-in build downloads the pinned DirectStorage SDK headers; it does not +introduce a link-time dependency on `dstorage.dll`. Deploy the SDK's matching x64 `dstorage.dll` and +`dstoragecore.dll` beside the application executable, following Microsoft's +[DirectStorage deployment guidance](https://github.com/microsoft/DirectStorage/blob/main/Docs/DeveloperGuidance.md#sdk-path). +For Python, the application executable is `python.exe`, not the CUDA provider DLL. +The application must make `dstorage.dll` discoverable through the Windows application/system/user DLL search +directories; the current working directory is not searched. + +Set the CUDA EP option `external_data_loader_use_directstorage` to `"1"` to enable this path. +It requires a Windows D3D12-capable NVIDIA adapter that supports CUDA external memory and fence import. +The D3D12 adapter is selected by the CUDA device's LUID, not by assuming that both APIs enumerate GPUs +in the same order. Linked D3D12 adapters are not supported. + +```text +external-data file -> DirectStorage -> shared D3D12 GPU buffer -> CUDA arena initializer + internal staging CUDA external memory device-to-device copy +``` + +The loader reuses a 32 MiB shared GPU buffer and a DirectStorage queue. A shared D3D12 fence establishes +completion and visibility to CUDA; request errors are checked before copying bytes into the initializer. +CUDA copies complete before the next DirectStorage write reuses the buffer. The additional GPU buffer +does not include DirectStorage's own internal staging allocations. This implementation loads uncompressed +ONNX external data, not GDeflate-compressed weights. + +**DirectStorage is not NVIDIA GPUDirect Storage:** Microsoft's uncompressed data flow can use system-memory +and upload-heap staging. Successful DirectStorage loading does not establish zero-copy storage-to-VRAM DMA +or prove that Windows BypassIO was used. Compare actual timings rather than assuming the API is faster. +See Microsoft's [uncompressed data flow documentation](https://github.com/microsoft/DirectStorage/blob/main/Docs/DeveloperGuidance.md#uncompressed-data-flow). + +DirectStorage accepts unaligned offsets and lengths. Because its API opens files by path, the loader compares +the opened file's identity with the handle already validated by ONNX Runtime before submitting reads. +Initialization or read failures produce a warning and disable DirectStorage for the remainder of that loader's +lifetime, using the configured host-memory fallback. Boolean tensors retain the host conversion path. +If both direct-storage options are enabled, DirectStorage is tried first, followed by GDS, then the host path; +normally enable only the option appropriate for the operating system. Unsupported builds/platforms report +unavailability and use the fallback instead of silently claiming DirectStorage support. + ### Pinned-buffer host loading The CUDA execution provider can load external initializers through two reusable 64 MiB pinned host buffers: @@ -110,19 +151,23 @@ The following CUDA execution provider options control the primary and fallback p | Option | Values | Default | Purpose | |---|---|---:|---| | `external_data_loader_use_gds` | `0` or `1` | `0` | Try GDS before another external-data loading path | +| `external_data_loader_use_directstorage` | `0` or `1` | `0` | Try Microsoft DirectStorage through D3D12/CUDA on Windows | | `external_data_loader_reading_threads` | `0` to `64` | `4` | Configure the pinned-buffer fallback; `0` disables it | -Keep `external_data_loader_reading_threads` greater than zero when enabling GDS to retain the pinned-buffer backup. +Keep `external_data_loader_reading_threads` greater than zero when enabling either direct-storage option to retain +the pinned-buffer backup. `1` uses synchronous reads into pinned memory. Values from `2` through `64` use that many parallel read tasks per 64 MiB pinned buffer. The default is `4`, so the pinned-buffer loader is enabled without additional configuration. Parallel reads are used for external tensors of at least 16 MiB; smaller tensors use one read. The optimal reader count depends on the storage device and filesystem. If pinned buffers or CUDA streams cannot be created, loading falls back -to a pageable buffer. If the value is `0`, the pageable path is used when GDS is disabled or unavailable. Models with +to a pageable buffer. If the value is `0`, the pageable path is used when direct storage is disabled or unavailable. Models with weights embedded in the ONNX file do not use this external-data path. Configure the CUDA EP in Python: ```python +import sys + import onnxruntime as ort session_options = ort.SessionOptions() @@ -130,7 +175,8 @@ providers = [ ( "CUDAExecutionProvider", { - "external_data_loader_use_gds": "1", + "external_data_loader_use_gds": "0" if sys.platform == "win32" else "1", + "external_data_loader_use_directstorage": "1" if sys.platform == "win32" else "0", "external_data_loader_reading_threads": "4", }, ), @@ -153,7 +199,11 @@ session_options.AddConfigEntry("session.prepack.enable_parallel", "1"); Ort::CUDAProviderOptions cuda_options; cuda_options.Update({ +#ifdef _WIN32 + {"external_data_loader_use_directstorage", "1"}, +#else {"external_data_loader_use_gds", "1"}, +#endif {"external_data_loader_reading_threads", "4"}, }); session_options.AppendExecutionProvider_CUDA_V2(*cuda_options); @@ -174,3 +224,71 @@ but use the configured fallback when native GDS is unavailable. A passing result or performance. `CudaGdsDriverTest.*` separately checks shared-driver lifetime synchronization without GPU hardware. `CApiTest.CUDAProviderOptions*Gds*` checks string-based configuration, invalid values, and serialization round trips. + +`CudaExternalDataLoaderTest.*DirectStorage*` covers option validation, aligned and unaligned reads, repeated +loads, multiple buffers, and configured fallbacks. `NativeDirectStorageWithoutFallback` calls the DirectStorage +backend directly, verifies loaded bytes, rejects out-of-range and mismatched-file requests, and cannot succeed +by using the host fallback. It explicitly skips when D3D12/CUDA/DirectStorage initialization is unavailable. +`CApiTest.CUDAProviderOptionsDirectStorageRoundTrip` covers the string-based option. + +Successful loads emit INFO records of the form `CUDA external data loader: path= bytes=`, +where `` is `pageable`, `pinned`, `gds`, or `directstorage`. Enable the **default** logger's INFO severity +(`ort.set_default_logger_severity(1)` in Python) as well as the session logger when collecting these records. +They identify the path actually used, including fallbacks; requested provider options alone are not proof. + +### Comparing all three loading paths + +`onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py` compares pageable CPU staging +(the loader-disabled baseline), pinned buffers, and the platform's direct-storage API, all targeting the +same CUDA device. Use the Python package from the build being evaluated, not an installed older wheel. +On Windows, enable the DirectStorage build option and deploy the SDK runtime as described above. + +For example, generate identical aligned external weights totaling 1 GiB and collect five measurements per path: + +```powershell +python onnxruntime\test\python\transformers\benchmark_cuda_model_loading.py ` + --generate-model .\cuda-loading-1gib --weight-count 16 --weight-dim 4096 ` + --threads 1 --repetitions 5 --output .\cuda-loading-1gib.json +``` + +Use `--weight-count 64` and a different output directory for 4 GiB of weights, provided enough GPU memory +is available. Generation refuses to overwrite an existing fixture. To benchmark a real model instead, supply +`--model model.onnx --inputs inputs.npz --expected-outputs expected.npz`; the NPZ keys must match tensor names, +and reference outputs must be computed independently. + +Each sample uses a fresh process. The timed interval covers session construction, including graph initialization +and completed weight transfers, but not imports, fixture generation, or output verification. A blocking inference +and comparison against independent reference values happen afterward. Reported GiB/s is therefore **effective +end-to-end initialization throughput**, not raw disk bandwidth. The JSON report includes timing distributions, +GPU/configuration metadata, actual loaded bytes by path, and separately classified fallback or mixed-path samples. +A direct-storage fallback is never counted as a successful direct-storage measurement. + +Caches remain OS-managed: fixture creation and warmups may warm the filesystem cache, and a fresh process does +not imply a cold disk. The optional POSIX per-file eviction hint is also not proof of cold-cache operation. +The script never performs privileged or global cache flushing. Record storage/filesystem details alongside results, +and do not compare Windows DirectStorage and Linux GDS numbers as if they came from an identical software stack. + +#### Windows measurements (2026-09-23) + +Measured on Windows 11 (build 26200), an NVIDIA RTX 4060 Laptop GPU (8 GiB, WDDM, driver 591.55), +and a local WD `SDCPNRZ-2T00-1124-WD` 2 TB SSD. The source build used MSVC 2022, CUDA 13.0.2, +DirectStorage 1.2.3, and CPython 3.13.14. It was a Release/quick build restricted to SM89 and +MatMul/Gather registrations, with contrib support enabled. The fixture used FP32 4096-by-4096 weights, +one intra-op thread, four pinned-buffer readers, disabled graph optimization/prepacking, and disabled TF32. + +Each row summarizes five fresh-process samples after one warmup per path, with OS-managed (potentially warm) +caches. The complete weight-byte counts were confirmed from actual-path logs and every sample's outputs +matched independent references. DirectStorage rows used the DirectStorage API, not the pinned fallback. + +| External weights | Loading path | Median initialization | Effective throughput | +|---|---|---:|---:| +| 1 GiB | CPU pageable | 1.517 s | 0.659 GiB/s | +| 1 GiB | Pinned buffers | 1.105 s | 0.905 GiB/s | +| 1 GiB | Microsoft DirectStorage | 1.853 s | 0.540 GiB/s | +| 4 GiB | CPU pageable | 6.201 s | 0.645 GiB/s | +| 4 GiB | Pinned buffers | 2.292 s | 1.745 GiB/s | +| 4 GiB | Microsoft DirectStorage | 7.870 s | 0.508 GiB/s | + +On this machine and workload, pinned buffers outperform both pageable loading and the current uncompressed +DirectStorage implementation. These are complete session-initialization measurements, not isolated I/O timings, +and are not evidence of cold-cache disk bandwidth or Linux GDS performance. diff --git a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h index 5d19f97ee1f9a..36433d0405234 100644 --- a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h +++ b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h @@ -44,4 +44,5 @@ struct OrtCUDAProviderOptionsV2 { int sdpa_kernel = 0; // Scaled Dot Product Attention kernel option size_t external_data_loader_reading_threads = 4; // Number of CPU read tasks per external-data staging buffer. 0 disables pinned-buffer fallback; 1 disables parallel reads. int external_data_loader_use_gds = 0; // Try GPUDirect Storage before the configured pinned/pageable host-memory fallback. + int external_data_loader_use_directstorage = 0; // Try Microsoft DirectStorage on Windows before the configured host-memory fallback. }; diff --git a/onnxruntime/core/framework/session_state_utils.cc b/onnxruntime/core/framework/session_state_utils.cc index db988873539ad..2b02480fe1112 100644 --- a/onnxruntime/core/framework/session_state_utils.cc +++ b/onnxruntime/core/framework/session_state_utils.cc @@ -149,11 +149,16 @@ static common::Status DeserializeTensorProto(const Env& env, const std::basic_st default_cpu_alloc, normalized_cpu_tensor)); utils::MakeCpuTensorCopy(cpu_staging_tensor, normalized_cpu_tensor); utils::NormalizeBoolTensorIfNeeded(normalized_cpu_tensor); - return CopyTensorFromCPUToDevice(data_transfer_mgr, normalized_cpu_tensor, std::move(tensor), ort_value); + ORT_RETURN_IF_ERROR( + CopyTensorFromCPUToDevice(data_transfer_mgr, normalized_cpu_tensor, std::move(tensor), ort_value)); + } else { + ORT_RETURN_IF_ERROR(CopyTensorFromCPUToDevice(data_transfer_mgr, deserialized_value.Get(), + std::move(tensor), ort_value)); } - - return CopyTensorFromCPUToDevice(data_transfer_mgr, deserialized_value.Get(), - std::move(tensor), ort_value); + if (device.Type() == OrtDevice::GPU && device.Vendor() == OrtDevice::VendorIds::NVIDIA) { + LOGS_DEFAULT(INFO) << "CUDA external data loader: path=pageable bytes=" << cpu_staging_tensor.SizeInBytes(); + } + return Status::OK(); } } else { if (device == default_cpu_device) { diff --git a/onnxruntime/core/platform/env.h b/onnxruntime/core/platform/env.h index de6eb1d083bca..dccd1ec497c14 100644 --- a/onnxruntime/core/platform/env.h +++ b/onnxruntime/core/platform/env.h @@ -145,6 +145,14 @@ class PosixFileDescriptorProvider { virtual int GetFileDescriptor() const = 0; }; +class WindowsFileHandleProvider { + public: + virtual ~WindowsFileHandleProvider() = default; + + // The handle remains owned by the provider and is valid only for its lifetime. + virtual void* GetFileHandle() const = 0; +}; + /// \brief An interface used by the onnxruntime implementation to /// access operating system functionality like the filesystem etc. /// diff --git a/onnxruntime/core/platform/windows/env.cc b/onnxruntime/core/platform/windows/env.cc index 33f8b3e20994d..1131c7cdf7874 100644 --- a/onnxruntime/core/platform/windows/env.cc +++ b/onnxruntime/core/platform/windows/env.cc @@ -360,11 +360,13 @@ common::Status WindowsEnv::GetFileLength(int fd, /*out*/ size_t& file_size) cons namespace { -class WindowsRandomAccessFile final : public RandomAccessFile { +class WindowsRandomAccessFile final : public RandomAccessFile, public WindowsFileHandleProvider { public: explicit WindowsRandomAccessFile(wil::unique_hfile file_handle) : file_handle_(std::move(file_handle)) {} ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(WindowsRandomAccessFile); + void* GetFileHandle() const override { return file_handle_.get(); } + Status GetLength(size_t& length) const override { LARGE_INTEGER file_size{}; if (!GetFileSizeEx(file_handle_.get(), &file_size)) { diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 007d714fdbfa4..be973b5a8d578 100755 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -3434,13 +3434,14 @@ std::unique_ptr CUDAExecutionProvider::GetDataTransf std::unique_ptr CUDAExecutionProvider::GetExternalDataLoader() const { if (info_.external_data_loader_reading_threads == 0 && - !info_.external_data_loader_use_gds) { + !info_.external_data_loader_use_gds && + !info_.external_data_loader_use_directstorage) { return nullptr; } return std::make_unique( info_.device_id, info_.external_data_loader_reading_threads, - info_.external_data_loader_use_gds); + info_.external_data_loader_use_gds, info_.external_data_loader_use_directstorage); } std::vector> diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc index 5dedcad1b711e..9a961e35b0680 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc @@ -40,6 +40,7 @@ constexpr const char* kFuseConvBias = "fuse_conv_bias"; constexpr const char* kSdpaKernel = "sdpa_kernel"; constexpr const char* kExternalDataLoaderReadingThreads = "external_data_loader_reading_threads"; constexpr const char* kExternalDataLoaderUseGds = "external_data_loader_use_gds"; +constexpr const char* kExternalDataLoaderUseDirectStorage = "external_data_loader_use_directstorage"; } // namespace provider_option_names } // namespace cuda @@ -150,6 +151,9 @@ CUDAExecutionProviderInfo CUDAExecutionProviderInfo::FromProviderOptions(const P .AddAssignmentToReference( cuda::provider_option_names::kExternalDataLoaderUseGds, info.external_data_loader_use_gds) + .AddAssignmentToReference( + cuda::provider_option_names::kExternalDataLoaderUseDirectStorage, + info.external_data_loader_use_directstorage) .AddValueParser( cuda::provider_option_names::kTunableOpEnable, [&info](const std::string& value_str) -> Status { @@ -209,6 +213,8 @@ ProviderOptions CUDAExecutionProviderInfo::ToProviderOptions(const CUDAExecution MakeStringWithClassicLocale(info.external_data_loader_reading_threads)}, {cuda::provider_option_names::kExternalDataLoaderUseGds, MakeStringWithClassicLocale(info.external_data_loader_use_gds)}, + {cuda::provider_option_names::kExternalDataLoaderUseDirectStorage, + MakeStringWithClassicLocale(info.external_data_loader_use_directstorage)}, }; return options; @@ -238,6 +244,8 @@ ProviderOptions CUDAExecutionProviderInfo::ToProviderOptions(const OrtCUDAProvid MakeStringWithClassicLocale(info.external_data_loader_reading_threads)}, {cuda::provider_option_names::kExternalDataLoaderUseGds, MakeStringWithClassicLocale(info.external_data_loader_use_gds)}, + {cuda::provider_option_names::kExternalDataLoaderUseDirectStorage, + MakeStringWithClassicLocale(info.external_data_loader_use_directstorage)}, }; return options; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h index 796217a2e7b36..2a6826e926b84 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h @@ -84,10 +84,11 @@ struct CUDAExecutionProviderInfo { int sdpa_kernel{0}; - // 0 disables the pinned-buffer loader and uses the pageable fallback if GDS is unavailable. + // 0 disables the pinned-buffer loader and uses the pageable fallback if direct storage is unavailable. // 1 uses the pinned-buffer loader with synchronous reads. 2..64 use that many parallel read tasks per block. size_t external_data_loader_reading_threads{4}; bool external_data_loader_use_gds{false}; + bool external_data_loader_use_directstorage{false}; static CUDAExecutionProviderInfo FromProviderOptions(const ProviderOptions& options); static ProviderOptions ToProviderOptions(const CUDAExecutionProviderInfo& info); @@ -124,6 +125,7 @@ struct std::hash<::onnxruntime::CUDAExecutionProviderInfo> { onnxruntime::HashCombine(info.enable_cudnn, value); onnxruntime::HashCombine(info.external_data_loader_reading_threads, value); onnxruntime::HashCombine(info.external_data_loader_use_gds, value); + onnxruntime::HashCombine(info.external_data_loader_use_directstorage, value); // Memory pointers onnxruntime::HashCombine(reinterpret_cast(info.user_compute_stream), value); diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc index 3976275637bd2..94b2fc049b3c4 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc @@ -116,10 +116,12 @@ common::Status LoadWithPageableBuffer(const RandomAccessFile& file, FileOffsetTy } // namespace -ExternalDataLoader::ExternalDataLoader(int device_id, size_t reading_thread_count, bool use_gds) +ExternalDataLoader::ExternalDataLoader(int device_id, size_t reading_thread_count, bool use_gds, + bool use_directstorage) : device_id_(device_id), reading_thread_count_(reading_thread_count), - use_gds_(use_gds) {} + use_gds_(use_gds), + use_directstorage_(use_directstorage) {} ExternalDataLoader::~ExternalDataLoader() { reader_pool_.reset(); @@ -164,6 +166,7 @@ void ExternalDataLoader::ReleaseResources() const noexcept { cudaSetDevice(device_id_) == cudaSuccess; gds_loader_.reset(); + directstorage_loader_.reset(); for (auto& stream : streams_) { if (stream != nullptr) { @@ -210,6 +213,32 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, CudaDeviceGuard device_guard; ORT_RETURN_IF_ERROR(device_guard.SetDevice(device_id_)); + if (use_directstorage_ && !directstorage_disabled_ && length != 0 && + std::endian::native == std::endian::little && !tensor.IsDataType()) { + Status status = Status::OK(); + if (!directstorage_loader_) { + status = DirectStorageLoader::Create(device_id_, directstorage_loader_); + } + if (status.IsOK()) { + void* handle = nullptr; +#if !defined(ORT_NO_RTTI) + const auto* provider = dynamic_cast(file.get()); + if (provider != nullptr) { + handle = provider->GetFileHandle(); + } +#endif + status = directstorage_loader_->Load(data_file_path, handle, data_offset, length, tensor); + } + if (status.IsOK()) { + LOGS_DEFAULT(INFO) << "CUDA external data loader: path=directstorage bytes=" << length; + return Status::OK(); + } + directstorage_disabled_ = true; + directstorage_loader_.reset(); + LOGS_DEFAULT(WARNING) << "Microsoft DirectStorage could not load external data; falling back to another CUDA " + << "external-data loading path. " << status.ErrorMessage(); + } + const bool gds_range_is_aligned = data_offset % static_cast(kGdsIoAlignment) == 0 && length % kGdsIoAlignment == 0; @@ -232,6 +261,7 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, file_descriptor, data_offset, length, tensor); } if (gds_status.IsOK()) { + LOGS_DEFAULT(INFO) << "CUDA external data loader: path=gds bytes=" << length; return Status::OK(); } @@ -244,14 +274,19 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, } if (reading_thread_count_ == 0) { - return LoadWithPageableBuffer(*file, data_offset, length, tensor, 1, reader_pool_); + ORT_RETURN_IF_ERROR(LoadWithPageableBuffer(*file, data_offset, length, tensor, 1, reader_pool_)); + LOGS_DEFAULT(INFO) << "CUDA external data loader: path=pageable bytes=" << length; + return Status::OK(); } const auto resource_status = EnsureResources(); if (!resource_status.IsOK()) { - // TODO: Remember setup failures during initialization and report the first CUDA error - // so later initializers do not repeatedly retry unavailable pinned buffers or streams. - return LoadWithPageableBuffer(*file, data_offset, length, tensor, reading_thread_count_, reader_pool_); + LOGS_DEFAULT(WARNING) << "CUDA pinned-buffer setup failed; falling back to pageable memory. " + << resource_status.ErrorMessage(); + ORT_RETURN_IF_ERROR( + LoadWithPageableBuffer(*file, data_offset, length, tensor, reading_thread_count_, reader_pool_)); + LOGS_DEFAULT(INFO) << "CUDA external data loader: path=pageable bytes=" << length; + return Status::OK(); } auto* destination = static_cast(tensor.MutableDataRaw()); @@ -315,7 +350,9 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, offset += chunk_size; } - return synchronize_streams(); + ORT_RETURN_IF_ERROR(synchronize_streams()); + LOGS_DEFAULT(INFO) << "CUDA external data loader: path=pinned bytes=" << length; + return Status::OK(); } } // namespace cuda diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader.h index 27f04ebd76882..06978072ae21a 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.h +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.h @@ -9,6 +9,7 @@ #include "core/framework/external_data_loader.h" #include "core/providers/cuda/cuda_external_data_loader_gds.h" +#include "core/providers/cuda/cuda_external_data_loader_directstorage.h" #include "cuda_pch.h" namespace onnxruntime { @@ -65,7 +66,8 @@ class ExternalDataLoaderThreadPool; */ class ExternalDataLoader final : public IExternalDataLoader { public: - ExternalDataLoader(int device_id, size_t reading_thread_count, bool use_gds = false); + ExternalDataLoader(int device_id, size_t reading_thread_count, bool use_gds = false, + bool use_directstorage = false); ~ExternalDataLoader() override; bool CanLoad(const OrtMemoryInfo& target_memory_info) const override; @@ -88,6 +90,9 @@ class ExternalDataLoader final : public IExternalDataLoader { const bool use_gds_; mutable bool gds_disabled_{false}; mutable std::unique_ptr gds_loader_; + const bool use_directstorage_; + mutable bool directstorage_disabled_{false}; + mutable std::unique_ptr directstorage_loader_; mutable std::unique_ptr reader_pool_; }; diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_directstorage.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader_directstorage.cc new file mode 100644 index 0000000000000..88b9155702ed4 --- /dev/null +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_directstorage.cc @@ -0,0 +1,255 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +// provider_api.h must be first to set SHARED_PROVIDER +#include "core/providers/shared_library/provider_api.h" + +#include "core/providers/cuda/cuda_external_data_loader_directstorage.h" + +#if defined(ORT_CUDA_DIRECTSTORAGE_AVAILABLE) +#include +#include +#include +#include +#include +#include + +#include "core/common/safeint.h" +#include "core/providers/cuda/cuda_common.h" +#endif + +namespace onnxruntime { +namespace cuda { +namespace { + +#if defined(ORT_CUDA_DIRECTSTORAGE_AVAILABLE) + +using Microsoft::WRL::ComPtr; +constexpr size_t kDirectStorageBufferSize = 32 * 1024 * 1024; + +common::Status CheckHResult(HRESULT result, const char* operation) { + ORT_RETURN_IF(FAILED(result), operation, " failed, HRESULT=", static_cast(result)); + return Status::OK(); +} + +class DirectStorageLibrary { + public: + DirectStorageLibrary() = default; + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(DirectStorageLibrary); + ~DirectStorageLibrary() { + if (handle != nullptr) { + FreeLibrary(handle); + } + } + HMODULE handle{nullptr}; +}; + +class WindowsDirectStorageLoader final : public DirectStorageLoader { + public: + ~WindowsDirectStorageLoader() override { + if (stream_ != nullptr) { + ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaStreamSynchronize(stream_))); + } + if (queue_) { + queue_->Close(); + } + if (stream_ != nullptr) { + ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaStreamDestroy(stream_))); + } + if (cuda_fence_ != nullptr) { + ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaDestroyExternalSemaphore(cuda_fence_))); + } + if (buffer_ != nullptr) { + ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaFree(buffer_))); + } + if (cuda_memory_ != nullptr) { + ORT_IGNORE_RETURN_VALUE(CUDA_CALL(cudaDestroyExternalMemory(cuda_memory_))); + } + } + + static common::Status Create(int device_id, std::unique_ptr& loader) { + auto candidate = std::unique_ptr(new WindowsDirectStorageLoader()); + ORT_RETURN_IF_ERROR(candidate->Initialize(device_id)); + loader = std::move(candidate); + return Status::OK(); + } + + common::Status Load(const std::filesystem::path& path, void* validated_file_handle, + int64_t data_offset, size_t data_length, Tensor& tensor) override { + ORT_RETURN_IF(validated_file_handle == nullptr, + "Microsoft DirectStorage requires the validated Windows file handle."); + BY_HANDLE_FILE_INFORMATION original{}; + ORT_RETURN_IF_NOT(GetFileInformationByHandle(validated_file_handle, &original), + "GetFileInformationByHandle failed: ", GetLastError()); + ComPtr file; + ORT_RETURN_IF_ERROR(CheckHResult(factory_->OpenFile(path.c_str(), IID_PPV_ARGS(&file)), + "DirectStorage OpenFile")); + bool pending = false; + auto drain_on_error = gsl::finally([&]() { + if (pending) { + queue_->Close(); + } + }); + BY_HANDLE_FILE_INFORMATION opened{}; + ORT_RETURN_IF_ERROR(CheckHResult(file->GetFileInformation(&opened), "DirectStorage GetFileInformation")); + // DirectStorage opens by path. Never read a replacement for the file validated by the caller. + ORT_RETURN_IF(original.dwVolumeSerialNumber != opened.dwVolumeSerialNumber || + original.nFileIndexHigh != opened.nFileIndexHigh || + original.nFileIndexLow != opened.nFileIndexLow, + "External-data file changed before DirectStorage opened it."); + const uint64_t file_size = (static_cast(opened.nFileSizeHigh) << 32) | opened.nFileSizeLow; + ORT_RETURN_IF(data_offset < 0 || static_cast(data_offset) > file_size || + data_length > file_size - static_cast(data_offset), + "DirectStorage external-data range is outside the file."); + + auto* destination = static_cast(tensor.MutableDataRaw()); + for (size_t offset = 0; offset < data_length;) { + const auto chunk = static_cast(std::min(kDirectStorageBufferSize, data_length - offset)); + DSTORAGE_REQUEST request{}; + request.Options.SourceType = DSTORAGE_REQUEST_SOURCE_FILE; + request.Options.DestinationType = DSTORAGE_REQUEST_DESTINATION_BUFFER; + request.Source.File.Source = file.Get(); + request.Source.File.Offset = SafeInt(data_offset) + offset; + request.Source.File.Size = chunk; + request.UncompressedSize = chunk; + request.Destination.Buffer.Resource = resource_.Get(); + request.Destination.Buffer.Size = chunk; + queue_->EnqueueRequest(&request); + queue_->EnqueueStatus(status_.Get(), 0); + queue_->EnqueueSignal(fence_.Get(), ++fence_value_); + queue_->Submit(); + pending = true; + + cudaExternalSemaphoreWaitParams wait{}; + wait.params.fence.value = fence_value_; + CUDA_RETURN_IF_ERROR(cudaWaitExternalSemaphoresAsync(&cuda_fence_, &wait, 1, stream_)); + CUDA_RETURN_IF_ERROR(cudaStreamSynchronize(stream_)); + pending = false; + ORT_RETURN_IF_ERROR(CheckHResult(status_->GetHResult(0), "DirectStorage read")); + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(destination + offset, buffer_, chunk, + cudaMemcpyDeviceToDevice, stream_)); + // Complete CUDA reads before DirectStorage writes the shared buffer again. + CUDA_RETURN_IF_ERROR(cudaStreamSynchronize(stream_)); + offset += chunk; + } + return Status::OK(); + } + + private: + WindowsDirectStorageLoader() = default; + + common::Status Initialize(int device_id) { + CUDA_RETURN_IF_ERROR(cudaSetDevice(device_id)); + cudaDeviceProp properties{}; + CUDA_RETURN_IF_ERROR(cudaGetDeviceProperties(&properties, device_id)); + const unsigned int node_mask = properties.luidDeviceNodeMask; + ORT_RETURN_IF(node_mask == 0 || (node_mask & (node_mask - 1)) != 0, + "DirectStorage requires a single CUDA device node."); + LUID adapter_luid{}; + static_assert(sizeof(properties.luid) == sizeof(adapter_luid)); + std::memcpy(&adapter_luid, properties.luid, sizeof(adapter_luid)); + ComPtr dxgi; + ORT_RETURN_IF_ERROR(CheckHResult(CreateDXGIFactory1(IID_PPV_ARGS(&dxgi)), "CreateDXGIFactory1")); + ComPtr adapter; + ORT_RETURN_IF_ERROR(CheckHResult(dxgi->EnumAdapterByLuid(adapter_luid, IID_PPV_ARGS(&adapter)), + "Find CUDA DXGI adapter")); + ORT_RETURN_IF_ERROR(CheckHResult( + D3D12CreateDevice(adapter.Get(), D3D_FEATURE_LEVEL_11_0, IID_PPV_ARGS(&device_)), "D3D12CreateDevice")); + ORT_RETURN_IF(device_->GetNodeCount() != 1, "DirectStorage does not support linked D3D12 adapters."); + + library_.handle = LoadLibraryExW(L"dstorage.dll", nullptr, LOAD_LIBRARY_SEARCH_DEFAULT_DIRS); + ORT_RETURN_IF(library_.handle == nullptr, "Microsoft DirectStorage is unavailable: dstorage.dll load failed: ", + GetLastError()); + const auto get_factory = + reinterpret_cast(GetProcAddress(library_.handle, "DStorageGetFactory")); + ORT_RETURN_IF(get_factory == nullptr, "DStorageGetFactory is unavailable: ", GetLastError()); + ORT_RETURN_IF_ERROR(CheckHResult(get_factory(IID_PPV_ARGS(&factory_)), "DStorageGetFactory")); + // The default DirectStorage staging size is 32 MiB. Do not reconfigure its process-wide factory. + DSTORAGE_QUEUE_DESC queue_desc{}; + queue_desc.SourceType = DSTORAGE_REQUEST_SOURCE_FILE; + queue_desc.Capacity = DSTORAGE_MIN_QUEUE_CAPACITY; + queue_desc.Priority = DSTORAGE_PRIORITY_NORMAL; + queue_desc.Device = device_.Get(); + ORT_RETURN_IF_ERROR(CheckHResult(factory_->CreateQueue(&queue_desc, IID_PPV_ARGS(&queue_)), + "DirectStorage CreateQueue")); + ORT_RETURN_IF_ERROR(CheckHResult(factory_->CreateStatusArray(1, "ORT external data", IID_PPV_ARGS(&status_)), + "DirectStorage CreateStatusArray")); + + D3D12_HEAP_PROPERTIES heap{}; + heap.Type = D3D12_HEAP_TYPE_DEFAULT; + heap.CreationNodeMask = node_mask; + heap.VisibleNodeMask = node_mask; + D3D12_RESOURCE_DESC desc{}; + desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER; + desc.Width = kDirectStorageBufferSize; + desc.Height = 1; + desc.DepthOrArraySize = 1; + desc.MipLevels = 1; + desc.SampleDesc.Count = 1; + desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR; + ORT_RETURN_IF_ERROR(CheckHResult( + device_->CreateCommittedResource(&heap, D3D12_HEAP_FLAG_SHARED, &desc, + D3D12_RESOURCE_STATE_COMMON, nullptr, IID_PPV_ARGS(&resource_)), + "Create shared DirectStorage buffer")); + HANDLE shared_memory = nullptr; + ORT_RETURN_IF_ERROR(CheckHResult( + device_->CreateSharedHandle(resource_.Get(), nullptr, GENERIC_ALL, nullptr, &shared_memory), + "Share DirectStorage buffer")); + auto close_memory = gsl::finally([&]() { CloseHandle(shared_memory); }); + cudaExternalMemoryHandleDesc memory_desc{}; + memory_desc.type = cudaExternalMemoryHandleTypeD3D12Resource; + memory_desc.handle.win32.handle = shared_memory; + memory_desc.size = device_->GetResourceAllocationInfo(node_mask, 1, &desc).SizeInBytes; + memory_desc.flags = cudaExternalMemoryDedicated; + CUDA_RETURN_IF_ERROR(cudaImportExternalMemory(&cuda_memory_, &memory_desc)); + cudaExternalMemoryBufferDesc buffer_desc{}; + buffer_desc.size = kDirectStorageBufferSize; + CUDA_RETURN_IF_ERROR(cudaExternalMemoryGetMappedBuffer(&buffer_, cuda_memory_, &buffer_desc)); + + ORT_RETURN_IF_ERROR(CheckHResult( + device_->CreateFence(0, D3D12_FENCE_FLAG_SHARED, IID_PPV_ARGS(&fence_)), "Create DirectStorage fence")); + HANDLE shared_fence = nullptr; + ORT_RETURN_IF_ERROR(CheckHResult( + device_->CreateSharedHandle(fence_.Get(), nullptr, GENERIC_ALL, nullptr, &shared_fence), + "Share DirectStorage fence")); + auto close_fence = gsl::finally([&]() { CloseHandle(shared_fence); }); + cudaExternalSemaphoreHandleDesc fence_desc{}; + fence_desc.type = cudaExternalSemaphoreHandleTypeD3D12Fence; + fence_desc.handle.win32.handle = shared_fence; + CUDA_RETURN_IF_ERROR(cudaImportExternalSemaphore(&cuda_fence_, &fence_desc)); + CUDA_RETURN_IF_ERROR(cudaStreamCreateWithFlags(&stream_, cudaStreamNonBlocking)); + return Status::OK(); + } + + DirectStorageLibrary library_; + ComPtr device_; + ComPtr factory_; + ComPtr resource_; + ComPtr fence_; + ComPtr status_; + ComPtr queue_; + cudaExternalMemory_t cuda_memory_{nullptr}; + cudaExternalSemaphore_t cuda_fence_{nullptr}; + cudaStream_t stream_{nullptr}; + void* buffer_{nullptr}; + uint64_t fence_value_{0}; +}; + +#endif + +} // namespace + +common::Status DirectStorageLoader::Create(int device_id, std::unique_ptr& loader) { +#if defined(ORT_CUDA_DIRECTSTORAGE_AVAILABLE) + return WindowsDirectStorageLoader::Create(device_id, loader); +#else + ORT_UNUSED_PARAMETER(device_id); + ORT_UNUSED_PARAMETER(loader); + return ORT_MAKE_STATUS(ONNXRUNTIME, NOT_IMPLEMENTED, + "Microsoft DirectStorage requires a Windows CUDA build with " + "onnxruntime_USE_CUDA_DIRECTSTORAGE=ON."); +#endif +} + +} // namespace cuda +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader_directstorage.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader_directstorage.h new file mode 100644 index 0000000000000..9821a46427adb --- /dev/null +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader_directstorage.h @@ -0,0 +1,36 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include +#include +#include +#include + +#include "core/common/common.h" +#include "core/common/status.h" + +namespace onnxruntime { +#ifndef SHARED_PROVIDER +class Tensor; +#endif + +namespace cuda { + +class DirectStorageLoader { + public: + virtual ~DirectStorageLoader() = default; + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(DirectStorageLoader); + + virtual common::Status Load(const std::filesystem::path& path, void* validated_file_handle, + int64_t data_offset, size_t data_length, Tensor& tensor) = 0; + + static common::Status Create(int device_id, std::unique_ptr& loader); + + protected: + DirectStorageLoader() = default; +}; + +} // namespace cuda +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc index bc216a792a481..ffee8f37c18eb 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc @@ -213,6 +213,9 @@ struct CUDA_Provider : Provider { ORT_ENFORCE(params->external_data_loader_use_gds == 0 || params->external_data_loader_use_gds == 1, "external_data_loader_use_gds must be 0 or 1."); + ORT_ENFORCE(params->external_data_loader_use_directstorage == 0 || + params->external_data_loader_use_directstorage == 1, + "external_data_loader_use_directstorage must be 0 or 1."); // Calling a function like ::cudaDeviceSynchronize will cause CUDA to ensure there is binary code for the current GPU architecture // Ideally this will be already part of the binary, but if not, CUDA will JIT it during this call. This can take a very long time @@ -255,6 +258,7 @@ struct CUDA_Provider : Provider { info.sdpa_kernel = params->sdpa_kernel; info.external_data_loader_reading_threads = params->external_data_loader_reading_threads; info.external_data_loader_use_gds = params->external_data_loader_use_gds != 0; + info.external_data_loader_use_directstorage = params->external_data_loader_use_directstorage != 0; return std::make_shared(info); } @@ -294,6 +298,8 @@ struct CUDA_Provider : Provider { internal_options.external_data_loader_reading_threads; cuda_options.external_data_loader_use_gds = internal_options.external_data_loader_use_gds; + cuda_options.external_data_loader_use_directstorage = + internal_options.external_data_loader_use_directstorage; } ProviderOptions GetProviderOptions(const void* provider_options) override { diff --git a/onnxruntime/core/session/provider_bridge_ort.cc b/onnxruntime/core/session/provider_bridge_ort.cc index c686485dcfafa..16720feb4400b 100644 --- a/onnxruntime/core/session/provider_bridge_ort.cc +++ b/onnxruntime/core/session/provider_bridge_ort.cc @@ -2949,6 +2949,14 @@ ORT_API(void, OrtApis::ReleaseTensorRTProviderOptions, _Frees_ptr_opt_ OrtTensor ORT_API_STATUS_IMPL(OrtApis::SessionOptionsAppendExecutionProvider_CUDA_V2, _In_ OrtSessionOptions* options, _In_ const OrtCUDAProviderOptionsV2* cuda_options) { API_IMPL_BEGIN + if (cuda_options->external_data_loader_use_directstorage != 0 && + cuda_options->external_data_loader_use_directstorage != 1) { + const auto message = onnxruntime::MakeString( + "external_data_loader_use_directstorage got ", + cuda_options->external_data_loader_use_directstorage, + "; must be 0 or 1."); + return OrtApis::CreateStatus(ORT_INVALID_ARGUMENT, message.c_str()); + } if (cuda_options->external_data_loader_reading_threads > OrtCUDAProviderOptionsV2::kMaxExternalDataLoaderReadingThreadCount) { const auto message = onnxruntime::MakeString( diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc index c9593f6333ecc..16e8f540a0f35 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_external_data_loader_test.cc @@ -66,12 +66,14 @@ void CreateExternalDataFile(size_t length, PathString& path, } void VerifyLoad(size_t length, size_t load_count = 1, size_t reading_thread_count = 4, - bool use_gds = false, size_t prefix_size = kFilePrefixSize) { + bool use_gds = false, size_t prefix_size = kFilePrefixSize, + bool use_directstorage = false) { OrtCUDAProviderOptionsV2 provider_options{}; provider_options.do_copy_in_default_stream = true; provider_options.use_tf32 = false; provider_options.external_data_loader_reading_threads = reading_thread_count; provider_options.external_data_loader_use_gds = use_gds; + provider_options.external_data_loader_use_directstorage = use_directstorage; auto execution_provider = CudaExecutionProviderWithOptions(&provider_options); ASSERT_NE(execution_provider, nullptr); auto loader = execution_provider->GetExternalDataLoader(); @@ -263,6 +265,74 @@ TEST(CudaExternalDataLoaderTest, LoadsUnalignedDataWithGdsEnabled) { VerifyLoad(cuda::kExternalDataLoaderParallelReadThreshold + 1, 2, 4, true); } +TEST(CudaExternalDataLoaderTest, LoadsWithDirectStorageOrConfiguredFallback) { + for (const size_t readers : {0, 1, 4}) { + SCOPED_TRACE(readers); + VerifyLoad(cuda::kGdsIoAlignment, 2, readers, false, cuda::kGdsIoAlignment, true); + VerifyLoad(cuda::kExternalDataLoaderParallelReadThreshold + 1, 2, readers, false, kFilePrefixSize, true); + } +} + +TEST(CudaExternalDataLoaderTest, LoadsMultipleBuffersWithDirectStorageEnabled) { + VerifyLoad(2 * cuda::kExternalDataLoaderBufferSize + 1, 2, 4, false, kFilePrefixSize, true); +} + +TEST(CudaExternalDataLoaderTest, RejectsInvalidDirectStorageOptionFromStructOptions) { + OrtCUDAProviderOptionsV2 provider_options{}; + provider_options.external_data_loader_use_directstorage = 2; + Ort::SessionOptions session_options; + Ort::Status status(Ort::GetApi().SessionOptionsAppendExecutionProvider_CUDA_V2(session_options, &provider_options)); + ASSERT_FALSE(status.IsOK()); + EXPECT_THAT(status.GetErrorMessage(), testing::HasSubstr("external_data_loader_use_directstorage")); + EXPECT_THAT(status.GetErrorMessage(), testing::HasSubstr("got 2")); + EXPECT_EQ(CudaExecutionProviderWithOptions(&provider_options), nullptr); +} + +#if defined(ORT_CUDA_DIRECTSTORAGE_AVAILABLE) && !defined(ORT_NO_RTTI) +TEST(CudaExternalDataLoaderTest, NativeDirectStorageWithoutFallback) { + std::unique_ptr loader; + const auto setup_status = cuda::DirectStorageLoader::Create(0, loader); + if (!setup_status.IsOK()) { + GTEST_SKIP() << "DirectStorage/D3D12/CUDA interoperability unavailable: " << setup_status.ErrorMessage(); + } + auto execution_provider = DefaultCudaExecutionProvider(); + ASSERT_NE(execution_provider, nullptr); + auto allocators = execution_provider->CreatePreferredAllocators(); + const auto allocator = std::find_if(allocators.begin(), allocators.end(), [](const AllocatorPtr& candidate) { + return candidate->Info().device.Type() == OrtDevice::GPU && + candidate->Info().mem_type == OrtMemTypeDefault; + }); + ASSERT_NE(allocator, allocators.end()); + constexpr size_t length = cuda::kExternalDataLoaderBufferSize + 17; + Tensor tensor(DataTypeImpl::GetType(), TensorShape({length}), *allocator); + for (size_t load = 0; load < 2; ++load) { + PathString path; + ASSERT_NO_FATAL_FAILURE(CreateExternalDataFile(length, path, {}, load)); + ScopedFileDeleter deleter{path}; + std::unique_ptr file; + ASSERT_STATUS_OK(Env::Default().OpenRandomAccessFile(path.c_str(), file)); + const auto* handles = dynamic_cast(file.get()); + ASSERT_NE(handles, nullptr); + ASSERT_STATUS_OK(loader->Load(path, handles->GetFileHandle(), kFilePrefixSize, length, tensor)); + std::vector output(length); + ASSERT_EQ(cudaSuccess, cudaMemcpy(output.data(), tensor.DataRaw(), length, cudaMemcpyDeviceToHost)); + for (size_t i = 0; i < length; ++i) { + ASSERT_EQ(TestValue(i, load), output[i]) << "byte=" << i << " load=" << load; + } + EXPECT_FALSE(loader->Load(path, handles->GetFileHandle(), -1, length, tensor).IsOK()); + EXPECT_FALSE(loader->Load(path, handles->GetFileHandle(), kFilePrefixSize + 1, length, tensor).IsOK()); + EXPECT_FALSE(loader->Load(path, nullptr, kFilePrefixSize, length, tensor).IsOK()); + + PathString other_path; + ASSERT_NO_FATAL_FAILURE(CreateExternalDataFile(length, other_path, {}, load + 1)); + ScopedFileDeleter other_deleter{other_path}; + const auto replaced = loader->Load(other_path, handles->GetFileHandle(), kFilePrefixSize, length, tensor); + EXPECT_FALSE(replaced.IsOK()); + EXPECT_THAT(replaced.ErrorMessage(), testing::HasSubstr("file changed")); + } +} +#endif + TEST(CudaExternalDataLoaderTest, NormalizesBoolWithPinnedAndPageableLoading) { const std::array input{0, 1, 2, 255}; const std::array expected{0, 1, 1, 1}; diff --git a/onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py b/onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py index 8b16d7e469368..f2490e7f8c540 100644 --- a/onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py +++ b/onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py @@ -3,25 +3,60 @@ # Licensed under the MIT License. # -------------------------------------------------------------------------- -"""Measure CUDA InferenceSession creation for a model with external data. +"""Compare pageable, pinned and direct-storage external weights on one CUDA device. -Example: - python benchmark_cuda_model_loading.py \ - --model /path/to/model.onnx \ - --external-data /path/to/model.onnx.data \ - --evict-file-cache +Generate a deterministic, 4-KiB-aligned fixture and benchmark it (no PyTorch needed): + python benchmark_cuda_model_loading.py --generate-model benchmark_data --repetitions 5 -Run this script in a fresh process for every sample. Cache eviction is applied -only to the paths passed on the command line and remains an operating-system -hint rather than a guarantee. +Reuse a model with independently computed reference outputs (NPZ keys are tensor names): + python benchmark_cuda_model_loading.py --model model.onnx --inputs inputs.npz \ + --expected-outputs expected.npz --output results.json + +Every warmup and measured sample runs in a fresh process. Only InferenceSession +construction is timed; imports, input loading, and correctness checks are excluded. +Dependencies are imported lazily so CLI help and reporting helpers need no ORT build. +Initialization completes weight transfers; a blocking inference and NumPy comparison +then verify the result. Loader INFO logs, not requested provider options, determine +the observed path. Direct storage means GDS on Linux or DirectStorage on Windows. +Uncompressed DirectStorage uses host/upload staging with a shared D3D12/CUDA +destination, not zero-host-copy NVMe-to-VRAM DMA. GDS requests native I/O with cuFile +compatibility mode disabled; the physical path remains platform-dependent. +ORT logs prove API use, not the physical storage-to-GPU DMA route. + +Caches remain OS-managed (warmups and fixture creation can warm them). The optional +legacy --evict-file-cache is only a per-file, unprivileged POSIX hint, never proof +of cold storage. No privileged or global cache flushing is performed. """ import argparse import json +import math import os +import platform +import random +import re +import statistics +import subprocess +import sys import time +from pathlib import Path -import onnxruntime +LOADER_LOG = re.compile(r"CUDA external data loader: path=(pageable|pinned|gds|directstorage) bytes=(\d+)\b") +UTF16LE_ASCII_LOG = re.compile(r"(?:[\t\r\n\x20-\x7e]\x00){2,}") +RESULT_PREFIX = "CUDA_LOADING_RESULT=" +PATHS = ("pageable", "pinned", "direct") +PATH_DESCRIPTIONS = { + "pageable": "CPU pageable buffer followed by H2D copy", + "pinned": "CPU pinned staging buffers followed by H2D copy", + "gds": ( + "NVIDIA cuFile API; native GDS requested with compatibility mode disabled, " + "physical storage-to-GPU path remains platform-dependent" + ), + "directstorage": ( + "Microsoft DirectStorage API; uncompressed host/upload staging with a shared D3D12/CUDA destination, " + "not zero-host-copy NVMe-to-VRAM DMA" + ), +} def reading_thread_count(value): @@ -31,81 +66,438 @@ def reading_thread_count(value): return count +def positive_int(value): + count = int(value) + if count <= 0: + raise argparse.ArgumentTypeError("must be positive") + return count + + +def nonnegative_int(value): + count = int(value) + if count < 0: + raise argparse.ArgumentTypeError("must be nonnegative") + return count + + def evict_file_cache(path): if not hasattr(os, "posix_fadvise"): raise RuntimeError("file-cache eviction requires os.posix_fadvise") - with open(path, "rb") as file: os.posix_fadvise(file.fileno(), 0, 0, os.POSIX_FADV_DONTNEED) -def main(): - parser = argparse.ArgumentParser(description="Benchmark CUDA model loading") - parser.add_argument("--model", required=True, help="Path to the ONNX model") - parser.add_argument( - "--external-data", - action="append", - default=[], - help="External-data file to evict with --evict-file-cache; may be specified more than once", - ) - parser.add_argument("--device-id", type=int, default=0, help="CUDA device ID") - parser.add_argument("--threads", type=int, default=96, help="Intra-op thread count") - parser.add_argument( - "--reading-threads", - type=reading_thread_count, - help="Override parallel reads per CUDA pinned staging buffer; 0 disables the loader (runtime default: 4)", - ) - parser.add_argument( - "--evict-file-cache", - action="store_true", - help="Advise the OS to evict the model and external-data files before creating the session", - ) - args = parser.parse_args() +def direct_path(platform_name): + if platform_name.startswith("linux"): + return "gds" + if platform_name == "win32": + return "directstorage" + return None + - paths = [args.model, *args.external_data] +def provider_options(path, device_id, reading_threads, platform_name): + options = { + "device_id": device_id, + "external_data_loader_reading_threads": 0 if path == "pageable" else reading_threads, + "external_data_loader_use_gds": 0, + "external_data_loader_use_directstorage": 0, + "use_tf32": 0, + } + if path == "direct": + api_path = direct_path(platform_name) + if api_path is None: + raise ValueError(f"Direct storage is unsupported on {platform_name}") + options[f"external_data_loader_use_{api_path}"] = 1 + return options + + +def parse_loader_logs(stderr): + """Parse mixed Python text and Windows native UTF-16LE ASCII log segments. + + Only paired ASCII/NUL runs are normalized, not arbitrary NUL characters. + The caller retains the original stderr as evidence in the sample. + """ + stderr = UTF16LE_ASCII_LOG.sub(lambda match: match[0][::2], stderr) + observed = {} + for path, size in LOADER_LOG.findall(stderr): + observed[path] = observed.get(path, 0) + int(size) + warnings = [ + line + for line in stderr.splitlines() + if "CUDA external data loader" in line + and ("fallback" in line.lower() or "falling back" in line.lower() or "[W:" in line) + ] + return observed, warnings + + +def classify_path(expected_path, observed, expected_bytes): + """Require complete byte accounting before claiming the requested path ran.""" + observed = {path: size for path, size in observed.items() if size > 0} + if not observed: + return "unverified" + if sum(observed.values()) != expected_bytes: + return "incomplete" + if len(observed) > 1: + return "mixed" + if expected_path not in observed: + return "fallback" + return "verified" + + +def percentile(values, fraction): + ordered = sorted(values) + position = (len(ordered) - 1) * fraction + lower = math.floor(position) + upper = math.ceil(position) + return ordered[lower] + (ordered[upper] - ordered[lower]) * (position - lower) + + +def summarize_samples(samples, expected_bytes): + """Keep verified, mixed, and fallback measurements in separate distributions.""" + summaries = [] + groups = {} + for sample in samples: + key = (sample["requested_path"], sample["status"], tuple(sorted(sample["observed_bytes"].items()))) + groups.setdefault(key, []).append(sample) + for (requested, status, observed), group in groups.items(): + seconds = [sample["seconds"] for sample in group] + median = statistics.median(seconds) + summaries.append( + { + "requested_path": requested, + "status": status, + "observed_bytes": dict(observed), + "observed_path_descriptions": {path: PATH_DESCRIPTIONS[path] for path, _ in observed}, + "throughput_definition": ( + "External weight bytes divided by end-to-end InferenceSession initialization time; " + "not storage bandwidth" + ), + "count": len(group), + "seconds": { + "min": min(seconds), + "median": median, + "mean": statistics.mean(seconds), + "stdev": statistics.stdev(seconds) if len(seconds) > 1 else 0.0, + "p90": percentile(seconds, 0.9), + "p95": percentile(seconds, 0.95), + "max": max(seconds), + }, + # End-to-end initialization throughput, not a device or disk bandwidth measurement. + "effective_gib_per_second": { + "median": statistics.median(expected_bytes / (2**30) / value for value in seconds), + "min": expected_bytes / (2**30) / max(seconds), + "max": expected_bytes / (2**30) / min(seconds), + }, + } + ) + return summaries + + +def generate_model(directory, weight_count, dimension, seed): + import numpy as np # noqa: PLC0415 + import onnx # noqa: PLC0415 + from onnx import TensorProto, helper # noqa: PLC0415 + + if dimension % 32: + raise ValueError("--weight-dim must be divisible by 32 for 4-KiB-aligned lengths") + directory = Path(directory).resolve() + directory.mkdir(parents=True, exist_ok=True) + paths = [directory / name for name in ("model.onnx", "weights.bin", "inputs.npz", "expected.npz")] for path in paths: - if not os.path.isfile(path): - raise FileNotFoundError(path) + if path.exists(): + raise FileExistsError(f"Refusing to overwrite {path}; use --model to reuse an existing fixture") + rng = np.random.default_rng(seed) + initializers, nodes, outputs, expected = [], [], [], {} + with paths[1].open("wb") as weights_file: + for index in range(weight_count): + name, output = f"weight_{index}", f"output_{index}" + # Binary fractions and a ones input yield stable sums without TF32 ambiguity. + weight = rng.integers(-8, 9, size=(dimension, dimension), dtype=np.int32).astype(" round_trip_options; + for (std::string entry; std::getline(stream, entry, ';');) { + const auto separator = entry.find('='); + ASSERT_NE(separator, std::string::npos); + round_trip_options.emplace(entry.substr(0, separator), entry.substr(separator + 1)); + } + ASSERT_EQ(round_trip_options.at(key), value); + Ort::CUDAProviderOptions restored; + restored.Update(round_trip_options); + EXPECT_EQ((*restored).external_data_loader_use_directstorage, value[0] - '0'); + EXPECT_EQ(restored.GetCUDAProviderOptionsAsString(), serialized); + Ort::SessionOptions session_options; + EXPECT_NO_THROW(session_options.AppendExecutionProvider_CUDA_V2(*options)); + } + options.Update({{key, "1"}}); + for (const char* value : {"2", "-1", "invalid", ""}) { + const char* keys[] = {key}; + const char* values[] = {value}; + Ort::Status status(Ort::GetApi().UpdateCUDAProviderOptions(options, keys, values, 1)); + ASSERT_FALSE(status.IsOK()); + EXPECT_THAT(status.GetErrorMessage(), + testing::HasSubstr(value[0] == '\0' ? "key/value cannot be empty" : key)); + EXPECT_EQ((*options).external_data_loader_use_directstorage, 1); + } +} + TEST(CApiTest, CUDAProviderOptionsGdsRoundTrip) { constexpr const char* key = "external_data_loader_use_gds"; Ort::CUDAProviderOptions cuda_options; From c40fe682722cd7a21fcbc6f608cea08382c2884d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Xavier=20Dupr=C3=A9?= Date: Wed, 23 Sep 2026 18:35:04 +0200 Subject: [PATCH 14/17] Address DirectStorage review feedback and enable Windows CI coverage Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Copilot-Session: 8af39836-82f0-478b-b1a0-338c518c1828 --- .github/workflows/windows_cuda.yml | 27 ++++++++++++++++--- cmake/CMakeLists.txt | 3 +++ cmake/external/directstorage.cmake | 5 ++++ docs/CUDA_cuDNN_Optional_Design.md | 7 ++--- docs/Model_Loading_Performance.md | 4 ++- .../benchmark_cuda_model_loading.py | 15 +++++------ .../test_benchmark_cuda_model_loading.py | 3 +-- 7 files changed, 47 insertions(+), 17 deletions(-) diff --git a/.github/workflows/windows_cuda.yml b/.github/workflows/windows_cuda.yml index b7209340deaa9..ed1d3baeaafef 100644 --- a/.github/workflows/windows_cuda.yml +++ b/.github/workflows/windows_cuda.yml @@ -121,13 +121,34 @@ jobs: exit $lastExitCode } # Execute the build process - python.exe ${{ github.workspace }}\tools\ci_build\build.py --update --build --config RelWithDebInfo --build_dir build --skip_submodule_sync --build_csharp --parallel --nvcc_threads 4 --flash_nvcc_threads 4 --use_binskim_compliant_compile_flags --cmake_generator "Visual Studio 17 2022" --build_shared_lib --build_wheel --build_java --use_cuda --cuda_home="$env:RUNNER_TEMP\v12.8" --enable_cuda_profiling --use_vcpkg --use_vcpkg_ms_internal_asset_cache --enable_transformers_tool_test --cmake_extra_defines onnxruntime_QUICK_BUILD=ON --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=86 --cmake_extra_defines onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON + python.exe ${{ github.workspace }}\tools\ci_build\build.py --update --build --config RelWithDebInfo --build_dir build --skip_submodule_sync --build_csharp --parallel --nvcc_threads 4 --flash_nvcc_threads 4 --use_binskim_compliant_compile_flags --cmake_generator "Visual Studio 17 2022" --build_shared_lib --build_wheel --build_java --use_cuda --cuda_home="$env:RUNNER_TEMP\v12.8" --enable_cuda_profiling --use_vcpkg --use_vcpkg_ms_internal_asset_cache --enable_transformers_tool_test --cmake_extra_defines onnxruntime_QUICK_BUILD=ON --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=86 --cmake_extra_defines onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON onnxruntime_USE_CUDA_DIRECTSTORAGE=ON if ($lastExitCode -ne 0) { exit $lastExitCode } - # Clean up the output directory before uploading artifacts $outputDir = "${{ runner.temp }}\build\RelWithDebInfo" + # DirectStorage resolves dstoragecore.dll relative to the process executable. + $testDir = Join-Path $outputDir "RelWithDebInfo" + $sdkManifest = Join-Path $outputDir "directstorage-source-dir.txt" + if (!(Test-Path -LiteralPath $sdkManifest -PathType Leaf)) { + throw "Missing DirectStorage SDK source manifest: $sdkManifest" + } + if (!(Test-Path -LiteralPath (Join-Path $testDir "onnxruntime_provider_test.exe") -PathType Leaf)) { + throw "Missing CUDA provider test executable in $testDir" + } + $sdkDir = (Get-Content -LiteralPath $sdkManifest -Raw -ErrorAction Stop).Trim() + if ([string]::IsNullOrWhiteSpace($sdkDir)) { + throw "Empty DirectStorage SDK source manifest: $sdkManifest" + } + foreach ($dll in @("dstorage.dll", "dstoragecore.dll")) { + $source = Join-Path $sdkDir "native\bin\x64\$dll" + if (!(Test-Path -LiteralPath $source -PathType Leaf)) { + throw "Missing DirectStorage runtime: $source" + } + Copy-Item -LiteralPath $source -Destination $testDir -ErrorAction Stop + } + + # Clean up the output directory before uploading artifacts Write-Host "Cleaning up files from $outputDir..." Remove-Item -Path "$outputDir\onnxruntime" -Recurse -Force -ErrorAction SilentlyContinue @@ -253,7 +274,7 @@ jobs: exit $lastExitCode } - python.exe ${{ github.workspace }}\tools\ci_build\build.py --test --config RelWithDebInfo --build_dir build --skip_submodule_sync --build_csharp --parallel --nvcc_threads 4 --flash_nvcc_threads 4 --use_binskim_compliant_compile_flags --cmake_generator "Visual Studio 17 2022" --build_shared_lib --build_wheel --build_java --use_cuda --cuda_home="$env:RUNNER_TEMP\v12.8" --enable_cuda_profiling --use_vcpkg --use_vcpkg_ms_internal_asset_cache --enable_transformers_tool_test --cmake_extra_defines onnxruntime_QUICK_BUILD=ON --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=86 --cmake_extra_defines onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON + python.exe ${{ github.workspace }}\tools\ci_build\build.py --test --config RelWithDebInfo --build_dir build --skip_submodule_sync --build_csharp --parallel --nvcc_threads 4 --flash_nvcc_threads 4 --use_binskim_compliant_compile_flags --cmake_generator "Visual Studio 17 2022" --build_shared_lib --build_wheel --build_java --use_cuda --cuda_home="$env:RUNNER_TEMP\v12.8" --enable_cuda_profiling --use_vcpkg --use_vcpkg_ms_internal_asset_cache --enable_transformers_tool_test --cmake_extra_defines onnxruntime_QUICK_BUILD=ON --cmake_extra_defines CMAKE_CUDA_ARCHITECTURES=86 --cmake_extra_defines onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS=ON onnxruntime_USE_CUDA_DIRECTSTORAGE=ON if ($lastExitCode -ne 0) { exit $lastExitCode } diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index 8832a9145dced..30ecf343e9c33 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -78,6 +78,9 @@ 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) +if(onnxruntime_USE_CUDA_DIRECTSTORAGE AND onnxruntime_BUILD_CUDA_EP_AS_PLUGIN) + message(FATAL_ERROR "onnxruntime_USE_CUDA_DIRECTSTORAGE is not supported with onnxruntime_BUILD_CUDA_EP_AS_PLUGIN.") +endif() option(onnxruntime_BUILD_CUDA_QUANT_PREPROCESS "Build CUDA weight-packing module onnxruntime_cuda_quant_preprocess.so" OFF) option(onnxruntime_CUDA_MINIMAL "Build CUDA without any operations apart from memcpy ops. Useful for a very minimal TRT build" OFF) option(onnxruntime_ENABLE_CUDA_LINE_NUMBER_INFO "When building with CUDA support, generate device code line number information." OFF) diff --git a/cmake/external/directstorage.cmake b/cmake/external/directstorage.cmake index 448a4d1fd1b06..ba4868a4a2d8e 100644 --- a/cmake/external/directstorage.cmake +++ b/cmake/external/directstorage.cmake @@ -9,3 +9,8 @@ onnxruntime_fetchcontent_declare( DOWNLOAD_NAME directstorage.zip ) onnxruntime_fetchcontent_makeavailable(directstorage) + +if(onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS) + # CI deploys the SDK runtime only beside test executables, including when the SDK source is overridden. + file(GENERATE OUTPUT "${CMAKE_BINARY_DIR}/directstorage-source-dir.txt" CONTENT "${directstorage_SOURCE_DIR}\n") +endif() diff --git a/docs/CUDA_cuDNN_Optional_Design.md b/docs/CUDA_cuDNN_Optional_Design.md index 1d761d31b72b0..6f37017b5258c 100644 --- a/docs/CUDA_cuDNN_Optional_Design.md +++ b/docs/CUDA_cuDNN_Optional_Design.md @@ -292,9 +292,10 @@ Implementation details: `CUDAExecutionProviderInfo::FromProviderOptions(...)`. - Emit it from `CUDAExecutionProviderInfo::ToProviderOptions(...)`. - Include it in `std::hash` because it changes the EP behavior. -- Do **not** add a field to `OrtCUDAProviderOptionsV2` for Phase 1. That struct is public C - ABI surface; string-key provider options are sufficient and can be set through existing - provider-options APIs. +- Phase 1 keeps this policy in `CUDAExecutionProviderInfo`. `OrtCUDAProviderOptionsV2` is opaque in the + public C API: callers obtain it through `CreateCUDAProviderOptions` and configure it through string keys. + Its definition in `core/providers/cuda/cuda_provider_options.h` is internal and may be extended for new + options. This differs from the publicly defined `OrtCUDAProviderOptions`, whose layout must remain stable. - Add an EP helper such as `CUDAExecutionProvider::IsCudnnEnabled()` or `CudaKernel::IsCudnnEnabled()` so kernels can distinguish: - cuDNN disabled by user (`enable_cudnn=0`), and diff --git a/docs/Model_Loading_Performance.md b/docs/Model_Loading_Performance.md index 8bb51bf58ad5d..7d0a3ac6d7762 100644 --- a/docs/Model_Loading_Performance.md +++ b/docs/Model_Loading_Performance.md @@ -96,7 +96,9 @@ enabling GDS in those builds logs a warning and uses the configured host-memory The built-in CUDA execution provider also supports Microsoft's DirectStorage API through D3D12/CUDA interoperability. Build with `--cmake_extra_defines onnxruntime_USE_CUDA_DIRECTSTORAGE=ON` in addition to the -usual CUDA build options. This opt-in build downloads the pinned DirectStorage SDK headers; it does not +usual CUDA build options. This option is not supported with `onnxruntime_BUILD_CUDA_EP_AS_PLUGIN=ON`; +CMake rejects that combination rather than silently compiling an unavailable backend. +This opt-in build downloads the pinned DirectStorage SDK headers; it does not introduce a link-time dependency on `dstorage.dll`. Deploy the SDK's matching x64 `dstorage.dll` and `dstoragecore.dll` beside the application executable, following Microsoft's [DirectStorage deployment guidance](https://github.com/microsoft/DirectStorage/blob/main/Docs/DeveloperGuidance.md#sdk-path). diff --git a/onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py b/onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py index f2490e7f8c540..6c916673edc6d 100644 --- a/onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py +++ b/onnxruntime/test/python/transformers/benchmark_cuda_model_loading.py @@ -196,7 +196,6 @@ def summarize_samples(samples, expected_bytes): def generate_model(directory, weight_count, dimension, seed): import numpy as np # noqa: PLC0415 import onnx # noqa: PLC0415 - from onnx import TensorProto, helper # noqa: PLC0415 if dimension % 32: raise ValueError("--weight-dim must be divisible by 32 for 4-KiB-aligned lengths") @@ -215,22 +214,22 @@ def generate_model(directory, weight_count, dimension, seed): weight = rng.integers(-8, 9, size=(dimension, dimension), dtype=np.int32).astype(" Date: Wed, 23 Sep 2026 19:07:06 +0200 Subject: [PATCH 15/17] Update CUDA cuDNN design documentation Clarify the behavior of CUDAExecutionProviderInfo and its options. Introduce a helper to check cuDNN status in kernels. Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com> --- docs/CUDA_cuDNN_Optional_Design.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/CUDA_cuDNN_Optional_Design.md b/docs/CUDA_cuDNN_Optional_Design.md index 6f37017b5258c..e3d80549b29c9 100644 --- a/docs/CUDA_cuDNN_Optional_Design.md +++ b/docs/CUDA_cuDNN_Optional_Design.md @@ -294,7 +294,7 @@ Implementation details: - Include it in `std::hash` because it changes the EP behavior. - Phase 1 keeps this policy in `CUDAExecutionProviderInfo`. `OrtCUDAProviderOptionsV2` is opaque in the public C API: callers obtain it through `CreateCUDAProviderOptions` and configure it through string keys. - Its definition in `core/providers/cuda/cuda_provider_options.h` is internal and may be extended for new + Its definition in `include/onnxruntime/core/providers/cuda/cuda_provider_options.h` is internal and may be extended for new options. This differs from the publicly defined `OrtCUDAProviderOptions`, whose layout must remain stable. - Add an EP helper such as `CUDAExecutionProvider::IsCudnnEnabled()` or `CudaKernel::IsCudnnEnabled()` so kernels can distinguish: From 126728e0a2f3812ba2a274660373cea33ea53b17 Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 08:42:18 +0000 Subject: [PATCH 16/17] Exclude direct storage from minimal builds Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/CMakeLists.txt | 2 +- cmake/external/directstorage.cmake | 2 +- cmake/onnxruntime_providers_cuda.cmake | 12 ++++++++++-- .../core/framework/session_state_utils.cc | 15 +++++++++++++++ onnxruntime/core/platform/env.h | 2 ++ onnxruntime/core/platform/posix/env.cc | 9 ++++++++- onnxruntime/core/platform/windows/env.cc | 9 ++++++++- .../providers/cuda/cuda_external_data_loader.cc | 17 +++++++++++++++-- .../providers/cuda/cuda_external_data_loader.h | 4 ++++ onnxruntime/core/session/provider_bridge_ort.cc | 2 ++ 10 files changed, 66 insertions(+), 8 deletions(-) diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index 30ecf343e9c33..a2a06d919c098 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -1549,7 +1549,7 @@ if (onnxruntime_USE_CUDA) endif() find_package(CUDAToolkit REQUIRED) - if(CMAKE_SYSTEM_NAME STREQUAL "Linux") + if(CMAKE_SYSTEM_NAME STREQUAL "Linux" AND NOT onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_CUDA_MINIMAL) include(CheckCXXSourceCompiles) include(CMakePushCheckState) cmake_push_check_state(RESET) diff --git a/cmake/external/directstorage.cmake b/cmake/external/directstorage.cmake index ba4868a4a2d8e..d5fcbf58b2f91 100644 --- a/cmake/external/directstorage.cmake +++ b/cmake/external/directstorage.cmake @@ -12,5 +12,5 @@ onnxruntime_fetchcontent_makeavailable(directstorage) if(onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS) # CI deploys the SDK runtime only beside test executables, including when the SDK source is overridden. - file(GENERATE OUTPUT "${CMAKE_BINARY_DIR}/directstorage-source-dir.txt" CONTENT "${directstorage_SOURCE_DIR}\n") + file(WRITE "${CMAKE_BINARY_DIR}/directstorage-source-dir.txt" "${directstorage_SOURCE_DIR}\n") endif() diff --git a/cmake/onnxruntime_providers_cuda.cmake b/cmake/onnxruntime_providers_cuda.cmake index 923c12819aad7..0af481707c674 100644 --- a/cmake/onnxruntime_providers_cuda.cmake +++ b/cmake/onnxruntime_providers_cuda.cmake @@ -24,6 +24,12 @@ endif() # Exclude plugin directory if it was picked up by GLOB_RECURSE list(FILTER onnxruntime_providers_cuda_cc_srcs EXCLUDE REGEX "core/providers/cuda/plugin/.*") + if(onnxruntime_MINIMAL_BUILD OR onnxruntime_CUDA_MINIMAL) + list(REMOVE_ITEM onnxruntime_providers_cuda_cc_srcs + "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_directstorage.cc" + "${ONNXRUNTIME_ROOT}/core/providers/cuda/cuda_external_data_loader_gds.cc" + ) + endif() # Remove pch files list(REMOVE_ITEM onnxruntime_providers_cuda_cc_srcs @@ -236,12 +242,14 @@ # config_cuda_provider_shared_module can be used to config onnxruntime_providers_cuda_obj, onnxruntime_providers_cuda & onnxruntime_providers_cuda_ut. # This function guarantees that all 3 targets have the same configurations. - if(onnxruntime_USE_CUDA_DIRECTSTORAGE) + if(onnxruntime_USE_CUDA_DIRECTSTORAGE AND + NOT onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_CUDA_MINIMAL) include(external/directstorage.cmake) endif() function(config_cuda_provider_shared_module target) - if(onnxruntime_USE_CUDA_DIRECTSTORAGE) + if(onnxruntime_USE_CUDA_DIRECTSTORAGE AND + NOT onnxruntime_MINIMAL_BUILD AND NOT onnxruntime_CUDA_MINIMAL) target_compile_definitions(${target} PRIVATE ORT_CUDA_DIRECTSTORAGE_AVAILABLE) target_include_directories(${target} PRIVATE "${directstorage_SOURCE_DIR}/native/include") target_link_libraries(${target} PRIVATE d3d12 dxgi) diff --git a/onnxruntime/core/framework/session_state_utils.cc b/onnxruntime/core/framework/session_state_utils.cc index 2b02480fe1112..bc99373e8bf92 100644 --- a/onnxruntime/core/framework/session_state_utils.cc +++ b/onnxruntime/core/framework/session_state_utils.cc @@ -142,6 +142,20 @@ static common::Status DeserializeTensorProto(const Env& env, const std::basic_st // Bool external initializers are copied verbatim and may carry bytes outside the canonical // {0, 1} set. The CPU staging tensor above can be backed by a read-only mmap, so normalize into // a writable CPU copy before copying to the device (see utils::NormalizeBoolTensorIfNeeded). +#if defined(ORT_MINIMAL_BUILD) + if (cpu_staging_tensor.IsDataType()) { + Tensor normalized_cpu_tensor; + ORT_RETURN_IF_ERROR(AllocateTensorOnDeviceOrMemory(/* use_device_allocator_for_initializers =*/true, + tensor_shape, type, + default_cpu_alloc, normalized_cpu_tensor)); + utils::MakeCpuTensorCopy(cpu_staging_tensor, normalized_cpu_tensor); + utils::NormalizeBoolTensorIfNeeded(normalized_cpu_tensor); + return CopyTensorFromCPUToDevice(data_transfer_mgr, normalized_cpu_tensor, std::move(tensor), ort_value); + } + + return CopyTensorFromCPUToDevice(data_transfer_mgr, deserialized_value.Get(), + std::move(tensor), ort_value); +#else if (cpu_staging_tensor.IsDataType()) { Tensor normalized_cpu_tensor; ORT_RETURN_IF_ERROR(AllocateTensorOnDeviceOrMemory(/* use_device_allocator_for_initializers =*/true, @@ -159,6 +173,7 @@ static common::Status DeserializeTensorProto(const Env& env, const std::basic_st LOGS_DEFAULT(INFO) << "CUDA external data loader: path=pageable bytes=" << cpu_staging_tensor.SizeInBytes(); } return Status::OK(); +#endif } } else { if (device == default_cpu_device) { diff --git a/onnxruntime/core/platform/env.h b/onnxruntime/core/platform/env.h index dccd1ec497c14..aa191f103dce7 100644 --- a/onnxruntime/core/platform/env.h +++ b/onnxruntime/core/platform/env.h @@ -137,6 +137,7 @@ class RandomAccessFile { ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(RandomAccessFile); }; +#if !defined(ORT_MINIMAL_BUILD) class PosixFileDescriptorProvider { public: virtual ~PosixFileDescriptorProvider() = default; @@ -152,6 +153,7 @@ class WindowsFileHandleProvider { // The handle remains owned by the provider and is valid only for its lifetime. virtual void* GetFileHandle() const = 0; }; +#endif /// \brief An interface used by the onnxruntime implementation to /// access operating system functionality like the filesystem etc. diff --git a/onnxruntime/core/platform/posix/env.cc b/onnxruntime/core/platform/posix/env.cc index 8c20cb5ce112c..8755acc613e6c 100644 --- a/onnxruntime/core/platform/posix/env.cc +++ b/onnxruntime/core/platform/posix/env.cc @@ -130,7 +130,12 @@ common::Status GetFileLength(int fd, size_t& file_size) { return common::Status::OK(); } -class PosixRandomAccessFile final : public RandomAccessFile, public PosixFileDescriptorProvider { +class PosixRandomAccessFile final : public RandomAccessFile +#if !defined(ORT_MINIMAL_BUILD) + , + public PosixFileDescriptorProvider +#endif +{ public: PosixRandomAccessFile(ScopedFileDescriptor descriptor, std::string path) : descriptor_(std::move(descriptor)), path_(std::move(path)) {} @@ -163,9 +168,11 @@ class PosixRandomAccessFile final : public RandomAccessFile, public PosixFileDes return common::Status::OK(); } +#if !defined(ORT_MINIMAL_BUILD) int GetFileDescriptor() const override { return descriptor_.Get(); } +#endif private: ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(PosixRandomAccessFile); diff --git a/onnxruntime/core/platform/windows/env.cc b/onnxruntime/core/platform/windows/env.cc index 1131c7cdf7874..250f5a116e8d9 100644 --- a/onnxruntime/core/platform/windows/env.cc +++ b/onnxruntime/core/platform/windows/env.cc @@ -360,12 +360,19 @@ common::Status WindowsEnv::GetFileLength(int fd, /*out*/ size_t& file_size) cons namespace { -class WindowsRandomAccessFile final : public RandomAccessFile, public WindowsFileHandleProvider { +class WindowsRandomAccessFile final : public RandomAccessFile +#if !defined(ORT_MINIMAL_BUILD) + , + public WindowsFileHandleProvider +#endif +{ public: explicit WindowsRandomAccessFile(wil::unique_hfile file_handle) : file_handle_(std::move(file_handle)) {} ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(WindowsRandomAccessFile); +#if !defined(ORT_MINIMAL_BUILD) void* GetFileHandle() const override { return file_handle_.get(); } +#endif Status GetLength(size_t& length) const override { LARGE_INTEGER file_size{}; diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc index 94b2fc049b3c4..34468355a2307 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.cc @@ -119,9 +119,18 @@ common::Status LoadWithPageableBuffer(const RandomAccessFile& file, FileOffsetTy ExternalDataLoader::ExternalDataLoader(int device_id, size_t reading_thread_count, bool use_gds, bool use_directstorage) : device_id_(device_id), - reading_thread_count_(reading_thread_count), + reading_thread_count_(reading_thread_count) +#if !defined(ORT_MINIMAL_BUILD) && !defined(USE_CUDA_MINIMAL) + , use_gds_(use_gds), - use_directstorage_(use_directstorage) {} + use_directstorage_(use_directstorage) +#endif +{ +#if defined(ORT_MINIMAL_BUILD) || defined(USE_CUDA_MINIMAL) + ORT_UNUSED_PARAMETER(use_gds); + ORT_UNUSED_PARAMETER(use_directstorage); +#endif +} ExternalDataLoader::~ExternalDataLoader() { reader_pool_.reset(); @@ -165,8 +174,10 @@ void ExternalDataLoader::ReleaseResources() const noexcept { previous_device != device_id_ && cudaSetDevice(device_id_) == cudaSuccess; +#if !defined(ORT_MINIMAL_BUILD) && !defined(USE_CUDA_MINIMAL) gds_loader_.reset(); directstorage_loader_.reset(); +#endif for (auto& stream : streams_) { if (stream != nullptr) { @@ -213,6 +224,7 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, CudaDeviceGuard device_guard; ORT_RETURN_IF_ERROR(device_guard.SetDevice(device_id_)); +#if !defined(ORT_MINIMAL_BUILD) && !defined(USE_CUDA_MINIMAL) if (use_directstorage_ && !directstorage_disabled_ && length != 0 && std::endian::native == std::endian::little && !tensor.IsDataType()) { Status status = Status::OK(); @@ -272,6 +284,7 @@ common::Status ExternalDataLoader::LoadTensor(const Env& env, << " loader. " << gds_status.ErrorMessage(); } +#endif if (reading_thread_count_ == 0) { ORT_RETURN_IF_ERROR(LoadWithPageableBuffer(*file, data_offset, length, tensor, 1, reader_pool_)); diff --git a/onnxruntime/core/providers/cuda/cuda_external_data_loader.h b/onnxruntime/core/providers/cuda/cuda_external_data_loader.h index 06978072ae21a..5cffc09642578 100644 --- a/onnxruntime/core/providers/cuda/cuda_external_data_loader.h +++ b/onnxruntime/core/providers/cuda/cuda_external_data_loader.h @@ -8,8 +8,10 @@ #include #include "core/framework/external_data_loader.h" +#if !defined(ORT_MINIMAL_BUILD) && !defined(USE_CUDA_MINIMAL) #include "core/providers/cuda/cuda_external_data_loader_gds.h" #include "core/providers/cuda/cuda_external_data_loader_directstorage.h" +#endif #include "cuda_pch.h" namespace onnxruntime { @@ -87,12 +89,14 @@ class ExternalDataLoader final : public IExternalDataLoader { mutable std::array buffers_{}; mutable std::array streams_{}; const size_t reading_thread_count_; +#if !defined(ORT_MINIMAL_BUILD) && !defined(USE_CUDA_MINIMAL) const bool use_gds_; mutable bool gds_disabled_{false}; mutable std::unique_ptr gds_loader_; const bool use_directstorage_; mutable bool directstorage_disabled_{false}; mutable std::unique_ptr directstorage_loader_; +#endif mutable std::unique_ptr reader_pool_; }; diff --git a/onnxruntime/core/session/provider_bridge_ort.cc b/onnxruntime/core/session/provider_bridge_ort.cc index 16720feb4400b..44555b8f83ec7 100644 --- a/onnxruntime/core/session/provider_bridge_ort.cc +++ b/onnxruntime/core/session/provider_bridge_ort.cc @@ -2949,6 +2949,7 @@ ORT_API(void, OrtApis::ReleaseTensorRTProviderOptions, _Frees_ptr_opt_ OrtTensor ORT_API_STATUS_IMPL(OrtApis::SessionOptionsAppendExecutionProvider_CUDA_V2, _In_ OrtSessionOptions* options, _In_ const OrtCUDAProviderOptionsV2* cuda_options) { API_IMPL_BEGIN +#if !defined(ORT_MINIMAL_BUILD) if (cuda_options->external_data_loader_use_directstorage != 0 && cuda_options->external_data_loader_use_directstorage != 1) { const auto message = onnxruntime::MakeString( @@ -2976,6 +2977,7 @@ ORT_API_STATUS_IMPL(OrtApis::SessionOptionsAppendExecutionProvider_CUDA_V2, _In_ ORT_INVALID_ARGUMENT, message.c_str()); } +#endif auto factory = onnxruntime::CudaProviderFactoryCreator::Create(cuda_options); if (!factory) { From 7116be9101e1887f861e94ff96c83113d00dc649 Mon Sep 17 00:00:00 2001 From: xadupre Date: Fri, 25 Sep 2026 10:42:47 +0000 Subject: [PATCH 17/17] Evaluate CUDA internal-test option after its prerequisites Declare BUILD_UNIT_TESTS before the dependent CUDA internal-test option so a fresh configure honors ENABLE_CUDA_EP_INTERNAL_TESTS=ON and writes the DirectStorage SDK manifest required by Windows CI. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- cmake/CMakeLists.txt | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/cmake/CMakeLists.txt b/cmake/CMakeLists.txt index a2a06d919c098..57fe08645dab1 100644 --- a/cmake/CMakeLists.txt +++ b/cmake/CMakeLists.txt @@ -71,11 +71,6 @@ option(onnxruntime_ENABLE_MEMLEAK_CHECKER "Experimental: Enable memory leak chec option(onnxruntime_ENABLE_CONVSYMKERNELAVX2_SAT_CHECKER "Experimental: Enable ConvSymKernelAvx2 assembly saturation checker in build" OFF) option(onnxruntime_USE_CUDA "Build with CUDA support" OFF) cmake_dependent_option(onnxruntime_USE_CUDA_DIRECTSTORAGE "Build Microsoft DirectStorage CUDA loading support" OFF "onnxruntime_USE_CUDA;WIN32" OFF) -# Enable ONNX Runtime CUDA EP's internal unit tests that directly access the EP's internal functions instead of through -# OpKernels. When the option is ON, we will have two copies of GTest library in the same process. It is not a typical -# use. If you hit any problem with that, please do not report it to GTest. Turn OFF the following build option instead. -cmake_dependent_option(onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS "Build with CUDA unit tests" OFF "onnxruntime_USE_CUDA;onnxruntime_BUILD_UNIT_TESTS" OFF) - 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) if(onnxruntime_USE_CUDA_DIRECTSTORAGE AND onnxruntime_BUILD_CUDA_EP_AS_PLUGIN) @@ -101,6 +96,12 @@ option(onnxruntime_USE_ARM_NEON_NCHWC "Build with ARM Neon NCHWc kernels in MLAS option(onnxruntime_USE_KLEIDIAI "Build with KleidiAI integration in MLAS" OFF) option(onnxruntime_USE_QMX_KLEIDIAI_COEXIST "Build with QMX and Arm KLEIDIAI libraries" OFF) option(onnxruntime_BUILD_UNIT_TESTS "Build ONNXRuntime unit tests" ON) +# Declare both prerequisites before evaluating this dependent option on a fresh configure. +# Enable ONNX Runtime CUDA EP's internal unit tests that directly access the EP's internal functions instead of through +# OpKernels. When the option is ON, we will have two copies of GTest library in the same process. It is not a typical +# use. If you hit any problem with that, please do not report it to GTest. Turn OFF the following build option instead. +cmake_dependent_option(onnxruntime_ENABLE_CUDA_EP_INTERNAL_TESTS "Build with CUDA unit tests" OFF "onnxruntime_USE_CUDA;onnxruntime_BUILD_UNIT_TESTS" OFF) + # Materialize the ONNX node-test corpus from ONNX's Python generators into the build tree at # configure/build time, instead of depending on the on-disk corpus shipped in the ONNX source # archive. This detaches ORT from ONNX PR #7959 (which deletes onnx/backend/test/data/node).