Skip to content

Commit

Permalink
Implement cudf::group_by (hash) for decimal32 and decimal64 (#7190
Browse files Browse the repository at this point in the history 
)

Follow up PR to #7169

This PR resolves a part of #3556.

Authors:
  - Conor Hoekstra (@codereport)

Approvers:
  - David (@davidwendt)
  - Jake Hemstad (@jrhemstad)
  - Devavret Makkar (@devavret)

URL: #7190
  • Loading branch information
@codereport
codereport authored Feb 5, 2021
1 parent 5a7c5d9 commit 5d151a7
Show file tree
Hide file tree
Showing 6 changed files with 249 additions and 78 deletions.
83 changes: 83 additions & 0 deletions cpp/include/cudf/detail/aggregation/aggregation.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand All @@ -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,
Expand All @@ -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.
Expand Down
77 changes: 0 additions & 77 deletions cpp/tests/groupby/group_count_test.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 {
};

Expand Down
51 changes: 51 additions & 0 deletions cpp/tests/groupby/group_max_test.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
52 changes: 52 additions & 0 deletions cpp/tests/groupby/group_min_test.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
61 changes: 61 additions & 0 deletions cpp/tests/groupby/group_sum_test.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
3 changes: 2 additions & 1 deletion cpp/tests/groupby/groupby_keys_test.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down

0 comments on commit 5d151a7

Please sign in to comment.