Skip to content

Commit

Permalink
Remove 10.2 workarounds in groupby functions for dictionary column ty…
Browse files Browse the repository at this point in the history
…pes (#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: #7949
  • Loading branch information
davidwendt authored Apr 27, 2021
1 parent 1a0d304 commit d08e041
Showing 8 changed files with 61 additions and 46 deletions.
57 changes: 51 additions & 6 deletions cpp/include/cudf/detail/aggregation/aggregation.cuh
Original file line number Diff line number Diff line change
@@ -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,
19 changes: 0 additions & 19 deletions cpp/src/groupby/groupby.cu
Original file line number Diff line number Diff line change
@@ -119,25 +119,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
3 changes: 1 addition & 2 deletions cpp/src/strings/convert/convert_floats.cu
Original file line number Diff line number Diff line change
@@ -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));
6 changes: 2 additions & 4 deletions cpp/tests/groupby/group_mean_test.cpp
Original file line number Diff line number Diff line change
@@ -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;
6 changes: 2 additions & 4 deletions cpp/tests/groupby/group_std_test.cpp
Original file line number Diff line number Diff line change
@@ -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;
4 changes: 1 addition & 3 deletions cpp/tests/groupby/group_sum_of_squares_test.cpp
Original file line number Diff line number Diff line change
@@ -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;
6 changes: 2 additions & 4 deletions cpp/tests/groupby/group_sum_test.cpp
Original file line number Diff line number Diff line change
@@ -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;
6 changes: 2 additions & 4 deletions cpp/tests/groupby/group_var_test.cpp
Original file line number Diff line number Diff line change
@@ -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;

0 comments on commit d08e041

Please sign in to comment.