From f2239016fe91e89a71ed9d3639bfaad081f98999 Mon Sep 17 00:00:00 2001 From: Conor Hoekstra <36027403+codereport@users.noreply.github.com> Date: Thu, 9 Sep 2021 00:49:20 -0400 Subject: [PATCH] Minor C++17 cleanup of `groupby.cu`: structured bindings, more concise lambda, etc (#9193) Due to the fact that my soul is slowly melting away due to my current internal PR I am working on, I set up this small cleanup in hopes of getting a dopamine or endorphin hit. Authors: - Conor Hoekstra (https://github.com/codereport) Approvers: - David Wendt (https://github.com/davidwendt) - Nghia Truong (https://github.com/ttnghia) - MithunR (https://github.com/mythrocks) - Jake Hemstad (https://github.com/jrhemstad) URL: https://github.com/rapidsai/cudf/pull/9193 --- cpp/src/groupby/hash/groupby.cu | 46 ++++++++++++++------------------- 1 file changed, 19 insertions(+), 27 deletions(-) diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu index 87f83c6edd6..694ca2afceb 100644 --- a/cpp/src/groupby/hash/groupby.cu +++ b/cpp/src/groupby/hash/groupby.cu @@ -537,27 +537,25 @@ void compute_single_pass_aggs(table_view const& keys, * `map`. */ template -std::pair, size_type> extract_populated_keys( - Map map, size_type num_keys, rmm::cuda_stream_view stream) +rmm::device_uvector extract_populated_keys(Map map, + size_type num_keys, + rmm::cuda_stream_view stream) { rmm::device_uvector populated_keys(num_keys, stream); - auto get_key = [] __device__(auto const& element) { - size_type key, value; - thrust::tie(key, value) = element; - return key; - }; + auto get_key = [] __device__(auto const& element) { return element.first; }; // first = key + auto get_key_it = thrust::make_transform_iterator(map.data(), get_key); + auto key_used = [unused = map.get_unused_key()] __device__(auto key) { return key != unused; }; - auto end_it = thrust::copy_if( - rmm::exec_policy(stream), - thrust::make_transform_iterator(map.data(), get_key), - thrust::make_transform_iterator(map.data() + map.capacity(), get_key), - populated_keys.begin(), - [unused_key = map.get_unused_key()] __device__(size_type key) { return key != unused_key; }); + auto end_it = thrust::copy_if(rmm::exec_policy(stream), + get_key_it, + get_key_it + map.capacity(), + populated_keys.begin(), + key_used); - size_type map_size = end_it - populated_keys.begin(); + populated_keys.resize(std::distance(populated_keys.begin(), end_it), stream); - return std::make_pair(std::move(populated_keys), map_size); + return populated_keys; } /** @@ -594,8 +592,8 @@ std::unique_ptr groupby_null_templated(table_view const& keys, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto d_keys = table_device_view::create(keys, stream); - auto map = create_hash_map(*d_keys, include_null_keys, stream); + auto d_keys_ptr = table_device_view::create(keys, stream); + auto map = create_hash_map(*d_keys_ptr, include_null_keys, stream); // Cache of sparse results where the location of aggregate value in each // column is indexed by the hash map @@ -607,9 +605,7 @@ std::unique_ptr
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. - auto map_and_size = extract_populated_keys(*map, keys.num_rows(), stream); - rmm::device_uvector gather_map{std::move(map_and_size.first)}; - size_type const map_size = map_and_size.second; + auto gather_map = extract_populated_keys(*map, keys.num_rows(), stream); // Compact all results from sparse_results and insert into cache sparse_to_dense_results(keys, @@ -617,19 +613,15 @@ std::unique_ptr
groupby_null_templated(table_view const& keys, &sparse_results, cache, gather_map, - map_size, + gather_map.size(), *map, keys_have_nulls, include_null_keys, stream, mr); - return cudf::detail::gather(keys, - gather_map.begin(), - gather_map.begin() + map_size, - out_of_bounds_policy::DONT_CHECK, - stream, - mr); + return cudf::detail::gather( + keys, gather_map.begin(), gather_map.end(), out_of_bounds_policy::DONT_CHECK, stream, mr); } } // namespace