diff --git a/cpp/include/cudf/detail/groupby/sort_helper.hpp b/cpp/include/cudf/detail/groupby/sort_helper.hpp
index 294978bf128..cadcb1265c4 100644
--- a/cpp/include/cudf/detail/groupby/sort_helper.hpp
+++ b/cpp/include/cudf/detail/groupby/sort_helper.hpp
@@ -22,7 +22,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_vector.hpp>
+#include <rmm/device_uvector.hpp>
 
 namespace cudf {
 namespace groupby {
@@ -40,8 +40,8 @@ namespace sort {
  *   value column
  */
 struct sort_groupby_helper {
-  using index_vector       = rmm::device_vector<size_type>;
-  using bitmask_vector     = rmm::device_vector<bitmask_type>;
+  using index_vector       = rmm::device_uvector<size_type>;
+  using bitmask_vector     = rmm::device_uvector<bitmask_type>;
   using column_ptr         = std::unique_ptr<column>;
   using index_vector_ptr   = std::unique_ptr<index_vector>;
   using bitmask_vector_ptr = std::unique_ptr<bitmask_vector>;
diff --git a/cpp/src/groupby/groupby.cu b/cpp/src/groupby/groupby.cu
index 3166b2be4d4..487aed4b411 100644
--- a/cpp/src/groupby/groupby.cu
+++ b/cpp/src/groupby/groupby.cu
@@ -164,18 +164,19 @@ groupby::groups groupby::get_groups(table_view values, rmm::mr::device_memory_re
   CUDF_FUNC_RANGE();
   auto grouped_keys = helper().sorted_keys(rmm::cuda_stream_default, mr);
 
-  auto group_offsets = helper().group_offsets(0);
+  auto const& group_offsets = helper().group_offsets(rmm::cuda_stream_default);
   std::vector<size_type> group_offsets_vector(group_offsets.size());
-  thrust::copy(group_offsets.begin(), group_offsets.end(), group_offsets_vector.begin());
+  thrust::copy(thrust::device_pointer_cast(group_offsets.begin()),
+               thrust::device_pointer_cast(group_offsets.end()),
+               group_offsets_vector.begin());
 
-  std::unique_ptr<table> grouped_values{nullptr};
   if (values.num_columns()) {
-    grouped_values = cudf::detail::gather(values,
-                                          helper().key_sort_order(),
-                                          cudf::out_of_bounds_policy::DONT_CHECK,
-                                          cudf::detail::negative_index_policy::NOT_ALLOWED,
-                                          rmm::cuda_stream_default,
-                                          mr);
+    auto grouped_values = cudf::detail::gather(values,
+                                               helper().key_sort_order(),
+                                               cudf::out_of_bounds_policy::DONT_CHECK,
+                                               cudf::detail::negative_index_policy::NOT_ALLOWED,
+                                               rmm::cuda_stream_default,
+                                               mr);
     return groupby::groups{
       std::move(grouped_keys), std::move(group_offsets_vector), std::move(grouped_values)};
   } else {
diff --git a/cpp/src/groupby/sort/group_argmax.cu b/cpp/src/groupby/sort/group_argmax.cu
index 94b5d3817a7..bed64c5147a 100644
--- a/cpp/src/groupby/sort/group_argmax.cu
+++ b/cpp/src/groupby/sort/group_argmax.cu
@@ -17,6 +17,7 @@
 #include <groupby/sort/group_single_pass_reduction_util.cuh>
 
 #include <cudf/detail/gather.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 
@@ -27,7 +28,7 @@ namespace groupby {
 namespace detail {
 std::unique_ptr<column> group_argmax(column_view const& values,
                                      size_type num_groups,
-                                     rmm::device_vector<size_type> const& group_labels,
+                                     cudf::device_span<size_type const> group_labels,
                                      column_view const& key_sort_order,
                                      rmm::cuda_stream_view stream,
                                      rmm::mr::device_memory_resource* mr)
diff --git a/cpp/src/groupby/sort/group_argmin.cu b/cpp/src/groupby/sort/group_argmin.cu
index 11a350ae1c4..ec97a609390 100644
--- a/cpp/src/groupby/sort/group_argmin.cu
+++ b/cpp/src/groupby/sort/group_argmin.cu
@@ -17,6 +17,7 @@
 #include <groupby/sort/group_single_pass_reduction_util.cuh>
 
 #include <cudf/detail/gather.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 
@@ -27,7 +28,7 @@ namespace groupby {
 namespace detail {
 std::unique_ptr<column> group_argmin(column_view const& values,
                                      size_type num_groups,
-                                     rmm::device_vector<size_type> const& group_labels,
+                                     cudf::device_span<size_type const> group_labels,
                                      column_view const& key_sort_order,
                                      rmm::cuda_stream_view stream,
                                      rmm::mr::device_memory_resource* mr)
diff --git a/cpp/src/groupby/sort/group_collect.cu b/cpp/src/groupby/sort/group_collect.cu
index 9c8ab92cc50..b7bcd05a72a 100644
--- a/cpp/src/groupby/sort/group_collect.cu
+++ b/cpp/src/groupby/sort/group_collect.cu
@@ -19,6 +19,7 @@
 #include <cudf/detail/aggregation/aggregation.hpp>
 #include <cudf/detail/gather.cuh>
 #include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 
@@ -26,13 +27,13 @@ namespace cudf {
 namespace groupby {
 namespace detail {
 std::unique_ptr<column> group_collect(column_view const &values,
-                                      rmm::device_vector<size_type> const &group_offsets,
+                                      cudf::device_span<size_type const> group_offsets,
                                       size_type num_groups,
                                       rmm::cuda_stream_view stream,
                                       rmm::mr::device_memory_resource *mr)
 {
   rmm::device_buffer offsets_data(
-    group_offsets.data().get(), group_offsets.size() * sizeof(cudf::size_type), stream, mr);
+    group_offsets.data(), group_offsets.size() * sizeof(cudf::size_type), stream, mr);
 
   auto offsets = std::make_unique<cudf::column>(
     cudf::data_type(cudf::type_to_id<cudf::size_type>()), num_groups + 1, std::move(offsets_data));
diff --git a/cpp/src/groupby/sort/group_count.cu b/cpp/src/groupby/sort/group_count.cu
index 352fc841d11..60e0ce31db1 100644
--- a/cpp/src/groupby/sort/group_count.cu
+++ b/cpp/src/groupby/sort/group_count.cu
@@ -18,6 +18,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
@@ -29,7 +30,7 @@ namespace cudf {
 namespace groupby {
 namespace detail {
 std::unique_ptr<column> group_count_valid(column_view const& values,
-                                          rmm::device_vector<size_type> const& group_labels,
+                                          cudf::device_span<size_type const> group_labels,
                                           size_type num_groups,
                                           rmm::cuda_stream_view stream,
                                           rmm::mr::device_memory_resource* mr)
@@ -70,7 +71,7 @@ std::unique_ptr<column> group_count_valid(column_view const& values,
   return result;
 }
 
-std::unique_ptr<column> group_count_all(rmm::device_vector<size_type> const& group_offsets,
+std::unique_ptr<column> group_count_all(cudf::device_span<size_type const> group_offsets,
                                         size_type num_groups,
                                         rmm::cuda_stream_view stream,
                                         rmm::mr::device_memory_resource* mr)
diff --git a/cpp/src/groupby/sort/group_max.cu b/cpp/src/groupby/sort/group_max.cu
index 06aa172d125..bd4e676b83d 100644
--- a/cpp/src/groupby/sort/group_max.cu
+++ b/cpp/src/groupby/sort/group_max.cu
@@ -23,7 +23,7 @@ namespace groupby {
 namespace detail {
 std::unique_ptr<column> group_max(column_view const& values,
                                   size_type num_groups,
-                                  rmm::device_vector<size_type> const& group_labels,
+                                  cudf::device_span<size_type const> group_labels,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr)
 {
diff --git a/cpp/src/groupby/sort/group_min.cu b/cpp/src/groupby/sort/group_min.cu
index 72bc3e6ba3d..a7c84ac639e 100644
--- a/cpp/src/groupby/sort/group_min.cu
+++ b/cpp/src/groupby/sort/group_min.cu
@@ -23,7 +23,7 @@ namespace groupby {
 namespace detail {
 std::unique_ptr<column> group_min(column_view const& values,
                                   size_type num_groups,
-                                  rmm::device_vector<size_type> const& group_labels,
+                                  cudf::device_span<size_type const> group_labels,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr)
 {
diff --git a/cpp/src/groupby/sort/group_nth_element.cu b/cpp/src/groupby/sort/group_nth_element.cu
index 95375f44605..5c8e8b790d4 100644
--- a/cpp/src/groupby/sort/group_nth_element.cu
+++ b/cpp/src/groupby/sort/group_nth_element.cu
@@ -23,6 +23,7 @@
 #include <cudf/detail/gather.cuh>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 
@@ -31,8 +32,8 @@ namespace groupby {
 namespace detail {
 std::unique_ptr<column> group_nth_element(column_view const &values,
                                           column_view const &group_sizes,
-                                          rmm::device_vector<size_type> const &group_labels,
-                                          rmm::device_vector<size_type> const &group_offsets,
+                                          cudf::device_span<size_type const> group_labels,
+                                          cudf::device_span<size_type const> group_offsets,
                                           size_type num_groups,
                                           size_type n,
                                           null_policy null_handling,
diff --git a/cpp/src/groupby/sort/group_nunique.cu b/cpp/src/groupby/sort/group_nunique.cu
index b14beb42435..09bbf13e8d4 100644
--- a/cpp/src/groupby/sort/group_nunique.cu
+++ b/cpp/src/groupby/sort/group_nunique.cu
@@ -19,6 +19,7 @@
 #include <cudf/detail/iterator.cuh>
 #include <cudf/table/row_operators.cuh>
 #include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
@@ -34,9 +35,9 @@ struct nunique_functor {
   template <typename T>
   typename std::enable_if_t<cudf::is_equality_comparable<T, T>(), std::unique_ptr<column>>
   operator()(column_view const& values,
-             rmm::device_vector<size_type> const& group_labels,
+             cudf::device_span<size_type const> group_labels,
              size_type const num_groups,
-             rmm::device_vector<size_type> const& group_offsets,
+             cudf::device_span<size_type const> group_offsets,
              null_policy null_handling,
              rmm::cuda_stream_view stream,
              rmm::mr::device_memory_resource* mr)
@@ -54,8 +55,8 @@ struct nunique_functor {
         [v = *values_view,
          equal,
          null_handling,
-         group_offsets = group_offsets.data().get(),
-         group_labels  = group_labels.data().get()] __device__(auto i) -> size_type {
+         group_offsets = group_offsets.data(),
+         group_labels  = group_labels.data()] __device__(auto i) -> size_type {
           bool is_input_countable =
             (null_handling == null_policy::INCLUDE || v.is_valid_nocheck(i));
           bool is_unique = is_input_countable &&
@@ -76,8 +77,8 @@ struct nunique_functor {
         thrust::make_counting_iterator<size_type>(0),
         [v = *values_view,
          equal,
-         group_offsets = group_offsets.data().get(),
-         group_labels  = group_labels.data().get()] __device__(auto i) -> size_type {
+         group_offsets = group_offsets.data(),
+         group_labels  = group_labels.data()] __device__(auto i) -> size_type {
           bool is_unique = group_offsets[group_labels[i]] == i ||  // first element or
                            (not equal.operator()<T>(i, i - 1));    // new unique value in sorted
           return static_cast<size_type>(is_unique);
@@ -95,9 +96,9 @@ struct nunique_functor {
   template <typename T>
   typename std::enable_if_t<!cudf::is_equality_comparable<T, T>(), std::unique_ptr<column>>
   operator()(column_view const& values,
-             rmm::device_vector<size_type> const& group_labels,
+             cudf::device_span<size_type const> group_labels,
              size_type const num_groups,
-             rmm::device_vector<size_type> const& group_offsets,
+             cudf::device_span<size_type const> group_offsets,
              null_policy null_handling,
              rmm::cuda_stream_view stream,
              rmm::mr::device_memory_resource* mr)
@@ -107,9 +108,9 @@ struct nunique_functor {
 };
 }  // namespace
 std::unique_ptr<column> group_nunique(column_view const& values,
-                                      rmm::device_vector<size_type> const& group_labels,
+                                      cudf::device_span<size_type const> group_labels,
                                       size_type const num_groups,
-                                      rmm::device_vector<size_type> const& group_offsets,
+                                      cudf::device_span<size_type const> group_offsets,
                                       null_policy null_handling,
                                       rmm::cuda_stream_view stream,
                                       rmm::mr::device_memory_resource* mr)
diff --git a/cpp/src/groupby/sort/group_quantiles.cu b/cpp/src/groupby/sort/group_quantiles.cu
index d5af2de0f7a..fcadb2e71fb 100644
--- a/cpp/src/groupby/sort/group_quantiles.cu
+++ b/cpp/src/groupby/sort/group_quantiles.cu
@@ -20,6 +20,7 @@
 #include <cudf/detail/aggregation/aggregation.hpp>
 #include <cudf/dictionary/detail/iterator.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <groupby/sort/group_reductions.hpp>
 #include <quantiles/quantiles_util.hpp>
@@ -77,7 +78,7 @@ struct quantiles_functor {
   std::enable_if_t<std::is_arithmetic<T>::value, std::unique_ptr<column>> operator()(
     column_view const& values,
     column_view const& group_sizes,
-    rmm::device_vector<size_type> const& group_offsets,
+    cudf::device_span<size_type const> group_offsets,
     size_type const num_groups,
     rmm::device_vector<double> const& quantile,
     interpolation interpolation,
@@ -110,7 +111,7 @@ struct quantiles_functor {
                            values_iter,
                            *group_size_view,
                            *result_view,
-                           group_offsets.data().get(),
+                           group_offsets.data(),
                            quantile.data().get(),
                            static_cast<size_type>(quantile.size()),
                            interpolation});
@@ -123,7 +124,7 @@ struct quantiles_functor {
                            values_iter,
                            *group_size_view,
                            *result_view,
-                           group_offsets.data().get(),
+                           group_offsets.data(),
                            quantile.data().get(),
                            static_cast<size_type>(quantile.size()),
                            interpolation});
@@ -145,7 +146,7 @@ struct quantiles_functor {
 // TODO: add optional check for is_sorted. Use context.flag_sorted
 std::unique_ptr<column> group_quantiles(column_view const& values,
                                         column_view const& group_sizes,
-                                        rmm::device_vector<size_type> const& group_offsets,
+                                        cudf::device_span<size_type const> group_offsets,
                                         size_type const num_groups,
                                         std::vector<double> const& quantiles,
                                         interpolation interp,
diff --git a/cpp/src/groupby/sort/group_reductions.hpp b/cpp/src/groupby/sort/group_reductions.hpp
index f5c21f1289e..b69fe6a0291 100644
--- a/cpp/src/groupby/sort/group_reductions.hpp
+++ b/cpp/src/groupby/sort/group_reductions.hpp
@@ -18,9 +18,9 @@
 
 #include <cudf/aggregation.hpp>
 #include <cudf/column/column.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_vector.hpp>
 
 #include <memory>
 
@@ -38,7 +38,7 @@ namespace detail {
  */
 std::unique_ptr<column> group_sum(column_view const& values,
                                   size_type num_groups,
-                                  rmm::device_vector<size_type> const& group_labels,
+                                  cudf::device_span<size_type const> group_labels,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr);
 
@@ -53,7 +53,7 @@ std::unique_ptr<column> group_sum(column_view const& values,
  */
 std::unique_ptr<column> group_min(column_view const& values,
                                   size_type num_groups,
-                                  rmm::device_vector<size_type> const& group_labels,
+                                  cudf::device_span<size_type const> group_labels,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr);
 
@@ -68,7 +68,7 @@ std::unique_ptr<column> group_min(column_view const& values,
  */
 std::unique_ptr<column> group_max(column_view const& values,
                                   size_type num_groups,
-                                  rmm::device_vector<size_type> const& group_labels,
+                                  cudf::device_span<size_type const> group_labels,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr);
 
@@ -84,7 +84,7 @@ std::unique_ptr<column> group_max(column_view const& values,
  */
 std::unique_ptr<column> group_argmax(column_view const& values,
                                      size_type num_groups,
-                                     rmm::device_vector<size_type> const& group_labels,
+                                     cudf::device_span<size_type const> group_labels,
                                      column_view const& key_sort_order,
                                      rmm::cuda_stream_view stream,
                                      rmm::mr::device_memory_resource* mr);
@@ -101,7 +101,7 @@ std::unique_ptr<column> group_argmax(column_view const& values,
  */
 std::unique_ptr<column> group_argmin(column_view const& values,
                                      size_type num_groups,
-                                     rmm::device_vector<size_type> const& group_labels,
+                                     cudf::device_span<size_type const> group_labels,
                                      column_view const& key_sort_order,
                                      rmm::cuda_stream_view stream,
                                      rmm::mr::device_memory_resource* mr);
@@ -117,7 +117,7 @@ std::unique_ptr<column> group_argmin(column_view const& values,
  * @param stream CUDA stream used for device memory operations and kernel launches.
  */
 std::unique_ptr<column> group_count_valid(column_view const& values,
-                                          rmm::device_vector<size_type> const& group_labels,
+                                          cudf::device_span<size_type const> group_labels,
                                           size_type num_groups,
                                           rmm::cuda_stream_view stream,
                                           rmm::mr::device_memory_resource* mr);
@@ -130,7 +130,7 @@ std::unique_ptr<column> group_count_valid(column_view const& values,
  * @param mr Device memory resource used to allocate the returned column's device memory
  * @param stream CUDA stream used for device memory operations and kernel launches.
  */
-std::unique_ptr<column> group_count_all(rmm::device_vector<size_type> const& group_offsets,
+std::unique_ptr<column> group_count_all(cudf::device_span<size_type const> group_offsets,
                                         size_type num_groups,
                                         rmm::cuda_stream_view stream,
                                         rmm::mr::device_memory_resource* mr);
@@ -150,7 +150,7 @@ std::unique_ptr<column> group_count_all(rmm::device_vector<size_type> const& gro
 std::unique_ptr<column> group_var(column_view const& values,
                                   column_view const& group_means,
                                   column_view const& group_sizes,
-                                  rmm::device_vector<size_type> const& group_labels,
+                                  cudf::device_span<size_type const> group_labels,
                                   size_type ddof,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr);
@@ -168,7 +168,7 @@ std::unique_ptr<column> group_var(column_view const& values,
  */
 std::unique_ptr<column> group_quantiles(column_view const& values,
                                         column_view const& group_sizes,
-                                        rmm::device_vector<size_type> const& group_offsets,
+                                        cudf::device_span<size_type const> group_offsets,
                                         size_type const num_groups,
                                         std::vector<double> const& quantiles,
                                         interpolation interp,
@@ -190,9 +190,9 @@ std::unique_ptr<column> group_quantiles(column_view const& values,
  * @param stream CUDA stream used for device memory operations and kernel launches.
  */
 std::unique_ptr<column> group_nunique(column_view const& values,
-                                      rmm::device_vector<size_type> const& group_labels,
+                                      cudf::device_span<size_type const> group_labels,
                                       size_type const num_groups,
-                                      rmm::device_vector<size_type> const& group_offsets,
+                                      cudf::device_span<size_type const> group_offsets,
                                       null_policy null_handling,
                                       rmm::cuda_stream_view stream,
                                       rmm::mr::device_memory_resource* mr);
@@ -213,8 +213,8 @@ std::unique_ptr<column> group_nunique(column_view const& values,
  */
 std::unique_ptr<column> group_nth_element(column_view const& values,
                                           column_view const& group_sizes,
-                                          rmm::device_vector<size_type> const& group_labels,
-                                          rmm::device_vector<size_type> const& group_offsets,
+                                          cudf::device_span<size_type const> group_labels,
+                                          cudf::device_span<size_type const> group_offsets,
                                           size_type num_groups,
                                           size_type n,
                                           null_policy null_handling,
@@ -230,7 +230,7 @@ std::unique_ptr<column> group_nth_element(column_view const& values,
  * @param stream CUDA stream used for device memory operations and kernel launches.
  */
 std::unique_ptr<column> group_collect(column_view const& values,
-                                      rmm::device_vector<size_type> const& group_offsets,
+                                      cudf::device_span<size_type const> group_offsets,
                                       size_type num_groups,
                                       rmm::cuda_stream_view stream,
                                       rmm::mr::device_memory_resource* mr);
diff --git a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh
index 5b26b7bf108..63a68974d6b 100644
--- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh
+++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh
@@ -25,9 +25,9 @@
 #include <cudf/detail/valid_if.cuh>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
 
 #include <thrust/iterator/discard_iterator.h>
@@ -54,7 +54,7 @@ struct reduce_functor {
   std::enable_if_t<is_supported<T>(), std::unique_ptr<column>> operator()(
     column_view const& values,
     size_type num_groups,
-    rmm::device_vector<cudf::size_type> const& group_labels,
+    cudf::device_span<size_type const> group_labels,
     rmm::cuda_stream_view stream,
     rmm::mr::device_memory_resource* mr)
   {
@@ -87,7 +87,7 @@ struct reduce_functor {
                          values.size(),
                          [d_values     = *valuesview,
                           d_result     = *resultview,
-                          dest_indices = group_labels.data().get()] __device__(auto i) {
+                          dest_indices = group_labels.data()] __device__(auto i) {
                            cudf::detail::update_target_element<DeviceType, K, true, true>{}(
                              d_result, dest_indices[i], d_values, i);
                          });
@@ -97,7 +97,7 @@ struct reduce_functor {
                          values.size(),
                          [d_values     = *valuesview,
                           d_result     = *resultview,
-                          dest_indices = group_labels.data().get()] __device__(auto i) {
+                          dest_indices = group_labels.data()] __device__(auto i) {
                            cudf::detail::update_target_element<dictionary32, K, true, true>{}(
                              d_result, dest_indices[i], d_values, i);
                          });
diff --git a/cpp/src/groupby/sort/group_std.cu b/cpp/src/groupby/sort/group_std.cu
index eb85932d8eb..e49d5ebb4aa 100644
--- a/cpp/src/groupby/sort/group_std.cu
+++ b/cpp/src/groupby/sort/group_std.cu
@@ -22,10 +22,10 @@
 #include <cudf/detail/aggregation/aggregation.hpp>
 #include <cudf/dictionary/detail/iterator.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
+#include <cudf/utilities/span.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
 
 #include <thrust/iterator/discard_iterator.h>
@@ -65,7 +65,7 @@ struct var_transform {
 template <typename ResultType, typename Iterator>
 void reduce_by_key_fn(column_device_view const& values,
                       Iterator values_iter,
-                      rmm::device_vector<size_type> const& group_labels,
+                      cudf::device_span<size_type const> group_labels,
                       ResultType const* d_means,
                       size_type const* d_group_sizes,
                       size_type ddof,
@@ -75,7 +75,7 @@ void reduce_by_key_fn(column_device_view const& values,
   auto var_iter = thrust::make_transform_iterator(
     thrust::make_counting_iterator(0),
     var_transform<ResultType, decltype(values_iter)>{
-      values, values_iter, d_means, d_group_sizes, group_labels.data().get(), ddof});
+      values, values_iter, d_means, d_group_sizes, group_labels.data(), ddof});
 
   thrust::reduce_by_key(rmm::exec_policy(stream),
                         group_labels.begin(),
@@ -91,7 +91,7 @@ struct var_functor {
     column_view const& values,
     column_view const& group_means,
     column_view const& group_sizes,
-    rmm::device_vector<size_type> const& group_labels,
+    cudf::device_span<size_type const> group_labels,
     size_type ddof,
     rmm::cuda_stream_view stream,
     rmm::mr::device_memory_resource* mr)
@@ -110,7 +110,7 @@ struct var_functor {
     auto values_view = column_device_view::create(values, stream);
     auto d_values    = *values_view;
 
-    auto d_group_labels = group_labels.data().get();
+    auto d_group_labels = group_labels.data();
     auto d_means        = group_means.data<ResultType>();
     auto d_group_sizes  = group_sizes.data<size_type>();
     auto d_result       = result->mutable_view().data<ResultType>();
@@ -157,7 +157,7 @@ struct var_functor {
 std::unique_ptr<column> group_var(column_view const& values,
                                   column_view const& group_means,
                                   column_view const& group_sizes,
-                                  rmm::device_vector<size_type> const& group_labels,
+                                  cudf::device_span<size_type const> group_labels,
                                   size_type ddof,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr)
diff --git a/cpp/src/groupby/sort/group_sum.cu b/cpp/src/groupby/sort/group_sum.cu
index f0b50e910c4..e9e6e985c54 100644
--- a/cpp/src/groupby/sort/group_sum.cu
+++ b/cpp/src/groupby/sort/group_sum.cu
@@ -15,6 +15,7 @@
  */
 
 #include <cudf/dictionary/dictionary_column_view.hpp>
+#include <cudf/utilities/span.hpp>
 #include <groupby/sort/group_single_pass_reduction_util.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -24,7 +25,7 @@ namespace groupby {
 namespace detail {
 std::unique_ptr<column> group_sum(column_view const& values,
                                   size_type num_groups,
-                                  rmm::device_vector<size_type> const& group_labels,
+                                  cudf::device_span<size_type const> group_labels,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr)
 {
diff --git a/cpp/src/groupby/sort/sort_helper.cu b/cpp/src/groupby/sort/sort_helper.cu
index 8f3070c3497..6a9da36e21b 100644
--- a/cpp/src/groupby/sort/sort_helper.cu
+++ b/cpp/src/groupby/sort/sort_helper.cu
@@ -35,6 +35,7 @@
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/sequence.h>
+#include <thrust/uninitialized_fill.h>
 #include <thrust/unique.h>
 
 #include <algorithm>
@@ -160,7 +161,7 @@ sort_groupby_helper::index_vector const& sort_groupby_helper::group_offsets(
 {
   if (_group_offsets) return *_group_offsets;
 
-  _group_offsets = std::make_unique<index_vector>(num_keys(stream) + 1);
+  _group_offsets = std::make_unique<index_vector>(num_keys(stream) + 1, stream);
 
   auto device_input_table = table_device_view::create(_keys, stream);
   auto sorted_order       = key_sort_order().data<size_type>();
@@ -182,9 +183,9 @@ sort_groupby_helper::index_vector const& sort_groupby_helper::group_offsets(
       permuted_row_equality_comparator<false>(*device_input_table, sorted_order));
   }
 
-  size_type num_groups          = thrust::distance(_group_offsets->begin(), result_end);
-  (*_group_offsets)[num_groups] = num_keys(stream);
-  _group_offsets->resize(num_groups + 1);
+  size_type num_groups = thrust::distance(_group_offsets->begin(), result_end);
+  _group_offsets->set_element(num_groups, num_keys(stream), stream);
+  _group_offsets->resize(num_groups + 1, stream);
 
   return *_group_offsets;
 }
@@ -195,12 +196,16 @@ sort_groupby_helper::index_vector const& sort_groupby_helper::group_labels(
   if (_group_labels) return *_group_labels;
 
   // Get group labels for future use in segmented sorting
-  _group_labels = std::make_unique<index_vector>(num_keys(stream));
+  _group_labels = std::make_unique<index_vector>(num_keys(stream), stream);
 
   auto& group_labels = *_group_labels;
 
   if (num_keys(stream) == 0) return group_labels;
 
+  thrust::uninitialized_fill(rmm::exec_policy(stream),
+                             group_labels.begin(),
+                             group_labels.end(),
+                             index_vector::value_type{0});
   thrust::scatter(rmm::exec_policy(stream),
                   thrust::make_constant_iterator(1, decltype(num_groups())(1)),
                   thrust::make_constant_iterator(1, num_groups()),
@@ -221,7 +226,7 @@ column_view sort_groupby_helper::unsorted_keys_labels(rmm::cuda_stream_view stre
     data_type(type_to_id<size_type>()), _keys.num_rows(), mask_state::ALL_NULL, stream);
 
   auto group_labels_view = cudf::column_view(
-    data_type(type_to_id<size_type>()), group_labels().size(), group_labels().data().get());
+    data_type(type_to_id<size_type>()), group_labels().size(), group_labels().data());
 
   auto scatter_map = key_sort_order();
 
diff --git a/cpp/src/rolling/grouped_rolling.cu b/cpp/src/rolling/grouped_rolling.cu
index 2dca4f608fd..135df6bdfe2 100644
--- a/cpp/src/rolling/grouped_rolling.cu
+++ b/cpp/src/rolling/grouped_rolling.cu
@@ -110,7 +110,7 @@ std::unique_ptr<column> grouped_rolling_window(table_view const& group_keys,
   using sort_groupby_helper = cudf::groupby::detail::sort::sort_groupby_helper;
 
   sort_groupby_helper helper{group_keys, cudf::null_policy::INCLUDE, cudf::sorted::YES};
-  auto group_offsets{helper.group_offsets()};
+  auto const& group_offsets{helper.group_offsets()};
   auto const& group_labels{helper.group_labels()};
 
   // `group_offsets` are interpreted in adjacent pairs, each pair representing the offsets
@@ -131,8 +131,8 @@ std::unique_ptr<column> grouped_rolling_window(table_view const& group_keys,
          group_offsets[group_offsets.size() - 1] == input.size() &&
          "Must have at least one group.");
 
-  auto preceding_calculator = [d_group_offsets = group_offsets.data().get(),
-                               d_group_labels  = group_labels.data().get(),
+  auto preceding_calculator = [d_group_offsets = group_offsets.data(),
+                               d_group_labels  = group_labels.data(),
                                preceding_window] __device__(size_type idx) {
     auto group_label = d_group_labels[idx];
     auto group_start = d_group_offsets[group_label];
@@ -140,8 +140,8 @@ std::unique_ptr<column> grouped_rolling_window(table_view const& group_keys,
                                         idx - group_start + 1);  // Preceding includes current row.
   };
 
-  auto following_calculator = [d_group_offsets = group_offsets.data().get(),
-                               d_group_labels  = group_labels.data().get(),
+  auto following_calculator = [d_group_offsets = group_offsets.data(),
+                               d_group_labels  = group_labels.data(),
                                following_window] __device__(size_type idx) {
     auto group_label = d_group_labels[idx];
     auto group_end =
@@ -152,10 +152,10 @@ std::unique_ptr<column> grouped_rolling_window(table_view const& group_keys,
 
   if (aggr->kind == aggregation::CUDA || aggr->kind == aggregation::PTX) {
     cudf::detail::preceding_window_wrapper grouped_preceding_window{
-      group_offsets.data().get(), group_labels.data().get(), preceding_window};
+      group_offsets.data(), group_labels.data(), preceding_window};
 
     cudf::detail::following_window_wrapper grouped_following_window{
-      group_offsets.data().get(), group_labels.data().get(), following_window};
+      group_offsets.data(), group_labels.data(), following_window};
 
     return cudf::detail::rolling_window_udf(input,
                                             grouped_preceding_window,
@@ -371,7 +371,7 @@ std::unique_ptr<column> time_range_window_ASC(column_view const& input,
 /// If there are no nulls for any given group, (nulls_begin, nulls_end) == (0,0).
 std::tuple<rmm::device_vector<size_type>, rmm::device_vector<size_type>>
 get_null_bounds_for_timestamp_column(column_view const& timestamp_column,
-                                     rmm::device_vector<size_type> const& group_offsets)
+                                     rmm::device_uvector<size_type> const& group_offsets)
 {
   // For each group, the null values are themselves clustered
   // at the beginning or the end of the group.
@@ -392,7 +392,7 @@ get_null_bounds_for_timestamp_column(column_view const& timestamp_column,
       thrust::make_counting_iterator(static_cast<size_type>(0)),
       thrust::make_counting_iterator(static_cast<size_type>(num_groups)),
       [d_timestamps    = *p_timestamps_device_view,
-       d_group_offsets = group_offsets.data().get(),
+       d_group_offsets = group_offsets.data(),
        d_null_start    = null_start.data(),
        d_null_end      = null_end.data()] __device__(auto group_label) {
         auto group_start           = d_group_offsets[group_label];
@@ -434,8 +434,8 @@ get_null_bounds_for_timestamp_column(column_view const& timestamp_column,
 std::unique_ptr<column> time_range_window_ASC(
   column_view const& input,
   column_view const& timestamp_column,
-  rmm::device_vector<cudf::size_type> const& group_offsets,
-  rmm::device_vector<cudf::size_type> const& group_labels,
+  rmm::device_uvector<cudf::size_type> const& group_offsets,
+  rmm::device_uvector<cudf::size_type> const& group_labels,
   TimeT preceding_window,
   bool preceding_window_is_unbounded,
   TimeT following_window,
@@ -450,8 +450,8 @@ std::unique_ptr<column> time_range_window_ASC(
     get_null_bounds_for_timestamp_column(timestamp_column, group_offsets);
 
   auto preceding_calculator =
-    [d_group_offsets = group_offsets.data().get(),
-     d_group_labels  = group_labels.data().get(),
+    [d_group_offsets = group_offsets.data(),
+     d_group_labels  = group_labels.data(),
      d_timestamps    = timestamp_column.data<TimeT>(),
      d_nulls_begin   = null_start.data().get(),
      d_nulls_end     = null_end.data().get(),
@@ -490,8 +490,8 @@ std::unique_ptr<column> time_range_window_ASC(
   auto preceding_column = expand_to_column(preceding_calculator, input.size(), stream, mr);
 
   auto following_calculator =
-    [d_group_offsets = group_offsets.data().get(),
-     d_group_labels  = group_labels.data().get(),
+    [d_group_offsets = group_offsets.data(),
+     d_group_labels  = group_labels.data(),
      d_timestamps    = timestamp_column.data<TimeT>(),
      d_nulls_begin   = null_start.data().get(),
      d_nulls_end     = null_end.data().get(),
@@ -633,8 +633,8 @@ std::unique_ptr<column> time_range_window_DESC(column_view const& input,
 std::unique_ptr<column> time_range_window_DESC(
   column_view const& input,
   column_view const& timestamp_column,
-  rmm::device_vector<cudf::size_type> const& group_offsets,
-  rmm::device_vector<cudf::size_type> const& group_labels,
+  rmm::device_uvector<cudf::size_type> const& group_offsets,
+  rmm::device_uvector<cudf::size_type> const& group_labels,
   TimeT preceding_window,
   bool preceding_window_is_unbounded,
   TimeT following_window,
@@ -649,8 +649,8 @@ std::unique_ptr<column> time_range_window_DESC(
     get_null_bounds_for_timestamp_column(timestamp_column, group_offsets);
 
   auto preceding_calculator =
-    [d_group_offsets = group_offsets.data().get(),
-     d_group_labels  = group_labels.data().get(),
+    [d_group_offsets = group_offsets.data(),
+     d_group_labels  = group_labels.data(),
      d_timestamps    = timestamp_column.data<TimeT>(),
      d_nulls_begin   = null_start.data().get(),
      d_nulls_end     = null_end.data().get(),
@@ -691,8 +691,8 @@ std::unique_ptr<column> time_range_window_DESC(
   auto preceding_column = expand_to_column(preceding_calculator, input.size(), stream, mr);
 
   auto following_calculator =
-    [d_group_offsets = group_offsets.data().get(),
-     d_group_labels  = group_labels.data().get(),
+    [d_group_offsets = group_offsets.data(),
+     d_group_labels  = group_labels.data(),
      d_timestamps    = timestamp_column.data<TimeT>(),
      d_nulls_begin   = null_start.data().get(),
      d_nulls_end     = null_end.data().get(),
@@ -745,8 +745,8 @@ std::unique_ptr<column> grouped_time_range_rolling_window_impl(
   column_view const& input,
   column_view const& timestamp_column,
   cudf::order const& timestamp_ordering,
-  rmm::device_vector<cudf::size_type> const& group_offsets,
-  rmm::device_vector<cudf::size_type> const& group_labels,
+  rmm::device_uvector<cudf::size_type> const& group_offsets,
+  rmm::device_uvector<cudf::size_type> const& group_labels,
   window_bounds preceding_window_in_days,  // TODO: Consider taking offset-type as type_id. Assumes
                                            // days for now.
   window_bounds following_window_in_days,
@@ -758,7 +758,7 @@ std::unique_ptr<column> grouped_time_range_rolling_window_impl(
   TimeT mult_factor{static_cast<TimeT>(multiplication_factor(timestamp_column.type()))};
 
   if (timestamp_ordering == cudf::order::ASCENDING) {
-    return group_offsets.empty()
+    return group_offsets.is_empty()
              ? time_range_window_ASC(input,
                                      timestamp_column,
                                      preceding_window_in_days.value * mult_factor,
@@ -782,7 +782,7 @@ std::unique_ptr<column> grouped_time_range_rolling_window_impl(
                                      stream,
                                      mr);
   } else {
-    return group_offsets.empty()
+    return group_offsets.is_empty()
              ? time_range_window_DESC(input,
                                       timestamp_column,
                                       preceding_window_in_days.value * mult_factor,
@@ -835,11 +835,11 @@ std::unique_ptr<column> grouped_time_range_rolling_window(table_view const& grou
   using sort_groupby_helper = cudf::groupby::detail::sort::sort_groupby_helper;
   using index_vector        = sort_groupby_helper::index_vector;
 
-  index_vector group_offsets, group_labels;
+  index_vector group_offsets(0, stream), group_labels(0, stream);
   if (group_keys.num_columns() > 0) {
     sort_groupby_helper helper{group_keys, cudf::null_policy::INCLUDE, cudf::sorted::YES};
-    group_offsets = helper.group_offsets();
-    group_labels  = helper.group_labels();
+    group_offsets = index_vector(helper.group_offsets(), stream);
+    group_labels  = index_vector(helper.group_labels(), stream);
   }
 
   // Assumes that `timestamp_column` is actually of a timestamp type.