diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index f7ed66cf386..3ed887e1269 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -129,6 +129,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, @@ -151,6 +179,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; + + atomicMax(&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, @@ -173,6 +229,33 @@ 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; + + 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 + } +}; + /** * @brief Function object to update a single element in a target column using * the dictionary key addressed by the specific index. diff --git a/cpp/tests/groupby/group_count_test.cpp b/cpp/tests/groupby/group_count_test.cpp index aca19a668e4..5147e150c58 100644 --- a/cpp/tests/groupby/group_count_test.cpp +++ b/cpp/tests/groupby/group_count_test.cpp @@ -201,83 +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}; - - 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); - } -} - -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..0a13510a948 100644 --- a/cpp/tests/groupby/group_max_test.cpp +++ b/cpp/tests/groupby/group_max_test.cpp @@ -196,5 +196,56 @@ 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)); + } +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/group_min_test.cpp b/cpp/tests/groupby/group_min_test.cpp index ca095f08547..4cd0c9864ad 100644 --- a/cpp/tests/groupby/group_min_test.cpp +++ b/cpp/tests/groupby/group_min_test.cpp @@ -196,5 +196,57 @@ 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)); + } +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/group_sum_test.cpp b/cpp/tests/groupby/group_sum_test.cpp index fb1d4184f62..f07c695b17a 100644 --- a/cpp/tests/groupby/group_sum_test.cpp +++ b/cpp/tests/groupby/group_sum_test.cpp @@ -157,5 +157,66 @@ 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); + } +} + } // 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);