diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index 193ac1f006c..fa1e61e26fd 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -157,7 +157,7 @@ ConfigureBench(FILL_BENCH filling/repeat_benchmark.cpp) # * groupby benchmark ----------------------------------------------------------------------------- ConfigureBench( GROUPBY_BENCH groupby/group_sum_benchmark.cu groupby/group_nth_benchmark.cu - groupby/group_shift_benchmark.cu + groupby/group_shift_benchmark.cu groupby/group_struct_benchmark.cu ) # ################################################################################################## diff --git a/cpp/benchmarks/groupby/group_struct_benchmark.cu b/cpp/benchmarks/groupby/group_struct_benchmark.cu new file mode 100644 index 00000000000..702983a63bf --- /dev/null +++ b/cpp/benchmarks/groupby/group_struct_benchmark.cu @@ -0,0 +1,107 @@ +/* + * Copyright (c) 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. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include + +#include + +#include +#include +#include +#include +#include + +#include + +static constexpr cudf::size_type num_struct_members = 8; +static constexpr cudf::size_type max_int = 100; +static constexpr cudf::size_type max_str_length = 32; + +static auto create_data_table(cudf::size_type n_rows) +{ + data_profile table_profile; + table_profile.set_distribution_params(cudf::type_id::INT32, distribution_id::UNIFORM, 0, max_int); + table_profile.set_distribution_params( + cudf::type_id::STRING, distribution_id::NORMAL, 0, max_str_length); + + // The first two struct members are int32 and string. + // The first column is also used as keys in groupby. + auto col_ids = std::vector{cudf::type_id::INT32, cudf::type_id::STRING}; + + // The subsequent struct members are int32 and string again. + for (cudf::size_type i = 3; i <= num_struct_members; ++i) { + if (i % 2) { + col_ids.push_back(cudf::type_id::INT32); + } else { + col_ids.push_back(cudf::type_id::STRING); + } + } + + return create_random_table(col_ids, num_struct_members, row_count{n_rows}, table_profile); +} + +// Max aggregation/scan technically has the same performance as min. +template +void BM_groupby_min_struct(benchmark::State& state) +{ + auto const n_rows = static_cast(state.range(0)); + auto data_cols = create_data_table(n_rows)->release(); + + auto const keys_view = data_cols.front()->view(); + auto const values = + cudf::make_structs_column(keys_view.size(), std::move(data_cols), 0, rmm::device_buffer()); + + using RequestType = std::conditional_t, + cudf::groupby::aggregation_request, + cudf::groupby::scan_request>; + + auto gb_obj = cudf::groupby::groupby(cudf::table_view({keys_view})); + auto requests = std::vector(); + requests.emplace_back(RequestType()); + requests.front().values = values->view(); + requests.front().aggregations.push_back(cudf::make_min_aggregation()); + + for (auto _ : state) { + [[maybe_unused]] auto const timer = cuda_event_timer(state, true); + if constexpr (std::is_same_v) { + [[maybe_unused]] auto const result = gb_obj.aggregate(requests); + } else { + [[maybe_unused]] auto const result = gb_obj.scan(requests); + } + } +} + +class Groupby : public cudf::benchmark { +}; + +#define MIN_RANGE 10'000 +#define MAX_RANGE 10'000'000 + +#define REGISTER_BENCHMARK(name, op_type) \ + BENCHMARK_DEFINE_F(Groupby, name)(::benchmark::State & state) \ + { \ + BM_groupby_min_struct(state); \ + } \ + BENCHMARK_REGISTER_F(Groupby, name) \ + ->UseManualTime() \ + ->Unit(benchmark::kMillisecond) \ + ->RangeMultiplier(4) \ + ->Ranges({{MIN_RANGE, MAX_RANGE}}); + +REGISTER_BENCHMARK(Aggregation, cudf::groupby_aggregation) +REGISTER_BENCHMARK(Scan, cudf::groupby_scan_aggregation) diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu index e7024c80a68..ef640256927 100644 --- a/cpp/src/groupby/hash/groupby.cu +++ b/cpp/src/groupby/hash/groupby.cu @@ -632,11 +632,22 @@ std::unique_ptr groupby_null_templated(table_view const& keys, */ bool can_use_hash_groupby(table_view const& keys, host_span requests) { - return std::all_of(requests.begin(), requests.end(), [](aggregation_request const& r) { - return std::all_of(r.aggregations.begin(), r.aggregations.end(), [](auto const& a) { - return is_hash_aggregation(a->kind); + auto const all_hash_aggregations = + std::all_of(requests.begin(), requests.end(), [](aggregation_request const& r) { + return std::all_of(r.aggregations.begin(), r.aggregations.end(), [](auto const& a) { + return is_hash_aggregation(a->kind); + }); }); - }); + + // Currently, structs are not supported in any of hash-based aggregations. + // Therefore, if any request contains structs then we must fallback to sort-based aggregations. + // TODO: Support structs in hash-based aggregations. + auto const has_struct = + std::all_of(requests.begin(), requests.end(), [](aggregation_request const& r) { + return r.values.type().id() == type_id::STRUCT; + }); + + return all_hash_aggregations && !has_struct; } // Hash-based groupby diff --git a/cpp/src/groupby/sort/group_argmax.cu b/cpp/src/groupby/sort/group_argmax.cu index 6ce23ffc35b..466171ec80b 100644 --- a/cpp/src/groupby/sort/group_argmax.cu +++ b/cpp/src/groupby/sort/group_argmax.cu @@ -34,7 +34,7 @@ std::unique_ptr group_argmax(column_view const& values, rmm::mr::device_memory_resource* mr) { auto indices = type_dispatcher(values.type(), - reduce_functor{}, + group_reduction_dispatcher{}, values, num_groups, group_labels, diff --git a/cpp/src/groupby/sort/group_argmin.cu b/cpp/src/groupby/sort/group_argmin.cu index ab91c2c0d29..4f7b2b713e6 100644 --- a/cpp/src/groupby/sort/group_argmin.cu +++ b/cpp/src/groupby/sort/group_argmin.cu @@ -34,7 +34,7 @@ std::unique_ptr group_argmin(column_view const& values, rmm::mr::device_memory_resource* mr) { auto indices = type_dispatcher(values.type(), - reduce_functor{}, + group_reduction_dispatcher{}, values, num_groups, group_labels, diff --git a/cpp/src/groupby/sort/group_max.cu b/cpp/src/groupby/sort/group_max.cu index 7dd0e43ad28..5da15266233 100644 --- a/cpp/src/groupby/sort/group_max.cu +++ b/cpp/src/groupby/sort/group_max.cu @@ -30,8 +30,13 @@ std::unique_ptr group_max(column_view const& values, auto values_type = cudf::is_dictionary(values.type()) ? dictionary_column_view(values).keys().type() : values.type(); - return type_dispatcher( - values_type, reduce_functor{}, values, num_groups, group_labels, stream, mr); + return type_dispatcher(values_type, + group_reduction_dispatcher{}, + values, + num_groups, + group_labels, + stream, + mr); } } // namespace detail diff --git a/cpp/src/groupby/sort/group_max_scan.cu b/cpp/src/groupby/sort/group_max_scan.cu index 303d606be9d..1551dc00a04 100644 --- a/cpp/src/groupby/sort/group_max_scan.cu +++ b/cpp/src/groupby/sort/group_max_scan.cu @@ -27,8 +27,13 @@ std::unique_ptr max_scan(column_view const& values, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return type_dispatcher( - values.type(), scan_functor{}, values, num_groups, group_labels, stream, mr); + return type_dispatcher(values.type(), + group_scan_dispatcher{}, + values, + num_groups, + group_labels, + stream, + mr); } } // namespace detail diff --git a/cpp/src/groupby/sort/group_min.cu b/cpp/src/groupby/sort/group_min.cu index 4124ec0f6f6..c42a0b94de0 100644 --- a/cpp/src/groupby/sort/group_min.cu +++ b/cpp/src/groupby/sort/group_min.cu @@ -30,8 +30,13 @@ std::unique_ptr group_min(column_view const& values, auto values_type = cudf::is_dictionary(values.type()) ? dictionary_column_view(values).keys().type() : values.type(); - return type_dispatcher( - values_type, reduce_functor{}, values, num_groups, group_labels, stream, mr); + return type_dispatcher(values_type, + group_reduction_dispatcher{}, + values, + num_groups, + group_labels, + stream, + mr); } } // namespace detail diff --git a/cpp/src/groupby/sort/group_min_scan.cu b/cpp/src/groupby/sort/group_min_scan.cu index 4a692cdf0bd..daaeb6bb6f7 100644 --- a/cpp/src/groupby/sort/group_min_scan.cu +++ b/cpp/src/groupby/sort/group_min_scan.cu @@ -27,8 +27,13 @@ std::unique_ptr min_scan(column_view const& values, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return type_dispatcher( - values.type(), scan_functor{}, values, num_groups, group_labels, stream, mr); + return type_dispatcher(values.type(), + group_scan_dispatcher{}, + values, + num_groups, + group_labels, + stream, + mr); } } // namespace detail diff --git a/cpp/src/groupby/sort/group_product.cu b/cpp/src/groupby/sort/group_product.cu index e9cf8611b58..74f5cbed041 100644 --- a/cpp/src/groupby/sort/group_product.cu +++ b/cpp/src/groupby/sort/group_product.cu @@ -33,7 +33,7 @@ std::unique_ptr group_product(column_view const& values, ? dictionary_column_view(values).keys().type() : values.type(); return type_dispatcher(values_type, - reduce_functor{}, + group_reduction_dispatcher{}, values, num_groups, group_labels, diff --git a/cpp/src/groupby/sort/group_scan_util.cuh b/cpp/src/groupby/sort/group_scan_util.cuh index ef9df937fc5..013ea924cce 100644 --- a/cpp/src/groupby/sort/group_scan_util.cuh +++ b/cpp/src/groupby/sort/group_scan_util.cuh @@ -16,12 +16,18 @@ #pragma once +#include + #include #include #include +#include #include +#include #include #include +#include +#include #include #include #include @@ -36,27 +42,51 @@ namespace cudf { namespace groupby { namespace detail { +// Error case when no other overload or specialization is available +template +struct group_scan_functor { + template + static std::unique_ptr invoke(Args&&...) + { + CUDF_FAIL("Unsupported groupby scan type-agg combination."); + } +}; + template -struct scan_functor { +struct group_scan_dispatcher { template - static constexpr bool is_supported() + std::unique_ptr operator()(column_view const& values, + size_type num_groups, + cudf::device_span group_labels, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { - if (K == aggregation::SUM) - return cudf::is_numeric() || cudf::is_duration() || cudf::is_fixed_point(); - else if (K == aggregation::MIN or K == aggregation::MAX) - return !cudf::is_dictionary() and is_relationally_comparable(); - else - return false; + return group_scan_functor::invoke(values, num_groups, group_labels, stream, mr); } +}; - template - std::enable_if_t() and not std::is_same_v, - std::unique_ptr> - operator()(column_view const& values, - size_type num_groups, - cudf::device_span group_labels, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +/** + * @brief Check if the given aggregation K with data type T is supported in groupby scan. + */ +template +static constexpr bool is_group_scan_supported() +{ + if (K == aggregation::SUM) + return cudf::is_numeric() || cudf::is_duration() || cudf::is_fixed_point(); + else if (K == aggregation::MIN or K == aggregation::MAX) + return not cudf::is_dictionary() and + (is_relationally_comparable() or std::is_same_v); + else + return false; +} + +template +struct group_scan_functor()>> { + static std::unique_ptr invoke(column_view const& values, + size_type num_groups, + cudf::device_span group_labels, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { using DeviceType = device_storage_type_t; using OpType = cudf::detail::corresponding_operator_t; @@ -78,70 +108,68 @@ struct scan_functor { auto result_view = mutable_column_device_view::create(result->mutable_view(), stream); auto values_view = column_device_view::create(values, stream); - if (values.has_nulls()) { - auto input = thrust::make_transform_iterator( - make_null_replacement_iterator(*values_view, OpType::template identity()), - thrust::identity{}); + // 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(), - input, - result_view->begin(), + inp_iter, + out_iter, thrust::equal_to{}, - OpType{}); + binop); + }; + + if (values.has_nulls()) { + auto input = thrust::make_transform_iterator( + make_null_replacement_iterator(*values_view, OpType::template identity()), + thrust::identity{}); + do_scan(input, result_view->begin(), OpType{}); result->set_null_mask(cudf::detail::copy_bitmask(values, stream)); } else { auto input = thrust::make_transform_iterator(values_view->begin(), thrust::identity{}); - thrust::inclusive_scan_by_key(rmm::exec_policy(stream), - group_labels.begin(), - group_labels.end(), - input, - result_view->begin(), - thrust::equal_to{}, - OpType{}); + do_scan(input, result_view->begin(), OpType{}); } return result; } +}; - template - std::enable_if_t() and std::is_same_v, - std::unique_ptr> - operator()(column_view const& values, - size_type num_groups, - cudf::device_span group_labels, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) +template +struct group_scan_functor()>> { + static std::unique_ptr invoke(column_view const& values, + size_type num_groups, + cudf::device_span group_labels, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { using OpType = cudf::detail::corresponding_operator_t; - if (values.is_empty()) { - return cudf::make_empty_column(cudf::data_type{cudf::type_id::STRING}); - } + if (values.is_empty()) { return cudf::make_empty_column(cudf::type_id::STRING); } // create an empty output vector we can fill with string_view instances auto results_vector = rmm::device_uvector(values.size(), stream); auto values_view = column_device_view::create(values, stream); - if (values.has_nulls()) { - auto input = make_null_replacement_iterator( - *values_view, OpType::template identity(), values.has_nulls()); + // 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(), - input, - results_vector.begin(), + inp_iter, + out_iter, thrust::equal_to{}, - OpType{}); + binop); + }; + + if (values.has_nulls()) { + auto input = make_null_replacement_iterator( + *values_view, OpType::template identity(), values.has_nulls()); + do_scan(input, results_vector.begin(), OpType{}); } else { - thrust::inclusive_scan_by_key(rmm::exec_policy(stream), - group_labels.begin(), - group_labels.end(), - values_view->begin(), - results_vector.begin(), - thrust::equal_to{}, - OpType{}); + do_scan(values_view->begin(), results_vector.begin(), OpType{}); } // turn the string_view vector into a strings column @@ -150,11 +178,89 @@ struct scan_functor { results->set_null_mask(cudf::detail::copy_bitmask(values, stream), values.null_count()); return results; } +}; - template - std::enable_if_t(), std::unique_ptr> operator()(Args&&... args) +template +struct group_scan_functor()>> { + static std::unique_ptr invoke(column_view const& values, + size_type num_groups, + cudf::device_span group_labels, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { - CUDF_FAIL("Unsupported groupby scan type-agg combination"); + 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_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(0, stream); + + // Create a gather map contaning indices of the prefix min/max elements. + auto gather_map = rmm::device_uvector(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(0); + if (values.has_nulls()) { + auto const binop = row_arg_minmax_fn(values.size(), + *d_flattened_values_ptr, + flattened_null_precedences.data(), + K == aggregation::MIN); + do_scan(count_iter, map_begin, binop); + } else { + auto const binop = row_arg_minmax_fn(values.size(), + *d_flattened_values_ptr, + 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()}, gather_map.size(), gather_map.data()); + + // Gather the children elements of the prefix min/max struct elements first. + auto scanned_children = + cudf::detail::gather( + table_view(std::vector{values.child_begin(), values.child_end()}), + gather_map_view, + cudf::out_of_bounds_policy::DONT_CHECK, + cudf::detail::negative_index_policy::NOT_ALLOWED, + stream, + mr) + ->release(); + + // After gathering the children elements, we need to push down nulls from the root structs + // column to them. + if (values.has_nulls()) { + for (std::unique_ptr& child : scanned_children) { + structs::detail::superimpose_parent_nulls( + values.null_mask(), values.null_count(), *child, stream, mr); + } + } + + return make_structs_column(values.size(), + std::move(scanned_children), + values.null_count(), + cudf::detail::copy_bitmask(values, stream, mr)); } }; diff --git a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh index db2ae5b5d8e..4e0820af236 100644 --- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh +++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh @@ -16,12 +16,17 @@ #pragma once +#include + #include #include #include #include #include +#include +#include #include +#include #include #include @@ -38,40 +43,33 @@ namespace groupby { namespace detail { /** - * @brief ArgMin binary operator with index values into input column. + * @brief Binary operator with index values into the input column. * * @tparam T Type of the underlying column. Must support '<' operator. */ template -struct ArgMin { +struct element_arg_minmax_fn { column_device_view const d_col; - CUDA_DEVICE_CALLABLE auto operator()(size_type const& lhs, size_type const& rhs) const - { - // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and - // github.com/NVIDIA/thrust/issues/1525 - // where invalid random values may be passed here by thrust::reduce_by_key - if (lhs < 0 || lhs >= d_col.size() || d_col.is_null(lhs)) { return rhs; } - if (rhs < 0 || rhs >= d_col.size() || d_col.is_null(rhs)) { return lhs; } - return d_col.element(lhs) < d_col.element(rhs) ? lhs : rhs; - } -}; + bool const has_nulls; + bool const arg_min; -/** - * @brief ArgMax binary operator with index values into input column. - * - * @tparam T Type of the underlying column. Must support '<' operator. - */ -template -struct ArgMax { - column_device_view const d_col; - CUDA_DEVICE_CALLABLE auto operator()(size_type const& lhs, size_type const& rhs) const + CUDA_DEVICE_CALLABLE auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const { // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and // github.com/NVIDIA/thrust/issues/1525 // where invalid random values may be passed here by thrust::reduce_by_key - if (lhs < 0 || lhs >= d_col.size() || d_col.is_null(lhs)) { return rhs; } - if (rhs < 0 || rhs >= d_col.size() || d_col.is_null(rhs)) { return lhs; } - return d_col.element(rhs) < d_col.element(lhs) ? lhs : rhs; + if (lhs_idx < 0 || lhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(lhs_idx))) { + return rhs_idx; + } + if (rhs_idx < 0 || rhs_idx >= d_col.size() || (has_nulls && d_col.is_null_nocheck(rhs_idx))) { + return lhs_idx; + } + + // Return `lhs_idx` iff: + // row(lhs_idx) < row(rhs_idx) and finding ArgMin, or + // row(lhs_idx) >= row(rhs_idx) and finding ArgMax. + auto const less = d_col.element(lhs_idx) < d_col.element(rhs_idx); + return less == arg_min ? lhs_idx : rhs_idx; } }; @@ -121,34 +119,58 @@ struct null_replaced_value_accessor : value_accessor { } }; +// Error case when no other overload or specialization is available +template +struct group_reduction_functor { + template + static std::unique_ptr invoke(Args&&...) + { + CUDF_FAIL("Unsupported groupby reduction type-agg combination."); + } +}; + template -struct reduce_functor { +struct group_reduction_dispatcher { template - static constexpr bool is_supported() + std::unique_ptr operator()(column_view const& values, + size_type num_groups, + cudf::device_span group_labels, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { - switch (K) { - case aggregation::SUM: - return cudf::is_numeric() || cudf::is_duration() || cudf::is_fixed_point(); - case aggregation::PRODUCT: return cudf::detail::is_product_supported(); - case aggregation::MIN: - case aggregation::MAX: - return cudf::is_fixed_width() and is_relationally_comparable(); - case aggregation::ARGMIN: - case aggregation::ARGMAX: return is_relationally_comparable(); - default: return false; - } + return group_reduction_functor::invoke(values, num_groups, group_labels, stream, mr); } +}; + +/** + * @brief Check if the given aggregation K with data type T is supported in groupby reduction. + */ +template +static constexpr bool is_group_reduction_supported() +{ + switch (K) { + case aggregation::SUM: + return cudf::is_numeric() || cudf::is_duration() || cudf::is_fixed_point(); + case aggregation::PRODUCT: return cudf::detail::is_product_supported(); + case aggregation::MIN: + case aggregation::MAX: return cudf::is_fixed_width() and is_relationally_comparable(); + case aggregation::ARGMIN: + case aggregation::ARGMAX: + return is_relationally_comparable() or std::is_same_v; + default: return false; + } +} + +template +struct group_reduction_functor()>> { + static std::unique_ptr invoke(column_view const& values, + size_type num_groups, + cudf::device_span group_labels, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) - template - std::enable_if_t(), std::unique_ptr> operator()( - column_view const& values, - size_type num_groups, - cudf::device_span group_labels, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr) { using DeviceType = device_storage_type_t; - using OpType = cudf::detail::corresponding_operator_t; using ResultType = cudf::detail::target_type_t; using ResultDType = device_storage_type_t; @@ -161,55 +183,119 @@ struct reduce_functor { if (values.is_empty()) { return result; } - auto resultview = mutable_column_device_view::create(result->mutable_view(), stream); - auto valuesview = column_device_view::create(values, stream); - if constexpr (K == aggregation::ARGMAX || K == aggregation::ARGMIN) { - using OpType = - std::conditional_t<(K == aggregation::ARGMAX), ArgMax, ArgMin>; + // Perform segmented reduction. + auto const do_reduction = [&](auto const& inp_iter, auto const& out_iter, auto const& binop) { thrust::reduce_by_key(rmm::exec_policy(stream), group_labels.data(), group_labels.data() + group_labels.size(), - thrust::make_counting_iterator(0), + inp_iter, thrust::make_discard_iterator(), - resultview->begin(), - thrust::equal_to{}, - OpType{*valuesview}); - } else { - auto init = OpType::template identity(); - auto begin = cudf::detail::make_counting_transform_iterator( - 0, null_replaced_value_accessor{*valuesview, init, values.has_nulls()}); - thrust::reduce_by_key(rmm::exec_policy(stream), - group_labels.data(), - group_labels.data() + group_labels.size(), - begin, - thrust::make_discard_iterator(), - resultview->begin(), + out_iter, thrust::equal_to{}, - OpType{}); + binop); + }; + + auto const d_values_ptr = column_device_view::create(values, stream); + auto const result_begin = result->mutable_view().template begin(); + + if constexpr (K == aggregation::ARGMAX || K == aggregation::ARGMIN) { + auto const count_iter = thrust::make_counting_iterator(0); + auto const binop = + element_arg_minmax_fn{*d_values_ptr, values.has_nulls(), K == aggregation::ARGMIN}; + do_reduction(count_iter, result_begin, binop); + } else { + using OpType = cudf::detail::corresponding_operator_t; + auto init = OpType::template identity(); + auto inp_values = cudf::detail::make_counting_transform_iterator( + 0, null_replaced_value_accessor{*d_values_ptr, init, values.has_nulls()}); + do_reduction(inp_values, result_begin, OpType{}); } if (values.has_nulls()) { rmm::device_uvector validity(num_groups, stream); + do_reduction(cudf::detail::make_validity_iterator(*d_values_ptr), + validity.begin(), + thrust::logical_or{}); + + auto [null_mask, null_count] = cudf::detail::valid_if( + validity.begin(), validity.end(), thrust::identity{}, stream, mr); + result->set_null_mask(std::move(null_mask), null_count); + } + return result; + } +}; + +template +struct group_reduction_functor< + K, + cudf::struct_view, + std::enable_if_t()>> { + static std::unique_ptr invoke(column_view const& values, + size_type num_groups, + cudf::device_span group_labels, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) + { + // This is be expected to be size_type. + using ResultType = cudf::detail::target_type_t; + + auto result = make_fixed_width_column( + data_type{type_to_id()}, num_groups, mask_state::UNALLOCATED, stream, mr); + + 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_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(0, stream); + + // Perform segmented reduction to find ARGMIN/ARGMAX. + auto const do_reduction = [&](auto const& inp_iter, auto const& out_iter, auto const& binop) { thrust::reduce_by_key(rmm::exec_policy(stream), group_labels.data(), group_labels.data() + group_labels.size(), - cudf::detail::make_validity_iterator(*valuesview), + inp_iter, thrust::make_discard_iterator(), - validity.begin(), + out_iter, thrust::equal_to{}, - thrust::logical_or{}); + binop); + }; + + auto const count_iter = thrust::make_counting_iterator(0); + auto const result_begin = result->mutable_view().template begin(); + if (values.has_nulls()) { + auto const binop = row_arg_minmax_fn(values.size(), + *d_flattened_values_ptr, + flattened_null_precedences.data(), + K == aggregation::ARGMIN); + do_reduction(count_iter, result_begin, binop); + + // Generate bitmask for the output by segmented reduction of the input bitmask. + auto const d_values_ptr = column_device_view::create(values, stream); + auto validity = rmm::device_uvector(num_groups, stream); + do_reduction(cudf::detail::make_validity_iterator(*d_values_ptr), + validity.begin(), + thrust::logical_or{}); + auto [null_mask, null_count] = cudf::detail::valid_if( validity.begin(), validity.end(), thrust::identity{}, stream, mr); - result->set_null_mask(std::move(null_mask)); - result->set_null_count(null_count); + result->set_null_mask(std::move(null_mask), null_count); + } else { + auto const binop = row_arg_minmax_fn(values.size(), + *d_flattened_values_ptr, + flattened_null_precedences.data(), + K == aggregation::ARGMIN); + do_reduction(count_iter, result_begin, binop); } - return result; - } - template - std::enable_if_t(), std::unique_ptr> operator()(Args&&... args) - { - CUDF_FAIL("Unsupported type-agg combination"); + return result; } }; diff --git a/cpp/src/groupby/sort/group_sum.cu b/cpp/src/groupby/sort/group_sum.cu index e9e6e985c54..e3c2ce7c864 100644 --- a/cpp/src/groupby/sort/group_sum.cu +++ b/cpp/src/groupby/sort/group_sum.cu @@ -32,8 +32,13 @@ std::unique_ptr group_sum(column_view const& values, auto values_type = cudf::is_dictionary(values.type()) ? dictionary_column_view(values).keys().type() : values.type(); - return type_dispatcher( - values_type, reduce_functor{}, values, num_groups, group_labels, stream, mr); + return type_dispatcher(values_type, + group_reduction_dispatcher{}, + values, + num_groups, + group_labels, + stream, + mr); } } // namespace detail diff --git a/cpp/src/groupby/sort/group_sum_scan.cu b/cpp/src/groupby/sort/group_sum_scan.cu index ae9b1c321d4..632fde3b9d5 100644 --- a/cpp/src/groupby/sort/group_sum_scan.cu +++ b/cpp/src/groupby/sort/group_sum_scan.cu @@ -27,8 +27,13 @@ std::unique_ptr sum_scan(column_view const& values, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - return type_dispatcher( - values.type(), scan_functor{}, values, num_groups, group_labels, stream, mr); + return type_dispatcher(values.type(), + group_scan_dispatcher{}, + values, + num_groups, + group_labels, + stream, + mr); } } // namespace detail diff --git a/cpp/src/groupby/sort/group_util.cuh b/cpp/src/groupby/sort/group_util.cuh new file mode 100644 index 00000000000..31ff29ed4c3 --- /dev/null +++ b/cpp/src/groupby/sort/group_util.cuh @@ -0,0 +1,66 @@ +/* + * Copyright (c) 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. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include + +namespace cudf { +namespace groupby { +namespace detail { + +/** + * @brief Binary operator ArgMin/ArgMax with index values into the input table. + * + * @tparam T Type of the underlying data. This is the fallback for the cases when T does not support + * '<' operator. + */ +template +struct row_arg_minmax_fn { + size_type const num_rows; + row_lexicographic_comparator const comp; + bool const arg_min; + + row_arg_minmax_fn(size_type const num_rows_, + table_device_view const& table_, + null_order const* null_precedence_, + bool const arg_min_) + : num_rows(num_rows_), comp(table_, table_, nullptr, null_precedence_), arg_min(arg_min_) + { + } + + // This function is explicitly prevented from inlining, because it calls to + // `row_lexicographic_comparator::operator()` which is inlined and very heavy-weight. As a result, + // instantiating this functor will result in huge code, and objects of this functor used with + // `thrust::reduce_by_key` or `thrust::scan_by_key` will result in significant compile time. + __attribute__((noinline)) __device__ auto operator()(size_type lhs_idx, size_type rhs_idx) const + { + // The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and + // github.com/NVIDIA/thrust/issues/1525 + // where invalid random values may be passed here by thrust::reduce_by_key + if (lhs_idx < 0 || lhs_idx >= num_rows) { return rhs_idx; } + if (rhs_idx < 0 || rhs_idx >= num_rows) { return lhs_idx; } + + // Return `lhs_idx` iff: + // row(lhs_idx) < row(rhs_idx) and finding ArgMin, or + // row(lhs_idx) >= row(rhs_idx) and finding ArgMax. + return comp(lhs_idx, rhs_idx) == arg_min ? lhs_idx : rhs_idx; + } +}; + +} // namespace detail +} // namespace groupby +} // namespace cudf diff --git a/cpp/tests/groupby/argmax_tests.cpp b/cpp/tests/groupby/argmax_tests.cpp index 7cf693f7b08..0b06c184b75 100644 --- a/cpp/tests/groupby/argmax_tests.cpp +++ b/cpp/tests/groupby/argmax_tests.cpp @@ -32,7 +32,7 @@ struct groupby_argmax_test : public cudf::test::BaseFixture { }; using K = int32_t; -TYPED_TEST_CASE(groupby_argmax_test, cudf::test::FixedWidthTypes); +TYPED_TEST_SUITE(groupby_argmax_test, cudf::test::FixedWidthTypes); TYPED_TEST(groupby_argmax_test, basic) { @@ -182,6 +182,78 @@ TEST_F(groupby_dictionary_argmax_test, basic) force_use_sort_impl::YES); } +struct groupby_argmax_struct_test : public cudf::test::BaseFixture { +}; + +TEST_F(groupby_argmax_struct_test, basic) +{ + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = [] { + auto child1 = + strings_column_wrapper{"año", "bit", "₹1", "aaa", "zit", "bat", "aab", "$1", "€1", "wut"}; + auto child2 = fixed_width_column_wrapper{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_indices = fixed_width_column_wrapper{0, 4, 2}; + + auto agg = cudf::make_argmax_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_indices, std::move(agg)); +} + +TEST_F(groupby_argmax_struct_test, slice_input) +{ + constexpr int32_t dont_care{1}; + auto const keys_original = fixed_width_column_wrapper{ + dont_care, dont_care, 1, 2, 3, 1, 2, 2, 1, 3, 3, 2, dont_care}; + auto const vals_original = [] { + auto child1 = strings_column_wrapper{"dont_care", + "dont_care", + "año", + "bit", + "₹1", + "aaa", + "zit", + "bat", + "aab", + "$1", + "€1", + "wut", + "dont_care"}; + auto child2 = fixed_width_column_wrapper{ + dont_care, dont_care, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, dont_care}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const keys = cudf::slice(keys_original, {2, 12})[0]; + auto const vals = cudf::slice(vals_original, {2, 12})[0]; + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_indices = fixed_width_column_wrapper{0, 4, 2}; + + auto agg = cudf::make_argmax_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_indices, std::move(agg)); +} + +TEST_F(groupby_argmax_struct_test, null_keys_and_values) +{ + constexpr int32_t null{0}; + auto const keys = + fixed_width_column_wrapper{{1, 2, 3, 1, 2, 2, 1, null, 3, 2, 4}, null_at(7)}; + auto const vals = [] { + auto child1 = strings_column_wrapper{ + "año", "bit", "₹1", "aaa", "zit", "" /*NULL*/, "" /*NULL*/, "$1", "€1", "wut", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 8, 7, 6, 5, null, null, 2, 1, 0, null}; + return structs_column_wrapper{{child1, child2}, nulls_at({5, 6, 10})}; + }(); + + auto const expect_keys = fixed_width_column_wrapper{{1, 2, 3, 4}, no_nulls()}; + auto const expect_indices = fixed_width_column_wrapper{{0, 4, 2, null}, null_at(3)}; + + auto agg = cudf::make_argmax_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_indices, std::move(agg)); +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/argmin_tests.cpp b/cpp/tests/groupby/argmin_tests.cpp index 915575546c9..67235a64066 100644 --- a/cpp/tests/groupby/argmin_tests.cpp +++ b/cpp/tests/groupby/argmin_tests.cpp @@ -32,7 +32,7 @@ struct groupby_argmin_test : public cudf::test::BaseFixture { }; using K = int32_t; -TYPED_TEST_CASE(groupby_argmin_test, cudf::test::FixedWidthTypes); +TYPED_TEST_SUITE(groupby_argmin_test, cudf::test::FixedWidthTypes); TYPED_TEST(groupby_argmin_test, basic) { @@ -183,5 +183,77 @@ TEST_F(groupby_dictionary_argmin_test, basic) force_use_sort_impl::YES); } +struct groupby_argmin_struct_test : public cudf::test::BaseFixture { +}; + +TEST_F(groupby_argmin_struct_test, basic) +{ + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = [] { + auto child1 = + strings_column_wrapper{"año", "bit", "₹1", "aaa", "zit", "bat", "aab", "$1", "€1", "wut"}; + auto child2 = fixed_width_column_wrapper{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_indices = fixed_width_column_wrapper{3, 5, 7}; + + auto agg = cudf::make_argmin_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_indices, std::move(agg)); +} + +TEST_F(groupby_argmin_struct_test, slice_input) +{ + constexpr int32_t dont_care{1}; + auto const keys_original = fixed_width_column_wrapper{ + dont_care, dont_care, 1, 2, 3, 1, 2, 2, 1, 3, 3, 2, dont_care}; + auto const vals_original = [] { + auto child1 = strings_column_wrapper{"dont_care", + "dont_care", + "año", + "bit", + "₹1", + "aaa", + "zit", + "bat", + "aab", + "$1", + "€1", + "wut", + "dont_care"}; + auto child2 = fixed_width_column_wrapper{ + dont_care, dont_care, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, dont_care}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const keys = cudf::slice(keys_original, {2, 12})[0]; + auto const vals = cudf::slice(vals_original, {2, 12})[0]; + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_indices = fixed_width_column_wrapper{3, 5, 7}; + + auto agg = cudf::make_argmin_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_indices, std::move(agg)); +} + +TEST_F(groupby_argmin_struct_test, null_keys_and_values) +{ + constexpr int32_t null{0}; + auto const keys = + fixed_width_column_wrapper{{1, 2, 3, 1, 2, 2, 1, null, 3, 2, 4}, null_at(7)}; + auto const vals = [] { + auto child1 = strings_column_wrapper{ + "año", "bit", "₹1", "aaa", "zit", "" /*NULL*/, "" /*NULL*/, "$1", "€1", "wut", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 8, 7, 6, 5, null, null, 2, 1, 0, null}; + return structs_column_wrapper{{child1, child2}, nulls_at({5, 6, 10})}; + }(); + + auto const expect_keys = fixed_width_column_wrapper{{1, 2, 3, 4}, no_nulls()}; + auto const expect_indices = fixed_width_column_wrapper{{3, 1, 8, null}, null_at(3)}; + + auto agg = cudf::make_argmin_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_indices, std::move(agg)); +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/max_scan_tests.cpp b/cpp/tests/groupby/max_scan_tests.cpp index 5108f91b31c..bb2f87fd424 100644 --- a/cpp/tests/groupby/max_scan_tests.cpp +++ b/cpp/tests/groupby/max_scan_tests.cpp @@ -39,7 +39,7 @@ struct groupby_max_scan_test : public cudf::test::BaseFixture { using result_wrapper = fixed_width_column_wrapper; }; -TYPED_TEST_CASE(groupby_max_scan_test, cudf::test::FixedWidthTypesWithoutFixedPoint); +TYPED_TEST_SUITE(groupby_max_scan_test, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(groupby_max_scan_test, basic) { @@ -148,7 +148,7 @@ template struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); +TYPED_TEST_SUITE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, GroupBySortMaxScanDecimalAsValue) { @@ -173,5 +173,90 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySortMaxScanDecimalAsValue) } } +struct groupby_max_scan_struct_test : public cudf::test::BaseFixture { +}; + +TEST_F(groupby_max_scan_struct_test, basic) +{ + auto const keys = key_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = [] { + auto child1 = + strings_column_wrapper{"año", "bit", "₹1", "aaa", "zit", "bat", "aab", "$1", "€1", "wut"}; + auto child2 = fixed_width_column_wrapper{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const expect_keys = key_wrapper{1, 1, 1, 2, 2, 2, 2, 3, 3, 3}; + auto const expect_vals = [] { + auto child1 = + strings_column_wrapper{"año", "año", "año", "bit", "zit", "zit", "zit", "₹1", "₹1", "₹1"}; + auto child2 = fixed_width_column_wrapper{1, 1, 1, 2, 5, 5, 5, 3, 3, 3}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto agg = cudf::make_max_aggregation(); + test_single_scan(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + +TEST_F(groupby_max_scan_struct_test, slice_input) +{ + constexpr int32_t dont_care{1}; + auto const keys_original = + key_wrapper{dont_care, dont_care, 1, 2, 3, 1, 2, 2, 1, 3, 3, 2, dont_care}; + auto const vals_original = [] { + auto child1 = strings_column_wrapper{"dont_care", + "dont_care", + "año", + "bit", + "₹1", + "aaa", + "zit", + "bat", + "aab", + "$1", + "€1", + "wut", + "dont_care"}; + auto child2 = key_wrapper{dont_care, dont_care, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, dont_care}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const keys = cudf::slice(keys_original, {2, 12})[0]; + auto const vals = cudf::slice(vals_original, {2, 12})[0]; + auto const expect_keys = key_wrapper{1, 1, 1, 2, 2, 2, 2, 3, 3, 3}; + auto const expect_vals = [] { + auto child1 = + strings_column_wrapper{"año", "año", "año", "bit", "zit", "zit", "zit", "₹1", "₹1", "₹1"}; + auto child2 = fixed_width_column_wrapper{1, 1, 1, 2, 5, 5, 5, 3, 3, 3}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto agg = cudf::make_max_aggregation(); + test_single_scan(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + +TEST_F(groupby_max_scan_struct_test, null_keys_and_values) +{ + constexpr int32_t null{0}; + auto const keys = key_wrapper{{1, 2, 3, 1, 2, 2, 1, null, 3, 2, 4}, null_at(7)}; + auto const vals = [] { + auto child1 = strings_column_wrapper{ + "año", "bit", "₹1", "aaa", "zit", "" /*NULL*/, "" /*NULL*/, "$1", "€1", "wut", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 8, 7, 6, 5, null, null, 2, 1, 0, null}; + return structs_column_wrapper{{child1, child2}, nulls_at({5, 6, 10})}; + }(); + + auto const expect_keys = key_wrapper{{1, 1, 1, 2, 2, 2, 2, 3, 3, 4}, no_nulls()}; + auto const expect_vals = [] { + auto child1 = strings_column_wrapper{ + "año", "año", "" /*NULL*/, "bit", "zit", "" /*NULL*/, "zit", "₹1", "₹1", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 9, null, 8, 5, null, 5, 7, 7, null}; + return structs_column_wrapper{{child1, child2}, nulls_at({2, 5, 9})}; + }(); + + auto agg = cudf::make_max_aggregation(); + test_single_scan(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/max_tests.cpp b/cpp/tests/groupby/max_tests.cpp index 491e6927304..8d15401aa09 100644 --- a/cpp/tests/groupby/max_tests.cpp +++ b/cpp/tests/groupby/max_tests.cpp @@ -33,7 +33,7 @@ struct groupby_max_test : public cudf::test::BaseFixture { }; using K = int32_t; -TYPED_TEST_CASE(groupby_max_test, cudf::test::FixedWidthTypesWithoutFixedPoint); +TYPED_TEST_SUITE(groupby_max_test, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(groupby_max_test, basic) { @@ -255,7 +255,7 @@ template struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); +TYPED_TEST_SUITE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, GroupBySortMaxDecimalAsValue) { @@ -304,5 +304,89 @@ TYPED_TEST(FixedPointTestBothReps, GroupByHashMaxDecimalAsValue) } } +struct groupby_max_struct_test : public cudf::test::BaseFixture { +}; + +TEST_F(groupby_max_struct_test, basic) +{ + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = [] { + auto child1 = + strings_column_wrapper{"año", "bit", "₹1", "aaa", "zit", "bat", "aab", "$1", "€1", "wut"}; + auto child2 = fixed_width_column_wrapper{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals = [] { + auto child1 = strings_column_wrapper{"año", "zit", "₹1"}; + auto child2 = fixed_width_column_wrapper{1, 5, 3}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto agg = cudf::make_max_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + +TEST_F(groupby_max_struct_test, slice_input) +{ + constexpr int32_t dont_care{1}; + auto const keys_original = fixed_width_column_wrapper{ + dont_care, dont_care, 1, 2, 3, 1, 2, 2, 1, 3, 3, 2, dont_care}; + auto const vals_original = [] { + auto child1 = strings_column_wrapper{"dont_care", + "dont_care", + "año", + "bit", + "₹1", + "aaa", + "zit", + "bat", + "aab", + "$1", + "€1", + "wut", + "dont_care"}; + auto child2 = fixed_width_column_wrapper{ + dont_care, dont_care, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, dont_care}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const keys = cudf::slice(keys_original, {2, 12})[0]; + auto const vals = cudf::slice(vals_original, {2, 12})[0]; + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals = [] { + auto child1 = strings_column_wrapper{"año", "zit", "₹1"}; + auto child2 = fixed_width_column_wrapper{1, 5, 3}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto agg = cudf::make_max_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + +TEST_F(groupby_max_struct_test, null_keys_and_values) +{ + constexpr int32_t null{0}; + auto const keys = + fixed_width_column_wrapper{{1, 2, 3, 1, 2, 2, 1, null, 3, 2, 4}, null_at(7)}; + auto const vals = [] { + auto child1 = strings_column_wrapper{ + "año", "bit", "₹1", "aaa", "zit", "" /*NULL*/, "" /*NULL*/, "$1", "€1", "wut", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 8, 7, 6, 5, null, null, 2, 1, 0, null}; + return structs_column_wrapper{{child1, child2}, nulls_at({5, 6, 10})}; + }(); + + auto const expect_keys = fixed_width_column_wrapper{{1, 2, 3, 4}, no_nulls()}; + auto const expect_vals = [] { + auto child1 = strings_column_wrapper{"año", "zit", "₹1", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 5, 7, null}; + return structs_column_wrapper{{child1, child2}, null_at(3)}; + }(); + + auto agg = cudf::make_max_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/min_scan_tests.cpp b/cpp/tests/groupby/min_scan_tests.cpp index 59b6382616f..06c0f5ceb3b 100644 --- a/cpp/tests/groupby/min_scan_tests.cpp +++ b/cpp/tests/groupby/min_scan_tests.cpp @@ -38,7 +38,7 @@ struct groupby_min_scan_test : public cudf::test::BaseFixture { using result_wrapper = fixed_width_column_wrapper; }; -TYPED_TEST_CASE(groupby_min_scan_test, cudf::test::FixedWidthTypesWithoutFixedPoint); +TYPED_TEST_SUITE(groupby_min_scan_test, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(groupby_min_scan_test, basic) { @@ -146,7 +146,7 @@ template struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); +TYPED_TEST_SUITE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, GroupBySortMinScanDecimalAsValue) { @@ -172,5 +172,90 @@ TYPED_TEST(FixedPointTestBothReps, GroupBySortMinScanDecimalAsValue) } } +struct groupby_min_scan_struct_test : public cudf::test::BaseFixture { +}; + +TEST_F(groupby_min_scan_struct_test, basic) +{ + auto const keys = key_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = [] { + auto child1 = + strings_column_wrapper{"año", "bit", "₹1", "aaa", "zit", "bat", "aab", "$1", "€1", "wut"}; + auto child2 = fixed_width_column_wrapper{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const expect_keys = key_wrapper{1, 1, 1, 2, 2, 2, 2, 3, 3, 3}; + auto const expect_vals = [] { + auto child1 = + strings_column_wrapper{"año", "aaa", "aaa", "bit", "bit", "bat", "bat", "₹1", "$1", "$1"}; + auto child2 = fixed_width_column_wrapper{1, 4, 4, 2, 2, 6, 6, 3, 8, 8}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto agg = cudf::make_min_aggregation(); + test_single_scan(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + +TEST_F(groupby_min_scan_struct_test, slice_input) +{ + constexpr int32_t dont_care{1}; + auto const keys_original = + key_wrapper{dont_care, dont_care, 1, 2, 3, 1, 2, 2, 1, 3, 3, 2, dont_care}; + auto const vals_original = [] { + auto child1 = strings_column_wrapper{"dont_care", + "dont_care", + "año", + "bit", + "₹1", + "aaa", + "zit", + "bat", + "aab", + "$1", + "€1", + "wut", + "dont_care"}; + auto child2 = key_wrapper{dont_care, dont_care, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, dont_care}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const keys = cudf::slice(keys_original, {2, 12})[0]; + auto const vals = cudf::slice(vals_original, {2, 12})[0]; + auto const expect_keys = key_wrapper{1, 1, 1, 2, 2, 2, 2, 3, 3, 3}; + auto const expect_vals = [] { + auto child1 = + strings_column_wrapper{"año", "aaa", "aaa", "bit", "bit", "bat", "bat", "₹1", "$1", "$1"}; + auto child2 = fixed_width_column_wrapper{1, 4, 4, 2, 2, 6, 6, 3, 8, 8}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto agg = cudf::make_min_aggregation(); + test_single_scan(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + +TEST_F(groupby_min_scan_struct_test, null_keys_and_values) +{ + constexpr int32_t null{0}; + auto const keys = key_wrapper{{1, 2, 3, 1, 2, 2, 1, null, 3, 2, 4}, null_at(7)}; + auto const vals = [] { + auto child1 = strings_column_wrapper{ + "año", "bit", "₹1", "aaa", "zit", "" /*NULL*/, "" /*NULL*/, "$1", "€1", "wut", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 8, 7, 6, 5, null, null, 2, 1, 0, null}; + return structs_column_wrapper{{child1, child2}, nulls_at({5, 6, 10})}; + }(); + + auto const expect_keys = key_wrapper{{1, 1, 1, 2, 2, 2, 2, 3, 3, 4}, no_nulls()}; + auto const expect_vals = [] { + auto child1 = strings_column_wrapper{ + "año", "aaa", "" /*NULL*/, "bit", "bit", "" /*NULL*/, "bit", "₹1", "€1", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 6, null, 8, 8, null, 8, 7, 1, null}; + return structs_column_wrapper{{child1, child2}, nulls_at({2, 5, 9})}; + }(); + + auto agg = cudf::make_min_aggregation(); + test_single_scan(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + } // namespace test } // namespace cudf diff --git a/cpp/tests/groupby/min_tests.cpp b/cpp/tests/groupby/min_tests.cpp index 4f8db1d750c..c2cfca83b29 100644 --- a/cpp/tests/groupby/min_tests.cpp +++ b/cpp/tests/groupby/min_tests.cpp @@ -33,7 +33,7 @@ struct groupby_min_test : public cudf::test::BaseFixture { }; using K = int32_t; -TYPED_TEST_CASE(groupby_min_test, cudf::test::FixedWidthTypesWithoutFixedPoint); +TYPED_TEST_SUITE(groupby_min_test, cudf::test::FixedWidthTypesWithoutFixedPoint); TYPED_TEST(groupby_min_test, basic) { @@ -255,7 +255,7 @@ template struct FixedPointTestBothReps : public cudf::test::BaseFixture { }; -TYPED_TEST_CASE(FixedPointTestBothReps, cudf::test::FixedPointTypes); +TYPED_TEST_SUITE(FixedPointTestBothReps, cudf::test::FixedPointTypes); TYPED_TEST(FixedPointTestBothReps, GroupBySortMinDecimalAsValue) { @@ -303,5 +303,89 @@ TYPED_TEST(FixedPointTestBothReps, GroupByHashMinDecimalAsValue) } } +struct groupby_min_struct_test : public cudf::test::BaseFixture { +}; + +TEST_F(groupby_min_struct_test, basic) +{ + auto const keys = fixed_width_column_wrapper{1, 2, 3, 1, 2, 2, 1, 3, 3, 2}; + auto const vals = [] { + auto child1 = + strings_column_wrapper{"año", "bit", "₹1", "aaa", "zit", "bat", "aab", "$1", "€1", "wut"}; + auto child2 = fixed_width_column_wrapper{1, 2, 3, 4, 5, 6, 7, 8, 9, 10}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals = [] { + auto child1 = strings_column_wrapper{"aaa", "bat", "$1"}; + auto child2 = fixed_width_column_wrapper{4, 6, 8}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto agg = cudf::make_min_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + +TEST_F(groupby_min_struct_test, slice_input) +{ + constexpr int32_t dont_care{1}; + auto const keys_original = fixed_width_column_wrapper{ + dont_care, dont_care, 1, 2, 3, 1, 2, 2, 1, 3, 3, 2, dont_care}; + auto const vals_original = [] { + auto child1 = strings_column_wrapper{"dont_care", + "dont_care", + "año", + "bit", + "₹1", + "aaa", + "zit", + "bat", + "aab", + "$1", + "€1", + "wut", + "dont_care"}; + auto child2 = fixed_width_column_wrapper{ + dont_care, dont_care, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, dont_care}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto const keys = cudf::slice(keys_original, {2, 12})[0]; + auto const vals = cudf::slice(vals_original, {2, 12})[0]; + auto const expect_keys = fixed_width_column_wrapper{1, 2, 3}; + auto const expect_vals = [] { + auto child1 = strings_column_wrapper{"aaa", "bat", "$1"}; + auto child2 = fixed_width_column_wrapper{4, 6, 8}; + return structs_column_wrapper{{child1, child2}}; + }(); + + auto agg = cudf::make_min_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + +TEST_F(groupby_min_struct_test, null_keys_and_values) +{ + constexpr int32_t null{0}; + auto const keys = + fixed_width_column_wrapper{{1, 2, 3, 1, 2, 2, 1, null, 3, 2, 4}, null_at(7)}; + auto const vals = [] { + auto child1 = strings_column_wrapper{ + "año", "bit", "₹1", "aaa", "zit", "" /*NULL*/, "" /*NULL*/, "$1", "€1", "wut", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{9, 8, 7, 6, 5, null, null, 2, 1, 0, null}; + return structs_column_wrapper{{child1, child2}, nulls_at({5, 6, 10})}; + }(); + + auto const expect_keys = fixed_width_column_wrapper{{1, 2, 3, 4}, no_nulls()}; + auto const expect_vals = [] { + auto child1 = strings_column_wrapper{"aaa", "bit", "€1", "" /*NULL*/}; + auto child2 = fixed_width_column_wrapper{6, 8, 1, null}; + return structs_column_wrapper{{child1, child2}, null_at(3)}; + }(); + + auto agg = cudf::make_min_aggregation(); + test_single_agg(keys, vals, expect_keys, expect_vals, std::move(agg)); +} + } // namespace test } // namespace cudf