diff --git a/CHANGELOG.md b/CHANGELOG.md
index 3fedbd2da..4a0ef8b32 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -7,6 +7,8 @@
 
 - PR #332 fix directed_hausdorff_distance's space_offsets name + documentation
 
+- PR #331 Use simplified `rmm::exec_policy`
+
 ## Bug Fixes
 
 # cuSpatial 0.17.0 (Date TBD)
diff --git a/cpp/src/indexing/construction/detail/phase_1.cuh b/cpp/src/indexing/construction/detail/phase_1.cuh
index aeba1dc5a..00be5e594 100644
--- a/cpp/src/indexing/construction/detail/phase_1.cuh
+++ b/cpp/src/indexing/construction/detail/phase_1.cuh
@@ -27,6 +27,7 @@
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/functional.h>
 #include <thrust/iterator/constant_iterator.h>
@@ -63,7 +64,7 @@ compute_point_keys_and_sorted_indices(cudf::column_view const &x,
                                       rmm::mr::device_memory_resource *mr)
 {
   rmm::device_uvector<uint32_t> keys(x.size(), stream);
-  thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::transform(rmm::exec_policy(stream),
                     make_zip_iterator(x.begin<T>(), y.begin<T>()),
                     make_zip_iterator(x.begin<T>(), y.begin<T>()) + x.size(),
                     keys.begin(),
@@ -79,15 +80,13 @@ compute_point_keys_and_sorted_indices(cudf::column_view const &x,
 
   auto indices = make_fixed_width_column<uint32_t>(keys.size(), stream, mr);
 
-  thrust::sequence(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::sequence(rmm::exec_policy(stream),
                    indices->mutable_view().begin<uint32_t>(),
                    indices->mutable_view().end<uint32_t>());
 
   // Sort the codes and point indices
-  thrust::stable_sort_by_key(rmm::exec_policy(stream)->on(stream.value()),
-                             keys.begin(),
-                             keys.end(),
-                             indices->mutable_view().begin<int32_t>());
+  thrust::stable_sort_by_key(
+    rmm::exec_policy(stream), keys.begin(), keys.end(), indices->mutable_view().begin<int32_t>());
 
   return std::make_pair(std::move(keys), std::move(indices));
 }
@@ -110,7 +109,7 @@ inline cudf::size_type build_tree_level(InputIterator1 keys_begin,
                                         BinaryOp binary_op,
                                         rmm::cuda_stream_view stream)
 {
-  auto result = thrust::reduce_by_key(rmm::exec_policy(stream)->on(stream.value()),
+  auto result = thrust::reduce_by_key(rmm::exec_policy(stream),
                                       keys_begin,
                                       keys_end,
                                       vals_in,
@@ -207,19 +206,19 @@ reverse_tree_levels(rmm::device_uvector<uint32_t> const &quad_keys_in,
     cudf::size_type level_end   = end_pos[level];
     cudf::size_type level_begin = begin_pos[level];
     cudf::size_type num_quads   = level_end - level_begin;
-    thrust::fill(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::fill(rmm::exec_policy(stream),
                  quad_levels.begin() + offset,
                  quad_levels.begin() + offset + num_quads,
                  level);
-    thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::copy(rmm::exec_policy(stream),
                  quad_keys_in.begin() + level_begin,
                  quad_keys_in.begin() + level_end,
                  quad_keys.begin() + offset);
-    thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::copy(rmm::exec_policy(stream),
                  quad_point_count_in.begin() + level_begin,
                  quad_point_count_in.begin() + level_end,
                  quad_point_count.begin() + offset);
-    thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::copy(rmm::exec_policy(stream),
                  quad_child_count_in.begin() + level_begin,
                  quad_child_count_in.begin() + level_end,
                  quad_child_count.begin() + offset);
@@ -293,10 +292,7 @@ inline auto make_full_levels(cudf::column_view const &x,
   quad_child_count.resize(num_bottom_quads * (max_depth + 1), stream);
 
   // Zero out the quad_child_count vector because we're reusing the point_keys vector
-  thrust::fill(rmm::exec_policy(stream)->on(stream.value()),
-               quad_child_count.begin(),
-               quad_child_count.end(),
-               0);
+  thrust::fill(rmm::exec_policy(stream), quad_child_count.begin(), quad_child_count.end(), 0);
 
   //
   // Compute "full" quads for the tree at each level. Starting from the quadrant
diff --git a/cpp/src/indexing/construction/detail/phase_2.cuh b/cpp/src/indexing/construction/detail/phase_2.cuh
index bd3dfc233..62002e44a 100644
--- a/cpp/src/indexing/construction/detail/phase_2.cuh
+++ b/cpp/src/indexing/construction/detail/phase_2.cuh
@@ -25,6 +25,7 @@
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
@@ -47,7 +48,7 @@ inline rmm::device_uvector<uint32_t> compute_leaf_positions(cudf::column_view co
                                                             rmm::cuda_stream_view stream)
 {
   rmm::device_uvector<uint32_t> leaf_pos(num_valid_nodes, stream);
-  auto result = thrust::copy_if(rmm::exec_policy(stream)->on(stream.value()),
+  auto result = thrust::copy_if(rmm::exec_policy(stream),
                                 thrust::make_counting_iterator(0),
                                 thrust::make_counting_iterator(0) + num_valid_nodes,
                                 indicator.begin<bool>(),
@@ -70,7 +71,7 @@ inline rmm::device_uvector<uint32_t> flatten_point_keys(
   rmm::device_uvector<uint32_t> flattened_keys(num_valid_nodes, stream);
   auto keys_and_levels =
     make_zip_iterator(quad_keys.begin(), quad_level.begin(), indicator.begin<bool>());
-  thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::transform(rmm::exec_policy(stream),
                     keys_and_levels,
                     keys_and_levels + num_valid_nodes,
                     flattened_keys.begin(),
@@ -107,24 +108,23 @@ inline rmm::device_uvector<uint32_t> compute_flattened_first_point_positions(
       flatten_point_keys(quad_keys, quad_level, indicator, num_valid_nodes, max_depth, stream);
 
     rmm::device_uvector<uint32_t> initial_sort_indices(num_valid_nodes, stream);
-    thrust::sequence(rmm::exec_policy(stream)->on(stream.value()),
-                     initial_sort_indices.begin(),
-                     initial_sort_indices.end());
+    thrust::sequence(
+      rmm::exec_policy(stream), initial_sort_indices.begin(), initial_sort_indices.end());
 
     rmm::device_uvector<uint32_t> quad_point_count_tmp(num_valid_nodes, stream);
-    thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::copy(rmm::exec_policy(stream),
                  quad_point_count.begin(),
                  quad_point_count.end(),
                  quad_point_count_tmp.begin());
 
     // sort indices and temporary point counts
     thrust::stable_sort_by_key(
-      rmm::exec_policy(stream)->on(stream.value()),
+      rmm::exec_policy(stream),
       flattened_keys.begin(),
       flattened_keys.end(),
       make_zip_iterator(initial_sort_indices.begin(), quad_point_count_tmp.begin()));
 
-    thrust::remove_if(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::remove_if(rmm::exec_policy(stream),
                       quad_point_count_tmp.begin(),
                       quad_point_count_tmp.begin() + num_valid_nodes,
                       quad_point_count_tmp.begin(),
@@ -151,7 +151,7 @@ inline rmm::device_uvector<uint32_t> compute_flattened_first_point_positions(
 
   rmm::device_uvector<uint32_t> quad_point_offsets_tmp(leaf_offsets.size(), stream);
 
-  thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::exclusive_scan(rmm::exec_policy(stream),
                          quad_point_count_tmp.begin(),
                          quad_point_count_tmp.end(),
                          quad_point_offsets_tmp.begin());
@@ -159,14 +159,14 @@ inline rmm::device_uvector<uint32_t> compute_flattened_first_point_positions(
   auto counts_and_offsets =
     make_zip_iterator(quad_point_count_tmp.begin(), quad_point_offsets_tmp.begin());
 
-  thrust::stable_sort_by_key(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::stable_sort_by_key(rmm::exec_policy(stream),
                              initial_sort_indices.begin(),
                              initial_sort_indices.end(),
                              counts_and_offsets);
 
   rmm::device_uvector<uint32_t> quad_point_offsets(num_valid_nodes, stream);
 
-  thrust::scatter(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::scatter(rmm::exec_policy(stream),
                   counts_and_offsets,
                   counts_and_offsets + leaf_offsets.size(),
                   leaf_offsets.begin(),
@@ -188,15 +188,14 @@ inline rmm::device_uvector<uint32_t> compute_parent_positions(
   auto parent_pos = [&]() {
     rmm::device_uvector<uint32_t> position_map(num_parent_nodes, stream);
     // line 1 of algorithm in Fig. 5 in ref.
-    thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::exclusive_scan(rmm::exec_policy(stream),
                            quad_child_count.begin(),
                            quad_child_count.begin() + num_parent_nodes,
                            position_map.begin());
     // line 2 of algorithm in Fig. 5 in ref.
     rmm::device_uvector<uint32_t> parent_pos(num_child_nodes, stream);
-    thrust::uninitialized_fill(
-      rmm::exec_policy(stream)->on(stream.value()), parent_pos.begin(), parent_pos.end(), 0);
-    thrust::scatter(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::uninitialized_fill(rmm::exec_policy(stream), parent_pos.begin(), parent_pos.end(), 0);
+    thrust::scatter(rmm::exec_policy(stream),
                     thrust::make_counting_iterator(0),
                     thrust::make_counting_iterator(0) + num_parent_nodes,
                     position_map.begin(),
@@ -205,7 +204,7 @@ inline rmm::device_uvector<uint32_t> compute_parent_positions(
   }();
 
   // line 3 of algorithm in Fig. 5 in ref.
-  thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::inclusive_scan(rmm::exec_policy(stream),
                          parent_pos.begin(),
                          parent_pos.begin() + num_child_nodes,
                          parent_pos.begin(),
@@ -249,7 +248,7 @@ inline std::pair<uint32_t, uint32_t> remove_unqualified_quads(
   // Start counting nodes at level 2, since children of the root node should not
   // be discarded.
   auto num_invalid_parent_nodes =
-    thrust::count_if(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::count_if(rmm::exec_policy(stream),
                      parent_point_counts,
                      parent_point_counts + (num_parent_nodes - level_1_size),
                      // i.e. quad_point_count[parent_pos] <= min_size
@@ -268,7 +267,7 @@ inline std::pair<uint32_t, uint32_t> remove_unqualified_quads(
                                 quad_levels.begin() + level_1_size);
 
   auto last_valid =
-    thrust::remove_if(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::remove_if(rmm::exec_policy(stream),
                       tree,
                       tree + num_child_nodes,
                       parent_point_counts,
@@ -314,14 +313,14 @@ inline std::unique_ptr<cudf::column> construct_non_leaf_indicator(
   auto is_quad = make_fixed_width_column<bool>(num_valid_nodes, stream, mr);
 
   // line 6 of algorithm in Fig. 5 in ref.
-  thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::transform(rmm::exec_policy(stream),
                     quad_point_count.begin(),
                     quad_point_count.begin() + num_parent_nodes,
                     is_quad->mutable_view().begin<bool>(),
                     thrust::placeholders::_1 > min_size);
 
   // line 7 of algorithm in Fig. 5 in ref.
-  thrust::replace_if(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::replace_if(rmm::exec_policy(stream),
                      quad_point_count.begin(),
                      quad_point_count.begin() + num_parent_nodes,
                      is_quad->view().begin<bool>(),
@@ -331,7 +330,7 @@ inline std::unique_ptr<cudf::column> construct_non_leaf_indicator(
   if (num_valid_nodes > num_parent_nodes) {
     // zero-fill the rest of the indicator column because
     // device_memory_resources aren't required to initialize allocations
-    thrust::fill(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::fill(rmm::exec_policy(stream),
                  is_quad->mutable_view().begin<bool>() + num_parent_nodes,
                  is_quad->mutable_view().end<bool>(),
                  0);
diff --git a/cpp/src/indexing/construction/point_quadtree.cu b/cpp/src/indexing/construction/point_quadtree.cu
index 2d8f2cccb..5bf86c2b0 100644
--- a/cpp/src/indexing/construction/point_quadtree.cu
+++ b/cpp/src/indexing/construction/point_quadtree.cu
@@ -25,6 +25,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 /*
  * quadtree indexing on points using the bottom-up algorithm described at ref.
@@ -54,7 +55,7 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
                                                    rmm::mr::device_memory_resource *mr)
 {
   // count the number of child nodes
-  auto num_child_nodes = thrust::reduce(rmm::exec_policy(stream)->on(stream.value()),
+  auto num_child_nodes = thrust::reduce(rmm::exec_policy(stream),
                                         quad_child_count.begin(),
                                         quad_child_count.begin() + num_parent_nodes);
 
@@ -91,7 +92,7 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
 
     auto quad_child_pos = make_fixed_width_column<uint32_t>(num_valid_nodes, stream, mr);
     // line 9 of algorithm in Fig. 5 in ref.
-    thrust::replace_if(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::replace_if(rmm::exec_policy(stream),
                        quad_child_count.begin(),
                        quad_child_count.begin() + num_valid_nodes,
                        is_quad->view().begin<uint8_t>(),
@@ -99,7 +100,7 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
                        0);
 
     // line 10 of algorithm in Fig. 5 in ref.
-    thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::exclusive_scan(rmm::exec_policy(stream),
                            quad_child_count.begin(),
                            quad_child_count.end(),
                            quad_child_pos->mutable_view().begin<uint32_t>(),
@@ -112,7 +113,7 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
 
     // for each value in `is_quad` copy from `quad_child_pos` if true, else
     // `quad_point_pos`
-    thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::transform(rmm::exec_policy(stream),
                       offsets_iter,
                       offsets_iter + num_valid_nodes,
                       offsets->mutable_view().template begin<uint32_t>(),
@@ -131,7 +132,7 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
   auto lengths_iter = make_zip_iterator(is_quad->view().begin<bool>(),  //
                                         quad_child_count.begin(),
                                         quad_point_count.begin());
-  thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::transform(rmm::exec_policy(stream),
                     lengths_iter,
                     lengths_iter + num_valid_nodes,
                     lengths->mutable_view().template begin<uint32_t>(),
@@ -144,7 +145,7 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
   auto keys = make_fixed_width_column<uint32_t>(num_valid_nodes, stream, mr);
 
   // Copy quad keys to keys output column
-  thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::copy(rmm::exec_policy(stream),
                quad_keys.begin(),
                quad_keys.end(),
                keys->mutable_view().begin<uint32_t>());
@@ -153,7 +154,7 @@ inline std::unique_ptr<cudf::table> make_quad_tree(rmm::device_uvector<uint32_t>
   auto levels = make_fixed_width_column<uint8_t>(num_valid_nodes, stream, mr);
 
   // Copy quad levels to levels output column
-  thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::copy(rmm::exec_policy(stream),
                quad_levels.begin(),
                quad_levels.end(),
                levels->mutable_view().begin<uint8_t>());
@@ -185,31 +186,31 @@ inline std::unique_ptr<cudf::table> make_leaf_tree(
   auto offsets = make_fixed_width_column<uint32_t>(num_top_quads, stream, mr);
 
   // copy quad keys from the front of the quad_keys list
-  thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::copy(rmm::exec_policy(stream),
                quad_keys.begin(),
                quad_keys.begin() + num_top_quads,
                keys->mutable_view().begin<uint32_t>());
 
   // copy point counts from the front of the quad_point_count list
-  thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::copy(rmm::exec_policy(stream),
                quad_point_count.begin(),
                quad_point_count.begin() + num_top_quads,
                lengths->mutable_view().begin<uint32_t>());
 
   // All leaves are children of the root node (level 0)
-  thrust::fill(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::fill(rmm::exec_policy(stream),
                levels->mutable_view().begin<uint8_t>(),
                levels->mutable_view().end<uint8_t>(),
                0);
 
   // Quad node indicators are false for leaf nodes
-  thrust::fill(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::fill(rmm::exec_policy(stream),
                is_quad->mutable_view().begin<bool>(),
                is_quad->mutable_view().end<bool>(),
                false);
 
   // compute offsets from lengths
-  thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::exclusive_scan(rmm::exec_policy(stream),
                          lengths->view().begin<uint32_t>(),
                          lengths->view().end<uint32_t>(),
                          offsets->mutable_view().begin<uint32_t>());
diff --git a/cpp/src/interpolate/cubic_spline.cu b/cpp/src/interpolate/cubic_spline.cu
index 500d5c2d6..e6f7577e2 100644
--- a/cpp/src/interpolate/cubic_spline.cu
+++ b/cpp/src/interpolate/cubic_spline.cu
@@ -24,6 +24,7 @@
 #include <cuspatial/error.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <cusparse.h>
 
@@ -48,7 +49,7 @@ struct parallel_search {
       curve_ids.type(), t.size(), cudf::mask_state::UNALLOCATED, stream, mr);
     int32_t* p_result = result->mutable_view().data<int32_t>();
     thrust::for_each(
-      rmm::exec_policy(stream)->on(stream.value()),
+      rmm::exec_policy(stream),
       thrust::make_counting_iterator<int>(0),
       thrust::make_counting_iterator<int>(query_coords.size()),
       [p_t, p_curve_ids, p_prefixes, p_query_coords, p_result] __device__(int index) {
@@ -103,7 +104,7 @@ struct interpolate {
       cudf::make_numeric_column(t.type(), t.size(), cudf::mask_state::UNALLOCATED, stream, mr);
     T* p_result = result->mutable_view().data<T>();
     thrust::for_each(
-      rmm::exec_policy(stream)->on(stream.value()),
+      rmm::exec_policy(stream),
       thrust::make_counting_iterator<int>(0),
       thrust::make_counting_iterator<int>(t.size()),
       [p_t, p_ids, p_coef_indices, p_d3, p_d2, p_d1, p_d0, p_result] __device__(int index) {
@@ -150,7 +151,7 @@ struct coefficients_compute {
     T* p_d1                   = d1.data<T>();
     T* p_d0                   = d0.data<T>();
     thrust::for_each(
-      rmm::exec_policy(stream)->on(stream.value()),
+      rmm::exec_policy(stream),
       thrust::make_counting_iterator<int>(1),
       thrust::make_counting_iterator<int>(prefixes.size()),
       [p_t, p_y, p_prefixes, p_h, p_i, p_z, p_d3, p_d2, p_d1, p_d0] __device__(int index) {
@@ -214,7 +215,7 @@ struct compute_spline_tridiagonals {
     T* p_u                    = u.data<T>();
     T* p_h                    = h.data<T>();
     T* p_i                    = i.data<T>();
-    thrust::for_each(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::for_each(rmm::exec_policy(stream),
                      thrust::make_counting_iterator<int>(1),
                      thrust::make_counting_iterator<int>(prefixes.size()),
                      [p_t, p_y, p_prefixes, p_d, p_dlu, p_u, p_h, p_i] __device__(int index) {
diff --git a/cpp/src/io/shp/polygon_shapefile_reader.cu b/cpp/src/io/shp/polygon_shapefile_reader.cu
index 4c2cd8697..b645bbb24 100644
--- a/cpp/src/io/shp/polygon_shapefile_reader.cu
+++ b/cpp/src/io/shp/polygon_shapefile_reader.cu
@@ -20,9 +20,9 @@
 #include <cudf/types.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 
-#include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_buffer.hpp>
+#include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
 
 #include <thrust/scan.h>
@@ -63,13 +63,13 @@ std::vector<std::unique_ptr<cudf::column>> read_polygon_shapefile(
   auto ys              = make_column<double>(std::get<3>(poly_vectors), stream, mr);
 
   // transform polygon lengths to polygon offsets
-  thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::exclusive_scan(rmm::exec_policy(stream),
                          polygon_offsets->view().begin<cudf::size_type>(),
                          polygon_offsets->view().end<cudf::size_type>(),
                          polygon_offsets->mutable_view().begin<cudf::size_type>());
 
   // transform ring lengths to ring offsets
-  thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::exclusive_scan(rmm::exec_policy(stream),
                          ring_offsets->view().begin<cudf::size_type>(),
                          ring_offsets->view().end<cudf::size_type>(),
                          ring_offsets->mutable_view().begin<cudf::size_type>());
diff --git a/cpp/src/join/detail/intersection.cuh b/cpp/src/join/detail/intersection.cuh
index cfc61a211..efdac370a 100644
--- a/cpp/src/join/detail/intersection.cuh
+++ b/cpp/src/join/detail/intersection.cuh
@@ -25,6 +25,7 @@
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/copy.h>
 #include <thrust/iterator/transform_iterator.h>
@@ -48,11 +49,10 @@ inline cudf::size_type copy_leaf_intersections(InputIterator input_begin,
 {
   return thrust::distance(
     output_begin,
-    thrust::copy_if(rmm::exec_policy(stream)->on(stream.value()),
-                    input_begin,
-                    input_end,
-                    output_begin,
-                    [] __device__(auto const &t) { return thrust::get<0>(t) == leaf_indicator; }));
+    thrust::copy_if(
+      rmm::exec_policy(stream), input_begin, input_end, output_begin, [] __device__(auto const &t) {
+        return thrust::get<0>(t) == leaf_indicator;
+      }));
 }
 
 template <typename InputIterator, typename OutputIterator>
@@ -61,14 +61,12 @@ inline cudf::size_type remove_non_quad_intersections(InputIterator input_begin,
                                                      OutputIterator output_begin,
                                                      rmm::cuda_stream_view stream)
 {
-  return thrust::distance(output_begin,
-                          thrust::remove_if(rmm::exec_policy(stream)->on(stream.value()),
-                                            input_begin,
-                                            input_end,
-                                            output_begin,
-                                            [] __device__(auto const &t) {
-                                              return thrust::get<0>(t) != quad_indicator;
-                                            }));
+  return thrust::distance(
+    output_begin,
+    thrust::remove_if(
+      rmm::exec_policy(stream), input_begin, input_end, output_begin, [] __device__(auto const &t) {
+        return thrust::get<0>(t) != quad_indicator;
+      }));
 }
 
 template <typename T,
@@ -98,7 +96,7 @@ inline std::pair<cudf::size_type, cudf::size_type> find_intersections(
   auto d_poly_x_max = cudf::column_device_view::create(poly_bbox.column(2), stream);
   auto d_poly_y_max = cudf::column_device_view::create(poly_bbox.column(3), stream);
 
-  thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::transform(rmm::exec_policy(stream),
                     make_zip_iterator(node_indices, poly_indices),
                     make_zip_iterator(node_indices, poly_indices) + num_pairs,
                     node_pairs,
diff --git a/cpp/src/join/detail/traversal.cuh b/cpp/src/join/detail/traversal.cuh
index 5dc0ba9b6..a88b9b0de 100644
--- a/cpp/src/join/detail/traversal.cuh
+++ b/cpp/src/join/detail/traversal.cuh
@@ -22,6 +22,7 @@
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/gather.h>
 #include <thrust/scan.h>
@@ -68,10 +69,8 @@ descend_quadtree(LengthsIter counts,
   // scan on the number of child nodes to compute the offsets
   // note: size is `num_quads + 1` so the last element is `num_children`
   rmm::device_uvector<uint32_t> parent_offsets(num_quads + 1, stream);
-  thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
-                         parent_counts,
-                         parent_counts + num_quads,
-                         parent_offsets.begin() + 1);
+  thrust::inclusive_scan(
+    rmm::exec_policy(stream), parent_counts, parent_counts + num_quads, parent_offsets.begin() + 1);
 
   parent_offsets.set_element_async(0, 0, stream);
 
@@ -79,10 +78,9 @@ descend_quadtree(LengthsIter counts,
 
   rmm::device_uvector<uint32_t> parent_indices(num_children, stream);
   // fill with zeroes
-  thrust::fill(
-    rmm::exec_policy(stream)->on(stream.value()), parent_indices.begin(), parent_indices.end(), 0);
+  thrust::fill(rmm::exec_policy(stream), parent_indices.begin(), parent_indices.end(), 0);
   // use the parent_offsets as the map to scatter sequential parent_indices
-  thrust::scatter(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::scatter(rmm::exec_policy(stream),
                   thrust::make_counting_iterator(0),
                   thrust::make_counting_iterator(0) + num_quads,
                   parent_offsets.begin(),
@@ -90,7 +88,7 @@ descend_quadtree(LengthsIter counts,
 
   // inclusive scan with maximum functor to fill the empty elements with their left-most non-empty
   // elements. `parent_indices` is now a full array of the sequence index of each quadrant's parent
-  thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::inclusive_scan(rmm::exec_policy(stream),
                          parent_indices.begin(),
                          parent_indices.begin() + num_children,
                          parent_indices.begin(),
@@ -103,7 +101,7 @@ descend_quadtree(LengthsIter counts,
   rmm::device_uvector<uint32_t> child_poly_indices(num_children, stream);
 
   // `parent_indices` is a gather map to retrieve non-leaf quads' respective child nodes
-  thrust::gather(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::gather(rmm::exec_policy(stream),
                  parent_indices.begin(),
                  parent_indices.begin() + num_children,
                  // curr level iterator
@@ -118,7 +116,7 @@ descend_quadtree(LengthsIter counts,
                                    child_poly_indices.begin()));
 
   rmm::device_uvector<uint32_t> relative_child_offsets(num_children, stream);
-  thrust::exclusive_scan_by_key(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::exclusive_scan_by_key(rmm::exec_policy(stream),
                                 parent_indices.begin(),
                                 parent_indices.begin() + num_children,
                                 thrust::constant_iterator<uint32_t>(1),
@@ -126,7 +124,7 @@ descend_quadtree(LengthsIter counts,
 
   // compute child quad indices using parent and relative child offsets
   auto child_offsets_iter = thrust::make_permutation_iterator(offsets, child_quad_indices.begin());
-  thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::transform(rmm::exec_policy(stream),
                     child_offsets_iter,
                     child_offsets_iter + num_children,
                     relative_child_offsets.begin(),
diff --git a/cpp/src/join/quadtree_point_in_polygon.cu b/cpp/src/join/quadtree_point_in_polygon.cu
index e6a7b0323..4301595c4 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_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/binary_search.h>
 #include <thrust/copy.h>
@@ -127,7 +128,7 @@ struct compute_quadtree_point_in_polygon {
     // `inclusive_scan` is the total number of points to be tested against any polygon.
     rmm::device_uvector<uint32_t> local_point_offsets(num_poly_quad_pairs + 1, stream);
 
-    thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::inclusive_scan(rmm::exec_policy(stream),
                            quad_lengths_iter,
                            quad_lengths_iter + num_poly_quad_pairs,
                            local_point_offsets.begin() + 1);
@@ -155,7 +156,7 @@ struct compute_quadtree_point_in_polygon {
     //     pp_pairs.append((polygon, point))
     // ```
     //
-    thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::transform(rmm::exec_policy(stream),
                       counting_iter,
                       counting_iter + num_total_points,
                       poly_and_point_indices,
@@ -172,7 +173,7 @@ struct compute_quadtree_point_in_polygon {
     // Compute the number of intersections by removing (poly, point) pairs that don't intersect
     auto num_intersections = thrust::distance(
       poly_and_point_indices,
-      thrust::remove_if(rmm::exec_policy(stream)->on(stream.value()),
+      thrust::remove_if(rmm::exec_policy(stream),
                         poly_and_point_indices,
                         poly_and_point_indices + num_total_points,
                         test_poly_point_intersection<T, decltype(point_xys_iter)>{
@@ -190,13 +191,13 @@ struct compute_quadtree_point_in_polygon {
     // `idxs.begin() + num_intersections`.
 
     // populate the polygon indices column
-    thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::copy(rmm::exec_policy(stream),
                  poly_idxs.begin(),
                  poly_idxs.begin() + num_intersections,
                  poly_idx_col->mutable_view().template begin<uint32_t>());
 
     // populate the point indices column
-    thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::copy(rmm::exec_policy(stream),
                  point_idxs.begin(),
                  point_idxs.begin() + num_intersections,
                  point_idx_col->mutable_view().template begin<uint32_t>());
diff --git a/cpp/src/join/quadtree_point_to_nearest_polyline.cu b/cpp/src/join/quadtree_point_to_nearest_polyline.cu
index 3cf688ef0..8d9c025b9 100644
--- a/cpp/src/join/quadtree_point_to_nearest_polyline.cu
+++ b/cpp/src/join/quadtree_point_to_nearest_polyline.cu
@@ -28,6 +28,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/gather.h>
 #include <thrust/iterator/constant_iterator.h>
@@ -152,23 +153,21 @@ struct compute_quadtree_point_to_nearest_polyline {
 
     rmm::device_uvector<uint32_t> d_poly_idx(num_pairs, stream);
 
-    thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::copy(rmm::exec_policy(stream),
                  poly_idx.begin<uint32_t>(),
                  poly_idx.end<uint32_t>(),
                  d_poly_idx.begin());
 
     rmm::device_uvector<uint32_t> d_quad_idx(num_pairs, stream);
 
-    thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::copy(rmm::exec_policy(stream),
                  quad_idx.begin<uint32_t>(),
                  quad_idx.end<uint32_t>(),
                  d_quad_idx.begin());
 
     // sort (d_poly_idx, d_quad_idx) using d_quad_idx as key => (quad_idxs, poly_idxs)
-    thrust::sort_by_key(rmm::exec_policy(stream)->on(stream.value()),
-                        d_quad_idx.begin(),
-                        d_quad_idx.end(),
-                        d_poly_idx.begin());
+    thrust::sort_by_key(
+      rmm::exec_policy(stream), d_quad_idx.begin(), d_quad_idx.end(), d_poly_idx.begin());
 
     // reduce_by_key using d_quad_idx as the key
     // exclusive_scan on numbers of polygons associated with a quadrant to create d_poly_idx_offsets
@@ -177,7 +176,7 @@ struct compute_quadtree_point_to_nearest_polyline {
 
     uint32_t num_quads =
       thrust::distance(d_poly_idx_offsets.begin(),
-                       thrust::reduce_by_key(rmm::exec_policy(stream)->on(stream.value()),
+                       thrust::reduce_by_key(rmm::exec_policy(stream),
                                              d_quad_idx.begin(),
                                              d_quad_idx.end(),
                                              thrust::constant_iterator<uint32_t>(1),
@@ -188,7 +187,7 @@ struct compute_quadtree_point_to_nearest_polyline {
     d_quad_idx.resize(num_quads, stream);
     d_poly_idx_offsets.resize(num_quads, stream);
 
-    thrust::exclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::exclusive_scan(rmm::exec_policy(stream),
                            d_poly_idx_offsets.begin(),
                            d_poly_idx_offsets.end(),
                            d_poly_idx_offsets.begin());
diff --git a/cpp/src/join/quadtree_poly_filtering.cu b/cpp/src/join/quadtree_poly_filtering.cu
index 98f0a5674..74a86a662 100644
--- a/cpp/src/join/quadtree_poly_filtering.cu
+++ b/cpp/src/join/quadtree_poly_filtering.cu
@@ -28,6 +28,7 @@
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <tuple>
 
@@ -61,7 +62,7 @@ inline std::unique_ptr<cudf::table> join_quadtree_and_bboxes(cudf::table_view co
 
   // Count the number of top-level nodes to start.
   // This could be provided explicitly, but count_if should be fast enough.
-  auto num_top_level_leaves = thrust::count_if(rmm::exec_policy(stream)->on(stream.value()),
+  auto num_top_level_leaves = thrust::count_if(rmm::exec_policy(stream),
                                                node_levels.begin<uint8_t>(),
                                                node_levels.end<uint8_t>(),
                                                thrust::placeholders::_1 == 0);
@@ -178,12 +179,12 @@ inline std::unique_ptr<cudf::table> join_quadtree_and_bboxes(cudf::table_view co
   cols.push_back(make_fixed_width_column<uint32_t>(num_results, stream, mr));
   cols.push_back(make_fixed_width_column<uint32_t>(num_results, stream, mr));
 
-  thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::copy(rmm::exec_policy(stream),
                out_poly_idxs.begin(),
                out_poly_idxs.begin() + num_results,
                cols.at(0)->mutable_view().begin<uint32_t>());
 
-  thrust::copy(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::copy(rmm::exec_policy(stream),
                out_node_idxs.begin(),
                out_node_idxs.begin() + num_results,
                cols.at(1)->mutable_view().begin<uint32_t>());
diff --git a/cpp/src/spatial/hausdorff.cu b/cpp/src/spatial/hausdorff.cu
index 39305cedd..2e7c50e00 100644
--- a/cpp/src/spatial/hausdorff.cu
+++ b/cpp/src/spatial/hausdorff.cu
@@ -29,6 +29,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/binary_search.h>
 #include <thrust/functional.h>
@@ -129,14 +130,14 @@ struct hausdorff_functor {
 
     auto num_cartesian = num_points * num_points;
 
-    thrust::inclusive_scan_by_key(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::inclusive_scan_by_key(rmm::exec_policy(stream),
                                   gpc_key_iter,
                                   gpc_key_iter + num_cartesian,
                                   hausdorff_acc_iter,
                                   scatter_out,
                                   thrust::equal_to<thrust::pair<int32_t, int32_t>>());
 
-    thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::transform(rmm::exec_policy(stream),
                       result_temp_iter,
                       result_temp_iter + num_results,
                       result->mutable_view().begin<T>(),
diff --git a/cpp/src/spatial/haversine.cu b/cpp/src/spatial/haversine.cu
index b76ec1584..4e09035d2 100644
--- a/cpp/src/spatial/haversine.cu
+++ b/cpp/src/spatial/haversine.cu
@@ -26,6 +26,7 @@
 
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 #include <rmm/mr/device/device_memory_resource.hpp>
 
 #include <memory>
@@ -82,7 +83,7 @@ struct haversine_functor {
 
     auto input_iter = thrust::make_zip_iterator(input_tuple);
 
-    thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::transform(rmm::exec_policy(stream),
                       input_iter,
                       input_iter + result->size(),
                       result->mutable_view().begin<T>(),
diff --git a/cpp/src/spatial/lonlat_to_cartesian.cu b/cpp/src/spatial/lonlat_to_cartesian.cu
index 2887765ff..b599858ae 100644
--- a/cpp/src/spatial/lonlat_to_cartesian.cu
+++ b/cpp/src/spatial/lonlat_to_cartesian.cu
@@ -22,6 +22,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <type_traits>
 #include <utility>
@@ -85,11 +86,8 @@ struct lonlat_to_cartesian_functor {
                                lat_to_y(origin_lat - lat));
     };
 
-    thrust::transform(rmm::exec_policy(stream)->on(stream.value()),
-                      input_zip,
-                      input_zip + input_lon.size(),
-                      output_zip,
-                      to_cartesian);
+    thrust::transform(
+      rmm::exec_policy(stream), input_zip, input_zip + input_lon.size(), output_zip, to_cartesian);
 
     return std::make_pair(std::move(output_x), std::move(output_y));
   }
diff --git a/cpp/src/spatial/polygon_bounding_box.cu b/cpp/src/spatial/polygon_bounding_box.cu
index c3ba056b3..5801e644c 100644
--- a/cpp/src/spatial/polygon_bounding_box.cu
+++ b/cpp/src/spatial/polygon_bounding_box.cu
@@ -26,6 +26,7 @@
 
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/functional.h>
 #include <thrust/gather.h>
@@ -69,20 +70,20 @@ std::unique_ptr<cudf::table> compute_polygon_bounding_boxes(cudf::column_view co
     rmm::device_vector<int32_t> first_ring_offsets(num_polygons);
 
     // Gather the first ring offset for each polygon
-    thrust::gather(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::gather(rmm::exec_policy(stream),
                    poly_offsets.begin<int32_t>(),
                    poly_offsets.end<int32_t>(),
                    ring_offsets.begin<int32_t>(),
                    first_ring_offsets.begin());
 
     // Scatter the first ring offset into a list of point_ids for reduction
-    thrust::scatter(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::scatter(rmm::exec_policy(stream),
                     thrust::make_counting_iterator(0),
                     thrust::make_counting_iterator(0) + num_polygons,
                     first_ring_offsets.begin(),
                     point_ids.begin());
 
-    thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+    thrust::inclusive_scan(rmm::exec_policy(stream),
                            point_ids.begin(),
                            point_ids.end(),
                            point_ids.begin(),
@@ -113,7 +114,7 @@ std::unique_ptr<cudf::table> compute_polygon_bounding_boxes(cudf::column_view co
   auto points_iter = thrust::make_zip_iterator(thrust::make_tuple(x.begin<T>(), y.begin<T>()));
   auto points_squared_iter = thrust::make_transform_iterator(points_iter, point_to_square<T>{});
 
-  thrust::reduce_by_key(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::reduce_by_key(rmm::exec_policy(stream),
                         point_ids.begin(),
                         point_ids.end(),
                         points_squared_iter,
diff --git a/cpp/src/spatial/polyline_bounding_box.cu b/cpp/src/spatial/polyline_bounding_box.cu
index 0f3edc313..1514c626f 100644
--- a/cpp/src/spatial/polyline_bounding_box.cu
+++ b/cpp/src/spatial/polyline_bounding_box.cu
@@ -24,6 +24,7 @@
 
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/functional.h>
 #include <thrust/gather.h>
@@ -65,13 +66,13 @@ std::unique_ptr<cudf::table> compute_polyline_bounding_boxes(cudf::column_view c
   rmm::device_vector<int32_t> point_ids(x.size());
 
   // Scatter the polyline offsets into a list of point_ids for reduction
-  thrust::scatter(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::scatter(rmm::exec_policy(stream),
                   thrust::make_counting_iterator(0),
                   thrust::make_counting_iterator(0) + num_polygons,
                   poly_offsets.begin<int32_t>(),
                   point_ids.begin());
 
-  thrust::inclusive_scan(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::inclusive_scan(rmm::exec_policy(stream),
                          point_ids.begin(),
                          point_ids.end(),
                          point_ids.begin(),
@@ -100,7 +101,7 @@ std::unique_ptr<cudf::table> compute_polyline_bounding_boxes(cudf::column_view c
   auto points_squared_iter =
     thrust::make_transform_iterator(points_iter, point_to_square<T>{expansion_radius});
 
-  thrust::reduce_by_key(rmm::exec_policy(stream)->on(stream.value()),
+  thrust::reduce_by_key(rmm::exec_policy(stream),
                         point_ids.begin(),
                         point_ids.end(),
                         points_squared_iter,
diff --git a/cpp/src/trajectory/derive_trajectories.cu b/cpp/src/trajectory/derive_trajectories.cu
index 7a63d74a1..5fbbd2e09 100644
--- a/cpp/src/trajectory/derive_trajectories.cu
+++ b/cpp/src/trajectory/derive_trajectories.cu
@@ -25,6 +25,7 @@
 
 #include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
@@ -50,10 +51,9 @@ std::pair<std::unique_ptr<cudf::table>, std::unique_ptr<cudf::column>> derive_tr
                                           stream,
                                           mr);
 
-  auto policy    = rmm::exec_policy(stream);
   auto sorted_id = sorted->get_column(0).view();
   rmm::device_vector<int32_t> lengths(object_id.size());
-  auto grouped = thrust::reduce_by_key(policy->on(stream.value()),
+  auto grouped = thrust::reduce_by_key(rmm::exec_policy(stream),
                                        sorted_id.begin<int32_t>(),
                                        sorted_id.end<int32_t>(),
                                        thrust::make_constant_iterator(1),
@@ -66,7 +66,7 @@ std::pair<std::unique_ptr<cudf::table>, std::unique_ptr<cudf::column>> derive_tr
                                            stream,
                                            mr);
 
-  thrust::exclusive_scan(policy->on(stream.value()),
+  thrust::exclusive_scan(rmm::exec_policy(stream),
                          lengths.begin(),
                          lengths.end(),
                          offsets->mutable_view().begin<int32_t>());
diff --git a/cpp/src/trajectory/trajectory_bounding_boxes.cu b/cpp/src/trajectory/trajectory_bounding_boxes.cu
index ec1528168..280c3fab9 100644
--- a/cpp/src/trajectory/trajectory_bounding_boxes.cu
+++ b/cpp/src/trajectory/trajectory_bounding_boxes.cu
@@ -27,6 +27,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/iterator/discard_iterator.h>
 
@@ -44,8 +45,6 @@ struct dispatch_element {
     rmm::cuda_stream_view stream,
     rmm::mr::device_memory_resource* mr)
   {
-    auto policy = rmm::exec_policy(stream);
-
     // Construct output columns
     auto type = cudf::data_type{cudf::type_to_id<Element>()};
     std::vector<std::unique_ptr<cudf::column>> cols{};
@@ -73,7 +72,7 @@ struct dispatch_element {
                          cols.at(3)->mutable_view().begin<Element>())  // bbox_y2
     );
 
-    thrust::fill(policy->on(stream.value()),
+    thrust::fill(rmm::exec_policy(stream),
                  bboxes,
                  bboxes + num_trajectories,
                  thrust::make_tuple(std::numeric_limits<Element>::max(),
@@ -82,7 +81,7 @@ struct dispatch_element {
                                     std::numeric_limits<Element>::min()));
 
     thrust::reduce_by_key(
-      policy->on(stream.value()),       // execution policy
+      rmm::exec_policy(stream),         // execution policy
       object_id.begin<int32_t>(),       // keys_first
       object_id.end<int32_t>(),         // keys_last
       points,                           // values_first
diff --git a/cpp/src/trajectory/trajectory_distances_and_speeds.cu b/cpp/src/trajectory/trajectory_distances_and_speeds.cu
index 90bb8a025..fb74626e2 100644
--- a/cpp/src/trajectory/trajectory_distances_and_speeds.cu
+++ b/cpp/src/trajectory/trajectory_distances_and_speeds.cu
@@ -27,6 +27,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <thrust/iterator/discard_iterator.h>
 
@@ -63,8 +64,6 @@ struct dispatch_timestamp {
     rmm::cuda_stream_view stream,
     rmm::mr::device_memory_resource* mr)
   {
-    auto policy = rmm::exec_policy(stream);
-
     // Construct output columns
     std::vector<std::unique_ptr<cudf::column>> cols{};
     cols.reserve(2);
@@ -104,7 +103,7 @@ struct dispatch_timestamp {
 
     // Compute duration and distance difference between adjacent elements that
     // share the same object id
-    thrust::adjacent_difference(policy->on(stream.value()),
+    thrust::adjacent_difference(rmm::exec_policy(stream),
                                 timestamp_point_and_id,                     // first
                                 timestamp_point_and_id + durations.size(),  // last
                                 duration_and_distance_1,                    // result
@@ -148,7 +147,7 @@ struct dispatch_timestamp {
 
     // Reduce the intermediate durations and kilometer distances into meter
     // distances and speeds in meters/second
-    thrust::reduce_by_key(policy->on(stream.value()),
+    thrust::reduce_by_key(rmm::exec_policy(stream),
                           object_id.begin<int32_t>(),       // keys_first
                           object_id.end<int32_t>(),         // keys_last
                           duration_and_distance_2 + 1,      // values_first