diff --git a/cpp/include/cudf/reduction.hpp b/cpp/include/cudf/reduction.hpp index 4d5959f99cbf..f93b2b709a0a 100644 --- a/cpp/include/cudf/reduction.hpp +++ b/cpp/include/cudf/reduction.hpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -33,8 +33,8 @@ enum class scan_type : bool { INCLUSIVE, EXCLUSIVE }; * `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 - * overflow during summation of `int64_t` values and returns a struct containing - * both the sum result and an overflow flag. + * overflow during summation of signed integer or decimal values and returns a struct + * containing both the sum result and an overflow flag. * * Only `min` and `max` ops are supported for reduction of non-arithmetic * types (e.g. timestamp or string). @@ -53,7 +53,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{INT64,BOOL8} | yes | {null,false} | {sum, overflow_flag}, input must be INT64 | + * | SUM_WITH_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 | @@ -78,7 +78,7 @@ enum class scan_type : bool { INCLUSIVE, EXCLUSIVE }; * @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 - * input column type is not `INT64` or the `output_dtype` is not `STRUCT`. + * input column type is not a signed integer or decimal, or the `output_type` is not `STRUCT`. * * @param col Input column view * @param agg Aggregation operator applied by the reduction diff --git a/cpp/include/cudf/reduction/detail/reduction_functions.hpp b/cpp/include/cudf/reduction/detail/reduction_functions.hpp index 1c16c05f93c9..aa9cc559ef52 100644 --- a/cpp/include/cudf/reduction/detail/reduction_functions.hpp +++ b/cpp/include/cudf/reduction/detail/reduction_functions.hpp @@ -40,15 +40,16 @@ std::unique_ptr sum(column_view const& col, rmm::device_async_resource_ref mr); /** - * @brief Computes sum with overflow detection of int64_t elements in input column + * @brief Computes sum with overflow detection of signed integer or decimal elements in input column * - * Returns a struct scalar with {sum: int64_t, overflow: bool} fields. - * Only supports int64_t input columns. + * Returns a struct scalar with {sum: same type as input, overflow: bool} fields. + * Supported input types: signed integers (int8/16/32/64) and decimals (decimal32/64/128). * - * @throw std::invalid_argument if input column type is not int64_t + * @throw std::invalid_argument if input column type is not a supported signed integer or decimal + * @throw std::invalid_argument if `output_type` is not STRUCT * - * @param col input column to compute sum with overflow detection (must be int64_t) - * @param output_type data type of return type (must be struct) + * @param col input column to compute sum with overflow detection + * @param output_type data type of return type (must be STRUCT) * @param init initial value of the sum * @param stream CUDA stream used for device memory operations and kernel launches * @param mr Device memory resource used to allocate the returned scalar's device memory diff --git a/cpp/src/reductions/reductions.cpp b/cpp/src/reductions/reductions.cpp index 05a54cb5a934..d58f145d8a7e 100644 --- a/cpp/src/reductions/reductions.cpp +++ b/cpp/src/reductions/reductions.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -96,7 +96,8 @@ struct reduction_function : public base_reductio }; template - requires(std::is_same_v) // only int64_t is supported for SUM_WITH_OVERFLOW + requires((cudf::is_integral_not_bool() && cudf::is_signed()) || + cudf::is_fixed_point()) struct reduction_function : public base_reduction_function { [[nodiscard]] std::unique_ptr reduce(reduction_parameters const& params) const diff --git a/cpp/src/reductions/sum_with_overflow.cu b/cpp/src/reductions/sum_with_overflow.cu index 693ad40ac1e0..28a74cfdafa9 100644 --- a/cpp/src/reductions/sum_with_overflow.cu +++ b/cpp/src/reductions/sum_with_overflow.cu @@ -8,171 +8,193 @@ #include #include #include +#include #include #include #include #include +#include +#include #include #include #include -#include -#include +#include #include namespace cudf::reduction::detail { -// Simple pair to hold sum and overflow flag +namespace { + +// `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 struct sum_overflow_result { - int64_t sum; - bool overflow; + DeviceType sum; + cudf::size_type wraps; - CUDF_HOST_DEVICE sum_overflow_result() : sum(0), overflow(false) {} - CUDF_HOST_DEVICE sum_overflow_result(int64_t s, bool o) : sum(s), overflow(o) {} + 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} {} }; -// Binary operator for combining sum_overflow_result values +template struct overflow_sum_op { - __device__ sum_overflow_result operator()(sum_overflow_result const& lhs, - sum_overflow_result const& rhs) const + __device__ sum_overflow_result operator()( + sum_overflow_result const& lhs, sum_overflow_result const& rhs) const { - // If either operand already has overflow, result has overflow - if (lhs.overflow || rhs.overflow) { - // Still compute the sum for consistency, but mark as overflow - // This addition may wrap but we've already detected overflow - return sum_overflow_result{lhs.sum + rhs.sum, true}; - } - - // Check for overflow BEFORE performing the addition to avoid UB - bool overflow_detected = false; - - // Check for positive overflow: would the addition exceed INT64_MAX? - if (rhs.sum > 0 && lhs.sum > cuda::std::numeric_limits::max() - rhs.sum) { - overflow_detected = true; - } - // Check for negative overflow: would the addition go below INT64_MIN? - else if (rhs.sum < 0 && lhs.sum < cuda::std::numeric_limits::min() - rhs.sum) { - overflow_detected = true; - } - - // Perform the addition (safe if no overflow detected) - int64_t const result_sum = lhs.sum + rhs.sum; - - return sum_overflow_result{result_sum, overflow_detected}; + auto const r = cuda::add_overflow(lhs.sum, rhs.sum); + auto const carry = r.overflow ? (rhs.sum > DeviceType{0} ? 1 : -1) : 0; + return sum_overflow_result{r.value, lhs.wraps + rhs.wraps + carry}; } }; -// Transform function to convert int64_t values to sum_overflow_result +template struct to_sum_overflow { - __device__ sum_overflow_result operator()(int64_t value) const + __device__ sum_overflow_result operator()(DeviceType value) const { - return sum_overflow_result{value, false}; + return sum_overflow_result{value, 0}; } }; -// Transform functor for null-aware conversion using index +template struct null_aware_to_sum_overflow { - cudf::column_device_view const* dcol_ptr; + cudf::column_device_view dcol; - CUDF_HOST_DEVICE null_aware_to_sum_overflow(cudf::column_device_view const* dcol) : dcol_ptr(dcol) - { - } + CUDF_HOST_DEVICE null_aware_to_sum_overflow(cudf::column_device_view const& d) : dcol{d} {} - __device__ sum_overflow_result operator()(cudf::size_type idx) const + __device__ sum_overflow_result operator()(cudf::size_type idx) const { - return dcol_ptr->is_valid(idx) ? sum_overflow_result{dcol_ptr->element(idx), false} - : sum_overflow_result{0, false}; + return dcol.is_valid(idx) ? sum_overflow_result{dcol.element(idx), 0} + : sum_overflow_result{DeviceType{0}, 0}; } }; -std::unique_ptr sum_with_overflow( +template +std::unique_ptr make_sum_overflow_struct_scalar( + device_storage_type_t sum_value, + bool overflow_value, + bool sum_is_valid, + cudf::data_type const& source_type, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + auto const temp_mr = cudf::get_current_device_resource_ref(); + + std::unique_ptr sum_scalar; + if constexpr (cudf::is_fixed_point()) { + sum_scalar = cudf::make_fixed_point_scalar( + sum_value, numeric::scale_type{source_type.scale()}, stream, temp_mr); + } else { + sum_scalar = + cudf::make_fixed_width_scalar(static_cast(sum_value), stream, temp_mr); + } + sum_scalar->set_valid_async(sum_is_valid, stream); + auto overflow_scalar = cudf::make_fixed_width_scalar(overflow_value, stream, temp_mr); + + std::vector> children; + children.push_back(cudf::make_column_from_scalar(*sum_scalar, 1, stream, temp_mr)); + children.push_back(cudf::make_column_from_scalar(*overflow_scalar, 1, stream, temp_mr)); + + std::vector child_views{children[0]->view(), children[1]->view()}; + return cudf::make_struct_scalar( + cudf::host_span{child_views}, stream, mr); +} + +template +std::unique_ptr sum_with_overflow_impl( column_view const& col, - cudf::data_type const output_dtype, std::optional> init, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { - CUDF_FUNC_RANGE(); + using DeviceType = device_storage_type_t; - // SUM_WITH_OVERFLOW only supports int64_t input - CUDF_EXPECTS(col.type().id() == cudf::type_id::INT64, - "SUM_WITH_OVERFLOW only supports int64_t input types", - std::invalid_argument); + if (init.has_value() && !init.value().get().is_valid(stream)) { + return make_sum_overflow_struct_scalar( + DeviceType{0}, false, false, col.type(), stream, mr); + } - // Handle empty column if (col.size() == 0 || col.size() == col.null_count()) { - // Create struct with {null sum, false overflow} - auto sum_scalar = - cudf::make_default_constructed_scalar(cudf::data_type{cudf::type_id::INT64}, stream, mr); - sum_scalar->set_valid_async(false, stream); - auto overflow_scalar = cudf::make_fixed_width_scalar(false, stream, mr); - - std::vector> children; - children.push_back(cudf::make_column_from_scalar(*sum_scalar, 1, stream, mr)); - children.push_back(cudf::make_column_from_scalar(*overflow_scalar, 1, stream, mr)); - - // Use host_span of column_views instead of table_view to avoid double wrapping - std::vector child_views; - child_views.push_back(children[0]->view()); - child_views.push_back(children[1]->view()); - - return cudf::make_struct_scalar( - cudf::host_span{child_views}, stream, mr); + return make_sum_overflow_struct_scalar( + DeviceType{0}, false, false, col.type(), stream, mr); } - // Create device view auto dcol = cudf::column_device_view::create(col, stream); - // Set up initial value - sum_overflow_result initial_value{0, false}; - if (init.has_value() && init.value().get().is_valid(stream)) { - auto const& init_scalar = static_cast const&>(init.value().get()); - initial_value.sum = init_scalar.value(stream); + sum_overflow_result initial_value{DeviceType{0}, 0}; + if (init.has_value()) { + auto const& init_scalar = static_cast const&>(init.value().get()); + initial_value.sum = static_cast(init_scalar.value(stream)); } - // Perform the reduction using thrust::transform_reduce auto counting_iter = cuda::counting_iterator{0}; - auto dcol_ptr = dcol.get(); - sum_overflow_result result; + sum_overflow_result result; if (col.has_nulls()) { - // Use null-aware transform functor result = thrust::transform_reduce( rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), counting_iter, counting_iter + col.size(), - null_aware_to_sum_overflow{dcol_ptr}, + null_aware_to_sum_overflow{*dcol}, initial_value, - overflow_sum_op{}); + overflow_sum_op{}); } else { - // Use direct iterator for non-null case - auto input_iter = dcol->begin(); + auto input_iter = dcol->begin(); result = thrust::transform_reduce( rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), input_iter, input_iter + col.size(), - to_sum_overflow{}, + to_sum_overflow{}, initial_value, - overflow_sum_op{}); + overflow_sum_op{}); } - // Create result struct scalar with {sum: int64_t, overflow: bool} - auto sum_scalar = cudf::make_fixed_width_scalar(result.sum, stream, mr); - auto overflow_scalar = cudf::make_fixed_width_scalar(result.overflow, stream, mr); + // On overflow, zero the sum value; the boolean flag is the source of truth. + auto const overflowed = result.wraps != 0; + return make_sum_overflow_struct_scalar( + overflowed ? DeviceType{0} : result.sum, overflowed, true, col.type(), stream, mr); +} - // Create struct scalar using cudf::make_struct_scalar with host_span of column_views - std::vector> children; - children.push_back(cudf::make_column_from_scalar(*sum_scalar, 1, stream, mr)); - children.push_back(cudf::make_column_from_scalar(*overflow_scalar, 1, stream, mr)); +struct sum_with_overflow_dispatcher { + template + requires((cudf::is_integral_not_bool() && cudf::is_signed()) || + cudf::is_fixed_point()) + std::unique_ptr operator()(column_view const& col, + std::optional> init, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) const + { + return sum_with_overflow_impl(col, init, stream, mr); + } - // Use host_span of column_views instead of table_view to avoid double wrapping - std::vector child_views; - child_views.push_back(children[0]->view()); - child_views.push_back(children[1]->view()); + template + requires(!((cudf::is_integral_not_bool() && cudf::is_signed()) || + cudf::is_fixed_point())) + std::unique_ptr operator()(column_view const&, + std::optional>, + rmm::cuda_stream_view, + rmm::device_async_resource_ref) const + { + CUDF_FAIL("SUM_WITH_OVERFLOW reduction supports only signed integer and decimal types.", + std::invalid_argument); + } +}; - return cudf::make_struct_scalar( - cudf::host_span{child_views}, stream, mr); +} // namespace + +std::unique_ptr sum_with_overflow( + column_view const& col, + cudf::data_type const output_dtype, + std::optional> init, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + CUDF_EXPECTS(output_dtype.id() == type_id::STRUCT, + "SUM_WITH_OVERFLOW output dtype must be STRUCT.", + std::invalid_argument); + return cudf::type_dispatcher(col.type(), sum_with_overflow_dispatcher{}, col, init, stream, mr); } } // namespace cudf::reduction::detail diff --git a/cpp/tests/reductions/reduction_tests.cpp b/cpp/tests/reductions/reduction_tests.cpp index 8cab946174be..5b878d9173d7 100644 --- a/cpp/tests/reductions/reduction_tests.cpp +++ b/cpp/tests/reductions/reduction_tests.cpp @@ -3458,9 +3458,78 @@ TEST_F(StructReductionTest, StructReductionMinMaxAndArgMinMaxWithNulls) } } -// Test for SUM_WITH_OVERFLOW aggregation using regular reduce() function +// Helpers to map a parameter type to its storage rep (T for integers, T::rep for decimals). +template ()> +struct sum_overflow_rep { + using type = T; +}; +template +struct sum_overflow_rep { + using type = typename T::rep; +}; +template +using sum_overflow_rep_t = typename sum_overflow_rep::type; + +// Test for SUM_WITH_OVERFLOW aggregation using regular reduce() function. +// Parametrized on signed integer types AND decimal types. +template struct ReduceWithOverflowTest : public cudf::test::BaseFixture { - // Helper function to extract sum and overflow from struct scalar returned by reduce() + using Rep = sum_overflow_rep_t; + static constexpr numeric::scale_type scale{0}; + + cudf::test::fixed_width_column_wrapper make_col(std::initializer_list values) + requires(!cudf::is_fixed_point()) + { + return cudf::test::fixed_width_column_wrapper(values); + } + cudf::test::fixed_point_column_wrapper make_col(std::initializer_list values) + requires(cudf::is_fixed_point()) + { + return cudf::test::fixed_point_column_wrapper(values, scale); + } + + cudf::test::fixed_width_column_wrapper make_null_col(std::initializer_list values, + std::initializer_list validity) + requires(!cudf::is_fixed_point()) + { + return cudf::test::fixed_width_column_wrapper(values, std::cbegin(validity)); + } + cudf::test::fixed_point_column_wrapper make_null_col(std::initializer_list values, + std::initializer_list validity) + requires(cudf::is_fixed_point()) + { + return cudf::test::fixed_point_column_wrapper(values, std::cbegin(validity), scale); + } + + cudf::test::fixed_width_column_wrapper make_col_from_vec(std::vector const& values) + requires(!cudf::is_fixed_point()) + { + return cudf::test::fixed_width_column_wrapper(values.begin(), values.end()); + } + cudf::test::fixed_point_column_wrapper make_col_from_vec(std::vector const& values) + requires(cudf::is_fixed_point()) + { + return cudf::test::fixed_point_column_wrapper(values.begin(), values.end(), scale); + } + + std::unique_ptr make_init_scalar(Rep value) + { + if constexpr (cudf::is_fixed_point()) { + return cudf::make_fixed_point_scalar(value, scale); + } else { + return cudf::make_fixed_width_scalar(static_cast(value)); + } + } + + Rep get_sum_value(std::unique_ptr const& sum_result) + { + if constexpr (cudf::is_fixed_point()) { + return static_cast const*>(sum_result.get())->value(); + } else { + return static_cast const*>(sum_result.get())->value(); + } + } + std::pair, std::unique_ptr> extract_sum_overflow( std::unique_ptr const& result) { @@ -3474,233 +3543,252 @@ struct ReduceWithOverflowTest : public cudf::test::BaseFixture { EXPECT_EQ(table_view.column(0).size(), 1); EXPECT_EQ(table_view.column(1).size(), 1); - auto sum_result = cudf::get_element(table_view.column(0), 0); - auto overflow_flag = cudf::get_element(table_view.column(1), 0); - return std::make_pair(std::move(sum_result), std::move(overflow_flag)); + return std::make_pair(cudf::get_element(table_view.column(0), 0), + cudf::get_element(table_view.column(1), 0)); } }; +using ReduceWithOverflowTypes = ::testing::Types; +TYPED_TEST_SUITE(ReduceWithOverflowTest, ReduceWithOverflowTypes); -TEST_F(ReduceWithOverflowTest, SumWithoutOverflow) +TYPED_TEST(ReduceWithOverflowTest, SumWithoutOverflow) { - std::vector values{1, 2, 3, 4, 5}; - cudf::test::fixed_width_column_wrapper col(values.begin(), values.end()); + using Rep = typename TestFixture::Rep; + auto col = this->make_col({Rep{1}, Rep{2}, Rep{3}, Rep{4}, Rep{5}}); auto result = cudf::reduce(col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}); - auto [sum_result, overflow_flag] = extract_sum_overflow(result); + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); EXPECT_TRUE(sum_result->is_valid()); - EXPECT_TRUE(overflow_flag->is_valid()); - - auto sum_value = static_cast const*>(sum_result.get())->value(); - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); - - EXPECT_EQ(sum_value, 15); // 1+2+3+4+5 = 15 - EXPECT_FALSE(overflow_value); // No overflow expected + EXPECT_EQ(sum_result->type().id(), cudf::type_to_id()); + EXPECT_EQ(this->get_sum_value(sum_result), Rep{15}); + EXPECT_FALSE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, PositiveOverflow) +TYPED_TEST(ReduceWithOverflowTest, PositiveOverflow) { - std::vector positive_overflow_values{std::numeric_limits::max(), - 1}; // max + 1 should overflow - cudf::test::fixed_width_column_wrapper col(positive_overflow_values.begin(), - positive_overflow_values.end()); + using Rep = typename TestFixture::Rep; + auto col = this->make_col({std::numeric_limits::max(), Rep{1}}); auto result = cudf::reduce(col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}); - auto [sum_result, overflow_flag] = extract_sum_overflow(result); - - EXPECT_TRUE(sum_result->is_valid()); - EXPECT_TRUE(overflow_flag->is_valid()); - - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); - - EXPECT_TRUE(overflow_value); // Should detect positive overflow + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_TRUE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, NegativeOverflow) +TYPED_TEST(ReduceWithOverflowTest, NegativeOverflow) { - std::vector negative_overflow_values{std::numeric_limits::min(), - -1}; // min - 1 should overflow - cudf::test::fixed_width_column_wrapper col(negative_overflow_values.begin(), - negative_overflow_values.end()); + using Rep = typename TestFixture::Rep; + auto col = this->make_col({std::numeric_limits::min(), Rep{-1}}); auto result = cudf::reduce(col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}); - auto [sum_result, overflow_flag] = extract_sum_overflow(result); - - EXPECT_TRUE(sum_result->is_valid()); - EXPECT_TRUE(overflow_flag->is_valid()); - - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); - - EXPECT_TRUE(overflow_value); // Should detect negative overflow + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_TRUE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, AccumulatingOverflow) +TYPED_TEST(ReduceWithOverflowTest, AccumulatingOverflow) { - // Use large values that when accumulated could cause overflow - std::vector accumulating_overflow{ - std::numeric_limits::max() / 3, - std::numeric_limits::max() / 3, - std::numeric_limits::max() / 3, - std::numeric_limits::max() / 3}; // This should overflow - cudf::test::fixed_width_column_wrapper col(accumulating_overflow.begin(), - accumulating_overflow.end()); + using Rep = typename TestFixture::Rep; + auto const big = static_cast(std::numeric_limits::max() / Rep{3}); + auto col = this->make_col({big, big, big, big}); // 4 * (max/3) > max auto result = cudf::reduce(col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}); - auto [sum_result, overflow_flag] = extract_sum_overflow(result); - - EXPECT_TRUE(sum_result->is_valid()); - EXPECT_TRUE(overflow_flag->is_valid()); - - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); - - // Should detect overflow since we're adding 4 * (max/3) which > max - EXPECT_TRUE(overflow_value); // Should detect accumulating overflow + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_TRUE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, EmptyColumn) +TYPED_TEST(ReduceWithOverflowTest, EmptyColumn) { - cudf::test::fixed_width_column_wrapper empty_col{}; + auto empty_col = this->make_col({}); auto result = cudf::reduce(empty_col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}); - auto [sum_result, overflow_flag] = extract_sum_overflow(result); - - EXPECT_FALSE(sum_result->is_valid()); // Should be null for empty input - EXPECT_TRUE(overflow_flag->is_valid()); - - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); - EXPECT_FALSE(overflow_value); // No overflow for empty input + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_FALSE(sum_result->is_valid()); + EXPECT_EQ(sum_result->type().id(), cudf::type_to_id()); + EXPECT_FALSE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, AllNullColumn) +TYPED_TEST(ReduceWithOverflowTest, AllNullColumn) { - std::vector values{1, 2, 3}; - std::vector validity{false, false, false}; - cudf::test::fixed_width_column_wrapper null_col( - values.begin(), values.end(), validity.begin()); + using Rep = typename TestFixture::Rep; + auto null_col = this->make_null_col({Rep{1}, Rep{2}, Rep{3}}, {false, false, false}); auto result = cudf::reduce(null_col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}); - auto [sum_result, overflow_flag] = extract_sum_overflow(result); - - EXPECT_FALSE(sum_result->is_valid()); // Should be null for all-null input - EXPECT_TRUE(overflow_flag->is_valid()); - - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); - EXPECT_FALSE(overflow_value); // No overflow for all-null input + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_FALSE(sum_result->is_valid()); + EXPECT_FALSE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, WithInitialValue) +TYPED_TEST(ReduceWithOverflowTest, WithInitialValue) { - std::vector values{1, 2, 3}; - cudf::test::fixed_width_column_wrapper col(values.begin(), values.end()); - auto init_scalar = cudf::make_fixed_width_scalar(10); + using Rep = typename TestFixture::Rep; + auto col = this->make_col({Rep{1}, Rep{2}, Rep{3}}); + auto init_scalar = this->make_init_scalar(Rep{10}); auto result = cudf::reduce(col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}, *init_scalar); - auto [sum_result, overflow_flag] = extract_sum_overflow(result); - + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); EXPECT_TRUE(sum_result->is_valid()); - EXPECT_TRUE(overflow_flag->is_valid()); + EXPECT_EQ(this->get_sum_value(sum_result), Rep{16}); // 10 + 1 + 2 + 3 + EXPECT_FALSE(static_cast const*>(overflow_flag.get())->value()); +} - auto sum_value = static_cast const*>(sum_result.get())->value(); - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); +TYPED_TEST(ReduceWithOverflowTest, InitialValuePositiveOverflow) +{ + using Rep = typename TestFixture::Rep; + auto col = this->make_col({Rep{1}, Rep{2}, Rep{3}}); + // (max - 3) + 1 + 2 + 3 = max + 3 overflows + auto init_scalar = + this->make_init_scalar(static_cast(std::numeric_limits::max() - Rep{3})); + + auto result = cudf::reduce(col, + *cudf::make_sum_with_overflow_aggregation(), + cudf::data_type{cudf::type_id::STRUCT}, + *init_scalar); - EXPECT_EQ(sum_value, 16); // 10 + 1 + 2 + 3 = 16 - EXPECT_FALSE(overflow_value); // No overflow expected + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_TRUE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, InitialValuePositiveOverflow) +TYPED_TEST(ReduceWithOverflowTest, InitialValueNegativeOverflow) { - std::vector values{1, 2, 3}; - cudf::test::fixed_width_column_wrapper col(values.begin(), values.end()); - auto init_scalar = cudf::make_fixed_width_scalar(std::numeric_limits::max() - - 3); // max - 3 + 6 = max + 3 (overflow) + using Rep = typename TestFixture::Rep; + auto col = this->make_col({Rep{-1}, Rep{-2}, Rep{-3}}); + // (min + 3) + (-1) + (-2) + (-3) = min - 3 overflows + auto init_scalar = + this->make_init_scalar(static_cast(std::numeric_limits::min() + Rep{3})); auto result = cudf::reduce(col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}, *init_scalar); - auto [sum_result, overflow_flag] = extract_sum_overflow(result); + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_TRUE(static_cast const*>(overflow_flag.get())->value()); +} + +TYPED_TEST(ReduceWithOverflowTest, SlicedColumn) +{ + using Rep = typename TestFixture::Rep; + auto full = + this->make_col({Rep{50}, Rep{60}, Rep{1}, Rep{2}, Rep{3}, Rep{4}, Rep{5}, Rep{70}, Rep{80}}); + auto sliced = cudf::slice(full, {2, 7}); + ASSERT_EQ(sliced.size(), 1); + + auto result = cudf::reduce(sliced.front(), + *cudf::make_sum_with_overflow_aggregation(), + cudf::data_type{cudf::type_id::STRUCT}); + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); EXPECT_TRUE(sum_result->is_valid()); - EXPECT_TRUE(overflow_flag->is_valid()); + EXPECT_EQ(this->get_sum_value(sum_result), Rep{15}); + EXPECT_FALSE(static_cast const*>(overflow_flag.get())->value()); +} + +TYPED_TEST(ReduceWithOverflowTest, SlicedColumnWithNulls) +{ + using Rep = typename TestFixture::Rep; + auto full = this->make_null_col( + {Rep{50}, Rep{60}, Rep{1}, Rep{2}, Rep{3}, Rep{4}, Rep{5}, Rep{70}, Rep{80}}, + {true, true, true, false, true, false, true, true, true}); + auto sliced = cudf::slice(full, {2, 7}); + ASSERT_EQ(sliced.size(), 1); - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); + auto result = cudf::reduce(sliced.front(), + *cudf::make_sum_with_overflow_aggregation(), + cudf::data_type{cudf::type_id::STRUCT}); - // (max - 3) + 1 + 2 + 3 = max + 3, which should overflow - EXPECT_TRUE(overflow_value); // Should detect overflow with initial value + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_TRUE(sum_result->is_valid()); + EXPECT_EQ(this->get_sum_value(sum_result), Rep{9}); // 1 + 3 + 5 + EXPECT_FALSE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, InitialValueNegativeOverflow) +TYPED_TEST(ReduceWithOverflowTest, MultiBlockInputNoOverflow) { - std::vector values{-1, -2, -3}; - cudf::test::fixed_width_column_wrapper col(values.begin(), values.end()); - auto init_scalar = cudf::make_fixed_width_scalar(std::numeric_limits::min() + - 3); // min + 3 - 6 = min - 3 (overflow) + using Rep = typename TestFixture::Rep; + // Alternating +1/-1 exercises both pairwise-overflow branches on every combine; the + // closed-form sum (-1 for odd N, 0 for even N) fits in every supported DeviceType. + constexpr cudf::size_type N = 100'001; + std::vector data; + data.reserve(N); + for (cudf::size_type i = 0; i < N; ++i) { + data.push_back(static_cast(i % 2 == 0 ? 1 : -1)); + } + auto col = this->make_col_from_vec(data); auto result = cudf::reduce(col, *cudf::make_sum_with_overflow_aggregation(), - cudf::data_type{cudf::type_id::STRUCT}, - *init_scalar); - - auto [sum_result, overflow_flag] = extract_sum_overflow(result); + cudf::data_type{cudf::type_id::STRUCT}); + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); EXPECT_TRUE(sum_result->is_valid()); - EXPECT_TRUE(overflow_flag->is_valid()); + EXPECT_EQ(this->get_sum_value(sum_result), Rep{1}); + EXPECT_FALSE(static_cast const*>(overflow_flag.get())->value()); +} + +TYPED_TEST(ReduceWithOverflowTest, InvalidInit) +{ + using Rep = typename TestFixture::Rep; + auto col = this->make_col({Rep{1}, Rep{2}, Rep{3}}); - auto overflow_value = - static_cast const*>(overflow_flag.get())->value(); + auto init_scalar = this->make_init_scalar(Rep{0}); + init_scalar->set_valid_async(false, cudf::get_default_stream()); + + auto result = cudf::reduce(col, + *cudf::make_sum_with_overflow_aggregation(), + cudf::data_type{cudf::type_id::STRUCT}, + *init_scalar); - // (min + 3) + (-1) + (-2) + (-3) = min - 3, which should overflow - EXPECT_TRUE(overflow_value); // Should detect negative overflow with initial value + auto [sum_result, overflow_flag] = this->extract_sum_overflow(result); + EXPECT_FALSE(sum_result->is_valid()); + EXPECT_FALSE(static_cast const*>(overflow_flag.get())->value()); } -TEST_F(ReduceWithOverflowTest, ErrorHandlingNonInt64) -{ - std::vector int32_values{1, 2, 3}; - cudf::test::fixed_width_column_wrapper int32_col(int32_values.begin(), - int32_values.end()); +// Non-typed fixture used by the error-handling test. +struct ReduceWithOverflowErrorTest : public cudf::test::BaseFixture {}; - EXPECT_THROW(cudf::reduce(int32_col, +TEST_F(ReduceWithOverflowErrorTest, UnsupportedTypes) +{ + // Unsigned, floating-point, and string columns must be rejected. + cudf::test::fixed_width_column_wrapper uint_col{1u, 2u, 3u}; + EXPECT_THROW(cudf::reduce(uint_col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}), std::invalid_argument); -} -TEST_F(ReduceWithOverflowTest, ErrorHandlingNonArithmetic) -{ - std::vector string_values{"a", "b", "c"}; - cudf::test::strings_column_wrapper string_col(string_values.begin(), string_values.end()); + cudf::test::fixed_width_column_wrapper float_col{1.0f, 2.0f, 3.0f}; + EXPECT_THROW(cudf::reduce(float_col, + *cudf::make_sum_with_overflow_aggregation(), + cudf::data_type{cudf::type_id::STRUCT}), + std::invalid_argument); + cudf::test::strings_column_wrapper string_col{"a", "b", "c"}; EXPECT_THROW(cudf::reduce(string_col, *cudf::make_sum_with_overflow_aggregation(), cudf::data_type{cudf::type_id::STRUCT}), @@ -3749,6 +3837,8 @@ TEST_F(ReductionIsValidTest, IsValidAggregation) EXPECT_TRUE(cudf::reduction::is_valid_aggregation(int64_type, cudf::aggregation::MERGE_TDIGEST)); EXPECT_TRUE(cudf::reduction::is_valid_aggregation(decimal_type, cudf::aggregation::SUM)); + EXPECT_TRUE( + cudf::reduction::is_valid_aggregation(decimal_type, cudf::aggregation::SUM_WITH_OVERFLOW)); EXPECT_TRUE( cudf::reduction::is_valid_aggregation(decimal_type, cudf::aggregation::SUM_OF_SQUARES)); EXPECT_TRUE(cudf::reduction::is_valid_aggregation(decimal_type, cudf::aggregation::MEDIAN)); diff --git a/java/src/main/java/ai/rapids/cudf/Aggregation.java b/java/src/main/java/ai/rapids/cudf/Aggregation.java index 48dad2978649..ef14955af766 100644 --- a/java/src/main/java/ai/rapids/cudf/Aggregation.java +++ b/java/src/main/java/ai/rapids/cudf/Aggregation.java @@ -542,10 +542,12 @@ private SumWithOverflowAggregation() { /** * Sum aggregation that also reports overflow. The result is a struct with - * children {sum, overflow: BOOL8}. For column reductions the input must be - * INT64. For hash-based groupby the input may be any signed integer type - * (INT8/16/32/64) or fixed-point decimal. Sort-based groupby, scan, - * segmented reduce, and rolling are not supported by cudf. + * children {sum: same type as input, overflow: BOOL8}. The input may be any + * signed integer type (INT8/16/32/64) or fixed-point decimal + * (DECIMAL32/64/128), for both column reductions and hash-based groupby. + * On overflow the sum value is zeroed; the boolean flag is the source of + * truth. Sort-based groupby, scan, segmented reduce, and rolling are not + * supported by cudf. */ static SumWithOverflowAggregation sumWithOverflow() { return new SumWithOverflowAggregation(); diff --git a/java/src/main/java/ai/rapids/cudf/ReductionAggregation.java b/java/src/main/java/ai/rapids/cudf/ReductionAggregation.java index 05833601fe8e..36b4e728d39f 100644 --- a/java/src/main/java/ai/rapids/cudf/ReductionAggregation.java +++ b/java/src/main/java/ai/rapids/cudf/ReductionAggregation.java @@ -53,8 +53,11 @@ public static ReductionAggregation sum() { } /** - * Sum reduction that also reports int64 overflow. Result is a struct scalar - * with children {sum: INT64, overflow: BOOL8}. Input column must be INT64. + * Sum reduction that also reports overflow. The result is a struct scalar + * with children {sum: same type as input, overflow: BOOL8}. The input may + * be any signed integer type (INT8/16/32/64) or fixed-point decimal + * (DECIMAL32/64/128). On overflow the sum value is zeroed; the boolean flag + * is the source of truth. */ public static ReductionAggregation sumWithOverflow() { return new ReductionAggregation(Aggregation.sumWithOverflow()); diff --git a/java/src/test/java/ai/rapids/cudf/ReductionTest.java b/java/src/test/java/ai/rapids/cudf/ReductionTest.java index f63eddfe415e..e8cad867da68 100644 --- a/java/src/test/java/ai/rapids/cudf/ReductionTest.java +++ b/java/src/test/java/ai/rapids/cudf/ReductionTest.java @@ -13,7 +13,6 @@ import org.junit.jupiter.params.provider.Arguments; import org.junit.jupiter.params.provider.MethodSource; -import java.math.BigInteger; import java.util.EnumSet; import java.util.List; import java.util.stream.Stream; @@ -695,24 +694,24 @@ void testSumWithOverflowNoOverflow() { @Test void testSumWithOverflowPositiveOverflow() { - // Long.MAX_VALUE + 1 wraps via two's complement to Long.MIN_VALUE. + // Sum is zeroed when overflow is detected; the flag is the source of truth. try (ColumnVector cv = ColumnVector.fromLongs(Long.MAX_VALUE, 1L); Scalar result = cv.reduce(ReductionAggregation.sumWithOverflow(), DType.STRUCT)) { SumWithOverflowResult r = readSumWithOverflow(result); assertTrue(r.sumValid); - assertEquals(Long.MIN_VALUE, r.sumValue); + assertEquals(0L, r.sumValue); assertTrue(r.overflow); } } @Test void testSumWithOverflowNegativeOverflow() { - // Long.MIN_VALUE + (-1) wraps via two's complement to Long.MAX_VALUE. + // Sum is zeroed when overflow is detected; the flag is the source of truth. try (ColumnVector cv = ColumnVector.fromLongs(Long.MIN_VALUE, -1L); Scalar result = cv.reduce(ReductionAggregation.sumWithOverflow(), DType.STRUCT)) { SumWithOverflowResult r = readSumWithOverflow(result); assertTrue(r.sumValid); - assertEquals(Long.MAX_VALUE, r.sumValue); + assertEquals(0L, r.sumValue); assertTrue(r.overflow); } } @@ -738,38 +737,9 @@ void testSumWithOverflowAllNullColumn() { } @Test - void testSumWithOverflowRejectsNonInt64() { - try (ColumnVector cv = ColumnVector.fromInts(1, 2, 3)) { - assertThrows(CudfException.class, () -> - cv.reduce(ReductionAggregation.sumWithOverflow(), DType.STRUCT).close()); - } - } - - // The reduction path of cudf::SUM_WITH_OVERFLOW is INT64-only (see - // cpp/src/reductions/reductions.cpp's `requires(std::is_same_v)`). - // The three tests below pin down the throw contract for fixed-point inputs; - // if a future cudf change broadens that requires-clause to accept decimals, - // these tests will start failing — the signal to expose decimal reduction. - @Test - void testSumWithOverflowReductionRejectsDecimal32() { - try (ColumnVector cv = ColumnVector.decimalFromInts(-2, 100, 200, 300)) { - assertThrows(CudfException.class, () -> - cv.reduce(ReductionAggregation.sumWithOverflow(), DType.STRUCT).close()); - } - } - - @Test - void testSumWithOverflowReductionRejectsDecimal64() { - try (ColumnVector cv = ColumnVector.decimalFromLongs(-4, 100L, 200L, 300L)) { - assertThrows(CudfException.class, () -> - cv.reduce(ReductionAggregation.sumWithOverflow(), DType.STRUCT).close()); - } - } - - @Test - void testSumWithOverflowReductionRejectsDecimal128() { - try (ColumnVector cv = ColumnVector.decimalFromBigInt(-10, - BigInteger.valueOf(100), BigInteger.valueOf(200), BigInteger.valueOf(300))) { + void testSumWithOverflowRejectsUnsupportedTypes() { + // SUM_WITH_OVERFLOW supports signed integers and decimals; float remains unsupported. + try (ColumnVector cv = ColumnVector.fromFloats(1.0f, 2.0f, 3.0f)) { assertThrows(CudfException.class, () -> cv.reduce(ReductionAggregation.sumWithOverflow(), DType.STRUCT).close()); }