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.

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

Cleanup groupby to use host_span, device_span, device_uvector #7698

Merged
merged 5 commits into from
Mar 26, 2021
Merged
Show file tree
Hide file tree
Changes from 3 commits
Commits
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
7 changes: 4 additions & 3 deletions cpp/include/cudf/detail/groupby.hpp
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.
Expand All @@ -16,6 +16,7 @@

#include <cudf/groupby.hpp>
#include <cudf/types.hpp>
#include <cudf/utilities/span.hpp>

#include <rmm/cuda_stream_view.hpp>

Expand All @@ -36,12 +37,12 @@ namespace hash {
* @return true A hash-based groupby can be used
* @return false A hash-based groupby cannot be used
*/
bool can_use_hash_groupby(table_view const& keys, std::vector<aggregation_request> const& requests);
bool can_use_hash_groupby(table_view const& keys, host_span<aggregation_request const> requests);

// Hash-based groupby
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby(
table_view const& keys,
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
null_policy include_null_keys,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr);
Expand Down
11 changes: 6 additions & 5 deletions cpp/include/cudf/groupby.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,7 @@
#include <cudf/aggregation.hpp>
#include <cudf/table/table_view.hpp>
#include <cudf/types.hpp>
#include <cudf/utilities/span.hpp>

#include <rmm/cuda_stream_view.hpp>

Expand Down Expand Up @@ -163,7 +164,7 @@ class groupby {
* specified in `requests`.
*/
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> aggregate(
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

/**
Expand Down Expand Up @@ -218,7 +219,7 @@ class groupby {
* specified in `requests`.
*/
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> scan(
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

/**
Expand Down Expand Up @@ -277,18 +278,18 @@ class groupby {
* aggregation requests.
*/
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> dispatch_aggregation(
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr);

// Sort-based groupby
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> sort_aggregate(
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr);

std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> sort_scan(
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr);
};
Expand Down
3 changes: 2 additions & 1 deletion cpp/src/groupby/common/utils.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -18,13 +18,14 @@

#include <cudf/detail/aggregation/result_cache.hpp>
#include <cudf/detail/groupby.hpp>
#include <cudf/utilities/span.hpp>
#include <vector>

namespace cudf {
namespace groupby {
namespace detail {
inline std::vector<aggregation_result> extract_results(
std::vector<aggregation_request> const& requests, cudf::detail::result_cache& cache)
host_span<aggregation_request const> requests, cudf::detail::result_cache& cache)
{
std::vector<aggregation_result> results(requests.size());

Expand Down
10 changes: 5 additions & 5 deletions cpp/src/groupby/groupby.cu
Original file line number Diff line number Diff line change
Expand Up @@ -55,7 +55,7 @@ groupby::groupby(table_view const& keys,

// Select hash vs. sort groupby implementation
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby::dispatch_aggregation(
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
Expand All @@ -79,7 +79,7 @@ groupby::~groupby() = default;

namespace {
/// Make an empty table with appropriate types for requested aggs
auto empty_results(std::vector<aggregation_request> const& requests)
auto empty_results(host_span<aggregation_request const> requests)
{
std::vector<aggregation_result> empty_results;

Expand All @@ -102,7 +102,7 @@ auto empty_results(std::vector<aggregation_request> const& requests)
}

/// Verifies the agg requested on the request's values is valid
void verify_valid_requests(std::vector<aggregation_request> const& requests)
void verify_valid_requests(host_span<aggregation_request const> requests)
{
CUDF_EXPECTS(
std::all_of(
Expand Down Expand Up @@ -143,7 +143,7 @@ void verify_valid_requests(std::vector<aggregation_request> const& requests)

// Compute aggregation requests
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby::aggregate(
std::vector<aggregation_request> const& requests, rmm::mr::device_memory_resource* mr)
host_span<aggregation_request const> requests, rmm::mr::device_memory_resource* mr)
{
CUDF_FUNC_RANGE();
CUDF_EXPECTS(
Expand All @@ -161,7 +161,7 @@ std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby::aggr

// Compute scan requests
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby::scan(
std::vector<aggregation_request> const& requests, rmm::mr::device_memory_resource* mr)
host_span<aggregation_request const> requests, rmm::mr::device_memory_resource* mr)
{
CUDF_FUNC_RANGE();
CUDF_EXPECTS(
Expand Down
24 changes: 12 additions & 12 deletions cpp/src/groupby/hash/groupby.cu
Original file line number Diff line number Diff line change
Expand Up @@ -110,7 +110,7 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final
data_type result_type;
cudf::detail::result_cache* sparse_results;
cudf::detail::result_cache* dense_results;
rmm::device_vector<size_type> const& gather_map;
device_span<size_type const> gather_map;
size_type const map_size;
Map const& map;
bitmask_type const* __restrict__ row_bitmask;
Expand All @@ -122,7 +122,7 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final
column_view col,
cudf::detail::result_cache* sparse_results,
cudf::detail::result_cache* dense_results,
rmm::device_vector<size_type> const& gather_map,
device_span<size_type const> gather_map,
size_type map_size,
Map const& map,
bitmask_type const* row_bitmask,
Expand Down Expand Up @@ -272,7 +272,7 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final

// flatten aggs to filter in single pass aggs
std::tuple<table_view, std::vector<aggregation::Kind>, std::vector<size_t>>
flatten_single_pass_aggs(std::vector<aggregation_request> const& requests)
flatten_single_pass_aggs(host_span<aggregation_request const> requests)
{
std::vector<column_view> columns;
std::vector<aggregation::Kind> agg_kinds;
Expand Down Expand Up @@ -311,10 +311,10 @@ flatten_single_pass_aggs(std::vector<aggregation_request> const& requests)
*/
template <typename Map>
void sparse_to_dense_results(table_view const& keys,
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
cudf::detail::result_cache* sparse_results,
cudf::detail::result_cache* dense_results,
rmm::device_vector<size_type> const& gather_map,
device_span<size_type const> gather_map,
size_type map_size,
Map const& map,
bool keys_have_nulls,
Expand Down Expand Up @@ -421,7 +421,7 @@ auto create_sparse_results_table(table_view const& flattened_values,
*/
template <bool keys_have_nulls, typename Map>
void compute_single_pass_aggs(table_view const& keys,
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
cudf::detail::result_cache* sparse_results,
Map& map,
null_policy include_null_keys,
Expand Down Expand Up @@ -469,10 +469,10 @@ void compute_single_pass_aggs(table_view const& keys,
* `map`.
*/
template <typename Map>
std::pair<rmm::device_vector<size_type>, size_type> extract_populated_keys(
std::pair<rmm::device_uvector<size_type>, size_type> extract_populated_keys(
Map map, size_type num_keys, rmm::cuda_stream_view stream)
{
rmm::device_vector<size_type> populated_keys(num_keys);
rmm::device_uvector<size_type> populated_keys(num_keys, stream);

auto get_key = [] __device__(auto const& element) {
size_type key, value;
Expand Down Expand Up @@ -520,7 +520,7 @@ std::pair<rmm::device_vector<size_type>, size_type> extract_populated_keys(
*/
template <bool keys_have_nulls>
std::unique_ptr<table> groupby_null_templated(table_view const& keys,
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
cudf::detail::result_cache* cache,
null_policy include_null_keys,
rmm::cuda_stream_view stream,
Expand All @@ -539,7 +539,7 @@ std::unique_ptr<table> groupby_null_templated(table_view const& keys,

// Extract the populated indices from the hash map and create a gather map.
// Gathering using this map from sparse results will give dense results.
rmm::device_vector<size_type> gather_map;
rmm::device_uvector<size_type> gather_map(0, stream);
size_type map_size;
std::tie(gather_map, map_size) = extract_populated_keys(*map, keys.num_rows(), stream);
karthikeyann marked this conversation as resolved.
Show resolved Hide resolved

Expand Down Expand Up @@ -576,7 +576,7 @@ std::unique_ptr<table> groupby_null_templated(table_view const& keys,
* @return true A hash-based groupby should be used
* @return false A hash-based groupby should not be used
*/
bool can_use_hash_groupby(table_view const& keys, std::vector<aggregation_request> const& requests)
bool can_use_hash_groupby(table_view const& keys, host_span<aggregation_request const> 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) {
Expand All @@ -588,7 +588,7 @@ bool can_use_hash_groupby(table_view const& keys, std::vector<aggregation_reques
// Hash-based groupby
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby(
table_view const& keys,
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
null_policy include_null_keys,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/groupby/sort/aggregate.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -385,7 +385,7 @@ void aggregrate_result_functor::operator()<aggregation::COLLECT_SET>(aggregation

// Sort-based groupby
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby::sort_aggregate(
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
Expand Down
10 changes: 5 additions & 5 deletions cpp/src/groupby/sort/group_nth_element.cu
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* Copyright (c) 2020, NVIDIA CORPORATION.
* Copyright (c) 2020-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.
Expand Down Expand Up @@ -69,17 +69,17 @@ std::unique_ptr<column> group_nth_element(column_view const &values,
auto bitmask_iterator =
thrust::make_transform_iterator(cudf::detail::make_validity_iterator(*values_view),
[] __device__(auto b) { return static_cast<size_type>(b); });
rmm::device_vector<size_type> intra_group_index(values.size());
rmm::device_uvector<size_type> intra_group_index(values.size(), stream);
// intra group index for valids only.
thrust::exclusive_scan_by_key(rmm::exec_policy(stream),
group_labels.begin(),
group_labels.end(),
bitmask_iterator,
intra_group_index.begin());
// group_size to recalculate n if n<0
rmm::device_vector<size_type> group_count = [&] {
rmm::device_uvector<size_type> group_count = [&] {
if (n < 0) {
rmm::device_vector<size_type> group_count(num_groups);
rmm::device_uvector<size_type> group_count(num_groups, stream);
thrust::reduce_by_key(rmm::exec_policy(stream),
group_labels.begin(),
group_labels.end(),
Expand All @@ -88,7 +88,7 @@ std::unique_ptr<column> group_nth_element(column_view const &values,
group_count.begin());
return group_count;
} else {
return rmm::device_vector<size_type>();
return rmm::device_uvector<size_type>(0, stream);
}
}();
// gather the valid index == n
Expand Down
8 changes: 4 additions & 4 deletions cpp/src/groupby/sort/group_quantiles.cu
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.
Expand Down Expand Up @@ -80,7 +80,7 @@ struct quantiles_functor {
column_view const& group_sizes,
cudf::device_span<size_type const> group_offsets,
size_type const num_groups,
rmm::device_vector<double> const& quantile,
device_span<double const> quantile,
interpolation interpolation,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
Expand Down Expand Up @@ -112,7 +112,7 @@ struct quantiles_functor {
*group_size_view,
*result_view,
group_offsets.data(),
quantile.data().get(),
quantile.data(),
static_cast<size_type>(quantile.size()),
interpolation});
} else {
Expand All @@ -125,7 +125,7 @@ struct quantiles_functor {
*group_size_view,
*result_view,
group_offsets.data(),
quantile.data().get(),
quantile.data(),
static_cast<size_type>(quantile.size()),
interpolation});
}
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/groupby/sort/scan.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -107,7 +107,7 @@ void scan_result_functor::operator()<aggregation::COUNT_ALL>(aggregation const&

// Sort-based groupby
std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby::sort_scan(
std::vector<aggregation_request> const& requests,
host_span<aggregation_request const> requests,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
Expand Down