Copilot · Sep 16, 2026 · Sep 16, 2026 · Sep 16, 2026 · Sep 16, 2026 · Sep 16, 2026
diff --git a/docs/ContribOperators.md b/docs/ContribOperators.md
@@ -3480,11 +3480,11 @@ This version of the operator has been available since version 1 of the 'com.micr
   The weight tensor B has shape [N, K] with one FP32 scale per `block_size` consecutive K values
   (`b_scale` of shape [N, ceil(K / block_size)]). The scaled weight value is
   `B_scaled[n, k] = fp8_e4m3(B[n, k]) * b_scale[n, k / block_size]`.
-  
+
   When the optional scalar `a_scale` is provided, the activation values used in the multiplication
   are `A_scaled = fp8_e4m3(A / a_scale) * a_scale` (W8A8). Otherwise, A retains its FP16/BF16
   precision (weight-only W8A16).
-  
+
   The operator multiplies the activation by the transpose of B_scaled and adds the optional bias.
   The output has shape [..., N] and the same element type as A.
 
@@ -7742,4 +7742,3 @@ No versioning maintained for experimental ops.
 <dt><tt>T</tt> : tensor(float)</dt>
 <dd>Constrain input and output types to float32 tensors.</dd>
 </dl>
-
diff --git a/docs/cuda_host_pageable_gather.md b/docs/cuda_host_pageable_gather.md
@@ -0,0 +1,27 @@
+# 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`. If either device capability is unavailable, ONNX Runtime emits a
+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,
+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.
+
+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.
+
+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/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc b/onnxruntime/contrib_ops/cuda/quantization/gather_block_quantized.cc
@@ -71,7 +71,8 @@
 #endif  // !defined(DISABLE_FLOAT4_TYPES)
 
 template <typename T1, typename T2, typename Tind>
-GatherBlockQuantized<T1, T2, Tind>::GatherBlockQuantized(const OpKernelInfo& info) : CudaKernel(info) {
+GatherBlockQuantized<T1, T2, Tind>::GatherBlockQuantized(const OpKernelInfo& info)
+    : CudaKernel(info), direct_host_data_(false), data_is_constant_(false) {
   if constexpr (IsFpQuantizedV<T1>) {
     bits_ = 0;  // Not applicable for FP8/FP4 data.
   } else {
@@ -82,6 +83,36 @@
   gather_axis_ = info.GetAttrOrDefault<int64_t>("gather_axis", 0);
   quantize_axis_ = info.GetAttrOrDefault<int64_t>("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, IsFp8QuantizedV<T1>,
+                                           data_is_constant_) ==
+      GatherBlockQuantizedDataPolicy::DirectHost;
+  if (option_enabled && IsFp8QuantizedV<T1> && !direct_host_data_) {
+    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. 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.
   // 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,15 +124,65 @@
   }
 }
 
+template <typename T1, typename T2, typename Tind>
+Status GatherBlockQuantized<T1, T2, Tind>::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.");
+
+  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();
+  }
+
+  auto device_data = IAllocator::MakeUniquePtr<void>(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();
+}
+
+template <typename T1, typename T2, typename Tind>
+Status GatherBlockQuantized<T1, T2, Tind>::PrePack(
+    const Tensor& tensor, int input_idx, AllocatorPtr alloc,
+    bool& is_packed, PrePackedWeights* prepacked_weights) {
+  is_packed = false;
+  if (input_idx != 0) {
+    return Status::OK();
+  }
+
+  if (!direct_host_data_ && tensor.Location().device.Type() != OrtDevice::CPU) {
+    return Status::OK();
+  }
+
+  std::lock_guard<std::mutex> 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<T1>();
+    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;
+  }
+  return Status::OK();
+}
+
 template <typename T1, typename T2, typename Tind>
 Status GatherBlockQuantized<T1, T2, Tind>::ComputeInternal(OpKernelContext* ctx) const {
   const Tensor* data = ctx->Input<Tensor>(0);
   const Tensor* indices = ctx->Input<Tensor>(1);
   const Tensor* scales = ctx->Input<Tensor>(2);
   const Tensor* zero_points = ctx->Input<Tensor>(3);
 
-  auto data_shape = data->Shape().GetDims();
-  int64_t data_rank = data->Shape().NumDimensions();
+  const gsl::span<const int64_t> data_shape =
+      data != nullptr ? data->Shape().GetDims() : gsl::span<const int64_t>{data_shape_};
+  int64_t data_rank = static_cast<int64_t>(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 +236,33 @@
     return Status::OK();
   }
 
-  const auto* data_ptr = data->Data<T1>();
+  const T1* data_ptr = nullptr;
+  if (direct_host_data_) {
+    data_ptr = direct_host_data_ptr_ != nullptr ? direct_host_data_ptr_
+               : data == nullptr                ? nullptr
+                                                : data->Data<T1>();
+  } else if (data_is_constant_) {
+    {
+      std::lock_guard<std::mutex> lock(device_data_mutex_);
+      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(
+            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 = device_data_ != nullptr ? static_cast<const T1*>(device_data_.get())
+                 : data == nullptr       ? nullptr
+                                         : data->Data<T1>();
+    }
+  } else {
+    data_ptr = data->Data<T1>();
+  }
+  ORT_RETURN_IF_NOT(N == 0 || data_ptr != nullptr,
+                    "GatherBlockQuantized fallback has no device-resident input 0.");
   const auto* indices_ptr = indices->Data<Tind>();
   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
@@ -43,6 +43,15 @@ struct IsFpQuantized<Float4E2M1x2> : std::true_type {};
 template <typename T1>
 inline constexpr bool IsFpQuantizedV = IsFpQuantized<T1>::value;
 
+template <typename T1>
+inline constexpr bool IsFp8QuantizedV =
+#if !defined(DISABLE_FLOAT8_TYPES)
+    std::is_same_v<T1, Float8E4M3FN> || std::is_same_v<T1, Float8E4M3FNUZ> ||
+    std::is_same_v<T1, Float8E5M2> || std::is_same_v<T1, Float8E5M2FNUZ>;
+#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
@@ -3,9 +3,10 @@
 
 #pragma once
 
+#include "contrib_ops/cuda/quantization/gather_block_quantized_data_policy.h"
 #include "core/providers/cuda/cuda_kernel.h"
 
-#include <iostream>
+#include <mutex>
 
 using namespace onnxruntime::cuda;
 
@@ -20,12 +21,22 @@
  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_;
+  const T1* direct_host_data_ptr_{};
+  mutable IAllocatorUniquePtr<void> device_data_;
+  mutable TensorShapeVector data_shape_;
 };
 
 }  // namespace cuda

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
@@ -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_execution_provider.h b/onnxruntime/core/providers/cuda/cuda_execution_provider.h
@@ -90,6 +90,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
@@ -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";
 constexpr const char* kExternalDataLoaderReadingThreads = "external_data_loader_reading_threads";
 
 }  // namespace provider_option_names
@@ -132,6 +133,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::kExternalDataLoaderReadingThreads,
@@ -201,6 +204,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)},
       {cuda::provider_option_names::kExternalDataLoaderReadingThreads,
        MakeStringWithClassicLocale(info.external_data_loader_reading_threads)},
   };

diff --git 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};
 
   // 0 disables the custom external-data loader and retains the framework's existing path.
   // 1 uses the pinned-buffer loader with synchronous reads. 2..64 use that many parallel read tasks per block.
@@ -121,6 +122,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);
     onnxruntime::HashCombine(info.external_data_loader_reading_threads, value);
 
     // Memory pointers

diff --git 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
@@ -251,7 +251,6 @@ struct CUDA_Provider : Provider {
     info.use_tf32 = params->use_tf32 != 0;
     info.sdpa_kernel = params->sdpa_kernel;
     info.external_data_loader_reading_threads = params->external_data_loader_reading_threads;
-
     return std::make_shared<CUDAProviderFactory>(info);
   }
 

diff --git 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
@@ -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
@@ -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);