Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix null handling for structs min and arg_min in groupby, groupby scan, reduction, and inclusive_scan #9864

Merged
merged 14 commits into from
Dec 16, 2021
Merged

Fix null handling for structs min and arg_min in groupby, groupby scan, reduction, and inclusive_scan #9864

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
2705661
Fix reduction, groupby and groupby scan
ttnghia Dec 7, 2021
d7b6b76
Fix aggregation type bug
ttnghia Dec 7, 2021
3a7de44
Fix comment
ttnghia Dec 7, 2021
3e56124
Fix inclusive_scan
ttnghia Dec 7, 2021
f4a235a
Extract common code, rewrite doxygen to clarify the purpose
ttnghia Dec 8, 2021
b6ca544
Further extract common code
ttnghia Dec 9, 2021
aa03b57
Merge branch 'branch-22.02' into fix_null_order_in_structs_min_max
ttnghia Dec 9, 2021
b1b3000
Fix typo
ttnghia Dec 15, 2021
831d60d
Remove redundant variable
ttnghia Dec 15, 2021
956788b
Reimplement comparison_binop_generator, producing `is_min_op` internally
ttnghia Dec 15, 2021
f4c36ee
Merge branch 'branch-22.02' into fix_null_order_in_structs_min_max
ttnghia Dec 15, 2021
8f35bfd
Merge branch 'branch-22.02' into fix_null_order_in_structs_min_max
ttnghia Dec 15, 2021
d910868
Change comments to clarify null at different levels
ttnghia Dec 15, 2021
51c14a6
Merge branch 'branch-22.02' into fix_null_order_in_structs_min_max
ttnghia Dec 15, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
54 changes: 14 additions & 40 deletions cpp/src/groupby/sort/group_scan_util.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@

#pragma once

#include <reductions/arg_minmax_util.cuh>
#include <reductions/struct_minmax_util.cuh>

#include <cudf/column/column.hpp>
#include <cudf/column/column_factories.hpp>
Expand All @@ -26,8 +26,6 @@
#include <cudf/detail/gather.hpp>
#include <cudf/detail/iterator.cuh>
#include <cudf/detail/null_mask.hpp>
#include <cudf/detail/structs/utilities.hpp>
#include <cudf/detail/utilities/vector_factories.hpp>
#include <cudf/table/table_device_view.cuh>
#include <cudf/types.hpp>
#include <cudf/utilities/span.hpp>
Expand Down Expand Up @@ -192,43 +190,19 @@ struct group_scan_functor<K,
{
if (values.is_empty()) { return cudf::empty_like(values); }

// When finding MIN, we need to consider nulls as larger than non-null elements.
// Thing is opposite when finding MAX.
auto const null_precedence = (K == aggregation::MIN) ? null_order::AFTER : null_order::BEFORE;
auto const flattened_values = structs::detail::flatten_nested_columns(
table_view{{values}}, {}, std::vector<null_order>{null_precedence});
auto const d_flattened_values_ptr = table_device_view::create(flattened_values, stream);
auto const flattened_null_precedences =
(K == aggregation::MIN)
? cudf::detail::make_device_uvector_async(flattened_values.null_orders(), stream)
: rmm::device_uvector<null_order>(0, stream);
auto constexpr is_min_op = K == aggregation::MIN;
auto const binop_generator =
cudf::reduction::detail::comparison_binop_generator(values, is_min_op, stream);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be a little cleaner to pass the aggregation to the comparison_binop_generator and have it do the is_min_op logic internally. That way the user doesn't have to type it out every time.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done.


// Create a gather map contaning indices of the prefix min/max elements.
auto gather_map = rmm::device_uvector<size_type>(values.size(), stream);
auto const map_begin = gather_map.begin();

// Perform segmented scan.
auto const do_scan = [&](auto const& inp_iter, auto const& out_iter, auto const& binop) {
thrust::inclusive_scan_by_key(rmm::exec_policy(stream),
group_labels.begin(),
group_labels.end(),
inp_iter,
out_iter,
thrust::equal_to{},
binop);
};

// Find the indices of the prefix min/max elements within each group.
auto const count_iter = thrust::make_counting_iterator<size_type>(0);
auto const binop = cudf::reduction::detail::row_arg_minmax_fn(values.size(),
*d_flattened_values_ptr,
values.has_nulls(),
flattened_null_precedences.data(),
K == aggregation::MIN);
do_scan(count_iter, map_begin, binop);

auto gather_map_view =
column_view(data_type{type_to_id<offset_type>()}, gather_map.size(), gather_map.data());
// Create a gather map contaning indices of the prefix min/max elements within each group.
ttnghia marked this conversation as resolved.
Show resolved Hide resolved
auto gather_map = rmm::device_uvector<size_type>(values.size(), stream);
thrust::inclusive_scan_by_key(rmm::exec_policy(stream),
group_labels.begin(),
group_labels.end(),
thrust::make_counting_iterator<size_type>(0),
gather_map.begin(),
thrust::equal_to{},
binop_generator.binop());

//
// Gather the children elements of the prefix min/max struct elements first.
Expand All @@ -240,7 +214,7 @@ struct group_scan_functor<K,
auto scanned_children =
cudf::detail::gather(
table_view(std::vector<column_view>{values.child_begin(), values.child_end()}),
gather_map_view,
gather_map,
cudf::out_of_bounds_policy::DONT_CHECK,
cudf::detail::negative_index_policy::NOT_ALLOWED,
stream,
Expand Down
25 changes: 5 additions & 20 deletions cpp/src/groupby/sort/group_single_pass_reduction_util.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,15 +16,13 @@

#pragma once

#include <reductions/arg_minmax_util.cuh>
#include <reductions/struct_minmax_util.cuh>

#include <cudf/column/column.hpp>
#include <cudf/column/column_factories.hpp>
#include <cudf/column/column_view.hpp>
#include <cudf/detail/aggregation/aggregation.cuh>
#include <cudf/detail/iterator.cuh>
#include <cudf/detail/structs/utilities.hpp>
#include <cudf/detail/utilities/vector_factories.hpp>
#include <cudf/detail/valid_if.cuh>
#include <cudf/table/row_operators.cuh>
#include <cudf/types.hpp>
Expand Down Expand Up @@ -244,17 +242,9 @@ struct group_reduction_functor<

if (values.is_empty()) { return result; }

// When finding ARGMIN, we need to consider nulls as larger than non-null elements.
// Thing is opposite for ARGMAX.
auto const null_precedence =
(K == aggregation::ARGMIN) ? null_order::AFTER : null_order::BEFORE;
auto const flattened_values = structs::detail::flatten_nested_columns(
table_view{{values}}, {}, std::vector<null_order>{null_precedence});
auto const d_flattened_values_ptr = table_device_view::create(flattened_values, stream);
auto const flattened_null_precedences =
(K == aggregation::ARGMIN)
? cudf::detail::make_device_uvector_async(flattened_values.null_orders(), stream)
: rmm::device_uvector<null_order>(0, stream);
auto constexpr is_min_op = K == aggregation::ARGMIN;
auto const binop_generator =
cudf::reduction::detail::comparison_binop_generator(values, is_min_op, stream);

// Perform segmented reduction to find ARGMIN/ARGMAX.
auto const do_reduction = [&](auto const& inp_iter, auto const& out_iter, auto const& binop) {
Expand All @@ -270,12 +260,7 @@ struct group_reduction_functor<

auto const count_iter = thrust::make_counting_iterator<ResultType>(0);
auto const result_begin = result->mutable_view().template begin<ResultType>();
auto const binop = cudf::reduction::detail::row_arg_minmax_fn(values.size(),
*d_flattened_values_ptr,
values.has_nulls(),
flattened_null_precedences.data(),
K == aggregation::ARGMIN);
do_reduction(count_iter, result_begin, binop);
do_reduction(count_iter, result_begin, binop_generator.binop());

if (values.has_nulls()) {
// Generate bitmask for the output by segmented reduction of the input bitmask.
Expand Down
65 changes: 0 additions & 65 deletions cpp/src/reductions/arg_minmax_util.cuh
View file
Open in desktop

This file was deleted.

42 changes: 11 additions & 31 deletions cpp/src/reductions/scan/scan_inclusive.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -14,17 +14,15 @@
* limitations under the License.
*/

#include <reductions/arg_minmax_util.cuh>
#include <reductions/scan/scan.cuh>
#include <reductions/struct_minmax_util.cuh>

#include <cudf/column/column_device_view.cuh>
#include <cudf/column/column_factories.hpp>
#include <cudf/detail/copy.hpp>
#include <cudf/detail/gather.hpp>
#include <cudf/detail/iterator.cuh>
#include <cudf/detail/null_mask.hpp>
#include <cudf/detail/structs/utilities.hpp>
#include <cudf/detail/utilities/vector_factories.hpp>
#include <cudf/reduction.hpp>

#include <rmm/cuda_stream_view.hpp>
Expand Down Expand Up @@ -159,35 +157,17 @@ struct scan_functor<Op, cudf::struct_view> {
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
// Op is used only to determined if we want to find the min or max element.
auto constexpr is_min_op = std::is_same_v<Op, DeviceMin>;

// Build indices of the scan operation results (ARGMIN/ARGMAX).
// When finding ARGMIN, we need to consider nulls as larger than non-null elements, and the
// opposite for ARGMAX.
auto gather_map = rmm::device_uvector<size_type>(input.size(), stream);
auto const do_scan = [&](auto const& binop) {
thrust::inclusive_scan(rmm::exec_policy(stream),
thrust::counting_iterator<size_type>(0),
thrust::counting_iterator<size_type>(input.size()),
gather_map.begin(),
binop);
};

auto constexpr null_precedence = is_min_op ? cudf::null_order::AFTER : cudf::null_order::BEFORE;
auto const flattened_input = cudf::structs::detail::flatten_nested_columns(
table_view{{input}}, {}, std::vector<null_order>{null_precedence});
auto const d_flattened_input_ptr = table_device_view::create(flattened_input, stream);
auto const flattened_null_precedences =
is_min_op ? cudf::detail::make_device_uvector_async(flattened_input.null_orders(), stream)
: rmm::device_uvector<cudf::null_order>(0, stream);

auto const binop = cudf::reduction::detail::row_arg_minmax_fn(input.size(),
*d_flattened_input_ptr,
input.has_nulls(),
flattened_null_precedences.data(),
is_min_op);
do_scan(binop);
auto const binop_generator =
cudf::reduction::detail::comparison_binop_generator(input, is_min_op, stream);

// Create a gather map contaning indices of the prefix min/max elements.
auto gather_map = rmm::device_uvector<size_type>(input.size(), stream);
thrust::inclusive_scan(rmm::exec_policy(stream),
thrust::counting_iterator<size_type>(0),
thrust::counting_iterator<size_type>(input.size()),
gather_map.begin(),
binop_generator.binop());

// Gather the children columns of the input column. Must use `get_sliced_child` to properly
// handle input in case it is a sliced view.
Expand Down
38 changes: 8 additions & 30 deletions cpp/src/reductions/simple.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,13 +16,12 @@

#pragma once

#include <reductions/arg_minmax_util.cuh>
#include <reductions/struct_minmax_util.cuh>

#include <cudf/detail/copy.hpp>
#include <cudf/detail/reduction.cuh>
#include <cudf/detail/structs/utilities.hpp>
#include <cudf/detail/utilities/cuda.cuh>
#include <cudf/detail/utilities/vector_factories.hpp>
#include <cudf/dictionary/detail/iterator.cuh>
#include <cudf/dictionary/dictionary_column_view.hpp>
#include <cudf/scalar/scalar_device_view.cuh>
Expand Down Expand Up @@ -295,36 +294,15 @@ struct same_element_type_dispatcher {
if (input.is_empty()) { return cudf::make_empty_scalar_like(input, stream, mr); }

auto constexpr is_min_op = std::is_same_v<Op, cudf::reduction::op::min>;
auto const binop_generator =
cudf::reduction::detail::comparison_binop_generator(input, is_min_op, stream);

// We will do reduction to find the ARGMIN/ARGMAX index, then return the element at that index.
// When finding ARGMIN, we need to consider nulls as larger than non-null elements, and the
// opposite for ARGMAX.
auto constexpr null_precedence = is_min_op ? cudf::null_order::AFTER : cudf::null_order::BEFORE;
auto const flattened_input = cudf::structs::detail::flatten_nested_columns(
table_view{{input}}, {}, std::vector<null_order>{null_precedence});
auto const d_flattened_input_ptr = table_device_view::create(flattened_input, stream);
auto const flattened_null_precedences =
is_min_op ? cudf::detail::make_device_uvector_async(flattened_input.null_orders(), stream)
: rmm::device_uvector<cudf::null_order>(0, stream);

// Perform reduction to find ARGMIN/ARGMAX.
auto const do_reduction = [&](auto const& binop) {
return thrust::reduce(rmm::exec_policy(stream),
thrust::make_counting_iterator(0),
thrust::make_counting_iterator(input.size()),
size_type{0},
binop);
};

auto const minmax_idx = [&] {
auto const binop =
cudf::reduction::detail::row_arg_minmax_fn(input.size(),
*d_flattened_input_ptr,
input.has_nulls(),
flattened_null_precedences.data(),
is_min_op);
return do_reduction(binop);
}();
auto const minmax_idx = thrust::reduce(rmm::exec_policy(stream),
thrust::make_counting_iterator(0),
thrust::make_counting_iterator(input.size()),
size_type{0},
binop_generator.binop());

return cudf::detail::get_element(input, minmax_idx, stream, mr);
}
Expand Down
Loading