From 5c3ea93ee5107a5e867b4623cfa5be49d5f78e9e Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Mon, 9 Jan 2023 11:33:12 -0800 Subject: [PATCH 1/3] Fix SUM/MEAN aggregation type support. --- .../cudf/detail/aggregation/aggregation.cuh | 6 +++--- .../cudf/detail/aggregation/aggregation.hpp | 21 +++++++++---------- cpp/src/rolling/detail/rolling.cuh | 11 ++-------- 3 files changed, 15 insertions(+), 23 deletions(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index 818e8cd7cc6..f13166d5321 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -232,8 +232,8 @@ struct update_target_element< aggregation::SUM, target_has_nulls, source_has_nulls, - std::enable_if_t() && cudf::has_atomic_support() && - !is_fixed_point()>> { + std::enable_if_t() && cudf::has_atomic_support() && + !cudf::is_fixed_point() && !cudf::is_timestamp()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, diff --git a/cpp/include/cudf/detail/aggregation/aggregation.hpp b/cpp/include/cudf/detail/aggregation/aggregation.hpp index 75027c78a68..360c314f2db 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.hpp +++ b/cpp/include/cudf/detail/aggregation/aggregation.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -1154,9 +1154,7 @@ struct target_type_impl { using type = bool; }; -// Always use `double` for MEAN -// Except for chrono types where result is chrono. (Use FloorDiv) -// TODO: MEAN should be only be enabled for duration types - not for timestamps +// Always use `double` for MEAN except for durations and fixed point types. template struct target_type_impl< Source, @@ -1167,10 +1165,10 @@ struct target_type_impl< }; template -struct target_type_impl< - Source, - k, - std::enable_if_t<(is_chrono() or is_fixed_point()) && (k == aggregation::MEAN)>> { +struct target_type_impl() or is_fixed_point()) && + (k == aggregation::MEAN)>> { using type = Source; }; @@ -1206,10 +1204,11 @@ struct target_type_impl< using type = Source; }; -// Summing/Multiplying chrono types, use same type accumulator -// TODO: Sum/Product should only be enabled for duration types - not for timestamps +// Summing duration types, use same type accumulator template -struct target_type_impl() && is_sum_product_agg(k)>> { +struct target_type_impl() && (k == aggregation::SUM)>> { using type = Source; }; diff --git a/cpp/src/rolling/detail/rolling.cuh b/cpp/src/rolling/detail/rolling.cuh index fcc85b4f913..d996f88ca49 100644 --- a/cpp/src/rolling/detail/rolling.cuh +++ b/cpp/src/rolling/detail/rolling.cuh @@ -84,16 +84,9 @@ struct DeviceRolling { static constexpr bool is_supported() { return cudf::detail::is_valid_aggregation() && has_corresponding_operator() && - // TODO: Delete all this extra logic once is_valid_aggregation<> cleans up some edge - // cases it isn't handling. - // MIN/MAX supports all fixed width types + // MIN/MAX only supports fixed width types (((O == aggregation::MIN || O == aggregation::MAX) && cudf::is_fixed_width()) || - - // SUM supports all fixed width types except timestamps - ((O == aggregation::SUM) && (cudf::is_fixed_width() && !cudf::is_timestamp())) || - - // MEAN supports numeric and duration - ((O == aggregation::MEAN) && (cudf::is_numeric() || cudf::is_duration()))); + (O == aggregation::SUM) || (O == aggregation::MEAN)); } // operations we do support From dc09c3f468f41d02284a6df07efec6e8f9ae9f0b Mon Sep 17 00:00:00 2001 From: Karthikeyan Natarajan Date: Tue, 10 Jan 2023 19:14:19 +0530 Subject: [PATCH 2/3] update TypedRollingEmptyInputTest for timestamp types --- cpp/tests/rolling/empty_input_test.cpp | 20 ++++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/cpp/tests/rolling/empty_input_test.cpp b/cpp/tests/rolling/empty_input_test.cpp index 001a4b1279e..b200ccf981f 100644 --- a/cpp/tests/rolling/empty_input_test.cpp +++ b/cpp/tests/rolling/empty_input_test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2021-2022, NVIDIA CORPORATION. + * Copyright (c) 2021-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -14,6 +14,7 @@ * limitations under the License. */ +#include "cudf/utilities/traits.hpp" #include #include #include @@ -24,6 +25,7 @@ #include #include #include +#include namespace { // Helper functions to construct rolling window operators. @@ -187,22 +189,32 @@ TYPED_TEST(TypedRollingEmptyInputTest, EmptyFixedWidthInputs) /// `SUM` returns 64-bit promoted types for integral/decimal input. /// For other fixed-width input types, the same type is returned. + /// Timestamp types are not supported. { auto aggs = agg_vector_t{}; aggs.emplace_back(sum()); using expected_type = cudf::detail::target_type_t; - rolling_output_type_matches(empty_input, aggs, type_to_id()); + if constexpr (cudf::is_timestamp()) + EXPECT_THROW(rolling_output_type_matches(empty_input, aggs, type_to_id()), + cudf::logic_error); + else + rolling_output_type_matches(empty_input, aggs, type_to_id()); } /// `MEAN` returns float64 for all numeric types, - /// except for chrono-types, which yield the same chrono-type. + /// except for duration-types, which yield the same duration-type. + /// Timestamp types are not supported. { auto aggs = agg_vector_t{}; aggs.emplace_back(mean()); using expected_type = cudf::detail::target_type_t; - rolling_output_type_matches(empty_input, aggs, type_to_id()); + if constexpr (cudf::is_timestamp()) + EXPECT_THROW(rolling_output_type_matches(empty_input, aggs, type_to_id()), + cudf::logic_error); + else + rolling_output_type_matches(empty_input, aggs, type_to_id()); } /// For an input type `T`, `COLLECT_LIST` returns a column of type `list`. From fad248fac70acbb01429d4223ee27b15b0b22dea Mon Sep 17 00:00:00 2001 From: Karthikeyan <6488848+karthikeyann@users.noreply.github.com> Date: Thu, 19 Jan 2023 23:52:38 +0530 Subject: [PATCH 3/3] Apply suggestions from code review --- cpp/tests/rolling/empty_input_test.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/cpp/tests/rolling/empty_input_test.cpp b/cpp/tests/rolling/empty_input_test.cpp index f275f8db0fa..aca1cbf40b7 100644 --- a/cpp/tests/rolling/empty_input_test.cpp +++ b/cpp/tests/rolling/empty_input_test.cpp @@ -14,7 +14,6 @@ * limitations under the License. */ -#include "cudf/utilities/traits.hpp" #include #include #include @@ -25,7 +24,6 @@ #include #include #include -#include namespace { // Helper functions to construct rolling window operators.