From c6bc111a18803b4a56e8090f237af00e5b5296ae Mon Sep 17 00:00:00 2001 From: Michael Wang Date: Wed, 6 Oct 2021 18:28:30 -0700 Subject: [PATCH] Add detail interface for `split` and `slice(table_view)`, refactors both function with `host_span` (#9226) `cudf::detail::slice` performs a `segmented_count_unset_bits` that requires stream ordering. The depending `split` interface does not have an internal version that accepts a `stream` argument. Similarly for `slice(table_view)`. This PR fixes that. Besides, slice/split interface is refactored to accept `host_span` to specify indices/splits, and is overloaded with `std::initializer_list`. This allows specifying the argument with both a container and a brace-init-list. Authors: - Michael Wang (https://github.com/isVoid) Approvers: - Nghia Truong (https://github.com/ttnghia) - H. Thomson Comer (https://github.com/thomcom) - David Wendt (https://github.com/davidwendt) URL: https://github.com/rapidsai/cudf/pull/9226 --- cpp/include/cudf/copying.hpp | 36 +++++++--- cpp/include/cudf/detail/copy.hpp | 61 +++++++++++++++-- cpp/src/copying/sample.cu | 5 +- cpp/src/copying/slice.cu | 39 ++++++++--- cpp/src/copying/split.cpp | 67 ++++++++++++++++--- cpp/src/dictionary/encode.cu | 6 +- cpp/src/io/csv/writer_impl.cu | 4 +- cpp/src/reductions/reductions.cpp | 6 +- cpp/src/strings/copying/copying.cu | 10 ++- cpp/src/strings/copying/shift.cu | 4 +- cpp/src/transpose/transpose.cu | 2 +- cpp/tests/quantiles/percentile_approx_test.cu | 2 +- 12 files changed, 199 insertions(+), 43 deletions(-) diff --git a/cpp/include/cudf/copying.hpp b/cpp/include/cudf/copying.hpp index 06768bdeb35..ba5043fb261 100644 --- a/cpp/include/cudf/copying.hpp +++ b/cpp/include/cudf/copying.hpp @@ -407,10 +407,15 @@ std::unique_ptr shift( * the range [0, input.size()). * * @param input View of column to slice - * @param indices A vector of indices used to take slices of `input`. + * @param indices Indices used to take slices of `input`. * @return Vector of views of `input` indicated by the ranges in `indices`. */ -std::vector slice(column_view const& input, std::vector const& indices); +std::vector slice(column_view const& input, host_span indices); +/** + * @ingroup copy_slice + * @copydoc cudf::slice(column_view const&, host_span) + */ +std::vector slice(column_view const& input, std::initializer_list indices); /** * @brief Slices a `table_view` into a set of `table_view`s according to a set of indices. @@ -441,10 +446,15 @@ std::vector slice(column_view const& input, std::vector * the range [0, input.size()). * * @param input View of table to slice - * @param indices A vector of indices used to take slices of `input`. + * @param indices Indices used to take slices of `input`. * @return Vector of views of `input` indicated by the ranges in `indices`. */ -std::vector slice(table_view const& input, std::vector const& indices); +std::vector slice(table_view const& input, host_span indices); +/** + * @ingroup copy_slice + * @copydoc cudf::slice(table_view const&, host_span) + */ +std::vector slice(table_view const& input, std::initializer_list indices); /** * @brief Splits a `column_view` into a set of `column_view`s according to a set of indices @@ -475,10 +485,15 @@ std::vector slice(table_view const& input, std::vector co * @throws cudf::logic_error When the values in the `splits` are 'strictly decreasing'. * * @param input View of column to split - * @param splits A vector of indices where the view will be split + * @param splits Indices where the view will be split * @return The set of requested views of `input` indicated by the `splits`. */ -std::vector split(column_view const& input, std::vector const& splits); +std::vector split(column_view const& input, host_span splits); +/** + * @ingroup copy_split + * @copydoc cudf::split(column_view const&, host_span) + */ +std::vector split(column_view const& input, std::initializer_list splits); /** * @brief Splits a `table_view` into a set of `table_view`s according to a set of indices @@ -511,10 +526,15 @@ std::vector split(column_view const& input, std::vector * @throws cudf::logic_error When the values in the `splits` are 'strictly decreasing'. * * @param input View of a table to split - * @param splits A vector of indices where the view will be split + * @param splits Indices where the view will be split * @return The set of requested views of `input` indicated by the `splits`. */ -std::vector split(table_view const& input, std::vector const& splits); +std::vector split(table_view const& input, host_span splits); +/** + * @ingroup copy_split + * @copydoc cudf::split(table_view const&, host_span) + */ +std::vector split(table_view const& input, std::initializer_list splits); /** * @brief Column data in a serialized format diff --git a/cpp/include/cudf/detail/copy.hpp b/cpp/include/cudf/detail/copy.hpp index 9f06661c8d1..50157d16876 100644 --- a/cpp/include/cudf/detail/copy.hpp +++ b/cpp/include/cudf/detail/copy.hpp @@ -19,10 +19,13 @@ #include #include #include +#include #include #include +#include + namespace cudf { namespace detail { /** @@ -67,21 +70,71 @@ ColumnView slice(ColumnView const& input, cudf::size_type begin, cudf::size_type } /** - * @copydoc cudf::slice(column_view const&,std::vector const&) + * @copydoc cudf::slice(column_view const&, host_span) * * @param stream CUDA stream used for device memory operations and kernel launches. */ std::vector slice(column_view const& input, - std::vector const& indices, + host_span indices, + rmm::cuda_stream_view stream = rmm::cuda_stream_default); +/** + * @copydoc cudf::slice(column_view const&, std::initializer_list) + * + * @param stream CUDA stream used for device memory operations and kernel launches. + */ +std::vector slice(column_view const& input, + std::initializer_list indices, rmm::cuda_stream_view stream = rmm::cuda_stream_default); /** - * @copydoc cudf::slice(table_view const&,std::vector const&) + * @copydoc cudf::slice(table_view const&, host_span) + * + * @param stream CUDA stream used for device memory operations and kernel launches. + */ +std::vector slice(table_view const& input, + host_span indices, + rmm::cuda_stream_view stream = rmm::cuda_stream_default); +/** + * @copydoc cudf::slice(table_view const&, std::initializer_list) * * @param stream CUDA stream used for device memory operations and kernel launches. */ std::vector slice(table_view const& input, - std::vector const& indices, + std::initializer_list indices, + rmm::cuda_stream_view stream = rmm::cuda_stream_default); + +/** + * @copydoc cudf::split(column_view const&, host_span) + * + * @param stream CUDA stream used for device memory operations and kernel launches. + */ +std::vector split(column_view const& input, + host_span splits, + rmm::cuda_stream_view stream = rmm::cuda_stream_default); +/** + * @copydoc cudf::split(column_view const&, std::initializer_list) + * + * @param stream CUDA stream used for device memory operations and kernel launches. + */ +std::vector split(column_view const& input, + std::initializer_list splits, + rmm::cuda_stream_view stream = rmm::cuda_stream_default); + +/** + * @copydoc cudf::split(table_view const&, host_span) + * + * @param stream CUDA stream used for device memory operations and kernel launches. + */ +std::vector split(table_view const& input, + host_span splits, + rmm::cuda_stream_view stream = rmm::cuda_stream_default); +/** + * @copydoc cudf::split(table_view const&, std::initializer_list) + * + * @param stream CUDA stream used for device memory operations and kernel launches. + */ +std::vector split(table_view const& input, + std::initializer_list splits, rmm::cuda_stream_view stream = rmm::cuda_stream_default); /** diff --git a/cpp/src/copying/sample.cu b/cpp/src/copying/sample.cu index 42dc9f76b18..3e0b27e9f19 100644 --- a/cpp/src/copying/sample.cu +++ b/cpp/src/copying/sample.cu @@ -70,8 +70,9 @@ std::unique_ptr sample(table_view const& input, gather_map_mutable_view.begin(), thrust::default_random_engine(seed)); - auto gather_map_view = - (n == num_rows) ? gather_map->view() : cudf::slice(gather_map->view(), {0, n})[0]; + auto gather_map_view = (n == num_rows) + ? gather_map->view() + : cudf::detail::slice(gather_map->view(), {0, n}, stream)[0]; return detail::gather(input, gather_map_view.begin(), gather_map_view.end(), diff --git a/cpp/src/copying/slice.cu b/cpp/src/copying/slice.cu index d1c12056393..9a3e349b907 100644 --- a/cpp/src/copying/slice.cu +++ b/cpp/src/copying/slice.cu @@ -29,7 +29,7 @@ namespace cudf { namespace detail { std::vector slice(column_view const& input, - std::vector const& indices, + host_span indices, rmm::cuda_stream_view stream) { CUDF_EXPECTS(indices.size() % 2 == 0, "indices size must be even"); @@ -64,16 +64,15 @@ std::vector slice(column_view const& input, } std::vector slice(table_view const& input, - std::vector const& indices, + host_span indices, rmm::cuda_stream_view stream) { - CUDF_FUNC_RANGE(); CUDF_EXPECTS(indices.size() % 2 == 0, "indices size must be even"); if (indices.empty()) { return {}; } // 2d arrangement of column_views that represent the outgoing table_views sliced_table[i][j] // where i is the i'th column of the j'th table_view - auto op = [&indices, stream](auto const& c) { return cudf::detail::slice(c, indices, stream); }; + auto op = [&indices, &stream](auto const& c) { return cudf::detail::slice(c, indices, stream); }; auto f = thrust::make_transform_iterator(input.begin(), op); auto sliced_table = std::vector>(f, f + input.num_columns()); @@ -93,20 +92,44 @@ std::vector slice(table_view const& input, return result; } +std::vector slice(column_view const& input, + std::initializer_list indices, + rmm::cuda_stream_view stream) +{ + return slice(input, host_span(indices.begin(), indices.size()), stream); +} + +std::vector slice(table_view const& input, + std::initializer_list indices, + rmm::cuda_stream_view stream) +{ + return slice(input, host_span(indices.begin(), indices.size()), stream); +}; + } // namespace detail -std::vector slice(cudf::column_view const& input, - std::vector const& indices) +std::vector slice(column_view const& input, host_span indices) { CUDF_FUNC_RANGE(); return detail::slice(input, indices, rmm::cuda_stream_default); } -std::vector slice(cudf::table_view const& input, - std::vector const& indices) +std::vector slice(table_view const& input, host_span indices) +{ + CUDF_FUNC_RANGE(); + return detail::slice(input, indices, rmm::cuda_stream_default); +}; + +std::vector slice(column_view const& input, std::initializer_list indices) { CUDF_FUNC_RANGE(); return detail::slice(input, indices, rmm::cuda_stream_default); } +std::vector slice(table_view const& input, std::initializer_list indices) +{ + CUDF_FUNC_RANGE(); + return detail::slice(input, indices, rmm::cuda_stream_default); +}; + } // namespace cudf diff --git a/cpp/src/copying/split.cpp b/cpp/src/copying/split.cpp index 97520800408..0fa802eb4b2 100644 --- a/cpp/src/copying/split.cpp +++ b/cpp/src/copying/split.cpp @@ -15,16 +15,22 @@ */ #include -#include +#include #include #include +#include + #include namespace cudf { +namespace detail { namespace { template -std::vector split(T const& input, size_type column_size, std::vector const& splits) +std::vector split(T const& input, + size_type column_size, + host_span splits, + rmm::cuda_stream_view stream) { if (splits.empty() or column_size == 0) { return std::vector{input}; } CUDF_EXPECTS(splits.back() <= column_size, "splits can't exceed size of input columns"); @@ -38,24 +44,67 @@ std::vector split(T const& input, size_type column_size, std::vector split(cudf::column_view const& input, - std::vector const& splits) + host_span splits, + rmm::cuda_stream_view stream) { - CUDF_FUNC_RANGE(); - return split(input, input.size(), splits); + return split(input, input.size(), splits, stream); } std::vector split(cudf::table_view const& input, - std::vector const& splits) + host_span splits, + rmm::cuda_stream_view stream) { - CUDF_FUNC_RANGE(); std::vector result{}; if (input.num_columns() == 0) { return result; } - return split(input, input.column(0).size(), splits); + return split(input, input.column(0).size(), splits, stream); +} + +std::vector split(column_view const& input, + std::initializer_list splits, + rmm::cuda_stream_view stream) +{ + return split(input, host_span(splits.begin(), splits.size()), stream); +} + +std::vector split(table_view const& input, + std::initializer_list splits, + rmm::cuda_stream_view stream) +{ + return detail::split(input, host_span(splits.begin(), splits.size()), stream); +} + +} // namespace detail + +std::vector split(cudf::column_view const& input, + host_span splits) +{ + CUDF_FUNC_RANGE(); + return detail::split(input, splits, rmm::cuda_stream_default); +} + +std::vector split(cudf::table_view const& input, + host_span splits) +{ + CUDF_FUNC_RANGE(); + return detail::split(input, splits, rmm::cuda_stream_default); +} + +std::vector split(column_view const& input, std::initializer_list splits) +{ + CUDF_FUNC_RANGE(); + return detail::split(input, splits, rmm::cuda_stream_default); +} + +std::vector split(table_view const& input, std::initializer_list splits) +{ + CUDF_FUNC_RANGE(); + return detail::split(input, splits, rmm::cuda_stream_default); } } // namespace cudf diff --git a/cpp/src/dictionary/encode.cu b/cpp/src/dictionary/encode.cu index 501e034c5fe..839b28413a6 100644 --- a/cpp/src/dictionary/encode.cu +++ b/cpp/src/dictionary/encode.cu @@ -15,7 +15,7 @@ */ #include -#include +#include #include #include #include @@ -53,7 +53,9 @@ std::unique_ptr encode(column_view const& input_column, if (keys_column->has_nulls()) { keys_column = std::make_unique( - slice(keys_column->view(), std::vector{0, keys_column->size() - 1}).front(), + cudf::detail::slice( + keys_column->view(), std::vector{0, keys_column->size() - 1}, stream) + .front(), stream, mr); keys_column->set_null_mask(rmm::device_buffer{0, stream, mr}, 0); // remove the null-mask diff --git a/cpp/src/io/csv/writer_impl.cu b/cpp/src/io/csv/writer_impl.cu index 9a0c701ea49..3a7b4bace9c 100644 --- a/cpp/src/io/csv/writer_impl.cu +++ b/cpp/src/io/csv/writer_impl.cu @@ -22,7 +22,7 @@ #include "writer_impl.hpp" #include -#include +#include #include #include #include @@ -423,7 +423,7 @@ void writer::impl::write(table_view const& table, }); // split table_view into chunks: - vector_views = cudf::split(table, splits); + vector_views = cudf::detail::split(table, splits, stream); } // convert each chunk to CSV: diff --git a/cpp/src/reductions/reductions.cpp b/cpp/src/reductions/reductions.cpp index 0d3ac2d366f..6f9149a47e2 100644 --- a/cpp/src/reductions/reductions.cpp +++ b/cpp/src/reductions/reductions.cpp @@ -68,7 +68,8 @@ struct reduce_dispatch_functor { } break; case aggregation::MEDIAN: { auto sorted_indices = sorted_order(table_view{{col}}, {}, {null_order::AFTER}, stream); - auto valid_sorted_indices = split(*sorted_indices, {col.size() - col.null_count()})[0]; + auto valid_sorted_indices = + split(*sorted_indices, {col.size() - col.null_count()}, stream)[0]; auto col_ptr = quantile(col, {0.5}, interpolation::LINEAR, valid_sorted_indices, true, stream); return get_element(*col_ptr, 0, stream, mr); @@ -78,7 +79,8 @@ struct reduce_dispatch_functor { CUDF_EXPECTS(quantile_agg->_quantiles.size() == 1, "Reduction quantile accepts only one quantile value"); auto sorted_indices = sorted_order(table_view{{col}}, {}, {null_order::AFTER}, stream); - auto valid_sorted_indices = split(*sorted_indices, {col.size() - col.null_count()})[0]; + auto valid_sorted_indices = + split(*sorted_indices, {col.size() - col.null_count()}, stream)[0]; auto col_ptr = quantile(col, quantile_agg->_quantiles, diff --git a/cpp/src/strings/copying/copying.cu b/cpp/src/strings/copying/copying.cu index 24572576498..cd95ab65898 100644 --- a/cpp/src/strings/copying/copying.cu +++ b/cpp/src/strings/copying/copying.cu @@ -15,7 +15,7 @@ */ #include -#include +#include #include #include #include @@ -42,7 +42,9 @@ std::unique_ptr copy_slice(strings_column_view const& strings, // slice the offsets child column auto offsets_column = std::make_unique( - cudf::slice(strings.offsets(), {offsets_offset, offsets_offset + strings_count + 1}).front(), + cudf::detail::slice( + strings.offsets(), {offsets_offset, offsets_offset + strings_count + 1}, stream) + .front(), stream, mr); auto const chars_offset = @@ -61,7 +63,9 @@ std::unique_ptr copy_slice(strings_column_view const& strings, auto const data_size = cudf::detail::get_value(offsets_column->view(), strings_count, stream); auto chars_column = std::make_unique( - cudf::slice(strings.chars(), {chars_offset, chars_offset + data_size}).front(), stream, mr); + cudf::detail::slice(strings.chars(), {chars_offset, chars_offset + data_size}, stream).front(), + stream, + mr); // slice the null mask auto null_mask = cudf::detail::copy_bitmask( diff --git a/cpp/src/strings/copying/shift.cu b/cpp/src/strings/copying/shift.cu index b4219585b78..024c8d2924d 100644 --- a/cpp/src/strings/copying/shift.cu +++ b/cpp/src/strings/copying/shift.cu @@ -96,7 +96,9 @@ std::unique_ptr shift(strings_column_view const& input, // output offsets column is the same size as the input auto const input_offsets = - cudf::slice(input.offsets(), {input.offset(), input.offset() + input.size() + 1}).front(); + cudf::detail::slice( + input.offsets(), {input.offset(), input.offset() + input.size() + 1}, stream) + .front(); auto const offsets_size = input_offsets.size(); auto offsets_column = cudf::detail::allocate_like( input_offsets, offsets_size, mask_allocation_policy::NEVER, stream, mr); diff --git a/cpp/src/transpose/transpose.cu b/cpp/src/transpose/transpose.cu index 5bc2cb21ac7..d119bc36c73 100644 --- a/cpp/src/transpose/transpose.cu +++ b/cpp/src/transpose/transpose.cu @@ -49,7 +49,7 @@ std::pair, table_view> transpose(table_view const& input auto splits_iter = thrust::make_transform_iterator( one_iter, [width = input.num_columns()](size_type idx) { return idx * width; }); auto splits = std::vector(splits_iter, splits_iter + input.num_rows() - 1); - auto output_column_views = cudf::split(output_column->view(), splits); + auto output_column_views = split(output_column->view(), splits, stream); return std::make_pair(std::move(output_column), table_view(output_column_views)); } diff --git a/cpp/tests/quantiles/percentile_approx_test.cu b/cpp/tests/quantiles/percentile_approx_test.cu index 39f7cc593d6..2b19699d870 100644 --- a/cpp/tests/quantiles/percentile_approx_test.cu +++ b/cpp/tests/quantiles/percentile_approx_test.cu @@ -432,4 +432,4 @@ TEST_F(PercentileApproxTest, NullPercentiles) {{99, 99, 8, 8}, valids.begin()}}; CUDF_TEST_EXPECT_COLUMNS_EQUAL(*result, expected); -} \ No newline at end of file +}