diff --git a/cub/benchmarks/bench/reduce/by_key.cu b/cub/benchmarks/bench/reduce/by_key.cu
index 870aedea540..0228a58c465 100644
--- a/cub/benchmarks/bench/reduce/by_key.cu
+++ b/cub/benchmarks/bench/reduce/by_key.cu
@@ -3,6 +3,11 @@
 
 #include <cub/device/device_reduce.cuh>
 
+#include <cuda/buffer>
+#include <cuda/std/execution>
+#include <cuda/std/functional>
+#include <cuda/stream>
+
 #include <look_back_helper.cuh>
 #include <nvbench_helper.cuh>
 
@@ -40,19 +45,22 @@ static void reduce_by_key(nvbench::state& state, nvbench::type_list<KeyT, ValueT
   constexpr std::size_t min_segment_size = 1;
   const std::size_t max_segment_size     = static_cast<std::size_t>(state.get_int64("MaxSegSize"));
 
-  thrust::device_vector<OffsetT> num_runs_out(1);
-  thrust::device_vector<ValueT> in_vals(elements);
-  thrust::device_vector<ValueT> out_vals(elements);
-  thrust::device_vector<KeyT> out_keys(elements);
-  thrust::device_vector<KeyT> in_keys = generate.uniform.key_segments(elements, min_segment_size, max_segment_size);
+  const auto stream = get_stream_ref(state);
+  const auto device = stream.device();
+  caching_allocator_t alloc;
 
-  const KeyT* d_in_keys   = thrust::raw_pointer_cast(in_keys.data());
-  KeyT* d_out_keys        = thrust::raw_pointer_cast(out_keys.data());
-  const ValueT* d_in_vals = thrust::raw_pointer_cast(in_vals.data());
-  ValueT* d_out_vals      = thrust::raw_pointer_cast(out_vals.data());
-  OffsetT* d_num_runs_out = thrust::raw_pointer_cast(num_runs_out.data());
+  auto num_runs_out  = cuda::make_buffer<OffsetT>(stream, pinned_memory_resource(), 1, cuda::no_init);
+  const auto in_vals = cuda::make_device_buffer<ValueT>(stream, device, elements, ValueT{});
+  auto out_vals      = cuda::make_device_buffer<ValueT>(stream, device, elements, cuda::no_init);
+  auto out_keys      = cuda::make_device_buffer<KeyT>(stream, device, elements, cuda::no_init);
+  const auto in_keys =
+    generate.uniform.key_segments(elements, min_segment_size, max_segment_size).device_buffer<KeyT>(stream, device);
 
-  caching_allocator_t alloc;
+  const KeyT* d_in_keys   = in_keys.data();
+  KeyT* d_out_keys        = out_keys.data();
+  const ValueT* d_in_vals = in_vals.data();
+  ValueT* d_out_vals      = out_vals.data();
+  OffsetT* d_num_runs_out = num_runs_out.data();
 
   // Run once to get the number of runs for reporting
   _CCCL_TRY_CUDA_API(
@@ -65,8 +73,8 @@ static void reduce_by_key(nvbench::state& state, nvbench::type_list<KeyT, ValueT
     d_num_runs_out,
     reduction_op_t{},
     static_cast<OffsetT>(elements),
-    alloc);
-  cudaDeviceSynchronize();
+    cub_bench_env(alloc, stream));
+  stream.sync();
   const OffsetT num_runs = num_runs_out[0];
 
   state.add_element_count(elements);
@@ -79,9 +87,9 @@ static void reduce_by_key(nvbench::state& state, nvbench::type_list<KeyT, ValueT
   state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch& launch) {
     auto env = cub_bench_env(
       alloc,
-      launch
+      get_stream_ref(launch)
 #if !TUNE_BASE
-      ,
+        ,
       cuda::execution::tune(bench_reduce_by_key_policy_selector{})
 #endif // !TUNE_BASE
     );
diff --git a/cub/benchmarks/bench/run_length_encode/encode.cu b/cub/benchmarks/bench/run_length_encode/encode.cu
index 37f748e7307..63520796fe5 100644
--- a/cub/benchmarks/bench/run_length_encode/encode.cu
+++ b/cub/benchmarks/bench/run_length_encode/encode.cu
@@ -3,6 +3,10 @@
 
 #include <cub/device/device_run_length_encode.cuh>
 
+#include <cuda/buffer>
+#include <cuda/std/execution>
+#include <cuda/stream>
+
 #include <look_back_helper.cuh>
 #include <nvbench_helper.cuh>
 
@@ -42,20 +46,32 @@ static void rle(nvbench::state& state, nvbench::type_list<T, OffsetT, RunLengthT
   constexpr std::size_t min_segment_size = 1;
   const std::size_t max_segment_size     = static_cast<std::size_t>(state.get_int64("MaxSegSize"));
 
-  thrust::device_vector<offset_t> num_runs_out(1);
-  thrust::device_vector<RunLengthT> out_counts(elements);
-  thrust::device_vector<T> out_keys(elements);
-  thrust::device_vector<T> in_keys = generate.uniform.key_segments(elements, min_segment_size, max_segment_size);
+  const auto stream = get_stream_ref(state);
+  const auto device = stream.device();
+  caching_allocator_t alloc;
+
+  auto num_runs_out = cuda::make_buffer<offset_t>(stream, pinned_memory_resource(), 1, cuda::no_init);
+  auto out_counts   = cuda::make_device_buffer<RunLengthT>(stream, device, elements, cuda::no_init);
+  auto out_keys     = cuda::make_device_buffer<T>(stream, device, elements, cuda::no_init);
+  const auto in_keys =
+    generate.uniform.key_segments(elements, min_segment_size, max_segment_size).device_buffer<T>(stream, device);
 
-  const T* d_in_keys       = thrust::raw_pointer_cast(in_keys.data());
-  T* d_out_keys            = thrust::raw_pointer_cast(out_keys.data());
-  RunLengthT* d_out_counts = thrust::raw_pointer_cast(out_counts.data());
-  offset_t* d_num_runs_out = thrust::raw_pointer_cast(num_runs_out.data());
+  const T* d_in_keys       = in_keys.data();
+  T* d_out_keys            = out_keys.data();
+  RunLengthT* d_out_counts = out_counts.data();
+  offset_t* d_num_runs_out = num_runs_out.data();
 
   // Run once to get num_runs for memory accounting
-  (void) cub::DeviceRunLengthEncode::Encode(
-    d_in_keys, d_out_keys, d_out_counts, d_num_runs_out, static_cast<OffsetT>(elements));
-  cudaDeviceSynchronize();
+  _CCCL_TRY_CUDA_API(
+    cub::DeviceRunLengthEncode::Encode,
+    "Encode failed",
+    d_in_keys,
+    d_out_keys,
+    d_out_counts,
+    d_num_runs_out,
+    static_cast<OffsetT>(elements),
+    cub_bench_env(alloc, stream));
+  stream.sync();
   const offset_t num_runs = num_runs_out[0];
 
   state.add_element_count(elements);
@@ -64,13 +80,12 @@ static void rle(nvbench::state& state, nvbench::type_list<T, OffsetT, RunLengthT
   state.add_global_memory_writes<RunLengthT>(num_runs);
   state.add_global_memory_writes<offset_t>(1);
 
-  caching_allocator_t alloc;
   state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch& launch) {
     auto env = cub_bench_env(
       alloc,
-      launch
+      get_stream_ref(launch)
 #if !TUNE_BASE
-      ,
+        ,
       cuda::execution::tune(bench_encode_policy_selector{})
 #endif // !TUNE_BASE
     );
diff --git a/cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu b/cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu
index 5e9e2bda2d3..1d3dc4c9f4e 100644
--- a/cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu
+++ b/cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu
@@ -3,6 +3,10 @@
 
 #include <cub/device/device_run_length_encode.cuh>
 
+#include <cuda/buffer>
+#include <cuda/std/execution>
+#include <cuda/stream>
+
 #include <look_back_helper.cuh>
 #include <nvbench_helper.cuh>
 
@@ -43,23 +47,23 @@ static void rle(nvbench::state& state, nvbench::type_list<T, OffsetT, RunLengthT
   constexpr std::size_t min_segment_size = 1;
   const std::size_t max_segment_size     = static_cast<std::size_t>(state.get_int64("MaxSegSize"));
 
-  thrust::device_vector<offset_t> num_runs_out(1);
-  thrust::device_vector<offset_t> out_offsets(elements);
-  thrust::device_vector<RunLengthT> out_lengths(elements);
-  thrust::device_vector<T> in_keys = generate.uniform.key_segments(elements, min_segment_size, max_segment_size);
+  const auto stream = get_stream_ref(state);
+  const auto device = stream.device();
+  caching_allocator_t alloc;
+
+  auto num_runs_out = cuda::make_buffer<offset_t>(stream, pinned_memory_resource(), 1, cuda::no_init);
+  auto out_offsets  = cuda::make_device_buffer<offset_t>(stream, device, elements, cuda::no_init);
+  auto out_lengths  = cuda::make_device_buffer<RunLengthT>(stream, device, elements, cuda::no_init);
+  const auto in_keys =
+    generate.uniform.key_segments(elements, min_segment_size, max_segment_size).device_buffer<T>(stream, device);
 
-  const T* d_in_keys        = thrust::raw_pointer_cast(in_keys.data());
-  offset_t* d_out_offsets   = thrust::raw_pointer_cast(out_offsets.data());
-  RunLengthT* d_out_lengths = thrust::raw_pointer_cast(out_lengths.data());
-  offset_t* d_num_runs_out  = thrust::raw_pointer_cast(num_runs_out.data());
+  const T* d_in_keys        = in_keys.data();
+  offset_t* d_out_offsets   = out_offsets.data();
+  RunLengthT* d_out_lengths = out_lengths.data();
+  offset_t* d_num_runs_out  = num_runs_out.data();
 
   {
     // Run once to get num_runs for memory accounting
-    auto memory_env = cuda::std::execution::env{
-#if !TUNE_BASE
-      cuda::execution::tune(bench_rle_policy_selector{})
-#endif // !TUNE_BASE
-    };
     _CCCL_TRY_CUDA_API(
       cub::DeviceRunLengthEncode::NonTrivialRuns,
       "NonTrivialRuns failed",
@@ -68,8 +72,14 @@ static void rle(nvbench::state& state, nvbench::type_list<T, OffsetT, RunLengthT
       d_out_lengths,
       d_num_runs_out,
       static_cast<OffsetT>(elements),
-      memory_env);
-    cudaDeviceSynchronize();
+      cub_bench_env(alloc,
+                    stream
+#if !TUNE_BASE
+                    ,
+                    cuda::execution::tune(bench_rle_policy_selector{})
+#endif // !TUNE_BASE
+                      ));
+    stream.sync();
   }
   const OffsetT num_runs = num_runs_out[0];
 
@@ -79,13 +89,12 @@ static void rle(nvbench::state& state, nvbench::type_list<T, OffsetT, RunLengthT
   state.add_global_memory_writes<OffsetT>(num_runs);
   state.add_global_memory_writes<OffsetT>(1);
 
-  caching_allocator_t alloc;
   state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch& launch) {
     auto env = cub_bench_env(
       alloc,
-      launch
+      get_stream_ref(launch)
 #if !TUNE_BASE
-      ,
+        ,
       cuda::execution::tune(bench_rle_policy_selector{})
 #endif // !TUNE_BASE
     );
diff --git a/cub/benchmarks/bench/select/unique.cu b/cub/benchmarks/bench/select/unique.cu
index f3486b8be29..9a33fa20679 100644
--- a/cub/benchmarks/bench/select/unique.cu
+++ b/cub/benchmarks/bench/select/unique.cu
@@ -3,7 +3,10 @@
 
 #include <cub/device/device_select.cuh>
 
-#include <cuda/std/algorithm>
+#include <cuda/buffer>
+#include <cuda/std/execution>
+#include <cuda/std/functional>
+#include <cuda/stream>
 
 #include <limits>
 
@@ -43,13 +46,18 @@ static void unique(nvbench::state& state, nvbench::type_list<T, InPlace>)
   const auto elements         = state.get_int64("Elements{io}");
   const auto max_segment_size = state.get_int64("MaxSegSize");
 
-  thrust::device_vector<T> in = generate.uniform.key_segments(elements, /* min_segmented_size */ 1, max_segment_size);
-  thrust::device_vector<T> out(elements, thrust::no_init);
-  thrust::device_vector<offset_t> num_unique_out(1);
+  const auto stream = get_stream_ref(state);
+  const auto device = stream.device();
+  caching_allocator_t alloc;
+
+  auto in = generate.uniform.key_segments(elements, /* min_segmented_size */ 1, max_segment_size)
+              .device_buffer<T>(stream, device);
+  auto out            = cuda::make_device_buffer<T>(stream, device, elements, cuda::no_init);
+  auto num_unique_out = cuda::make_buffer<offset_t>(stream, pinned_memory_resource(), 1, cuda::no_init);
 
-  T* d_in                = thrust::raw_pointer_cast(in.data());
-  T* d_out               = thrust::raw_pointer_cast(out.data());
-  offset_t* d_num_unique = thrust::raw_pointer_cast(num_unique_out.data());
+  T* d_in                = in.data();
+  T* d_out               = out.data();
+  offset_t* d_num_unique = num_unique_out.data();
 
   // Get number of unique elements for metrics
   _CCCL_TRY_CUDA_API(
@@ -59,8 +67,9 @@ static void unique(nvbench::state& state, nvbench::type_list<T, InPlace>)
     d_out,
     d_num_unique,
     static_cast<offset_t>(elements),
-    ::cuda::std::equal_to<>{});
-  cudaDeviceSynchronize();
+    ::cuda::std::equal_to<>{},
+    cub_bench_env(alloc, stream));
+  stream.sync();
   const offset_t num_unique = num_unique_out[0];
 
   state.add_element_count(elements);
@@ -68,13 +77,12 @@ static void unique(nvbench::state& state, nvbench::type_list<T, InPlace>)
   state.add_global_memory_writes<T>(num_unique);
   state.add_global_memory_writes<offset_t>(1);
 
-  caching_allocator_t alloc;
   state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch& launch) {
     auto env = cub_bench_env(
       alloc,
-      launch
+      get_stream_ref(launch)
 #if !TUNE_BASE
-      ,
+        ,
       cuda::execution::tune(bench_policy_selector<T>{})
 #endif // !TUNE_BASE
     );
diff --git a/cub/benchmarks/bench/select/unique_by_key.cu b/cub/benchmarks/bench/select/unique_by_key.cu
index 01dc2e5a354..f09632da22d 100644
--- a/cub/benchmarks/bench/select/unique_by_key.cu
+++ b/cub/benchmarks/bench/select/unique_by_key.cu
@@ -3,6 +3,11 @@
 
 #include <cub/device/device_select.cuh>
 
+#include <cuda/buffer>
+#include <cuda/std/execution>
+#include <cuda/std/functional>
+#include <cuda/stream>
+
 #include <look_back_helper.cuh>
 #include <nvbench_helper.cuh>
 
@@ -50,17 +55,22 @@ static void select(nvbench::state& state, nvbench::type_list<KeyT, ValueT, Offse
   constexpr std::size_t min_segment_size = 1;
   const std::size_t max_segment_size     = static_cast<std::size_t>(state.get_int64("MaxSegSize"));
 
-  thrust::device_vector<OffsetT> num_runs_out(1);
-  thrust::device_vector<ValueT> in_vals(elements);
-  thrust::device_vector<ValueT> out_vals(elements);
-  thrust::device_vector<KeyT> out_keys(elements);
-  thrust::device_vector<KeyT> in_keys = generate.uniform.key_segments(elements, min_segment_size, max_segment_size);
+  const auto stream = get_stream_ref(state);
+  const auto device = stream.device();
+  caching_allocator_t alloc;
+
+  auto num_runs_out  = cuda::make_buffer<OffsetT>(stream, pinned_memory_resource(), 1, cuda::no_init);
+  const auto in_vals = cuda::make_device_buffer<ValueT>(stream, device, elements, ValueT{});
+  auto out_vals      = cuda::make_device_buffer<ValueT>(stream, device, elements, cuda::no_init);
+  auto out_keys      = cuda::make_device_buffer<KeyT>(stream, device, elements, cuda::no_init);
+  const auto in_keys =
+    generate.uniform.key_segments(elements, min_segment_size, max_segment_size).device_buffer<KeyT>(stream, device);
 
-  const KeyT* d_in_keys   = thrust::raw_pointer_cast(in_keys.data());
-  KeyT* d_out_keys        = thrust::raw_pointer_cast(out_keys.data());
-  const ValueT* d_in_vals = thrust::raw_pointer_cast(in_vals.data());
-  ValueT* d_out_vals      = thrust::raw_pointer_cast(out_vals.data());
-  OffsetT* d_num_runs_out = thrust::raw_pointer_cast(num_runs_out.data());
+  const KeyT* d_in_keys   = in_keys.data();
+  KeyT* d_out_keys        = out_keys.data();
+  const ValueT* d_in_vals = in_vals.data();
+  ValueT* d_out_vals      = out_vals.data();
+  OffsetT* d_num_runs_out = num_runs_out.data();
 
   const auto num_items = static_cast<OffsetT>(elements);
 
@@ -74,8 +84,9 @@ static void select(nvbench::state& state, nvbench::type_list<KeyT, ValueT, Offse
     d_out_vals,
     d_num_runs_out,
     num_items,
-    equality_op_t{});
-  _CCCL_TRY_CUDA_API(cudaDeviceSynchronize, "Sync failed");
+    equality_op_t{},
+    cub_bench_env(alloc, stream));
+  stream.sync();
   const OffsetT num_runs = num_runs_out[0];
 
   state.add_element_count(elements);
@@ -85,13 +96,12 @@ static void select(nvbench::state& state, nvbench::type_list<KeyT, ValueT, Offse
   state.add_global_memory_writes<KeyT>(num_runs);
   state.add_global_memory_writes<OffsetT>(1);
 
-  caching_allocator_t alloc;
   state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch& launch) {
     auto env = cub_bench_env(
       alloc,
-      launch
+      get_stream_ref(launch)
 #if !TUNE_BASE
-      ,
+        ,
       cuda::execution::tune(bench_unique_by_key_policy_selector{})
 #endif // !TUNE_BASE
     );
diff --git a/nvbench_helper/nvbench_helper/nvbench_helper.cu b/nvbench_helper/nvbench_helper/nvbench_helper.cu
index f7a3102f709..3ee294696e9 100644
--- a/nvbench_helper/nvbench_helper/nvbench_helper.cu
+++ b/nvbench_helper/nvbench_helper/nvbench_helper.cu
@@ -110,6 +110,13 @@ public:
   template <typename T>
   void generate(seed_t seed, cuda::std::span<T> device_span, bit_entropy entropy, T min, T max);
 
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+  void set_stream(cuda::stream_ref stream)
+  {
+    curandSetStream(m_gen, stream.get());
+  }
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+
   const double* new_uniform_distribution(seed_t seed, std::size_t num_items);
   const double* new_lognormal_distribution(seed_t seed, std::size_t num_items);
   const double* new_constant(std::size_t num_items, double val);
@@ -284,6 +291,16 @@ public:
     }
   }
 
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+  template <typename T>
+  void generate(cuda::stream_ref stream, seed_t seed, cuda::std::span<T> span, bit_entropy entropy, T min, T max)
+  {
+    construct_guard(executor::device);
+    m_device_generator->set_stream(stream);
+    this->generate(thrust::cuda::par.on(stream.get()), *m_device_generator, seed, span, entropy, min, max);
+  }
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+
   template <typename T>
   void power_law_segment_offsets(executor exec, seed_t seed, cuda::std::span<T> span, std::size_t total_elements)
   {
@@ -387,7 +404,7 @@ void generator_t::generate(
 
       for (int i = 0; i < number_of_steps; i++, ++seed)
       {
-        this->generate(is_device ? executor::device : executor::host, seed, tmp, bit_entropy::_1_000, min, max);
+        this->generate(exec, dist, seed, tmp, bit_entropy::_1_000, min, max);
 
         thrust::transform(
           exec,
@@ -466,7 +483,7 @@ void generator_t::generate(
 
       for (int i = 0; i < number_of_steps; i++, ++seed)
       {
-        this->generate(is_device ? executor::device : executor::host, seed, tmp, bit_entropy::_1_000, min, max);
+        this->generate(exec, dist, seed, tmp, bit_entropy::_1_000, min, max);
 
         thrust::transform(
           exec,
@@ -592,6 +609,14 @@ void gen_device(seed_t seed, cuda::std::span<T> device_span, bit_entropy entropy
   gen(executor::device, seed, device_span, entropy, min, max);
 }
 
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+template <typename T>
+void gen_device(seed_t seed, cuda::stream_ref stream, cuda::std::span<T> device_span, bit_entropy entropy, T min, T max)
+{
+  generator_t{}.generate(stream, seed, device_span, entropy, min, max);
+}
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+
 template <class T>
 struct offset_to_iterator_t
 {
@@ -797,8 +822,17 @@ INSTANTIATE(uint64_t);
 #undef INSTANTIATE
 
 // Instantiates only the uniform data generators used by non-segmented benchmarks (e.g. Reduce/Scan/RadixSort).
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+#  define INSTANTIATE_STREAM_GEN(TYPE)      \
+    template void detail::gen_device<TYPE>( \
+      seed_t, cuda::stream_ref, cuda::std::span<TYPE>, bit_entropy, TYPE min, TYPE max)
+#else // THRUST_DEVICE_SYSTEM != THRUST_DEVICE_SYSTEM_CUDA
+#  define INSTANTIATE_STREAM_GEN(TYPE)
+#endif // THRUST_DEVICE_SYSTEM != THRUST_DEVICE_SYSTEM_CUDA
+
 #define INSTANTIATE_GEN(TYPE)                                                                             \
   template void detail::gen_device<TYPE>(seed_t, cuda::std::span<TYPE>, bit_entropy, TYPE min, TYPE max); \
+  INSTANTIATE_STREAM_GEN(TYPE);                                                                           \
   template void detail::gen_host<TYPE>(seed_t, cuda::std::span<TYPE>, bit_entropy, TYPE min, TYPE max)
 
 #define INSTANTIATE(TYPE)                                                                                               \
@@ -838,3 +872,4 @@ INSTANTIATE_GEN(__nv_bfloat16);
 
 #undef INSTANTIATE
 #undef INSTANTIATE_GEN
+#undef INSTANTIATE_STREAM_GEN
diff --git a/nvbench_helper/nvbench_helper/nvbench_helper.cuh b/nvbench_helper/nvbench_helper/nvbench_helper.cuh
index d9e0ccd5ed9..668ec22d72e 100644
--- a/nvbench_helper/nvbench_helper/nvbench_helper.cuh
+++ b/nvbench_helper/nvbench_helper/nvbench_helper.cuh
@@ -13,6 +13,8 @@
 #include <cuda/std/type_traits>
 
 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+#  include <cuda/buffer>
+#  include <cuda/devices>
 #  include <cuda/memory_resource>
 #  include <cuda/std/execution>
 #  include <cuda/stream>
@@ -212,6 +214,24 @@ NVBENCH_DECLARE_TYPE_STRINGS(bit_entropy, "BE", "bit entropy");
   throw std::runtime_error("Can't convert string to bit entropy");
 }
 
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+[[nodiscard]] inline ::cuda::stream_ref get_stream_ref(nvbench::state& state)
+{
+  return ::cuda::stream_ref{state.get_cuda_stream()};
+}
+
+[[nodiscard]] inline ::cuda::stream_ref get_stream_ref(nvbench::launch& launch)
+{
+  return ::cuda::stream_ref{launch.get_stream()};
+}
+
+[[nodiscard]] inline auto pinned_memory_resource()
+{
+  return ::cuda::mr::synchronous_resource_adapter<::cuda::mr::legacy_pinned_memory_resource>{
+    ::cuda::mr::legacy_pinned_memory_resource{}};
+}
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+
 // Creates an interpolated value of type T between min (at = 0.0) and max (at = 1.0).
 template <typename T>
 [[nodiscard]] T lerp_min_max(double at) noexcept
@@ -229,12 +249,45 @@ namespace detail
 {
 void do_not_optimize(const void* ptr);
 
+template <typename T>
+struct default_gen_bounds
+{
+  [[nodiscard]] static T min()
+  {
+    return ::cuda::std::numeric_limits<T>::min();
+  }
+
+  [[nodiscard]] static T max()
+  {
+    return ::cuda::std::numeric_limits<T>::max();
+  }
+};
+
+template <typename T>
+struct default_gen_bounds<::cuda::std::complex<T>>
+{
+  [[nodiscard]] static ::cuda::std::complex<T> min()
+  {
+    return {::cuda::std::numeric_limits<T>::min(), ::cuda::std::numeric_limits<T>::min()};
+  }
+
+  [[nodiscard]] static ::cuda::std::complex<T> max()
+  {
+    return {::cuda::std::numeric_limits<T>::max(), ::cuda::std::numeric_limits<T>::max()};
+  }
+};
+
 template <typename T>
 void gen_host(seed_t seed, cuda::std::span<T> data, bit_entropy entropy, T min, T max);
 
 template <typename T>
 void gen_device(seed_t seed, cuda::std::span<T> data, bit_entropy entropy, T min, T max);
 
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+template <typename T>
+void gen_device(seed_t seed, cuda::stream_ref stream, cuda::std::span<T> data, bit_entropy entropy, T min, T max);
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+
 template <typename T>
 void gen_uniform_key_segments_host(
   seed_t seed, cuda::std::span<T> data, std::size_t min_segment_size, std::size_t max_segment_size);
@@ -278,6 +331,21 @@ struct generator_base_t
     ++m_seed;
     return vec;
   }
+
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+  template <typename T>
+  auto device_buffer(cuda::stream_ref stream, cuda::device_ref device, T min, T max)
+  {
+    auto buffer = cuda::make_device_buffer<T>(stream, device, m_elements, cuda::no_init);
+    stream.sync();
+
+    cuda::std::span<T> span(buffer.data(), buffer.size());
+    gen_device(m_seed, stream, span, m_entropy, min, max);
+
+    ++m_seed;
+    return buffer;
+  }
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
 };
 
 template <class T>
@@ -290,6 +358,13 @@ struct vector_generator_t : generator_base_t
   {
     return generator_base_t::generate(m_min, m_max);
   }
+
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+  auto device_buffer(cuda::stream_ref stream, cuda::device_ref device)
+  {
+    return generator_base_t::device_buffer(stream, device, m_min, m_max);
+  }
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
 };
 
 template <>
@@ -298,21 +373,16 @@ struct vector_generator_t<void> : generator_base_t
   template <typename T>
   operator thrust::device_vector<T>()
   {
-    return generator_base_t::generate(::cuda::std::numeric_limits<T>::min(), ::cuda::std::numeric_limits<T>::max());
+    return generator_base_t::generate(default_gen_bounds<T>::min(), default_gen_bounds<T>::max());
   }
 
-  // This overload is needed because numeric limits is not specialized for complex, making
-  // the min and max values for complex equal zero.
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
   template <typename T>
-  operator thrust::device_vector<::cuda::std::complex<T>>()
+  auto device_buffer(cuda::stream_ref stream, cuda::device_ref device)
   {
-    const auto min =
-      ::cuda::std::complex<T>{::cuda::std::numeric_limits<T>::min(), ::cuda::std::numeric_limits<T>::min()};
-    const auto max =
-      ::cuda::std::complex<T>{::cuda::std::numeric_limits<T>::max(), ::cuda::std::numeric_limits<T>::max()};
-
-    return generator_base_t::generate(min, max);
+    return generator_base_t::device_buffer(stream, device, default_gen_bounds<T>::min(), default_gen_bounds<T>::max());
   }
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
 };
 
 struct uniform_key_segments_generator_t
@@ -335,6 +405,21 @@ struct uniform_key_segments_generator_t
     ++m_seed;
     return keys_vec;
   }
+
+#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+  template <class KeyT>
+  auto device_buffer(cuda::stream_ref stream, cuda::device_ref device)
+  {
+    auto keys_buffer = cuda::make_device_buffer<KeyT>(stream, device, m_total_elements, cuda::no_init);
+    stream.sync();
+
+    cuda::std::span<KeyT> keys(keys_buffer.data(), keys_buffer.size());
+    gen_uniform_key_segments_device(m_seed, keys, m_min_segment_size, m_max_segment_size);
+
+    ++m_seed;
+    return keys_buffer;
+  }
+#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
 };
 
 struct uniform_segment_offsets_generator_t
@@ -753,8 +838,7 @@ auto policy(caching_allocator_t& alloc, nvbench::launch& launch)
 }
 auto cuda_policy(caching_allocator_t& alloc, nvbench::launch& launch)
 {
-  return cuda::execution::gpu.with(cuda::mr::get_memory_resource, alloc)
-    .with(cuda::get_stream, launch.get_stream().get_stream());
+  return cuda::execution::gpu.with(cuda::mr::get_memory_resource, alloc).with(cuda::get_stream, get_stream_ref(launch));
 }
 #else
 auto policy(caching_allocator_t&, nvbench::launch&)
@@ -764,12 +848,27 @@ auto policy(caching_allocator_t&, nvbench::launch&)
 #endif
 
 #if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
+// Returns an environment for benchmarking using memory_resource as MR, stream, and any additional envs passed in.
+template <typename MemoryResource, typename... MoreEnvs>
+auto cub_bench_env(MemoryResource& memory_resource, ::cuda::stream_ref stream, MoreEnvs... envs)
+{
+  return cuda::std::execution::env{stream, memory_resource, envs...};
+}
+
+// Returns an environment for benchmarking using alloc as MR, stream, and any additional envs passed in.
+template <typename... MoreEnvs>
+auto cub_bench_env(caching_allocator_t& alloc, ::cuda::stream_ref stream, MoreEnvs... envs)
+{
+  return cuda::std::execution::env{
+    stream, ::cuda::std::execution::prop{cuda::mr::get_memory_resource, ::cuda::mr::resource_ref<>{alloc}}, envs...};
+}
+
 // Returns an environment for benchmarking using alloc as MR, launch's stream, and any additional envs passed in.
 template <typename... MoreEnvs>
 auto cub_bench_env(caching_allocator_t& alloc, nvbench::launch& launch, MoreEnvs... envs)
 {
   return cuda::std::execution::env{
-    ::cuda::stream_ref{launch.get_stream().get_stream()},
+    get_stream_ref(launch),
     ::cuda::std::execution::prop{cuda::mr::get_memory_resource, ::cuda::mr::resource_ref<>{alloc}},
     envs...};
 }