Fix temporary host source lifetime with a separate copy API

.github/workflows/test.yaml

@@ -54,8 +54,6 @@ jobs:
       date: ${{ inputs.date }}
       script: ci/test_cpp.sh
       sha: ${{ inputs.sha }}
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-cpp-benchmark-tests:
     permissions:
       actions: read
@@ -123,8 +121,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: "ci/test_python_cudf.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-python-other-tests:
     # Tests for dask_cudf, custreamz, cudf_kafka are separated for CI parallelism
     permissions:
@@ -141,8 +137,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: "ci/test_python_other.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   conda-java-tests:
     permissions:
       actions: read
@@ -224,8 +218,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: ci/cudf_pandas_scripts/run_tests.sh
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   third-party-integration-tests-cudf-pandas:
     permissions:
       actions: read
@@ -279,8 +271,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: "ci/test_wheel_cudf_polars.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   cudf-polars-polars-tests:
     permissions:
       actions: read
@@ -296,8 +286,6 @@ jobs:
       date: ${{ inputs.date }}
       sha: ${{ inputs.sha }}
       script: "ci/test_cudf_polars_polars_tests.sh"
-      # https://github.com/rapidsai/cudf/issues/23498
-      matrix_filter: map(select(.GPU != "gb300" and .GPU != "gh200"))
   narwhals-tests:
     permissions:
       actions: read

cpp/examples/parquet_inspect/parquet_inspect_utils.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -8,6 +8,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda_memcpy.hpp>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/io/experimental/hybrid_scan.hpp>
 #include <cudf/io/parquet_io_utils.hpp>
 #include <cudf/io/parquet_schema.hpp>
@@ -120,7 +121,8 @@ auto make_index_column(cudf::size_type num_rows, rmm::cuda_stream_view stream)
 {
   std::vector<cudf::size_type> data(num_rows);
   std::iota(data.begin(), data.end(), 0);
-  auto buffer = rmm::device_buffer(data.data(), num_rows * sizeof(int64_t), stream);
+  auto buffer = cudf::detail::make_device_buffer_async(
+    cudf::host_span<cudf::size_type const>{data}, stream, cudf::get_current_device_resource_ref());
   return std::make_unique<cudf::column>(cudf::data_type{cudf::type_to_id<cudf::size_type>()},
                                         num_rows,
                                         std::move(buffer),
@@ -140,7 +142,8 @@ auto make_index_column(cudf::size_type num_rows, rmm::cuda_stream_view stream)
 template <typename T>
 auto make_column(cudf::host_span<T const> host_data, rmm::cuda_stream_view stream)
 {
-  auto device_buffer = rmm::device_buffer(host_data.data(), host_data.size() * sizeof(T), stream);
+  auto device_buffer = cudf::detail::make_device_buffer_async(
+    host_data, stream, cudf::get_current_device_resource_ref());
   return std::make_unique<cudf::column>(cudf::data_type{cudf::type_to_id<T>()},
                                         host_data.size(),
                                         std::move(device_buffer),
@@ -172,8 +175,8 @@ auto make_page_data_list_column(cudf::host_span<T const> data,
 
   auto offsets_column = make_column<cudf::size_type>(col_page_offsets, stream);
 
-  auto page_data_buffer =
-    rmm::device_buffer(data.data(), num_pages_this_column * sizeof(int64_t), stream);
+  auto page_data_buffer = cudf::detail::make_device_buffer_async(
+    data.subspan(0, num_pages_this_column), stream, cudf::get_current_device_resource_ref());
 
   auto page_data_column =
     std::make_unique<cudf::column>(cudf::data_type{cudf::type_to_id<int64_t>()},
@@ -279,11 +282,17 @@ void write_rowgroup_metadata(cudf::io::parquet::FileMetaData const& metadata,
   columns.emplace_back(make_index_column(num_row_groups, stream));
 
   auto row_offsets_buffer =
-    rmm::device_buffer(row_group_row_offsets.data(), num_row_groups * sizeof(int64_t), stream);
+    cudf::detail::make_device_buffer_async(cudf::host_span<int64_t const>{row_group_row_offsets},
+                                           stream,
+                                           cudf::get_current_device_resource_ref());
   auto row_counts_buffer =
-    rmm::device_buffer(row_group_row_counts.data(), num_row_groups * sizeof(int64_t), stream);
+    cudf::detail::make_device_buffer_async(cudf::host_span<int64_t const>{row_group_row_counts},
+                                           stream,
+                                           cudf::get_current_device_resource_ref());
   auto byte_offsets_buffer =
-    rmm::device_buffer(row_group_byte_offsets.data(), num_row_groups * sizeof(int64_t), stream);
+    cudf::detail::make_device_buffer_async(cudf::host_span<int64_t const>{row_group_byte_offsets},
+                                           stream,
+                                           cudf::get_current_device_resource_ref());
 
   columns.emplace_back(std::make_unique<cudf::column>(cudf::data_type{cudf::type_to_id<int64_t>()},
                                                       num_row_groups,

cpp/include/cudf/detail/utilities/cuda_memcpy.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -19,6 +19,11 @@ enum class host_memory_kind : uint8_t { PINNED, PAGEABLE };
 void cuda_memcpy_async_impl(
   void* dst, void const* src, size_t size, host_memory_kind kind, rmm::cuda_stream_view stream);
 
+[[nodiscard]] cudaError_t memcpy_h2d_async(void* dst,
+                                           void const* src,
+                                           size_t size,
+                                           rmm::cuda_stream_view stream);
+
 /**
  * @brief Wrapper around cudaMemcpyBatchAsync
  *
@@ -70,9 +75,10 @@ void cuda_memcpy_async_impl(
                                        rmm::cuda_stream_view stream);
 
 /**
- * @brief Asynchronously copies data from host to device memory.
+ * @brief Asynchronously copies data from host to device memory
  *
- * Implementation may use different strategies depending on the size and type of host data.
+ * The destination copy remains stream ordered. The host source may be released or changed as soon
+ * as this function returns.
  *
  * @param dst Destination device memory
  * @param src Source host memory
@@ -82,12 +88,7 @@ template <typename T>
 void cuda_memcpy_async(device_span<T> dst, host_span<T const> src, rmm::cuda_stream_view stream)
 {
   CUDF_EXPECTS(dst.size() == src.size(), "Mismatched sizes in cuda_memcpy_async");
-  auto const is_pinned = src.is_device_accessible();
-  cuda_memcpy_async_impl(dst.data(),
-                         src.data(),
-                         src.size_bytes(),
-                         is_pinned ? host_memory_kind::PINNED : host_memory_kind::PAGEABLE,
-                         stream);
+  CUDF_CUDA_TRY(memcpy_h2d_async(dst.data(), src.data(), src.size_bytes(), stream));
 }
 
 /**

cpp/include/cudf/detail/utilities/vector_factories.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -21,6 +21,7 @@
 #include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/device_buffer.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/resource_ref.hpp>
 
@@ -77,7 +78,8 @@ rmm::device_uvector<T> make_zeroed_device_uvector(std::size_t size,
  * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a
  * `host_span`
  *
- * @note This function does not synchronize `stream`.
+ * @note The destination copy remains stream ordered. The host source may be released or changed as
+ * soon as this function returns.
  *
  * @tparam T The type of the data to copy (may be const-qualified)
  * @param source_data The host_span of data to deep copy
@@ -99,7 +101,8 @@ rmm::device_uvector<std::remove_cv_t<T>> make_device_uvector_async(
  * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a host
  * container
  *
- * @note This function does not synchronize `stream`.
+ * @note The destination copy remains stream ordered. The host source may be released or changed as
+ * soon as this function returns.
  *
  * @tparam Container The type of the container to copy from
  * @tparam T The type of the data to copy
@@ -119,7 +122,8 @@ rmm::device_uvector<typename Container::value_type> make_device_uvector_async(
 /**
  * @brief Asynchronously construct a `device_uvector` from a `std::vector`
  *
- * @note This function does not synchronize `stream`.
+ * @note The destination copy remains stream ordered. The host source may be released or changed as
+ * soon as this function returns.
  *
  * @tparam T The type of the data to copy
  * @tparam Allocator The allocator type of the std::vector
@@ -136,6 +140,32 @@ rmm::device_uvector<T> make_device_uvector_async(std::vector<T, Allocator> const
   return make_device_uvector_async(host_span<T const>{source_data}, stream, mr);
 }
 
+/**
+ * @brief Asynchronously construct a `device_buffer` containing a deep copy of host data
+ *
+ * @note The destination copy remains stream ordered. The host source may be released or changed as
+ * soon as this function returns.
+ *
+ * @tparam T The type of the data to copy (may be const-qualified)
+ * @param source_data The host data to deep copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_buffer
+ * @return A device_buffer containing the copied data
+ */
+template <typename T>
+rmm::device_buffer make_device_buffer_async(host_span<T> source_data,
+                                            rmm::cuda_stream_view stream,
+                                            rmm::device_async_resource_ref mr)
+{
+  using value_type = std::remove_cv_t<T>;
+  rmm::device_buffer ret(source_data.size_bytes(), stream, mr);
+  cuda_memcpy_async<value_type>(
+    device_span<value_type>{static_cast<value_type*>(ret.data()), source_data.size()},
+    host_span<value_type const>{source_data},
+    stream);
+  return ret;
+}
+
 /**
  * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a
  * `device_span`

cpp/include/cudf_test/column_wrapper.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -160,8 +160,10 @@ rmm::device_buffer make_elements(InputIterator begin, InputIterator end)
   auto transform_begin = thrust::make_transform_iterator(begin, transformer);
   auto const size      = cudf::distance(begin, end);
   auto const elements  = thrust::host_vector<ElementTo>(transform_begin, transform_begin + size);
-  return rmm::device_buffer{
-    elements.data(), size * sizeof(ElementTo), cudf::test::get_default_stream()};
+  return cudf::detail::make_device_buffer_async(
+    cudf::host_span<ElementTo const>{elements.data(), elements.size()},
+    cudf::test::get_default_stream(),
+    cudf::get_current_device_resource_ref());
 }
 
 // The two signatures below are identical to the above overload apart from
@@ -190,8 +192,10 @@ rmm::device_buffer make_elements(InputIterator begin, InputIterator end)
   auto transform_begin = thrust::make_transform_iterator(begin, transformer);
   auto const size      = cudf::distance(begin, end);
   auto const elements  = thrust::host_vector<RepType>(transform_begin, transform_begin + size);
-  return rmm::device_buffer{
-    elements.data(), size * sizeof(RepType), cudf::test::get_default_stream()};
+  return cudf::detail::make_device_buffer_async(
+    cudf::host_span<RepType const>{elements.data(), elements.size()},
+    cudf::test::get_default_stream(),
+    cudf::get_current_device_resource_ref());
 }
 
 /**
@@ -221,8 +225,10 @@ rmm::device_buffer make_elements(InputIterator begin, InputIterator end)
   auto transformer_begin = thrust::make_transform_iterator(begin, to_rep);
   auto const size        = cudf::distance(begin, end);
   auto const elements = thrust::host_vector<RepType>(transformer_begin, transformer_begin + size);
-  return rmm::device_buffer{
-    elements.data(), size * sizeof(RepType), cudf::test::get_default_stream()};
+  return cudf::detail::make_device_buffer_async(
+    cudf::host_span<RepType const>{elements.data(), elements.size()},
+    cudf::test::get_default_stream(),
+    cudf::get_current_device_resource_ref());
 }
 //! @endcond
 
@@ -277,9 +283,10 @@ std::pair<rmm::device_buffer, cudf::size_type> make_null_mask(ValidityIterator b
                                                               ValidityIterator end)
 {
   auto [null_mask, null_count] = make_null_mask_vector(begin, end);
-  auto d_mask                  = rmm::device_buffer{null_mask.data(),
-                                   cudf::bitmask_allocation_size_bytes(cudf::distance(begin, end)),
-                                   cudf::test::get_default_stream()};
+  auto d_mask                  = cudf::detail::make_device_buffer_async(
+    cudf::host_span<bitmask_type const>{null_mask.data(), null_mask.size()},
+    cudf::test::get_default_stream(),
+    cudf::get_current_device_resource_ref());
   return {std::move(d_mask), null_count};
 }
 
@@ -563,12 +570,14 @@ class fixed_point_column_wrapper : public detail::column_wrapper {
     auto const id        = type_to_id<numeric::fixed_point<Rep, numeric::Radix::BASE_10>>();
     auto const data_type = cudf::data_type{id, static_cast<int32_t>(scale)};
 
-    wrapped.reset(new cudf::column{
-      data_type,
-      size,
-      rmm::device_buffer{elements.data(), size * sizeof(Rep), cudf::test::get_default_stream()},
-      rmm::device_buffer{},
-      0});
+    wrapped.reset(new cudf::column{data_type,
+                                   size,
+                                   cudf::detail::make_device_buffer_async(
+                                     cudf::host_span<Rep const>{elements.data(), elements.size()},
+                                     cudf::test::get_default_stream(),
+                                     cudf::get_current_device_resource_ref()),
+                                   rmm::device_buffer{},
+                                   0});
   }
 
   /**
@@ -629,12 +638,14 @@ class fixed_point_column_wrapper : public detail::column_wrapper {
     auto const id                = type_to_id<numeric::fixed_point<Rep, numeric::Radix::BASE_10>>();
     auto const data_type         = cudf::data_type{id, static_cast<int32_t>(scale)};
     auto [null_mask, null_count] = detail::make_null_mask(v, v + size);
-    wrapped.reset(new cudf::column{
-      data_type,
-      size,
-      rmm::device_buffer{elements.data(), size * sizeof(Rep), cudf::test::get_default_stream()},
-      std::move(null_mask),
-      null_count});
+    wrapped.reset(new cudf::column{data_type,
+                                   size,
+                                   cudf::detail::make_device_buffer_async(
+                                     cudf::host_span<Rep const>{elements.data(), elements.size()},
+                                     cudf::test::get_default_stream(),
+                                     cudf::get_current_device_resource_ref()),
+                                   std::move(null_mask),
+                                   null_count});
   }
 
   /**

cpp/libcudf_streaming/tests/streaming/test_channel_metadata.cpp

@@ -6,6 +6,7 @@
 #include <cudf_test/cudf_gtest.hpp>
 
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/utilities/default_stream.hpp>
 
@@ -166,7 +167,8 @@ class StreamingChannelMetadataGPU : public ::testing::Test {
 
   std::shared_ptr<table_chunk> make_chunk(std::vector<int32_t> vals)
   {
-    rmm::device_buffer buf(vals.data(), vals.size() * sizeof(int32_t), stream);
+    auto buf = cudf::detail::make_device_buffer_async(
+      cudf::host_span<int32_t const>{vals}, stream, cudf::get_current_device_resource_ref());
     auto col = std::make_unique<cudf::column>(cudf::data_type{cudf::type_id::INT32},
                                               static_cast<cudf::size_type>(vals.size()),
                                               std::move(buf),

cpp/src/interop/from_arrow_host.cu

@@ -110,7 +110,8 @@ std::pair<std::unique_ptr<rmm::device_buffer>, size_type> get_mask_buffer(
   auto const copy_size    = cudf::util::div_rounding_up_safe(num_rows + bit_index, bits_in_byte);
 
   auto mask = rmm::device_uvector<bitmask_type>(padded_words, stream, mr);
-  CUDF_CUDA_TRY(cudf::detail::memcpy_async(mask.data(), bitmap + offset_index, copy_size, stream));
+  CUDF_CUDA_TRY(
+    cudf::detail::memcpy_h2d_async(mask.data(), bitmap + offset_index, copy_size, stream));
 
   if (mask_words > 0 && bit_index > 0) {
     auto dest_mask = rmm::device_uvector<bitmask_type>(padded_words, stream, mr);
@@ -158,10 +159,10 @@ struct dispatch_copy_from_arrow_host {
 
     auto col = make_fixed_width_column(type, num_rows, mask_state::UNALLOCATED, stream, mr);
     auto mutable_column_view = col->mutable_view();
-    CUDF_CUDA_TRY(cudf::detail::memcpy_async(mutable_column_view.data<DeviceType>(),
-                                             data_buffer + input->offset,
-                                             sizeof(DeviceType) * num_rows,
-                                             stream));
+    CUDF_CUDA_TRY(cudf::detail::memcpy_h2d_async(mutable_column_view.data<DeviceType>(),
+                                                 data_buffer + input->offset,
+                                                 sizeof(DeviceType) * num_rows,
+                                                 stream));
 
     if (!skip_mask) {
       auto [mask, null_count] = get_mask_buffer(input, stream, mr);
@@ -190,7 +191,7 @@ std::unique_ptr<column> dispatch_copy_from_arrow_host::operator()<bool>(ArrowSch
 
   auto data = rmm::device_uvector<bitmask_type>(data_words, stream, mr);
   CUDF_CUDA_TRY(
-    cudf::detail::memcpy_async(data.data(), data_buffer + offset_index, copy_size, stream));
+    cudf::detail::memcpy_h2d_async(data.data(), data_buffer + offset_index, copy_size, stream));
 
   if (data_words > 0 && bit_index > 0) {
     auto dest_data = rmm::device_uvector<bitmask_type>(data_words, stream, mr);