Unbind cudf::size_type from offsets used by list columns

cpp/include/cudf/column/column_factories.hpp

@@ -485,7 +485,10 @@ std::unique_ptr<column> make_strings_column(size_type num_strings,
  * @endcode
  *
  * @param num_rows The number of lists the column represents.
- * @param offsets_column The column of offset values for this column. Each value should
+ * @param offsets_column The column of offset values for this column. Must be of type
+ * `type_id::INT32` -- per the Arrow columnar format, a LIST column's offsets are always 32-bit
+ * (a 64-bit offsets variant would be a distinct LARGE_LIST type, which cudf does not have).
+ * This is deliberately independent of `cudf::size_type`. Each value should
  * represent the starting offset into the child elements that corresponds to the beginning of the
  * row, with the first row starting at 0. The length of row N can be determined by subtracting
  * `offsets[N+1] - offsets[N]`. The total number of offsets should be 1 longer than the

cpp/include/cudf/detail/sizes_to_offsets_iterator.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -286,8 +286,10 @@ auto sizes_to_offsets(SizesIterator begin,
  * The return also includes the total number of elements -- the last element value from the
  * scan.
  *
+ * The returned column is always `type_id::INT32` since offsets children are 32-bit.
+ *
  * @throw std::overflow_error if the total size of the scan (last element) greater than maximum
- * value of `size_type`
+ * value of `int32_t`
  *
  * @tparam InputIterator Used as input to scan to set the offset values
  * @param begin The beginning of the input sequence
@@ -303,26 +305,32 @@ std::pair<std::unique_ptr<column>, size_type> make_offsets_child_column(
   rmm::cuda_stream_view stream,
   rmm::device_async_resource_ref mr)
 {
-  auto count          = static_cast<size_type>(std::distance(begin, end));
-  auto offsets_column = make_numeric_column(
-    data_type{type_to_id<size_type>()}, count + 1, mask_state::UNALLOCATED, stream, mr);
+  auto count = static_cast<size_type>(std::distance(begin, end));
+  auto offsets_column =
+    make_numeric_column(data_type{type_id::INT32}, count + 1, mask_state::UNALLOCATED, stream, mr);
   auto offsets_view = offsets_column->mutable_view();
-  auto d_offsets    = offsets_view.template data<size_type>();
+  auto d_offsets    = offsets_view.template data<int32_t>();
 
   // The number of offsets is count+1 so to build the offsets from the sizes
   // using exclusive-scan technically requires count+1 input values even though
   // the final input value is never used.
   // The input iterator is wrapped here to allow the last value to be safely read.
+  // The input sizes are deliberately not narrowed to the 32-bit offsets type here -- doing so
+  // would corrupt individual sizes larger than `int32_t` before they reach the scan, so the
+  // overflow check below could no longer detect the overflow. Narrowing happens only on write
+  // to `d_offsets`, after the accumulated total has been validated.
+  using SizeType = cuda::std::iter_value_t<InputIterator>;
   auto map_fn =
-    cuda::proclaim_return_type<size_type>([begin, count] __device__(size_type idx) -> size_type {
-      return idx < count ? static_cast<size_type>(begin[idx]) : size_type{0};
+    cuda::proclaim_return_type<SizeType>([begin, count] __device__(size_type idx) -> SizeType {
+      return idx < count ? begin[idx] : SizeType{0};
     });
   auto input_itr = cudf::detail::make_counting_transform_iterator(0, map_fn);
   // Use the sizes-to-offsets iterator to compute the total number of elements
   auto const total_elements =
     sizes_to_offsets(input_itr, input_itr + count + 1, d_offsets, 0, stream);
+  // the offsets are 32-bit so the total must fit in an int32_t
   CUDF_EXPECTS(
-    total_elements <= static_cast<decltype(total_elements)>(std::numeric_limits<size_type>::max()),
+    total_elements <= static_cast<decltype(total_elements)>(std::numeric_limits<int32_t>::max()),
     "Size of output exceeds the column size limit",
     std::overflow_error);
 

cpp/include/cudf/lists/detail/gather.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -100,7 +100,7 @@ gather_data make_gather_data(cudf::lists_column_view const& source_column,
   // handle sliced columns
   size_type const shift =
     source_column.offset() > 0
-      ? cudf::detail::get_value<size_type>(source_column.offsets(), source_column.offset(), stream)
+      ? cudf::detail::get_value<int32_t>(source_column.offsets(), source_column.offset(), stream)
       : 0;
 
   // generate the base offsets

cpp/include/cudf/lists/detail/scatter.cuh

@@ -231,12 +231,12 @@ std::unique_ptr<column> scatter(scalar const& slr,
                                    ? cudf::create_null_mask(1, mask_state::UNALLOCATED, stream, mr)
                                    : cudf::create_null_mask(1, mask_state::ALL_NULL, stream, mr);
   auto offset_column =
-    make_numeric_column(data_type{type_to_id<size_type>()}, 2, mask_state::UNALLOCATED, stream, mr);
+    make_numeric_column(data_type{type_id::INT32}, 2, mask_state::UNALLOCATED, stream, mr);
   thrust::sequence(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
-                   offset_column->mutable_view().begin<size_type>(),
-                   offset_column->mutable_view().end<size_type>(),
-                   0,
-                   lv->view().size());
+                   offset_column->mutable_view().begin<int32_t>(),
+                   offset_column->mutable_view().end<int32_t>(),
+                   int32_t{0},
+                   static_cast<int32_t>(lv->view().size()));
   auto wrapped = column_view(data_type{type_id::LIST},
                              1,
                              nullptr,

cpp/include/cudf/lists/list_device_view.cuh

@@ -40,10 +40,10 @@ class list_device_view {
     cudf_assert(row_index >= 0 && row_index < lists_column.size() && row_index < offsets.size() &&
                 "row_index out of bounds");
 
-    begin_offset = offsets.element<size_type>(row_index + lists_column.offset());
+    begin_offset = offsets.element<int32_t>(row_index + lists_column.offset());
     cudf_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() &&
                 "begin_offset out of bounds.");
-    _size = offsets.element<size_type>(row_index + 1 + lists_column.offset()) - begin_offset;
+    _size = offsets.element<int32_t>(row_index + 1 + lists_column.offset()) - begin_offset;
   }
 
   ~list_device_view() = default;

cpp/include/cudf/lists/lists_column_device_view.cuh

@@ -79,7 +79,7 @@ class lists_column_device_view : private column_device_view {
    */
   [[nodiscard]] __device__ inline size_type offset_at(size_type idx) const
   {
-    return offsets().size() > 0 ? offsets().element<size_type>(offset() + idx) : 0;
+    return offsets().size() > 0 ? offsets().element<int32_t>(offset() + idx) : 0;
   }
 
   /**

cpp/include/cudf/lists/lists_column_view.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -62,7 +62,7 @@ class lists_column_view : private column_view {
   using column_view::null_mask;
   using column_view::offset;
   using column_view::size;
-  using offset_iterator = size_type const*;  ///< Iterator type for offsets
+  using offset_iterator = int32_t const*;  ///< Iterator type for offsets
 
   /**
    * @brief Returns the parent column.
@@ -108,7 +108,7 @@ class lists_column_view : private column_view {
    */
   [[nodiscard]] offset_iterator offsets_begin() const noexcept
   {
-    return offsets().begin<size_type>() + offset();
+    return offsets().begin<int32_t>() + offset();
   }
 
   /**

cpp/include/cudf/tdigest/tdigest_column_view.hpp

@@ -66,7 +66,7 @@ class tdigest_column_view : private column_view {
   tdigest_column_view& operator=(tdigest_column_view&&) = default;
 
   using column_view::size;
-  using offset_iterator = size_type const*;  ///< Iterator over offsets
+  using offset_iterator = int32_t const*;  ///< Iterator over offsets
 
   // mean and weight column indices within tdigest inner struct columns
   static constexpr size_type mean_column_index{0};    ///< Mean column index

cpp/include/cudf_test/column_wrapper.hpp

@@ -327,16 +327,15 @@ template <typename StringsIterator, typename ValidityIterator>
 auto make_chars_and_offsets(StringsIterator begin, StringsIterator end, ValidityIterator v)
 {
   std::vector<char> chars{};
-  std::vector<cudf::size_type> offsets(1, 0);
+  std::vector<int32_t> offsets(1, 0);
   for (auto str = begin; str < end; ++str) {
     std::string tmp = (*v++) ? std::string(*str) : std::string{};
     chars.insert(chars.end(), std::cbegin(tmp), std::cend(tmp));
     auto const last_offset = static_cast<std::size_t>(offsets.back());
     auto const next_offset = last_offset + tmp.length();
-    CUDF_EXPECTS(
-      next_offset < static_cast<std::size_t>(std::numeric_limits<cudf::size_type>::max()),
-      "Cannot use strings_column_wrapper to build a large strings column");
-    offsets.push_back(static_cast<cudf::size_type>(next_offset));
+    CUDF_EXPECTS(next_offset < static_cast<std::size_t>(std::numeric_limits<int32_t>::max()),
+                 "Cannot use strings_column_wrapper to build a large strings column");
+    offsets.push_back(static_cast<int32_t>(next_offset));
   }
   return std::pair(std::move(chars), std::move(offsets));
 };
@@ -1771,7 +1770,7 @@ class lists_column_wrapper : public detail::column_wrapper {
     rmm::cuda_stream_view stream = cudf::test::get_default_stream(),
     cudf::memory_resources mr    = cudf::get_current_device_resource_ref())
   {
-    cudf::test::fixed_width_column_wrapper<cudf::size_type> offsets({0, 0}, stream, mr);
+    cudf::test::fixed_width_column_wrapper<int32_t> offsets({0, 0}, stream, mr);
     cudf::test::fixed_width_column_wrapper<int> values{};
     return lists_column_wrapper<T>(
       1,
@@ -1860,7 +1859,7 @@ class lists_column_wrapper : public detail::column_wrapper {
     // add the final offset
     offsetv.push_back(count);
     auto offsets =
-      cudf::test::fixed_width_column_wrapper<size_type>(offsetv.begin(), offsetv.end(), stream, mr)
+      cudf::test::fixed_width_column_wrapper<int32_t>(offsetv.begin(), offsetv.end(), stream, mr)
         .release();
 
     // concatenate them together, skipping children that are null.
@@ -1909,7 +1908,7 @@ class lists_column_wrapper : public detail::column_wrapper {
       offsetv.push_back(c->size());
     }
     auto offsets =
-      cudf::test::fixed_width_column_wrapper<size_type>(offsetv.begin(), offsetv.end(), stream, mr)
+      cudf::test::fixed_width_column_wrapper<int32_t>(offsetv.begin(), offsetv.end(), stream, mr)
         .release();
 
     // construct the list column. mark this as a root

cpp/src/groupby/groupby.cu

@@ -104,7 +104,7 @@ struct empty_column_constructor {
 
     if constexpr (k == aggregation::Kind::HISTOGRAM) {
       return make_lists_column(0,
-                               make_empty_column(type_to_id<size_type>()),
+                               make_empty_column(type_id::INT32),
                                cudf::reduction::detail::make_empty_histogram_like(values),
                                0,
                                {});

cpp/src/groupby/sort/group_collect.cu

@@ -63,7 +63,7 @@ std::pair<std::unique_ptr<column>, std::unique_ptr<column>> purge_null_entries(
     cuda::counting_iterator<size_type>{0},
     cuda::counting_iterator<size_type>{num_groups},
     null_purged_sizes.begin(),
-    [d_offsets = offsets.template begin<size_type>(), not_null_pred] __device__(auto i) {
+    [d_offsets = offsets.template begin<int32_t>(), not_null_pred] __device__(auto i) {
       return thrust::count_if(thrust::seq,
                               cuda::counting_iterator<size_type>{d_offsets[i]},
                               cuda::counting_iterator<size_type>{d_offsets[i + 1]},
@@ -86,12 +86,12 @@ std::unique_ptr<column> group_collect(column_view const& values,
   auto [child_column,
         offsets_column] = [null_handling, num_groups, &values, &group_offsets, stream, mr] {
     auto offsets_column = make_numeric_column(
-      data_type(type_to_id<size_type>()), num_groups + 1, mask_state::UNALLOCATED, stream, mr);
+      data_type(type_id::INT32), num_groups + 1, mask_state::UNALLOCATED, stream, mr);
 
     thrust::copy(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
                  group_offsets.begin(),
                  group_offsets.end(),
-                 offsets_column->mutable_view().template begin<size_type>());
+                 offsets_column->mutable_view().template begin<int32_t>());
 
     // If column of grouped values contains null elements, and null_policy == EXCLUDE,
     // those elements must be filtered out, and offsets recomputed.

cpp/src/groupby/sort/group_merge_lists.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -28,7 +28,7 @@ std::unique_ptr<column> group_merge_lists(column_view const& values,
                "Input to `group_merge_lists` must be a non-nullable lists column.");
 
   auto offsets_column = make_numeric_column(
-    data_type(type_to_id<size_type>()), num_groups + 1, mask_state::UNALLOCATED, stream, mr);
+    data_type(type_id::INT32), num_groups + 1, mask_state::UNALLOCATED, stream, mr);
 
   // Generate offsets of the output lists column by gathering from the provided group offsets and
   // the input list offsets.
@@ -44,7 +44,7 @@ std::unique_ptr<column> group_merge_lists(column_view const& values,
                  group_offsets.begin(),
                  group_offsets.end(),
                  lists_column_view(values).offsets_begin(),
-                 offsets_column->mutable_view().template begin<size_type>());
+                 offsets_column->mutable_view().template begin<int32_t>());
 
   // The child column of the output lists column is just copied from the input column.
   auto child_column =

cpp/src/hash/md5_hash.cu

@@ -323,8 +323,8 @@ std::unique_ptr<column> md5(table_view const& input,
             if (data_col.type().id() == type_id::LIST) {
               CUDF_UNREACHABLE("Nested list unsupported");
             }
-            auto const offset_begin = offsets.element<size_type>(row_index);
-            auto const offset_end   = offsets.element<size_type>(row_index + 1);
+            auto const offset_begin = offsets.element<int32_t>(row_index + col.offset());
+            auto const offset_end   = offsets.element<int32_t>(row_index + 1 + col.offset());
             cudf::type_dispatcher<dispatch_storage_type>(
               data_col.type(), ListHasherDispatcher(&hasher, data_col), offset_begin, offset_end);
           } else {

cpp/src/io/json/host_tree_algorithms.cu

@@ -220,7 +220,7 @@ struct json_column_data {
   using row_offset_t = json_column::row_offset_t;
   row_offset_t* string_offsets;
   row_offset_t* string_lengths;
-  row_offset_t* child_offsets;
+  int32_t* child_offsets;
   bitmask_type* validity;
 };
 
@@ -1161,7 +1161,7 @@ void scatter_offsets(tree_meta_t const& tree,
         col.child_offsets.begin(),
         col.child_offsets.end(),
         col.child_offsets.begin(),
-        cuda::maximum<json_column::row_offset_t>{});
+        cuda::maximum<int32_t>{});
     }
   }
   stream.synchronize();

cpp/src/io/json/nested_json.hpp

@@ -72,8 +72,8 @@ struct json_column {
   std::vector<row_offset_t> string_offsets;
   std::vector<row_offset_t> string_lengths;
 
-  // Row offsets
-  std::vector<row_offset_t> child_offsets;
+  // Row offsets (LIST offsets child data; always 32-bit)
+  std::vector<int32_t> child_offsets;
 
   // Validity bitmap
   std::vector<bitmask_type> validity;
@@ -148,8 +148,8 @@ struct device_json_column {
   rmm::device_uvector<row_offset_t> string_offsets;
   rmm::device_uvector<row_offset_t> string_lengths;
 
-  // Row offsets
-  rmm::device_uvector<row_offset_t> child_offsets;
+  // Row offsets (LIST offsets child data; always 32-bit)
+  rmm::device_uvector<int32_t> child_offsets;
 
   // Validity bitmap
   rmm::device_buffer validity;

cpp/src/io/json/parser_features.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -75,7 +75,7 @@ struct empty_column_functor {
     auto const& child_name        = schema.child_types.begin()->first;
     std::unique_ptr<column> child = cudf::type_dispatcher(
       schema.child_types.at(child_name).type, *this, schema.child_types.at(child_name));
-    auto offsets = make_empty_column(data_type(type_to_id<size_type>()));
+    auto offsets = make_empty_column(data_type(type_id::INT32));
     std::vector<std::unique_ptr<column>> child_columns;
     child_columns.push_back(std::move(offsets));
     child_columns.push_back(std::move(child));
@@ -117,7 +117,7 @@ struct allnull_column_functor {
   [[nodiscard]] auto make_zeroed_offsets(size_type size) const
   {
     auto offsets_buff =
-      cudf::detail::make_zeroed_device_uvector_async<size_type>(size + 1, stream, mr);
+      cudf::detail::make_zeroed_device_uvector_async<int32_t>(size + 1, stream, mr);
     return std::make_unique<column>(std::move(offsets_buff), rmm::device_buffer{}, 0);
   }
 

cpp/src/io/orc/stripe_enc.cu

@@ -845,8 +845,8 @@ CUDF_KERNEL void __launch_bounds__(block_size)
           case MAP: {
             auto const& offsets = column.child(lists_column_view::offsets_column_index);
             // Compute list length from the offsets
-            s->lengths.u32[nz_idx] = offsets.element<size_type>(row + 1 + column.offset()) -
-                                     offsets.element<size_type>(row + column.offset());
+            s->lengths.u32[nz_idx] = offsets.element<int32_t>(row + 1 + column.offset()) -
+                                     offsets.element<int32_t>(row + column.offset());
           } break;
           default: break;
         }

cpp/src/io/orc/writer_impl.cu

@@ -1837,8 +1837,8 @@ void pushdown_lists_null_mask(orc_column_view const& col,
       auto const is_row_valid = d_col.is_valid(idx) and bit_value_or(parent_pd_mask, idx, true);
       if (not is_row_valid) {
         auto offsets                = d_col.child(lists_column_view::offsets_column_index);
-        auto const child_rows_begin = offsets.element<size_type>(idx + d_col.offset());
-        auto const child_rows_end   = offsets.element<size_type>(idx + 1 + d_col.offset());
+        auto const child_rows_begin = offsets.element<int32_t>(idx + d_col.offset());
+        auto const child_rows_end   = offsets.element<int32_t>(idx + 1 + d_col.offset());
         for (auto child_row = child_rows_begin; child_row < child_rows_end; ++child_row)
           clear_bit(out_mask.data(), child_row);
       }
@@ -2087,9 +2087,9 @@ hostdevice_2dvector<rowgroup_rows> calculate_rowgroup_bounds(orc_table_view cons
 
             auto offsets = parent_col.child(lists_column_view::offsets_column_index);
             auto const rows_begin =
-              offsets.element<size_type>(parent_rg.begin + parent_col.offset()) - col.offset();
+              offsets.element<int32_t>(parent_rg.begin + parent_col.offset()) - col.offset();
             auto const rows_end =
-              offsets.element<size_type>(parent_rg.end + parent_col.offset()) - col.offset();
+              offsets.element<int32_t>(parent_rg.end + parent_col.offset()) - col.offset();
 
             return rowgroup_rows{rows_begin, rows_end};
           }

cpp/src/io/statistics/statistics.cuh

@@ -118,7 +118,7 @@ __device__ T get_element(column_device_view const& col, uint32_t row)
 {
   using et              = typename T::element_type;
   size_type const index = row + col.offset();  // account for this view's _offset
-  auto const* d_offsets = col.child(lists_column_view::offsets_column_index).data<size_type>();
+  auto const* d_offsets = col.child(lists_column_view::offsets_column_index).data<int32_t>();
   auto const* d_data    = col.child(lists_column_view::child_column_index).data<et>();
   auto const offset     = d_offsets[index];
   return T(d_data + offset, d_offsets[index + 1] - offset);