fix merge conflict

rapidsai · Oct 4, 2021 · fe429a6 · fe429a6
2 parents c3d1a24 + d68e626
commit fe429a6
Show file tree

Hide file tree

Showing 24 changed files with 185 additions and 94 deletions.
diff --git a/cpp/include/cudf/detail/copy_if.cuh b/cpp/include/cudf/detail/copy_if.cuh
@@ -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.
@@ -19,7 +19,7 @@
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/copy.hpp>
-#include <cudf/detail/gather.cuh>
+#include <cudf/detail/gather.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/device_atomics.cuh>
@@ -36,12 +36,15 @@
 #include <rmm/device_buffer.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
+#include <rmm/exec_policy.hpp>
 
 #include <cub/cub.cuh>
 
 #include <algorithm>
 
-namespace {
+namespace cudf {
+namespace detail {
+
 // Compute the count of elements that pass the mask within each block
 template <typename Filter, int block_size>
 __global__ void compute_block_counts(cudf::size_type* __restrict__ block_counts,
@@ -293,9 +296,9 @@ struct scatter_gather_functor {
                     filter);
 
     auto output_table = cudf::detail::gather(cudf::table_view{{input}},
-                                             indices.begin(),
-                                             indices.end(),
+                                             indices,
                                              cudf::out_of_bounds_policy::DONT_CHECK,
+                                             cudf::detail::negative_index_policy::NOT_ALLOWED,
                                              stream,
                                              mr);
 
@@ -304,10 +307,6 @@ struct scatter_gather_functor {
   }
 };
 
-}  // namespace
-
-namespace cudf {
-namespace detail {
 /**
  * @brief Filters `input` using a Filter function object
  *

diff --git a/cpp/include/cudf/detail/gather.hpp b/cpp/include/cudf/detail/gather.hpp
@@ -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.
@@ -16,10 +16,10 @@
 #pragma once
 
 #include <cudf/column/column_view.hpp>
-#include <cudf/table/table_view.hpp>
-
 #include <cudf/copying.hpp>
 #include <cudf/table/table.hpp>
+#include <cudf/table/table_view.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 
@@ -55,10 +55,10 @@ enum class negative_index_policy : bool { ALLOWED, NOT_ALLOWED };
  * indices. If `policy` is set to `DONT_CHECK` and there are out-of-bounds indices in `gather_map`,
  * the behavior is undefined.
  * @param[in] negative_index_policy Interpret each negative index `i` in the
- * gathermap as the positive index `i+num_source_rows`.
+ * `gather_map` as the positive index `i+num_source_rows`.
  * @param[in] stream CUDA stream used for device memory operations and kernel launches.
  * @param[in] mr Device memory resource used to allocate the returned table's device memory
- * @return cudf::table Result of the gather
+ * @return Result of the gather
  */
 std::unique_ptr<table> gather(
   table_view const& source_table,
@@ -67,5 +67,21 @@ std::unique_ptr<table> gather(
   negative_index_policy neg_indices,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @copydoc cudf::detail::gather(table_view const&,column_view const&,table_view
+ * const&,cudf::out_of_bounds_policy,cudf::detail::negative_index_policy,rmm::cuda_stream_view,
+ * rmm::mr::device_memory_resource*)
+ *
+ * @throws cudf::logic_error if `gather_map` span size is larger than max of `size_type`.
+ */
+std::unique_ptr<table> gather(
+  table_view const& source_table,
+  device_span<size_type const> const gather_map,
+  out_of_bounds_policy bounds_policy,
+  negative_index_policy neg_indices,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 }  // namespace detail
 }  // namespace cudf
diff --git a/cpp/include/cudf/detail/scatter.hpp b/cpp/include/cudf/detail/scatter.hpp
@@ -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.
@@ -18,7 +18,7 @@
 #include <cudf/column/column_view.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/table/table_view.hpp>
-#include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 
@@ -70,6 +70,20 @@ std::unique_ptr<table> scatter(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @copydoc cudf::detail::scatter(table_view const&,column_view const&,table_view
+ * const&,bool,rmm::cuda_stream_view,rmm::mr::device_memory_resource*)
+ *
+ * @throws cudf::logic_error if `scatter_map` span size is larger than max of `size_type`.
+ */
+std::unique_ptr<table> scatter(
+  table_view const& source,
+  device_span<size_type const> const scatter_map,
+  table_view const& target,
+  bool check_bounds                   = false,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 /**
  * @brief Scatters a row of scalar values into a copy of the target table
  * according to a scatter map.

diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu
@@ -14,13 +14,12 @@
  * limitations under the License.
  */
 
-#include <cudf/copying.hpp>
 #include <cudf/detail/copy.hpp>
 #include <cudf/detail/copy_if_else.cuh>
-#include <cudf/detail/gather.cuh>
+#include <cudf/detail/gather.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
-#include <cudf/detail/scatter.cuh>
+#include <cudf/detail/scatter.hpp>
 #include <cudf/strings/detail/copy_if_else.cuh>
 #include <cudf/strings/string_view.cuh>
 #include <cudf/utilities/traits.hpp>
@@ -168,23 +167,24 @@ std::unique_ptr<column> scatter_gather_based_if_else(cudf::column_view const& lh
                                                      rmm::cuda_stream_view stream,
                                                      rmm::mr::device_memory_resource* mr)
 {
-  auto scatter_map = rmm::device_uvector<size_type>{static_cast<std::size_t>(size), stream};
-  auto const scatter_map_end = thrust::copy_if(rmm::exec_policy(stream),
-                                               thrust::make_counting_iterator(size_type{0}),
-                                               thrust::make_counting_iterator(size_type{size}),
-                                               scatter_map.begin(),
-                                               is_left);
+  auto gather_map = rmm::device_uvector<size_type>{static_cast<std::size_t>(size), stream};
+  auto const gather_map_end = thrust::copy_if(rmm::exec_policy(stream),
+                                              thrust::make_counting_iterator(size_type{0}),
+                                              thrust::make_counting_iterator(size_type{size}),
+                                              gather_map.begin(),
+                                              is_left);
+
+  gather_map.resize(thrust::distance(gather_map.begin(), gather_map_end), stream);
 
   auto const scatter_src_lhs = cudf::detail::gather(table_view{std::vector<column_view>{lhs}},
-                                                    scatter_map.begin(),
-                                                    scatter_map_end,
+                                                    gather_map,
                                                     out_of_bounds_policy::DONT_CHECK,
+                                                    negative_index_policy::NOT_ALLOWED,
                                                     stream);
 
   auto result = cudf::detail::scatter(
     table_view{std::vector<column_view>{scatter_src_lhs->get_column(0).view()}},
-    scatter_map.begin(),
-    scatter_map_end,
+    gather_map,
     table_view{std::vector<column_view>{rhs}},
     false,
     stream,
@@ -214,8 +214,12 @@ std::unique_ptr<column> scatter_gather_based_if_else(cudf::scalar const& lhs,
                                                    static_cast<cudf::size_type>(scatter_map_size),
                                                    scatter_map.begin()};
 
-  auto result = cudf::scatter(
-    scatter_source, scatter_map_column_view, table_view{std::vector<column_view>{rhs}}, false, mr);
+  auto result = cudf::detail::scatter(scatter_source,
+                                      scatter_map_column_view,
+                                      table_view{std::vector<column_view>{rhs}},
+                                      false,
+                                      stream,
+                                      mr);
 
   return std::move(result->release()[0]);
 }

diff --git a/cpp/src/copying/gather.cu b/cpp/src/copying/gather.cu
@@ -54,6 +54,21 @@ std::unique_ptr<table> gather(table_view const& source_table,
   return gather(source_table, map_begin, map_end, bounds_policy, stream, mr);
 }
 
+std::unique_ptr<table> gather(table_view const& source_table,
+                              device_span<size_type const> const gather_map,
+                              out_of_bounds_policy bounds_policy,
+                              negative_index_policy neg_indices,
+                              rmm::cuda_stream_view stream,
+                              rmm::mr::device_memory_resource* mr)
+{
+  CUDF_EXPECTS(gather_map.size() <= std::numeric_limits<size_type>::max(),
+               "invalid gather map size");
+  auto map_col = column_view(data_type{type_to_id<size_type>()},
+                             static_cast<size_type>(gather_map.size()),
+                             gather_map.data());
+  return gather(source_table, map_col, bounds_policy, neg_indices, stream, mr);
+}
+
 }  // namespace detail
 
 std::unique_ptr<table> gather(table_view const& source_table,

diff --git a/cpp/src/copying/scatter.cu b/cpp/src/copying/scatter.cu
@@ -16,8 +16,6 @@
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/copying.hpp>
 #include <cudf/detail/copy.hpp>
-#include <cudf/detail/gather.cuh>
-#include <cudf/detail/gather.hpp>
 #include <cudf/detail/indexalator.cuh>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
@@ -305,6 +303,21 @@ std::unique_ptr<table> scatter(table_view const& source,
   return detail::scatter(source, map_begin, map_end, target, check_bounds, stream, mr);
 }
 
+std::unique_ptr<table> scatter(table_view const& source,
+                               device_span<size_type const> const scatter_map,
+                               table_view const& target,
+                               bool check_bounds,
+                               rmm::cuda_stream_view stream,
+                               rmm::mr::device_memory_resource* mr)
+{
+  CUDF_EXPECTS(scatter_map.size() <= std::numeric_limits<size_type>::max(),
+               "invalid scatter map size");
+  auto map_col = column_view(data_type{type_to_id<size_type>()},
+                             static_cast<size_type>(scatter_map.size()),
+                             scatter_map.data());
+  return scatter(source, map_col, target, check_bounds, stream, mr);
+}
+
 std::unique_ptr<table> scatter(std::vector<std::reference_wrapper<const scalar>> const& source,
                                column_view const& indices,
                                table_view const& target,

diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu
@@ -26,13 +26,14 @@
 #include <cudf/detail/aggregation/aggregation.hpp>
 #include <cudf/detail/aggregation/result_cache.hpp>
 #include <cudf/detail/binaryop.hpp>
-#include <cudf/detail/gather.cuh>
 #include <cudf/detail/gather.hpp>
 #include <cudf/detail/groupby.hpp>
+#include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/replace.hpp>
 #include <cudf/detail/unary.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
 #include <cudf/detail/utilities/hash_functions.cuh>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <cudf/groupby.hpp>
 #include <cudf/scalar/scalar.hpp>
@@ -167,7 +168,6 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final
   cudf::detail::result_cache* sparse_results;
   cudf::detail::result_cache* dense_results;
   device_span<size_type const> gather_map;
-  size_type const map_size;
   Map const& map;
   bitmask_type const* __restrict__ row_bitmask;
   rmm::cuda_stream_view stream;
@@ -181,7 +181,6 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final
                               cudf::detail::result_cache* sparse_results,
                               cudf::detail::result_cache* dense_results,
                               device_span<size_type const> gather_map,
-                              size_type map_size,
                               Map const& map,
                               bitmask_type const* row_bitmask,
                               rmm::cuda_stream_view stream,
@@ -191,7 +190,6 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final
       sparse_results(sparse_results),
       dense_results(dense_results),
       gather_map(gather_map),
-      map_size(map_size),
       map(map),
       row_bitmask(row_bitmask),
       stream(stream),
@@ -203,11 +201,12 @@ class hash_compound_agg_finalizer final : public cudf::detail::aggregation_final
 
   auto to_dense_agg_result(cudf::aggregation const& agg)
   {
-    auto s                  = sparse_results->get_result(col_idx, agg);
+    auto s = sparse_results->get_result(col_idx, agg);
+
     auto dense_result_table = cudf::detail::gather(table_view({std::move(s)}),
-                                                   gather_map.begin(),
-                                                   gather_map.begin() + map_size,
+                                                   gather_map,
                                                    out_of_bounds_policy::DONT_CHECK,
+                                                   cudf::detail::negative_index_policy::NOT_ALLOWED,
                                                    stream,
                                                    mr);
     return std::move(dense_result_table->release()[0]);
@@ -385,7 +384,7 @@ void sparse_to_dense_results(table_view const& keys,
                              cudf::detail::result_cache* sparse_results,
                              cudf::detail::result_cache* dense_results,
                              device_span<size_type const> gather_map,
-                             size_type map_size,
+                             // size_type map_size,
                              Map const& map,
                              bool keys_have_nulls,
                              null_policy include_null_keys,
@@ -403,16 +402,8 @@ void sparse_to_dense_results(table_view const& keys,
 
     // Given an aggregation, this will get the result from sparse_results and
     // convert and return dense, compacted result
-    auto finalizer = hash_compound_agg_finalizer<Map>(i,
-                                                      col,
-                                                      sparse_results,
-                                                      dense_results,
-                                                      gather_map,
-                                                      map_size,
-                                                      map,
-                                                      row_bitmask_ptr,
-                                                      stream,
-                                                      mr);
+    auto finalizer = hash_compound_agg_finalizer<Map>(
+      i, col, sparse_results, dense_results, gather_map, map, row_bitmask_ptr, stream, mr);
     for (auto&& agg : agg_v) {
       agg->finalize(finalizer);
     }
@@ -613,15 +604,18 @@ std::unique_ptr<table> groupby_null_templated(table_view const& keys,
                           &sparse_results,
                           cache,
                           gather_map,
-                          gather_map.size(),
                           *map,
                           keys_have_nulls,
                           include_null_keys,
                           stream,
                           mr);
 
-  return cudf::detail::gather(
-    keys, gather_map.begin(), gather_map.end(), out_of_bounds_policy::DONT_CHECK, stream, mr);
+  return cudf::detail::gather(keys,
+                              gather_map,
+                              out_of_bounds_policy::DONT_CHECK,
+                              cudf::detail::negative_index_policy::NOT_ALLOWED,
+                              stream,
+                              mr);
 }
 
 }  // namespace

diff --git a/cpp/src/groupby/sort/group_collect.cu b/cpp/src/groupby/sort/group_collect.cu
@@ -18,7 +18,6 @@
 #include <cudf/column/column_view.hpp>
 #include <cudf/detail/aggregation/aggregation.hpp>
 #include <cudf/detail/copy_if.cuh>
-#include <cudf/detail/gather.cuh>
 #include <cudf/strings/detail/utilities.cuh>
 #include <cudf/types.hpp>
 #include <cudf/utilities/span.hpp>