@@ -17,9 +17,9 @@
 #include <cudf/utilities/concepts.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 
-#include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/buffer>
 #include <cuda/functional>
 #include <cuda/iterator>
 #include <cuda/std/concepts>
@@ -234,7 +234,8 @@ struct null_considering_binop {
                  "Output column type should match input column type");
 
     // Shallow copy of the resultant strings
-    rmm::device_uvector<cudf::string_view> out_col_strings(col_size, stream);
+    cuda::device_buffer<cudf::string_view> out_col_strings(
+      stream, cudf::get_current_device_resource_ref(), col_size, cuda::no_init);
 
     // Invalid output column strings - null rows
     cudf::string_view const invalid_str{nullptr, 0};
 
@@ -7,6 +7,7 @@
 #include <cudf/detail/null_mask.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/utilities/buffer_factories.hpp>
 #include <cudf/detail/utilities/grid_1d.cuh>
 #include <cudf/detail/utilities/integer_utils.hpp>
 #include <cudf/detail/utilities/vector_factories.hpp>
@@ -25,6 +26,7 @@
 #include <cub/block/block_reduce.cuh>
 #include <cuda/atomic>
 #include <cuda/bit>
+#include <cuda/buffer>
 #include <cuda/memory_resource>
 #include <cuda/numeric>
 #include <cuda/std/execution>
@@ -500,7 +502,7 @@ std::vector<size_type> batch_count_set_bits(host_span<bitmask_type const* const>
   auto const d_bitmasks = cudf::detail::make_device_uvector_async(bitmasks, stream, tmp_mr);
   auto const num_words  = num_bitmask_words(num_bits_to_count);
   auto d_non_zero_count =
-    cudf::detail::make_zeroed_device_uvector_async<size_type>(num_bitmasks, stream, tmp_mr);
+    cudf::detail::make_zeroed_device_buffer_async<size_type>(num_bitmasks, stream, tmp_mr);
 
   constexpr size_type block_size{256};
   // We use a 2D grid to launch the kernel, where the first dimension is to access elements in each
 
@@ -8,7 +8,9 @@
 #include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
+#include <cudf/utilities/memory_resource.hpp>
 
+#include <cuda/buffer>
 #include <cuda/stream>
 
 #include <functional>
@@ -75,19 +77,26 @@ create_device_view_from_view(ColumnView const& source,
   // Each ColumnDeviceView instance may have child objects that
   // require setting some internal device pointers before being copied
   // from CPU to device.
-  auto const descendant_storage =
-    new rmm::device_uvector<char>(descendant_storage_bytes, stream, mr);
-
-  auto deleter = [descendant_storage](ColumnDeviceView* v) {
+  auto descendant_storage = std::make_unique<cuda::device_buffer<char>>(
+    stream,
+    mr,
+    descendant_storage_bytes,
+    cuda::no_init,
+    cuda::std::execution::prop{cuda::allocation_alignment, alignof(ColumnDeviceView)});
+  cudf::device_span<char> descendant_storage_data = *descendant_storage;
+
+  auto tmp_result = std::make_unique<ColumnDeviceView>(
+    source, staging_buffer.data(), descendant_storage_data.data());
+
+  auto deleter = [descendant_storage = descendant_storage.release()](ColumnDeviceView* v) {
     v->destroy();
     delete descendant_storage;
   };
 
-  std::unique_ptr<ColumnDeviceView, decltype(deleter)> result{
-    new ColumnDeviceView(source, staging_buffer.data(), descendant_storage->data()), deleter};
+  std::unique_ptr<ColumnDeviceView, decltype(deleter)> result{tmp_result.release(), deleter};
 
   // copy the CPU memory with all the children into device memory
-  detail::cuda_memcpy<char>(*descendant_storage, staging_buffer, stream);
+  detail::cuda_memcpy<char>(descendant_storage_data, staging_buffer, stream);
   return result;
 }
 
 
@@ -14,6 +14,7 @@
 #include <cudf/strings/detail/strings_column_factories.cuh>
 #include <cudf/utilities/memory_resource.hpp>
 
+#include <cuda/buffer>
 #include <cuda/iterator>
 #include <thrust/uninitialized_fill.h>
 
@@ -67,7 +68,8 @@ std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::stri
   auto const d_str = ss.value(stream);  // no actual data is copied
 
   // fill the column with the scalar
-  rmm::device_uvector<cudf::strings::detail::string_index_pair> indices(size, stream);
+  cuda::device_buffer<cudf::strings::detail::string_index_pair> indices(
+    stream, cudf::get_current_device_resource_ref(), size, cuda::no_init);
   auto const row_value =
     d_str.empty() ? cudf::strings::detail::string_index_pair{"", 0}
                   : cudf::strings::detail::string_index_pair{d_str.data(), d_str.size_bytes()};
 
@@ -9,7 +9,9 @@
 #include <cudf/column/scalar_column_view.hpp>
 #include <cudf/detail/utilities/cuda_memcpy.hpp>
 #include <cudf/detail/utilities/vector_factories.hpp>
+#include <cudf/utilities/memory_resource.hpp>
 
+#include <cuda/buffer>
 #include <cuda/iterator>
 
 #include <jit/cache.hpp>
@@ -42,19 +44,9 @@ std::map<uint32_t, std::string> build_ptx_params(std::span<std::string const> ou
                                                  std::span<std::string const> input_typenames,
                                                  bool has_user_data);
 
-template <typename T>
-rmm::device_uvector<T> to_device_vector(std::vector<T> const& host,
-                                        cuda::stream_ref stream,
-                                        rmm::device_async_resource_ref mr)
-{
-  rmm::device_uvector<T> device{host.size(), stream, mr};
-  cudf::detail::cuda_memcpy_async<T>(device, host, stream);
-  return device;
-}
-
 template <typename DeviceView, typename ColumnView>
 std::tuple<std::vector<std::unique_ptr<DeviceView, std::function<void(DeviceView*)>>>,
-           rmm::device_uvector<DeviceView>>
+           cuda::device_buffer<DeviceView>>
 column_views_to_device(std::span<ColumnView const> views,
                        cuda::stream_ref stream,
                        rmm::device_async_resource_ref mr)
@@ -72,7 +64,7 @@ column_views_to_device(std::span<ColumnView const> views,
     host_array.push_back(*h);
   }
 
-  rmm::device_uvector<DeviceView> device_array{handles.size(), stream, mr};
+  cuda::device_buffer<DeviceView> device_array{stream, mr, handles.size(), cuda::no_init};
   cudf::detail::cuda_memcpy<DeviceView>(device_array, host_array, stream);
 
   return std::make_tuple(std::move(handles), std::move(device_array));
 
@@ -10,6 +10,7 @@
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/row_operator/hashing.cuh>
 #include <cudf/detail/scatter.hpp>
+#include <cudf/detail/utilities/buffer_factories.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/grid_1d.cuh>
 #include <cudf/detail/utilities/integer_utils.hpp>
@@ -368,7 +369,7 @@ void copy_block_partitions_impl(InputIter const input,
   CUDF_CUDA_TRY(cudaGetLastError());
 }
 
-rmm::device_uvector<size_type> compute_gather_map(size_type num_rows,
+cuda::device_buffer<size_type> compute_gather_map(size_type num_rows,
                                                   size_type num_partitions,
                                                   size_type const* row_partition_numbers,
                                                   size_type const* row_partition_offset,
@@ -378,10 +379,11 @@ rmm::device_uvector<size_type> compute_gather_map(size_type num_rows,
                                                   cuda::stream_ref stream)
 {
   auto sequence = cuda::counting_iterator<cudf::size_type>{0};
-  rmm::device_uvector<size_type> gather_map(num_rows, stream);
+  cuda::device_buffer<size_type> gather_map(
+    stream, cudf::get_current_device_resource_ref(), num_rows, cuda::no_init);
 
   copy_block_partitions_impl(sequence,
-                             gather_map.begin(),
+                             gather_map.data(),
                              num_rows,
                              num_partitions,
                              row_partition_numbers,
@@ -480,7 +482,8 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table_g
   CUDF_EXPECTS(num_partitions < std::numeric_limits<size_type>::max(),
                "num_partitions exceeds cudf's supported limit");
 
-  auto row_partition_numbers = rmm::device_uvector<size_type>(num_rows, stream);
+  cuda::device_buffer<size_type> row_partition_numbers(
+    stream, cudf::get_current_device_resource_ref(), num_rows, cuda::no_init);
 
   // Compute partition number for each row
   if (is_power_two(num_partitions)) {
@@ -506,7 +509,7 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table_g
   // Build histogram via cub::DeviceHistogram::HistogramEven.
   // HistogramEven writes num_partitions bins; the extra element is used by the exclusive scan
   // below to produce the total row count as the last offset. Zero-initialize to avoid UB.
-  auto histogram = cudf::detail::make_zeroed_device_uvector_async<size_type>(
+  auto histogram = cudf::detail::make_zeroed_device_buffer_async<size_type>(
     num_partitions + 1, stream, cudf::get_current_device_resource_ref());
   {
     auto const num_levels  = num_partitions + 1;
@@ -547,7 +550,8 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table_g
   auto const pinned_offsets = cudf::detail::make_pinned_vector_async(histogram, stream);
 
   // Build scatter map: atomically increment partition offsets
-  rmm::device_uvector<size_type> scatter_map(num_rows, stream);
+  cuda::device_buffer<size_type> scatter_map(
+    stream, cudf::get_current_device_resource_ref(), num_rows, cuda::no_init);
   thrust::transform(
     rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
     row_partition_numbers.begin(),
@@ -606,7 +610,7 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
   std::size_t const grid_size = util::div_rounding_up_safe(num_rows, rows_per_block);
 
   // Allocate array to hold which partition each row belongs to
-  auto row_partition_numbers = rmm::device_uvector<size_type>(num_rows, stream, temp_mr);
+  cuda::device_buffer<size_type> row_partition_numbers(stream, temp_mr, num_rows, cuda::no_init);
 
   // Array to hold the size of each partition computed by each block
   //  i.e., { {block0 partition0 size, block1 partition0 size, ...},
@@ -615,17 +619,17 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
   //          {block0 partition(num_partitions-1) size, block1
   //          partition(num_partitions -1) size, ...} }
   auto block_partition_sizes =
-    rmm::device_uvector<size_type>(grid_size * num_partitions, stream, temp_mr);
+    cuda::device_buffer<size_type>(stream, temp_mr, grid_size * num_partitions, cuda::no_init);
 
   auto scanned_block_partition_sizes =
-    rmm::device_uvector<size_type>(grid_size * num_partitions, stream, temp_mr);
+    cuda::device_buffer<size_type>(stream, temp_mr, grid_size * num_partitions, cuda::no_init);
 
   // Holds the total number of rows in each partition
   auto global_partition_sizes =
-    cudf::detail::make_zeroed_device_uvector_async<size_type>(num_partitions, stream, temp_mr);
+    cudf::detail::make_zeroed_device_buffer_async<size_type>(num_partitions, stream, temp_mr);
 
   auto row_partition_offset =
-    cudf::detail::make_zeroed_device_uvector_async<size_type>(num_rows, stream, temp_mr);
+    cudf::detail::make_zeroed_device_buffer_async<size_type>(num_rows, stream, temp_mr);
 
   // If the number of partitions is a power of two, we can compute the partition
   // number of each row more efficiently with bitwise operations
@@ -784,7 +788,8 @@ struct dispatch_map_type {
     requires(is_index_type<MapType>())
   {
     // Build a histogram of the number of rows in each partition
-    rmm::device_uvector<size_type> histogram(num_partitions + 1, stream);
+    cuda::device_buffer<size_type> histogram(
+      stream, cudf::get_current_device_resource_ref(), num_partitions + 1, cuda::no_init);
     std::size_t temp_storage_bytes{};
     std::size_t const num_levels = num_partitions + 1;
     size_type const lower_level  = 0;
@@ -824,7 +829,8 @@ struct dispatch_map_type {
 
     // Unfortunately need to materialize the scatter map because
     // `detail::scatter` requires multiple passes through the iterator
-    rmm::device_uvector<size_type> scatter_map(partition_map.size(), stream);
+    cuda::device_buffer<size_type> scatter_map(
+      stream, cudf::get_current_device_resource_ref(), partition_map.size(), cuda::no_init);
 
     // For each `partition_map[i]`, atomically increment the corresponding
     // partition offset to determine `i`s location in the output
 
@@ -22,6 +22,7 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/buffer>
 #include <cuda/functional>
 #include <cuda/iterator>
 #include <cuda/stream>
@@ -84,7 +85,8 @@ std::pair<std::unique_ptr<cudf::table>, std::vector<cudf::size_type>> degenerate
                                }));
 
   if (num_partitions == nrows) {
-    rmm::device_uvector<cudf::size_type> partition_offsets(num_partitions + 1, stream);
+    cuda::device_buffer<cudf::size_type> partition_offsets(
+      stream, cudf::get_current_device_resource_ref(), num_partitions + 1, cuda::no_init);
     thrust::sequence(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
                      partition_offsets.begin(),
                      partition_offsets.end());
@@ -98,7 +100,8 @@ std::pair<std::unique_ptr<cudf::table>, std::vector<cudf::size_type>> degenerate
 
     return std::pair{std::move(uniq_tbl), cudf::detail::make_std_vector(partition_offsets, stream)};
   } else {  //( num_partitions > nrows )
-    rmm::device_uvector<cudf::size_type> d_row_indices(nrows, stream);
+    cuda::device_buffer<cudf::size_type> d_row_indices(
+      stream, cudf::get_current_device_resource_ref(), nrows, cuda::no_init);
 
     // copy rotated right partition indexes that
     // fall in the interval [0, nrows):
@@ -128,7 +131,8 @@ std::pair<std::unique_ptr<cudf::table>, std::vector<cudf::size_type>> degenerate
         [nrows] __device__(auto index) { return (index < nrows ? 1 : 0); }));
 
     // offsets (part 2: compute partition offsets):
-    rmm::device_uvector<cudf::size_type> partition_offsets(num_partitions + 1, stream);
+    cuda::device_buffer<cudf::size_type> partition_offsets(
+      stream, cudf::get_current_device_resource_ref(), num_partitions + 1, cuda::no_init);
     thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
                            nedges_iter_begin,
                            nedges_iter_begin + num_partitions,
 
@@ -20,6 +20,7 @@
 
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/buffer>
 #include <cuda/functional>
 #include <cuda/iterator>
 #include <cuda/std/iterator>
@@ -158,19 +159,23 @@ std::unique_ptr<column> compute_lead_lag_for_nested(aggregation::Kind op,
   if (default_outputs.is_empty()) { return std::move(output_with_nulls->release()[0]); }
 
   // Must scatter defaults.
-  auto scatter_map = rmm::device_uvector<size_type>(input.size(), stream);
+  // TODO: Use a resizable buffer when available: https://github.com/NVIDIA/cccl/issues/12051
+  auto scatter_map_storage = cuda::device_buffer<size_type>(
+    stream, cudf::get_current_device_resource_ref(), input.size(), cuda::no_init);
 
   // Find all indices at which LEAD/LAG computed nulls previously.
   auto scatter_map_end =
     cudf::detail::copy_if(cuda::counting_iterator<size_type>{0},
                           cuda::counting_iterator<size_type>{input.size()},
-                          scatter_map.begin(),
+                          scatter_map_storage.begin(),
                           is_null_index_predicate(input.size(), gather_map.begin<size_type>()),
                           stream);
 
-  scatter_map.resize(cuda::std::distance(scatter_map.begin(), scatter_map_end), stream);
+  auto scatter_map = cudf::device_span<size_type>{
+    scatter_map_storage.data(),
+    static_cast<std::size_t>(cuda::std::distance(scatter_map_storage.begin(), scatter_map_end))};
   // Bail early, if all LEAD/LAG computations succeeded. No defaults need be substituted.
-  if (scatter_map.is_empty()) { return std::move(output_with_nulls->release()[0]); }
+  if (scatter_map.empty()) { return std::move(output_with_nulls->release()[0]); }
 
   // Gather only those default values that are to be substituted.
   auto gathered_defaults =
 
@@ -14,6 +14,7 @@
 
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/buffer>
 #include <cuda/iterator>
 #include <cuda/std/iterator>
 #include <cuda/std/limits>
@@ -146,7 +147,8 @@ std::unique_ptr<column> nth_element(size_type n,
     gather_index_calculator<null_handling, PrecedingIter, FollowingIter>{
       n, input, preceding, following, min_periods, stream});
 
-  auto gather_map = rmm::device_uvector<size_type>(input.size(), stream);
+  auto gather_map = cuda::device_buffer<size_type>(
+    stream, cudf::get_current_device_resource_ref(), input.size(), cuda::no_init);
   thrust::copy(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
                gather_iter,
                gather_iter + input.size(),
 
@@ -22,6 +22,7 @@
 
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/buffer>
 #include <cuda/stream>
 #include <thrust/transform.h>
 
@@ -114,7 +115,7 @@ struct contains_scalar_dispatch {
     // Using a temporary buffer for intermediate transform results from the lambda containing
     // the comparator speeds up compile-time significantly without much degradation in
     // runtime performance over using the comparator in a transform iterator with thrust::count_if.
-    auto d_results = rmm::device_uvector<bool>(haystack.size(), stream, temp_mr);
+    auto d_results = cuda::device_buffer<bool>(stream, temp_mr, haystack.size(), cuda::no_init);
     thrust::transform(
       rmm::exec_policy_nosync(stream, temp_mr),
       begin,