From d08e0413839f31dd6ba19582e6fbb20b2205966c Mon Sep 17 00:00:00 2001 From: David Wendt <45795991+davidwendt@users.noreply.github.com> Date: Mon, 26 Apr 2021 21:51:07 -0400 Subject: [PATCH] Remove 10.2 workarounds in groupby functions for dictionary column types (#7949) Reference #7913 This removes the compiler restrictions in the libcudf groupby code that was disabled because of compiler issues on nvcc 10.2 when adding dictionary columns support. The corresponding disabled gtests have been enabled as well. Details of the compiler bug are documented here: https://nvbugswb.nvidia.com/NvBugs5/SWBug.aspx?bugid=3186317 This also fixes an issue where the sum-of-squares aggregation was not working for values specified as a dictionary column. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Karthikeyan (https://github.com/karthikeyann) - Nghia Truong (https://github.com/ttnghia) URL: https://github.com/rapidsai/cudf/pull/7949 --- .../cudf/detail/aggregation/aggregation.cuh | 57 +++++++++++++++++-- cpp/src/groupby/groupby.cu | 19 ------- cpp/src/strings/convert/convert_floats.cu | 3 +- cpp/tests/groupby/group_mean_test.cpp | 6 +- cpp/tests/groupby/group_std_test.cpp | 6 +- .../groupby/group_sum_of_squares_test.cpp | 4 +- cpp/tests/groupby/group_sum_test.cpp | 6 +- cpp/tests/groupby/group_var_test.cpp | 6 +- 8 files changed, 61 insertions(+), 46 deletions(-) diff --git a/cpp/include/cudf/detail/aggregation/aggregation.cuh b/cpp/include/cudf/detail/aggregation/aggregation.cuh index 4d78f5ef05a..1084ad5810d 100644 --- a/cpp/include/cudf/detail/aggregation/aggregation.cuh +++ b/cpp/include/cudf/detail/aggregation/aggregation.cuh @@ -262,7 +262,7 @@ struct update_target_element() && !is_fixed_point()>* = nullptr> __device__ void operator()(mutable_column_device_view& target, @@ -270,15 +270,11 @@ struct update_target_from_dictionary { column_device_view& d_dictionary, size_type source_index) const noexcept { -// This code will segfault in nvcc/ptxas 10.2 only -// https://nvbugswb.nvidia.com/NvBugs5/SWBug.aspx?bugid=3186317 -#if (__CUDACC_VER_MAJOR__ != 10) or (__CUDACC_VER_MINOR__ != 2) auto const keys = d_dictionary.child(cudf::dictionary_column_view::keys_column_index); auto const value = keys.element( static_cast(d_dictionary.element(source_index))); using Target = target_type_t; atomicAdd(&target.element(target_index), static_cast(value)); -#endif } template () || is_fixed_point()>* = nullptr> @@ -304,7 +300,7 @@ struct update_target_element()>* = nullptr> + __device__ void operator()(mutable_column_device_view& target, + size_type target_index, + column_device_view& d_dictionary, + size_type source_index) const noexcept + { + auto const keys = d_dictionary.child(cudf::dictionary_column_view::keys_column_index); + auto const value = keys.element( + static_cast(d_dictionary.element(source_index))); + using Target = target_type_t; + atomicAdd(&target.element(target_index), static_cast(value * value)); + } + template ()>* = nullptr> + __device__ void operator()(mutable_column_device_view& target, + size_type target_index, + column_device_view& d_dictionary, + size_type source_index) const noexcept {}; +}; + +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 (source_has_nulls and source.is_null(source_index)) { return; } + + type_dispatcher(source.child(cudf::dictionary_column_view::keys_column_index).type(), + update_target_from_dictionary_squares{}, + target, + target_index, + source, + source_index); + + if (target_has_nulls and target.is_null(target_index)) { target.set_valid(target_index); } + } +}; + template struct update_target_element requests) }); }), "Invalid type/aggregation combination."); - -// The aggregations listed in the lambda below will not work with a values column of type -// dictionary if this is compiled with nvcc/ptxas 10.2. -// https://nvbugswb.nvidia.com/NvBugs5/SWBug.aspx?bugid=3186317&cp= -#if (__CUDACC_VER_MAJOR__ == 10) and (__CUDACC_VER_MINOR__ == 2) - CUDF_EXPECTS( - std::all_of( - requests.begin(), - requests.end(), - [](auto const& request) { - return std::all_of( - request.aggregations.begin(), request.aggregations.end(), [&request](auto const& agg) { - return (!cudf::is_dictionary(request.values.type()) || - !(agg->kind == aggregation::SUM or agg->kind == aggregation::MEAN or - agg->kind == aggregation::STD or agg->kind == aggregation::VARIANCE)); - }); - }), - "dictionary type not supported for this aggregation"); -#endif } } // namespace diff --git a/cpp/src/strings/convert/convert_floats.cu b/cpp/src/strings/convert/convert_floats.cu index 2eb64e65d96..1b8d9d7366e 100644 --- a/cpp/src/strings/convert/convert_floats.cu +++ b/cpp/src/strings/convert/convert_floats.cu @@ -124,8 +124,7 @@ __device__ inline double stod(string_view const& d_str) else if (exp_ten < std::numeric_limits::min_exponent10) return double{0}; - // using exp10() since the pow(10.0,exp_ten) function is - // very inaccurate in 10.2: http://nvbugs/2971187 + // exp10() is faster than pow(10.0,exp_ten) double const base = sign * static_cast(digits) * exp10(static_cast(1 - num_digits)); double const exponent = exp10(static_cast(exp_ten)); diff --git a/cpp/tests/groupby/group_mean_test.cpp b/cpp/tests/groupby/group_mean_test.cpp index 6d12516a8cf..37a8ce4e2d6 100644 --- a/cpp/tests/groupby/group_mean_test.cpp +++ b/cpp/tests/groupby/group_mean_test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -141,9 +141,7 @@ TYPED_TEST(groupby_mean_test, null_keys_and_values) struct groupby_dictionary_mean_test : public cudf::test::BaseFixture { }; -// This tests will not work until the following ptxas bug is fixed in 10.2 -// https://nvbugswb.nvidia.com/NvBugs5/SWBug.aspx?bugid=3186317&cp= -TEST_F(groupby_dictionary_mean_test, DISABLED_basic) +TEST_F(groupby_dictionary_mean_test, basic) { using K = int32_t; using V = int16_t; diff --git a/cpp/tests/groupby/group_std_test.cpp b/cpp/tests/groupby/group_std_test.cpp index 776e24f1fc1..7cd034d309f 100644 --- a/cpp/tests/groupby/group_std_test.cpp +++ b/cpp/tests/groupby/group_std_test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -144,9 +144,7 @@ TYPED_TEST(groupby_std_test, ddof_non_default) } // clang-format on -// 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(groupby_std_test, DISABLED_dictionary) +TYPED_TEST(groupby_std_test, dictionary) { using K = int32_t; using V = TypeParam; diff --git a/cpp/tests/groupby/group_sum_of_squares_test.cpp b/cpp/tests/groupby/group_sum_of_squares_test.cpp index 24601d2b246..aad1df1bf27 100644 --- a/cpp/tests/groupby/group_sum_of_squares_test.cpp +++ b/cpp/tests/groupby/group_sum_of_squares_test.cpp @@ -124,9 +124,7 @@ TYPED_TEST(groupby_sum_of_squares_test, null_keys_and_values) } // clang-format on -// 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(groupby_sum_of_squares_test, DISABLED_dictionary) +TYPED_TEST(groupby_sum_of_squares_test, dictionary) { using K = int32_t; using V = TypeParam; diff --git a/cpp/tests/groupby/group_sum_test.cpp b/cpp/tests/groupby/group_sum_test.cpp index f07c695b17a..662634d9450 100644 --- a/cpp/tests/groupby/group_sum_test.cpp +++ b/cpp/tests/groupby/group_sum_test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -136,9 +136,7 @@ TYPED_TEST(groupby_sum_test, null_keys_and_values) } // clang-format on -// These tests 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(groupby_sum_test, DISABLED_dictionary) +TYPED_TEST(groupby_sum_test, dictionary) { using K = int32_t; using V = TypeParam; diff --git a/cpp/tests/groupby/group_var_test.cpp b/cpp/tests/groupby/group_var_test.cpp index 5fa52bad22a..62d6beaa906 100644 --- a/cpp/tests/groupby/group_var_test.cpp +++ b/cpp/tests/groupby/group_var_test.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2020, NVIDIA CORPORATION. + * Copyright (c) 2019-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -144,9 +144,7 @@ TYPED_TEST(groupby_var_test, ddof_non_default) } // clang-format on -// 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(groupby_var_test, DISABLED_dictionary) +TYPED_TEST(groupby_var_test, dictionary) { using K = int32_t; using V = TypeParam;