Use CCCL overflow utils instead of custom overflow logic

cpp/include/cudf/detail/aggregation/device_aggregators.cuh

@@ -16,7 +16,7 @@
 #include <cudf/utilities/traits.cuh>
 #include <cudf/utilities/type_dispatcher.hpp>
 
-#include <cuda/std/limits>
+#include <cuda/numeric>
 #include <cuda/std/type_traits>
 
 namespace cudf::detail {
@@ -134,9 +134,7 @@ template <typename Source>
     (cudf::is_fixed_point<Source>() && cudf::has_atomic_support<device_storage_type_t<Source>>()) ||
     cuda::std::is_same_v<Source, numeric::decimal128>)
 struct update_target_element<Source, aggregation::SUM_WITH_OVERFLOW> {
-  using DeviceType               = device_storage_type_t<Source>;
-  static constexpr auto type_max = cuda::std::numeric_limits<DeviceType>::max();
-  static constexpr auto type_min = cuda::std::numeric_limits<DeviceType>::min();
+  using DeviceType = device_storage_type_t<Source>;
 
   __device__ void operator()(mutable_column_device_view target,
                              size_type target_index,
@@ -155,12 +153,9 @@ struct update_target_element<Source, aggregation::SUM_WITH_OVERFLOW> {
       *(overflow_column.data<bool>() + target_index)};
     if (bool_ref.load(cuda::memory_order_relaxed)) { return; }
 
-    // TODO: to be replaced by CCCL equivalents once https://github.com/NVIDIA/cccl/pull/3755 is
-    // ready
-    auto const overflow =
-      source_value > 0 ? old_sum > type_max - source_value : old_sum < type_min - source_value;
-
-    if (overflow) { cudf::detail::atomic_max(&overflow_column.element<bool>(target_index), true); }
+    if (cuda::add_overflow<DeviceType>(old_sum, source_value).overflow) {
+      cudf::detail::atomic_max(&overflow_column.element<bool>(target_index), true);
+    }
   }
 };
 

cpp/include/cudf/fixed_point/fixed_point.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2020-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2020-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -9,6 +9,7 @@
 #include <cudf/fixed_point/temporary.hpp>
 #include <cudf/types.hpp>
 
+#include <cuda/numeric>
 #include <cuda/std/functional>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
@@ -603,8 +604,7 @@ class fixed_point {
 template <typename Rep, typename T>
 CUDF_HOST_DEVICE inline auto addition_overflow(T lhs, T rhs)
 {
-  return rhs > 0 ? lhs > cuda::std::numeric_limits<Rep>::max() - rhs
-                 : lhs < cuda::std::numeric_limits<Rep>::min() - rhs;
+  return cuda::add_overflow<Rep>(lhs, rhs).overflow;
 }
 
 /** @brief Function for identifying integer overflow when subtracting
@@ -618,8 +618,7 @@ CUDF_HOST_DEVICE inline auto addition_overflow(T lhs, T rhs)
 template <typename Rep, typename T>
 CUDF_HOST_DEVICE inline auto subtraction_overflow(T lhs, T rhs)
 {
-  return rhs > 0 ? lhs < cuda::std::numeric_limits<Rep>::min() + rhs
-                 : lhs > cuda::std::numeric_limits<Rep>::max() + rhs;
+  return cuda::sub_overflow<Rep>(lhs, rhs).overflow;
 }
 
 /** @brief Function for identifying integer overflow when dividing
@@ -633,7 +632,7 @@ CUDF_HOST_DEVICE inline auto subtraction_overflow(T lhs, T rhs)
 template <typename Rep, typename T>
 CUDF_HOST_DEVICE inline auto division_overflow(T lhs, T rhs)
 {
-  return lhs == cuda::std::numeric_limits<Rep>::min() && rhs == -1;
+  return cuda::div_overflow<Rep>(lhs, rhs).overflow;
 }
 
 /** @brief Function for identifying integer overflow when multiplying
@@ -647,11 +646,7 @@ CUDF_HOST_DEVICE inline auto division_overflow(T lhs, T rhs)
 template <typename Rep, typename T>
 CUDF_HOST_DEVICE inline auto multiplication_overflow(T lhs, T rhs)
 {
-  auto const min = cuda::std::numeric_limits<Rep>::min();
-  auto const max = cuda::std::numeric_limits<Rep>::max();
-  if (rhs > 0) { return lhs > max / rhs || lhs < min / rhs; }
-  if (rhs < -1) { return lhs > min / rhs || lhs < max / rhs; }
-  return rhs == -1 && lhs == min;
+  return cuda::mul_overflow<Rep>(lhs, rhs).overflow;
 }
 
 // PLUS Operation

cpp/src/bitmask/null_mask.cu

@@ -25,11 +25,11 @@
 #include <cooperative_groups/reduce.h>
 #include <cub/cub.cuh>
 #include <cuda/atomic>
+#include <cuda/numeric>
 #include <thrust/execution_policy.h>
 #include <thrust/tabulate.h>
 
 #include <algorithm>
-#include <limits>
 #include <numeric>
 
 namespace cudf {
@@ -211,7 +211,7 @@ void set_null_masks(cudf::host_span<bitmask_type*> bitmasks,
       auto const num_words =
         num_bitmask_words(end_bits[i]) - begin_bits[i] / detail::size_in_bits<bitmask_type>();
       // Handle overflow if any
-      if (num_words >= std::numeric_limits<size_t>::max() - cumulative_null_mask_words) {
+      if (cuda::add_overflow<size_t>(cumulative_null_mask_words, num_words).overflow) {
         average_nullmask_words +=
           cudf::util::div_rounding_up_safe<size_t>(cumulative_null_mask_words, num_bitmasks);
         cumulative_null_mask_words = 0;

cpp/src/io/orc/orc.cpp

@@ -12,9 +12,9 @@
 #include <cudf/io/orc.hpp>
 #include <cudf/lists/lists_column_view.hpp>
 
+#include <cuda/numeric>
 #include <thrust/tabulate.h>
 
-#include <limits>
 #include <stdexcept>
 #include <string>
 
@@ -432,9 +432,9 @@ host_span<uint8_t const> orc_decompressor::decompress_blocks(host_span<uint8_t c
       // Uncompressed block
       max_dst_length += block_len;
     } else {
-      CUDF_EXPECTS(max_dst_length <= std::numeric_limits<size_t>::max() - m_blockSize,
-                   "ORC decompression: compression block size overflow");
-      max_dst_length += m_blockSize;
+      auto const next = cuda::add_overflow<size_t>(max_dst_length, m_blockSize);
+      CUDF_EXPECTS(!next.overflow, "ORC decompression: compression block size overflow");
+      max_dst_length = next.value;
     }
     i += block_len;
     CUDF_EXPECTS(i <= src.size() and block_len <= m_blockSize, "Error in decompression");

cpp/src/io/parquet/reader_impl_helpers.cpp

@@ -20,6 +20,7 @@
 #include <cudf/logger.hpp>
 
 #include <cuda/iterator>
+#include <cuda/numeric>
 #include <cuda/std/tuple>
 #include <thrust/iterator/zip_iterator.h>
 
@@ -1368,36 +1369,38 @@ std::vector<std::vector<size_type>> aggregate_reader_metadata::apply_byte_bounds
   auto filtered_row_group_indices =
     std::vector<std::vector<size_type>>(input_row_group_indices.size());
 
-  std::for_each(
-    input_row_group_indices.front().begin(),
-    input_row_group_indices.front().end(),
-    [&](auto const& rg_idx) {
-      // Get the file offset of this row group
-      auto const row_group_file_offset = [&]() {
-        auto const& rg = per_file_metadata.front().row_groups[rg_idx];
-        if (rg.file_offset.has_value()) {
-          return rg.file_offset.value();
-        } else if (rg.columns.front().file_offset != 0) {
-          return rg.columns.front().file_offset;
-        } else {
-          auto const& col_meta = rg.columns.front().meta_data;
-          return col_meta.dictionary_page_offset != 0
-                   ? std::min(col_meta.dictionary_page_offset, col_meta.data_page_offset)
-                   : col_meta.data_page_offset;
-        }
-      }();
-
-      // Check if the row group starts within the byte range: row group file offset is >=
-      // bytes_to_skip AND (bytes_to_read is not specified OR the max byte offset overflows
-      // size_t OR row group file offset is < bytes_to_skip + bytes_to_read)
-      auto const is_within_byte_range =
-        std::cmp_greater_equal(row_group_file_offset, bytes_to_skip) and
-        (not bytes_to_read.has_value() or
-         (std::numeric_limits<size_t>::max() - bytes_to_read.value() <= bytes_to_skip) or
-         std::cmp_less(row_group_file_offset, bytes_to_skip + bytes_to_read.value()));
-
-      if (is_within_byte_range) { filtered_row_group_indices.front().emplace_back(rg_idx); }
-    });
+  std::for_each(input_row_group_indices.front().begin(),
+                input_row_group_indices.front().end(),
+                [&](auto const& rg_idx) {
+                  // Get the file offset of this row group
+                  auto const row_group_file_offset = [&]() {
+                    auto const& rg = per_file_metadata.front().row_groups[rg_idx];
+                    if (rg.file_offset.has_value()) {
+                      return rg.file_offset.value();
+                    } else if (rg.columns.front().file_offset != 0) {
+                      return rg.columns.front().file_offset;
+                    } else {
+                      auto const& col_meta = rg.columns.front().meta_data;
+                      return col_meta.dictionary_page_offset != 0
+                               ? std::min(col_meta.dictionary_page_offset,
+                                          col_meta.data_page_offset)
+                               : col_meta.data_page_offset;
+                    }
+                  }();
+
+                  // Check if the row group starts within the byte range: row group file offset is
+                  // >= bytes_to_skip AND (bytes_to_read is not specified OR the max byte offset
+                  // overflows size_t OR row group file offset is < bytes_to_skip + bytes_to_read)
+                  auto const is_within_byte_range =
+                    std::cmp_greater_equal(row_group_file_offset, bytes_to_skip) and
+                    (not bytes_to_read.has_value() or
+                     cuda::add_overflow<size_t>(bytes_to_skip, bytes_to_read.value()).overflow or
+                     std::cmp_less(row_group_file_offset, bytes_to_skip + bytes_to_read.value()));
+
+                  if (is_within_byte_range) {
+                    filtered_row_group_indices.front().emplace_back(rg_idx);
+                  }
+                });
 
   return filtered_row_group_indices;
 }