diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh
index 32f6cc6db7a..d737af9c532 100644
--- a/cpp/include/cudf/lists/detail/scatter.cuh
+++ b/cpp/include/cudf/lists/detail/scatter.cuh
@@ -19,8 +19,8 @@
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/copying.hpp>
+#include <cudf/detail/get_value.cuh>
 #include <cudf/detail/valid_if.cuh>
-#include <cudf/lists/detail/utilities.cuh>
 #include <cudf/lists/list_device_view.cuh>
 #include <cudf/null_mask.hpp>
 #include <cudf/strings/detail/utilities.cuh>
@@ -333,7 +333,8 @@ struct list_child_constructor {
     auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view);
     auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view);
 
-    auto const num_child_rows{get_num_child_rows(list_offsets, stream)};
+    auto const num_child_rows{
+      cudf::detail::get_value<size_type>(list_offsets, list_offsets.size() - 1, stream)};
 
     auto const child_null_mask =
       source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable()
@@ -427,7 +428,8 @@ struct list_child_constructor {
     auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view);
     auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view);
 
-    int32_t num_child_rows{get_num_child_rows(list_offsets, stream)};
+    auto const num_child_rows{
+      cudf::detail::get_value<size_type>(list_offsets, list_offsets.size() - 1, stream)};
 
     auto string_views = rmm::device_vector<string_view>(num_child_rows);
 
@@ -516,7 +518,8 @@ struct list_child_constructor {
     auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view);
     auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view);
 
-    auto num_child_rows = get_num_child_rows(list_offsets, stream);
+    auto const num_child_rows{
+      cudf::detail::get_value<size_type>(list_offsets, list_offsets.size() - 1, stream)};
 
     auto child_list_views = rmm::device_uvector<unbound_list_view>(num_child_rows, stream, mr);
 
@@ -621,7 +624,8 @@ struct list_child_constructor {
     auto const source_structs = source_lists_column_view.child();
     auto const target_structs = target_lists_column_view.child();
 
-    auto const num_child_rows = get_num_child_rows(list_offsets, stream);
+    auto const num_child_rows{
+      cudf::detail::get_value<size_type>(list_offsets, list_offsets.size() - 1, stream)};
 
     auto const num_struct_members =
       std::distance(source_structs.child_begin(), source_structs.child_end());
diff --git a/cpp/include/cudf/lists/detail/utilities.cuh b/cpp/include/cudf/lists/detail/utilities.cuh
deleted file mode 100644
index ccee9b0d5d9..00000000000
--- a/cpp/include/cudf/lists/detail/utilities.cuh
+++ /dev/null
@@ -1,46 +0,0 @@
-/*
- * Copyright (c) 2021, NVIDIA CORPORATION.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-#pragma once
-
-#include <cudf/column/column_view.hpp>
-#include <rmm/cuda_stream_view.hpp>
-
-namespace cudf {
-namespace detail {
-
-/**
- * @brief Fetch the number of rows in a lists column's child given its offsets column.
- *
- * @param[in] list_offsets Offsets child of a lists column
- * @param[in] stream The cuda-stream to synchronize on, when reading from device memory
- * @return cudf::size_type The number of child rows in the lists column
- */
-static cudf::size_type get_num_child_rows(cudf::column_view const& list_offsets,
-                                          rmm::cuda_stream_view stream)
-{
-  // Number of rows in child-column == last offset value.
-  cudf::size_type num_child_rows{};
-  CUDA_TRY(cudaMemcpyAsync(&num_child_rows,
-                           list_offsets.data<cudf::size_type>() + list_offsets.size() - 1,
-                           sizeof(cudf::size_type),
-                           cudaMemcpyDeviceToHost,
-                           stream.value()));
-  stream.synchronize();
-  return num_child_rows;
-}
-
-}  // namespace detail
-}  // namespace cudf
diff --git a/cpp/src/rolling/rolling_detail.cuh b/cpp/src/rolling/rolling_detail.cuh
index 8a0f5f8002d..2ede50b468a 100644
--- a/cpp/src/rolling/rolling_detail.cuh
+++ b/cpp/src/rolling/rolling_detail.cuh
@@ -29,6 +29,7 @@
 #include <cudf/detail/aggregation/aggregation.hpp>
 #include <cudf/detail/copy.hpp>
 #include <cudf/detail/gather.hpp>
+#include <cudf/detail/get_value.cuh>
 #include <cudf/detail/groupby/sort_helper.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
@@ -36,7 +37,6 @@
 #include <cudf/detail/valid_if.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <cudf/dictionary/dictionary_factories.hpp>
-#include <cudf/lists/detail/utilities.cuh>
 #include <cudf/rolling.hpp>
 #include <cudf/strings/detail/utilities.cuh>
 #include <cudf/types.hpp>
@@ -983,7 +983,8 @@ struct rolling_window_launcher {
     // This accounts for the `0` added by default to the offsets
     // column, marking the beginning of the column.
 
-    auto num_child_rows = get_num_child_rows(offsets, stream);
+    auto const num_child_rows{
+      cudf::detail::get_value<size_type>(offsets, offsets.size() - 1, stream)};
 
     auto scatter_values =
       make_fixed_width_column(size_data_type, offsets.size(), mask_state::UNALLOCATED, stream, mr);