Remove 10.2 workarounds in groupby functions for dictionary column types

cpp/include/cudf/detail/aggregation/aggregation.cuh

@@ -262,23 +262,19 @@ struct update_target_element<Source,
  *
  * `target[target_index] = d_dictionary.keys[d_dictionary.indices[source_index]]`
  */
-struct update_target_from_dictionary {
+struct update_target_from_dictionary_sum {
   template <typename KeyType,
             std::enable_if_t<is_fixed_width<KeyType>() && !is_fixed_point<KeyType>()>* = nullptr>
   __device__ void operator()(mutable_column_device_view& target,
                              size_type target_index,
                              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<KeyType>(
       static_cast<cudf::size_type>(d_dictionary.element<dictionary32>(source_index)));
     using Target = target_type_t<KeyType, aggregation::SUM>;
     atomicAdd(&target.element<Target>(target_index), static_cast<Target>(value));
-#endif
   }
   template <typename KeyType,
             std::enable_if_t<!is_fixed_width<KeyType>() || is_fixed_point<KeyType>()>* = nullptr>
@@ -304,7 +300,7 @@ struct update_target_element<dictionary32, aggregation::SUM, target_has_nulls, s
     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{},
+                    update_target_from_dictionary_sum{},
                     target,
                     target_index,
                     source,
@@ -340,6 +336,55 @@ struct update_target_element<Source,
   }
 };
 
+/**
+ * @brief Function object to update a single element in a target column using
+ * the dictionary key addressed by the specific index.
+ *
+ * `target[target_index] = d_dictionary.keys[d_dictionary.indices[source_index]]^2`
+ */
+struct update_target_from_dictionary_squares {
+  template <typename KeyType, std::enable_if_t<is_product_supported<KeyType>()>* = 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<KeyType>(
+      static_cast<cudf::size_type>(d_dictionary.element<dictionary32>(source_index)));
+    using Target = target_type_t<KeyType, aggregation::SUM_OF_SQUARES>;
+    atomicAdd(&target.element<Target>(target_index), static_cast<Target>(value * value));
+  }
+  template <typename KeyType, std::enable_if_t<!is_product_supported<KeyType>()>* = 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 <bool target_has_nulls, bool source_has_nulls>
+struct update_target_element<dictionary32,
+                             aggregation::SUM_OF_SQUARES,
+                             target_has_nulls,
+                             source_has_nulls> {
+  __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 <typename Source, bool target_has_nulls, bool source_has_nulls>
 struct update_target_element<Source,
                              aggregation::PRODUCT,

cpp/src/groupby/groupby.cu

@@ -118,25 +118,6 @@ void verify_valid_requests(host_span<aggregation_request const> 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

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<double>::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<double>(digits) * exp10(static_cast<double>(1 - num_digits));
   double const exponent = exp10(static_cast<double>(exp_ten));

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;

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;

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;

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;

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;