From 34f02c2a94ec9b96c7fad19aac98fc97bdf0a3c1 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 20 Jan 2021 16:00:29 -0500 Subject: [PATCH 01/11] Initial changes (10.2 not working) --- .../cudf/detail/aggregation/aggregation.cuh | 3 ++- cpp/tests/groupby/group_count_test.cpp | 17 +++++++++++++++++ 2 files changed, 19 insertions(+), 1 deletion(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index f7ed66cf386..3d18c20dd41 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -341,7 +341,8 @@ struct elementwise_aggregator { column_device_view source, size_type source_index) const noexcept { - update_target_element{}( + using DeviceType = device_storage_type_t; + update_target_element{}( target, target_index, source, source_index); } }; diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index aca19a668e4..948033da0fe 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -221,6 +221,8 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) auto const expect_vals_min = fp_wrapper{{0, 1, 2}, scale}; auto const expect_vals_max = fp_wrapper{{6, 9, 8}, scale}; + // group_by sort tests + auto agg1 = cudf::make_sum_aggregation(); test_single_agg( keys, vals, expect_keys, expect_vals_sum, std::move(agg1), force_use_sort_impl::YES); @@ -237,6 +239,21 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) EXPECT_THROW( test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), cudf::logic_error); + + // group_by hash tests + + auto agg5 = cudf::make_sum_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); + + // auto agg6 = cudf::make_min_aggregation(); + // test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); + + // auto agg7 = cudf::make_max_aggregation(); + // test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); + + // auto agg8 = cudf::make_product_aggregation(); + // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), + // cudf::logic_error); } } From a13b306964abad0f8338ecb0c7228e66407fc161 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 27 Jan 2021 22:31:59 -0500 Subject: [PATCH 02/11] Disable 10.2 --- .../cudf/detail/aggregation/aggregation.cuh | 6 ++++++ cpp/tests/groupby/group_count_test.cpp | 17 ++++++++++------- 2 files changed, 16 insertions(+), 7 deletions(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index 3d18c20dd41..b6cea150f23 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -341,7 +341,13 @@ struct elementwise_aggregator { column_device_view source, size_type source_index) const noexcept { +#if (__CUDACC_VER_MAJOR__ == 10) and (__CUDACC_VER_MINOR__ == 2) + release_assert(not cudf::is_fixed_point() && + "CUDA 10.2 does not support decimal32 or decimal64 hash groupby."); + using DeviceType = Source; +#else using DeviceType = device_storage_type_t; +#endif update_target_element{}( target, target_index, source, source_index); } diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index 948033da0fe..7b897d1f8f4 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -240,20 +240,23 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), cudf::logic_error); +#if !((__CUDACC_VER_MAJOR__ == 10) and (__CUDACC_VER_MINOR__ == 2)) + // group_by hash tests auto agg5 = cudf::make_sum_aggregation(); test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); - // auto agg6 = cudf::make_min_aggregation(); - // test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); + auto agg6 = cudf::make_min_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); + + auto agg7 = cudf::make_max_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); - // auto agg7 = cudf::make_max_aggregation(); - // test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); + auto agg8 = cudf::make_product_aggregation(); + EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), cudf::logic_error); - // auto agg8 = cudf::make_product_aggregation(); - // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), - // cudf::logic_error); +#endif } } From 2e6bb533f56ae2f4632073e267e966898dab2928 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 2 Feb 2021 11:10:14 -0500 Subject: [PATCH 03/11] Push code down to update_target_element --- .../cudf/detail/aggregation/aggregation.cuh | 38 +++++++++++-------- cpp/tests/groupby/group_count_test.cpp | 12 +++--- 2 files changed, 28 insertions(+), 22 deletions(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index b6cea150f23..a55b98336f9 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -26,6 +26,7 @@ #include #include +#include "cudf/utilities/type_dispatcher.hpp" namespace cudf { namespace detail { @@ -152,12 +153,11 @@ struct update_target_element< }; template -struct update_target_element< - Source, - aggregation::SUM, - target_has_nulls, - source_has_nulls, - std::enable_if_t() && !is_fixed_point()>> { +struct update_target_element()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, @@ -166,8 +166,21 @@ struct update_target_element< if (source_has_nulls and source.is_null(source_index)) { return; } using Target = target_type_t; - atomicAdd(&target.element(target_index), - static_cast(source.element(source_index))); + + // #if (__CUDACC_VER_MAJOR__ != 10) or (__CUDACC_VER_MINOR__ != 2) + + using DeviceTarget = device_storage_type_t; + using DeviceSource = device_storage_type_t; + + // #else + + // using DeviceTarget = Target; + // using DeviceSource = Source; + + // #endif + + atomicAdd(&target.element(target_index), + static_cast(source.element(source_index))); if (target_has_nulls and target.is_null(target_index)) { target.set_valid(target_index); } } @@ -341,14 +354,7 @@ struct elementwise_aggregator { column_device_view source, size_type source_index) const noexcept { -#if (__CUDACC_VER_MAJOR__ == 10) and (__CUDACC_VER_MINOR__ == 2) - release_assert(not cudf::is_fixed_point() && - "CUDA 10.2 does not support decimal32 or decimal64 hash groupby."); - using DeviceType = Source; -#else - using DeviceType = device_storage_type_t; -#endif - update_target_element{}( + update_target_element{}( target, target_index, source, source_index); } }; diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index 7b897d1f8f4..4750f6d947e 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -247,14 +247,14 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) auto agg5 = cudf::make_sum_aggregation(); test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); - auto agg6 = cudf::make_min_aggregation(); - test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); + // auto agg6 = cudf::make_min_aggregation(); + // test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); - auto agg7 = cudf::make_max_aggregation(); - test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); + // auto agg7 = cudf::make_max_aggregation(); + // test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); - auto agg8 = cudf::make_product_aggregation(); - EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), cudf::logic_error); + // auto agg8 = cudf::make_product_aggregation(); + // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), cudf::logic_error); #endif } From cd04a8e0f129f39f152bab0880290264ae6d21a3 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 2 Feb 2021 11:24:59 -0500 Subject: [PATCH 04/11] Clang format --- cpp/tests/groupby/group_count_test.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index 4750f6d947e..cfbcb2fb1d0 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -254,7 +254,8 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) // test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); // auto agg8 = cudf::make_product_aggregation(); - // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), cudf::logic_error); + // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), + // cudf::logic_error); #endif } From d22fb16935bfee0851bcd6470250b369239718e1 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 2 Feb 2021 23:51:26 -0500 Subject: [PATCH 05/11] Final changes --- .../cudf/detail/aggregation/aggregation.cuh | 91 +++++++++++++++++-- cpp/tests/groupby/group_count_test.cpp | 6 +- 2 files changed, 83 insertions(+), 14 deletions(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index a55b98336f9..b9866e72e75 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -130,6 +130,34 @@ struct update_target_element< } }; +template +struct update_target_element()>> { + __device__ void operator()(mutable_column_device_view target, + size_type target_index, + column_device_view source, + size_type source_index) const noexcept + { +#if (__CUDACC_VER_MAJOR__ != 10) or (__CUDACC_VER_MINOR__ != 2) + + if (source_has_nulls and source.is_null(source_index)) { return; } + + using Target = target_type_t; + using DeviceTarget = device_storage_type_t; + using DeviceSource = device_storage_type_t; + + atomicMin(&target.element(target_index), + static_cast(source.element(source_index))); + + if (target_has_nulls and target.is_null(target_index)) { target.set_valid(target_index); } + +#endif + } +}; + template struct update_target_element< Source, @@ -154,35 +182,78 @@ struct update_target_element< template struct update_target_element()>> { + std::enable_if_t()>> { __device__ void operator()(mutable_column_device_view target, size_type target_index, column_device_view source, size_type source_index) const noexcept { - if (source_has_nulls and source.is_null(source_index)) { return; } - - using Target = target_type_t; +#if (__CUDACC_VER_MAJOR__ != 10) or (__CUDACC_VER_MINOR__ != 2) - // #if (__CUDACC_VER_MAJOR__ != 10) or (__CUDACC_VER_MINOR__ != 2) + if (source_has_nulls and source.is_null(source_index)) { return; } + using Target = target_type_t; using DeviceTarget = device_storage_type_t; using DeviceSource = device_storage_type_t; - // #else + atomicMax(&target.element(target_index), + static_cast(source.element(source_index))); - // using DeviceTarget = Target; - // using DeviceSource = Source; + if (target_has_nulls and target.is_null(target_index)) { target.set_valid(target_index); } - // #endif +#endif + } +}; + +template +struct update_target_element< + Source, + aggregation::SUM, + target_has_nulls, + source_has_nulls, + std::enable_if_t() && !is_fixed_point()>> { + __device__ void operator()(mutable_column_device_view target, + size_type target_index, + column_device_view source, + size_type source_index) const noexcept + { + if (source_has_nulls and source.is_null(source_index)) { return; } + + using Target = target_type_t; + atomicAdd(&target.element(target_index), + static_cast(source.element(source_index))); + + if (target_has_nulls and target.is_null(target_index)) { target.set_valid(target_index); } + } +}; + +template +struct update_target_element()>> { + __device__ void operator()(mutable_column_device_view target, + size_type target_index, + column_device_view source, + size_type source_index) const noexcept + { +#if (__CUDACC_VER_MAJOR__ != 10) or (__CUDACC_VER_MINOR__ != 2) + + if (source_has_nulls and source.is_null(source_index)) { return; } + + using Target = target_type_t; + using DeviceTarget = device_storage_type_t; + using DeviceSource = device_storage_type_t; atomicAdd(&target.element(target_index), static_cast(source.element(source_index))); if (target_has_nulls and target.is_null(target_index)) { target.set_valid(target_index); } +#endif } }; diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index cfbcb2fb1d0..d9b468265db 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -240,12 +240,11 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), cudf::logic_error); -#if !((__CUDACC_VER_MAJOR__ == 10) and (__CUDACC_VER_MINOR__ == 2)) // group_by hash tests - auto agg5 = cudf::make_sum_aggregation(); - test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); + // auto agg5 = cudf::make_sum_aggregation(); + // test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); // auto agg6 = cudf::make_min_aggregation(); // test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); @@ -257,7 +256,6 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), // cudf::logic_error); -#endif } } From 4c361d4dab21b9dbb126cbf4c6ce274e5541d0f3 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Tue, 2 Feb 2021 23:57:08 -0500 Subject: [PATCH 06/11] Clang format --- cpp/tests/groupby/group_count_test.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index d9b468265db..69be50d18c8 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -240,7 +240,6 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), cudf::logic_error); - // group_by hash tests // auto agg5 = cudf::make_sum_aggregation(); @@ -255,7 +254,6 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) // auto agg8 = cudf::make_product_aggregation(); // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), // cudf::logic_error); - } } From 5b6dcd1b2442efbcedad248b3b8f1e8257530ac4 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Wed, 3 Feb 2021 20:33:54 -0500 Subject: [PATCH 07/11] Clean up --- cpp/include/cudf/detail/aggregation/aggregation.cuh | 1 - cpp/tests/groupby/group_count_test.cpp | 1 + 2 files changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index b9866e72e75..3ed887e1269 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -26,7 +26,6 @@ #include #include -#include "cudf/utilities/type_dispatcher.hpp" namespace cudf { namespace detail { diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index 69be50d18c8..d23f26bc6e2 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -240,6 +240,7 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), cudf::logic_error); + // commented out until we drop support for CUDA 10.2 // group_by hash tests // auto agg5 = cudf::make_sum_aggregation(); From ac86d89d9f8344705e2fa5c4e04cbf5349eca850 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Feb 2021 02:26:50 -0500 Subject: [PATCH 08/11] Temporarily turn on test (shoud fail on 10.2) --- cpp/tests/groupby/group_count_test.cpp | 17 ++++++++--------- 1 file changed, 8 insertions(+), 9 deletions(-) diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index d23f26bc6e2..b2642758448 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -243,18 +243,17 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) // commented out until we drop support for CUDA 10.2 // group_by hash tests - // auto agg5 = cudf::make_sum_aggregation(); - // test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); + auto agg5 = cudf::make_sum_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); - // auto agg6 = cudf::make_min_aggregation(); - // test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); + auto agg6 = cudf::make_min_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); - // auto agg7 = cudf::make_max_aggregation(); - // test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); + auto agg7 = cudf::make_max_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); - // auto agg8 = cudf::make_product_aggregation(); - // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), - // cudf::logic_error); + auto agg8 = cudf::make_product_aggregation(); + EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), cudf::logic_error); } } From 364135eb59a7d8ef2303ed01e9a70532a1175b74 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Feb 2021 02:27:37 -0500 Subject: [PATCH 09/11] Revert "Temporarily turn on test (shoud fail on 10.2)" This reverts commit ac86d89d9f8344705e2fa5c4e04cbf5349eca850. --- cpp/tests/groupby/group_count_test.cpp | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index b2642758448..d23f26bc6e2 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -243,17 +243,18 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) // commented out until we drop support for CUDA 10.2 // group_by hash tests - auto agg5 = cudf::make_sum_aggregation(); - test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); + // auto agg5 = cudf::make_sum_aggregation(); + // test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); - auto agg6 = cudf::make_min_aggregation(); - test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); + // auto agg6 = cudf::make_min_aggregation(); + // test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); - auto agg7 = cudf::make_max_aggregation(); - test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); + // auto agg7 = cudf::make_max_aggregation(); + // test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); - auto agg8 = cudf::make_product_aggregation(); - EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), cudf::logic_error); + // auto agg8 = cudf::make_product_aggregation(); + // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), + // cudf::logic_error); } } From 44d7ca6fe663677ec3e1b2f8c8a63073fa491f42 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Feb 2021 11:15:09 -0500 Subject: [PATCH 10/11] Split up tests, disable correctly --- cpp/tests/groupby/group_count_test.cpp | 95 -------------------------- cpp/tests/groupby/group_max_test.cpp | 72 +++++++++++++++++++ cpp/tests/groupby/group_min_test.cpp | 73 ++++++++++++++++++++ cpp/tests/groupby/group_sum_test.cpp | 88 ++++++++++++++++++++++++ 4 files changed, 233 insertions(+), 95 deletions(-) diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index d23f26bc6e2..5147e150c58 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -201,101 +201,6 @@ TYPED_TEST(FixedPointTestBothReps, GroupByCount) test_single_agg(keys, vals, expect_keys, expect_vals, std::move(agg2)); } -TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValue) -{ - using namespace numeric; - using decimalXX = TypeParam; - using RepType = cudf::device_storage_type_t; - using fp_wrapper = cudf::test::fixed_point_column_wrapper; - using fp64_wrapper = cudf::test::fixed_point_column_wrapper; - using K = int32_t; - - for (auto const i : {2, 1, 0, -1, -2}) { - auto const scale = scale_type{i}; - auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; - auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; - - auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; - - auto const expect_vals_sum = fp64_wrapper{{9, 19, 17}, scale}; - auto const expect_vals_min = fp_wrapper{{0, 1, 2}, scale}; - auto const expect_vals_max = fp_wrapper{{6, 9, 8}, scale}; - - // group_by sort tests - - auto agg1 = cudf::make_sum_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_sum, std::move(agg1), force_use_sort_impl::YES); - - auto agg2 = cudf::make_min_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_min, std::move(agg2), force_use_sort_impl::YES); - - auto agg3 = cudf::make_max_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_max, std::move(agg3), force_use_sort_impl::YES); - - auto agg4 = cudf::make_product_aggregation(); - EXPECT_THROW( - test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), - cudf::logic_error); - - // commented out until we drop support for CUDA 10.2 - // group_by hash tests - - // auto agg5 = cudf::make_sum_aggregation(); - // test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); - - // auto agg6 = cudf::make_min_aggregation(); - // test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); - - // auto agg7 = cudf::make_max_aggregation(); - // test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); - - // auto agg8 = cudf::make_product_aggregation(); - // EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), - // cudf::logic_error); - } -} - -TYPED_TEST(FixedPointTestBothReps, GroupBySumProductMinMaxDecimalAsValueAndKey) -{ - using namespace numeric; - using decimalXX = TypeParam; - using RepType = cudf::device_storage_type_t; - using fp_wrapper = cudf::test::fixed_point_column_wrapper; - using fp64_wrapper = cudf::test::fixed_point_column_wrapper; - - for (auto const i : {2, 1, 0, -1, -2}) { - auto const scale = scale_type{i}; - auto const keys = fp_wrapper{{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}, scale}; - auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; - - auto const expect_keys = fp_wrapper{{1, 2, 3}, scale}; - - auto const expect_vals_sum = fp64_wrapper{{9, 19, 17}, scale}; - auto const expect_vals_min = fp_wrapper{{0, 1, 2}, scale}; - auto const expect_vals_max = fp_wrapper{{6, 9, 8}, scale}; - - auto agg1 = cudf::make_sum_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_sum, std::move(agg1), force_use_sort_impl::YES); - - auto agg2 = cudf::make_min_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_min, std::move(agg2), force_use_sort_impl::YES); - - auto agg3 = cudf::make_max_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_max, std::move(agg3), force_use_sort_impl::YES); - - auto agg4 = cudf::make_product_aggregation(); - EXPECT_THROW( - test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), - cudf::logic_error); - } -} - struct groupby_dictionary_count_test : public cudf::test::BaseFixture { }; diff --git a/cpp/tests/groupby/group_max_test.cpp b/cpp/tests/groupby/group_max_test.cpp index cc55e70910a..a786ebce001 100644 --- a/cpp/tests/groupby/group_max_test.cpp +++ b/cpp/tests/groupby/group_max_test.cpp @@ -196,5 +196,77 @@ TEST_F(groupby_dictionary_max_test, basic) force_use_sort_impl::YES); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, GroupBySortMaxDecimalAsValue) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + using K = int32_t; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals_max = fp_wrapper{{6, 9, 8}, scale}; + + auto agg3 = cudf::make_max_aggregation(); + test_single_agg( + keys, vals, expect_keys, expect_vals_max, std::move(agg3), force_use_sort_impl::YES); + } +} + +// This test will not work until the following ptxas bug is fixed in 10.2 +// https://nvbugswb.nvidia.com/NvBugs5/SWBug.aspx?bugid=3186317&cp= +TYPED_TEST(FixedPointTestBothReps, DISABLED_GroupByHashMaxDecimalAsValue) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + using K = int32_t; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals_max = fp_wrapper{{6, 9, 8}, scale}; + + auto agg7 = cudf::make_max_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_max, std::move(agg7)); + } +} + +TYPED_TEST(FixedPointTestBothReps, GroupBySortMaxDecimalAsValueAndKey) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fp_wrapper{{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}, scale}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fp_wrapper{{1, 2, 3}, scale}; + auto const expect_vals_max = fp_wrapper{{6, 9, 8}, scale}; + + auto agg3 = cudf::make_max_aggregation(); + test_single_agg( + keys, vals, expect_keys, expect_vals_max, std::move(agg3), force_use_sort_impl::YES); + } +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/group_min_test.cpp b/cpp/tests/groupby/group_min_test.cpp index ca095f08547..922185b4bac 100644 --- a/cpp/tests/groupby/group_min_test.cpp +++ b/cpp/tests/groupby/group_min_test.cpp @@ -196,5 +196,78 @@ TEST_F(groupby_dictionary_min_test, basic) force_use_sort_impl::YES); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, GroupBySortMinDecimalAsValue) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + + using K = int32_t; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals_min = fp_wrapper{{0, 1, 2}, scale}; + + auto agg2 = cudf::make_min_aggregation(); + test_single_agg( + keys, vals, expect_keys, expect_vals_min, std::move(agg2), force_use_sort_impl::YES); + } +} + +// This test will not work until the following ptxas bug is fixed in 10.2 +// https://nvbugswb.nvidia.com/NvBugs5/SWBug.aspx?bugid=3186317&cp= +TYPED_TEST(FixedPointTestBothReps, DISABLED_GroupByHashMinDecimalAsValue) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + using K = int32_t; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals_min = fp_wrapper{{0, 1, 2}, scale}; + + auto agg6 = cudf::make_min_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_min, std::move(agg6)); + } +} + +TYPED_TEST(FixedPointTestBothReps, GroupBySortMinDecimalAsValueAndKey) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fp_wrapper{{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}, scale}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fp_wrapper{{1, 2, 3}, scale}; + auto const expect_vals_min = fp_wrapper{{0, 1, 2}, scale}; + + auto agg2 = cudf::make_min_aggregation(); + test_single_agg( + keys, vals, expect_keys, expect_vals_min, std::move(agg2), force_use_sort_impl::YES); + } +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/group_sum_test.cpp b/cpp/tests/groupby/group_sum_test.cpp index fb1d4184f62..81150c2a718 100644 --- a/cpp/tests/groupby/group_sum_test.cpp +++ b/cpp/tests/groupby/group_sum_test.cpp @@ -157,5 +157,93 @@ TYPED_TEST(groupby_sum_test, DISABLED_dictionary) keys, vals, expect_keys, expect_vals, cudf::make_sum_aggregation(), force_use_sort_impl::YES); } +template +struct FixedPointTestBothReps : public cudf::test::BaseFixture { +}; + +TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); + +TYPED_TEST(FixedPointTestBothReps, GroupBySortSumDecimalAsValue) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + using fp64_wrapper = cudf::test::fixed_point_column_wrapper; + using K = int32_t; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals_sum = fp64_wrapper{{9, 19, 17}, scale}; + + auto agg1 = cudf::make_sum_aggregation(); + test_single_agg( + keys, vals, expect_keys, expect_vals_sum, std::move(agg1), force_use_sort_impl::YES); + + auto agg4 = cudf::make_product_aggregation(); + EXPECT_THROW( + test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), + cudf::logic_error); + } +} + +// This test will not work until the following ptxas bug is fixed in 10.2 +// https://nvbugswb.nvidia.com/NvBugs5/SWBug.aspx?bugid=3186317&cp= +TYPED_TEST(FixedPointTestBothReps, DISABLED_GroupByHashSumDecimalAsValue) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + using fp64_wrapper = cudf::test::fixed_point_column_wrapper; + using K = int32_t; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals_sum = fp64_wrapper{{9, 19, 17}, scale}; + + auto agg5 = cudf::make_sum_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals_sum, std::move(agg5)); + + auto agg8 = cudf::make_product_aggregation(); + EXPECT_THROW(test_single_agg(keys, vals, expect_keys, {}, std::move(agg8)), cudf::logic_error); + } +} + +TYPED_TEST(FixedPointTestBothReps, GroupBySortSumDecimalAsValueAndKey) +{ + using namespace numeric; + using decimalXX = TypeParam; + using RepType = cudf::device_storage_type_t; + using fp_wrapper = cudf::test::fixed_point_column_wrapper; + using fp64_wrapper = cudf::test::fixed_point_column_wrapper; + + for (auto const i : {2, 1, 0, -1, -2}) { + auto const scale = scale_type{i}; + auto const keys = fp_wrapper{{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}, scale}; + auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; + + auto const expect_keys = fp_wrapper{{1, 2, 3}, scale}; + auto const expect_vals_sum = fp64_wrapper{{9, 19, 17}, scale}; + + auto agg1 = cudf::make_sum_aggregation(); + test_single_agg( + keys, vals, expect_keys, expect_vals_sum, std::move(agg1), force_use_sort_impl::YES); + + auto agg4 = cudf::make_product_aggregation(); + EXPECT_THROW( + test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), + cudf::logic_error); + } +} + } // namespace test } // namespace cudf From 06f2add9f7b1a36708922f514d4c946f8aaea660 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra Date: Thu, 4 Feb 2021 16:40:43 -0500 Subject: [PATCH 11/11] Address PR comment --- cpp/tests/groupby/group_max_test.cpp | 21 ------------------- cpp/tests/groupby/group_min_test.cpp | 21 ------------------- cpp/tests/groupby/group_sum_test.cpp | 27 ------------------------- cpp/tests/groupby/groupby_keys_test.cpp | 3 ++- 4 files changed, 2 insertions(+), 70 deletions(-) diff --git a/cpp/tests/groupby/group_max_test.cpp b/cpp/tests/groupby/group_max_test.cpp index a786ebce001..0a13510a948 100644 --- a/cpp/tests/groupby/group_max_test.cpp +++ b/cpp/tests/groupby/group_max_test.cpp @@ -247,26 +247,5 @@ TYPED_TEST(FixedPointTestBothReps, DISABLED_GroupByHashMaxDecimalAsValue) } } -TYPED_TEST(FixedPointTestBothReps, GroupBySortMaxDecimalAsValueAndKey) -{ - using namespace numeric; - using decimalXX = TypeParam; - using RepType = cudf::device_storage_type_t; - using fp_wrapper = cudf::test::fixed_point_column_wrapper; - - for (auto const i : {2, 1, 0, -1, -2}) { - auto const scale = scale_type{i}; - auto const keys = fp_wrapper{{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}, scale}; - auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; - - auto const expect_keys = fp_wrapper{{1, 2, 3}, scale}; - auto const expect_vals_max = fp_wrapper{{6, 9, 8}, scale}; - - auto agg3 = cudf::make_max_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_max, std::move(agg3), force_use_sort_impl::YES); - } -} - } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/group_min_test.cpp b/cpp/tests/groupby/group_min_test.cpp index 922185b4bac..4cd0c9864ad 100644 --- a/cpp/tests/groupby/group_min_test.cpp +++ b/cpp/tests/groupby/group_min_test.cpp @@ -248,26 +248,5 @@ TYPED_TEST(FixedPointTestBothReps, DISABLED_GroupByHashMinDecimalAsValue) } } -TYPED_TEST(FixedPointTestBothReps, GroupBySortMinDecimalAsValueAndKey) -{ - using namespace numeric; - using decimalXX = TypeParam; - using RepType = cudf::device_storage_type_t; - using fp_wrapper = cudf::test::fixed_point_column_wrapper; - - for (auto const i : {2, 1, 0, -1, -2}) { - auto const scale = scale_type{i}; - auto const keys = fp_wrapper{{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}, scale}; - auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; - - auto const expect_keys = fp_wrapper{{1, 2, 3}, scale}; - auto const expect_vals_min = fp_wrapper{{0, 1, 2}, scale}; - - auto agg2 = cudf::make_min_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_min, std::move(agg2), force_use_sort_impl::YES); - } -} - } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/group_sum_test.cpp b/cpp/tests/groupby/group_sum_test.cpp index 81150c2a718..f07c695b17a 100644 --- a/cpp/tests/groupby/group_sum_test.cpp +++ b/cpp/tests/groupby/group_sum_test.cpp @@ -218,32 +218,5 @@ TYPED_TEST(FixedPointTestBothReps, DISABLED_GroupByHashSumDecimalAsValue) } } -TYPED_TEST(FixedPointTestBothReps, GroupBySortSumDecimalAsValueAndKey) -{ - using namespace numeric; - using decimalXX = TypeParam; - using RepType = cudf::device_storage_type_t; - using fp_wrapper = cudf::test::fixed_point_column_wrapper; - using fp64_wrapper = cudf::test::fixed_point_column_wrapper; - - for (auto const i : {2, 1, 0, -1, -2}) { - auto const scale = scale_type{i}; - auto const keys = fp_wrapper{{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}, scale}; - auto const vals = fp_wrapper{{0, 1, 2, 3, 4, 5, 6, 7, 8, 9}, scale}; - - auto const expect_keys = fp_wrapper{{1, 2, 3}, scale}; - auto const expect_vals_sum = fp64_wrapper{{9, 19, 17}, scale}; - - auto agg1 = cudf::make_sum_aggregation(); - test_single_agg( - keys, vals, expect_keys, expect_vals_sum, std::move(agg1), force_use_sort_impl::YES); - - auto agg4 = cudf::make_product_aggregation(); - EXPECT_THROW( - test_single_agg(keys, vals, expect_keys, {}, std::move(agg4), force_use_sort_impl::YES), - cudf::logic_error); - } -} - } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/groupby_keys_test.cpp b/cpp/tests/groupby/groupby_keys_test.cpp index 553c16e8349..06ec9eb8968 100644 --- a/cpp/tests/groupby/groupby_keys_test.cpp +++ b/cpp/tests/groupby/groupby_keys_test.cpp @@ -28,7 +28,8 @@ template struct groupby_keys_test : public cudf::test::BaseFixture { }; -using supported_types = cudf::test::Types; +using supported_types = cudf::test:: + Types; TYPED_TEST_CASE(groupby_keys_test, supported_types);