From 9461695daed08b3688443565136c88261b1f179b Mon Sep 17 00:00:00 2001
From: David Wendt <dwendt@nvidia.com>
Date: Thu, 13 Aug 2026 15:24:53 -0400
Subject: [PATCH] Check for output offsets int32 overflow in list APIs

---
 cpp/src/lists/combine/concatenate_rows.cu | 17 +++++++++++------
 cpp/src/lists/interleave_columns.cu       | 11 ++++++-----
 2 files changed, 17 insertions(+), 11 deletions(-)

diff --git a/cpp/src/lists/combine/concatenate_rows.cu b/cpp/src/lists/combine/concatenate_rows.cu
index 2964276a3fbc..d4b7d88d22c8 100644
--- a/cpp/src/lists/combine/concatenate_rows.cu
+++ b/cpp/src/lists/combine/concatenate_rows.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
  */
 
@@ -10,6 +10,7 @@
 #include <cudf/detail/gather.cuh>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/sizes_to_offsets_iterator.cuh>
 #include <cudf/lists/combine.hpp>
 #include <cudf/utilities/default_stream.hpp>
 #include <cudf/utilities/error.hpp>
@@ -110,11 +111,15 @@ generate_regrouped_offsets_and_null_mask(table_device_view const& input,
                                     stream);
 
   // convert to offsets
-  thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
-                         offsets->view().begin<size_type>(),
-                         offsets->view().begin<size_type>() + input.num_rows() + 1,
-                         offsets->mutable_view().begin<size_type>(),
-                         0);
+  auto total_size =
+    cudf::detail::sizes_to_offsets(offsets->view().begin<size_type>(),
+                                   offsets->view().begin<size_type>() + input.num_rows() + 1,
+                                   offsets->mutable_view().begin<size_type>(),
+                                   0,
+                                   stream);
+  CUDF_EXPECTS(total_size <= static_cast<decltype(total_size)>(std::numeric_limits<int32_t>::max()),
+               "Size of offsets exceeds maximum int32 limit",
+               std::overflow_error);
 
   // generate appropriate null mask
   auto [null_mask, null_count] = [&]() {
diff --git a/cpp/src/lists/interleave_columns.cu b/cpp/src/lists/interleave_columns.cu
index 202dc1183418..4361b401603f 100644
--- a/cpp/src/lists/interleave_columns.cu
+++ b/cpp/src/lists/interleave_columns.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
  */
 
@@ -76,10 +76,11 @@ generate_list_offsets_and_validities(table_view const& input,
     }));
 
   // Compute offsets from sizes.
-  thrust::exclusive_scan(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
-                         d_offsets,
-                         d_offsets + num_output_lists + 1,
-                         d_offsets);
+  auto total_size = cudf::detail::sizes_to_offsets(
+    d_offsets, d_offsets + num_output_lists + 1, d_offsets, 0, stream);
+  CUDF_EXPECTS(total_size <= static_cast<decltype(total_size)>(std::numeric_limits<int32_t>::max()),
+               "Size of offsets exceeds maximum int32 limit",
+               std::overflow_error);
 
   return {std::move(list_offsets), std::move(validities)};
 }