Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion cpp/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -881,7 +881,7 @@ add_library(
src/reductions/std.cu
src/reductions/sum.cu
src/reductions/sum_of_squares.cu
src/reductions/sum_with_overflow.cu
src/reductions/sum_overflow.cu
src/reductions/unique_count.cu
src/reductions/unique_count_column.cu
src/reductions/var.cu
Expand Down
3 changes: 2 additions & 1 deletion cpp/include/cudf/aggregation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -80,7 +80,7 @@ class aggregation {
SUM = 0, ///< sum reduction
SUM_OVERFLOW, ///< sum reduction with overflow detection
/// @deprecated Use SUM_OVERFLOW instead.
SUM_WITH_OVERFLOW = SUM_OVERFLOW,
SUM_WITH_OVERFLOW [[deprecated("Use SUM_OVERFLOW instead.")]] = SUM_OVERFLOW,
PRODUCT, ///< product reduction
MIN, ///< min reduction
MAX, ///< max reduction
Expand Down Expand Up @@ -223,6 +223,7 @@ std::unique_ptr<Base> make_sum_overflow_aggregation();
/// @return A SUM_WITH_OVERFLOW aggregation object
/// @deprecated Use make_sum_overflow_aggregation() instead.
template <typename Base = aggregation>
[[deprecated("Use make_sum_overflow_aggregation() instead.")]]
std::unique_ptr<Base> make_sum_with_overflow_aggregation();

/// Factory to create a PRODUCT aggregation
Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cudf/detail/aggregation/aggregation.cuh
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -65,7 +65,7 @@ struct corresponding_operator<aggregation::SUM> {
using type = DeviceSum;
};
template <>
struct corresponding_operator<aggregation::SUM_WITH_OVERFLOW> {
struct corresponding_operator<aggregation::SUM_OVERFLOW> {
using type = DeviceSum;
};
template <>
Expand Down
24 changes: 12 additions & 12 deletions cpp/include/cudf/detail/aggregation/aggregation.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -79,15 +79,15 @@ class sum_aggregation final
};

/**
* @brief Derived class for specifying a sum_with_overflow aggregation
* @brief Derived class for specifying a sum_overflow aggregation
*/
class sum_with_overflow_aggregation final
: public clonable<sum_with_overflow_aggregation>::derived_from<groupby_aggregation,
groupby_scan_aggregation,
reduce_aggregation,
segmented_reduce_aggregation> {
class sum_overflow_aggregation final
: public clonable<sum_overflow_aggregation>::derived_from<groupby_aggregation,
groupby_scan_aggregation,
reduce_aggregation,
segmented_reduce_aggregation> {
public:
sum_with_overflow_aggregation() : aggregation(SUM_WITH_OVERFLOW) {}
sum_overflow_aggregation() : aggregation(SUM_OVERFLOW) {}
};

/**
Expand Down Expand Up @@ -987,11 +987,11 @@ 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
// SUM_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
struct target_type_impl<Source, aggregation::SUM_OVERFLOW> {
using type = struct_view; // SUM_OVERFLOW outputs a struct with sum and overflow fields
};

// Always use `double` for M2
Expand Down Expand Up @@ -1196,8 +1196,8 @@ CUDF_HOST_DEVICE inline decltype(auto) aggregation_dispatcher(aggregation::Kind
switch (k) {
case aggregation::SUM:
return f.template operator()<aggregation::SUM>(std::forward<Ts>(args)...);
case aggregation::SUM_WITH_OVERFLOW:
return f.template operator()<aggregation::SUM_WITH_OVERFLOW>(std::forward<Ts>(args)...);
case aggregation::SUM_OVERFLOW:
return f.template operator()<aggregation::SUM_OVERFLOW>(std::forward<Ts>(args)...);
case aggregation::PRODUCT:
return f.template operator()<aggregation::PRODUCT>(std::forward<Ts>(args)...);
case aggregation::MIN:
Expand Down
4 changes: 2 additions & 2 deletions cpp/include/cudf/detail/aggregation/device_aggregators.cuh
Original file line number Diff line number Diff line change
@@ -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
*/
#pragma once
Expand Down Expand Up @@ -133,7 +133,7 @@ template <typename Source>
(cudf::is_integral_not_bool<Source>() && cudf::is_signed<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> {
struct update_target_element<Source, aggregation::SUM_OVERFLOW> {
using DeviceType = device_storage_type_t<Source>;

__device__ void operator()(mutable_column_device_view target,
Expand Down
16 changes: 8 additions & 8 deletions cpp/include/cudf/reduction.hpp
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -28,11 +28,11 @@ enum class scan_type : bool { INCLUSIVE, EXCLUSIVE };
/**
* @brief Computes the reduction of the values in all rows of a column.
*
* This function does not detect overflows in reductions except for the `SUM_WITH_OVERFLOW`
* This function does not detect overflows in reductions except for the `SUM_OVERFLOW`
* aggregation. When `output_type` does not match the `col.type()`, their values may be promoted to
* `int64_t` or `double` for computing aggregations and then cast to `output_type` before returning.
*
* The `SUM_WITH_OVERFLOW` aggregation is a special case that detects integer
* The `SUM_OVERFLOW` aggregation is a special case that detects integer
* overflow during summation of signed integer or decimal values and returns a struct
* containing both the sum result and an overflow flag. On overflow the sum value
* is unspecified; the boolean flag is the source of truth.
Expand All @@ -54,7 +54,7 @@ enum class scan_type : bool { INCLUSIVE, EXCLUSIVE };
* | Aggregation | Output Type | Init Value | Empty Input | Comments |
* | :---------: | ----------- | :--------: | ----------- | -------- |
* | SUM/PRODUCT | output_type | yes | NA | Input accumulated into output_type variable |
* | SUM_WITH_OVERFLOW | STRUCT{col.type,BOOL8} | yes | {null,false} | {sum, overflow_flag}, input must be signed integer or decimal |
* | SUM_OVERFLOW | STRUCT{col.type,BOOL8} | yes | {null,false} | {sum, overflow_flag}, input must be signed integer or decimal |
* | SUM_OF_SQUARES | output_type | no | NA | Input accumulated into output_type variable |
* | MIN/MAX | col.type | yes | NA | Supports arithmetic, timestamp, duration, string types only |
* | ANY/ALL | BOOL8 | yes | True for ALL only | Checks for non-zero elements |
Expand All @@ -78,7 +78,7 @@ enum class scan_type : bool { INCLUSIVE, EXCLUSIVE };
* @throw std::invalid_argument if `any` or `all` reduction is called and the output type is not BOOL8.
* @throw std::invalid_argument if `mean`, `var`, or `std` reduction is called and
* the `output_type` is not floating point.
* @throw std::invalid_argument if `sum_with_overflow` reduction is called and the
* @throw std::invalid_argument if `sum_overflow` reduction is called and the
* input column type is not a signed integer or decimal, or the `output_type` is not `STRUCT`.
*
* @param col Input column view
Expand All @@ -99,15 +99,15 @@ std::unique_ptr<scalar> reduce(
/**
* @brief Computes the reduction of the values in all rows of a column with an initial value
*
* Only `sum`, `product`, `min`, `max`, `any`, `all`, and `sum_with_overflow` reductions are
* supported. For `sum_with_overflow`, the initial value is added to the sum and overflow
* Only `sum`, `product`, `min`, `max`, `any`, `all`, and `sum_overflow` reductions are
* supported. For `sum_overflow`, the initial value is added to the sum and overflow
* detection is performed throughout the entire computation.
*
* @see cudf::reduce(column_view const&,reduce_aggregation
* const&,data_type,rmm::cuda_stream_view,rmm::device_async_resource_ref) for more details
*
* @throw std::invalid_argument if reduction is not `sum`, `product`, `min`, `max`, `any`, `all`,
* or `sum_with_overflow` and `init` is specified.
* or `sum_overflow` and `init` is specified.
*
* @param col Input column view
* @param agg Aggregation operator applied by the reduction
Expand Down
12 changes: 6 additions & 6 deletions cpp/include/cudf/reduction/detail/reduction_functions.hpp
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -55,11 +55,11 @@ std::unique_ptr<scalar> sum(column_view const& col,
* @param mr Device memory resource used to allocate the returned scalar's device memory
* @return Struct scalar with sum and overflow flag
*/
std::unique_ptr<scalar> sum_with_overflow(column_view const& col,
data_type const output_type,
std::optional<std::reference_wrapper<scalar const>> init,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr);
std::unique_ptr<scalar> sum_overflow(column_view const& col,
data_type const output_type,
std::optional<std::reference_wrapper<scalar const>> init,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr);

/**
* @brief Computes minimum of elements in input column
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/aggregation/aggregation.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -34,11 +34,11 @@ template CUDF_EXPORT std::unique_ptr<scan_aggregation> make_sum_aggregation<scan
template CUDF_EXPORT std::unique_ptr<segmented_reduce_aggregation>
make_sum_aggregation<segmented_reduce_aggregation>();

/// Factory to create a SUM_WITH_OVERFLOW aggregation
/// Factory to create a SUM_OVERFLOW aggregation
template <typename Base>
std::unique_ptr<Base> make_sum_overflow_aggregation()
{
return std::make_unique<detail::sum_with_overflow_aggregation>();
return std::make_unique<detail::sum_overflow_aggregation>();
}
template CUDF_EXPORT std::unique_ptr<aggregation> make_sum_overflow_aggregation<aggregation>();
template CUDF_EXPORT std::unique_ptr<groupby_aggregation>
Expand All @@ -53,7 +53,7 @@ make_sum_overflow_aggregation<segmented_reduce_aggregation>();
template <typename Base>
std::unique_ptr<Base> make_sum_with_overflow_aggregation()
{
return std::make_unique<detail::sum_with_overflow_aggregation>();
return make_sum_overflow_aggregation<Base>();
}
template CUDF_EXPORT std::unique_ptr<aggregation> make_sum_with_overflow_aggregation<aggregation>();
template CUDF_EXPORT std::unique_ptr<groupby_aggregation>
Expand Down
12 changes: 6 additions & 6 deletions cpp/src/aggregation/aggregation.cu
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -42,16 +42,16 @@ struct identity_initializer {
static constexpr bool is_supported()
{
return is_identity_supported<T, k>() or
(k == aggregation::SUM_WITH_OVERFLOW and std::is_same_v<T, cudf::struct_view>);
(k == aggregation::SUM_OVERFLOW and std::is_same_v<T, cudf::struct_view>);
}

public:
template <typename T, aggregation::Kind k>
void operator()(mutable_column_view const& col, rmm::cuda_stream_view stream)
requires(is_supported<T, k>())
{
if constexpr (k == aggregation::SUM_WITH_OVERFLOW) {
// SUM_WITH_OVERFLOW uses a struct with sum and overflow children
if constexpr (k == aggregation::SUM_OVERFLOW) {
// SUM_OVERFLOW uses a struct with sum and overflow children
auto sum_col = col.child(0);
auto overflow_col = col.child(1);

Expand All @@ -64,8 +64,8 @@ struct identity_initializer {
col.size(),
false);
} else if constexpr (std::is_same_v<T, cudf::struct_view>) {
// This should only happen for SUM_WITH_OVERFLOW, but handle it just in case
CUDF_FAIL("Struct columns are only supported for SUM_WITH_OVERFLOW aggregation");
// This should only happen for SUM_OVERFLOW, but handle it just in case
CUDF_FAIL("Struct columns are only supported for SUM_OVERFLOW aggregation");
} else {
using DeviceType = device_storage_type_t<T>;
thrust::fill(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/groupby/common/utils.hpp
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -65,7 +65,7 @@ constexpr bool is_hash_aggregation(aggregation::Kind k)
{
switch (k) {
case aggregation::SUM:
case aggregation::SUM_WITH_OVERFLOW:
case aggregation::SUM_OVERFLOW:
case aggregation::SUM_OF_SQUARES:
case aggregation::PRODUCT:
case aggregation::MIN:
Expand Down
14 changes: 7 additions & 7 deletions cpp/src/groupby/groupby.cu
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -111,8 +111,8 @@ struct empty_column_constructor {
}
if constexpr (k == aggregation::Kind::MERGE_HISTOGRAM) { return empty_like(values); }

if constexpr (k == aggregation::Kind::SUM_WITH_OVERFLOW) {
// SUM_WITH_OVERFLOW returns a struct with sum (same type as input) and overflow (bool)
if constexpr (k == aggregation::Kind::SUM_OVERFLOW) {
// SUM_OVERFLOW returns a struct with sum (same type as input) and overflow (bool)
// children
std::vector<std::unique_ptr<cudf::column>> children;
children.push_back(make_empty_column(values.type()));
Expand Down Expand Up @@ -201,13 +201,13 @@ void verify_valid_requests(host_span<RequestType const> requests)
}),
"Invalid type/aggregation combination.");

// Additional validation for SUM_WITH_OVERFLOW: only signed integers and decimals are supported
// Additional validation for SUM_OVERFLOW: only signed integers and decimals are supported
for (auto const& request : requests) {
for (auto const& agg : request.aggregations) {
if (agg->kind == aggregation::SUM_WITH_OVERFLOW) {
if (agg->kind == aggregation::SUM_OVERFLOW) {
CUDF_EXPECTS(
cudf::detail::is_valid_aggregation(request.values.type(), aggregation::SUM_WITH_OVERFLOW),
"SUM_WITH_OVERFLOW aggregation only supports signed integer types and decimal types. "
cudf::detail::is_valid_aggregation(request.values.type(), aggregation::SUM_OVERFLOW),
"SUM_OVERFLOW aggregation only supports signed integer types and decimal types. "
"Unsigned integers, bool, dictionary columns, and other types are not supported.");
}
}
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/groupby/hash/compute_single_pass_aggs.cu
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2024-2025, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2024-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: Apache-2.0
*/

Expand All @@ -15,10 +15,10 @@ std::pair<bool, size_type> is_shared_memory_compatible(host_span<aggregation::Ki
table_view const& values,
size_type grid_size)
{
// If any aggregation has values type is dictionary, or the aggregation is SUM_WITH_OVERFLOW,
// If any aggregation has values type is dictionary, or the aggregation is SUM_OVERFLOW,
// we should always use global memory code path.
for (std::size_t i = 0; i < agg_kinds.size(); ++i) {
if (is_dictionary(values.column(i).type()) || agg_kinds[i] == aggregation::SUM_WITH_OVERFLOW) {
if (is_dictionary(values.column(i).type()) || agg_kinds[i] == aggregation::SUM_OVERFLOW) {
return {false, 0};
}
}
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/groupby/hash/compute_single_pass_aggs.cuh
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -47,7 +47,7 @@ std::pair<rmm::device_uvector<size_type>, bool> compute_single_pass_aggs(

// Performs naive global memory aggregations when the workload is not compatible with shared
// memory, such as when aggregating dictionary columns, when there is insufficient dynamic
// shared memory for shared memory aggregations, or when SUM_WITH_OVERFLOW aggregations are
// shared memory for shared memory aggregations, or when SUM_OVERFLOW aggregations are
// present.
auto const run_aggs_by_global_mem_kernel = [&] {
auto [agg_results, unique_key_indices] = compute_global_memory_aggs(
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/groupby/hash/output_utils.cu
Original file line number Diff line number Diff line change
@@ -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
*/

Expand Down Expand Up @@ -42,7 +42,7 @@ namespace {
* @brief Functor to create the result columns for hash-based groupby aggregations
*
* This functor handles the creation of appropriately typed and sized columns for each
* aggregation, including special handling for SUM_WITH_OVERFLOW which requires a struct column.
* aggregation, including special handling for SUM_OVERFLOW which requires a struct column.
* For data types smaller than 4 bytes, the buffer size is adjusted to be a multiple of 4 to
* ensure memory safety when atomic operations use 4-byte CAS loops to emulate smaller atomics.
*/
Expand Down Expand Up @@ -80,7 +80,7 @@ struct result_column_creator {
}
return make_fixed_width_column(d_type, size, state, stream, mr);
};
if (agg != aggregation::SUM_WITH_OVERFLOW) {
if (agg != aggregation::SUM_OVERFLOW) {
auto const target_type = cudf::detail::target_type(col_type, agg);
auto const mask_flag = nullable ? mask_state::ALL_NULL : mask_state::UNALLOCATED;
return make_uninitialized_column(target_type, output_size, mask_flag);
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/groupby/streaming_groupby/impl.cu
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: Apache-2.0
*/

Expand Down Expand Up @@ -149,8 +149,8 @@ void streaming_groupby::impl::initialize(table_view const& data, rmm::cuda_strea
"Streaming groupby does not support MIN/MAX on variable-width types "
"(internally decomposed to ARGMIN/ARGMAX).",
std::invalid_argument);
CUDF_EXPECTS(k != aggregation::SUM_WITH_OVERFLOW,
"Streaming groupby does not support SUM_WITH_OVERFLOW "
CUDF_EXPECTS(k != aggregation::SUM_OVERFLOW,
"Streaming groupby does not support SUM_OVERFLOW "
"(struct intermediate cannot be merged across batches).",
std::invalid_argument);
}
Expand Down
Loading
Loading