Implement cudf::group_by (hash) for decimal32 and decimal64

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

@@ -129,6 +129,34 @@ struct update_target_element<
   }
 };
 
+template <typename Source, bool target_has_nulls, bool source_has_nulls>
+struct update_target_element<Source,
+                             aggregation::MIN,
+                             target_has_nulls,
+                             source_has_nulls,
+                             std::enable_if_t<is_fixed_point<Source>()>> {
+  __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<Source, aggregation::MIN>;
+    using DeviceTarget = device_storage_type_t<Target>;
+    using DeviceSource = device_storage_type_t<Source>;
+
+    atomicMin(&target.element<DeviceTarget>(target_index),
+              static_cast<DeviceTarget>(source.element<DeviceSource>(source_index)));
+
+    if (target_has_nulls and target.is_null(target_index)) { target.set_valid(target_index); }
+
+#endif
+  }
+};
+
 template <typename Source, bool target_has_nulls, bool source_has_nulls>
 struct update_target_element<
   Source,
@@ -151,6 +179,34 @@ struct update_target_element<
   }
 };
 
+template <typename Source, bool target_has_nulls, bool source_has_nulls>
+struct update_target_element<Source,
+                             aggregation::MAX,
+                             target_has_nulls,
+                             source_has_nulls,
+                             std::enable_if_t<is_fixed_point<Source>()>> {
+  __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<Source, aggregation::MAX>;
+    using DeviceTarget = device_storage_type_t<Target>;
+    using DeviceSource = device_storage_type_t<Source>;
+
+    atomicMax(&target.element<DeviceTarget>(target_index),
+              static_cast<DeviceTarget>(source.element<DeviceSource>(source_index)));
+
+    if (target_has_nulls and target.is_null(target_index)) { target.set_valid(target_index); }
+
+#endif
+  }
+};
+
 template <typename Source, bool target_has_nulls, bool source_has_nulls>
 struct update_target_element<
   Source,
@@ -173,6 +229,33 @@ struct update_target_element<
   }
 };
 
+template <typename Source, bool target_has_nulls, bool source_has_nulls>
+struct update_target_element<Source,
+                             aggregation::SUM,
+                             target_has_nulls,
+                             source_has_nulls,
+                             std::enable_if_t<is_fixed_point<Source>()>> {
+  __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<Source, aggregation::SUM>;
+    using DeviceTarget = device_storage_type_t<Target>;
+    using DeviceSource = device_storage_type_t<Source>;
+
+    atomicAdd(&target.element<DeviceTarget>(target_index),
+              static_cast<DeviceTarget>(source.element<DeviceSource>(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.

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<decimalXX>;
-  using fp_wrapper   = cudf::test::fixed_point_column_wrapper<RepType>;
-  using fp64_wrapper = cudf::test::fixed_point_column_wrapper<int64_t>;
-  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<K>{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<K>{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<decimalXX>;
-  using fp_wrapper   = cudf::test::fixed_point_column_wrapper<RepType>;
-  using fp64_wrapper = cudf::test::fixed_point_column_wrapper<int64_t>;
-
-  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 {
 };
 

cpp/tests/groupby/group_max_test.cpp

@@ -196,5 +196,56 @@ TEST_F(groupby_dictionary_max_test, basic)
                   force_use_sort_impl::YES);
 }
 
+template <typename T>
+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<decimalXX>;
+  using fp_wrapper = cudf::test::fixed_point_column_wrapper<RepType>;
+  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<K>{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<K>{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<decimalXX>;
+  using fp_wrapper = cudf::test::fixed_point_column_wrapper<RepType>;
+  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<K>{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<K>{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

cpp/tests/groupby/group_min_test.cpp

@@ -196,5 +196,57 @@ TEST_F(groupby_dictionary_min_test, basic)
                   force_use_sort_impl::YES);
 }
 
+template <typename T>
+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<decimalXX>;
+  using fp_wrapper = cudf::test::fixed_point_column_wrapper<RepType>;
+
+  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<K>{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<K>{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<decimalXX>;
+  using fp_wrapper = cudf::test::fixed_point_column_wrapper<RepType>;
+  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<K>{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<K>{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

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 <typename T>
+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<decimalXX>;
+  using fp_wrapper   = cudf::test::fixed_point_column_wrapper<RepType>;
+  using fp64_wrapper = cudf::test::fixed_point_column_wrapper<int64_t>;
+  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<K>{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<K>{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<decimalXX>;
+  using fp_wrapper   = cudf::test::fixed_point_column_wrapper<RepType>;
+  using fp64_wrapper = cudf::test::fixed_point_column_wrapper<int64_t>;
+  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<K>{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<K>{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

cpp/tests/groupby/groupby_keys_test.cpp

@@ -28,7 +28,8 @@ template <typename V>
 struct groupby_keys_test : public cudf::test::BaseFixture {
 };
 
-using supported_types = cudf::test::Types<int8_t, int16_t, int32_t, int64_t, float, double>;
+using supported_types = cudf::test::
+  Types<int8_t, int16_t, int32_t, int64_t, float, double, numeric::decimal32, numeric::decimal64>;
 
 TYPED_TEST_CASE(groupby_keys_test, supported_types);