Repository: rapidsai/cuspatial (language: python)

Issue:
Replace device_memory_resource* with device_async_resource_ref
Update occurences of `rmm::mr::device_memory_resource*` to
`rmm::device_async_resource_ref` in this repo. The rationale is explained
in [this document](https://docs.google.com/document/d/1QsLQn17qexczFGbzjLzTkIiGztfaQvLu-DpNu2EPw1E/edit#heading=h.wdftxmu1lxt).



Discussion / hints:

Fix patch:
diff --git a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md
index 74db9350f..325c9ce79 100644
--- a/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md
+++ b/cpp/doxygen/developer_guide/DEVELOPER_GUIDE.md
@@ -84,7 +84,7 @@ Examples:
 
 ```c++
 template <typename IteratorType>
-void algorithm_function(int x, rmm::cuda_stream_view s, rmm::device_memory_resource* mr)
+void algorithm_function(int x, rmm::cuda_stream_view s, rmm::device_async_resource_ref mr)
 {
   ...
 }
@@ -233,9 +233,10 @@ std::unique_ptr<cudf::table> points_in_spatial_window(
   cudf::column_view const& y);
 ```
 
-## RMM Memory Resources (`rmm::device_memory_resource`)
+## Memory Resources (`rmm::device_memory_resource`)
 
-libcuspatial allocates all device memory via RMM memory resources (MR). See the
+libcuspatial allocates all device memory via RMM memory resources (MR) or CUDA MRs. Either type
+can be passed to libcuspatial functions via `rmm::device_async_resource_ref` parameters. See the
 [RMM documentation](https://github.com/rapidsai/rmm/blob/main/README.md) for details.
 
 ### Current Device Memory Resource
@@ -245,6 +246,27 @@ RMM provides a "default" memory resource for each device that can be accessed an
 respectively. All memory resource parameters should be defaulted to use the return value of
 `rmm::mr::get_current_device_resource()`.
 
+### Resource Refs
+
+Memory resources are passed via resource ref parameters. A resource ref is memory resource wrapper
+that enables consumers to specify properties of resources that they expect. These are defined
+in the `cuda::mr` namespace of libcu++, but RMM provides some convenience wrappers in
+`rmm/resource_ref.hpp`:
+ - `rmm::device_resource_ref` accepts a memory resource that provides synchronous allocation
+    of device-accessible memory.
+ - `rmm::device_async_resource_ref` accepts a memory resource that provides stream-ordered allocation
+    of device-accessible memory.
+ - `rmm::host_resource_ref` accepts a memory resource that provides synchronous allocation of host- 
+    accessible memory.
+ - `rmm::host_async_resource_ref` accepts a memory resource that provides stream-ordered allocation
+    of host-accessible memory.
+ - `rmm::host_device_resource_ref` accepts a memory resource that provides synchronous allocation of
+    host- and device-accessible memory.
+ - `rmm::host_device_async_resource_ref` accepts a memory resource that provides stream-ordered
+    allocation of host- and device-accessible memory.
+
+See the libcu++ [docs on `resource_ref`](https://nvidia.github.io/cccl/libcudacxx/extended_api/memory_resource/resource_ref.html) for more information.
+
 # libcuspatial API and Implementation
 
 This section provides specifics about the structure and implementation of cuSpatial API functions.
@@ -439,8 +461,8 @@ There are a few key points to notice.
      cuSpatial APIs will not need to use this returned iterator.
   9. All APIs that run CUDA device code (including Thrust algorithms) or allocate memory take a CUDA
      stream on which to execute the device code and allocate memory.
-  10. Any API that allocate and return device data (not shown here) should also take an
-      `rmm::device_memory_resource` to use for output memory allocation.
+  10. Any API that allocates and returns device data (not shown here) should also take an
+      `rmm::device_async_resource_ref` to use for output memory allocation.
 
 ### (Multiple) Return Values
 
@@ -542,22 +564,28 @@ control how device memory is allocated.
 
 ### Output Memory
 
-Any libcuspatial API that allocates memory that is *returned* to a user must accept a pointer to a
-`device_memory_resource` as the last parameter. Inside the API, this memory resource must be used
-to allocate any memory for returned objects. It should therefore be passed into functions whose
-outputs will be returned. Example:
+Any libcuspatial API that allocates memory that is *returned* to a user must accept a
+`rmm::device_async_resource_ref` as the last parameter. Inside the API, this memory resource must
+be used to allocate any memory for returned objects. It should therefore be passed into functions
+whose outputs will be returned. Example:
 
 ```c++
 // Returned `column` contains newly allocated memory,
 // therefore the API must accept a memory resource pointer
 std::unique_ptr<column> returns_output_memory(
-  ..., rmm::device_memory_resource * mr = rmm::mr::get_current_device_resource());
+  ..., rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 // This API does not allocate any new *output* memory, therefore
 // a memory resource is unnecessary
 void does_not_allocate_output_memory(...);
 ```
 
+This rule automatically applies to all detail APIs that allocate memory. Any detail API may be
+called by any public API, and therefore could be allocating memory that is returned to the user.
+To support such uses cases, all detail APIs allocating memory resources should accept an `mr`
+parameter. Callers are responsible for either passing through a provided `mr` or
+`rmm::mr::get_current_device_resource()` as needed.
+
 ### Temporary Memory
 
 Not all memory allocated within a libcuspatial API is returned to the caller. Often algorithms must
@@ -566,7 +594,7 @@ obtained from `rmm::mr::get_current_device_resource()` for temporary memory allo
 
 ```c++
 rmm::device_buffer some_function(
-  ..., rmm::mr::device_memory_resource mr * = rmm::mr::get_current_device_resource()) {
+  ..., rmm::mr::device_async_resource_ref mr = rmm::mr::get_current_device_resource()) {
     rmm::device_buffer returned_buffer(..., mr); // Returned buffer uses the passed in MR
     ...
     rmm::device_buffer temporary_buffer(...); // Temporary buffer uses default MR
@@ -578,12 +606,12 @@ rmm::device_buffer some_function(
 ### Memory Management
 
 libcuspatial code eschews raw pointers and direct memory allocation. Use RMM classes built to
-use [`device_memory_resource`](https://github.com/rapidsai/rmm/#device_memory_resource) for device
+use [memory resources](https://github.com/rapidsai/rmm/#device_memory_resource) for device
 memory allocation with automated lifetime management.
 
 #### rmm::device_buffer
 Allocates a specified number of bytes of untyped, uninitialized device memory using a
-`device_memory_resource`. If no resource is explicitly provided, uses
+memory resource. If no `rmm::device_async_resource_ref` is explicitly provided, uses
 `rmm::mr::get_current_device_resource()`.
 
 `rmm::device_buffer` is movable and copyable on a stream. A copy performs a deep copy of the
diff --git a/cpp/include/cuspatial/bounding_boxes.hpp b/cpp/include/cuspatial/bounding_boxes.hpp
index 9229eb7c2..2134bc1d1 100644
--- a/cpp/include/cuspatial/bounding_boxes.hpp
+++ b/cpp/include/cuspatial/bounding_boxes.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -53,7 +54,7 @@ std::unique_ptr<cudf::table> linestring_bounding_boxes(
   cudf::column_view const& x,
   cudf::column_view const& y,
   double expansion_radius,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute minimum bounding box for each polygon in a list.
@@ -80,8 +81,8 @@ std::unique_ptr<cudf::table> polygon_bounding_boxes(
   cudf::column_view const& ring_offsets,
   cudf::column_view const& x,
   cudf::column_view const& y,
-  double expansion_radius             = 0.0,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  double expansion_radius           = 0.0,
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/detail/index/construction/phase_1.cuh b/cpp/include/cuspatial/detail/index/construction/phase_1.cuh
index 3d06c93b1..a444ae64a 100644
--- a/cpp/include/cuspatial/detail/index/construction/phase_1.cuh
+++ b/cpp/include/cuspatial/detail/index/construction/phase_1.cuh
@@ -24,6 +24,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <cuda/functional>
 #include <thrust/copy.h>
@@ -65,7 +66,7 @@ compute_point_keys_and_sorted_indices(PointIt points_first,
                                       T scale,
                                       int8_t max_depth,
                                       rmm::cuda_stream_view stream,
-                                      rmm::mr::device_memory_resource* mr)
+                                      rmm::device_async_resource_ref mr)
 {
   auto num_points = thrust::distance(points_first, points_last);
   rmm::device_uvector<uint32_t> keys(num_points, stream);
@@ -259,7 +260,7 @@ inline auto make_full_levels(PointIt points_first,
                              T scale,
                              int8_t max_depth,
                              rmm::cuda_stream_view stream,
-                             rmm::mr::device_memory_resource* mr)
+                             rmm::device_async_resource_ref mr)
 {
   auto num_points = thrust::distance(points_first, points_last);
   // Compute point keys and sort into bottom-level quadrants
diff --git a/cpp/include/cuspatial/detail/index/construction/phase_2.cuh b/cpp/include/cuspatial/detail/index/construction/phase_2.cuh
index ffbb79c0d..88c4947e0 100644
--- a/cpp/include/cuspatial/detail/index/construction/phase_2.cuh
+++ b/cpp/include/cuspatial/detail/index/construction/phase_2.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,6 +21,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/copy.h>
 #include <thrust/count.h>
@@ -309,7 +310,7 @@ inline rmm::device_uvector<bool> construct_non_leaf_indicator(
   int32_t num_parent_nodes,
   int32_t num_valid_nodes,
   int32_t max_size,
-  rmm::mr::device_memory_resource* mr,
+  rmm::device_async_resource_ref mr,
   rmm::cuda_stream_view stream)
 {
   //
diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh
index 1d88e5e36..818c6560b 100644
--- a/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh
+++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection.cuh
@@ -34,6 +34,7 @@
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <cuda/atomic>
 #include <thrust/binary_search.h>
@@ -142,7 +143,7 @@ std::unique_ptr<rmm::device_uvector<types_t>> compute_types_buffer(
   OffsetRangeB points_offset,
   OffsetRangeB segments_offset,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   auto types_buffer = std::make_unique<rmm::device_uvector<types_t>>(union_column_size, stream, mr);
   thrust::tabulate(rmm::exec_policy(stream),
@@ -162,7 +163,7 @@ std::unique_ptr<rmm::device_uvector<types_t>> compute_types_buffer(
 template <typename index_t, typename types_t>
 std::unique_ptr<rmm::device_uvector<index_t>> compute_offset_buffer(
   rmm::device_uvector<types_t> const& types_buffer,
-  rmm::mr::device_memory_resource* mr,
+  rmm::device_async_resource_ref mr,
   rmm::cuda_stream_view stream)
 {
   auto N                                = types_buffer.size();
@@ -202,7 +203,7 @@ template <typename T,
 linestring_intersection_result<T, index_t> pairwise_linestring_intersection(
   MultiLinestringRange1 multilinestrings1,
   MultiLinestringRange2 multilinestrings2,
-  rmm::mr::device_memory_resource* mr,
+  rmm::device_async_resource_ref mr,
   rmm::cuda_stream_view stream)
 {
   using types_t = typename linestring_intersection_result<T, index_t>::types_t;
diff --git a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh
index e513cf651..f1a365555 100644
--- a/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh
+++ b/cpp/include/cuspatial/detail/intersection/linestring_intersection_with_duplicates.cuh
@@ -29,6 +29,7 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/binary_search.h>
 #include <thrust/distance.h>
@@ -228,7 +229,7 @@ struct linestring_intersection_intermediates {
   /** @brief Construct a zero-pair, zero-geometry intermediate object
    */
   linestring_intersection_intermediates(rmm::cuda_stream_view stream,
-                                        rmm::mr::device_memory_resource* mr)
+                                        rmm::device_async_resource_ref mr)
     : offsets(std::make_unique<rmm::device_uvector<IndexType>>(1, stream)),
       geoms(std::make_unique<rmm::device_uvector<GeomType>>(0, stream, mr)),
       lhs_linestring_ids(std::make_unique<rmm::device_uvector<IndexType>>(0, stream)),
@@ -244,7 +245,7 @@ struct linestring_intersection_intermediates {
                                         std::size_t num_geoms,
                                         rmm::device_uvector<IndexType> const& num_geoms_per_pair,
                                         rmm::cuda_stream_view stream,
-                                        rmm::mr::device_memory_resource* mr)
+                                        rmm::device_async_resource_ref mr)
     : offsets(std::make_unique<rmm::device_uvector<IndexType>>(num_pairs + 1, stream)),
       geoms(std::make_unique<rmm::device_uvector<GeomType>>(num_geoms, stream, mr)),
       lhs_linestring_ids(std::make_unique<rmm::device_uvector<IndexType>>(num_geoms, stream)),
@@ -472,7 +473,7 @@ std::pair<linestring_intersection_intermediates<vec_2d<T>, index_t>,
           linestring_intersection_intermediates<segment<T>, index_t>>
 pairwise_linestring_intersection_with_duplicates(MultiLinestringRange1 multilinestrings1,
                                                  MultiLinestringRange2 multilinestrings2,
-                                                 rmm::mr::device_memory_resource* mr,
+                                                 rmm::device_async_resource_ref mr,
                                                  rmm::cuda_stream_view stream)
 {
   static_assert(std::is_integral_v<index_t>, "Index type must be integral.");
diff --git a/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh b/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh
index 43ff7b0f6..472bb29b0 100644
--- a/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh
+++ b/cpp/include/cuspatial/detail/join/quadtree_bbox_filtering.cuh
@@ -22,6 +22,8 @@
 #include <cuspatial/point_quadtree.cuh>
 #include <cuspatial/traits.hpp>
 
+#include <rmm/resource_ref.hpp>
+
 #include <cuda/functional>
 #include <thrust/iterator/discard_iterator.h>
 
@@ -39,7 +41,7 @@ join_quadtree_and_bounding_boxes(point_quadtree_ref quadtree,
                                  T scale,
                                  int8_t max_depth,
                                  rmm::cuda_stream_view stream,
-                                 rmm::mr::device_memory_resource* mr)
+                                 rmm::device_async_resource_ref mr)
 {
   static_assert(is_same<T, cuspatial::iterator_vec_base_type<BoundingBoxIterator>>(),
                 "Iterator value_type mismatch");
diff --git a/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh b/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh
index 1ae0ff1bf..0b8f23e19 100644
--- a/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh
+++ b/cpp/include/cuspatial/detail/join/quadtree_point_in_polygon.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,6 +26,7 @@
 
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/permutation_iterator.h>
 #include <thrust/scan.h>
@@ -108,7 +109,7 @@ std::pair<rmm::device_uvector<IndexType>, rmm::device_uvector<IndexType>> quadtr
   PointIterator points_first,
   MultiPolygonRange polygons,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   using T = iterator_vec_base_type<PointIterator>;
 
diff --git a/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh b/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh
index 6335f133f..5b61882b6 100644
--- a/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh
+++ b/cpp/include/cuspatial/detail/join/quadtree_point_to_nearest_linestring.cuh
@@ -27,6 +27,7 @@
 
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <cuda/functional>
 #include <thrust/detail/raw_reference_cast.h>
@@ -156,7 +157,7 @@ quadtree_point_to_nearest_linestring(LinestringIndexIterator linestring_indices_
                                      PointIterator points_first,
                                      MultiLinestringRange linestrings,
                                      rmm::cuda_stream_view stream,
-                                     rmm::mr::device_memory_resource* mr)
+                                     rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(linestrings.num_multilinestrings() == linestrings.num_linestrings(),
                     "Only one linestring per multilinestring currently supported.");
diff --git a/cpp/include/cuspatial/detail/point_quadtree.cuh b/cpp/include/cuspatial/detail/point_quadtree.cuh
index 7e0d70dac..19922af35 100644
--- a/cpp/include/cuspatial/detail/point_quadtree.cuh
+++ b/cpp/include/cuspatial/detail/point_quadtree.cuh
@@ -24,6 +24,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <cuda/functional>
 #include <thrust/distance.h>
@@ -48,7 +49,7 @@ inline point_quadtree make_quad_tree(rmm::device_uvector<uint32_t>& keys,
                                      int32_t max_size,
                                      int32_t level_1_size,
                                      rmm::cuda_stream_view stream,
-                                     rmm::mr::device_memory_resource* mr)
+                                     rmm::device_async_resource_ref mr)
 {
   // count the number of child nodes
   auto num_child_nodes = thrust::reduce(rmm::exec_policy(stream),
@@ -155,7 +156,7 @@ inline point_quadtree make_leaf_tree(rmm::device_uvector<uint32_t>& keys,
                                      rmm::device_uvector<uint32_t>& lengths,
                                      int32_t num_top_quads,
                                      rmm::cuda_stream_view stream,
-                                     rmm::mr::device_memory_resource* mr)
+                                     rmm::device_async_resource_ref mr)
 {
   rmm::device_uvector<uint8_t> levels(num_top_quads, stream, mr);
   rmm::device_uvector<bool> is_internal_node(num_top_quads, stream, mr);
@@ -195,7 +196,7 @@ inline std::pair<rmm::device_uvector<uint32_t>, point_quadtree> construct_quadtr
   T scale,
   int8_t max_depth,
   int32_t max_size,
-  rmm::mr::device_memory_resource* mr,
+  rmm::device_async_resource_ref mr,
   rmm::cuda_stream_view stream)
 {
   // Construct the full set of non-empty subquadrants starting from the lowest level.
@@ -243,7 +244,7 @@ std::pair<rmm::device_uvector<uint32_t>, point_quadtree> quadtree_on_points(
   int8_t max_depth,
   int32_t max_size,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   auto num_points = thrust::distance(points_first, points_last);
   if (num_points <= 0) {
diff --git a/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh b/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh
index 6da914181..11b8a594d 100644
--- a/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh
+++ b/cpp/include/cuspatial/detail/trajectory/derive_trajectories.cuh
@@ -22,6 +22,7 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <cub/device/device_merge_sort.cuh>
 #include <thrust/gather.h>
@@ -63,7 +64,7 @@ void order_trajectories(IdInputIt ids_first,
                         PointOutputIt points_out_first,
                         TimestampOutputIt timestamps_out_first,
                         rmm::cuda_stream_view stream,
-                        rmm::mr::device_memory_resource* mr)
+                        rmm::device_async_resource_ref mr)
 {
   using id_type        = iterator_value_type<IdInputIt>;
   using timestamp_type = iterator_value_type<TimestampInputIt>;
@@ -116,7 +117,7 @@ std::unique_ptr<rmm::device_uvector<OffsetType>> derive_trajectories(
   PointOutputIt points_out_first,
   TimestampOutputIt timestamps_out_first,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   detail::order_trajectories(ids_first,
                              ids_last,
diff --git a/cpp/include/cuspatial/distance.hpp b/cpp/include/cuspatial/distance.hpp
index 11cec4fae..76c920b11 100644
--- a/cpp/include/cuspatial/distance.hpp
+++ b/cpp/include/cuspatial/distance.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -25,6 +25,7 @@
 #include <cudf/utilities/span.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <optional>
 
@@ -53,8 +54,8 @@ std::unique_ptr<cudf::column> haversine_distance(
   cudf::column_view const& a_lat,
   cudf::column_view const& b_lon,
   cudf::column_view const& b_lat,
-  double const radius                 = EARTH_RADIUS_KM,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  double const radius               = EARTH_RADIUS_KM,
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief computes Hausdorff distances for all pairs in a collection of spaces
@@ -124,7 +125,7 @@ std::pair<std::unique_ptr<cudf::column>, cudf::table_view> directed_hausdorff_di
   cudf::column_view const& xs,
   cudf::column_view const& ys,
   cudf::column_view const& space_offsets,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute pairwise (multi)point-to-(multi)point Cartesian distance
@@ -144,7 +145,7 @@ std::pair<std::unique_ptr<cudf::column>, cudf::table_view> directed_hausdorff_di
 std::unique_ptr<cudf::column> pairwise_point_distance(
   geometry_column_view const& multipoints1,
   geometry_column_view const& multipoints2,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute pairwise (multi)points-to-(multi)linestrings Cartesian distance
@@ -166,7 +167,7 @@ std::unique_ptr<cudf::column> pairwise_point_distance(
 std::unique_ptr<cudf::column> pairwise_point_linestring_distance(
   geometry_column_view const& multipoints,
   geometry_column_view const& multilinestrings,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute pairwise (multi)point-to-(multi)polygon Cartesian distance
@@ -190,7 +191,7 @@ std::unique_ptr<cudf::column> pairwise_point_linestring_distance(
 std::unique_ptr<cudf::column> pairwise_point_polygon_distance(
   geometry_column_view const& multipoints,
   geometry_column_view const& multipolygons,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute pairwise (multi)linestring-to-(multi)linestring Cartesian distance
@@ -212,7 +213,7 @@ std::unique_ptr<cudf::column> pairwise_point_polygon_distance(
 std::unique_ptr<cudf::column> pairwise_linestring_distance(
   geometry_column_view const& multilinestrings1,
   geometry_column_view const& multilinestrings2,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute pairwise (multi)linestring-to-(multi)polygon Cartesian distance
@@ -237,7 +238,7 @@ std::unique_ptr<cudf::column> pairwise_linestring_distance(
 std::unique_ptr<cudf::column> pairwise_linestring_polygon_distance(
   geometry_column_view const& multilinestrings,
   geometry_column_view const& multipolygons,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute pairwise (multi)polygon-to-(multi)polygon Cartesian distance
@@ -256,7 +257,7 @@ std::unique_ptr<cudf::column> pairwise_linestring_polygon_distance(
 std::unique_ptr<cudf::column> pairwise_polygon_distance(
   geometry_column_view const& multipolygons1,
   geometry_column_view const& multipolygons2,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/intersection.cuh b/cpp/include/cuspatial/intersection.cuh
index 331e6eb0a..c086d84cc 100644
--- a/cpp/include/cuspatial/intersection.cuh
+++ b/cpp/include/cuspatial/intersection.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -22,6 +22,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/pair.h>
 
@@ -87,8 +88,8 @@ template <typename T,
 linestring_intersection_result<T, index_t> pairwise_linestring_intersection(
   MultiLinestringRange1 multilinestrings1,
   MultiLinestringRange2 multilinestrings2,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource(),
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default);
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource(),
+  rmm::cuda_stream_view stream      = rmm::cuda_stream_default);
 
 }  // namespace cuspatial
 
diff --git a/cpp/include/cuspatial/intersection.hpp b/cpp/include/cuspatial/intersection.hpp
index aedaeb521..1869768b1 100644
--- a/cpp/include/cuspatial/intersection.hpp
+++ b/cpp/include/cuspatial/intersection.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,6 +21,7 @@
 #include <cudf/column/column.hpp>
 
 #include <rmm/mr/device/device_memory_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 namespace cuspatial {
 /**
@@ -54,6 +55,6 @@ struct linestring_intersection_column_result {
 linestring_intersection_column_result pairwise_linestring_intersection(
   geometry_column_view const& multilinestrings_lhs,
   geometry_column_view const& multilinestrings_rhs,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 }  // namespace cuspatial
diff --git a/cpp/include/cuspatial/nearest_points.hpp b/cpp/include/cuspatial/nearest_points.hpp
index 640a1c9b6..46f0c14a2 100644
--- a/cpp/include/cuspatial/nearest_points.hpp
+++ b/cpp/include/cuspatial/nearest_points.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,8 @@
 #include <cudf/column/column_view.hpp>
 #include <cudf/utilities/span.hpp>
 
+#include <rmm/resource_ref.hpp>
+
 #include <optional>
 
 namespace cuspatial {
@@ -165,7 +167,7 @@ point_linestring_nearest_points_result pairwise_point_linestring_nearest_points(
   std::optional<cudf::device_span<cudf::size_type const>> multilinestring_geometry_offsets,
   cudf::device_span<cudf::size_type const> linestring_part_offsets,
   cudf::column_view linestring_points_xy,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp b/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp
index 9bfb13700..0df60c269 100644
--- a/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp
+++ b/cpp/include/cuspatial/pairwise_multipoint_equals_count.hpp
@@ -21,6 +21,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -69,6 +70,6 @@ namespace cuspatial {
 std::unique_ptr<cudf::column> pairwise_multipoint_equals_count(
   geometry_column_view const& lhs,
   geometry_column_view const& rhs,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 }  // namespace cuspatial
diff --git a/cpp/include/cuspatial/point_in_polygon.hpp b/cpp/include/cuspatial/point_in_polygon.hpp
index 11e7381c3..c66e2e260 100644
--- a/cpp/include/cuspatial/point_in_polygon.hpp
+++ b/cpp/include/cuspatial/point_in_polygon.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,6 +21,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -78,7 +79,7 @@ std::unique_ptr<cudf::column> point_in_polygon(
   cudf::column_view const& poly_ring_offsets,
   cudf::column_view const& poly_points_x,
   cudf::column_view const& poly_points_y,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Given (point, polygon pairs), tests whether the point of each pair is inside the polygon
@@ -127,7 +128,7 @@ std::unique_ptr<cudf::column> pairwise_point_in_polygon(
   cudf::column_view const& poly_ring_offsets,
   cudf::column_view const& poly_points_x,
   cudf::column_view const& poly_points_y,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/point_quadtree.cuh b/cpp/include/cuspatial/point_quadtree.cuh
index a84b07558..f64cee6f0 100644
--- a/cpp/include/cuspatial/point_quadtree.cuh
+++ b/cpp/include/cuspatial/point_quadtree.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -22,6 +22,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/tuple.h>
@@ -175,8 +176,8 @@ std::pair<rmm::device_uvector<uint32_t>, point_quadtree> quadtree_on_points(
   T scale,
   int8_t max_depth,
   int32_t max_size,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::cuda_stream_view stream      = rmm::cuda_stream_default,
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/point_quadtree.hpp b/cpp/include/cuspatial/point_quadtree.hpp
index 17e156ea6..72306730f 100644
--- a/cpp/include/cuspatial/point_quadtree.hpp
+++ b/cpp/include/cuspatial/point_quadtree.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -74,7 +75,7 @@ std::pair<std::unique_ptr<cudf::column>, std::unique_ptr<cudf::table>> quadtree_
   double scale,
   int8_t max_depth,
   cudf::size_type max_size,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/points_in_range.hpp b/cpp/include/cuspatial/points_in_range.hpp
index 7bd23c2c0..8637e2b24 100644
--- a/cpp/include/cuspatial/points_in_range.hpp
+++ b/cpp/include/cuspatial/points_in_range.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,6 +21,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -60,7 +61,7 @@ std::unique_ptr<cudf::table> points_in_range(
   double range_max_y,
   cudf::column_view const& x,
   cudf::column_view const& y,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/projection.hpp b/cpp/include/cuspatial/projection.hpp
index 1158d8c7c..e71164148 100644
--- a/cpp/include/cuspatial/projection.hpp
+++ b/cpp/include/cuspatial/projection.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -50,7 +51,7 @@ std::pair<std::unique_ptr<cudf::column>, std::unique_ptr<cudf::column>> sinusoid
   double origin_lat,
   cudf::column_view const& input_lon,
   cudf::column_view const& input_lat,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/spatial_join.cuh b/cpp/include/cuspatial/spatial_join.cuh
index 9815d5f4f..096af5fb2 100644
--- a/cpp/include/cuspatial/spatial_join.cuh
+++ b/cpp/include/cuspatial/spatial_join.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -23,6 +23,7 @@
 #include <cuspatial/traits.hpp>
 
 #include <rmm/device_uvector.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <iterator>
 #include <utility>
@@ -69,8 +70,8 @@ join_quadtree_and_bounding_boxes(
   vec_2d<T> const& v_min,
   T scale,
   int8_t max_depth,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::cuda_stream_view stream      = rmm::cuda_stream_default,
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Test whether the specified points are inside any of the specified polygons.
@@ -127,8 +128,8 @@ std::pair<rmm::device_uvector<IndexType>, rmm::device_uvector<IndexType>> quadtr
   PointIndexIterator point_indices_last,
   PointIterator points_first,
   MultiPolygonRange polygons,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::cuda_stream_view stream      = rmm::cuda_stream_default,
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Finds the nearest linestring to each point in a quadrant, and computes the distances
@@ -184,8 +185,8 @@ quadtree_point_to_nearest_linestring(
   PointIndexIterator point_indices_last,
   PointIterator points_first,
   MultiLinestringRange linestrings,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::cuda_stream_view stream      = rmm::cuda_stream_default,
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 }  // namespace cuspatial
 
diff --git a/cpp/include/cuspatial/spatial_join.hpp b/cpp/include/cuspatial/spatial_join.hpp
index 579a3c4f1..e3bec92d8 100644
--- a/cpp/include/cuspatial/spatial_join.hpp
+++ b/cpp/include/cuspatial/spatial_join.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -71,7 +72,7 @@ std::unique_ptr<cudf::table> join_quadtree_and_bounding_boxes(
   double y_max,
   double scale,
   int8_t max_depth,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Test whether the specified points are inside any of the specified polygons.
@@ -122,7 +123,7 @@ std::unique_ptr<cudf::table> quadtree_point_in_polygon(
   cudf::column_view const& ring_offsets,
   cudf::column_view const& poly_points_x,
   cudf::column_view const& poly_points_y,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Finds the nearest linestring to each point in a quadrant, and computes the distances
@@ -171,7 +172,7 @@ std::unique_ptr<cudf::table> quadtree_point_to_nearest_linestring(
   cudf::column_view const& linestring_offsets,
   cudf::column_view const& linestring_points_x,
   cudf::column_view const& linestring_points_y,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/include/cuspatial/trajectory.cuh b/cpp/include/cuspatial/trajectory.cuh
index f6c5d76f4..dc5ec9708 100644
--- a/cpp/include/cuspatial/trajectory.cuh
+++ b/cpp/include/cuspatial/trajectory.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <iterator>
 #include <memory>
@@ -85,8 +86,8 @@ std::unique_ptr<rmm::device_uvector<OffsetType>> derive_trajectories(
   IdOutputIt ids_output_first,
   PointOutputIt points_output_first,
   TimestampOutputIt timestamps_output_first,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::cuda_stream_view stream      = rmm::cuda_stream_default,
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute the total distance (in meters) and average speed (in m/s) of objects in
diff --git a/cpp/include/cuspatial/trajectory.hpp b/cpp/include/cuspatial/trajectory.hpp
index 5eb5eaf0d..308063569 100644
--- a/cpp/include/cuspatial/trajectory.hpp
+++ b/cpp/include/cuspatial/trajectory.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -22,6 +22,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -62,7 +63,7 @@ std::pair<std::unique_ptr<cudf::table>, std::unique_ptr<cudf::column>> derive_tr
   cudf::column_view const& x,
   cudf::column_view const& y,
   cudf::column_view const& timestamp,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute the distance and speed of objects in a trajectory. Groups the
@@ -95,7 +96,7 @@ std::unique_ptr<cudf::table> trajectory_distances_and_speeds(
   cudf::column_view const& x,
   cudf::column_view const& y,
   cudf::column_view const& timestamp,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Compute the spatial bounding boxes of trajectories. Groups the x, y,
@@ -127,7 +128,7 @@ std::unique_ptr<cudf::table> trajectory_bounding_boxes(
   cudf::column_view const& object_id,
   cudf::column_view const& x,
   cudf::column_view const& y,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource());
 
 /**
  * @} // end of doxygen group
diff --git a/cpp/src/bounding_boxes/linestring_bounding_boxes.cu b/cpp/src/bounding_boxes/linestring_bounding_boxes.cu
index 86b6aeb75..7716d231c 100644
--- a/cpp/src/bounding_boxes/linestring_bounding_boxes.cu
+++ b/cpp/src/bounding_boxes/linestring_bounding_boxes.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,6 +26,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/zip_iterator.h>
 
@@ -43,7 +44,7 @@ std::unique_ptr<cudf::table> compute_linestring_bounding_boxes(
   cudf::column_view const& y,
   T expansion_radius,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   auto num_linestrings = linestring_offsets.size() > 0 ? linestring_offsets.size() - 1 : 0;
 
@@ -93,7 +94,7 @@ struct dispatch_compute_linestring_bounding_boxes {
              cudf::column_view const& y,
              double expansion_radius,
              rmm::cuda_stream_view stream,
-             rmm::mr::device_memory_resource* mr)
+             rmm::device_async_resource_ref mr)
   {
     return compute_linestring_bounding_boxes<T>(
       linestring_offsets, x, y, static_cast<T>(expansion_radius), stream, mr);
@@ -109,7 +110,7 @@ std::unique_ptr<cudf::table> linestring_bounding_boxes(cudf::column_view const&
                                                        cudf::column_view const& y,
                                                        double expansion_radius,
                                                        rmm::cuda_stream_view stream,
-                                                       rmm::mr::device_memory_resource* mr)
+                                                       rmm::device_async_resource_ref mr)
 {
   return cudf::type_dispatcher(x.type(),
                                dispatch_compute_linestring_bounding_boxes{},
@@ -127,7 +128,7 @@ std::unique_ptr<cudf::table> linestring_bounding_boxes(cudf::column_view const&
                                                        cudf::column_view const& x,
                                                        cudf::column_view const& y,
                                                        double expansion_radius,
-                                                       rmm::mr::device_memory_resource* mr)
+                                                       rmm::device_async_resource_ref mr)
 {
   auto num_linestrings = linestring_offsets.size() > 0 ? linestring_offsets.size() - 1 : 0;
 
diff --git a/cpp/src/bounding_boxes/polygon_bounding_boxes.cu b/cpp/src/bounding_boxes/polygon_bounding_boxes.cu
index e38e12dfd..c4c092e55 100644
--- a/cpp/src/bounding_boxes/polygon_bounding_boxes.cu
+++ b/cpp/src/bounding_boxes/polygon_bounding_boxes.cu
@@ -27,6 +27,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/zip_iterator.h>
 
@@ -43,7 +44,7 @@ std::unique_ptr<cudf::table> compute_polygon_bounding_boxes(cudf::column_view co
                                                             cudf::column_view const& y,
                                                             T expansion_radius,
                                                             rmm::cuda_stream_view stream,
-                                                            rmm::mr::device_memory_resource* mr)
+                                                            rmm::device_async_resource_ref mr)
 {
   auto num_polygons = poly_offsets.size() > 0 ? poly_offsets.size() - 1 : 0;
 
@@ -96,7 +97,7 @@ struct dispatch_compute_polygon_bounding_boxes {
              cudf::column_view const& y,
              T expansion_radius,
              rmm::cuda_stream_view stream,
-             rmm::mr::device_memory_resource* mr)
+             rmm::device_async_resource_ref mr)
   {
     return compute_polygon_bounding_boxes<T>(
       poly_offsets, ring_offsets, x, y, expansion_radius, stream, mr);
@@ -113,7 +114,7 @@ std::unique_ptr<cudf::table> polygon_bounding_boxes(cudf::column_view const& pol
                                                     cudf::column_view const& y,
                                                     double expansion_radius,
                                                     rmm::cuda_stream_view stream,
-                                                    rmm::mr::device_memory_resource* mr)
+                                                    rmm::device_async_resource_ref mr)
 {
   return cudf::type_dispatcher(x.type(),
                                dispatch_compute_polygon_bounding_boxes{},
@@ -133,7 +134,7 @@ std::unique_ptr<cudf::table> polygon_bounding_boxes(cudf::column_view const& pol
                                                     cudf::column_view const& x,
                                                     cudf::column_view const& y,
                                                     double expansion_radius,
-                                                    rmm::mr::device_memory_resource* mr)
+                                                    rmm::device_async_resource_ref mr)
 {
   auto num_polys = poly_offsets.size() > 0 ? poly_offsets.size() - 1 : 0;
   auto num_rings = ring_offsets.size() > 0 ? ring_offsets.size() - 1 : 0;
diff --git a/cpp/src/distance/hausdorff.cu b/cpp/src/distance/hausdorff.cu
index 9706b5bf6..b4f2a62d1 100644
--- a/cpp/src/distance/hausdorff.cu
+++ b/cpp/src/distance/hausdorff.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -32,6 +32,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/binary_search.h>
 #include <thrust/distance.h>
@@ -81,7 +82,7 @@ struct hausdorff_functor {
              cudf::column_view const& ys,
              cudf::column_view const& space_offsets,
              rmm::cuda_stream_view stream,
-             rmm::mr::device_memory_resource* mr)
+             rmm::device_async_resource_ref mr)
   {
     auto const num_points = static_cast<uint32_t>(xs.size());
     auto const num_spaces = static_cast<uint32_t>(space_offsets.size());
@@ -120,7 +121,7 @@ std::pair<std::unique_ptr<cudf::column>, cudf::table_view> directed_hausdorff_di
   cudf::column_view const& xs,
   cudf::column_view const& ys,
   cudf::column_view const& space_offsets,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(xs.type() == ys.type(), "Inputs `xs` and `ys` must have same type.");
   CUSPATIAL_EXPECTS(xs.size() == ys.size(), "Inputs `xs` and `ys` must have same length.");
diff --git a/cpp/src/distance/haversine.cu b/cpp/src/distance/haversine.cu
index 3c4e631da..18a55cd2d 100644
--- a/cpp/src/distance/haversine.cu
+++ b/cpp/src/distance/haversine.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -27,6 +27,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 #include <type_traits>
@@ -49,7 +50,7 @@ struct haversine_functor {
     cudf::column_view const& b_lat,
     T radius,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     if (a_lon.is_empty()) { return cudf::empty_like(a_lon); }
 
@@ -81,7 +82,7 @@ std::unique_ptr<cudf::column> haversine_distance(cudf::column_view const& a_lon,
                                                  cudf::column_view const& b_lat,
                                                  double radius,
                                                  rmm::cuda_stream_view stream,
-                                                 rmm::mr::device_memory_resource* mr)
+                                                 rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(radius > 0, "radius must be positive.");
 
@@ -108,7 +109,7 @@ std::unique_ptr<cudf::column> haversine_distance(cudf::column_view const& a_lon,
                                                  cudf::column_view const& b_lon,
                                                  cudf::column_view const& b_lat,
                                                  double radius,
-                                                 rmm::mr::device_memory_resource* mr)
+                                                 rmm::device_async_resource_ref mr)
 {
   return cuspatial::detail::haversine_distance(
     a_lon, a_lat, b_lon, b_lat, radius, rmm::cuda_stream_default, mr);
diff --git a/cpp/src/distance/linestring_distance.cu b/cpp/src/distance/linestring_distance.cu
index fe7b99fda..840d82b43 100644
--- a/cpp/src/distance/linestring_distance.cu
+++ b/cpp/src/distance/linestring_distance.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -28,6 +28,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 #include <type_traits>
@@ -42,7 +43,7 @@ struct pairwise_linestring_distance_launch {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multilinestrings1,
                                            geometry_column_view const& multilinestrings2,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     auto size = multilinestrings1.size();
 
@@ -71,7 +72,7 @@ struct pairwise_linestring_distance_functor {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multilinestrings1,
                                            geometry_column_view const& multilinestrings2,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     CUSPATIAL_EXPECTS(multilinestrings1.geometry_type() == geometry_type_id::LINESTRING &&
                         multilinestrings2.geometry_type() == geometry_type_id::LINESTRING,
@@ -96,7 +97,7 @@ struct pairwise_linestring_distance_functor {
 std::unique_ptr<cudf::column> pairwise_linestring_distance(
   geometry_column_view const& multilinestrings1,
   geometry_column_view const& multilinestrings2,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   return multi_geometry_double_dispatch<detail::pairwise_linestring_distance_functor>(
     multilinestrings1.collection_type(),
diff --git a/cpp/src/distance/linestring_polygon_distance.cu b/cpp/src/distance/linestring_polygon_distance.cu
index b6d370b90..552ab792d 100644
--- a/cpp/src/distance/linestring_polygon_distance.cu
+++ b/cpp/src/distance/linestring_polygon_distance.cu
@@ -34,6 +34,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/counting_iterator.h>
 
@@ -54,7 +55,7 @@ struct pairwise_linestring_polygon_distance_impl {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multilinestrings,
                                            geometry_column_view const& multipolygons,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     auto multilinestrings_range =
       make_multilinestring_range<is_multi_linestring, T, cudf::size_type>(multilinestrings);
@@ -87,7 +88,7 @@ struct pairwise_linestring_polygon_distance {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multilinestrings,
                                            geometry_column_view const& multipolygons,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     CUSPATIAL_EXPECTS(multilinestrings.geometry_type() == geometry_type_id::LINESTRING &&
                         multipolygons.geometry_type() == geometry_type_id::POLYGON,
@@ -114,7 +115,7 @@ struct pairwise_linestring_polygon_distance {
 std::unique_ptr<cudf::column> pairwise_linestring_polygon_distance(
   geometry_column_view const& multilinestrings,
   geometry_column_view const& multipolygons,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   return multi_geometry_double_dispatch<detail::pairwise_linestring_polygon_distance>(
     multilinestrings.collection_type(),
diff --git a/cpp/src/distance/point_distance.cu b/cpp/src/distance/point_distance.cu
index d2349b7a2..960c500d1 100644
--- a/cpp/src/distance/point_distance.cu
+++ b/cpp/src/distance/point_distance.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -27,6 +27,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 #include <type_traits>
@@ -40,7 +41,7 @@ struct pairwise_point_distance_impl {
     geometry_column_view const& multipoints1,
     geometry_column_view const& multipoints2,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     auto size = multipoints1.size();
 
@@ -68,7 +69,7 @@ struct pairwise_point_distance_functor {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multipoints1,
                                            geometry_column_view const& multipoints2,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     CUSPATIAL_EXPECTS(multipoints1.geometry_type() == geometry_type_id::POINT &&
                         multipoints2.geometry_type() == geometry_type_id::POINT,
@@ -94,7 +95,7 @@ struct pairwise_point_distance_functor {
 
 std::unique_ptr<cudf::column> pairwise_point_distance(geometry_column_view const& multipoints1,
                                                       geometry_column_view const& multipoints2,
-                                                      rmm::mr::device_memory_resource* mr)
+                                                      rmm::device_async_resource_ref mr)
 {
   return multi_geometry_double_dispatch<detail::pairwise_point_distance_functor>(
     multipoints1.collection_type(),
diff --git a/cpp/src/distance/point_linestring_distance.cu b/cpp/src/distance/point_linestring_distance.cu
index 59cb1b460..42d45b152 100644
--- a/cpp/src/distance/point_linestring_distance.cu
+++ b/cpp/src/distance/point_linestring_distance.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -30,6 +30,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 #include <type_traits>
@@ -48,7 +49,7 @@ struct pairwise_point_linestring_distance_impl {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multipoints,
                                            geometry_column_view const& multilinestrings,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     auto size = multipoints.size();
 
@@ -79,7 +80,7 @@ struct pairwise_point_linestring_distance_functor {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multipoints,
                                            geometry_column_view const& multilinestrings,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     CUSPATIAL_EXPECTS(multipoints.geometry_type() == geometry_type_id::POINT &&
                         multilinestrings.geometry_type() == geometry_type_id::LINESTRING,
@@ -106,7 +107,7 @@ struct pairwise_point_linestring_distance_functor {
 std::unique_ptr<cudf::column> pairwise_point_linestring_distance(
   geometry_column_view const& multipoints,
   geometry_column_view const& multilinestrings,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   return multi_geometry_double_dispatch<detail::pairwise_point_linestring_distance_functor>(
     multipoints.collection_type(),
diff --git a/cpp/src/distance/point_polygon_distance.cu b/cpp/src/distance/point_polygon_distance.cu
index dd7364fbd..b75b328bc 100644
--- a/cpp/src/distance/point_polygon_distance.cu
+++ b/cpp/src/distance/point_polygon_distance.cu
@@ -34,6 +34,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/counting_iterator.h>
 
@@ -54,7 +55,7 @@ struct pairwise_point_polygon_distance_impl {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multipoints,
                                            geometry_column_view const& multipolygons,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     auto multipoints_range = make_multipoint_range<is_multi_point, T, cudf::size_type>(multipoints);
     auto multipolygons_range =
@@ -83,7 +84,7 @@ struct pairwise_point_polygon_distance {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& multipoints,
                                            geometry_column_view const& multipolygons,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     return cudf::type_dispatcher(
       multipoints.coordinate_type(),
@@ -100,7 +101,7 @@ struct pairwise_point_polygon_distance {
 std::unique_ptr<cudf::column> pairwise_point_polygon_distance(
   geometry_column_view const& multipoints,
   geometry_column_view const& multipolygons,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(multipoints.geometry_type() == geometry_type_id::POINT &&
                       multipolygons.geometry_type() == geometry_type_id::POLYGON,
diff --git a/cpp/src/distance/polygon_distance.cu b/cpp/src/distance/polygon_distance.cu
index 678aa3f56..c5a6b0ed1 100644
--- a/cpp/src/distance/polygon_distance.cu
+++ b/cpp/src/distance/polygon_distance.cu
@@ -32,6 +32,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/counting_iterator.h>
 
@@ -52,7 +53,7 @@ struct pairwise_polygon_distance_impl {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& lhs,
                                            geometry_column_view const& rhs,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     auto lhs_range = make_multipolygon_range<is_multi_polygon_lhs, T, cudf::size_type>(lhs);
     auto rhs_range = make_multipolygon_range<is_multi_polygon_rhs, T, cudf::size_type>(rhs);
@@ -80,7 +81,7 @@ struct pairwise_polygon_distance {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& lhs,
                                            geometry_column_view const& rhs,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::POLYGON &&
                         rhs.geometry_type() == geometry_type_id::POLYGON,
@@ -106,7 +107,7 @@ struct pairwise_polygon_distance {
 
 std::unique_ptr<cudf::column> pairwise_polygon_distance(geometry_column_view const& lhs,
                                                         geometry_column_view const& rhs,
-                                                        rmm::mr::device_memory_resource* mr)
+                                                        rmm::device_async_resource_ref mr)
 {
   return multi_geometry_double_dispatch<detail::pairwise_polygon_distance>(
     lhs.collection_type(), rhs.collection_type(), lhs, rhs, rmm::cuda_stream_default, mr);
diff --git a/cpp/src/equality/pairwise_multipoint_equals_count.cu b/cpp/src/equality/pairwise_multipoint_equals_count.cu
index 71e2b6d2b..abfde9e9c 100644
--- a/cpp/src/equality/pairwise_multipoint_equals_count.cu
+++ b/cpp/src/equality/pairwise_multipoint_equals_count.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,6 +26,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/pair.h>
@@ -47,7 +48,7 @@ struct pairwise_multipoint_equals_count_impl {
   std::unique_ptr<cudf::column> operator()(geometry_column_view const& lhs,
                                            geometry_column_view const& rhs,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     auto size = lhs.size();  // lhs is a buffer of xy coords
     auto type = cudf::data_type(cudf::type_to_id<uint32_t>());
@@ -78,7 +79,7 @@ struct pairwise_multipoint_equals_count {
   std::unique_ptr<cudf::column> operator()(geometry_column_view lhs,
                                            geometry_column_view rhs,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     return cudf::type_dispatcher(
       lhs.coordinate_type(),
@@ -94,7 +95,7 @@ struct pairwise_multipoint_equals_count {
 
 std::unique_ptr<cudf::column> pairwise_multipoint_equals_count(geometry_column_view const& lhs,
                                                                geometry_column_view const& rhs,
-                                                               rmm::mr::device_memory_resource* mr)
+                                                               rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::POINT &&
                       rhs.geometry_type() == geometry_type_id::POINT,
diff --git a/cpp/src/indexing/point_quadtree.cu b/cpp/src/indexing/point_quadtree.cu
index 4e2f5987f..a8e53f1e1 100644
--- a/cpp/src/indexing/point_quadtree.cu
+++ b/cpp/src/indexing/point_quadtree.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,6 +26,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 
@@ -72,7 +73,7 @@ struct dispatch_construct_quadtree {
     int8_t max_depth,
     cudf::size_type max_size,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     auto points = cuspatial::make_vec_2d_iterator(x.begin<T>(), y.begin<T>());
     auto [point_indices, tree] =
@@ -131,7 +132,7 @@ std::pair<std::unique_ptr<cudf::column>, std::unique_ptr<cudf::table>> quadtree_
   int8_t max_depth,
   cudf::size_type max_size,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   return cudf::type_dispatcher(x.type(),
                                dispatch_construct_quadtree{},
@@ -160,7 +161,7 @@ std::pair<std::unique_ptr<cudf::column>, std::unique_ptr<cudf::table>> quadtree_
   double scale,
   int8_t max_depth,
   cudf::size_type max_size,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(x.size() == y.size(), "x and y columns must have the same length");
   if (x.is_empty() || y.is_empty()) {
diff --git a/cpp/src/intersection/linestring_intersection.cu b/cpp/src/intersection/linestring_intersection.cu
index 30e60e798..46734f69d 100644
--- a/cpp/src/intersection/linestring_intersection.cu
+++ b/cpp/src/intersection/linestring_intersection.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -34,6 +34,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 #include <type_traits>
@@ -43,7 +44,7 @@ namespace detail {
 
 std::unique_ptr<cudf::column> even_sequence(cudf::size_type size,
                                             rmm::cuda_stream_view stream,
-                                            rmm::mr::device_memory_resource* mr)
+                                            rmm::device_async_resource_ref mr)
 {
   auto res = cudf::make_numeric_column(cudf::data_type{cudf::type_to_id<cudf::size_type>()},
                                        size,
@@ -67,7 +68,7 @@ struct pairwise_linestring_intersection_launch {
   operator()(geometry_column_view const& multilinestrings1,
              geometry_column_view const& multilinestrings2,
              rmm::cuda_stream_view stream,
-             rmm::mr::device_memory_resource* mr)
+             rmm::device_async_resource_ref mr)
   {
     using index_t = cudf::size_type;
 
@@ -155,7 +156,7 @@ struct pairwise_linestring_intersection {
   linestring_intersection_column_result operator()(geometry_column_view const& linestrings1,
                                                    geometry_column_view const& linestrings2,
                                                    rmm::cuda_stream_view stream,
-                                                   rmm::mr::device_memory_resource* mr)
+                                                   rmm::device_async_resource_ref mr)
   {
     CUSPATIAL_EXPECTS(linestrings1.coordinate_type() == linestrings2.coordinate_type(),
                       "Input linestring coordinates must be the same type.");
@@ -176,7 +177,7 @@ struct pairwise_linestring_intersection {
 linestring_intersection_column_result pairwise_linestring_intersection(
   geometry_column_view const& lhs,
   geometry_column_view const& rhs,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(lhs.geometry_type() == geometry_type_id::LINESTRING &&
                       rhs.geometry_type() == geometry_type_id::LINESTRING,
diff --git a/cpp/src/join/quadtree_bbox_filtering.cu b/cpp/src/join/quadtree_bbox_filtering.cu
index 59bd369cd..1757938f3 100644
--- a/cpp/src/join/quadtree_bbox_filtering.cu
+++ b/cpp/src/join/quadtree_bbox_filtering.cu
@@ -27,6 +27,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <tuple>
 
@@ -42,7 +43,7 @@ struct dispatch_quadtree_bounding_box_join {
                                                  double y_min,
                                                  double scale,
                                                  int8_t max_depth,
-                                                 rmm::mr::device_memory_resource* mr,
+                                                 rmm::device_async_resource_ref mr,
                                                  rmm::cuda_stream_view stream)
   {
     auto bbox_min = cuspatial::make_vec_2d_iterator(bbox.column(0).template begin<T>(),
@@ -94,7 +95,7 @@ std::unique_ptr<cudf::table> join_quadtree_and_bounding_boxes(cudf::table_view c
                                                               double y_max,
                                                               double scale,
                                                               int8_t max_depth,
-                                                              rmm::mr::device_memory_resource* mr)
+                                                              rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(quadtree.num_columns() == 5, "quadtree table must have 5 columns");
   CUSPATIAL_EXPECTS(bbox.num_columns() == 4, "bbox table must have 4 columns");
diff --git a/cpp/src/join/quadtree_point_in_polygon.cu b/cpp/src/join/quadtree_point_in_polygon.cu
index 75b723aa8..48f9055ba 100644
--- a/cpp/src/join/quadtree_point_in_polygon.cu
+++ b/cpp/src/join/quadtree_point_in_polygon.cu
@@ -28,6 +28,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_buffer.hpp>
+#include <rmm/resource_ref.hpp>
 
 namespace cuspatial {
 namespace detail {
@@ -52,7 +53,7 @@ struct compute_quadtree_point_in_polygon {
     cudf::column_view const& poly_points_x,
     cudf::column_view const& poly_points_y,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     auto poly_indices = poly_quad_pairs.column(0);
     auto quad_indices = poly_quad_pairs.column(1);
@@ -120,7 +121,7 @@ std::unique_ptr<cudf::table> quadtree_point_in_polygon(cudf::table_view const& p
                                                        cudf::column_view const& poly_points_x,
                                                        cudf::column_view const& poly_points_y,
                                                        rmm::cuda_stream_view stream,
-                                                       rmm::mr::device_memory_resource* mr)
+                                                       rmm::device_async_resource_ref mr)
 {
   return cudf::type_dispatcher(point_x.type(),
                                compute_quadtree_point_in_polygon{},
@@ -148,7 +149,7 @@ std::unique_ptr<cudf::table> quadtree_point_in_polygon(cudf::table_view const& p
                                                        cudf::column_view const& ring_offsets,
                                                        cudf::column_view const& poly_points_x,
                                                        cudf::column_view const& poly_points_y,
-                                                       rmm::mr::device_memory_resource* mr)
+                                                       rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(poly_quad_pairs.num_columns() == 2,
                     "a quadrant-polygon table must have 2 columns");
diff --git a/cpp/src/join/quadtree_point_to_nearest_linestring.cu b/cpp/src/join/quadtree_point_to_nearest_linestring.cu
index c423c1cb4..8288ecbc3 100644
--- a/cpp/src/join/quadtree_point_to_nearest_linestring.cu
+++ b/cpp/src/join/quadtree_point_to_nearest_linestring.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -28,6 +28,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <limits>
 #include <memory>
@@ -54,7 +55,7 @@ struct compute_quadtree_point_to_nearest_linestring {
     cudf::column_view const& linestring_points_x,
     cudf::column_view const& linestring_points_y,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     auto linestring_indices = linestring_quad_pairs.column(0);
     auto quad_indices       = linestring_quad_pairs.column(1);
@@ -120,7 +121,7 @@ std::unique_ptr<cudf::table> quadtree_point_to_nearest_linestring(
   cudf::column_view const& linestring_points_x,
   cudf::column_view const& linestring_points_y,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   return cudf::type_dispatcher(point_x.type(),
                                compute_quadtree_point_to_nearest_linestring{},
@@ -147,7 +148,7 @@ std::unique_ptr<cudf::table> quadtree_point_to_nearest_linestring(
   cudf::column_view const& linestring_offsets,
   cudf::column_view const& linestring_points_x,
   cudf::column_view const& linestring_points_y,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(linestring_quad_pairs.num_columns() == 2,
                     "a quadrant-linestring table must have 2 columns");
diff --git a/cpp/src/nearest_points/point_linestring_nearest_points.cu b/cpp/src/nearest_points/point_linestring_nearest_points.cu
index 9e15cc225..00a4d365f 100644
--- a/cpp/src/nearest_points/point_linestring_nearest_points.cu
+++ b/cpp/src/nearest_points/point_linestring_nearest_points.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -30,6 +30,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
@@ -56,7 +57,7 @@ struct pairwise_point_linestring_nearest_points_impl {
     cudf::device_span<cudf::size_type const> linestring_offsets,
     cudf::column_view linestring_points_xy,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     auto num_points            = static_cast<SizeType>(points_xy.size() / 2);
     auto num_linestring_points = static_cast<SizeType>(linestring_points_xy.size() / 2);
@@ -219,7 +220,7 @@ struct pairwise_point_linestring_nearest_points_functor {
     cudf::device_span<cudf::size_type const> linestring_part_offsets,
     cudf::column_view linestring_points_xy,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     CUSPATIAL_EXPECTS(points_xy.size() % 2 == 0 && linestring_points_xy.size() % 2 == 0,
                       "Points array must contain even number of coordinates.");
@@ -258,7 +259,7 @@ point_linestring_nearest_points_result pairwise_point_linestring_nearest_points(
   std::optional<cudf::device_span<cudf::size_type const>> multilinestring_geometry_offsets,
   cudf::device_span<cudf::size_type const> linestring_part_offsets,
   cudf::column_view linestring_points_xy,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   return double_boolean_dispatch<detail::pairwise_point_linestring_nearest_points_functor>(
     multipoint_geometry_offsets.has_value(),
diff --git a/cpp/src/point_in_polygon/point_in_polygon.cu b/cpp/src/point_in_polygon/point_in_polygon.cu
index e6eb56d08..c22a8da94 100644
--- a/cpp/src/point_in_polygon/point_in_polygon.cu
+++ b/cpp/src/point_in_polygon/point_in_polygon.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -28,6 +28,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 #include <type_traits>
@@ -56,7 +57,7 @@ struct point_in_polygon_functor {
                                            cudf::column_view const& poly_points_y,
                                            bool pairwise,
                                            rmm::cuda_stream_view stream,
-                                           rmm::mr::device_memory_resource* mr)
+                                           rmm::device_async_resource_ref mr)
   {
     auto size = test_points_x.size();
     auto tid  = pairwise ? cudf::type_to_id<uint8_t>() : cudf::type_to_id<int32_t>();
@@ -112,7 +113,7 @@ std::unique_ptr<cudf::column> point_in_polygon(cudf::column_view const& test_poi
                                                cudf::column_view const& poly_points_y,
                                                bool pairwise,
                                                rmm::cuda_stream_view stream,
-                                               rmm::mr::device_memory_resource* mr)
+                                               rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(
     test_points_x.size() == test_points_y.size() and poly_points_x.size() == poly_points_y.size(),
@@ -155,7 +156,7 @@ std::unique_ptr<cudf::column> point_in_polygon(cudf::column_view const& test_poi
                                                cudf::column_view const& poly_ring_offsets,
                                                cudf::column_view const& poly_points_x,
                                                cudf::column_view const& poly_points_y,
-                                               rmm::mr::device_memory_resource* mr)
+                                               rmm::device_async_resource_ref mr)
 {
   return cuspatial::detail::point_in_polygon(test_points_x,
                                              test_points_y,
@@ -174,7 +175,7 @@ std::unique_ptr<cudf::column> pairwise_point_in_polygon(cudf::column_view const&
                                                         cudf::column_view const& poly_ring_offsets,
                                                         cudf::column_view const& poly_points_x,
                                                         cudf::column_view const& poly_points_y,
-                                                        rmm::mr::device_memory_resource* mr)
+                                                        rmm::device_async_resource_ref mr)
 {
   return cuspatial::detail::point_in_polygon(test_points_x,
                                              test_points_y,
diff --git a/cpp/src/points_in_range/points_in_range.cu b/cpp/src/points_in_range/points_in_range.cu
index df6f60c6a..edb843a93 100644
--- a/cpp/src/points_in_range/points_in_range.cu
+++ b/cpp/src/points_in_range/points_in_range.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -25,6 +25,7 @@
 #include <cudf/table/table.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 #include <type_traits>
@@ -42,7 +43,7 @@ struct points_in_range_dispatch {
                                           cudf::column_view const& x,
                                           cudf::column_view const& y,
                                           rmm::cuda_stream_view stream,
-                                          rmm::mr::device_memory_resource* mr)
+                                          rmm::device_async_resource_ref mr)
   {
     auto points_begin = cuspatial::make_vec_2d_iterator(x.begin<T>(), y.begin<T>());
 
@@ -98,7 +99,7 @@ std::unique_ptr<cudf::table> points_in_range(double range_min_x,
                                              cudf::column_view const& x,
                                              cudf::column_view const& y,
                                              rmm::cuda_stream_view stream,
-                                             rmm::mr::device_memory_resource* mr)
+                                             rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(x.type() == y.type(), "Type mismatch between x and y arrays");
   CUSPATIAL_EXPECTS(x.size() == y.size(), "Size mismatch between x and y arrays");
@@ -129,7 +130,7 @@ std::unique_ptr<cudf::table> points_in_range(double range_min_x,
                                              double range_max_y,
                                              cudf::column_view const& x,
                                              cudf::column_view const& y,
-                                             rmm::mr::device_memory_resource* mr)
+                                             rmm::device_async_resource_ref mr)
 {
   return detail::points_in_range(
     range_min_x, range_max_x, range_min_y, range_max_y, x, y, rmm::cuda_stream_default, mr);
diff --git a/cpp/src/projection/sinusoidal_projection.cu b/cpp/src/projection/sinusoidal_projection.cu
index 1be4ade40..2552601a4 100644
--- a/cpp/src/projection/sinusoidal_projection.cu
+++ b/cpp/src/projection/sinusoidal_projection.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,6 +26,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/pair.h>
@@ -53,7 +54,7 @@ struct dispatch_sinusoidal_projection {
     cudf::column_view const& input_lon,
     cudf::column_view const& input_lat,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     auto size = input_lon.size();
     auto type = cudf::data_type{cudf::type_to_id<T>()};
@@ -87,7 +88,7 @@ pair_of_columns sinusoidal_projection(double origin_lon,
                                       cudf::column_view const& input_lon,
                                       cudf::column_view const& input_lat,
                                       rmm::cuda_stream_view stream,
-                                      rmm::mr::device_memory_resource* mr)
+                                      rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(
     origin_lon >= -180 && origin_lon <= 180 && origin_lat >= -90 && origin_lat <= 90,
@@ -116,7 +117,7 @@ pair_of_columns sinusoidal_projection(double origin_lon,
                                       double origin_lat,
                                       cudf::column_view const& input_lon,
                                       cudf::column_view const& input_lat,
-                                      rmm::mr::device_memory_resource* mr)
+                                      rmm::device_async_resource_ref mr)
 {
   return detail::sinusoidal_projection(
     origin_lon, origin_lat, input_lon, input_lat, rmm::cuda_stream_default, mr);
diff --git a/cpp/src/trajectory/derive_trajectories.cu b/cpp/src/trajectory/derive_trajectories.cu
index 8356369ef..e4e6f3745 100644
--- a/cpp/src/trajectory/derive_trajectories.cu
+++ b/cpp/src/trajectory/derive_trajectories.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -25,6 +25,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <memory>
 #include <vector>
@@ -43,7 +44,7 @@ struct derive_trajectories_dispatch {
     cudf::column_view const& y,
     cudf::column_view const& timestamp,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     auto cols = std::vector<std::unique_ptr<cudf::column>>{};
     cols.reserve(4);
@@ -95,7 +96,7 @@ std::pair<std::unique_ptr<cudf::table>, std::unique_ptr<cudf::column>> derive_tr
   cudf::column_view const& y,
   cudf::column_view const& timestamp,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   return cudf::double_type_dispatcher(x.type(),
                                       timestamp.type(),
@@ -114,7 +115,7 @@ std::pair<std::unique_ptr<cudf::table>, std::unique_ptr<cudf::column>> derive_tr
   cudf::column_view const& x,
   cudf::column_view const& y,
   cudf::column_view const& timestamp,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(
     x.size() == y.size() && x.size() == object_id.size() && x.size() == timestamp.size(),
diff --git a/cpp/src/trajectory/trajectory_bounding_boxes.cu b/cpp/src/trajectory/trajectory_bounding_boxes.cu
index 8d441b9c4..baeea23e4 100644
--- a/cpp/src/trajectory/trajectory_bounding_boxes.cu
+++ b/cpp/src/trajectory/trajectory_bounding_boxes.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -27,6 +27,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/zip_iterator.h>
 
@@ -45,7 +46,7 @@ struct dispatch_element {
     cudf::column_view const& x,
     cudf::column_view const& y,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     // Construct output columns
     auto type = cudf::data_type{cudf::type_to_id<T>()};
@@ -101,7 +102,7 @@ std::unique_ptr<cudf::table> trajectory_bounding_boxes(cudf::size_type num_traje
                                                        cudf::column_view const& x,
                                                        cudf::column_view const& y,
                                                        rmm::cuda_stream_view stream,
-                                                       rmm::mr::device_memory_resource* mr)
+                                                       rmm::device_async_resource_ref mr)
 {
   return cudf::type_dispatcher(
     x.type(), dispatch_element{}, num_trajectories, object_id, x, y, stream, mr);
@@ -112,7 +113,7 @@ std::unique_ptr<cudf::table> trajectory_bounding_boxes(cudf::size_type num_traje
                                                        cudf::column_view const& object_id,
                                                        cudf::column_view const& x,
                                                        cudf::column_view const& y,
-                                                       rmm::mr::device_memory_resource* mr)
+                                                       rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(object_id.size() == x.size() && x.size() == y.size(), "Data size mismatch");
   CUSPATIAL_EXPECTS(x.type().id() == y.type().id(), "Data type mismatch");
diff --git a/cpp/src/trajectory/trajectory_distances_and_speeds.cu b/cpp/src/trajectory/trajectory_distances_and_speeds.cu
index a27f8b7f8..89474e925 100644
--- a/cpp/src/trajectory/trajectory_distances_and_speeds.cu
+++ b/cpp/src/trajectory/trajectory_distances_and_speeds.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -29,6 +29,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/adjacent_difference.h>
 #include <thrust/functional.h>
@@ -54,7 +55,7 @@ struct dispatch_timestamp {
     cudf::column_view const& y,
     cudf::column_view const& timestamp,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     // Construct output columns
     auto type = cudf::data_type{cudf::type_to_id<T>()};
@@ -102,7 +103,7 @@ struct dispatch_element {
     cudf::column_view const& y,
     cudf::column_view const& timestamp,
     rmm::cuda_stream_view stream,
-    rmm::mr::device_memory_resource* mr)
+    rmm::device_async_resource_ref mr)
   {
     return cudf::type_dispatcher(timestamp.type(),
                                  dispatch_timestamp<Element>{},
@@ -132,7 +133,7 @@ std::unique_ptr<cudf::table> trajectory_distances_and_speeds(cudf::size_type num
                                                              cudf::column_view const& y,
                                                              cudf::column_view const& timestamp,
                                                              rmm::cuda_stream_view stream,
-                                                             rmm::mr::device_memory_resource* mr)
+                                                             rmm::device_async_resource_ref mr)
 {
   return cudf::type_dispatcher(
     x.type(), dispatch_element{}, num_trajectories, object_id, x, y, timestamp, stream, mr);
@@ -144,7 +145,7 @@ std::unique_ptr<cudf::table> trajectory_distances_and_speeds(cudf::size_type num
                                                              cudf::column_view const& x,
                                                              cudf::column_view const& y,
                                                              cudf::column_view const& timestamp,
-                                                             rmm::mr::device_memory_resource* mr)
+                                                             rmm::device_async_resource_ref mr)
 {
   CUSPATIAL_EXPECTS(
     x.size() == y.size() && x.size() == object_id.size() && x.size() == timestamp.size(),


Test patch:
diff --git a/cpp/include/cuspatial_test/base_fixture.hpp b/cpp/include/cuspatial_test/base_fixture.hpp
index c9927b489..820e54e77 100644
--- a/cpp/include/cuspatial_test/base_fixture.hpp
+++ b/cpp/include/cuspatial_test/base_fixture.hpp
@@ -20,6 +20,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 namespace cuspatial {
 namespace test {
@@ -28,7 +29,7 @@ namespace test {
  * @brief Mixin to supply rmm resources for fixtures
  */
 class RMMResourceMixin {
-  rmm::mr::device_memory_resource* _mr{rmm::mr::get_current_device_resource()};
+  rmm::device_async_resource_ref _mr{rmm::mr::get_current_device_resource()};
   rmm::cuda_stream_view _stream{rmm::cuda_stream_default};
 
  public:
@@ -37,7 +38,7 @@ class RMMResourceMixin {
    * all tests inheriting from this fixture
    * @return pointer to memory resource
    */
-  rmm::mr::device_memory_resource* mr() { return _mr; }
+  rmm::device_async_resource_ref mr() { return _mr; }
 
   /**
    * @brief Returns `cuda_stream_view` that should be used for computation in
diff --git a/cpp/include/cuspatial_test/vector_factories.cuh b/cpp/include/cuspatial_test/vector_factories.cuh
index 139d411c5..a772cd6cf 100644
--- a/cpp/include/cuspatial_test/vector_factories.cuh
+++ b/cpp/include/cuspatial_test/vector_factories.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -28,6 +28,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 #include <rmm/device_vector.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/device_vector.h>
 #include <thrust/host_vector.h>
@@ -56,7 +57,7 @@ auto make_device_vector(std::initializer_list<T> inl)
 template <typename T>
 auto make_device_uvector(std::initializer_list<T> inl,
                          rmm::cuda_stream_view stream,
-                         rmm::mr::device_memory_resource* mr)
+                         rmm::device_async_resource_ref mr)
 {
   std::vector<T> hvec(inl.begin(), inl.end());
   auto res = rmm::device_uvector<T>(inl.size(), stream, mr);
diff --git a/cpp/tests/distance/linestring_polygon_distance_test.cu b/cpp/tests/distance/linestring_polygon_distance_test.cu
index 7cc3b6b55..9a09fd84b 100644
--- a/cpp/tests/distance/linestring_polygon_distance_test.cu
+++ b/cpp/tests/distance/linestring_polygon_distance_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -26,6 +26,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <initializer_list>
 
@@ -35,7 +36,7 @@ using namespace cuspatial::test;
 template <typename T>
 struct PairwiseLinestringPolygonDistanceTest : public BaseFixture {
   rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; }
-  rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); }
+  rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); }
 
   void run_single(std::initializer_list<std::size_t> multilinestring_geometry_offsets,
                   std::initializer_list<std::size_t> multilinestring_part_offsets,
diff --git a/cpp/tests/distance/point_polygon_distance_test.cu b/cpp/tests/distance/point_polygon_distance_test.cu
index c16d324e0..244ac5ee6 100644
--- a/cpp/tests/distance/point_polygon_distance_test.cu
+++ b/cpp/tests/distance/point_polygon_distance_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -25,6 +25,7 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <initializer_list>
 
@@ -36,7 +37,7 @@ double constexpr PI = 3.14159265358979323846;
 template <typename T>
 struct PairwisePointPolygonDistanceTest : public ::testing::Test {
   rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; }
-  rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); }
+  rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); }
 
   void run_single(std::initializer_list<std::initializer_list<vec_2d<T>>> multipoints,
                   std::initializer_list<std::size_t> multipolygon_geometry_offsets,
diff --git a/cpp/tests/intersection/intersection_test_utils.cuh b/cpp/tests/intersection/intersection_test_utils.cuh
index ab0971648..3a2112d07 100644
--- a/cpp/tests/intersection/intersection_test_utils.cuh
+++ b/cpp/tests/intersection/intersection_test_utils.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -22,6 +22,8 @@
 #include <cuspatial/geometry/segment.cuh>
 #include <cuspatial/intersection.cuh>
 
+#include <rmm/resource_ref.hpp>
+
 #include <thrust/copy.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/scatter.h>
@@ -89,7 +91,7 @@ struct order_key_value_pairs {
 template <typename T, typename IndexType, typename type_t>
 linestring_intersection_result<T, IndexType> segment_sort_intersection_result(
   linestring_intersection_result<T, IndexType>& result,
-  rmm::mr::device_memory_resource* mr,
+  rmm::device_async_resource_ref mr,
   rmm::cuda_stream_view stream)
 {
   auto const num_points   = result.points_coords->size();
@@ -189,7 +191,7 @@ auto make_linestring_intersection_result(
   std::initializer_list<IndexType> rhs_linestring_ids,
   std::initializer_list<IndexType> rhs_segment_ids,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+  rmm::device_async_resource_ref mr)
 {
   auto d_geometry_collection_offset =
     make_device_uvector<IndexType>(geometry_collection_offset, stream, mr);
diff --git a/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu b/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu
index eb36d5e35..760ab72a4 100644
--- a/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu
+++ b/cpp/tests/intersection/linestring_intersection_intermediates_remove_if_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022-2023, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2024, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -27,6 +27,7 @@
 #include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/iterator/zip_iterator.h>
 
@@ -44,7 +45,7 @@ auto make_linestring_intersection_intermediates(std::initializer_list<IndexType>
                                                 std::initializer_list<IndexType> rhs_linestring_ids,
                                                 std::initializer_list<IndexType> rhs_segment_ids,
                                                 rmm::cuda_stream_view stream,
-                                                rmm::mr::device_memory_resource* mr)
+                                                rmm::device_async_resource_ref mr)
 {
   auto d_offset             = make_device_uvector<IndexType>(offsets, stream, mr);
   auto d_geoms              = make_device_uvector<Geom>(geoms, stream, mr);
@@ -66,7 +67,7 @@ auto make_linestring_intersection_intermediates(std::initializer_list<IndexType>
 template <typename T>
 struct LinestringIntersectionIntermediatesRemoveIfTest : public ::testing::Test {
   rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; }
-  rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); }
+  rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); }
 
   template <typename FlagType, typename IntermediateType>
   void run_single(IntermediateType& intermediates,
diff --git a/cpp/tests/intersection/linestring_intersection_test.cpp b/cpp/tests/intersection/linestring_intersection_test.cpp
index 9b9318104..d91d99ddb 100644
--- a/cpp/tests/intersection/linestring_intersection_test.cpp
+++ b/cpp/tests/intersection/linestring_intersection_test.cpp
@@ -37,6 +37,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <initializer_list>
 #include <memory>
@@ -116,7 +117,7 @@ std::pair<collection_type_id, std::unique_ptr<cudf::column>> make_linestring_col
 
 struct LinestringIntersectionTestBase : public BaseFixture {
   rmm::cuda_stream_view stream{rmm::cuda_stream_default};
-  rmm::mr::device_memory_resource* mr{rmm::mr::get_current_device_resource()};
+  rmm::device_async_resource_ref mr{rmm::mr::get_current_device_resource()};
 };
 
 template <typename T>
diff --git a/cpp/tests/intersection/linestring_intersection_test.cu b/cpp/tests/intersection/linestring_intersection_test.cu
index 7f758821c..d6ec4affd 100644
--- a/cpp/tests/intersection/linestring_intersection_test.cu
+++ b/cpp/tests/intersection/linestring_intersection_test.cu
@@ -30,6 +30,7 @@
 
 #include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/binary_search.h>
 #include <thrust/sort.h>
@@ -43,7 +44,7 @@ using namespace cuspatial::test;
 template <typename T>
 struct LinestringIntersectionTest : public ::testing::Test {
   rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; }
-  rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); }
+  rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); }
 
   template <typename IndexType, typename MultiLinestringRange, typename IntersectionResult>
   void run_single_test(MultiLinestringRange lhs,
diff --git a/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu b/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu
index 359b45700..720eb2efa 100644
--- a/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu
+++ b/cpp/tests/intersection/linestring_intersection_with_duplicates_test.cu
@@ -31,6 +31,7 @@
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <cub/device/device_segmented_sort.cuh>
 #include <thrust/gather.h>
@@ -72,7 +73,7 @@ using namespace cuspatial::test;
 template <typename Intermediate>
 Intermediate segmented_sort_intersection_intermediates(Intermediate& intermediate,
                                                        rmm::cuda_stream_view stream,
-                                                       rmm::mr::device_memory_resource* mr)
+                                                       rmm::device_async_resource_ref mr)
 {
   using GeomType  = typename Intermediate::geometry_t;
   using IndexType = typename Intermediate::index_t;
@@ -126,7 +127,7 @@ Intermediate segmented_sort_intersection_intermediates(Intermediate& intermediat
 template <typename T>
 struct LinestringIntersectionDuplicatesTest : public ::testing::Test {
   rmm::cuda_stream_view stream() { return rmm::cuda_stream_default; }
-  rmm::mr::device_memory_resource* mr() { return rmm::mr::get_current_device_resource(); }
+  rmm::device_async_resource_ref mr() { return rmm::mr::get_current_device_resource(); }
 
   template <typename IndexType, typename MultilinestringRange1, typename MultilinestringRange2>
   void run_single(MultilinestringRange1 lhs,
diff --git a/cpp/tests/join/quadtree_point_in_polygon_test_large.cu b/cpp/tests/join/quadtree_point_in_polygon_test_large.cu
index d80709486..e1f0c86f6 100644
--- a/cpp/tests/join/quadtree_point_in_polygon_test_large.cu
+++ b/cpp/tests/join/quadtree_point_in_polygon_test_large.cu
@@ -31,6 +31,7 @@
 #include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
 #include <rmm/mr/device/per_device_resource.hpp>
+#include <rmm/resource_ref.hpp>
 
 #include <thrust/gather.h>
 #include <thrust/sort.h>
@@ -54,7 +55,7 @@ inline auto generate_points(
   uint32_t points_per_quad,
   std::size_t seed,
   rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+  rmm::device_async_resource_ref mr = rmm::mr::get_current_device_resource())
 {
   auto engine     = cuspatial::test::deterministic_engine(0);
   auto uniform    = cuspatial::test::make_normal_dist<T>(0.0, 1.0);