Add SUM_OVERFLOW in sort groupby

cpp/CMakeLists.txt

@@ -620,6 +620,7 @@ add_library(
   src/groupby/sort/group_replace_nulls.cu
   src/groupby/sort/group_std.cu
   src/groupby/sort/group_sum.cu
+  src/groupby/sort/group_sum_overflow.cu
   src/groupby/sort/group_sum_scan.cu
   src/groupby/sort/group_topk.cu
   src/groupby/sort/host_udf_aggregation.cpp

cpp/include/cudf/aggregation.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -77,48 +77,50 @@ class aggregation {
    * @brief Possible aggregation operations.
    */
   enum Kind : int32_t {
-    SUM = 0,            ///< sum reduction
-    SUM_WITH_OVERFLOW,  ///< sum reduction with overflow detection
-    PRODUCT,            ///< product reduction
-    MIN,                ///< min reduction
-    MAX,                ///< max reduction
-    COUNT_VALID,        ///< count number of valid elements
-    COUNT_ALL,          ///< count number of elements
-    ANY,                ///< any reduction
-    ALL,                ///< all reduction
-    SUM_OF_SQUARES,     ///< sum of squares reduction
-    MEAN,               ///< arithmetic mean reduction
-    M2,                 ///< sum of squares of differences from the mean
-    VARIANCE,           ///< variance
-    STD,                ///< standard deviation
-    MEDIAN,             ///< median reduction
-    QUANTILE,           ///< compute specified quantile(s)
-    ARGMAX,             ///< Index of max element
-    ARGMIN,             ///< Index of min element
-    NUNIQUE,            ///< count number of unique elements
-    NTH_ELEMENT,        ///< get the nth element
-    ROW_NUMBER,         ///< get row-number of current index (relative to rolling window)
-    EWMA,               ///< get exponential weighted moving average at current index
-    RANK,               ///< get rank of current index
-    COLLECT_LIST,       ///< collect values into a list
-    COLLECT_SET,        ///< collect values into a list without duplicate entries
-    LEAD,               ///< window function, accesses row at specified offset following current row
-    LAG,                ///< window function, accesses row at specified offset preceding current row
-    PTX,                ///< PTX  based UDF aggregation
-    CUDA,               ///< CUDA based UDF aggregation
-    HOST_UDF,           ///< host based UDF aggregation
-    MERGE_LISTS,        ///< merge multiple lists values into one list
-    MERGE_SETS,         ///< merge multiple lists values into one list then drop duplicate entries
-    MERGE_M2,           ///< merge partial values of M2 aggregation,
-    COVARIANCE,         ///< covariance between two sets of elements
-    CORRELATION,        ///< correlation between two sets of elements
-    TDIGEST,            ///< create a tdigest from a set of input values
-    MERGE_TDIGEST,      ///< create a tdigest by merging multiple tdigests together
-    HISTOGRAM,          ///< compute frequency of each element
-    MERGE_HISTOGRAM,    ///< merge partial values of HISTOGRAM aggregation
-    BITWISE_AGG,        ///< bitwise aggregation on numeric columns
-    TOP_K,              ///< top k elements in a group
-    INVALID             ///< invalid aggregation, used as a placeholder when default-constructed
+    SUM = 0,       ///< sum reduction
+    SUM_OVERFLOW,  ///< sum reduction with overflow detection
+    /// @deprecated Use SUM_OVERFLOW instead.
+    SUM_WITH_OVERFLOW = SUM_OVERFLOW,
+    PRODUCT,          ///< product reduction
+    MIN,              ///< min reduction
+    MAX,              ///< max reduction
+    COUNT_VALID,      ///< count number of valid elements
+    COUNT_ALL,        ///< count number of elements
+    ANY,              ///< any reduction
+    ALL,              ///< all reduction
+    SUM_OF_SQUARES,   ///< sum of squares reduction
+    MEAN,             ///< arithmetic mean reduction
+    M2,               ///< sum of squares of differences from the mean
+    VARIANCE,         ///< variance
+    STD,              ///< standard deviation
+    MEDIAN,           ///< median reduction
+    QUANTILE,         ///< compute specified quantile(s)
+    ARGMAX,           ///< Index of max element
+    ARGMIN,           ///< Index of min element
+    NUNIQUE,          ///< count number of unique elements
+    NTH_ELEMENT,      ///< get the nth element
+    ROW_NUMBER,       ///< get row-number of current index (relative to rolling window)
+    EWMA,             ///< get exponential weighted moving average at current index
+    RANK,             ///< get rank of current index
+    COLLECT_LIST,     ///< collect values into a list
+    COLLECT_SET,      ///< collect values into a list without duplicate entries
+    LEAD,             ///< window function, accesses row at specified offset following current row
+    LAG,              ///< window function, accesses row at specified offset preceding current row
+    PTX,              ///< PTX  based UDF aggregation
+    CUDA,             ///< CUDA based UDF aggregation
+    HOST_UDF,         ///< host based UDF aggregation
+    MERGE_LISTS,      ///< merge multiple lists values into one list
+    MERGE_SETS,       ///< merge multiple lists values into one list then drop duplicate entries
+    MERGE_M2,         ///< merge partial values of M2 aggregation,
+    COVARIANCE,       ///< covariance between two sets of elements
+    CORRELATION,      ///< correlation between two sets of elements
+    TDIGEST,          ///< create a tdigest from a set of input values
+    MERGE_TDIGEST,    ///< create a tdigest by merging multiple tdigests together
+    HISTOGRAM,        ///< compute frequency of each element
+    MERGE_HISTOGRAM,  ///< merge partial values of HISTOGRAM aggregation
+    BITWISE_AGG,      ///< bitwise aggregation on numeric columns
+    TOP_K,            ///< top k elements in a group
+    INVALID           ///< invalid aggregation, used as a placeholder when default-constructed
   };
 
   /**
@@ -212,8 +214,14 @@ enum class ewm_history : int32_t { INFINITE, FINITE };
 template <typename Base = aggregation>
 std::unique_ptr<Base> make_sum_aggregation();
 
+/// Factory to create a SUM_OVERFLOW aggregation
+/// @return A SUM_OVERFLOW aggregation object
+template <typename Base = aggregation>
+std::unique_ptr<Base> make_sum_overflow_aggregation();
+
 /// Factory to create a SUM_WITH_OVERFLOW aggregation
 /// @return A SUM_WITH_OVERFLOW aggregation object
+/// @deprecated Use make_sum_overflow_aggregation() instead.
 template <typename Base = aggregation>
 std::unique_ptr<Base> make_sum_with_overflow_aggregation();
 

cpp/include/cudf/detail/aggregation/aggregation.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -977,11 +977,19 @@ struct target_type_impl<Source,
   using type = Source;
 };
 
-// SUM_WITH_OVERFLOW outputs a struct {sum: Source, overflow: bool} where sum type matches input
-// type, only supports signed integral types (excluding bool) and decimal types
+/**
+ * @brief Whether `Source` is a valid input type for the SUM_OVERFLOW aggregation.
+ *
+ * Supports signed integral types (excluding bool) and fixed-point (decimal) types.
+ */
 template <typename Source>
-  requires((cudf::is_integral_not_bool<Source>() && cudf::is_signed<Source>()) ||
-           cudf::is_fixed_point<Source>())
+concept sum_overflow_supported =
+  (cudf::is_integral_not_bool<Source>() && cudf::is_signed<Source>()) ||
+  cudf::is_fixed_point<Source>();
+
+// SUM_WITH_OVERFLOW outputs a struct {sum: Source, overflow: bool} where the sum matches the input
+// type
+template <sum_overflow_supported Source>
 struct target_type_impl<Source, aggregation::SUM_WITH_OVERFLOW> {
   using type = struct_view;  // SUM_WITH_OVERFLOW outputs a struct with sum and overflow fields
 };

cpp/include/cudf/reduction/detail/sum_overflow.cuh

@@ -0,0 +1,66 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#pragma once
+
+#include <cudf/column/column_device_view.cuh>
+#include <cudf/types.hpp>
+
+#include <cuda/numeric>
+
+namespace cudf {
+namespace reduction::detail {
+
+/**
+ * @brief Running accumulator for a sum that detects signed-integer overflow.
+ *
+ * `wraps` is the net number of times the running sum has stepped outside [MIN, MAX].
+ * A final `wraps == 0` means the true sum fits in `DeviceType`, i.e. no overflow.
+ */
+template <typename DeviceType>
+struct sum_overflow_result {
+  DeviceType sum;
+  cudf::size_type wraps;
+
+  CUDF_HOST_DEVICE sum_overflow_result() : sum{0}, wraps{0} {}
+  CUDF_HOST_DEVICE sum_overflow_result(DeviceType s, cudf::size_type w) : sum{s}, wraps{w} {}
+};
+
+/// @brief Associative combine: wrap the sums and track the net carry direction.
+template <typename DeviceType>
+struct overflow_sum_op {
+  __device__ sum_overflow_result<DeviceType> operator()(
+    sum_overflow_result<DeviceType> const& lhs, sum_overflow_result<DeviceType> const& rhs) const
+  {
+    auto const r     = cuda::add_overflow<DeviceType>(lhs.sum, rhs.sum);
+    auto const carry = r.overflow ? (rhs.sum > DeviceType{0} ? 1 : -1) : 0;
+    return sum_overflow_result<DeviceType>{r.value, lhs.wraps + rhs.wraps + carry};
+  }
+};
+
+/// @brief Maps a value to a zero-wrap accumulator.
+template <typename DeviceType>
+struct to_sum_overflow {
+  __device__ sum_overflow_result<DeviceType> operator()(DeviceType value) const
+  {
+    return sum_overflow_result<DeviceType>{value, 0};
+  }
+};
+
+/// @brief Maps a row index to an accumulator, treating nulls as a zero contribution.
+template <typename DeviceType>
+struct null_aware_to_sum_overflow {
+  cudf::column_device_view dcol;
+
+  CUDF_HOST_DEVICE null_aware_to_sum_overflow(cudf::column_device_view const& d) : dcol{d} {}
+
+  __device__ sum_overflow_result<DeviceType> operator()(cudf::size_type idx) const
+  {
+    return dcol.is_valid(idx) ? sum_overflow_result<DeviceType>{dcol.element<DeviceType>(idx), 0}
+                              : sum_overflow_result<DeviceType>{DeviceType{0}, 0};
+  }
+};
+
+}  // namespace reduction::detail
+}  // namespace cudf

cpp/src/aggregation/aggregation.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -35,6 +35,21 @@ template CUDF_EXPORT std::unique_ptr<segmented_reduce_aggregation>
 make_sum_aggregation<segmented_reduce_aggregation>();
 
 /// Factory to create a SUM_WITH_OVERFLOW aggregation
+template <typename Base>
+std::unique_ptr<Base> make_sum_overflow_aggregation()
+{
+  return std::make_unique<detail::sum_with_overflow_aggregation>();
+}
+template CUDF_EXPORT std::unique_ptr<aggregation> make_sum_overflow_aggregation<aggregation>();
+template CUDF_EXPORT std::unique_ptr<groupby_aggregation>
+make_sum_overflow_aggregation<groupby_aggregation>();
+template CUDF_EXPORT std::unique_ptr<groupby_scan_aggregation>
+make_sum_overflow_aggregation<groupby_scan_aggregation>();
+template CUDF_EXPORT std::unique_ptr<reduce_aggregation>
+make_sum_overflow_aggregation<reduce_aggregation>();
+template CUDF_EXPORT std::unique_ptr<segmented_reduce_aggregation>
+make_sum_overflow_aggregation<segmented_reduce_aggregation>();
+
 template <typename Base>
 std::unique_ptr<Base> make_sum_with_overflow_aggregation()
 {

cpp/src/groupby/sort/aggregate.cpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -140,6 +140,18 @@ void aggregate_result_functor::operator()<aggregation::SUM>(aggregation const& a
       get_grouped_values(), helper.num_groups(stream), helper.group_labels(stream), stream, mr));
 }
 
+template <>
+void aggregate_result_functor::operator()<aggregation::SUM_OVERFLOW>(aggregation const& agg)
+{
+  if (cache.has_result(values, agg)) return;
+
+  cache.add_result(
+    values,
+    agg,
+    detail::group_sum_overflow(
+      get_grouped_values(), helper.num_groups(stream), helper.group_labels(stream), stream, mr));
+}
+
 template <>
 void aggregate_result_functor::operator()<aggregation::PRODUCT>(aggregation const& agg)
 {
@@ -878,11 +890,6 @@ std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby::sort
     auto store_functor =
       detail::aggregate_result_functor(request.values, helper(), cache, stream, mr);
     for (auto const& agg : request.aggregations) {
-      // SUM_WITH_OVERFLOW is only supported with hash-based groupby, not sort-based
-      CUDF_EXPECTS(agg->kind != aggregation::SUM_WITH_OVERFLOW,
-                   "SUM_WITH_OVERFLOW aggregation is only supported with hash-based groupby, not "
-                   "sort-based groupby");
-
       // TODO (dm): single pass compute all supported reductions
       cudf::detail::aggregation_dispatcher(agg->kind, store_functor, *agg);
     }

cpp/src/groupby/sort/group_reductions.hpp

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -44,6 +44,26 @@ std::unique_ptr<column> group_sum(column_view const& values,
                                   rmm::cuda_stream_view stream,
                                   rmm::device_async_resource_ref mr);
 
+/**
+ * @brief Internal API to calculate groupwise sum with overflow detection.
+ *
+ * Returns a STRUCT column with two children: the (wrapping) sum and a BOOL8 overflow flag that is
+ * true when the true sum does not fit in the value type. On overflow the sum value is unspecified;
+ * the flag is the meaningful output. A group is null only when all of its values are null.
+ *
+ * @param values Grouped values to sum
+ * @param num_groups Number of groups
+ * @param group_labels ID of group that the corresponding value belongs to
+ * @param stream CUDA stream used for device memory operations and kernel launches.
+ * @param mr Device memory resource used to allocate the returned column's device memory
+ */
+[[nodiscard]] std::unique_ptr<column> group_sum_overflow(
+  column_view const& values,
+  size_type num_groups,
+  cudf::device_span<size_type const> group_labels,
+  rmm::cuda_stream_view stream,
+  rmm::device_async_resource_ref mr);
+
 /**
  * @brief Internal API to calculate groupwise product
  *