From ef87ffd25371833e51df9ad99256a2a95c7893c1 Mon Sep 17 00:00:00 2001
From: David Wendt <dwendt@nvidia.com>
Date: Wed, 7 Feb 2024 19:52:08 -0500
Subject: [PATCH 1/3] Use offsetalator in cudf::row_bit_count()

---
 cpp/src/transform/row_bit_count.cu | 10 ++++++----
 1 file changed, 6 insertions(+), 4 deletions(-)

diff --git a/cpp/src/transform/row_bit_count.cu b/cpp/src/transform/row_bit_count.cu
index a91dc8fbbc6..17409eb8759 100644
--- a/cpp/src/transform/row_bit_count.cu
+++ b/cpp/src/transform/row_bit_count.cu
@@ -20,6 +20,7 @@
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/vector_factories.hpp>
+#include <cudf/lists/detail/interleave_columns.hpp>
 #include <cudf/lists/lists_column_view.hpp>
 #include <cudf/structs/structs_column_view.hpp>
 #include <cudf/table/table_device_view.cuh>
@@ -352,11 +353,12 @@ __device__ size_type row_size_functor::operator()<string_view>(column_device_vie
     return 0;
   }
 
-  auto const offsets_size  = sizeof(size_type) * CHAR_BIT;
+  auto const offsets_size =
+    (offsets.type().id() == type_id::INT32 ? sizeof(int32_t) : sizeof(int64_t)) * CHAR_BIT;
   auto const validity_size = col.nullable() ? 1 : 0;
-  auto const chars_size =
-    (offsets.data<size_type>()[row_end] - offsets.data<size_type>()[row_start]) * CHAR_BIT;
-  return ((offsets_size + validity_size) * num_rows) + chars_size;
+  auto const d_offsets     = cudf::detail::input_offsetalator(offsets.head(), offsets.type());
+  auto const chars_size    = (d_offsets[row_end] - d_offsets[row_start]) * CHAR_BIT;
+  return static_cast<size_type>(((offsets_size + validity_size) * num_rows) + chars_size);
 }
 
 /**

From f76c75e006ba142686498d39e438dd1ed6176875 Mon Sep 17 00:00:00 2001
From: David Wendt <dwendt@nvidia.com>
Date: Mon, 12 Feb 2024 11:01:20 -0500
Subject: [PATCH 2/3] fix include statement

---
 cpp/src/transform/row_bit_count.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/src/transform/row_bit_count.cu b/cpp/src/transform/row_bit_count.cu
index 17409eb8759..e4698fb1262 100644
--- a/cpp/src/transform/row_bit_count.cu
+++ b/cpp/src/transform/row_bit_count.cu
@@ -18,9 +18,9 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/offsets_iterator.cuh>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/vector_factories.hpp>
-#include <cudf/lists/detail/interleave_columns.hpp>
 #include <cudf/lists/lists_column_view.hpp>
 #include <cudf/structs/structs_column_view.hpp>
 #include <cudf/table/table_device_view.cuh>

From 75a9a92e29bb3e9913c5850be3ea99cea5ea07ca Mon Sep 17 00:00:00 2001
From: David Wendt <dwendt@nvidia.com>
Date: Wed, 14 Feb 2024 09:52:52 -0500
Subject: [PATCH 3/3] empty; force rebuild