From 20ac107b360473cd2eec4824e0399d6deaa65d81 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 16 Sep 2026 01:46:47 +0000 Subject: [PATCH 01/10] Add CUDA host-pageable gather implementation Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- docs/cuda_host_pageable_gather.md | 18 +++ .../providers/cuda/cuda_provider_options.h | 1 + .../quantization/gather_block_quantized.cc | 112 +++++++++++++++- .../quantization/gather_block_quantized.cuh | 9 ++ .../quantization/gather_block_quantized.h | 25 +++- .../providers/cuda/cuda_execution_provider.h | 1 + .../cuda/cuda_execution_provider_info.cc | 7 + .../cuda/cuda_execution_provider_info.h | 2 + onnxruntime/core/providers/cuda/cuda_kernel.h | 2 + .../providers/cuda/cuda_provider_factory.cc | 2 + .../core/providers/cuda/plugin/cuda_ep.cc | 2 + .../core/providers/cuda/plugin/cuda_ep.h | 1 + .../providers/cuda/plugin/cuda_ep_factory.cc | 4 + .../cuda/plugin/cuda_kernel_adapter.h | 5 + .../gather_block_quantized_op_test.cc | 120 ++++++++++++++++++ 15 files changed, 305 insertions(+), 6 deletions(-) create mode 100644 docs/cuda_host_pageable_gather.md diff --git a/docs/cuda_host_pageable_gather.md b/docs/cuda_host_pageable_gather.md new file mode 100644 index 0000000000000..f30638a079f1b --- /dev/null +++ b/docs/cuda_host_pageable_gather.md @@ -0,0 +1,18 @@ +# CUDA host-pageable `GatherBlockQuantized` + +The CUDA Execution Provider option `enable_host_pageable_gather` enables direct access to CPU-resident FP8 +`com.microsoft::GatherBlockQuantized` input data. It is disabled by default. + +Direct access requires a CUDA device that reports both `cudaDevAttrPageableMemoryAccess` and +`cudaDevAttrPageableMemoryAccessUsesHostPageTables`. CUDA Graph capture is not currently supported with this mode. +If either device capability is unavailable or CUDA Graphs are enabled, ONNX Runtime emits a warning and makes one +persistent CUDA copy of a constant input instead. + +The option does not create or manage a file mapping. The model initializer must already be supplied as CPU memory, +such as a file-backed mapping, and that mapping remains live for the session. The direct path does not register, +prefetch, hash, scan, or copy the complete initializer, and the initializer is not accounted as CUDA-resident memory. +Fallback copies are allocated by the CUDA initializer allocator; they are not currently included in capacity-aware +partitioning estimates. + +This mode primarily reduces GPU memory capacity requirements. Performance depends on storage latency and operating +system page-cache state, so cold prefill can be slower and less predictable than using resident GPU memory. diff --git a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h index 7a44cb8e76daf..3df52b95a95bf 100644 --- a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h +++ b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h @@ -40,4 +40,5 @@ 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 + int enable_host_pageable_gather = 0; // Enable direct host-pageable FP8 GatherBlockQuantized access. }; diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc index a67514ac560b2..5a733b2abfaf0 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc @@ -3,6 +3,7 @@ #include +#include "core/common/logging/logging.h" #include "core/providers/cuda/cuda_common.h" #include "contrib_ops/cuda/quantization/gather_block_quantized.h" #include "contrib_ops/cuda/quantization/gather_block_quantized.cuh" @@ -21,7 +22,8 @@ using namespace onnxruntime::cuda; (*KernelDefBuilder::Create()) \ .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ .TypeConstraint("T2", DataTypeImpl::GetTensorType()) \ - .TypeConstraint("Tind", DataTypeImpl::GetTensorType()), \ + .TypeConstraint("Tind", DataTypeImpl::GetTensorType()) \ + .InputMemoryType(OrtMemTypeCPUInput, 0), \ GatherBlockQuantized); REGISTER_GATHERBLOCKQUANTIZED(uint8_t, float, int32_t); @@ -71,7 +73,8 @@ REGISTER_GATHERBLOCKQUANTIZED(Float4E2M1x2, BFloat16, int64_t); #endif // !defined(DISABLE_FLOAT4_TYPES) template -GatherBlockQuantized::GatherBlockQuantized(const OpKernelInfo& info) : CudaKernel(info) { +GatherBlockQuantized::GatherBlockQuantized(const OpKernelInfo& info) + : CudaKernel(info), direct_host_data_(false), data_is_constant_(false) { if constexpr (IsFpQuantizedV) { bits_ = 0; // Not applicable for FP8/FP4 data. } else { @@ -82,6 +85,36 @@ GatherBlockQuantized::GatherBlockQuantized(const OpKernelInfo& inf gather_axis_ = info.GetAttrOrDefault("gather_axis", 0); quantize_axis_ = info.GetAttrOrDefault("quantize_axis", 1); + const Tensor* constant_data = nullptr; + data_is_constant_ = info.TryGetConstantInput(0, &constant_data); + + int pageable_memory_access = 0; + int uses_host_page_tables = 0; +#if defined(CUDA_VERSION) && CUDA_VERSION >= 10020 + const bool attributes_available = + cudaDeviceGetAttribute(&pageable_memory_access, cudaDevAttrPageableMemoryAccess, GetDeviceId()) == cudaSuccess && + cudaDeviceGetAttribute(&uses_host_page_tables, cudaDevAttrPageableMemoryAccessUsesHostPageTables, + GetDeviceId()) == cudaSuccess; + if (!attributes_available) { + pageable_memory_access = 0; + uses_host_page_tables = 0; + cudaGetLastError(); + } +#endif + + const bool option_enabled = EnableHostPageableGather(); + direct_host_data_ = + SelectGatherBlockQuantizedDataPolicy(option_enabled, pageable_memory_access != 0, + uses_host_page_tables != 0, IsCudaGraphEnabled(), + IsFp8QuantizedV) == + GatherBlockQuantizedDataPolicy::DirectHost; + if (option_enabled && IsFp8QuantizedV && !direct_host_data_) { + LOGS_DEFAULT(WARNING) + << "enable_host_pageable_gather was requested, but direct host-pageable GatherBlockQuantized " + "access is unavailable because the CUDA device lacks pageable memory access through host page tables " + "or CUDA Graph capture is enabled. Falling back to a persistent CUDA copy."; + } + // If block size is set, it has to be no smaller than 16 and must be power of 2. // block_size_ & (block_size_ - 1) == 0 checks if block_size_ only has 1 bit set. // block_size_ == 0 is only valid for FP8/FP4 data, meaning the whole quantize_axis dimension @@ -93,6 +126,42 @@ GatherBlockQuantized::GatherBlockQuantized(const OpKernelInfo& inf } } +template +Status GatherBlockQuantized::CreateDeviceCopy(const Tensor& tensor, AllocatorPtr alloc) const { + ORT_RETURN_IF_NOT(tensor.Location().device.Type() == OrtDevice::CPU, + "GatherBlockQuantized input 0 must reside in CPU memory."); + + data_shape_.assign(tensor.Shape().GetDims().begin(), tensor.Shape().GetDims().end()); + const size_t bytes = tensor.SizeInBytes(); + if (bytes == 0) { + device_data_.reset(); + return Status::OK(); + } + + device_data_ = IAllocator::MakeUniquePtr(alloc, bytes); + ORT_RETURN_IF_NOT(device_data_ != nullptr, "Failed to allocate persistent CUDA storage for GatherBlockQuantized."); + CUDA_RETURN_IF_ERROR(cudaMemcpy(device_data_.get(), tensor.DataRaw(), bytes, cudaMemcpyHostToDevice)); + return Status::OK(); +} + +template +Status GatherBlockQuantized::PrePack( + const Tensor& tensor, int input_idx, AllocatorPtr alloc, + bool& is_packed, PrePackedWeights* prepacked_weights) { + is_packed = false; + if (input_idx != 0 || direct_host_data_) { + return Status::OK(); + } + + std::lock_guard lock(device_data_mutex_); + ORT_RETURN_IF_ERROR(CreateDeviceCopy(tensor, std::move(alloc))); + is_packed = true; + if (prepacked_weights != nullptr) { + prepacked_weights->has_kernel_owned_packed_weights_ = true; + } + return Status::OK(); +} + template Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) const { const Tensor* data = ctx->Input(0); @@ -100,8 +169,15 @@ Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) const Tensor* scales = ctx->Input(2); const Tensor* zero_points = ctx->Input(3); - auto data_shape = data->Shape().GetDims(); - int64_t data_rank = data->Shape().NumDimensions(); + TensorShapeVector data_shape_storage; + if (data != nullptr) { + data_shape_storage.assign(data->Shape().GetDims().begin(), data->Shape().GetDims().end()); + } else { + std::lock_guard lock(device_data_mutex_); + data_shape_storage = data_shape_; + } + const gsl::span data_shape{data_shape_storage}; + int64_t data_rank = static_cast(data_shape.size()); const int64_t gather_axis = HandleNegativeAxis(gather_axis_, data_rank); const int64_t quantize_axis = HandleNegativeAxis(quantize_axis_, data_rank); @@ -155,7 +231,33 @@ Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) return Status::OK(); } - const auto* data_ptr = data->Data(); + IAllocatorUniquePtr runtime_data; + const T1* data_ptr = nullptr; + if (direct_host_data_) { + ORT_RETURN_IF_NOT(data != nullptr && data->Location().device.Type() == OrtDevice::CPU, + "Direct host-pageable GatherBlockQuantized requires CPU-resident input 0."); + data_ptr = data->Data(); + } else if (data_is_constant_) { + { + std::lock_guard lock(device_data_mutex_); + if (device_data_ == nullptr && data != nullptr && data->SizeInBytes() != 0) { + ORT_RETURN_IF_ERROR(CreateDeviceCopy(*data, Info().GetAllocator(OrtMemTypeDefault))); + } + data_ptr = static_cast(device_data_.get()); + } + } else { + ORT_RETURN_IF_NOT(data != nullptr && data->Location().device.Type() == OrtDevice::CPU, + "GatherBlockQuantized input 0 must reside in CPU memory."); + if (data->SizeInBytes() != 0) { + runtime_data = GetScratchBuffer(data->SizeInBytes(), GetComputeStream(ctx)); + ORT_RETURN_IF_NOT(runtime_data != nullptr, "Failed to allocate CUDA storage for GatherBlockQuantized input 0."); + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(runtime_data.get(), data->DataRaw(), data->SizeInBytes(), + cudaMemcpyHostToDevice, Stream(ctx))); + data_ptr = static_cast(runtime_data.get()); + } + } + ORT_RETURN_IF_NOT(N == 0 || data_ptr != nullptr, + "GatherBlockQuantized fallback has no device-resident input 0."); const auto* indices_ptr = indices->Data(); const T1* zero_points_ptr = nullptr; if (zero_points != nullptr) { diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cuh b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cuh index b240455622474..97ae4688eca11 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cuh +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cuh @@ -43,6 +43,15 @@ struct IsFpQuantized : std::true_type {}; template inline constexpr bool IsFpQuantizedV = IsFpQuantized::value; +template +inline constexpr bool IsFp8QuantizedV = +#if !defined(DISABLE_FLOAT8_TYPES) + std::is_same_v || std::is_same_v || + std::is_same_v || std::is_same_v; +#else + false; +#endif + struct GatherBlockQuantizedParam { cudaStream_t stream; int64_t after_gather_dim; diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h index 7718b6dd06765..85588f6164b53 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h @@ -5,7 +5,7 @@ #include "core/providers/cuda/cuda_kernel.h" -#include +#include using namespace onnxruntime::cuda; @@ -15,17 +15,40 @@ namespace cuda { using namespace onnxruntime::cuda; +enum class GatherBlockQuantizedDataPolicy { + DeviceCopy, + DirectHost, +}; + +constexpr GatherBlockQuantizedDataPolicy SelectGatherBlockQuantizedDataPolicy( + bool option_enabled, bool pageable_memory_access, bool uses_host_page_tables, + bool cuda_graph_enabled, bool is_fp8) { + return option_enabled && pageable_memory_access && uses_host_page_tables && + !cuda_graph_enabled && is_fp8 + ? GatherBlockQuantizedDataPolicy::DirectHost + : GatherBlockQuantizedDataPolicy::DeviceCopy; +} + template class GatherBlockQuantized final : public CudaKernel { public: GatherBlockQuantized(const OpKernelInfo& info); Status ComputeInternal(OpKernelContext* context) const override; + Status PrePack(const Tensor& tensor, int input_idx, AllocatorPtr alloc, + bool& is_packed, PrePackedWeights* prepacked_weights) override; private: + Status CreateDeviceCopy(const Tensor& tensor, AllocatorPtr alloc) const; + int64_t bits_; int64_t block_size_; int64_t gather_axis_; int64_t quantize_axis_; + bool direct_host_data_; + bool data_is_constant_; + mutable std::mutex device_data_mutex_; + mutable IAllocatorUniquePtr device_data_; + mutable TensorShapeVector data_shape_; }; } // namespace cuda diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.h b/onnxruntime/core/providers/cuda/cuda_execution_provider.h index ef546273f3165..4479cd25cc883 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.h +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.h @@ -86,6 +86,7 @@ class CUDAExecutionProvider : public IExecutionProvider { bool IsNHWCPreferred() const { return info_.prefer_nhwc; } bool IsFuseConvBias() const { return info_.fuse_conv_bias; } bool UseTF32() const { return info_.use_tf32; } + bool EnableHostPageableGather() const { return info_.enable_host_pageable_gather; } // Attention kernel options parsed from sdpa_kernel cuda provider option. const AttentionKernelOptions* GetAttentionKernelOptions() const { diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc index 7949537eb1181..f1fb4396516d8 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc @@ -38,6 +38,7 @@ constexpr const char* kUseEPLevelUnifiedStream = "use_ep_level_unified_stream"; constexpr const char* kUseTF32 = "use_tf32"; constexpr const char* kFuseConvBias = "fuse_conv_bias"; constexpr const char* kSdpaKernel = "sdpa_kernel"; +constexpr const char* kEnableHostPageableGather = "enable_host_pageable_gather"; } // namespace provider_option_names } // namespace cuda @@ -131,6 +132,8 @@ CUDAExecutionProviderInfo CUDAExecutionProviderInfo::FromProviderOptions(const P .AddAssignmentToReference(cuda::provider_option_names::kUseEPLevelUnifiedStream, info.use_ep_level_unified_stream) .AddAssignmentToReference(cuda::provider_option_names::kUseTF32, info.use_tf32) .AddAssignmentToReference(cuda::provider_option_names::kSdpaKernel, info.sdpa_kernel) + .AddAssignmentToReference(cuda::provider_option_names::kEnableHostPageableGather, + info.enable_host_pageable_gather) .AddAssignmentToReference(cuda::provider_option_names::kFuseConvBias, info.fuse_conv_bias) .AddValueParser( cuda::provider_option_names::kTunableOpEnable, @@ -187,6 +190,8 @@ ProviderOptions CUDAExecutionProviderInfo::ToProviderOptions(const CUDAExecution {cuda::provider_option_names::kUseTF32, MakeStringWithClassicLocale(info.use_tf32)}, {cuda::provider_option_names::kSdpaKernel, MakeStringWithClassicLocale(info.sdpa_kernel)}, {cuda::provider_option_names::kFuseConvBias, MakeStringWithClassicLocale(info.fuse_conv_bias)}, + {cuda::provider_option_names::kEnableHostPageableGather, + MakeStringWithClassicLocale(info.enable_host_pageable_gather)}, }; return options; @@ -212,6 +217,8 @@ ProviderOptions CUDAExecutionProviderInfo::ToProviderOptions(const OrtCUDAProvid {cuda::provider_option_names::kUseTF32, MakeStringWithClassicLocale(info.use_tf32)}, {cuda::provider_option_names::kFuseConvBias, MakeStringWithClassicLocale(info.fuse_conv_bias)}, {cuda::provider_option_names::kSdpaKernel, MakeStringWithClassicLocale(info.sdpa_kernel)}, + {cuda::provider_option_names::kEnableHostPageableGather, + MakeStringWithClassicLocale(info.enable_host_pageable_gather)}, }; 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 8173828695aff..8d0286e41f92b 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h @@ -83,6 +83,7 @@ struct CUDAExecutionProviderInfo { bool fuse_conv_bias{false}; int sdpa_kernel{0}; + bool enable_host_pageable_gather{false}; static CUDAExecutionProviderInfo FromProviderOptions(const ProviderOptions& options); static ProviderOptions ToProviderOptions(const CUDAExecutionProviderInfo& info); @@ -117,6 +118,7 @@ struct std::hash<::onnxruntime::CUDAExecutionProviderInfo> { onnxruntime::HashCombine(info.tunable_op.max_tuning_duration_ms, value); onnxruntime::HashCombine(info.sdpa_kernel, value); onnxruntime::HashCombine(info.enable_cudnn, value); + onnxruntime::HashCombine(info.enable_host_pageable_gather, value); // Memory pointers onnxruntime::HashCombine(reinterpret_cast(info.user_compute_stream), value); diff --git a/onnxruntime/core/providers/cuda/cuda_kernel.h b/onnxruntime/core/providers/cuda/cuda_kernel.h index b9d48e4925334..b873cca995662 100644 --- a/onnxruntime/core/providers/cuda/cuda_kernel.h +++ b/onnxruntime/core/providers/cuda/cuda_kernel.h @@ -125,6 +125,8 @@ class CudaKernel : public OpKernel { bool GetCudnnConvUseMaxWorkspace() const { return provider_->GetCudnnConvUseMaxWorkspace(); } bool GetCudnnConv1dPadToNc1d() const { return provider_->GetCudnnConv1dPadToNc1d(); } bool IsFuseConvBias() const { return provider_->IsFuseConvBias(); } + bool EnableHostPageableGather() const { return provider_->EnableHostPageableGather(); } + bool IsCudaGraphEnabled() const { return provider_->IsGraphCaptureEnabled(); } // Compatibility helper used by kernels that need the underlying ORT stream object. inline onnxruntime::Stream* GetComputeStream(OpKernelContext* ctx) const { diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc index 66d3617c8da9a..9b789b8472a7e 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc @@ -245,6 +245,7 @@ struct CUDA_Provider : Provider { info.use_ep_level_unified_stream = params->use_ep_level_unified_stream != 0; info.use_tf32 = params->use_tf32 != 0; info.sdpa_kernel = params->sdpa_kernel; + info.enable_host_pageable_gather = params->enable_host_pageable_gather != 0; return std::make_shared(info); } @@ -280,6 +281,7 @@ struct CUDA_Provider : Provider { cuda_options.use_tf32 = internal_options.use_tf32; cuda_options.sdpa_kernel = internal_options.sdpa_kernel; cuda_options.fuse_conv_bias = internal_options.fuse_conv_bias; + cuda_options.enable_host_pageable_gather = internal_options.enable_host_pageable_gather; } ProviderOptions GetProviderOptions(const void* provider_options) override { diff --git a/onnxruntime/core/providers/cuda/plugin/cuda_ep.cc b/onnxruntime/core/providers/cuda/plugin/cuda_ep.cc index 4c946c872ba11..28a8dc1b80256 100644 --- a/onnxruntime/core/providers/cuda/plugin/cuda_ep.cc +++ b/onnxruntime/core/providers/cuda/plugin/cuda_ep.cc @@ -212,6 +212,8 @@ CudaEp::CudaEp(CudaEpFactory& factory, const Config& config, const OrtLogger& lo adapter_config.cudnn_conv1d_pad_to_nc1d = config_.cudnn_conv1d_pad_to_nc1d; adapter_config.enable_cudnn = config_.enable_cudnn; adapter_config.fuse_conv_bias = config_.fuse_conv_bias; + adapter_config.enable_cuda_graph = config_.enable_cuda_graph; + adapter_config.enable_host_pageable_gather = config_.enable_host_pageable_gather; adapter_config.sdpa_kernel = config_.sdpa_kernel; adapter_config.device_id = config_.device_id; adapter_config.do_copy_in_default_stream = config_.do_copy_in_default_stream; diff --git a/onnxruntime/core/providers/cuda/plugin/cuda_ep.h b/onnxruntime/core/providers/cuda/plugin/cuda_ep.h index fdfe3ff02f535..dafe90cc733ab 100644 --- a/onnxruntime/core/providers/cuda/plugin/cuda_ep.h +++ b/onnxruntime/core/providers/cuda/plugin/cuda_ep.h @@ -34,6 +34,7 @@ class CudaEp : public onnxruntime::ep::adapter::Ep { bool fuse_conv_bias = false; ///< Enable cuDNN frontend conv+bias fusion. int sdpa_kernel = 0; ///< Attention backend bitmask override. bool enable_cuda_graph = false; ///< Enable CUDA graph capture and replay. + bool enable_host_pageable_gather = false; ///< Enable direct host-pageable FP8 gather access. int min_num_runs_before_cuda_graph_capture = 2; ///< Warm-up runs before graph capture begins. bool has_user_compute_stream = false; ///< Whether user provided an external CUDA stream. void* user_compute_stream = nullptr; ///< User-provided CUDA stream (cudaStream_t cast to void*). diff --git a/onnxruntime/core/providers/cuda/plugin/cuda_ep_factory.cc b/onnxruntime/core/providers/cuda/plugin/cuda_ep_factory.cc index e675836508be2..d5fd7f80224de 100644 --- a/onnxruntime/core/providers/cuda/plugin/cuda_ep_factory.cc +++ b/onnxruntime/core/providers/cuda/plugin/cuda_ep_factory.cc @@ -652,6 +652,7 @@ OrtStatus* ORT_API_CALL CudaEpFactory::CreateEpImpl( const std::string fuse_conv_bias_key = ep_options_prefix + "fuse_conv_bias"; const std::string sdpa_kernel_key = ep_options_prefix + "sdpa_kernel"; const std::string enable_cuda_graph_key = ep_options_prefix + "enable_cuda_graph"; + const std::string enable_host_pageable_gather_key = ep_options_prefix + "enable_host_pageable_gather"; const std::string min_runs_key = ep_options_prefix + "min_num_runs_before_cuda_graph_capture"; const std::string has_user_compute_stream_key = ep_options_prefix + "has_user_compute_stream"; const std::string user_compute_stream_key = ep_options_prefix + "user_compute_stream"; @@ -688,6 +689,9 @@ OrtStatus* ORT_API_CALL CudaEpFactory::CreateEpImpl( read_session_config_bool( {enable_cuda_graph_key, "ep.cuda.enable_cuda_graph", "enable_cuda_graph"}, config.enable_cuda_graph); + read_session_config_bool( + {enable_host_pageable_gather_key, "ep.cuda.enable_host_pageable_gather", "enable_host_pageable_gather"}, + config.enable_host_pageable_gather); read_session_config_non_negative_int( {min_runs_key, "ep.cuda.min_num_runs_before_cuda_graph_capture"}, config.min_num_runs_before_cuda_graph_capture); diff --git a/onnxruntime/core/providers/cuda/plugin/cuda_kernel_adapter.h b/onnxruntime/core/providers/cuda/plugin/cuda_kernel_adapter.h index 4ccb919eb8bc9..09c523a8d3b28 100644 --- a/onnxruntime/core/providers/cuda/plugin/cuda_kernel_adapter.h +++ b/onnxruntime/core/providers/cuda/plugin/cuda_kernel_adapter.h @@ -534,6 +534,8 @@ struct CudaKernelAdapterRuntimeConfig { bool cudnn_conv1d_pad_to_nc1d = false; bool enable_cudnn = true; bool fuse_conv_bias = false; + bool enable_cuda_graph = false; + bool enable_host_pageable_gather = false; int sdpa_kernel = 0; int device_id = 0; bool do_copy_in_default_stream = true; @@ -1173,7 +1175,10 @@ class CudaKernel : public OpKernel { bool GetCudnnConv1dPadToNc1d() const { return runtime_config_->cudnn_conv1d_pad_to_nc1d; } bool UseTF32() const { return use_tf32_; } bool IsFuseConvBias() const { return runtime_config_->fuse_conv_bias; } + bool EnableHostPageableGather() const { return runtime_config_->enable_host_pageable_gather; } + bool IsCudaGraphEnabled() const { return runtime_config_->enable_cuda_graph; } bool IsArchAvailable(int arch) const { return GetDeviceProp().major >= arch; } + int GetDeviceId() const { return device_id_; } // Delegate to the base OpKernel::Info() which holds a safe copy of OpKernelInfo. // Do NOT store a reference to the constructor parameter — it becomes dangling. const OpKernelInfo& Info() const { return OpKernel::Info(); } diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index 2a878832dbd1d..f2d643eccbb87 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -17,6 +17,13 @@ #include "test/providers/provider_test_utils.h" #include "test/util/include/default_providers.h" +#ifdef USE_CUDA +#include "core/providers/cuda/cuda_execution_provider.h" +#include "core/providers/cuda/cuda_execution_provider_info.h" +#include "core/session/onnxruntime_session_options_config_keys.h" +#include "contrib_ops/cuda/quantization/gather_block_quantized.h" +#endif + namespace onnxruntime { namespace test { @@ -1365,6 +1372,119 @@ TEST(GatherBlockQuantizedOpTest, FpFloat16Output) { } #ifdef USE_CUDA +TEST(GatherBlockQuantizedOpTest, HostPageablePolicySelection) { + using cuda::GatherBlockQuantizedDataPolicy; + using cuda::SelectGatherBlockQuantizedDataPolicy; + + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, true, true, false, true), + GatherBlockQuantizedDataPolicy::DeviceCopy); + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, false, false, false, true), + GatherBlockQuantizedDataPolicy::DeviceCopy); + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, true), + GatherBlockQuantizedDataPolicy::DirectHost); + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, false, false, true), + GatherBlockQuantizedDataPolicy::DeviceCopy); + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, true, true), + GatherBlockQuantizedDataPolicy::DeviceCopy); + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, false), + GatherBlockQuantizedDataPolicy::DeviceCopy); +} + +TEST(GatherBlockQuantizedOpTest, HostPageableProviderOptionRoundTripAndHash) { + CUDAExecutionProviderInfo default_info = + CUDAExecutionProviderInfo::FromProviderOptions({}); + EXPECT_FALSE(default_info.enable_host_pageable_gather); + EXPECT_EQ(OrtCUDAProviderOptionsV2{}.enable_host_pageable_gather, 0); + + CUDAExecutionProviderInfo disabled_info = + CUDAExecutionProviderInfo::FromProviderOptions({{"enable_host_pageable_gather", "0"}}); + EXPECT_FALSE(disabled_info.enable_host_pageable_gather); + + CUDAExecutionProviderInfo enabled_info = + CUDAExecutionProviderInfo::FromProviderOptions({{"enable_host_pageable_gather", "1"}}); + EXPECT_TRUE(enabled_info.enable_host_pageable_gather); + const ProviderOptions serialized = CUDAExecutionProviderInfo::ToProviderOptions(enabled_info); + ASSERT_EQ(serialized.count("enable_host_pageable_gather"), 1u); + EXPECT_EQ(serialized.at("enable_host_pageable_gather"), "1"); + EXPECT_TRUE(CUDAExecutionProviderInfo::FromProviderOptions(serialized).enable_host_pageable_gather); + EXPECT_NE(std::hash{}(disabled_info), + std::hash{}(enabled_info)); +} + +#if !defined(DISABLE_FLOAT8_TYPES) +TEST(GatherBlockQuantizedOpTest, FpFallbackWithPrepackingDisabledCuda) { + if (!HasCudaEnvironment(0)) { + GTEST_SKIP() << "CUDA not available"; + } + + OpTester test("GatherBlockQuantized", 1, kMSDomain); + test.AddAttribute("gather_axis", 0); + test.AddAttribute("quantize_axis", 1); + test.AddAttribute("block_size", 0); + test.AddInput("data", {2, 2}, + {Float8E4M3FN(1.0f), Float8E4M3FN(2.0f), + Float8E4M3FN(3.0f), Float8E4M3FN(4.0f)}, + true); + test.AddInput("indices", {2}, {1, 0}); + test.AddInput("scales", {2, 1}, {2.0f, 0.5f}); + test.AddOutput("output", {2, 2}, {1.5f, 2.0f, 2.0f, 4.0f}); + + SessionOptions session_options; + ASSERT_STATUS_OK( + session_options.config_options.AddConfigEntry(kOrtSessionOptionsConfigDisablePrepacking, "1")); + test.Config(session_options); + + CUDAExecutionProviderInfo info; + info.enable_host_pageable_gather = false; + std::vector> providers; + providers.push_back(std::make_unique(info)); + test.ConfigEps(std::move(providers)); + test.RunWithConfig(); +} + +TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCuda) { + if (!HasCudaEnvironment(0)) { + GTEST_SKIP() << "CUDA not available"; + } + + int pageable_memory_access = 0; + int uses_host_page_tables = 0; +#if defined(CUDA_VERSION) && CUDA_VERSION >= 10020 + if (cudaDeviceGetAttribute(&pageable_memory_access, cudaDevAttrPageableMemoryAccess, 0) != cudaSuccess || + cudaDeviceGetAttribute(&uses_host_page_tables, cudaDevAttrPageableMemoryAccessUsesHostPageTables, 0) != + cudaSuccess) { + cudaGetLastError(); + GTEST_SKIP() << "CUDA pageable-memory attributes are unavailable"; + } +#endif + if (pageable_memory_access == 0 || uses_host_page_tables == 0) { + GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; + } + + OpTester test("GatherBlockQuantized", 1, kMSDomain); + test.AddAttribute("gather_axis", 0); + test.AddAttribute("quantize_axis", 1); + test.AddAttribute("block_size", 0); + test.AddInput("data", {2, 2}, + {Float8E4M3FN(1.0f), Float8E4M3FN(2.0f), + Float8E4M3FN(3.0f), Float8E4M3FN(4.0f)}, + true); + test.AddInput("indices", {3}, {-1, 0, 1}); + test.AddInput("scales", {2, 1}, {MLFloat16(2.0f), MLFloat16(0.5f)}); + test.AddOutput("output", {3, 2}, + {MLFloat16(1.5f), MLFloat16(2.0f), + MLFloat16(2.0f), MLFloat16(4.0f), + MLFloat16(1.5f), MLFloat16(2.0f)}); + + CUDAExecutionProviderInfo info; + info.enable_host_pageable_gather = true; + std::vector> providers; + providers.push_back(std::make_unique(info)); + test.ConfigEps(std::move(providers)); + test.RunWithConfig(); +} +#endif + TEST(GatherBlockQuantizedOpTest, FpBFloat16OutputCuda) { if (!HasCudaEnvironment(0)) { GTEST_SKIP() << "CUDA not available"; From e879ff661c1ae2bac19bdec0c8668ba13b9beb3b Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 16 Sep 2026 02:00:58 +0000 Subject: [PATCH 02/10] Harden CUDA gather fallback and tests Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- docs/cuda_host_pageable_gather.md | 5 ++++ .../quantization/gather_block_quantized.cc | 26 +++++++++++-------- .../gather_block_quantized_op_test.cc | 10 ++++--- 3 files changed, 27 insertions(+), 14 deletions(-) diff --git a/docs/cuda_host_pageable_gather.md b/docs/cuda_host_pageable_gather.md index f30638a079f1b..b4ab9af1dbae4 100644 --- a/docs/cuda_host_pageable_gather.md +++ b/docs/cuda_host_pageable_gather.md @@ -14,5 +14,10 @@ prefetch, hash, scan, or copy the complete initializer, and the initializer is n Fallback copies are allocated by the CUDA initializer allocator; they are not currently included in capacity-aware partitioning estimates. +Because the CPU memory contract is part of the static CUDA kernel registration, non-constant input data is staged to +CUDA on every run, including for non-FP8 types. Such runtime input is not supported during CUDA Graph capture. Multiple +nodes that use the same initializer can also create separate fallback copies; direct host access does not duplicate the +initializer. + This mode primarily reduces GPU memory capacity requirements. Performance depends on storage latency and operating system page-cache state, so cold prefill can be slower and less predictable than using resident GPU memory. diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc index 5a733b2abfaf0..2e02548d0a8f6 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc @@ -13,17 +13,17 @@ namespace contrib { namespace cuda { using namespace onnxruntime::cuda; -#define REGISTER_GATHERBLOCKQUANTIZED(T1, T2, Tind) \ - ONNX_OPERATOR_THREE_TYPED_KERNEL_EX( \ - GatherBlockQuantized, \ - kMSDomain, 1, \ - T1, T2, Tind, \ - kCudaExecutionProvider, \ - (*KernelDefBuilder::Create()) \ - .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ - .TypeConstraint("T2", DataTypeImpl::GetTensorType()) \ - .TypeConstraint("Tind", DataTypeImpl::GetTensorType()) \ - .InputMemoryType(OrtMemTypeCPUInput, 0), \ +#define REGISTER_GATHERBLOCKQUANTIZED(T1, T2, Tind) \ + ONNX_OPERATOR_THREE_TYPED_KERNEL_EX( \ + GatherBlockQuantized, \ + kMSDomain, 1, \ + T1, T2, Tind, \ + kCudaExecutionProvider, \ + (*KernelDefBuilder::Create()) \ + .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ + .TypeConstraint("T2", DataTypeImpl::GetTensorType()) \ + .TypeConstraint("Tind", DataTypeImpl::GetTensorType()) \ + .InputMemoryType(OrtMemTypeCPUInput, 0), \ GatherBlockQuantized); REGISTER_GATHERBLOCKQUANTIZED(uint8_t, float, int32_t); @@ -249,6 +249,10 @@ Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) ORT_RETURN_IF_NOT(data != nullptr && data->Location().device.Type() == OrtDevice::CPU, "GatherBlockQuantized input 0 must reside in CPU memory."); if (data->SizeInBytes() != 0) { + cudaStreamCaptureStatus capture_status = cudaStreamCaptureStatusNone; + CUDA_RETURN_IF_ERROR(cudaStreamIsCapturing(Stream(ctx), &capture_status)); + ORT_RETURN_IF_NOT(capture_status == cudaStreamCaptureStatusNone, + "CUDA Graph capture requires GatherBlockQuantized input 0 to be a constant initializer."); runtime_data = GetScratchBuffer(data->SizeInBytes(), GetComputeStream(ctx)); ORT_RETURN_IF_NOT(runtime_data != nullptr, "Failed to allocate CUDA storage for GatherBlockQuantized input 0."); CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(runtime_data.get(), data->DataRaw(), data->SizeInBytes(), diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index f2d643eccbb87..f861e6af8b2aa 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -18,10 +18,12 @@ #include "test/util/include/default_providers.h" #ifdef USE_CUDA +#include "contrib_ops/cuda/quantization/gather_block_quantized.h" +#ifndef BUILD_CUDA_EP_AS_PLUGIN #include "core/providers/cuda/cuda_execution_provider.h" #include "core/providers/cuda/cuda_execution_provider_info.h" #include "core/session/onnxruntime_session_options_config_keys.h" -#include "contrib_ops/cuda/quantization/gather_block_quantized.h" +#endif #endif namespace onnxruntime { @@ -1373,8 +1375,8 @@ TEST(GatherBlockQuantizedOpTest, FpFloat16Output) { #ifdef USE_CUDA TEST(GatherBlockQuantizedOpTest, HostPageablePolicySelection) { - using cuda::GatherBlockQuantizedDataPolicy; - using cuda::SelectGatherBlockQuantizedDataPolicy; + using contrib::cuda::GatherBlockQuantizedDataPolicy; + using contrib::cuda::SelectGatherBlockQuantizedDataPolicy; EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, true, true, false, true), GatherBlockQuantizedDataPolicy::DeviceCopy); @@ -1390,6 +1392,7 @@ TEST(GatherBlockQuantizedOpTest, HostPageablePolicySelection) { GatherBlockQuantizedDataPolicy::DeviceCopy); } +#ifndef BUILD_CUDA_EP_AS_PLUGIN TEST(GatherBlockQuantizedOpTest, HostPageableProviderOptionRoundTripAndHash) { CUDAExecutionProviderInfo default_info = CUDAExecutionProviderInfo::FromProviderOptions({}); @@ -1484,6 +1487,7 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCuda) { test.RunWithConfig(); } #endif +#endif TEST(GatherBlockQuantizedOpTest, FpBFloat16OutputCuda) { if (!HasCudaEnvironment(0)) { From 9d263675aa0259d96c396472a8e8e39cf3548a49 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 16 Sep 2026 02:05:10 +0000 Subject: [PATCH 03/10] Complete plugin and constant-input coverage Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- .../quantization/gather_block_quantized.cc | 2 +- .../quantization/gather_block_quantized.h | 4 +- .../gather_block_quantized_op_test.cc | 42 ++++++++++++------- onnxruntime/test/util/default_providers.cc | 2 + 4 files changed, 32 insertions(+), 18 deletions(-) diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc index 2e02548d0a8f6..13bf9c3e73c82 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc @@ -106,7 +106,7 @@ GatherBlockQuantized::GatherBlockQuantized(const OpKernelInfo& inf direct_host_data_ = SelectGatherBlockQuantizedDataPolicy(option_enabled, pageable_memory_access != 0, uses_host_page_tables != 0, IsCudaGraphEnabled(), - IsFp8QuantizedV) == + IsFp8QuantizedV, data_is_constant_) == GatherBlockQuantizedDataPolicy::DirectHost; if (option_enabled && IsFp8QuantizedV && !direct_host_data_) { LOGS_DEFAULT(WARNING) diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h index 85588f6164b53..1d3bea5f685eb 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h @@ -22,9 +22,9 @@ enum class GatherBlockQuantizedDataPolicy { constexpr GatherBlockQuantizedDataPolicy SelectGatherBlockQuantizedDataPolicy( bool option_enabled, bool pageable_memory_access, bool uses_host_page_tables, - bool cuda_graph_enabled, bool is_fp8) { + bool cuda_graph_enabled, bool is_fp8, bool is_constant_initializer) { return option_enabled && pageable_memory_access && uses_host_page_tables && - !cuda_graph_enabled && is_fp8 + !cuda_graph_enabled && is_fp8 && is_constant_initializer ? GatherBlockQuantizedDataPolicy::DirectHost : GatherBlockQuantizedDataPolicy::DeviceCopy; } diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index f861e6af8b2aa..7626e801e8f35 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -19,10 +19,9 @@ #ifdef USE_CUDA #include "contrib_ops/cuda/quantization/gather_block_quantized.h" +#include "core/session/onnxruntime_session_options_config_keys.h" #ifndef BUILD_CUDA_EP_AS_PLUGIN -#include "core/providers/cuda/cuda_execution_provider.h" #include "core/providers/cuda/cuda_execution_provider_info.h" -#include "core/session/onnxruntime_session_options_config_keys.h" #endif #endif @@ -1378,17 +1377,19 @@ TEST(GatherBlockQuantizedOpTest, HostPageablePolicySelection) { using contrib::cuda::GatherBlockQuantizedDataPolicy; using contrib::cuda::SelectGatherBlockQuantizedDataPolicy; - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, true, true, false, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, true, true, false, true, true), GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, false, false, false, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, false, false, false, true, true), GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, true, true), GatherBlockQuantizedDataPolicy::DirectHost); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, false, false, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, false, false, true, true), + GatherBlockQuantizedDataPolicy::DeviceCopy); + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, true, true, true), GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, true, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, false, true), GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, false), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, true, false), GatherBlockQuantizedDataPolicy::DeviceCopy); } @@ -1414,12 +1415,21 @@ TEST(GatherBlockQuantizedOpTest, HostPageableProviderOptionRoundTripAndHash) { std::hash{}(enabled_info)); } +#endif + #if !defined(DISABLE_FLOAT8_TYPES) TEST(GatherBlockQuantizedOpTest, FpFallbackWithPrepackingDisabledCuda) { if (!HasCudaEnvironment(0)) { GTEST_SKIP() << "CUDA not available"; } + OrtCUDAProviderOptionsV2 info; + info.enable_host_pageable_gather = 0; + auto cuda_ep = CudaExecutionProviderWithOptions(&info); + if (cuda_ep == nullptr) { + GTEST_SKIP() << "CUDA EP not available"; + } + OpTester test("GatherBlockQuantized", 1, kMSDomain); test.AddAttribute("gather_axis", 0); test.AddAttribute("quantize_axis", 1); @@ -1437,10 +1447,8 @@ TEST(GatherBlockQuantizedOpTest, FpFallbackWithPrepackingDisabledCuda) { session_options.config_options.AddConfigEntry(kOrtSessionOptionsConfigDisablePrepacking, "1")); test.Config(session_options); - CUDAExecutionProviderInfo info; - info.enable_host_pageable_gather = false; std::vector> providers; - providers.push_back(std::make_unique(info)); + providers.push_back(std::move(cuda_ep)); test.ConfigEps(std::move(providers)); test.RunWithConfig(); } @@ -1464,6 +1472,13 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCuda) { GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; } + OrtCUDAProviderOptionsV2 info; + info.enable_host_pageable_gather = 1; + auto cuda_ep = CudaExecutionProviderWithOptions(&info); + if (cuda_ep == nullptr) { + GTEST_SKIP() << "CUDA EP not available"; + } + OpTester test("GatherBlockQuantized", 1, kMSDomain); test.AddAttribute("gather_axis", 0); test.AddAttribute("quantize_axis", 1); @@ -1479,15 +1494,12 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCuda) { MLFloat16(2.0f), MLFloat16(4.0f), MLFloat16(1.5f), MLFloat16(2.0f)}); - CUDAExecutionProviderInfo info; - info.enable_host_pageable_gather = true; std::vector> providers; - providers.push_back(std::make_unique(info)); + providers.push_back(std::move(cuda_ep)); test.ConfigEps(std::move(providers)); test.RunWithConfig(); } #endif -#endif TEST(GatherBlockQuantizedOpTest, FpBFloat16OutputCuda) { if (!HasCudaEnvironment(0)) { diff --git a/onnxruntime/test/util/default_providers.cc b/onnxruntime/test/util/default_providers.cc index 26a65a44c74c5..cb4610c20c30a 100644 --- a/onnxruntime/test/util/default_providers.cc +++ b/onnxruntime/test/util/default_providers.cc @@ -55,6 +55,8 @@ std::unique_ptr CudaPluginExecutionProviderWithOptions(const AddCudaPluginOption(config_options, "use_tf32", std::to_string(provider_options->use_tf32)); AddCudaPluginOption(config_options, "fuse_conv_bias", std::to_string(provider_options->fuse_conv_bias)); AddCudaPluginOption(config_options, "sdpa_kernel", std::to_string(provider_options->sdpa_kernel)); + AddCudaPluginOption(config_options, "enable_host_pageable_gather", + std::to_string(provider_options->enable_host_pageable_gather)); } return dynamic_plugin_ep_infra::MakeEp(nullptr, &config_options); From e2126e315b288553c9206386d85043a9d32e05be Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 16 Sep 2026 03:50:44 +0000 Subject: [PATCH 04/10] Enable host-pageable gather in CUDA Graphs Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- docs/cuda_host_pageable_gather.md | 12 +- .../quantization/gather_block_quantized.cc | 24 +- .../quantization/gather_block_quantized.h | 4 +- .../gather_block_quantized_op_test.cc | 210 +++++++++++++++++- 4 files changed, 225 insertions(+), 25 deletions(-) diff --git a/docs/cuda_host_pageable_gather.md b/docs/cuda_host_pageable_gather.md index b4ab9af1dbae4..fc12d056f246d 100644 --- a/docs/cuda_host_pageable_gather.md +++ b/docs/cuda_host_pageable_gather.md @@ -4,9 +4,8 @@ The CUDA Execution Provider option `enable_host_pageable_gather` enables direct `com.microsoft::GatherBlockQuantized` input data. It is disabled by default. Direct access requires a CUDA device that reports both `cudaDevAttrPageableMemoryAccess` and -`cudaDevAttrPageableMemoryAccessUsesHostPageTables`. CUDA Graph capture is not currently supported with this mode. -If either device capability is unavailable or CUDA Graphs are enabled, ONNX Runtime emits a warning and makes one -persistent CUDA copy of a constant input instead. +`cudaDevAttrPageableMemoryAccessUsesHostPageTables`. If either device capability is unavailable, ONNX Runtime emits a +warning and makes one persistent CUDA copy of a constant input instead. The option does not create or manage a file mapping. The model initializer must already be supplied as CPU memory, such as a file-backed mapping, and that mapping remains live for the session. The direct path does not register, @@ -14,6 +13,13 @@ prefetch, hash, scan, or copy the complete initializer, and the initializer is n Fallback copies are allocated by the CUDA initializer allocator; they are not currently included in capacity-aware partitioning estimates. +CUDA Graph capture is supported for direct host-pageable access. The initializer mapping must remain alive at the same +virtual address until the graph executable is destroyed, and indices, scales, and outputs must retain their normal +CUDA Graph-stable addresses. The direct path performs no mapping, registration, allocation, copy, capability query, or +synchronization during capture. Persistent-copy fallbacks are prepared during prepacking or an uncaptured warmup run; +if lazy fallback initialization is still required when capture starts, the run fails instead of allocating or copying +during capture. + Because the CPU memory contract is part of the static CUDA kernel registration, non-constant input data is staged to CUDA on every run, including for non-FP8 types. Such runtime input is not supported during CUDA Graph capture. Multiple nodes that use the same initializer can also create separate fallback copies; direct host access does not duplicate the diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc index 13bf9c3e73c82..e6947c6a01c7f 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc @@ -105,14 +105,14 @@ GatherBlockQuantized::GatherBlockQuantized(const OpKernelInfo& inf const bool option_enabled = EnableHostPageableGather(); direct_host_data_ = SelectGatherBlockQuantizedDataPolicy(option_enabled, pageable_memory_access != 0, - uses_host_page_tables != 0, IsCudaGraphEnabled(), - IsFp8QuantizedV, data_is_constant_) == + uses_host_page_tables != 0, IsFp8QuantizedV, + data_is_constant_) == GatherBlockQuantizedDataPolicy::DirectHost; if (option_enabled && IsFp8QuantizedV && !direct_host_data_) { LOGS_DEFAULT(WARNING) << "enable_host_pageable_gather was requested, but direct host-pageable GatherBlockQuantized " - "access is unavailable because the CUDA device lacks pageable memory access through host page tables " - "or CUDA Graph capture is enabled. Falling back to a persistent CUDA copy."; + "access is unavailable because input 0 is not a constant initializer or the CUDA device lacks pageable " + "memory access through host page tables. Falling back to a persistent CUDA copy."; } // If block size is set, it has to be no smaller than 16 and must be power of 2. @@ -169,14 +169,8 @@ Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) const Tensor* scales = ctx->Input(2); const Tensor* zero_points = ctx->Input(3); - TensorShapeVector data_shape_storage; - if (data != nullptr) { - data_shape_storage.assign(data->Shape().GetDims().begin(), data->Shape().GetDims().end()); - } else { - std::lock_guard lock(device_data_mutex_); - data_shape_storage = data_shape_; - } - const gsl::span data_shape{data_shape_storage}; + const gsl::span data_shape = + data != nullptr ? data->Shape().GetDims() : gsl::span{data_shape_}; int64_t data_rank = static_cast(data_shape.size()); const int64_t gather_axis = HandleNegativeAxis(gather_axis_, data_rank); const int64_t quantize_axis = HandleNegativeAxis(quantize_axis_, data_rank); @@ -241,6 +235,12 @@ Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) { std::lock_guard lock(device_data_mutex_); if (device_data_ == nullptr && data != nullptr && data->SizeInBytes() != 0) { + cudaStreamCaptureStatus capture_status = cudaStreamCaptureStatusNone; + CUDA_RETURN_IF_ERROR(cudaStreamIsCapturing(Stream(ctx), &capture_status)); + ORT_RETURN_IF_NOT( + capture_status == cudaStreamCaptureStatusNone, + "GatherBlockQuantized cannot initialize its persistent CUDA fallback copy during CUDA Graph capture. " + "Enable prepacking or run an uncaptured warmup iteration before capture."); ORT_RETURN_IF_ERROR(CreateDeviceCopy(*data, Info().GetAllocator(OrtMemTypeDefault))); } data_ptr = static_cast(device_data_.get()); diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h index 1d3bea5f685eb..ad9f6bf30b48f 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h @@ -22,9 +22,9 @@ enum class GatherBlockQuantizedDataPolicy { constexpr GatherBlockQuantizedDataPolicy SelectGatherBlockQuantizedDataPolicy( bool option_enabled, bool pageable_memory_access, bool uses_host_page_tables, - bool cuda_graph_enabled, bool is_fp8, bool is_constant_initializer) { + bool is_fp8, bool is_constant_initializer) { return option_enabled && pageable_memory_access && uses_host_page_tables && - !cuda_graph_enabled && is_fp8 && is_constant_initializer + is_fp8 && is_constant_initializer ? GatherBlockQuantizedDataPolicy::DirectHost : GatherBlockQuantizedDataPolicy::DeviceCopy; } diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index 7626e801e8f35..0d68999f6da5d 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -2,6 +2,9 @@ // Licensed under the MIT License. #include +#include +#include +#include #include #include #include @@ -10,6 +13,10 @@ #include #include +#ifndef _WIN32 +#include +#endif + #include "core/common/common.h" #include "core/framework/execution_provider.h" #include "test/common/cuda_op_test_utils.h" @@ -19,7 +26,11 @@ #ifdef USE_CUDA #include "contrib_ops/cuda/quantization/gather_block_quantized.h" +#include "core/graph/model.h" +#include "core/platform/env.h" +#include "core/session/inference_session.h" #include "core/session/onnxruntime_session_options_config_keys.h" +#include "test/util/include/temp_dir.h" #ifndef BUILD_CUDA_EP_AS_PLUGIN #include "core/providers/cuda/cuda_execution_provider_info.h" #endif @@ -1377,19 +1388,17 @@ TEST(GatherBlockQuantizedOpTest, HostPageablePolicySelection) { using contrib::cuda::GatherBlockQuantizedDataPolicy; using contrib::cuda::SelectGatherBlockQuantizedDataPolicy; - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, true, true, false, true, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, true, true, true, true), GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, false, false, false, true, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(false, false, false, true, true), GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, true, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, true, true), GatherBlockQuantizedDataPolicy::DirectHost); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, false, false, true, true), - GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, true, true, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, false, true, true), GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, false, true), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, true), GatherBlockQuantizedDataPolicy::DeviceCopy); - EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, true, false), + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, true, false), GatherBlockQuantizedDataPolicy::DeviceCopy); } @@ -1499,6 +1508,191 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCuda) { test.ConfigEps(std::move(providers)); test.RunWithConfig(); } + +TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCudaGraph) { + if (!HasCudaEnvironment(0)) { + GTEST_SKIP() << "CUDA not available"; + } + + int pageable_memory_access = 0; + int uses_host_page_tables = 0; +#if defined(CUDA_VERSION) && CUDA_VERSION >= 10020 + if (cudaDeviceGetAttribute(&pageable_memory_access, cudaDevAttrPageableMemoryAccess, 0) != cudaSuccess || + cudaDeviceGetAttribute(&uses_host_page_tables, cudaDevAttrPageableMemoryAccessUsesHostPageTables, 0) != + cudaSuccess) { + cudaGetLastError(); + GTEST_SKIP() << "CUDA pageable-memory attributes are unavailable"; + } +#endif + if (pageable_memory_access == 0 || uses_host_page_tables == 0) { + GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; + } + + OrtCUDAProviderOptionsV2 info; + info.enable_cuda_graph = 1; + info.enable_host_pageable_gather = 1; + auto cuda_ep = CudaExecutionProviderWithOptions(&info); + if (cuda_ep == nullptr) { + GTEST_SKIP() << "CUDA EP not available"; + } + IExecutionProvider* cuda_ep_ptr = cuda_ep.get(); + + const std::vector data = { + Float8E4M3FN(1.0f), Float8E4M3FN(2.0f), + Float8E4M3FN(3.0f), Float8E4M3FN(4.0f), + Float8E4M3FN(5.0f), Float8E4M3FN(6.0f), + Float8E4M3FN(7.0f), Float8E4M3FN(8.0f)}; + const size_t data_bytes = data.size() * sizeof(data[0]); + const auto temp_dir_path = + std::filesystem::temp_directory_path() / + ("ort_gather_block_quantized_cuda_graph_" + + std::to_string(reinterpret_cast(&pageable_memory_access))); + TemporaryDirectory temp_dir(temp_dir_path.native()); + const auto data_path = temp_dir_path / "data.bin"; + { + std::ofstream data_file(data_path, std::ios::binary); + ASSERT_TRUE(data_file.good()); + data_file.write(reinterpret_cast(data.data()), static_cast(data_bytes)); + ASSERT_TRUE(data_file.good()); + } + + Env::MappedMemoryPtr mapped_memory; + ASSERT_STATUS_OK(Env::Default().MapFileIntoMemory(data_path.c_str(), 0, data_bytes, mapped_memory)); + const void* const mapped_address = mapped_memory.get(); + OrtMemoryInfo cpu_memory_info{CPU, OrtDeviceAllocator}; + Tensor mapped_tensor(DataTypeImpl::GetType(), TensorShape({4, 2}), + mapped_memory.get(), cpu_memory_info); + OrtValue mapped_data_value; + Tensor::InitOrtValue(std::move(mapped_tensor), mapped_data_value); + + std::unordered_map domain_to_version = {{onnxruntime::kMSDomain, 1}}; + std::vector model_specific_functions; + auto model = std::make_unique( + "gather_block_quantized_cuda_graph", true, ModelMetaData(), PathString(), + IOnnxRuntimeOpSchemaRegistryList(), domain_to_version, model_specific_functions, + DefaultLoggingManager().DefaultLogger(), ModelOptions(true, true)); + auto& graph = model->MainGraph(); + + std::vector tensor_types; + tensor_types.reserve(4); + auto add_tensor_type = [&](int elem_type, std::initializer_list dims) { + tensor_types.emplace_back(); + auto* type = &tensor_types.back(); + type->mutable_tensor_type()->set_elem_type(elem_type); + auto* shape = type->mutable_tensor_type()->mutable_shape(); + for (const int64_t dim : dims) { + shape->add_dim()->set_dim_value(dim); + } + return type; + }; + + auto& data_arg = graph.GetOrCreateNodeArg( + "data", add_tensor_type(ONNX_NAMESPACE::TensorProto_DataType_FLOAT8E4M3FN, {4, 2})); + auto& indices_arg = graph.GetOrCreateNodeArg( + "indices", add_tensor_type(ONNX_NAMESPACE::TensorProto_DataType_INT64, {2})); + auto& scales_arg = graph.GetOrCreateNodeArg( + "scales", add_tensor_type(ONNX_NAMESPACE::TensorProto_DataType_FLOAT, {4, 1})); + auto& output_arg = graph.GetOrCreateNodeArg( + "output", add_tensor_type(ONNX_NAMESPACE::TensorProto_DataType_FLOAT, {2, 2})); + + ONNX_NAMESPACE::TensorProto data_initializer; + data_initializer.set_name("data"); + data_initializer.set_data_type(ONNX_NAMESPACE::TensorProto_DataType_FLOAT8E4M3FN); + data_initializer.add_dims(4); + data_initializer.add_dims(2); + data_initializer.mutable_raw_data()->assign(data_bytes, '\0'); + graph.AddInitializedTensor(data_initializer); + + NodeAttributes attributes = { + {"block_size", utils::MakeAttribute("block_size", int64_t{0})}, + {"gather_axis", utils::MakeAttribute("gather_axis", int64_t{0})}, + {"quantize_axis", utils::MakeAttribute("quantize_axis", int64_t{1})}, + }; + auto& node = graph.AddNode("gather_block_quantized", "GatherBlockQuantized", + "CUDA Graph direct host-pageable test", + {&data_arg, &indices_arg, &scales_arg}, {&output_arg}, + &attributes, onnxruntime::kMSDomain); + node.SetExecutionProviderType(cuda_ep_ptr->Type()); + ASSERT_STATUS_OK(graph.Resolve()); + + std::string model_string; + ASSERT_TRUE(model->ToProto().SerializeToString(&model_string)); + std::stringstream model_stream(model_string); + + SessionOptions session_options; + ASSERT_STATUS_OK(session_options.AddInitializer("data", &mapped_data_value)); + { + InferenceSession session(session_options, GetEnvironment()); + ASSERT_STATUS_OK(session.RegisterExecutionProvider(std::move(cuda_ep))); + auto device_allocators = cuda_ep_ptr->CreatePreferredAllocators(); + const OrtMemoryInfo* device_memory_info = nullptr; + for (const auto& allocator : device_allocators) { + if (allocator->Info().device.Type() == OrtDevice::GPU && + allocator->Info().mem_type == OrtMemTypeDefault) { + device_memory_info = &allocator->Info(); + break; + } + } + ASSERT_NE(device_memory_info, nullptr); + ASSERT_STATUS_OK(session.Load(model_stream)); + ASSERT_STATUS_OK(session.Initialize()); + auto device_allocator = session.GetAllocator(*device_memory_info); + ASSERT_NE(device_allocator, nullptr); + + auto make_gpu_value = [&](const auto& values, const TensorShape& shape) { + using T = typename std::decay_t::value_type; + Tensor cpu_tensor(DataTypeImpl::GetType(), shape, const_cast(values.data()), cpu_memory_info); + Tensor gpu_tensor(DataTypeImpl::GetType(), shape, device_allocator); + ORT_THROW_IF_ERROR(cuda_ep_ptr->GetDataTransfer()->CopyTensor(cpu_tensor, gpu_tensor)); + OrtValue value; + Tensor::InitOrtValue(std::move(gpu_tensor), value); + return value; + }; + + std::vector indices = {0, 2}; + const std::vector scales = {1.0f, 0.5f, 2.0f, 0.25f}; + auto indices_value = make_gpu_value(indices, TensorShape({2})); + auto scales_value = make_gpu_value(scales, TensorShape({4, 1})); + auto output_value = make_gpu_value(std::vector(4), TensorShape({2, 2})); + + std::unique_ptr io_binding; + ASSERT_STATUS_OK(session.NewIOBinding(&io_binding)); + ASSERT_STATUS_OK(io_binding->BindInput("indices", indices_value)); + ASSERT_STATUS_OK(io_binding->BindInput("scales", scales_value)); + ASSERT_STATUS_OK(io_binding->BindOutput("output", output_value)); + + RunOptions run_options; + ASSERT_STATUS_OK(run_options.config_options.AddConfigEntry("gpu_graph_id", "1")); + for (int i = 0; i < 3; ++i) { + ASSERT_STATUS_OK(session.Run(run_options, *io_binding)); + } + ASSERT_TRUE(cuda_ep_ptr->IsGraphCaptured(1)); + + auto verify_output = [&](std::initializer_list expected) { + ASSERT_EQ(cudaSuccess, cudaDeviceSynchronize()); + std::vector actual(expected.size()); + Tensor cpu_output(DataTypeImpl::GetType(), TensorShape({2, 2}), actual.data(), cpu_memory_info); + ASSERT_STATUS_OK(cuda_ep_ptr->GetDataTransfer()->CopyTensor(output_value.Get(), cpu_output)); + EXPECT_EQ(actual, std::vector(expected.begin(), expected.end())); + }; + verify_output({1.0f, 2.0f, 10.0f, 12.0f}); + + indices = {3, 1}; + ASSERT_EQ(cudaSuccess, + cudaMemcpy(indices_value.GetMutable()->MutableData(), indices.data(), + indices.size() * sizeof(indices[0]), cudaMemcpyHostToDevice)); + ASSERT_STATUS_OK(session.Run(run_options, *io_binding)); + verify_output({1.75f, 2.0f, 1.5f, 2.0f}); + +#ifndef _WIN32 + ASSERT_EQ(0, madvise(mapped_memory.get(), data_bytes, MADV_DONTNEED)); + ASSERT_STATUS_OK(session.Run(run_options, *io_binding)); + verify_output({1.75f, 2.0f, 1.5f, 2.0f}); +#endif + + EXPECT_EQ(mapped_address, mapped_memory.get()); + } +} #endif TEST(GatherBlockQuantizedOpTest, FpBFloat16OutputCuda) { From 4d9c851f310c786db10888bc977e146d514ebd8f Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 16 Sep 2026 03:56:29 +0000 Subject: [PATCH 05/10] Strengthen direct host CUDA Graph coverage Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- .../contrib_ops/gather_block_quantized_op_test.cc | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index 0d68999f6da5d..56ea71eab90dd 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -1394,6 +1394,8 @@ TEST(GatherBlockQuantizedOpTest, HostPageablePolicySelection) { GatherBlockQuantizedDataPolicy::DeviceCopy); EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, true, true), GatherBlockQuantizedDataPolicy::DirectHost); + EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, false, true, true, true), + GatherBlockQuantizedDataPolicy::DeviceCopy); EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, false, true, true), GatherBlockQuantizedDataPolicy::DeviceCopy); EXPECT_EQ(SelectGatherBlockQuantizedDataPolicy(true, true, true, false, true), @@ -1558,7 +1560,6 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCudaGraph) { Env::MappedMemoryPtr mapped_memory; ASSERT_STATUS_OK(Env::Default().MapFileIntoMemory(data_path.c_str(), 0, data_bytes, mapped_memory)); - const void* const mapped_address = mapped_memory.get(); OrtMemoryInfo cpu_memory_info{CPU, OrtDeviceAllocator}; Tensor mapped_tensor(DataTypeImpl::GetType(), TensorShape({4, 2}), mapped_memory.get(), cpu_memory_info); @@ -1685,12 +1686,18 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCudaGraph) { verify_output({1.75f, 2.0f, 1.5f, 2.0f}); #ifndef _WIN32 + auto* mapped_data = reinterpret_cast(mapped_memory.get()); + mapped_data[2] = Float8E4M3FN(6.0f); + mapped_data[3] = Float8E4M3FN(8.0f); + mapped_data[6] = Float8E4M3FN(2.0f); + mapped_data[7] = Float8E4M3FN(4.0f); + ASSERT_STATUS_OK(session.Run(run_options, *io_binding)); + verify_output({0.5f, 1.0f, 3.0f, 4.0f}); + ASSERT_EQ(0, madvise(mapped_memory.get(), data_bytes, MADV_DONTNEED)); ASSERT_STATUS_OK(session.Run(run_options, *io_binding)); verify_output({1.75f, 2.0f, 1.5f, 2.0f}); #endif - - EXPECT_EQ(mapped_address, mapped_memory.get()); } } #endif From 4eb6a0816e6f939f54371af6c67f2522affb574e Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 16 Sep 2026 06:35:50 +0000 Subject: [PATCH 06/10] Address CUDA gather review feedback Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- docs/cuda_host_pageable_gather.md | 8 +-- .../providers/cuda/cuda_provider_options.h | 1 - .../quantization/gather_block_quantized.cc | 68 +++++++++---------- .../quantization/gather_block_quantized.h | 1 + .../cuda/cuda_execution_provider_info.cc | 2 - .../providers/cuda/cuda_provider_factory.cc | 3 - .../gather_block_quantized_op_test.cc | 18 ++--- onnxruntime/test/util/default_providers.cc | 24 ++++++- .../test/util/include/default_providers.h | 1 + 9 files changed, 65 insertions(+), 61 deletions(-) diff --git a/docs/cuda_host_pageable_gather.md b/docs/cuda_host_pageable_gather.md index fc12d056f246d..cc5e9b3260ed0 100644 --- a/docs/cuda_host_pageable_gather.md +++ b/docs/cuda_host_pageable_gather.md @@ -5,7 +5,7 @@ The CUDA Execution Provider option `enable_host_pageable_gather` enables direct Direct access requires a CUDA device that reports both `cudaDevAttrPageableMemoryAccess` and `cudaDevAttrPageableMemoryAccessUsesHostPageTables`. If either device capability is unavailable, ONNX Runtime emits a -warning and makes one persistent CUDA copy of a constant input instead. +warning and uses the standard CUDA input path. The option does not create or manage a file mapping. The model initializer must already be supplied as CPU memory, such as a file-backed mapping, and that mapping remains live for the session. The direct path does not register, @@ -20,10 +20,8 @@ synchronization during capture. Persistent-copy fallbacks are prepared during pr if lazy fallback initialization is still required when capture starts, the run fails instead of allocating or copying during capture. -Because the CPU memory contract is part of the static CUDA kernel registration, non-constant input data is staged to -CUDA on every run, including for non-FP8 types. Such runtime input is not supported during CUDA Graph capture. Multiple -nodes that use the same initializer can also create separate fallback copies; direct host access does not duplicate the -initializer. +Non-constant and non-FP8 input data retains the standard GPU-input contract. Multiple nodes that use the same constant +initializer can create separate fallback copies; direct host access does not duplicate the initializer. This mode primarily reduces GPU memory capacity requirements. Performance depends on storage latency and operating system page-cache state, so cold prefill can be slower and less predictable than using resident GPU memory. diff --git a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h index 3df52b95a95bf..7a44cb8e76daf 100644 --- a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h +++ b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h @@ -40,5 +40,4 @@ 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 - int enable_host_pageable_gather = 0; // Enable direct host-pageable FP8 GatherBlockQuantized access. }; diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc index e6947c6a01c7f..fca1bd093a0d0 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc @@ -13,17 +13,16 @@ namespace contrib { namespace cuda { using namespace onnxruntime::cuda; -#define REGISTER_GATHERBLOCKQUANTIZED(T1, T2, Tind) \ - ONNX_OPERATOR_THREE_TYPED_KERNEL_EX( \ - GatherBlockQuantized, \ - kMSDomain, 1, \ - T1, T2, Tind, \ - kCudaExecutionProvider, \ - (*KernelDefBuilder::Create()) \ - .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ - .TypeConstraint("T2", DataTypeImpl::GetTensorType()) \ - .TypeConstraint("Tind", DataTypeImpl::GetTensorType()) \ - .InputMemoryType(OrtMemTypeCPUInput, 0), \ +#define REGISTER_GATHERBLOCKQUANTIZED(T1, T2, Tind) \ + ONNX_OPERATOR_THREE_TYPED_KERNEL_EX( \ + GatherBlockQuantized, \ + kMSDomain, 1, \ + T1, T2, Tind, \ + kCudaExecutionProvider, \ + (*KernelDefBuilder::Create()) \ + .TypeConstraint("T1", DataTypeImpl::GetTensorType()) \ + .TypeConstraint("T2", DataTypeImpl::GetTensorType()) \ + .TypeConstraint("Tind", DataTypeImpl::GetTensorType()), \ GatherBlockQuantized); REGISTER_GATHERBLOCKQUANTIZED(uint8_t, float, int32_t); @@ -112,7 +111,7 @@ GatherBlockQuantized::GatherBlockQuantized(const OpKernelInfo& inf LOGS_DEFAULT(WARNING) << "enable_host_pageable_gather was requested, but direct host-pageable GatherBlockQuantized " "access is unavailable because input 0 is not a constant initializer or the CUDA device lacks pageable " - "memory access through host page tables. Falling back to a persistent CUDA copy."; + "memory access through host page tables. Using the standard CUDA input path."; } // If block size is set, it has to be no smaller than 16 and must be power of 2. @@ -131,16 +130,18 @@ Status GatherBlockQuantized::CreateDeviceCopy(const Tensor& tensor ORT_RETURN_IF_NOT(tensor.Location().device.Type() == OrtDevice::CPU, "GatherBlockQuantized input 0 must reside in CPU memory."); - data_shape_.assign(tensor.Shape().GetDims().begin(), tensor.Shape().GetDims().end()); const size_t bytes = tensor.SizeInBytes(); if (bytes == 0) { + data_shape_.assign(tensor.Shape().GetDims().begin(), tensor.Shape().GetDims().end()); device_data_.reset(); return Status::OK(); } - device_data_ = IAllocator::MakeUniquePtr(alloc, bytes); - ORT_RETURN_IF_NOT(device_data_ != nullptr, "Failed to allocate persistent CUDA storage for GatherBlockQuantized."); - CUDA_RETURN_IF_ERROR(cudaMemcpy(device_data_.get(), tensor.DataRaw(), bytes, cudaMemcpyHostToDevice)); + auto device_data = IAllocator::MakeUniquePtr(alloc, bytes); + ORT_RETURN_IF_NOT(device_data != nullptr, "Failed to allocate persistent CUDA storage for GatherBlockQuantized."); + CUDA_RETURN_IF_ERROR(cudaMemcpy(device_data.get(), tensor.DataRaw(), bytes, cudaMemcpyHostToDevice)); + data_shape_.assign(tensor.Shape().GetDims().begin(), tensor.Shape().GetDims().end()); + device_data_ = std::move(device_data); return Status::OK(); } @@ -149,12 +150,17 @@ Status GatherBlockQuantized::PrePack( const Tensor& tensor, int input_idx, AllocatorPtr alloc, bool& is_packed, PrePackedWeights* prepacked_weights) { is_packed = false; - if (input_idx != 0 || direct_host_data_) { + if (input_idx != 0) { return Status::OK(); } std::lock_guard lock(device_data_mutex_); - ORT_RETURN_IF_ERROR(CreateDeviceCopy(tensor, std::move(alloc))); + if (direct_host_data_) { + direct_host_data_ptr_ = tensor.Data(); + data_shape_.assign(tensor.Shape().GetDims().begin(), tensor.Shape().GetDims().end()); + } else { + ORT_RETURN_IF_ERROR(CreateDeviceCopy(tensor, std::move(alloc))); + } is_packed = true; if (prepacked_weights != nullptr) { prepacked_weights->has_kernel_owned_packed_weights_ = true; @@ -225,16 +231,14 @@ Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) return Status::OK(); } - IAllocatorUniquePtr runtime_data; const T1* data_ptr = nullptr; if (direct_host_data_) { - ORT_RETURN_IF_NOT(data != nullptr && data->Location().device.Type() == OrtDevice::CPU, - "Direct host-pageable GatherBlockQuantized requires CPU-resident input 0."); - data_ptr = data->Data(); + data_ptr = direct_host_data_ptr_; } else if (data_is_constant_) { { std::lock_guard lock(device_data_mutex_); - if (device_data_ == nullptr && data != nullptr && data->SizeInBytes() != 0) { + if (device_data_ == nullptr && data != nullptr && + data->Location().device.Type() == OrtDevice::CPU && data->SizeInBytes() != 0) { cudaStreamCaptureStatus capture_status = cudaStreamCaptureStatusNone; CUDA_RETURN_IF_ERROR(cudaStreamIsCapturing(Stream(ctx), &capture_status)); ORT_RETURN_IF_NOT( @@ -243,22 +247,12 @@ Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) "Enable prepacking or run an uncaptured warmup iteration before capture."); ORT_RETURN_IF_ERROR(CreateDeviceCopy(*data, Info().GetAllocator(OrtMemTypeDefault))); } - data_ptr = static_cast(device_data_.get()); + data_ptr = device_data_ != nullptr ? static_cast(device_data_.get()) + : data == nullptr ? nullptr + : data->Data(); } } else { - ORT_RETURN_IF_NOT(data != nullptr && data->Location().device.Type() == OrtDevice::CPU, - "GatherBlockQuantized input 0 must reside in CPU memory."); - if (data->SizeInBytes() != 0) { - cudaStreamCaptureStatus capture_status = cudaStreamCaptureStatusNone; - CUDA_RETURN_IF_ERROR(cudaStreamIsCapturing(Stream(ctx), &capture_status)); - ORT_RETURN_IF_NOT(capture_status == cudaStreamCaptureStatusNone, - "CUDA Graph capture requires GatherBlockQuantized input 0 to be a constant initializer."); - runtime_data = GetScratchBuffer(data->SizeInBytes(), GetComputeStream(ctx)); - ORT_RETURN_IF_NOT(runtime_data != nullptr, "Failed to allocate CUDA storage for GatherBlockQuantized input 0."); - CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(runtime_data.get(), data->DataRaw(), data->SizeInBytes(), - cudaMemcpyHostToDevice, Stream(ctx))); - data_ptr = static_cast(runtime_data.get()); - } + data_ptr = data->Data(); } ORT_RETURN_IF_NOT(N == 0 || data_ptr != nullptr, "GatherBlockQuantized fallback has no device-resident input 0."); diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h index ad9f6bf30b48f..60a3f3ffaba42 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h @@ -47,6 +47,7 @@ class GatherBlockQuantized final : public CudaKernel { bool direct_host_data_; bool data_is_constant_; mutable std::mutex device_data_mutex_; + const T1* direct_host_data_ptr_{}; mutable IAllocatorUniquePtr device_data_; mutable TensorShapeVector data_shape_; }; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc index f1fb4396516d8..03cf02bfb95b7 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc @@ -217,8 +217,6 @@ ProviderOptions CUDAExecutionProviderInfo::ToProviderOptions(const OrtCUDAProvid {cuda::provider_option_names::kUseTF32, MakeStringWithClassicLocale(info.use_tf32)}, {cuda::provider_option_names::kFuseConvBias, MakeStringWithClassicLocale(info.fuse_conv_bias)}, {cuda::provider_option_names::kSdpaKernel, MakeStringWithClassicLocale(info.sdpa_kernel)}, - {cuda::provider_option_names::kEnableHostPageableGather, - MakeStringWithClassicLocale(info.enable_host_pageable_gather)}, }; return options; diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc index 9b789b8472a7e..27c7c9ff078d4 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc @@ -245,8 +245,6 @@ struct CUDA_Provider : Provider { info.use_ep_level_unified_stream = params->use_ep_level_unified_stream != 0; info.use_tf32 = params->use_tf32 != 0; info.sdpa_kernel = params->sdpa_kernel; - info.enable_host_pageable_gather = params->enable_host_pageable_gather != 0; - return std::make_shared(info); } @@ -281,7 +279,6 @@ struct CUDA_Provider : Provider { cuda_options.use_tf32 = internal_options.use_tf32; cuda_options.sdpa_kernel = internal_options.sdpa_kernel; cuda_options.fuse_conv_bias = internal_options.fuse_conv_bias; - cuda_options.enable_host_pageable_gather = internal_options.enable_host_pageable_gather; } ProviderOptions GetProviderOptions(const void* provider_options) override { diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index 56ea71eab90dd..0aebbd8f5086d 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -1409,7 +1409,6 @@ TEST(GatherBlockQuantizedOpTest, HostPageableProviderOptionRoundTripAndHash) { CUDAExecutionProviderInfo default_info = CUDAExecutionProviderInfo::FromProviderOptions({}); EXPECT_FALSE(default_info.enable_host_pageable_gather); - EXPECT_EQ(OrtCUDAProviderOptionsV2{}.enable_host_pageable_gather, 0); CUDAExecutionProviderInfo disabled_info = CUDAExecutionProviderInfo::FromProviderOptions({{"enable_host_pageable_gather", "0"}}); @@ -1434,9 +1433,8 @@ TEST(GatherBlockQuantizedOpTest, FpFallbackWithPrepackingDisabledCuda) { GTEST_SKIP() << "CUDA not available"; } - OrtCUDAProviderOptionsV2 info; - info.enable_host_pageable_gather = 0; - auto cuda_ep = CudaExecutionProviderWithOptions(&info); + auto cuda_ep = CudaExecutionProviderWithOptions( + ProviderOptions{{"enable_host_pageable_gather", "0"}}); if (cuda_ep == nullptr) { GTEST_SKIP() << "CUDA EP not available"; } @@ -1483,9 +1481,8 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCuda) { GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; } - OrtCUDAProviderOptionsV2 info; - info.enable_host_pageable_gather = 1; - auto cuda_ep = CudaExecutionProviderWithOptions(&info); + auto cuda_ep = CudaExecutionProviderWithOptions( + ProviderOptions{{"enable_host_pageable_gather", "1"}}); if (cuda_ep == nullptr) { GTEST_SKIP() << "CUDA EP not available"; } @@ -1530,10 +1527,9 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCudaGraph) { GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; } - OrtCUDAProviderOptionsV2 info; - info.enable_cuda_graph = 1; - info.enable_host_pageable_gather = 1; - auto cuda_ep = CudaExecutionProviderWithOptions(&info); + auto cuda_ep = CudaExecutionProviderWithOptions( + ProviderOptions{{"enable_cuda_graph", "1"}, + {"enable_host_pageable_gather", "1"}}); if (cuda_ep == nullptr) { GTEST_SKIP() << "CUDA EP not available"; } diff --git a/onnxruntime/test/util/default_providers.cc b/onnxruntime/test/util/default_providers.cc index cb4610c20c30a..ef9b71720316b 100644 --- a/onnxruntime/test/util/default_providers.cc +++ b/onnxruntime/test/util/default_providers.cc @@ -13,6 +13,10 @@ #include "core/providers/coreml/coreml_provider_factory.h" #endif #ifdef USE_CUDA +#if !defined(ORT_UNIT_TEST_HAS_CUDA_PLUGIN_EP) || !defined(ORT_UNIT_TEST_ENABLE_DYNAMIC_PLUGIN_EP_USAGE) +#include "core/providers/cuda/cuda_execution_provider.h" +#include "core/providers/cuda/cuda_execution_provider_info.h" +#endif #include "core/providers/cuda/cuda_provider_options.h" #endif #if defined(USE_WEBGPU) @@ -55,8 +59,6 @@ std::unique_ptr CudaPluginExecutionProviderWithOptions(const AddCudaPluginOption(config_options, "use_tf32", std::to_string(provider_options->use_tf32)); AddCudaPluginOption(config_options, "fuse_conv_bias", std::to_string(provider_options->fuse_conv_bias)); AddCudaPluginOption(config_options, "sdpa_kernel", std::to_string(provider_options->sdpa_kernel)); - AddCudaPluginOption(config_options, "enable_host_pageable_gather", - std::to_string(provider_options->enable_host_pageable_gather)); } return dynamic_plugin_ep_infra::MakeEp(nullptr, &config_options); @@ -203,6 +205,24 @@ std::unique_ptr CudaExecutionProviderWithOptions(const OrtCU #endif } +std::unique_ptr CudaExecutionProviderWithOptions(const ProviderOptions& provider_options) { +#ifdef USE_CUDA +#if defined(ORT_UNIT_TEST_HAS_CUDA_PLUGIN_EP) && defined(ORT_UNIT_TEST_ENABLE_DYNAMIC_PLUGIN_EP_USAGE) + ConfigOptions config_options; + for (const auto& [key, value] : provider_options) { + ORT_THROW_IF_ERROR(config_options.AddConfigEntry(key.c_str(), value.c_str())); + } + return dynamic_plugin_ep_infra::MakeEp(nullptr, &config_options); +#else + return std::make_unique( + CUDAExecutionProviderInfo::FromProviderOptions(provider_options)); +#endif +#else + ORT_UNUSED_PARAMETER(provider_options); + return nullptr; +#endif +} + std::unique_ptr DefaultDnnlExecutionProvider() { #ifdef USE_DNNL OrtDnnlProviderOptions dnnl_options; diff --git a/onnxruntime/test/util/include/default_providers.h b/onnxruntime/test/util/include/default_providers.h index 306397d8745a2..345e9d4de529e 100644 --- a/onnxruntime/test/util/include/default_providers.h +++ b/onnxruntime/test/util/include/default_providers.h @@ -39,6 +39,7 @@ std::unique_ptr DefaultCudaExecutionProvider(); std::unique_ptr DefaultCudaNHWCExecutionProvider(); #endif std::unique_ptr CudaExecutionProviderWithOptions(const OrtCUDAProviderOptionsV2* provider_options); +std::unique_ptr CudaExecutionProviderWithOptions(const ProviderOptions& provider_options); std::unique_ptr DefaultDnnlExecutionProvider(); std::unique_ptr DnnlExecutionProviderWithOptions(const OrtDnnlProviderOptions* provider_options); // std::unique_ptr DefaultTvmExecutionProvider(); From 2651debc90aa16f3c7c8adbc4e437330f7bb6cab Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 16 Sep 2026 06:50:12 +0000 Subject: [PATCH 07/10] Handle disabled prepacking in direct mode Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- .../contrib_ops/cuda/quantization/gather_block_quantized.cc | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc index fca1bd093a0d0..72434e3b12aa6 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc @@ -156,6 +156,8 @@ Status GatherBlockQuantized::PrePack( std::lock_guard lock(device_data_mutex_); if (direct_host_data_) { + ORT_RETURN_IF_NOT(tensor.Location().device.Type() == OrtDevice::CPU, + "Direct host-pageable GatherBlockQuantized requires a CPU-resident initializer."); direct_host_data_ptr_ = tensor.Data(); data_shape_.assign(tensor.Shape().GetDims().begin(), tensor.Shape().GetDims().end()); } else { @@ -233,7 +235,9 @@ Status GatherBlockQuantized::ComputeInternal(OpKernelContext* ctx) const T1* data_ptr = nullptr; if (direct_host_data_) { - data_ptr = direct_host_data_ptr_; + data_ptr = direct_host_data_ptr_ != nullptr ? direct_host_data_ptr_ + : data == nullptr ? nullptr + : data->Data(); } else if (data_is_constant_) { { std::lock_guard lock(device_data_mutex_); From 2aa61aefe822fe3f7b465833c9c97e57f6c169a5 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Wed, 16 Sep 2026 23:46:31 +0000 Subject: [PATCH 08/10] Fix CUDA CI compilation and generated docs Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- docs/ContribOperators.md | 27 ++++++++++++++++--- .../quantization/gather_block_quantized.cc | 1 - .../gather_block_quantized_op_test.cc | 26 +++++++++++++++--- onnxruntime/test/util/default_providers.cc | 22 --------------- .../test/util/include/default_providers.h | 1 - 5 files changed, 46 insertions(+), 31 deletions(-) diff --git a/docs/ContribOperators.md b/docs/ContribOperators.md index 7c146b628a2dc..7659698814171 100644 --- a/docs/ContribOperators.md +++ b/docs/ContribOperators.md @@ -2938,6 +2938,25 @@ This version of the operator has been available since version 1 of the 'com.micr **Cache Format:** The past and present KV cache tensors are expected in a BNSH format: `(batch_size, num_heads, cache_sequence_length, head_size)`, where `cache_sequence_length` is the length of the cached key/value sequences, or the maximum sequence length when past and present buffer sharing is used. + **Windowed KV Cache (`sliding_window_cache` attribute):** + When `sliding_window_cache` is 1, the past/present buffers are window-sized instead of full-length and the operator evicts internally. Let `C` be the cache capacity (dimension 2 of `past_key`, which is also the sequence dimension of `present_key`), `W` be `local_window_size`, and `T` be the absolute number of tokens processed so far by this batch entry, i.e. `seqlens_k[b] + 1`. The scalar `total_sequence_length` input is only the batch maximum of `T`; the layout below is per batch entry, so a ragged batch gets a different resident range per entry. `C` must be at least `W`. + + After a step, rows `[0, L)` of `present_key` and `present_value` hold the `L` most recent positions in increasing position order, so row `i` holds absolute position `T - L + i`. The retained positions are always physically contiguous and start at row 0; the layout never wraps around, so a ring-buffer layout cannot be exposed through these outputs. Rows `[L, C)` are unspecified. The resident count `L` is a function of `T` alone: + + ``` + G = C - W + 1 + L(T) = T if T <= C + L(T) = T - G * ceil((T - C) / G) otherwise + ``` + + Hence `min(T, W) <= L(T) <= min(T, C)`: the whole window stays resident, and eviction reclaims `G` positions at once rather than one position per step, so consumers must not assume that the cache is kept full at `min(T, C)`. + + Because `L` depends only on `T`, the resulting layout is independent of how the tokens were split into steps: a multi-token step of `S` tokens (speculative decoding, chunked prefill) leaves exactly the layout that the same tokens would produce one at a time. Any `S >= 1` is accepted, including `S > C`; a step that would evict positions it still has to read is staged internally, so the capacity does not have to cover the step. When past context is present, the existing operator restriction still applies: `sequence_length > 1` requires `batch_size == 1`. + + An execution provider may accept only part of the `C >= W` range. A configuration with `C < W` (equivalently, `W > C`) is invalid and is rejected with `INVALID_ARGUMENT`. The CUDA implementation requires `C == W`, so there `G` is 1 and `L(T)` is `min(T, C)`; a larger capacity is rejected. The CPU implementation accepts any `C >= W`, and slack above the window amortizes compaction over `G` steps. + + To drop the last `k` tokens, for example after rejecting speculative draft tokens, re-run with the smaller `total_sequence_length` and `seqlens_k` and leave the buffer untouched. That is exact when `L(T - k) == L(T) - k`, which callers can evaluate with the formula above. Otherwise the shorter layout needs positions that have already been evicted, and the window has to be re-materialized. + **Quantization:** When quantization is enabled, `past_key` and `past_value` inputs can be of type `float8e4m3fn`, `uint8` or `int8`. The corresponding `k_scale` and `v_scale` tensors must be provided. The operator will output `present_key` and `present_value` in same format as the `past_key` and `past_value`. @@ -2981,7 +3000,7 @@ This version of the operator has been available since version 1 of the 'com.micr
scale : float
Custom scale will be used if specified. Default value is 1/sqrt(head_size)
sliding_window_cache : int
-
Set to 1 when the past/present KV buffers are window-sized instead of holding the whole sequence. The op then keeps only the min(total_sequence_length, cache_capacity) most recent tokens, contiguously, using cache-relative indexing and evicting from the front as needed. Requires local_window_size > 0 and a cache capacity of at least local_window_size. Multi-token steps may use a temporary staging buffer, so the capacity need not cover the entire step. Default value is 0 (full-length cache).
+
Set to 1 when the past/present KV buffers are window-sized instead of holding the whole sequence. The op then evicts internally and indexes the buffers in cache-relative coordinates, keeping the most recent positions contiguously at rows [0, L) with min(T, local_window_size) <= L <= min(T, capacity), where T is seqlens_k[b] + 1 for that batch entry. Requires local_window_size > 0 and a cache capacity of at least local_window_size; a smaller capacity (W > C) is rejected with INVALID_ARGUMENT. The CUDA implementation additionally requires the capacity to equal local_window_size. Multi-token steps of any length are supported and produce the same layout as single-token steps, so the capacity need not cover the entire step. When past context is present, sequence_length > 1 requires batch_size == 1. See the Windowed KV Cache section of the operator description for the exact resident-range, eviction and rollback contract. Default value is 0 (full-length cache).
smooth_softmax : int
Use a smooth factor in softmax.
softcap : float
@@ -3000,9 +3019,9 @@ This version of the operator has been available since version 1 of the 'com.micr
value (optional) : T
Value with shape (batch_size, kv_sequence_length, kv_hidden_size)
past_key (optional) : T_CACHE
-
past state key with support for format BNSH. When past_key uses same tensor as present_key(k-v cache), it is of length max_sequence_length... otherwise of length past_sequence_length.
+
past state key with support for format BNSH. When past_key uses same tensor as present_key(k-v cache), it is of length max_sequence_length... otherwise of length past_sequence_length. When sliding_window_cache is 1 this length is the window cache capacity C, which is chosen by the caller independently of the sequence length and must be at least local_window_size.
past_value (optional) : T_CACHE
-
past state value with support for format BNSH. When past_value uses same tensor as present_value(k-v cache), it is of length max_sequence_length... otherwise of length past_sequence_length.
+
past state value with support for format BNSH. When past_value uses same tensor as present_value(k-v cache), it is of length max_sequence_length... otherwise of length past_sequence_length. When sliding_window_cache is 1 this length is the window cache capacity C, which is chosen by the caller independently of the sequence length and must be at least local_window_size.
seqlens_k : M
1D Tensor of shape (batch_size). Equivalent to (total_sequence_lengths - 1).
total_sequence_length : M
@@ -3014,7 +3033,7 @@ This version of the operator has been available since version 1 of the 'com.micr
position_ids (optional) : tensor(int64)
2D tensor with shape (batch_size, sequence_length). When processing the first prompt the kernel uses only the first element
attention_bias (optional) : T
-
additional add to QxK' with shape (batch_size or 1, num_heads or 1, sequence_length, total_sequence_length)
+
additional add to QxK' with shape (batch_size or 1, num_heads or 1, sequence_length, total_sequence_length). The last dimension is indexed by absolute key position and stays total_sequence_length when sliding_window_cache is 1: it is not reduced to the cache capacity or to local_window_size. The operator reads the columns of the positions that are resident in the cache and ignores the rest. CPU supports this windowed absolute-column indexing; CUDA rejects attention_bias when sliding_window_cache is 1.
head_sink (optional) : T
1D tensor with shape (num_heads). Each head has a smooth factor adding to the denominator of softmax.
k_scale (optional) : T_KV_SCALE
diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc index 72434e3b12aa6..ab01a1701f979 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc @@ -3,7 +3,6 @@ #include -#include "core/common/logging/logging.h" #include "core/providers/cuda/cuda_common.h" #include "contrib_ops/cuda/quantization/gather_block_quantized.h" #include "contrib_ops/cuda/quantization/gather_block_quantized.cuh" diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index 0aebbd8f5086d..e253d47ea7abd 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -27,11 +27,15 @@ #ifdef USE_CUDA #include "contrib_ops/cuda/quantization/gather_block_quantized.h" #include "core/graph/model.h" +#include "core/graph/node_attr_utils.h" #include "core/platform/env.h" +#include "core/session/IOBinding.h" #include "core/session/inference_session.h" #include "core/session/onnxruntime_session_options_config_keys.h" +#include "test/unittest_util/test_dynamic_plugin_ep.h" #include "test/util/include/temp_dir.h" #ifndef BUILD_CUDA_EP_AS_PLUGIN +#include "core/providers/cuda/cuda_execution_provider.h" #include "core/providers/cuda/cuda_execution_provider_info.h" #endif #endif @@ -39,6 +43,22 @@ namespace onnxruntime { namespace test { +#ifdef USE_CUDA +std::unique_ptr CudaExecutionProviderWithStringOptions( + const ProviderOptions& provider_options) { +#ifdef BUILD_CUDA_EP_AS_PLUGIN + ConfigOptions config_options; + for (const auto& [key, value] : provider_options) { + ORT_THROW_IF_ERROR(config_options.AddConfigEntry(key.c_str(), value.c_str())); + } + return dynamic_plugin_ep_infra::MakeEp(nullptr, &config_options); +#else + return std::make_unique( + CUDAExecutionProviderInfo::FromProviderOptions(provider_options)); +#endif +} +#endif + // When uint8_t data type is used GatherBlockQuantize applies MatMulNBit's conventions for storing the data. // That is when no zero points are specified a default zero point of 8 (for 4 bits) or 128 (for 8 bits) is used. // This convertor hence compensates for that by adding it to the data values, so that the outputs match the results @@ -1433,7 +1453,7 @@ TEST(GatherBlockQuantizedOpTest, FpFallbackWithPrepackingDisabledCuda) { GTEST_SKIP() << "CUDA not available"; } - auto cuda_ep = CudaExecutionProviderWithOptions( + auto cuda_ep = CudaExecutionProviderWithStringOptions( ProviderOptions{{"enable_host_pageable_gather", "0"}}); if (cuda_ep == nullptr) { GTEST_SKIP() << "CUDA EP not available"; @@ -1481,7 +1501,7 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCuda) { GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; } - auto cuda_ep = CudaExecutionProviderWithOptions( + auto cuda_ep = CudaExecutionProviderWithStringOptions( ProviderOptions{{"enable_host_pageable_gather", "1"}}); if (cuda_ep == nullptr) { GTEST_SKIP() << "CUDA EP not available"; @@ -1527,7 +1547,7 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCudaGraph) { GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; } - auto cuda_ep = CudaExecutionProviderWithOptions( + auto cuda_ep = CudaExecutionProviderWithStringOptions( ProviderOptions{{"enable_cuda_graph", "1"}, {"enable_host_pageable_gather", "1"}}); if (cuda_ep == nullptr) { diff --git a/onnxruntime/test/util/default_providers.cc b/onnxruntime/test/util/default_providers.cc index ef9b71720316b..26a65a44c74c5 100644 --- a/onnxruntime/test/util/default_providers.cc +++ b/onnxruntime/test/util/default_providers.cc @@ -13,10 +13,6 @@ #include "core/providers/coreml/coreml_provider_factory.h" #endif #ifdef USE_CUDA -#if !defined(ORT_UNIT_TEST_HAS_CUDA_PLUGIN_EP) || !defined(ORT_UNIT_TEST_ENABLE_DYNAMIC_PLUGIN_EP_USAGE) -#include "core/providers/cuda/cuda_execution_provider.h" -#include "core/providers/cuda/cuda_execution_provider_info.h" -#endif #include "core/providers/cuda/cuda_provider_options.h" #endif #if defined(USE_WEBGPU) @@ -205,24 +201,6 @@ std::unique_ptr CudaExecutionProviderWithOptions(const OrtCU #endif } -std::unique_ptr CudaExecutionProviderWithOptions(const ProviderOptions& provider_options) { -#ifdef USE_CUDA -#if defined(ORT_UNIT_TEST_HAS_CUDA_PLUGIN_EP) && defined(ORT_UNIT_TEST_ENABLE_DYNAMIC_PLUGIN_EP_USAGE) - ConfigOptions config_options; - for (const auto& [key, value] : provider_options) { - ORT_THROW_IF_ERROR(config_options.AddConfigEntry(key.c_str(), value.c_str())); - } - return dynamic_plugin_ep_infra::MakeEp(nullptr, &config_options); -#else - return std::make_unique( - CUDAExecutionProviderInfo::FromProviderOptions(provider_options)); -#endif -#else - ORT_UNUSED_PARAMETER(provider_options); - return nullptr; -#endif -} - std::unique_ptr DefaultDnnlExecutionProvider() { #ifdef USE_DNNL OrtDnnlProviderOptions dnnl_options; diff --git a/onnxruntime/test/util/include/default_providers.h b/onnxruntime/test/util/include/default_providers.h index 345e9d4de529e..306397d8745a2 100644 --- a/onnxruntime/test/util/include/default_providers.h +++ b/onnxruntime/test/util/include/default_providers.h @@ -39,7 +39,6 @@ std::unique_ptr DefaultCudaExecutionProvider(); std::unique_ptr DefaultCudaNHWCExecutionProvider(); #endif std::unique_ptr CudaExecutionProviderWithOptions(const OrtCUDAProviderOptionsV2* provider_options); -std::unique_ptr CudaExecutionProviderWithOptions(const ProviderOptions& provider_options); std::unique_ptr DefaultDnnlExecutionProvider(); std::unique_ptr DnnlExecutionProviderWithOptions(const OrtDnnlProviderOptions* provider_options); // std::unique_ptr DefaultTvmExecutionProvider(); From 492ce32496d53bcabf34d1cca2a1e254741658d7 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Thu, 17 Sep 2026 01:28:53 +0000 Subject: [PATCH 09/10] Fix CUDA provider test option construction Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- .../quantization/gather_block_quantized.h | 15 +--------- .../gather_block_quantized_data_policy.h | 22 ++++++++++++++ .../providers/cuda/cuda_provider_factory.cc | 6 ++++ .../providers/cuda/cuda_provider_factory.h | 3 ++ .../cuda/cuda_provider_factory_creator.h | 2 ++ .../core/session/provider_bridge_ort.cc | 8 +++++ .../gather_block_quantized_op_test.cc | 29 +++---------------- onnxruntime/test/util/default_providers.cc | 19 ++++++++++++ .../test/util/include/default_providers.h | 1 + 9 files changed, 66 insertions(+), 39 deletions(-) create mode 100644 onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized_data_policy.h diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h index 60a3f3ffaba42..b7a5a84c9c1f1 100644 --- a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.h @@ -3,6 +3,7 @@ #pragma once +#include "contrib_ops/cuda/quantization/gather_block_quantized_data_policy.h" #include "core/providers/cuda/cuda_kernel.h" #include @@ -15,20 +16,6 @@ namespace cuda { using namespace onnxruntime::cuda; -enum class GatherBlockQuantizedDataPolicy { - DeviceCopy, - DirectHost, -}; - -constexpr GatherBlockQuantizedDataPolicy SelectGatherBlockQuantizedDataPolicy( - bool option_enabled, bool pageable_memory_access, bool uses_host_page_tables, - bool is_fp8, bool is_constant_initializer) { - return option_enabled && pageable_memory_access && uses_host_page_tables && - is_fp8 && is_constant_initializer - ? GatherBlockQuantizedDataPolicy::DirectHost - : GatherBlockQuantizedDataPolicy::DeviceCopy; -} - template class GatherBlockQuantized final : public CudaKernel { public: diff --git a/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized_data_policy.h b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized_data_policy.h new file mode 100644 index 0000000000000..a48f5e2ed3576 --- /dev/null +++ b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized_data_policy.h @@ -0,0 +1,22 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +namespace onnxruntime::contrib::cuda { + +enum class GatherBlockQuantizedDataPolicy { + DeviceCopy, + DirectHost, +}; + +constexpr GatherBlockQuantizedDataPolicy SelectGatherBlockQuantizedDataPolicy( + bool option_enabled, bool pageable_memory_access, bool uses_host_page_tables, + bool is_fp8, bool is_constant_initializer) { + return option_enabled && pageable_memory_access && uses_host_page_tables && + is_fp8 && is_constant_initializer + ? GatherBlockQuantizedDataPolicy::DirectHost + : GatherBlockQuantizedDataPolicy::DeviceCopy; +} + +} // namespace onnxruntime::contrib::cuda diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc index 27c7c9ff078d4..d622091500a0b 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc @@ -180,6 +180,12 @@ struct ProviderInfo_CUDA_Impl final : ProviderInfo_CUDA { return std::make_shared(info); } + std::shared_ptr CreateExecutionProviderFactory( + const ProviderOptions& provider_options) override { + return std::make_shared( + CUDAExecutionProviderInfo::FromProviderOptions(provider_options)); + } + std::shared_ptr CreateCudaAllocator(int16_t device_id, size_t gpu_mem_limit, onnxruntime::ArenaExtendStrategy arena_extend_strategy, onnxruntime::CUDAExecutionProviderExternalAllocatorInfo& external_allocator_info, const OrtArenaCfg* default_memory_arena_cfg) override { CUDAExecutionProvider::CUDAAllocatorParams params{}; params.device_id = device_id; diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.h b/onnxruntime/core/providers/cuda/cuda_provider_factory.h index 1a4b19cb100d3..8e3fc1745c448 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.h +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.h @@ -61,6 +61,9 @@ struct ProviderInfo_CUDA { ORT_NOT_IMPLEMENTED(__FUNCTION__, " is only implements in test code path."); } + virtual std::shared_ptr CreateExecutionProviderFactory( + const ProviderOptions& provider_options) = 0; + protected: ~ProviderInfo_CUDA() = default; // Can only be destroyed through a subclass instance }; diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory_creator.h b/onnxruntime/core/providers/cuda/cuda_provider_factory_creator.h index 850fd8af151b6..23a4197f90b16 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory_creator.h +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory_creator.h @@ -5,6 +5,7 @@ #include +#include "core/framework/provider_options.h" #include "core/providers/providers.h" struct OrtCUDAProviderOptions; @@ -15,5 +16,6 @@ namespace onnxruntime { struct CudaProviderFactoryCreator { static std::shared_ptr Create(const OrtCUDAProviderOptions* provider_options); static std::shared_ptr Create(const OrtCUDAProviderOptionsV2* provider_options); + static std::shared_ptr Create(const ProviderOptions& provider_options); }; } // namespace onnxruntime diff --git a/onnxruntime/core/session/provider_bridge_ort.cc b/onnxruntime/core/session/provider_bridge_ort.cc index c45d21639c008..7b3c99787263e 100644 --- a/onnxruntime/core/session/provider_bridge_ort.cc +++ b/onnxruntime/core/session/provider_bridge_ort.cc @@ -2147,6 +2147,14 @@ std::shared_ptr CudaProviderFactoryCreator::Create( return nullptr; } +std::shared_ptr CudaProviderFactoryCreator::Create( + const ProviderOptions& provider_options) try { + return GetProviderInfo_CUDA().CreateExecutionProviderFactory(provider_options); +} catch (const std::exception& exception) { + LOGS_DEFAULT(ERROR) << exception.what(); + return nullptr; +} + std::shared_ptr CannProviderFactoryCreator::Create(const OrtCANNProviderOptions* provider_options) { return s_library_cann.Get().CreateExecutionProviderFactory(provider_options); diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index e253d47ea7abd..632d43efd82ff 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -25,40 +25,19 @@ #include "test/util/include/default_providers.h" #ifdef USE_CUDA -#include "contrib_ops/cuda/quantization/gather_block_quantized.h" +#include "contrib_ops/cuda/quantization/gather_block_quantized_data_policy.h" #include "core/graph/model.h" #include "core/graph/node_attr_utils.h" #include "core/platform/env.h" #include "core/session/IOBinding.h" #include "core/session/inference_session.h" #include "core/session/onnxruntime_session_options_config_keys.h" -#include "test/unittest_util/test_dynamic_plugin_ep.h" #include "test/util/include/temp_dir.h" -#ifndef BUILD_CUDA_EP_AS_PLUGIN -#include "core/providers/cuda/cuda_execution_provider.h" -#include "core/providers/cuda/cuda_execution_provider_info.h" -#endif #endif namespace onnxruntime { namespace test { -#ifdef USE_CUDA -std::unique_ptr CudaExecutionProviderWithStringOptions( - const ProviderOptions& provider_options) { -#ifdef BUILD_CUDA_EP_AS_PLUGIN - ConfigOptions config_options; - for (const auto& [key, value] : provider_options) { - ORT_THROW_IF_ERROR(config_options.AddConfigEntry(key.c_str(), value.c_str())); - } - return dynamic_plugin_ep_infra::MakeEp(nullptr, &config_options); -#else - return std::make_unique( - CUDAExecutionProviderInfo::FromProviderOptions(provider_options)); -#endif -} -#endif - // When uint8_t data type is used GatherBlockQuantize applies MatMulNBit's conventions for storing the data. // That is when no zero points are specified a default zero point of 8 (for 4 bits) or 128 (for 8 bits) is used. // This convertor hence compensates for that by adding it to the data values, so that the outputs match the results @@ -1453,7 +1432,7 @@ TEST(GatherBlockQuantizedOpTest, FpFallbackWithPrepackingDisabledCuda) { GTEST_SKIP() << "CUDA not available"; } - auto cuda_ep = CudaExecutionProviderWithStringOptions( + auto cuda_ep = CudaExecutionProviderWithOptions( ProviderOptions{{"enable_host_pageable_gather", "0"}}); if (cuda_ep == nullptr) { GTEST_SKIP() << "CUDA EP not available"; @@ -1501,7 +1480,7 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCuda) { GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; } - auto cuda_ep = CudaExecutionProviderWithStringOptions( + auto cuda_ep = CudaExecutionProviderWithOptions( ProviderOptions{{"enable_host_pageable_gather", "1"}}); if (cuda_ep == nullptr) { GTEST_SKIP() << "CUDA EP not available"; @@ -1547,7 +1526,7 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCudaGraph) { GTEST_SKIP() << "CUDA device does not use host page tables for pageable memory"; } - auto cuda_ep = CudaExecutionProviderWithStringOptions( + auto cuda_ep = CudaExecutionProviderWithOptions( ProviderOptions{{"enable_cuda_graph", "1"}, {"enable_host_pageable_gather", "1"}}); if (cuda_ep == nullptr) { diff --git a/onnxruntime/test/util/default_providers.cc b/onnxruntime/test/util/default_providers.cc index 26a65a44c74c5..751284588f780 100644 --- a/onnxruntime/test/util/default_providers.cc +++ b/onnxruntime/test/util/default_providers.cc @@ -201,6 +201,25 @@ std::unique_ptr CudaExecutionProviderWithOptions(const OrtCU #endif } +std::unique_ptr CudaExecutionProviderWithOptions(const ProviderOptions& provider_options) { +#ifdef USE_CUDA +#if defined(ORT_UNIT_TEST_HAS_CUDA_PLUGIN_EP) && defined(ORT_UNIT_TEST_ENABLE_DYNAMIC_PLUGIN_EP_USAGE) + ConfigOptions config_options; + for (const auto& [key, value] : provider_options) { + ORT_THROW_IF_ERROR(config_options.AddConfigEntry(key.c_str(), value.c_str())); + } + return dynamic_plugin_ep_infra::MakeEp(nullptr, &config_options); +#else + if (auto factory = CudaProviderFactoryCreator::Create(provider_options)) + return factory->CreateProvider(); + return nullptr; +#endif +#else + ORT_UNUSED_PARAMETER(provider_options); + return nullptr; +#endif +} + std::unique_ptr DefaultDnnlExecutionProvider() { #ifdef USE_DNNL OrtDnnlProviderOptions dnnl_options; diff --git a/onnxruntime/test/util/include/default_providers.h b/onnxruntime/test/util/include/default_providers.h index 306397d8745a2..345e9d4de529e 100644 --- a/onnxruntime/test/util/include/default_providers.h +++ b/onnxruntime/test/util/include/default_providers.h @@ -39,6 +39,7 @@ std::unique_ptr DefaultCudaExecutionProvider(); std::unique_ptr DefaultCudaNHWCExecutionProvider(); #endif std::unique_ptr CudaExecutionProviderWithOptions(const OrtCUDAProviderOptionsV2* provider_options); +std::unique_ptr CudaExecutionProviderWithOptions(const ProviderOptions& provider_options); std::unique_ptr DefaultDnnlExecutionProvider(); std::unique_ptr DnnlExecutionProviderWithOptions(const OrtDnnlProviderOptions* provider_options); // std::unique_ptr DefaultTvmExecutionProvider(); From 2dff4726fb1a94edb291c13515fa8113679f9c72 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Fri, 18 Sep 2026 09:59:59 +0000 Subject: [PATCH 10/10] Fix remaining CUDA test compilation Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> --- .../gather_block_quantized_op_test.cc | 31 ++----------------- .../cuda/test_cases/cuda_test_provider.cc | 24 ++++++++++++++ 2 files changed, 27 insertions(+), 28 deletions(-) diff --git a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc index 632d43efd82ff..05c704781ffdb 100644 --- a/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc +++ b/onnxruntime/test/contrib_ops/gather_block_quantized_op_test.cc @@ -1403,29 +1403,6 @@ TEST(GatherBlockQuantizedOpTest, HostPageablePolicySelection) { GatherBlockQuantizedDataPolicy::DeviceCopy); } -#ifndef BUILD_CUDA_EP_AS_PLUGIN -TEST(GatherBlockQuantizedOpTest, HostPageableProviderOptionRoundTripAndHash) { - CUDAExecutionProviderInfo default_info = - CUDAExecutionProviderInfo::FromProviderOptions({}); - EXPECT_FALSE(default_info.enable_host_pageable_gather); - - CUDAExecutionProviderInfo disabled_info = - CUDAExecutionProviderInfo::FromProviderOptions({{"enable_host_pageable_gather", "0"}}); - EXPECT_FALSE(disabled_info.enable_host_pageable_gather); - - CUDAExecutionProviderInfo enabled_info = - CUDAExecutionProviderInfo::FromProviderOptions({{"enable_host_pageable_gather", "1"}}); - EXPECT_TRUE(enabled_info.enable_host_pageable_gather); - const ProviderOptions serialized = CUDAExecutionProviderInfo::ToProviderOptions(enabled_info); - ASSERT_EQ(serialized.count("enable_host_pageable_gather"), 1u); - EXPECT_EQ(serialized.at("enable_host_pageable_gather"), "1"); - EXPECT_TRUE(CUDAExecutionProviderInfo::FromProviderOptions(serialized).enable_host_pageable_gather); - EXPECT_NE(std::hash{}(disabled_info), - std::hash{}(enabled_info)); -} - -#endif - #if !defined(DISABLE_FLOAT8_TYPES) TEST(GatherBlockQuantizedOpTest, FpFallbackWithPrepackingDisabledCuda) { if (!HasCudaEnvironment(0)) { @@ -1665,18 +1642,16 @@ TEST(GatherBlockQuantizedOpTest, FpDirectHostPageableCudaGraph) { ASSERT_TRUE(cuda_ep_ptr->IsGraphCaptured(1)); auto verify_output = [&](std::initializer_list expected) { - ASSERT_EQ(cudaSuccess, cudaDeviceSynchronize()); std::vector actual(expected.size()); Tensor cpu_output(DataTypeImpl::GetType(), TensorShape({2, 2}), actual.data(), cpu_memory_info); ASSERT_STATUS_OK(cuda_ep_ptr->GetDataTransfer()->CopyTensor(output_value.Get(), cpu_output)); EXPECT_EQ(actual, std::vector(expected.begin(), expected.end())); }; verify_output({1.0f, 2.0f, 10.0f, 12.0f}); - indices = {3, 1}; - ASSERT_EQ(cudaSuccess, - cudaMemcpy(indices_value.GetMutable()->MutableData(), indices.data(), - indices.size() * sizeof(indices[0]), cudaMemcpyHostToDevice)); + indices = {3, 1}; + Tensor cpu_indices(DataTypeImpl::GetType(), TensorShape({2}), indices.data(), cpu_memory_info); + ASSERT_STATUS_OK(cuda_ep_ptr->GetDataTransfer()->CopyTensor(cpu_indices, *indices_value.GetMutable())); ASSERT_STATUS_OK(session.Run(run_options, *io_binding)); verify_output({1.75f, 2.0f, 1.5f, 2.0f}); diff --git a/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc b/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc index 01c7573b9de14..1c7d060e006d1 100644 --- a/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc +++ b/onnxruntime/test/providers/cuda/test_cases/cuda_test_provider.cc @@ -34,6 +34,26 @@ namespace onnxruntime { void InitializeRegistry(); void DeleteRegistry(); +TEST(CUDAProviderOptionsTest, HostPageableGatherRoundTripAndHash) { + const CUDAExecutionProviderInfo default_info = + CUDAExecutionProviderInfo::FromProviderOptions({}); + EXPECT_FALSE(default_info.enable_host_pageable_gather); + + const CUDAExecutionProviderInfo disabled_info = + CUDAExecutionProviderInfo::FromProviderOptions({{"enable_host_pageable_gather", "0"}}); + EXPECT_FALSE(disabled_info.enable_host_pageable_gather); + + const CUDAExecutionProviderInfo enabled_info = + CUDAExecutionProviderInfo::FromProviderOptions({{"enable_host_pageable_gather", "1"}}); + EXPECT_TRUE(enabled_info.enable_host_pageable_gather); + const ProviderOptions serialized = CUDAExecutionProviderInfo::ToProviderOptions(enabled_info); + ASSERT_EQ(serialized.count("enable_host_pageable_gather"), 1u); + EXPECT_EQ(serialized.at("enable_host_pageable_gather"), "1"); + EXPECT_TRUE(CUDAExecutionProviderInfo::FromProviderOptions(serialized).enable_host_pageable_gather); + EXPECT_NE(std::hash{}(disabled_info), + std::hash{}(enabled_info)); +} + struct ProviderInfo_CUDA_TestImpl : ProviderInfo_CUDA { OrtStatus* SetCurrentGpuDeviceId(_In_ int) override { return nullptr; @@ -99,6 +119,10 @@ struct ProviderInfo_CUDA_TestImpl : ProviderInfo_CUDA { return nullptr; } + std::shared_ptr CreateExecutionProviderFactory(const ProviderOptions&) override { + return nullptr; + } + std::shared_ptr CreateCudaAllocator(int16_t, size_t, onnxruntime::ArenaExtendStrategy, onnxruntime::CUDAExecutionProviderExternalAllocatorInfo&, const OrtArenaCfg*) override {