From 5741f695b7be17bf2cf7a02db6d14d5d2c1f2383 Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Mon, 2 Nov 2020 16:36:22 -0800 Subject: [PATCH 01/10] [scatter] Initial impl: Supports lists of: 1. fixed width data types 2. strings 3. lists (of all of the above, and lists(of ...)) --- cpp/include/cudf/detail/scatter.cuh | 14 + cpp/include/cudf/lists/detail/scatter.cuh | 769 ++++++++++++++++++ cpp/include/cudf/lists/list_device_view.cuh | 144 ++++ .../cudf/lists/lists_column_device_view.cuh | 90 ++ cpp/tests/CMakeLists.txt | 1 + cpp/tests/copying/scatter_list_tests.cu | 554 +++++++++++++ 6 files changed, 1572 insertions(+) create mode 100644 cpp/include/cudf/lists/detail/scatter.cuh create mode 100644 cpp/include/cudf/lists/list_device_view.cuh create mode 100644 cpp/include/cudf/lists/lists_column_device_view.cuh create mode 100644 cpp/tests/copying/scatter_list_tests.cu diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index 0e30ce603cf..37b4259c3a6 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -25,6 +25,7 @@ #include #include #include +#include #include #include @@ -118,6 +119,19 @@ struct column_scatterer_impl { } }; +template +struct column_scatterer_impl { + std::unique_ptr operator()(column_view const& source, + MapIterator scatter_map_begin, + MapIterator scatter_map_end, + column_view const& target, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) const + { + return cudf::lists::detail::scatter(source, scatter_map_begin, scatter_map_end, target, mr, stream); + } +}; + template struct column_scatterer { template diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh new file mode 100644 index 00000000000..b0b9a411717 --- /dev/null +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -0,0 +1,769 @@ +/* + * Copyright (c) 2020, 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 +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace cudf { +namespace lists { +namespace detail { + +namespace { + +/** + * @brief Holder for a list row's positional information, without + * also holding a reference to the list column. + * + * Analogous to the list_view, this class is default constructable, + * and can thus be stored in rmm::device_vector. It is used to represent + * the results of a `scatter()` operation; a device_vector may hold + * several instances of unbound_list_view, each with a flag indicating + * whether it came from the scatter source or target. Each instance + * may later be "bound" to the appropriate source/target column, to + * reconstruct the list_view. + */ +struct unbound_list_view +{ + /** + * @brief Flag type, indicating whether this list row originated from + * the source or target column, in `scatter()`. + */ + enum label_t : bool {SOURCE, TARGET}; + + using lists_column_device_view = cudf::detail::lists_column_device_view; + using list_device_view = cudf::list_device_view; + + unbound_list_view() = default; + unbound_list_view(unbound_list_view const&) = default; + unbound_list_view(unbound_list_view &&) = default; + unbound_list_view& operator = (unbound_list_view const&) = default; + unbound_list_view& operator = (unbound_list_view &&) = default; + + /** + * @brief (__device__) Constructor, for use from `scatter()`. + * + * @param scatter_source_label Whether the row came from source or target + * @param lists_column The actual source/target lists column + * @param row_index Index of the row in lists_column that this instance represents + */ + CUDA_DEVICE_CALLABLE unbound_list_view(label_t scatter_source_label, + cudf::detail::lists_column_device_view const& lists_column, + size_type const& row_index) + : _label{scatter_source_label}, + _row_index{row_index} + { + _size = list_device_view{lists_column, row_index}.size(); + } + + /** + * @brief (__device__) Constructor, for use when constructing the child column + * of a scattered list column + * + * @param scatter_source_label Whether the row came from source or target + * @param row_index Index of the row that this instance represents in the source/target column + * @param size The number of elements in this list row + */ + CUDA_DEVICE_CALLABLE unbound_list_view(label_t scatter_source_label, + size_type const& row_index, + size_type const& size) + : _label{scatter_source_label}, + _row_index{row_index}, + _size{size} + {} + + /** + * @brief Returns number of elements in this list-row. + */ + CUDA_DEVICE_CALLABLE size_type size() const { return _size; } + + /** + * @brief Returns whether this row came from the `scatter()` source or target + */ + CUDA_DEVICE_CALLABLE label_t label() const { return _label; } + + /** + * @brief Returns the index in the source/target column + */ + CUDA_DEVICE_CALLABLE size_type row_index() const { return _row_index; } + + /** + * @brief Binds to source/target column (depending on SOURCE/TARGET labels), + * to produce a bound list_view. + * + * @param scatter_source Source column for the scatter operation + * @param scatter_target Target column for the scatter operation + * @return A (bound) list_view for the row that this object represents + */ + CUDA_DEVICE_CALLABLE list_device_view bind_to_column( + lists_column_device_view const& scatter_source, + lists_column_device_view const& scatter_target) const + { + return list_device_view(_label == SOURCE? scatter_source : scatter_target, _row_index); + } + + private: + + // Note: Cannot store reference to list column, because of storage in device_vector. + // Only keep track of whether this list row came from the source or target of scatter. + + label_t _label {SOURCE}; // Whether this list row came from the scatter source or target. + size_type _row_index{}; // Row index in the Lists column. + size_type _size{}; // Number of elements in *this* list row. +}; + +rmm::device_vector list_vector_from_column( + unbound_list_view::label_t label, + cudf::detail::lists_column_device_view const& lists_column, + cudaStream_t stream +) +{ + auto n_rows = lists_column.size(); + + auto vector = rmm::device_vector(n_rows); + + thrust::for_each_n( + rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + n_rows, + [ + label, + lists_column, + output = vector.data().get() + ] __device__ (size_type row_index) + { + output[row_index] = unbound_list_view{label, lists_column, row_index}; + } + ); + + return vector; +} + +/** + * @brief Utility function to fetch the number of rows in a lists column's + * child column, given its offsets column. + * (This is simply the last value in the offsets column.) + * + * @param list_offsets Offsets child of a lists column + * @param stream The cuda-stream to synchronize on, when reading from device memory + * @return int32_t The last element in the list_offsets column, indicating + * the number of rows in the lists-column's child. + */ +int32_t get_num_child_rows(cudf::column_view const& list_offsets, cudaStream_t stream) +{ + // Number of rows in child-column == last offset value. + int32_t num_child_rows{}; + CUDA_TRY(cudaMemcpyAsync(&num_child_rows, + list_offsets.data()+list_offsets.size()-1, + sizeof(int32_t), + cudaMemcpyDeviceToHost, + stream)); + CUDA_TRY(cudaStreamSynchronize(stream)); + return num_child_rows; +} + +/** + * @brief Constructs null mask for a scattered list's child column + * + * @param parent_list_vector Vector of unbound_list_view, for parent lists column + * @param parent_list_offsets List column offsets for parent lists column + * @param source_lists Source lists column for scatter operation + * @param target_lists Target lists column for scatter operation + * @param num_child_rows Number of rows in child column + * @param mr Device memory resource used to allocate child column's null mask + * @param stream CUDA stream used for device memory operations and kernel launches + * @return std::pair Child column's null mask and null row count + */ +std::pair +construct_child_nullmask(rmm::device_vector const& parent_list_vector, + column_view const& parent_list_offsets, + cudf::detail::lists_column_device_view const& source_lists, + cudf::detail::lists_column_device_view const& target_lists, + size_type num_child_rows, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) +{ + auto is_valid_predicate = [ + d_list_vector = parent_list_vector.data().get(), + d_offsets = parent_list_offsets.template data(), + d_offsets_size = parent_list_offsets.size(), + source_lists, + target_lists + ] __device__ (auto const& i) { + + auto list_start = thrust::upper_bound(thrust::seq, + d_offsets, + d_offsets + d_offsets_size, + i) - 1; + auto list_index = list_start - d_offsets; + auto element_index = i - *list_start; + + auto list_row = d_list_vector[list_index]; + return !list_row.bind_to_column(source_lists, target_lists).is_null(element_index); + }; + + return cudf::detail::valid_if( + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(num_child_rows), + is_valid_predicate, + stream, + mr + ); +} + +#ifndef NDEBUG +void print(std::string const& msg, column_view const& col, cudaStream_t stream) +{ + if (col.type().id() != type_id::INT32) + { + std::cout << "[Cannot print non-INT32 column.]" << std::endl; + return; + } + + std::cout << msg << " = ["; + thrust::for_each_n( + rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + col.size(), + [c = col.template data()]__device__(auto const& i) { + printf("%d,", c[i]); + } + ); + std::cout << "]" << std::endl; +} + +void print(std::string const& msg, rmm::device_vector const& scatter, cudaStream_t stream) +{ + std::cout << msg << " == ["; + + thrust::for_each_n( + rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + scatter.size(), + [s = scatter.data().get()] __device__ (auto const& i) { + auto si = s[i]; + printf("%s[%d](%d), ", (si.label() == unbound_list_view::SOURCE? "S":"T"), si.row_index(), si.size()); + } + ); + std::cout << "]" << std::endl; +} +#endif // NDEBUG + +/** + * @brief (type_dispatch endpoint) Functor that constructs the child column result + * of `scatter()`ing a list column. + * + * The protocol is as follows: + * + * Inputs: + * 1. list_vector: A device_vector of unbound_list_view, with each element + * indicating the position, size, and which column the list + * row came from. + * 2. list_offsets: The offsets column for the (outer) lists column, each offset + * marking the beginning of a list row. + * 3. source_list: The lists-column that is the source of the scatter(). + * 4. target_list: The lists-column that is the target of the scatter(). + * + * Output: A (possibly non-list) child column, which may be used in combination + * with list_offsets to fully construct the outer list. + * + * Example: + * + * Consider the following scatter operation of two `list` columns: + * + * 1. Source: [{9,9,9,9}, {8,8,8}], i.e. + * a. Child: [9,9,9,9,8,8,8] + * b. Offsets: [0, 4, 7] + * + * 2. Target: [{1,1}, {2,2}, {3,3}], i.e. + * a. Child: [1,1,2,2,3,3] + * b. Offsets: [0, 2, 4, 6] + * + * 3. Scatter-map: [2, 0] + * + * 4. Expected output: [{8,8,8}, {2,2}, {9,9,9,9}], i.e. + * a. Child: [8,8,8,2,2,9,9,9,9] <--- THIS + * b. Offsets: [0, 3, 5, 9] + * + * It is the Expected Child column above that list_child_constructor attempts + * to construct. + * + * `list_child_constructor` expects to be called with the `Source`/`Target` + * lists columns, along with the following: + * + * 1. list_vector: [ S[1](3), T[1](2), S[0](4) ] + * Each unbound_list_view (e.g. S[1](3)) indicates: + * a. Which column the row is bound to: S == Source, T == Target + * b. The list index. E.g. S[1] indicates the 2nd list row of the Source column. + * c. The row size. E.g. S[1](3) indicates that the row has 3 elements. + * + * 2. list_offsets: [0, 3, 5, 9] + * The caller may construct this with an `inclusive_scan()` on `list_vector` + * element sizes. + */ +struct list_child_constructor +{ + private: + /** + * @brief Function to determine what types are supported as child column types, + * when scattering lists. + * + * @tparam T The data type of the child column of the list being scattered. + */ + template + struct is_supported_child_type + { + static const bool value = cudf::is_fixed_width() + || std::is_same::value + || std::is_same::value; + }; + + public: + + /** + * @brief SFINAE catch-all, for unsupported child column types. + */ + template + std::enable_if_t::value, std::unique_ptr> operator()( + rmm::device_vector const& list_vector, + cudf::column_view const& list_offsets, + cudf::lists_column_view const& source_list, + cudf::lists_column_view const& target_list, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) const + { + CUDF_FAIL("list_child_constructor unsupported!"); + } + + /** + * @brief Implementation for fixed_width child column types. + */ + template + std::enable_if_t(), std::unique_ptr> operator()( + rmm::device_vector const& list_vector, + cudf::column_view const& list_offsets, + cudf::lists_column_view const& source_lists_column_view, + cudf::lists_column_view const& target_lists_column_view, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) const + { + auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); + auto target_column_device_view = column_device_view::create(target_lists_column_view.parent(), stream); + auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view); + auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view); + + // Number of rows in child-column == last offset value. + int32_t num_child_rows{get_num_child_rows(list_offsets, stream)}; + + auto child_null_mask = + source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() + ? construct_child_nullmask(list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + : std::make_pair(rmm::device_buffer{}, 0); + +#ifndef NDEBUG + print("list_offsets ", list_offsets, stream); + print("source_lists.child() ", source_lists_column_view.child(), stream); + print("source_lists.offsets() ", source_lists_column_view.offsets(), stream); + print("target_lists.child() ", target_lists_column_view.child(), stream); + print("target_lists.offsets() ", target_lists_column_view.offsets(), stream); + print("scatter_rows ", list_vector, stream); +#endif // NDEBUG + + // Init child-column. + auto child_column = cudf::make_fixed_width_column( + cudf::data_type{cudf::type_to_id()}, + num_child_rows, + child_null_mask.first, + child_null_mask.second, + stream, + mr + ); + + // Function to copy child-values for specified index of unbound_list_view + // to the child column. + auto copy_child_values_for_list_index = [ + d_scattered_lists = list_vector.data().get(), // unbound_list_view* + d_child_column = child_column->mutable_view().data(), + d_offsets = list_offsets.template data(), + source_lists, + target_lists + ] __device__ (auto const& row_index) { + + auto unbound_list_row = d_scattered_lists[row_index]; + auto actual_list_row = unbound_list_row.bind_to_column(source_lists, target_lists); + auto const& bound_column= (unbound_list_row.label() == unbound_list_view::SOURCE? source_lists : target_lists); + auto list_begin_offset = bound_column.offsets().element(unbound_list_row.row_index()); + auto list_end_offset = bound_column.offsets().element(unbound_list_row.row_index()+1); + +#ifndef NDEBUG + printf("%d: Unbound == %s[%d](%d), Bound size == %d, calc_begin==%d, calc_end=%d, calc_size=%d\n", + row_index, + (unbound_list_row.label() == unbound_list_view::SOURCE? "S":"T"), + unbound_list_row.row_index(), + unbound_list_row.size(), + actual_list_row.size(), + list_begin_offset, + list_end_offset, + list_end_offset-list_begin_offset + ); +#endif // NDEBUG + + // Copy all elements in this list row, to "appropriate" offset in child-column. + auto destination_start_offset = d_offsets[row_index]; + thrust::for_each_n( + thrust::seq, + thrust::make_counting_iterator(0), + actual_list_row.size(), + [actual_list_row, d_child_column, destination_start_offset] __device__ (auto const& list_element_index) + { + d_child_column[destination_start_offset + list_element_index] = + actual_list_row.template element(list_element_index); + } + ); + }; + + // For each list-row, copy underlying elements to the child column. + thrust::for_each_n( + rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + list_vector.size(), + copy_child_values_for_list_index + ); + + return std::make_unique(child_column->view()); + } + + /** + * @brief Implementation for list child columns that contain strings. + */ + template + std::enable_if_t::value, + std::unique_ptr> + operator()( + rmm::device_vector const& list_vector, + cudf::column_view const& list_offsets, + cudf::lists_column_view const& source_lists_column_view, + cudf::lists_column_view const& target_lists_column_view, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) const + { + auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); + auto target_column_device_view = column_device_view::create(target_lists_column_view.parent(), stream); + auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view); + auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view); + + int32_t num_child_rows{get_num_child_rows(list_offsets, stream)}; + + auto string_views = rmm::device_vector(num_child_rows); + + auto populate_string_views = [ + d_scattered_lists = list_vector.data().get(), // unbound_list_view* + d_list_offsets = list_offsets.template data(), + d_string_views = string_views.data().get(), + source_lists, + target_lists + ] __device__ (auto const& row_index) { + + auto unbound_list_view = d_scattered_lists[row_index]; + auto actual_list_row = unbound_list_view.bind_to_column(source_lists, target_lists); + auto lists_column = actual_list_row.get_column(); + auto lists_offsets_column = lists_column.offsets(); + auto child_strings_column = lists_column.child(); + auto string_offsets_column = child_strings_column.child(cudf::strings_column_view::offsets_column_index); + auto string_chars_column = child_strings_column.child(cudf::strings_column_view::chars_column_index); + + auto output_start_offset = d_list_offsets[row_index]; // Offset in `string_views` at which string_views are + // to be written for this list row_index. + auto input_list_start = lists_offsets_column.template element(unbound_list_view.row_index()); + + thrust::for_each_n( + thrust::seq, + thrust::make_counting_iterator(0), + actual_list_row.size(), + [ + output_start_offset, + d_string_views, + input_list_start, + d_string_offsets = string_offsets_column.template data(), + d_string_chars = string_chars_column.template data() + ] __device__ (auto const& string_idx) + { + // auto string_offset = output_start_offset + string_idx; + auto string_start_idx = d_string_offsets[input_list_start + string_idx]; + auto string_end_idx = d_string_offsets[input_list_start + string_idx + 1]; + + d_string_views[output_start_offset + string_idx] = + string_view{d_string_chars + string_start_idx, string_end_idx - string_start_idx}; + } + ); + }; + + thrust::for_each_n( + rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + list_vector.size(), + populate_string_views + ); + + // string_views should now have been populated with source and target references. + + auto string_offsets = cudf::strings::detail::child_offsets_from_string_vector(string_views, mr, stream); + auto string_chars = cudf::strings::detail::child_chars_from_string_vector(string_views, string_offsets->view().data(), 0, mr, stream); + auto child_null_mask = + source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() + ? construct_child_nullmask(list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + : std::make_pair(rmm::device_buffer{}, 0); + + return cudf::make_strings_column(num_child_rows, + std::move(string_offsets), + std::move(string_chars), + child_null_mask.second, // Null count. + std::move(child_null_mask.first), // Null mask. + stream, + mr); + } + + /** + * @brief (Recursively) Constructs a child column that is itself a list column. + */ + template + std::enable_if_t::value, + std::unique_ptr> + operator() ( + rmm::device_vector const& list_vector, + cudf::column_view const& list_offsets, + cudf::lists_column_view const& source_lists_column_view, + cudf::lists_column_view const& target_lists_column_view, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) const + { + auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); + auto target_column_device_view = column_device_view::create(target_lists_column_view.parent(), stream); + auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view); + auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view); + + auto num_child_rows = get_num_child_rows(list_offsets, stream); + + auto child_list_views = rmm::device_vector(num_child_rows); + + // Function to convert from parent list_device_view instances to child list_device_views. + // For instance, if a parent list_device_view has 3 elements, it should have 3 corresponding + // child list_device_view instances. + auto populate_child_list_views = [ + d_scattered_lists = list_vector.data().get(), + d_list_offsets = list_offsets.template data(), + d_child_list_views = child_list_views.data().get(), + source_lists, + target_lists + ] __device__ (auto const& row_index) { + + auto scattered_row = d_scattered_lists[row_index]; + auto label = scattered_row.label(); + auto bound_list_row = scattered_row.bind_to_column(source_lists, target_lists); + auto lists_offsets_column = bound_list_row.get_column().offsets(); + + auto child_column = bound_list_row.get_column().child(); + auto child_offsets = child_column.child(cudf::lists_column_view::offsets_column_index); + + // For lists row at row_index, + // 1. Number of entries in child_list_views == bound_list_row.size(). + // 2. Offset of the first child list_view == d_list_offsets[row_index]. + auto output_start_offset = d_list_offsets[row_index]; + auto input_list_start = lists_offsets_column.template element(scattered_row.row_index()); + + thrust::for_each_n( + thrust::seq, + thrust::make_counting_iterator(0), + bound_list_row.size(), + [ + input_list_start, + output_start_offset, + label, + d_child_list_views, + d_child_offsets = child_offsets.template data() + ] __device__ (auto const& child_list_index) + { + auto child_start_idx = d_child_offsets[input_list_start + child_list_index]; + auto child_end_idx = d_child_offsets[input_list_start + child_list_index + 1]; + + d_child_list_views[output_start_offset + child_list_index] = + unbound_list_view{label, input_list_start + child_list_index, child_end_idx - child_start_idx}; + } + ); + }; + + thrust::for_each_n( + rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + list_vector.size(), + populate_child_list_views + ); + + // child_list_views should now have been populated, with source and target references. + + auto begin = thrust::make_transform_iterator( + child_list_views.begin(), + [] __device__ (auto const& row) { return row.size(); } + ); + + auto child_offsets = cudf::strings::detail::make_offsets_child_column( + begin, + begin + child_list_views.size(), + mr, + stream + ); + + auto child_column = cudf::type_dispatcher( + source_lists_column_view.child().child(1).type(), + list_child_constructor{}, + child_list_views, + child_offsets->view(), + cudf::lists_column_view(source_lists_column_view.child()), + cudf::lists_column_view(target_lists_column_view.child()), + mr, + stream + ); + + auto child_null_mask = + source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() + ? construct_child_nullmask(list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + : std::make_pair(rmm::device_buffer{}, 0); + + return cudf::make_lists_column( + num_child_rows, + std::move(child_offsets), + std::move(child_column), + child_null_mask.second, // Null count + std::move(child_null_mask.first), // Null mask + stream, + mr + ); + } +}; + +} // namespace; + +/** + * @brief Scatters lists into a copy of the target column + * according to a scatter map. + * + * The scatter is performed according to the scatter iterator such that row + * `scatter_map[i]` of the output column is replaced by the source list-row. + * All other rows of the output column equal corresponding rows of the target table. + * + * If the same index appears more than once in the scatter map, the result is + * undefined. + * + * The caller must update the null mask in the output column. + * + * @tparam SourceIterator must produce list_view objects + * @tparam MapIterator must produce index values within the target column. + * + * @param mr Device memory resource used to allocate the returned column's device memory + * @param stream CUDA stream used for device memory operations and kernel launches. + * @return New lists column. + */ +template +std::unique_ptr scatter( + column_view const& source, + MapIterator scatter_map_begin, + MapIterator scatter_map_end, + column_view const& target, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource(), + cudaStream_t stream = 0 + ) +{ + auto num_rows = target.size(); + + if (num_rows == 0) + { + return cudf::empty_like(target); + } + + auto child_column_type = lists_column_view(target).child().type(); + + // TODO: Deep(er) checks that source and target have identical types. + + using lists_column_device_view = cudf::detail::lists_column_device_view; + using unbound_list_view = cudf::lists::detail::unbound_list_view; + + auto source_lists_column_view = lists_column_view(source); // Checks that this is a list column. + auto source_device_view = column_device_view::create(source, stream); + // auto source_lists_column_device_view = lists_column_device_view(*source_device_view); + auto source_vector = list_vector_from_column(unbound_list_view::SOURCE, lists_column_device_view(*source_device_view), stream); + + auto target_lists_column_view = lists_column_view(target); // Checks that target is a list column. + auto target_device_view = column_device_view::create(target, stream); + // auto target_lists_column_device_view = lists_column_device_view(*target_device_view); + auto target_vector = list_vector_from_column(unbound_list_view::TARGET, lists_column_device_view(*target_device_view), stream); + + // Scatter. + thrust::scatter( + rmm::exec_policy(stream)->on(stream), + source_vector.begin(), + source_vector.end(), + scatter_map_begin, + target_vector.begin() + ); + + auto list_size_begin = thrust::make_transform_iterator(target_vector.begin(), [] __device__(unbound_list_view l) { return l.size(); }); + auto offsets_column = cudf::strings::detail::make_offsets_child_column( + list_size_begin, + list_size_begin + target.size(), + mr, + stream + ); + + auto child_column = cudf::type_dispatcher( + child_column_type, + list_child_constructor{}, + target_vector, + offsets_column->view(), + source_lists_column_view, + target_lists_column_view, + mr, + stream + ); + + rmm::device_buffer null_mask{0, stream, mr}; + if (target.has_nulls()) { + null_mask = copy_bitmask(target, stream, mr); + } + + return cudf::make_lists_column( + num_rows, + std::move(offsets_column), + std::move(child_column), + cudf::UNKNOWN_NULL_COUNT, + std::move(null_mask), + stream, + mr + ); +} + +} // namespace detail; +} // namespace lists; +} // namespace cudf; diff --git a/cpp/include/cudf/lists/list_device_view.cuh b/cpp/include/cudf/lists/list_device_view.cuh new file mode 100644 index 00000000000..f3b7d2e7ebe --- /dev/null +++ b/cpp/include/cudf/lists/list_device_view.cuh @@ -0,0 +1,144 @@ +/* + * Copyright (c) 2020, 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 +#include +#include + +namespace cudf { + +/** + * @brief A non-owning, immutable view of device data that represents + * a list of elements of arbitrary type (including further nested lists). + * + */ +class list_device_view { + + using lists_column_device_view = cudf::detail::lists_column_device_view; + + public: + + list_device_view() = default; + + CUDA_DEVICE_CALLABLE list_device_view(lists_column_device_view const& lists_column, size_type const& idx); + + ~list_device_view() = default; + + /** + * @brief Fetches the offset in the list column's child that corresponds to + * the element at the specified list index. + * + * Consider the following lists column: + * [ + * [0,1,2], + * [3,4,5], + * [6,7,8] + * ] + * + * The list's internals would look like: + * offsets: [0, 3, 6, 9] + * child : [0, 1, 2, 3, 4, 5, 6, 7, 8] + * + * The second list row (i.e. row_index=1) is [3,4,5]. + * The third element (i.e. idx=2) of the second list row is 5. + * + * The offset of this element as stored in the child column (i.e. 5) + * may be fetched using this method. + */ + CUDA_DEVICE_CALLABLE size_type element_offset(size_type idx) const; + + /** + * @brief Fetches the element at the specified index, within the list row. + * + * @tparam The type of the list's element. + * @param The index into the list row + * @return The element at the specified index of the list row. + */ + template + CUDA_DEVICE_CALLABLE T element(size_type idx) const; + + /** + * @brief Checks whether element is null at specified index in the list row. + */ + CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const; + + /** + * @brief Checks whether this list row is null. + */ + CUDA_DEVICE_CALLABLE bool is_null() const; + + /** + * @brief Fetches the number of elements in this list row. + */ + CUDA_DEVICE_CALLABLE size_type size() const { return _size; } + + /** + * @brief Fetches the lists_column_device_view that contains this list. + */ + CUDA_DEVICE_CALLABLE lists_column_device_view const& get_column() const { return lists_column; } + + private: + + lists_column_device_view const& lists_column; + size_type _row_index{}; // Row index in the Lists column vector. + size_type _size{}; // Number of elements in *this* list row. + + size_type begin_offset; // Offset in list_column_device_view where this list begins. + +}; + +CUDA_DEVICE_CALLABLE list_device_view::list_device_view( + lists_column_device_view const& lists_column, size_type const& row_index) + : lists_column(lists_column), _row_index(row_index) +{ + release_assert(row_index >= 0 && row_index < lists_column.size() && "row_index out of bounds"); + + column_device_view const& offsets = lists_column.offsets(); + release_assert(row_index < offsets.size() && "row_index should not have exceeded offset size"); + + begin_offset = offsets.element(row_index); + release_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() && + "begin_offset out of bounds."); + _size = offsets.element(row_index + 1) - begin_offset; +} + +CUDA_DEVICE_CALLABLE size_type list_device_view::element_offset(size_type idx) const +{ + release_assert(idx >= 0 && idx < size() && "idx out of bounds"); + // release_assert(!is_null() && !is_null(idx) && "Cannot read null element."); + return begin_offset + idx; +} + +template +CUDA_DEVICE_CALLABLE T list_device_view::element(size_type idx) const +{ + return lists_column.child().element(element_offset(idx)); +} + +CUDA_DEVICE_CALLABLE bool list_device_view::is_null(size_type idx) const +{ + release_assert(idx >= 0 && idx < size() && "Index out of bounds."); + auto element_offset = begin_offset + idx; + return lists_column.child().is_null(element_offset); +} + +CUDA_DEVICE_CALLABLE bool list_device_view::is_null() const +{ + return lists_column.is_null(_row_index); +} + +} // namespace cudf diff --git a/cpp/include/cudf/lists/lists_column_device_view.cuh b/cpp/include/cudf/lists/lists_column_device_view.cuh new file mode 100644 index 00000000000..5244140a783 --- /dev/null +++ b/cpp/include/cudf/lists/lists_column_device_view.cuh @@ -0,0 +1,90 @@ +/* + * Copyright (c) 2020, 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 +#include +#include +#include + +namespace cudf { + +namespace detail { + +/** + * @brief Given a column-device-view, an instance of this class provides a + * wrapper on this compound column for list operations. + * Analogous to list_column_view. + */ +class lists_column_device_view { + + public: + + ~lists_column_device_view() = default; + lists_column_device_view(lists_column_device_view const&) = default; + lists_column_device_view(lists_column_device_view&&) = default; + lists_column_device_view& operator= (lists_column_device_view const&) = default; + lists_column_device_view& operator= (lists_column_device_view &&) = default; + + lists_column_device_view(column_device_view const& underlying_) + : underlying(underlying_) + { + CUDF_EXPECTS(underlying_.type().id() == type_id::LIST, "lists_column_device_view only supports lists"); + } + + /** + * @brief Fetches number of rows in the lists column + */ + CUDA_HOST_DEVICE_CALLABLE cudf::size_type size() const + { + return underlying.size(); + } + + /** + * @brief Fetches the offsets column of the underlying list column. + */ + CUDA_DEVICE_CALLABLE column_device_view offsets() const + { + return underlying.child(lists_column_view::offsets_column_index); + } + + /** + * @brief Fetches the child column of the underlying list column. + */ + CUDA_DEVICE_CALLABLE column_device_view child() const + { + return underlying.child(lists_column_view::child_column_index); + } + + /** + * @brief Indicates whether the list column is nullable. + */ + CUDA_DEVICE_CALLABLE bool nullable() const { return underlying.nullable(); } + + /** + * @brief Indicates whether the row (i.e. list) at the specified + * index is null. + */ + CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const { return underlying.is_null(idx); } + + private: + + column_device_view underlying; +}; + +} // namespace detail + +} // namespace cudf; diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 954411b4716..4cab4277bae 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -378,6 +378,7 @@ set(COPYING_TEST_SRC "${CMAKE_CURRENT_SOURCE_DIR}/copying/gather_struct_tests.cu" "${CMAKE_CURRENT_SOURCE_DIR}/copying/detail_gather_tests.cu" "${CMAKE_CURRENT_SOURCE_DIR}/copying/scatter_tests.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/copying/scatter_list_tests.cu" "${CMAKE_CURRENT_SOURCE_DIR}/copying/copy_range_tests.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/copying/slice_tests.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/copying/split_tests.cpp" diff --git a/cpp/tests/copying/scatter_list_tests.cu b/cpp/tests/copying/scatter_list_tests.cu new file mode 100644 index 00000000000..8a86d2fd3ca --- /dev/null +++ b/cpp/tests/copying/scatter_list_tests.cu @@ -0,0 +1,554 @@ +/* + * Copyright (c) 2020, 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 +#include +#include +#include + +#include + +template +class TypedScatterListsTest : public cudf::test::BaseFixture { +}; +using FixedWidthTypes= cudf::test::Concat; +TYPED_TEST_CASE(TypedScatterListsTest, FixedWidthTypes); + +class ScatterListsTest : public cudf::test::BaseFixture { +}; + +TYPED_TEST(TypedScatterListsTest, ListsOfFixedWidth) +{ + using namespace cudf::test; + using T = TypeParam; + + auto src_list_column = lists_column_wrapper{ + {9, 9, 9, 9}, {8, 8, 8} + }; + + auto target_list_column = lists_column_wrapper{ + {0,0}, {1,1}, {2,2}, {3,3}, {4,4}, {5,5}, {6,6} + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), + scatter_map, + cudf::table_view({target_list_column})); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT ( + ret->get_column(0), + lists_column_wrapper {{8,8,8}, {1,1}, {9,9,9,9}, {3,3}, {4,4}, {5,5}, {6,6}} + ); +} + +TYPED_TEST(TypedScatterListsTest, EmptyListsOfFixedWidth) +{ + using namespace cudf::test; + using T = TypeParam; + + auto src_child = fixed_width_column_wrapper { + {9, 9, 9, 9, 8, 8, 8}, + }; + + // One null list row, and one row with nulls. + auto src_list_column = cudf::make_lists_column( + 3, + fixed_width_column_wrapper{0, 4, 7, 7}.release(), + src_child.release(), + 0, + {} + ); + + auto target_list_column = lists_column_wrapper{ + {0,0}, {1,1}, {2,2}, {3,3}, {4,4}, {5,5}, {6,6} + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_child_ints = fixed_width_column_wrapper { + {8,8,8, 1,1, 9,9,9,9, 3,3, 4,4, 6,6 } + }; + auto expected_lists_column = cudf::make_lists_column( + 7, + fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), + expected_child_ints.release(), + 0, + {} + ); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + expected_lists_column->view(), + ret->get_column(0) + ); +} + +TYPED_TEST(TypedScatterListsTest, EmptyListsOfNullableFixedWidth) +{ + using namespace cudf::test; + using T = TypeParam; + + auto src_child = fixed_width_column_wrapper { + {9, 9, 9, 9, 8, 8, 8}, + {1, 1, 1, 0, 1, 1, 1} + }; + + // One null list row, and one row with nulls. + auto src_list_column = cudf::make_lists_column( + 3, + fixed_width_column_wrapper{0, 4, 7, 7}.release(), + src_child.release(), + 0, + {} + ); + + auto target_list_column = lists_column_wrapper{ + {0,0}, {1,1}, {2,2}, {3,3}, {4,4}, {5,5}, {6,6} + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_child_ints = fixed_width_column_wrapper { + {8,8,8, 1,1, 9,9,9,9, 3,3, 4,4, 6,6 }, + {1,1,1, 1,1, 1,1,1,0, 1,1, 1,1, 1,1 } + }; + auto expected_lists_column = cudf::make_lists_column( + 7, + fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), + expected_child_ints.release(), + 0, + {} + ); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + expected_lists_column->view(), + ret->get_column(0) + ); +} + +TYPED_TEST(TypedScatterListsTest, NullableListsOfNullableFixedWidth) +{ + using namespace cudf::test; + using T = TypeParam; + + auto src_child = fixed_width_column_wrapper { + {9, 9, 9, 9, 8, 8, 8}, + {1, 1, 1, 0, 1, 1, 1} + }; + + auto src_list_validity = make_counting_transform_iterator(0, [](auto i) { return i != 2; }); + // One null list row, and one row with nulls. + auto src_list_column = cudf::make_lists_column( + 3, + fixed_width_column_wrapper{0, 4, 7, 7}.release(), + src_child.release(), + 1, + detail::make_null_mask(src_list_validity, src_list_validity + 3) + ); + + auto target_list_column = lists_column_wrapper{ + {0,0}, {1,1}, {2,2}, {3,3}, {4,4}, {5,5}, {6,6} + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_child_ints = fixed_width_column_wrapper { + {8,8,8, 1,1, 9,9,9,9, 3,3, 4,4, 6,6 }, + {1,1,1, 1,1, 1,1,1,0, 1,1, 1,1, 1,1 } + }; + + auto expected_validity = make_counting_transform_iterator(0, [](auto i) { return i != 5; }); + auto expected_lists_column = cudf::make_lists_column( + 7, + fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), + expected_child_ints.release(), + 1, + detail::make_null_mask(expected_validity, expected_validity + 7) + ); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + expected_lists_column->view(), + ret->get_column(0) + ); +} + +TEST_F(ScatterListsTest, ListsOfStrings) +{ + using namespace cudf::test; + + auto src_list_column = lists_column_wrapper { + {"all", "the", "leaves", "are", "brown"}, + {"california", "dreaming"} + }; + + auto target_list_column = lists_column_wrapper { + {"zero"}, + {"one", "one"}, + {"two", "two"}, + {"three", "three", "three"}, + {"four", "four", "four", "four"} + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), + scatter_map, + cudf::table_view({target_list_column}) + ); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + lists_column_wrapper{ + {"california", "dreaming"}, + {"one", "one"}, + {"all", "the", "leaves", "are", "brown"}, + {"three", "three", "three"}, + {"four", "four", "four", "four"} + }, + ret->get_column(0) + ); +} + +TEST_F(ScatterListsTest, ListsOfNullableStrings) +{ + using namespace cudf::test; + + auto src_strings_column = strings_column_wrapper{ + {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, + { 1, 1, 1, 0, 1, 0, 1} + }; + + auto src_list_column = cudf::make_lists_column( + 2, + fixed_width_column_wrapper{0, 5, 7}.release(), + src_strings_column.release(), + 0, + {} + ); + + auto target_list_column = lists_column_wrapper { + {"zero"}, + {"one", "one"}, + {"two", "two"}, + {"three", "three"}, + {"four", "four"}, + {"five", "five"} + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column}) + ); + + auto expected_strings = strings_column_wrapper { + {"california", "dreaming", "one", "one", "all", "the", "leaves", "are", "brown", + "three", "three", "four", "four", "five", "five"}, + make_counting_transform_iterator(0, [](auto i) {return i!=0 && i!=7;}) + }; + + auto expected_lists = cudf::make_lists_column( + 6, + fixed_width_column_wrapper{0,2,4,9,11,13,15}.release(), + expected_strings.release(), + 0, + {} + ); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + expected_lists->view(), + ret->get_column(0) + ); +} + +TEST_F(ScatterListsTest, EmptyListsOfNullableStrings) +{ + using namespace cudf::test; + + auto src_strings_column = strings_column_wrapper{ + {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, + { 1, 1, 1, 0, 1, 0, 1} + }; + + auto src_list_column = cudf::make_lists_column( + 3, + fixed_width_column_wrapper{0, 5, 5, 7}.release(), + src_strings_column.release(), + 0, + {} + ); + + auto target_list_column = lists_column_wrapper { + {"zero"}, + {"one", "one"}, + {"two", "two"}, + {"three", "three"}, + {"four", "four"}, + {"five", "five"} + }; + + auto scatter_map = fixed_width_column_wrapper{2, 4, 0}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column}) + ); + + auto expected_strings = strings_column_wrapper { + {"california", "dreaming", + "one", "one", + "all", "the", "leaves", "are", "brown", + "three", "three", + "five", "five"}, + make_counting_transform_iterator(0, [](auto i) {return i!=0 && i!=7;}) + }; + + auto expected_lists = cudf::make_lists_column( + 6, + fixed_width_column_wrapper{0,2,4,9,11,11,13}.release(), + expected_strings.release(), + 0, + {} + ); + + std::cout << "Expected: " << std::endl; print(expected_lists->view()); + std::cout << "Received: " << std::endl; print(ret->get_column(0)); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + expected_lists->view(), + ret->get_column(0) + ); +} + +TEST_F(ScatterListsTest, NullableListsOfNullableStrings) +{ + using namespace cudf::test; + + auto src_strings_column = strings_column_wrapper{ + {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, + { 1, 1, 1, 0, 1, 0, 1} + }; + + auto src_validity = make_counting_transform_iterator(0, [](auto i) { return i != 1;}); + auto src_list_column = cudf::make_lists_column( + 3, + fixed_width_column_wrapper{0, 5, 5, 7}.release(), + src_strings_column.release(), + 1, + detail::make_null_mask(src_validity, src_validity + 3) + ); + + auto target_list_column = lists_column_wrapper { + {"zero"}, + {"one", "one"}, + {"two", "two"}, + {"three", "three"}, + {"four", "four"}, + {"five", "five"} + }; + + auto scatter_map = fixed_width_column_wrapper{2, 4, 0}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column}) + ); + + auto expected_strings = strings_column_wrapper { + {"california", "dreaming", + "one", "one", + "all", "the", "leaves", "are", "brown", + "three", "three", + "five", "five"}, + make_counting_transform_iterator(0, [](auto i) {return i!=0 && i!=7;}) + }; + + auto expected_validity = make_counting_transform_iterator(0, [](auto i) { return i != 4; }); + auto expected_lists = cudf::make_lists_column( + 6, + fixed_width_column_wrapper{0,2,4,9,11,11,13}.release(), + expected_strings.release(), + 1, + detail::make_null_mask(expected_validity, expected_validity+6) + ); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + expected_lists->view(), + ret->get_column(0) + ); +} + +TYPED_TEST(TypedScatterListsTest, ListsOfLists) +{ + using namespace cudf::test; + using T = TypeParam; + + auto src_list_column = lists_column_wrapper { + { {1,1,1,1}, {2,2,2,2} }, + { {3,3,3,3}, {4,4,4,4} } + }; + + auto target_list_column = lists_column_wrapper { + { {9,9,9}, {8,8,8}, {7,7,7} }, + { {6,6,6}, {5,5,5}, {4,4,4} }, + { {3,3,3}, {2,2,2}, {1,1,1} }, + { {9,9}, {8,8}, {7,7} }, + { {6,6}, {5,5}, {4,4} }, + { {3,3}, {2,2}, {1,1} } + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), + scatter_map, + cudf::table_view({target_list_column}) + ); + + auto expected = lists_column_wrapper { + { {3,3,3,3}, {4,4,4,4} }, + { {6,6,6}, {5,5,5}, {4,4,4} }, + { {1,1,1,1}, {2,2,2,2} }, + { {9,9}, {8,8}, {7,7} }, + { {6,6}, {5,5}, {4,4} }, + { {3,3}, {2,2}, {1,1} } + }; + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + expected, + ret->get_column(0) + ); +} + +TYPED_TEST(TypedScatterListsTest, EmptyListsOfLists) +{ + using namespace cudf::test; + using T = TypeParam; + + auto src_list_column = lists_column_wrapper { + { {1,1,1,1}, {2,2,2,2} }, + { {3,3,3,3}, {} }, + {} + }; + + auto target_list_column = lists_column_wrapper { + { {9,9,9}, {8,8,8}, {7,7,7} }, + { {6,6,6}, {5,5,5}, {4,4,4} }, + { {3,3,3}, {2,2,2}, {1,1,1} }, + { {9,9}, {8,8}, {7,7} }, + { {6,6}, {5,5}, {4,4} }, + { {3,3}, {2,2}, {1,1} } + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 4}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), + scatter_map, + cudf::table_view({target_list_column}) + ); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + lists_column_wrapper { + { {3,3,3,3}, {} }, + { {6,6,6}, {5,5,5}, {4,4,4} }, + { {1,1,1,1}, {2,2,2,2} }, + { {9,9}, {8,8}, {7,7} }, + { }, + { {3,3}, {2,2}, {1,1} } + }, + ret->get_column(0) + ); +} + +TYPED_TEST(TypedScatterListsTest, NullListsOfLists) +{ + using namespace cudf::test; + using T = TypeParam; + + auto src_list_column = lists_column_wrapper { + { + { {1,1,1,1}, {2,2,2,2} }, + { {3,3,3,3}, {} }, + {} + }, + make_counting_transform_iterator(0, [](auto i) { return i != 2; }) + }; + + auto target_list_column = lists_column_wrapper { + { {9,9,9}, {8,8,8}, {7,7,7} }, + { {6,6,6}, {5,5,5}, {4,4,4} }, + { {3,3,3}, {2,2,2}, {1,1,1} }, + { {9,9}, {8,8}, {7,7} }, + { {6,6}, {5,5}, {4,4} }, + { {3,3}, {2,2}, {1,1} } + }; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 4}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), + scatter_map, + cudf::table_view({target_list_column}) + ); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + lists_column_wrapper { + { + { {3,3,3,3}, {} }, + { {6,6,6}, {5,5,5}, {4,4,4} }, + { {1,1,1,1}, {2,2,2,2} }, + { {9,9}, {8,8}, {7,7} }, + { }, + { {3,3}, {2,2}, {1,1} } + }, + make_counting_transform_iterator(0, [](auto i) { return i != 4; }) + }, + ret->get_column(0) + ); +} From 3cac246f1d0e78259998df90d993d5d4af8eb010 Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Fri, 13 Nov 2020 17:03:06 -0800 Subject: [PATCH 02/10] [scatter] Minor refactoring: 1. Additional type checking for child columns 2. Moved list_device_view functions inline --- cpp/include/cudf/lists/detail/scatter.cuh | 19 +++-- cpp/include/cudf/lists/list_device_view.cuh | 77 +++++++++------------ 2 files changed, 47 insertions(+), 49 deletions(-) diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index b0b9a411717..3e99ac553ee 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -656,7 +656,7 @@ struct list_child_constructor num_child_rows, std::move(child_offsets), std::move(child_column), - child_null_mask.second, // Null count + child_null_mask.second, // Null count std::move(child_null_mask.first), // Null mask stream, mr @@ -664,6 +664,19 @@ struct list_child_constructor } }; +/** + * @brief Checks that the specified columns have matching schemas, all the way down. + */ +void assert_same_data_type(column_view const& lhs, column_view const& rhs) +{ + CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Mismatched Data types."); + CUDF_EXPECTS(lhs.num_children() == rhs.num_children(), "Mismatched number of child columns."); + + for (int i{0}; i scatter( auto child_column_type = lists_column_view(target).child().type(); - // TODO: Deep(er) checks that source and target have identical types. + assert_same_data_type(source, target); using lists_column_device_view = cudf::detail::lists_column_device_view; using unbound_list_view = cudf::lists::detail::unbound_list_view; auto source_lists_column_view = lists_column_view(source); // Checks that this is a list column. auto source_device_view = column_device_view::create(source, stream); - // auto source_lists_column_device_view = lists_column_device_view(*source_device_view); auto source_vector = list_vector_from_column(unbound_list_view::SOURCE, lists_column_device_view(*source_device_view), stream); auto target_lists_column_view = lists_column_view(target); // Checks that target is a list column. auto target_device_view = column_device_view::create(target, stream); - // auto target_lists_column_device_view = lists_column_device_view(*target_device_view); auto target_vector = list_vector_from_column(unbound_list_view::TARGET, lists_column_device_view(*target_device_view), stream); // Scatter. diff --git a/cpp/include/cudf/lists/list_device_view.cuh b/cpp/include/cudf/lists/list_device_view.cuh index f3b7d2e7ebe..12764740ef8 100644 --- a/cpp/include/cudf/lists/list_device_view.cuh +++ b/cpp/include/cudf/lists/list_device_view.cuh @@ -34,7 +34,19 @@ class list_device_view { list_device_view() = default; - CUDA_DEVICE_CALLABLE list_device_view(lists_column_device_view const& lists_column, size_type const& idx); + CUDA_DEVICE_CALLABLE list_device_view(lists_column_device_view const& lists_column, size_type const& row_index) + : lists_column(lists_column), _row_index(row_index) + { + release_assert(row_index >= 0 && row_index < lists_column.size() && "row_index out of bounds"); + + column_device_view const& offsets = lists_column.offsets(); + release_assert(row_index < offsets.size() && "row_index should not have exceeded offset size"); + + begin_offset = offsets.element(row_index); + release_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() && + "begin_offset out of bounds."); + _size = offsets.element(row_index + 1) - begin_offset; + } ~list_device_view() = default; @@ -59,7 +71,11 @@ class list_device_view { * The offset of this element as stored in the child column (i.e. 5) * may be fetched using this method. */ - CUDA_DEVICE_CALLABLE size_type element_offset(size_type idx) const; + CUDA_DEVICE_CALLABLE size_type element_offset(size_type idx) const + { + release_assert(idx >= 0 && idx < size() && "idx out of bounds"); + return begin_offset + idx; + } /** * @brief Fetches the element at the specified index, within the list row. @@ -69,17 +85,28 @@ class list_device_view { * @return The element at the specified index of the list row. */ template - CUDA_DEVICE_CALLABLE T element(size_type idx) const; + CUDA_DEVICE_CALLABLE T element(size_type idx) const + { + return lists_column.child().element(element_offset(idx)); + } /** * @brief Checks whether element is null at specified index in the list row. */ - CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const; + CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const + { + release_assert(idx >= 0 && idx < size() && "Index out of bounds."); + auto element_offset = begin_offset + idx; + return lists_column.child().is_null(element_offset); + } /** * @brief Checks whether this list row is null. */ - CUDA_DEVICE_CALLABLE bool is_null() const; + CUDA_DEVICE_CALLABLE bool is_null() const + { + return lists_column.is_null(_row_index); + } /** * @brief Fetches the number of elements in this list row. @@ -101,44 +128,4 @@ class list_device_view { }; -CUDA_DEVICE_CALLABLE list_device_view::list_device_view( - lists_column_device_view const& lists_column, size_type const& row_index) - : lists_column(lists_column), _row_index(row_index) -{ - release_assert(row_index >= 0 && row_index < lists_column.size() && "row_index out of bounds"); - - column_device_view const& offsets = lists_column.offsets(); - release_assert(row_index < offsets.size() && "row_index should not have exceeded offset size"); - - begin_offset = offsets.element(row_index); - release_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() && - "begin_offset out of bounds."); - _size = offsets.element(row_index + 1) - begin_offset; -} - -CUDA_DEVICE_CALLABLE size_type list_device_view::element_offset(size_type idx) const -{ - release_assert(idx >= 0 && idx < size() && "idx out of bounds"); - // release_assert(!is_null() && !is_null(idx) && "Cannot read null element."); - return begin_offset + idx; -} - -template -CUDA_DEVICE_CALLABLE T list_device_view::element(size_type idx) const -{ - return lists_column.child().element(element_offset(idx)); -} - -CUDA_DEVICE_CALLABLE bool list_device_view::is_null(size_type idx) const -{ - release_assert(idx >= 0 && idx < size() && "Index out of bounds."); - auto element_offset = begin_offset + idx; - return lists_column.child().is_null(element_offset); -} - -CUDA_DEVICE_CALLABLE bool list_device_view::is_null() const -{ - return lists_column.is_null(_row_index); -} - } // namespace cudf From bce24b4356957312e88d79276740e34b39ba3d7a Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Sat, 14 Nov 2020 20:30:30 -0800 Subject: [PATCH 03/10] [scatter] Changelog, + clang-format. --- CHANGELOG.md | 1 + cpp/include/cudf/detail/scatter.cuh | 5 +- cpp/include/cudf/lists/detail/scatter.cuh | 703 +++++++-------- cpp/include/cudf/lists/list_device_view.cuh | 194 ++-- .../cudf/lists/lists_column_device_view.cuh | 34 +- cpp/tests/copying/scatter_list_tests.cu | 846 ++++++++---------- 6 files changed, 817 insertions(+), 966 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index 96da44e83e1..25e0098e8f4 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -17,6 +17,7 @@ - PR #6638 Add support for `pipe` API - PR #6675 Add DecimalDtype to cuDF - PR #6739 Add Java bindings for is_timestamp +- PR #6768 Add support for scatter() on list columns ## Improvements diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index 37b4259c3a6..dc16780c05b 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -24,8 +24,8 @@ #include #include #include -#include #include +#include #include #include @@ -128,7 +128,8 @@ struct column_scatterer_impl { rmm::mr::device_memory_resource* mr, cudaStream_t stream) const { - return cudf::lists::detail::scatter(source, scatter_map_begin, scatter_map_end, target, mr, stream); + return cudf::lists::detail::scatter( + source, scatter_map_begin, scatter_map_end, target, mr, stream); } }; diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 3e99ac553ee..36ed1c2d7db 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -16,9 +16,9 @@ #pragma once -#include #include -#include +#include +#include #include #include #include @@ -26,7 +26,7 @@ #include #include #include -#include +#include namespace cudf { namespace lists { @@ -37,44 +37,42 @@ namespace { /** * @brief Holder for a list row's positional information, without * also holding a reference to the list column. - * + * * Analogous to the list_view, this class is default constructable, * and can thus be stored in rmm::device_vector. It is used to represent - * the results of a `scatter()` operation; a device_vector may hold - * several instances of unbound_list_view, each with a flag indicating + * the results of a `scatter()` operation; a device_vector may hold + * several instances of unbound_list_view, each with a flag indicating * whether it came from the scatter source or target. Each instance * may later be "bound" to the appropriate source/target column, to * reconstruct the list_view. */ -struct unbound_list_view -{ +struct unbound_list_view { /** * @brief Flag type, indicating whether this list row originated from * the source or target column, in `scatter()`. */ - enum label_t : bool {SOURCE, TARGET}; + enum label_t : bool { SOURCE, TARGET }; using lists_column_device_view = cudf::detail::lists_column_device_view; - using list_device_view = cudf::list_device_view; + using list_device_view = cudf::list_device_view; - unbound_list_view() = default; + unbound_list_view() = default; unbound_list_view(unbound_list_view const&) = default; - unbound_list_view(unbound_list_view &&) = default; - unbound_list_view& operator = (unbound_list_view const&) = default; - unbound_list_view& operator = (unbound_list_view &&) = default; + unbound_list_view(unbound_list_view&&) = default; + unbound_list_view& operator=(unbound_list_view const&) = default; + unbound_list_view& operator=(unbound_list_view&&) = default; /** * @brief (__device__) Constructor, for use from `scatter()`. - * + * * @param scatter_source_label Whether the row came from source or target * @param lists_column The actual source/target lists column * @param row_index Index of the row in lists_column that this instance represents */ CUDA_DEVICE_CALLABLE unbound_list_view(label_t scatter_source_label, - cudf::detail::lists_column_device_view const& lists_column, - size_type const& row_index) - : _label{scatter_source_label}, - _row_index{row_index} + cudf::detail::lists_column_device_view const& lists_column, + size_type const& row_index) + : _label{scatter_source_label}, _row_index{row_index} { _size = list_device_view{lists_column, row_index}.size(); } @@ -82,18 +80,17 @@ struct unbound_list_view /** * @brief (__device__) Constructor, for use when constructing the child column * of a scattered list column - * + * * @param scatter_source_label Whether the row came from source or target * @param row_index Index of the row that this instance represents in the source/target column * @param size The number of elements in this list row */ CUDA_DEVICE_CALLABLE unbound_list_view(label_t scatter_source_label, - size_type const& row_index, - size_type const& size) - : _label{scatter_source_label}, - _row_index{row_index}, - _size{size} - {} + size_type const& row_index, + size_type const& size) + : _label{scatter_source_label}, _row_index{row_index}, _size{size} + { + } /** * @brief Returns number of elements in this list-row. @@ -113,33 +110,31 @@ struct unbound_list_view /** * @brief Binds to source/target column (depending on SOURCE/TARGET labels), * to produce a bound list_view. - * + * * @param scatter_source Source column for the scatter operation * @param scatter_target Target column for the scatter operation * @return A (bound) list_view for the row that this object represents */ - CUDA_DEVICE_CALLABLE list_device_view bind_to_column( - lists_column_device_view const& scatter_source, - lists_column_device_view const& scatter_target) const + CUDA_DEVICE_CALLABLE list_device_view + bind_to_column(lists_column_device_view const& scatter_source, + lists_column_device_view const& scatter_target) const { - return list_device_view(_label == SOURCE? scatter_source : scatter_target, _row_index); + return list_device_view(_label == SOURCE ? scatter_source : scatter_target, _row_index); } - private: + private: + // Note: Cannot store reference to list column, because of storage in device_vector. + // Only keep track of whether this list row came from the source or target of scatter. - // Note: Cannot store reference to list column, because of storage in device_vector. - // Only keep track of whether this list row came from the source or target of scatter. - - label_t _label {SOURCE}; // Whether this list row came from the scatter source or target. - size_type _row_index{}; // Row index in the Lists column. - size_type _size{}; // Number of elements in *this* list row. + label_t _label{SOURCE}; // Whether this list row came from the scatter source or target. + size_type _row_index{}; // Row index in the Lists column. + size_type _size{}; // Number of elements in *this* list row. }; rmm::device_vector list_vector_from_column( unbound_list_view::label_t label, cudf::detail::lists_column_device_view const& lists_column, - cudaStream_t stream -) + cudaStream_t stream) { auto n_rows = lists_column.size(); @@ -149,15 +144,9 @@ rmm::device_vector list_vector_from_column( rmm::exec_policy(stream)->on(stream), thrust::make_counting_iterator(0), n_rows, - [ - label, - lists_column, - output = vector.data().get() - ] __device__ (size_type row_index) - { + [label, lists_column, output = vector.data().get()] __device__(size_type row_index) { output[row_index] = unbound_list_view{label, lists_column, row_index}; - } - ); + }); return vector; } @@ -166,7 +155,7 @@ rmm::device_vector list_vector_from_column( * @brief Utility function to fetch the number of rows in a lists column's * child column, given its offsets column. * (This is simply the last value in the offsets column.) - * + * * @param list_offsets Offsets child of a lists column * @param stream The cuda-stream to synchronize on, when reading from device memory * @return int32_t The last element in the list_offsets column, indicating @@ -176,18 +165,18 @@ int32_t get_num_child_rows(cudf::column_view const& list_offsets, cudaStream_t s { // Number of rows in child-column == last offset value. int32_t num_child_rows{}; - CUDA_TRY(cudaMemcpyAsync(&num_child_rows, - list_offsets.data()+list_offsets.size()-1, - sizeof(int32_t), - cudaMemcpyDeviceToHost, - stream)); - CUDA_TRY(cudaStreamSynchronize(stream)); + CUDA_TRY(cudaMemcpyAsync(&num_child_rows, + list_offsets.data() + list_offsets.size() - 1, + sizeof(int32_t), + cudaMemcpyDeviceToHost, + stream)); + CUDA_TRY(cudaStreamSynchronize(stream)); return num_child_rows; } /** * @brief Constructs null mask for a scattered list's child column - * + * * @param parent_list_vector Vector of unbound_list_view, for parent lists column * @param parent_list_offsets List column offsets for parent lists column * @param source_lists Source lists column for scatter operation @@ -197,48 +186,40 @@ int32_t get_num_child_rows(cudf::column_view const& list_offsets, cudaStream_t s * @param stream CUDA stream used for device memory operations and kernel launches * @return std::pair Child column's null mask and null row count */ -std::pair -construct_child_nullmask(rmm::device_vector const& parent_list_vector, - column_view const& parent_list_offsets, - cudf::detail::lists_column_device_view const& source_lists, - cudf::detail::lists_column_device_view const& target_lists, - size_type num_child_rows, - rmm::mr::device_memory_resource* mr, - cudaStream_t stream) +std::pair construct_child_nullmask( + rmm::device_vector const& parent_list_vector, + column_view const& parent_list_offsets, + cudf::detail::lists_column_device_view const& source_lists, + cudf::detail::lists_column_device_view const& target_lists, + size_type num_child_rows, + rmm::mr::device_memory_resource* mr, + cudaStream_t stream) { - auto is_valid_predicate = [ - d_list_vector = parent_list_vector.data().get(), - d_offsets = parent_list_offsets.template data(), - d_offsets_size = parent_list_offsets.size(), - source_lists, - target_lists - ] __device__ (auto const& i) { - - auto list_start = thrust::upper_bound(thrust::seq, - d_offsets, - d_offsets + d_offsets_size, - i) - 1; - auto list_index = list_start - d_offsets; + auto is_valid_predicate = [d_list_vector = parent_list_vector.data().get(), + d_offsets = parent_list_offsets.template data(), + d_offsets_size = parent_list_offsets.size(), + source_lists, + target_lists] __device__(auto const& i) { + auto list_start = + thrust::upper_bound(thrust::seq, d_offsets, d_offsets + d_offsets_size, i) - 1; + auto list_index = list_start - d_offsets; auto element_index = i - *list_start; auto list_row = d_list_vector[list_index]; return !list_row.bind_to_column(source_lists, target_lists).is_null(element_index); }; - return cudf::detail::valid_if( - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(num_child_rows), - is_valid_predicate, - stream, - mr - ); + return cudf::detail::valid_if(thrust::make_counting_iterator(0), + thrust::make_counting_iterator(num_child_rows), + is_valid_predicate, + stream, + mr); } #ifndef NDEBUG void print(std::string const& msg, column_view const& col, cudaStream_t stream) { - if (col.type().id() != type_id::INT32) - { + if (col.type().id() != type_id::INT32) { std::cout << "[Cannot print non-INT32 column.]" << std::endl; return; } @@ -248,36 +229,36 @@ void print(std::string const& msg, column_view const& col, cudaStream_t stream) rmm::exec_policy(stream)->on(stream), thrust::make_counting_iterator(0), col.size(), - [c = col.template data()]__device__(auto const& i) { - printf("%d,", c[i]); - } - ); + [c = col.template data()] __device__(auto const& i) { printf("%d,", c[i]); }); std::cout << "]" << std::endl; } -void print(std::string const& msg, rmm::device_vector const& scatter, cudaStream_t stream) +void print(std::string const& msg, + rmm::device_vector const& scatter, + cudaStream_t stream) { std::cout << msg << " == ["; - thrust::for_each_n( - rmm::exec_policy(stream)->on(stream), - thrust::make_counting_iterator(0), - scatter.size(), - [s = scatter.data().get()] __device__ (auto const& i) { - auto si = s[i]; - printf("%s[%d](%d), ", (si.label() == unbound_list_view::SOURCE? "S":"T"), si.row_index(), si.size()); - } - ); + thrust::for_each_n(rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + scatter.size(), + [s = scatter.data().get()] __device__(auto const& i) { + auto si = s[i]; + printf("%s[%d](%d), ", + (si.label() == unbound_list_view::SOURCE ? "S" : "T"), + si.row_index(), + si.size()); + }); std::cout << "]" << std::endl; } -#endif // NDEBUG +#endif // NDEBUG /** * @brief (type_dispatch endpoint) Functor that constructs the child column result * of `scatter()`ing a list column. - * + * * The protocol is as follows: - * + * * Inputs: * 1. list_vector: A device_vector of unbound_list_view, with each element * indicating the position, size, and which column the list @@ -286,69 +267,65 @@ void print(std::string const& msg, rmm::device_vector const& * marking the beginning of a list row. * 3. source_list: The lists-column that is the source of the scatter(). * 4. target_list: The lists-column that is the target of the scatter(). - * + * * Output: A (possibly non-list) child column, which may be used in combination * with list_offsets to fully construct the outer list. - * + * * Example: - * + * * Consider the following scatter operation of two `list` columns: - * + * * 1. Source: [{9,9,9,9}, {8,8,8}], i.e. * a. Child: [9,9,9,9,8,8,8] * b. Offsets: [0, 4, 7] - * + * * 2. Target: [{1,1}, {2,2}, {3,3}], i.e. * a. Child: [1,1,2,2,3,3] * b. Offsets: [0, 2, 4, 6] - * + * * 3. Scatter-map: [2, 0] - * + * * 4. Expected output: [{8,8,8}, {2,2}, {9,9,9,9}], i.e. * a. Child: [8,8,8,2,2,9,9,9,9] <--- THIS * b. Offsets: [0, 3, 5, 9] - * + * * It is the Expected Child column above that list_child_constructor attempts * to construct. - * + * * `list_child_constructor` expects to be called with the `Source`/`Target` * lists columns, along with the following: - * + * * 1. list_vector: [ S[1](3), T[1](2), S[0](4) ] * Each unbound_list_view (e.g. S[1](3)) indicates: * a. Which column the row is bound to: S == Source, T == Target * b. The list index. E.g. S[1] indicates the 2nd list row of the Source column. * c. The row size. E.g. S[1](3) indicates that the row has 3 elements. - * + * * 2. list_offsets: [0, 3, 5, 9] - * The caller may construct this with an `inclusive_scan()` on `list_vector` + * The caller may construct this with an `inclusive_scan()` on `list_vector` * element sizes. */ -struct list_child_constructor -{ - private: +struct list_child_constructor { + private: /** * @brief Function to determine what types are supported as child column types, * when scattering lists. - * + * * @tparam T The data type of the child column of the list being scattered. */ template - struct is_supported_child_type - { - static const bool value = cudf::is_fixed_width() - || std::is_same::value - || std::is_same::value; + struct is_supported_child_type { + static const bool value = cudf::is_fixed_width() || std::is_same::value || + std::is_same::value; }; - public: - + public: /** * @brief SFINAE catch-all, for unsupported child column types. */ - template + template std::enable_if_t::value, std::unique_ptr> operator()( - rmm::device_vector const& list_vector, + rmm::device_vector const& list_vector, cudf::column_view const& list_offsets, cudf::lists_column_view const& source_list, cudf::lists_column_view const& target_list, @@ -363,25 +340,28 @@ struct list_child_constructor */ template std::enable_if_t(), std::unique_ptr> operator()( - rmm::device_vector const& list_vector, + rmm::device_vector const& list_vector, cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, rmm::mr::device_memory_resource* mr, cudaStream_t stream) const { - auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); - auto target_column_device_view = column_device_view::create(target_lists_column_view.parent(), stream); + auto source_column_device_view = + column_device_view::create(source_lists_column_view.parent(), stream); + auto target_column_device_view = + column_device_view::create(target_lists_column_view.parent(), stream); auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view); auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view); // Number of rows in child-column == last offset value. int32_t num_child_rows{get_num_child_rows(list_offsets, stream)}; - auto child_null_mask = + auto child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() - ? construct_child_nullmask(list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) - : std::make_pair(rmm::device_buffer{}, 0); + ? construct_child_nullmask( + list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + : std::make_pair(rmm::device_buffer{}, 0); #ifndef NDEBUG print("list_offsets ", list_offsets, stream); @@ -390,68 +370,64 @@ struct list_child_constructor print("target_lists.child() ", target_lists_column_view.child(), stream); print("target_lists.offsets() ", target_lists_column_view.offsets(), stream); print("scatter_rows ", list_vector, stream); -#endif // NDEBUG +#endif // NDEBUG // Init child-column. - auto child_column = cudf::make_fixed_width_column( - cudf::data_type{cudf::type_to_id()}, - num_child_rows, - child_null_mask.first, - child_null_mask.second, - stream, - mr - ); + auto child_column = cudf::make_fixed_width_column(cudf::data_type{cudf::type_to_id()}, + num_child_rows, + child_null_mask.first, + child_null_mask.second, + stream, + mr); // Function to copy child-values for specified index of unbound_list_view // to the child column. - auto copy_child_values_for_list_index = [ - d_scattered_lists = list_vector.data().get(), // unbound_list_view* - d_child_column = child_column->mutable_view().data(), - d_offsets = list_offsets.template data(), - source_lists, - target_lists - ] __device__ (auto const& row_index) { - - auto unbound_list_row = d_scattered_lists[row_index]; - auto actual_list_row = unbound_list_row.bind_to_column(source_lists, target_lists); - auto const& bound_column= (unbound_list_row.label() == unbound_list_view::SOURCE? source_lists : target_lists); - auto list_begin_offset = bound_column.offsets().element(unbound_list_row.row_index()); - auto list_end_offset = bound_column.offsets().element(unbound_list_row.row_index()+1); + auto copy_child_values_for_list_index = [d_scattered_lists = + list_vector.data().get(), // unbound_list_view* + d_child_column = + child_column->mutable_view().data(), + d_offsets = list_offsets.template data(), + source_lists, + target_lists] __device__(auto const& row_index) { + auto unbound_list_row = d_scattered_lists[row_index]; + auto actual_list_row = unbound_list_row.bind_to_column(source_lists, target_lists); + auto const& bound_column = + (unbound_list_row.label() == unbound_list_view::SOURCE ? source_lists : target_lists); + auto list_begin_offset = + bound_column.offsets().element(unbound_list_row.row_index()); + auto list_end_offset = + bound_column.offsets().element(unbound_list_row.row_index() + 1); #ifndef NDEBUG - printf("%d: Unbound == %s[%d](%d), Bound size == %d, calc_begin==%d, calc_end=%d, calc_size=%d\n", - row_index, - (unbound_list_row.label() == unbound_list_view::SOURCE? "S":"T"), - unbound_list_row.row_index(), - unbound_list_row.size(), - actual_list_row.size(), - list_begin_offset, - list_end_offset, - list_end_offset-list_begin_offset - ); -#endif // NDEBUG - + printf( + "%d: Unbound == %s[%d](%d), Bound size == %d, calc_begin==%d, calc_end=%d, calc_size=%d\n", + row_index, + (unbound_list_row.label() == unbound_list_view::SOURCE ? "S" : "T"), + unbound_list_row.row_index(), + unbound_list_row.size(), + actual_list_row.size(), + list_begin_offset, + list_end_offset, + list_end_offset - list_begin_offset); +#endif // NDEBUG + // Copy all elements in this list row, to "appropriate" offset in child-column. auto destination_start_offset = d_offsets[row_index]; - thrust::for_each_n( - thrust::seq, - thrust::make_counting_iterator(0), - actual_list_row.size(), - [actual_list_row, d_child_column, destination_start_offset] __device__ (auto const& list_element_index) - { - d_child_column[destination_start_offset + list_element_index] = - actual_list_row.template element(list_element_index); - } - ); + thrust::for_each_n(thrust::seq, + thrust::make_counting_iterator(0), + actual_list_row.size(), + [actual_list_row, d_child_column, destination_start_offset] __device__( + auto const& list_element_index) { + d_child_column[destination_start_offset + list_element_index] = + actual_list_row.template element(list_element_index); + }); }; // For each list-row, copy underlying elements to the child column. - thrust::for_each_n( - rmm::exec_policy(stream)->on(stream), - thrust::make_counting_iterator(0), - list_vector.size(), - copy_child_values_for_list_index - ); + thrust::for_each_n(rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + list_vector.size(), + copy_child_values_for_list_index); return std::make_unique(child_column->view()); } @@ -459,19 +435,19 @@ struct list_child_constructor /** * @brief Implementation for list child columns that contain strings. */ - template - std::enable_if_t::value, - std::unique_ptr> - operator()( - rmm::device_vector const& list_vector, + template + std::enable_if_t::value, std::unique_ptr> operator()( + rmm::device_vector const& list_vector, cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, rmm::mr::device_memory_resource* mr, cudaStream_t stream) const { - auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); - auto target_column_device_view = column_device_view::create(target_lists_column_view.parent(), stream); + auto source_column_device_view = + column_device_view::create(source_lists_column_view.parent(), stream); + auto target_column_device_view = + column_device_view::create(target_lists_column_view.parent(), stream); auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view); auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view); @@ -479,70 +455,70 @@ struct list_child_constructor auto string_views = rmm::device_vector(num_child_rows); - auto populate_string_views = [ - d_scattered_lists = list_vector.data().get(), // unbound_list_view* - d_list_offsets = list_offsets.template data(), - d_string_views = string_views.data().get(), - source_lists, - target_lists - ] __device__ (auto const& row_index) { - + auto populate_string_views = [d_scattered_lists = + list_vector.data().get(), // unbound_list_view* + d_list_offsets = list_offsets.template data(), + d_string_views = string_views.data().get(), + source_lists, + target_lists] __device__(auto const& row_index) { auto unbound_list_view = d_scattered_lists[row_index]; - auto actual_list_row = unbound_list_view.bind_to_column(source_lists, target_lists); - auto lists_column = actual_list_row.get_column(); - auto lists_offsets_column = lists_column.offsets(); - auto child_strings_column = lists_column.child(); - auto string_offsets_column = child_strings_column.child(cudf::strings_column_view::offsets_column_index); - auto string_chars_column = child_strings_column.child(cudf::strings_column_view::chars_column_index); - - auto output_start_offset = d_list_offsets[row_index]; // Offset in `string_views` at which string_views are - // to be written for this list row_index. - auto input_list_start = lists_offsets_column.template element(unbound_list_view.row_index()); + auto actual_list_row = unbound_list_view.bind_to_column(source_lists, target_lists); + auto lists_column = actual_list_row.get_column(); + auto lists_offsets_column = lists_column.offsets(); + auto child_strings_column = lists_column.child(); + auto string_offsets_column = + child_strings_column.child(cudf::strings_column_view::offsets_column_index); + auto string_chars_column = + child_strings_column.child(cudf::strings_column_view::chars_column_index); + + auto output_start_offset = + d_list_offsets[row_index]; // Offset in `string_views` at which string_views are + // to be written for this list row_index. + auto input_list_start = + lists_offsets_column.template element(unbound_list_view.row_index()); thrust::for_each_n( thrust::seq, thrust::make_counting_iterator(0), actual_list_row.size(), - [ - output_start_offset, - d_string_views, - input_list_start, - d_string_offsets = string_offsets_column.template data(), - d_string_chars = string_chars_column.template data() - ] __device__ (auto const& string_idx) - { + [output_start_offset, + d_string_views, + input_list_start, + d_string_offsets = string_offsets_column.template data(), + d_string_chars = + string_chars_column.template data()] __device__(auto const& string_idx) { // auto string_offset = output_start_offset + string_idx; - auto string_start_idx = d_string_offsets[input_list_start + string_idx]; - auto string_end_idx = d_string_offsets[input_list_start + string_idx + 1]; + auto string_start_idx = d_string_offsets[input_list_start + string_idx]; + auto string_end_idx = d_string_offsets[input_list_start + string_idx + 1]; - d_string_views[output_start_offset + string_idx] = + d_string_views[output_start_offset + string_idx] = string_view{d_string_chars + string_start_idx, string_end_idx - string_start_idx}; - } - ); + }); }; - thrust::for_each_n( - rmm::exec_policy(stream)->on(stream), - thrust::make_counting_iterator(0), - list_vector.size(), - populate_string_views - ); + thrust::for_each_n(rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + list_vector.size(), + populate_string_views); // string_views should now have been populated with source and target references. - auto string_offsets = cudf::strings::detail::child_offsets_from_string_vector(string_views, mr, stream); - auto string_chars = cudf::strings::detail::child_chars_from_string_vector(string_views, string_offsets->view().data(), 0, mr, stream); - auto child_null_mask = + auto string_offsets = + cudf::strings::detail::child_offsets_from_string_vector(string_views, mr, stream); + auto string_chars = cudf::strings::detail::child_chars_from_string_vector( + string_views, string_offsets->view().data(), 0, mr, stream); + auto child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() - ? construct_child_nullmask(list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) - : std::make_pair(rmm::device_buffer{}, 0); + ? construct_child_nullmask( + list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + : std::make_pair(rmm::device_buffer{}, 0); return cudf::make_strings_column(num_child_rows, std::move(string_offsets), std::move(string_chars), - child_null_mask.second, // Null count. - std::move(child_null_mask.first), // Null mask. - stream, + child_null_mask.second, // Null count. + std::move(child_null_mask.first), // Null mask. + stream, mr); } @@ -550,18 +526,18 @@ struct list_child_constructor * @brief (Recursively) Constructs a child column that is itself a list column. */ template - std::enable_if_t::value, - std::unique_ptr> - operator() ( - rmm::device_vector const& list_vector, + std::enable_if_t::value, std::unique_ptr> operator()( + rmm::device_vector const& list_vector, cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, rmm::mr::device_memory_resource* mr, cudaStream_t stream) const { - auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); - auto target_column_device_view = column_device_view::create(target_lists_column_view.parent(), stream); + auto source_column_device_view = + column_device_view::create(source_lists_column_view.parent(), stream); + auto target_column_device_view = + column_device_view::create(target_lists_column_view.parent(), stream); auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view); auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view); @@ -572,95 +548,80 @@ struct list_child_constructor // Function to convert from parent list_device_view instances to child list_device_views. // For instance, if a parent list_device_view has 3 elements, it should have 3 corresponding // child list_device_view instances. - auto populate_child_list_views = [ - d_scattered_lists = list_vector.data().get(), - d_list_offsets = list_offsets.template data(), - d_child_list_views = child_list_views.data().get(), - source_lists, - target_lists - ] __device__ (auto const& row_index) { - + auto populate_child_list_views = [d_scattered_lists = list_vector.data().get(), + d_list_offsets = list_offsets.template data(), + d_child_list_views = child_list_views.data().get(), + source_lists, + target_lists] __device__(auto const& row_index) { auto scattered_row = d_scattered_lists[row_index]; auto label = scattered_row.label(); auto bound_list_row = scattered_row.bind_to_column(source_lists, target_lists); auto lists_offsets_column = bound_list_row.get_column().offsets(); - auto child_column = bound_list_row.get_column().child(); - auto child_offsets = child_column.child(cudf::lists_column_view::offsets_column_index); + auto child_column = bound_list_row.get_column().child(); + auto child_offsets = child_column.child(cudf::lists_column_view::offsets_column_index); // For lists row at row_index, // 1. Number of entries in child_list_views == bound_list_row.size(). // 2. Offset of the first child list_view == d_list_offsets[row_index]. - auto output_start_offset = d_list_offsets[row_index]; - auto input_list_start = lists_offsets_column.template element(scattered_row.row_index()); + auto output_start_offset = d_list_offsets[row_index]; + auto input_list_start = + lists_offsets_column.template element(scattered_row.row_index()); thrust::for_each_n( thrust::seq, thrust::make_counting_iterator(0), bound_list_row.size(), - [ - input_list_start, - output_start_offset, - label, - d_child_list_views, - d_child_offsets = child_offsets.template data() - ] __device__ (auto const& child_list_index) - { + [input_list_start, + output_start_offset, + label, + d_child_list_views, + d_child_offsets = + child_offsets.template data()] __device__(auto const& child_list_index) { auto child_start_idx = d_child_offsets[input_list_start + child_list_index]; auto child_end_idx = d_child_offsets[input_list_start + child_list_index + 1]; - d_child_list_views[output_start_offset + child_list_index] = - unbound_list_view{label, input_list_start + child_list_index, child_end_idx - child_start_idx}; - } - ); + d_child_list_views[output_start_offset + child_list_index] = unbound_list_view{ + label, input_list_start + child_list_index, child_end_idx - child_start_idx}; + }); }; - thrust::for_each_n( - rmm::exec_policy(stream)->on(stream), - thrust::make_counting_iterator(0), - list_vector.size(), - populate_child_list_views - ); + thrust::for_each_n(rmm::exec_policy(stream)->on(stream), + thrust::make_counting_iterator(0), + list_vector.size(), + populate_child_list_views); // child_list_views should now have been populated, with source and target references. auto begin = thrust::make_transform_iterator( - child_list_views.begin(), - [] __device__ (auto const& row) { return row.size(); } - ); + child_list_views.begin(), [] __device__(auto const& row) { return row.size(); }); auto child_offsets = cudf::strings::detail::make_offsets_child_column( - begin, - begin + child_list_views.size(), - mr, - stream - ); - - auto child_column = cudf::type_dispatcher( - source_lists_column_view.child().child(1).type(), - list_child_constructor{}, - child_list_views, - child_offsets->view(), - cudf::lists_column_view(source_lists_column_view.child()), - cudf::lists_column_view(target_lists_column_view.child()), - mr, - stream - ); - - auto child_null_mask = + begin, begin + child_list_views.size(), mr, stream); + + auto child_column = + cudf::type_dispatcher(source_lists_column_view.child().child(1).type(), + list_child_constructor{}, + child_list_views, + child_offsets->view(), + cudf::lists_column_view(source_lists_column_view.child()), + cudf::lists_column_view(target_lists_column_view.child()), + mr, + stream); + + auto child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() - ? construct_child_nullmask(list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) - : std::make_pair(rmm::device_buffer{}, 0); - - return cudf::make_lists_column( - num_child_rows, - std::move(child_offsets), - std::move(child_column), - child_null_mask.second, // Null count - std::move(child_null_mask.first), // Null mask - stream, - mr - ); + ? construct_child_nullmask( + list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + : std::make_pair(rmm::device_buffer{}, 0); + + return cudf::make_lists_column(num_child_rows, + std::move(child_offsets), + std::move(child_column), + child_null_mask.second, // Null count + std::move(child_null_mask.first), // Null mask + stream, + mr); } }; @@ -672,12 +633,10 @@ void assert_same_data_type(column_view const& lhs, column_view const& rhs) CUDF_EXPECTS(lhs.type().id() == rhs.type().id(), "Mismatched Data types."); CUDF_EXPECTS(lhs.num_children() == rhs.num_children(), "Mismatched number of child columns."); - for (int i{0}; i scatter( MapIterator scatter_map_end, column_view const& target, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource(), - cudaStream_t stream = 0 - ) + cudaStream_t stream = 0) { - auto num_rows = target.size(); - - if (num_rows == 0) - { - return cudf::empty_like(target); - } - - auto child_column_type = lists_column_view(target).child().type(); - - assert_same_data_type(source, target); - - using lists_column_device_view = cudf::detail::lists_column_device_view; - using unbound_list_view = cudf::lists::detail::unbound_list_view; - - auto source_lists_column_view = lists_column_view(source); // Checks that this is a list column. - auto source_device_view = column_device_view::create(source, stream); - auto source_vector = list_vector_from_column(unbound_list_view::SOURCE, lists_column_device_view(*source_device_view), stream); - - auto target_lists_column_view = lists_column_view(target); // Checks that target is a list column. - auto target_device_view = column_device_view::create(target, stream); - auto target_vector = list_vector_from_column(unbound_list_view::TARGET, lists_column_device_view(*target_device_view), stream); - - // Scatter. - thrust::scatter( - rmm::exec_policy(stream)->on(stream), - source_vector.begin(), - source_vector.end(), - scatter_map_begin, - target_vector.begin() - ); - - auto list_size_begin = thrust::make_transform_iterator(target_vector.begin(), [] __device__(unbound_list_view l) { return l.size(); }); - auto offsets_column = cudf::strings::detail::make_offsets_child_column( - list_size_begin, - list_size_begin + target.size(), - mr, - stream - ); - - auto child_column = cudf::type_dispatcher( - child_column_type, - list_child_constructor{}, - target_vector, - offsets_column->view(), - source_lists_column_view, - target_lists_column_view, - mr, - stream - ); - - rmm::device_buffer null_mask{0, stream, mr}; - if (target.has_nulls()) { - null_mask = copy_bitmask(target, stream, mr); - } - - return cudf::make_lists_column( - num_rows, - std::move(offsets_column), - std::move(child_column), - cudf::UNKNOWN_NULL_COUNT, - std::move(null_mask), - stream, - mr - ); + auto num_rows = target.size(); + + if (num_rows == 0) { return cudf::empty_like(target); } + + auto child_column_type = lists_column_view(target).child().type(); + + assert_same_data_type(source, target); + + using lists_column_device_view = cudf::detail::lists_column_device_view; + using unbound_list_view = cudf::lists::detail::unbound_list_view; + + auto source_lists_column_view = lists_column_view(source); // Checks that this is a list column. + auto source_device_view = column_device_view::create(source, stream); + auto source_vector = list_vector_from_column( + unbound_list_view::SOURCE, lists_column_device_view(*source_device_view), stream); + + auto target_lists_column_view = + lists_column_view(target); // Checks that target is a list column. + auto target_device_view = column_device_view::create(target, stream); + auto target_vector = list_vector_from_column( + unbound_list_view::TARGET, lists_column_device_view(*target_device_view), stream); + + // Scatter. + thrust::scatter(rmm::exec_policy(stream)->on(stream), + source_vector.begin(), + source_vector.end(), + scatter_map_begin, + target_vector.begin()); + + auto list_size_begin = thrust::make_transform_iterator( + target_vector.begin(), [] __device__(unbound_list_view l) { return l.size(); }); + auto offsets_column = cudf::strings::detail::make_offsets_child_column( + list_size_begin, list_size_begin + target.size(), mr, stream); + + auto child_column = cudf::type_dispatcher(child_column_type, + list_child_constructor{}, + target_vector, + offsets_column->view(), + source_lists_column_view, + target_lists_column_view, + mr, + stream); + + rmm::device_buffer null_mask{0, stream, mr}; + if (target.has_nulls()) { null_mask = copy_bitmask(target, stream, mr); } + + return cudf::make_lists_column(num_rows, + std::move(offsets_column), + std::move(child_column), + cudf::UNKNOWN_NULL_COUNT, + std::move(null_mask), + stream, + mr); } -} // namespace detail; -} // namespace lists; -} // namespace cudf; +} // namespace detail +} // namespace lists +} // namespace cudf diff --git a/cpp/include/cudf/lists/list_device_view.cuh b/cpp/include/cudf/lists/list_device_view.cuh index 12764740ef8..5e904c49cf3 100644 --- a/cpp/include/cudf/lists/list_device_view.cuh +++ b/cpp/include/cudf/lists/list_device_view.cuh @@ -16,8 +16,8 @@ #pragma once #include -#include #include +#include namespace cudf { @@ -27,105 +27,99 @@ namespace cudf { * */ class list_device_view { - - using lists_column_device_view = cudf::detail::lists_column_device_view; - - public: - - list_device_view() = default; - - CUDA_DEVICE_CALLABLE list_device_view(lists_column_device_view const& lists_column, size_type const& row_index) - : lists_column(lists_column), _row_index(row_index) - { - release_assert(row_index >= 0 && row_index < lists_column.size() && "row_index out of bounds"); - - column_device_view const& offsets = lists_column.offsets(); - release_assert(row_index < offsets.size() && "row_index should not have exceeded offset size"); - - begin_offset = offsets.element(row_index); - release_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() && - "begin_offset out of bounds."); - _size = offsets.element(row_index + 1) - begin_offset; - } - - ~list_device_view() = default; - - /** - * @brief Fetches the offset in the list column's child that corresponds to - * the element at the specified list index. - * - * Consider the following lists column: - * [ - * [0,1,2], - * [3,4,5], - * [6,7,8] - * ] - * - * The list's internals would look like: - * offsets: [0, 3, 6, 9] - * child : [0, 1, 2, 3, 4, 5, 6, 7, 8] - * - * The second list row (i.e. row_index=1) is [3,4,5]. - * The third element (i.e. idx=2) of the second list row is 5. - * - * The offset of this element as stored in the child column (i.e. 5) - * may be fetched using this method. - */ - CUDA_DEVICE_CALLABLE size_type element_offset(size_type idx) const - { - release_assert(idx >= 0 && idx < size() && "idx out of bounds"); - return begin_offset + idx; - } - - /** - * @brief Fetches the element at the specified index, within the list row. - * - * @tparam The type of the list's element. - * @param The index into the list row - * @return The element at the specified index of the list row. - */ - template - CUDA_DEVICE_CALLABLE T element(size_type idx) const - { - return lists_column.child().element(element_offset(idx)); - } - - /** - * @brief Checks whether element is null at specified index in the list row. - */ - CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const - { - release_assert(idx >= 0 && idx < size() && "Index out of bounds."); - auto element_offset = begin_offset + idx; - return lists_column.child().is_null(element_offset); - } - - /** - * @brief Checks whether this list row is null. - */ - CUDA_DEVICE_CALLABLE bool is_null() const - { - return lists_column.is_null(_row_index); - } - - /** - * @brief Fetches the number of elements in this list row. - */ - CUDA_DEVICE_CALLABLE size_type size() const { return _size; } - - /** - * @brief Fetches the lists_column_device_view that contains this list. - */ - CUDA_DEVICE_CALLABLE lists_column_device_view const& get_column() const { return lists_column; } - - private: - - lists_column_device_view const& lists_column; - size_type _row_index{}; // Row index in the Lists column vector. - size_type _size{}; // Number of elements in *this* list row. - - size_type begin_offset; // Offset in list_column_device_view where this list begins. - + using lists_column_device_view = cudf::detail::lists_column_device_view; + + public: + list_device_view() = default; + + CUDA_DEVICE_CALLABLE list_device_view(lists_column_device_view const& lists_column, + size_type const& row_index) + : lists_column(lists_column), _row_index(row_index) + { + release_assert(row_index >= 0 && row_index < lists_column.size() && "row_index out of bounds"); + + column_device_view const& offsets = lists_column.offsets(); + release_assert(row_index < offsets.size() && "row_index should not have exceeded offset size"); + + begin_offset = offsets.element(row_index); + release_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() && + "begin_offset out of bounds."); + _size = offsets.element(row_index + 1) - begin_offset; + } + + ~list_device_view() = default; + + /** + * @brief Fetches the offset in the list column's child that corresponds to + * the element at the specified list index. + * + * Consider the following lists column: + * [ + * [0,1,2], + * [3,4,5], + * [6,7,8] + * ] + * + * The list's internals would look like: + * offsets: [0, 3, 6, 9] + * child : [0, 1, 2, 3, 4, 5, 6, 7, 8] + * + * The second list row (i.e. row_index=1) is [3,4,5]. + * The third element (i.e. idx=2) of the second list row is 5. + * + * The offset of this element as stored in the child column (i.e. 5) + * may be fetched using this method. + */ + CUDA_DEVICE_CALLABLE size_type element_offset(size_type idx) const + { + release_assert(idx >= 0 && idx < size() && "idx out of bounds"); + return begin_offset + idx; + } + + /** + * @brief Fetches the element at the specified index, within the list row. + * + * @tparam The type of the list's element. + * @param The index into the list row + * @return The element at the specified index of the list row. + */ + template + CUDA_DEVICE_CALLABLE T element(size_type idx) const + { + return lists_column.child().element(element_offset(idx)); + } + + /** + * @brief Checks whether element is null at specified index in the list row. + */ + CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const + { + release_assert(idx >= 0 && idx < size() && "Index out of bounds."); + auto element_offset = begin_offset + idx; + return lists_column.child().is_null(element_offset); + } + + /** + * @brief Checks whether this list row is null. + */ + CUDA_DEVICE_CALLABLE bool is_null() const { return lists_column.is_null(_row_index); } + + /** + * @brief Fetches the number of elements in this list row. + */ + CUDA_DEVICE_CALLABLE size_type size() const { return _size; } + + /** + * @brief Fetches the lists_column_device_view that contains this list. + */ + CUDA_DEVICE_CALLABLE lists_column_device_view const& get_column() const { return lists_column; } + + private: + lists_column_device_view const& lists_column; + size_type _row_index{}; // Row index in the Lists column vector. + size_type _size{}; // Number of elements in *this* list row. + + size_type begin_offset; // Offset in list_column_device_view where this list begins. }; } // namespace cudf diff --git a/cpp/include/cudf/lists/lists_column_device_view.cuh b/cpp/include/cudf/lists/lists_column_device_view.cuh index 5244140a783..187b9c2cf6a 100644 --- a/cpp/include/cudf/lists/lists_column_device_view.cuh +++ b/cpp/include/cudf/lists/lists_column_device_view.cuh @@ -16,9 +16,9 @@ #pragma once #include -#include #include #include +#include namespace cudf { @@ -30,43 +30,38 @@ namespace detail { * Analogous to list_column_view. */ class lists_column_device_view { - public: - ~lists_column_device_view() = default; lists_column_device_view(lists_column_device_view const&) = default; lists_column_device_view(lists_column_device_view&&) = default; - lists_column_device_view& operator= (lists_column_device_view const&) = default; - lists_column_device_view& operator= (lists_column_device_view &&) = default; + lists_column_device_view& operator=(lists_column_device_view const&) = default; + lists_column_device_view& operator=(lists_column_device_view&&) = default; - lists_column_device_view(column_device_view const& underlying_) - : underlying(underlying_) + lists_column_device_view(column_device_view const& underlying_) : underlying(underlying_) { - CUDF_EXPECTS(underlying_.type().id() == type_id::LIST, "lists_column_device_view only supports lists"); + CUDF_EXPECTS(underlying_.type().id() == type_id::LIST, + "lists_column_device_view only supports lists"); } /** * @brief Fetches number of rows in the lists column */ - CUDA_HOST_DEVICE_CALLABLE cudf::size_type size() const - { - return underlying.size(); - } + CUDA_HOST_DEVICE_CALLABLE cudf::size_type size() const { return underlying.size(); } /** * @brief Fetches the offsets column of the underlying list column. */ - CUDA_DEVICE_CALLABLE column_device_view offsets() const - { - return underlying.child(lists_column_view::offsets_column_index); + CUDA_DEVICE_CALLABLE column_device_view offsets() const + { + return underlying.child(lists_column_view::offsets_column_index); } /** * @brief Fetches the child column of the underlying list column. */ - CUDA_DEVICE_CALLABLE column_device_view child() const - { - return underlying.child(lists_column_view::child_column_index); + CUDA_DEVICE_CALLABLE column_device_view child() const + { + return underlying.child(lists_column_view::child_column_index); } /** @@ -81,10 +76,9 @@ class lists_column_device_view { CUDA_DEVICE_CALLABLE bool is_null(size_type idx) const { return underlying.is_null(idx); } private: - column_device_view underlying; }; } // namespace detail -} // namespace cudf; +} // namespace cudf diff --git a/cpp/tests/copying/scatter_list_tests.cu b/cpp/tests/copying/scatter_list_tests.cu index 8a86d2fd3ca..1d2691ed3a9 100644 --- a/cpp/tests/copying/scatter_list_tests.cu +++ b/cpp/tests/copying/scatter_list_tests.cu @@ -33,10 +33,10 @@ template class TypedScatterListsTest : public cudf::test::BaseFixture { }; -using FixedWidthTypes= cudf::test::Concat; +using FixedWidthTypes = cudf::test::Concat; TYPED_TEST_CASE(TypedScatterListsTest, FixedWidthTypes); class ScatterListsTest : public cudf::test::BaseFixture { @@ -44,511 +44,425 @@ class ScatterListsTest : public cudf::test::BaseFixture { TYPED_TEST(TypedScatterListsTest, ListsOfFixedWidth) { - using namespace cudf::test; - using T = TypeParam; + using namespace cudf::test; + using T = TypeParam; - auto src_list_column = lists_column_wrapper{ - {9, 9, 9, 9}, {8, 8, 8} - }; + auto src_list_column = lists_column_wrapper{{9, 9, 9, 9}, {8, 8, 8}}; - auto target_list_column = lists_column_wrapper{ - {0,0}, {1,1}, {2,2}, {3,3}, {4,4}, {5,5}, {6,6} - }; + auto target_list_column = + lists_column_wrapper{{0, 0}, {1, 1}, {2, 2}, {3, 3}, {4, 4}, {5, 5}, {6, 6}}; - auto scatter_map = fixed_width_column_wrapper{2, 0}; + auto scatter_map = fixed_width_column_wrapper{2, 0}; - auto ret = cudf::scatter( - cudf::table_view({src_list_column}), - scatter_map, - cudf::table_view({target_list_column})); + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), scatter_map, cudf::table_view({target_list_column})); - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT ( - ret->get_column(0), - lists_column_wrapper {{8,8,8}, {1,1}, {9,9,9,9}, {3,3}, {4,4}, {5,5}, {6,6}} - ); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + ret->get_column(0), + lists_column_wrapper{ + {8, 8, 8}, {1, 1}, {9, 9, 9, 9}, {3, 3}, {4, 4}, {5, 5}, {6, 6}}); } TYPED_TEST(TypedScatterListsTest, EmptyListsOfFixedWidth) { - using namespace cudf::test; - using T = TypeParam; - - auto src_child = fixed_width_column_wrapper { - {9, 9, 9, 9, 8, 8, 8}, - }; - - // One null list row, and one row with nulls. - auto src_list_column = cudf::make_lists_column( - 3, - fixed_width_column_wrapper{0, 4, 7, 7}.release(), - src_child.release(), - 0, - {} - ); - - auto target_list_column = lists_column_wrapper{ - {0,0}, {1,1}, {2,2}, {3,3}, {4,4}, {5,5}, {6,6} - }; - - auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column->view()}), - scatter_map, - cudf::table_view({target_list_column})); - - auto expected_child_ints = fixed_width_column_wrapper { - {8,8,8, 1,1, 9,9,9,9, 3,3, 4,4, 6,6 } - }; - auto expected_lists_column = cudf::make_lists_column( - 7, - fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), - expected_child_ints.release(), - 0, - {} - ); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - expected_lists_column->view(), - ret->get_column(0) - ); + using namespace cudf::test; + using T = TypeParam; + + auto src_child = fixed_width_column_wrapper{ + {9, 9, 9, 9, 8, 8, 8}, + }; + + // One null list row, and one row with nulls. + auto src_list_column = + cudf::make_lists_column(3, + fixed_width_column_wrapper{0, 4, 7, 7}.release(), + src_child.release(), + 0, + {}); + + auto target_list_column = + lists_column_wrapper{{0, 0}, {1, 1}, {2, 2}, {3, 3}, {4, 4}, {5, 5}, {6, 6}}; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; + + auto ret = cudf::scatter(cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_child_ints = + fixed_width_column_wrapper{{8, 8, 8, 1, 1, 9, 9, 9, 9, 3, 3, 4, 4, 6, 6}}; + auto expected_lists_column = cudf::make_lists_column( + 7, + fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), + expected_child_ints.release(), + 0, + {}); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_lists_column->view(), ret->get_column(0)); } TYPED_TEST(TypedScatterListsTest, EmptyListsOfNullableFixedWidth) { - using namespace cudf::test; - using T = TypeParam; - - auto src_child = fixed_width_column_wrapper { - {9, 9, 9, 9, 8, 8, 8}, - {1, 1, 1, 0, 1, 1, 1} - }; - - // One null list row, and one row with nulls. - auto src_list_column = cudf::make_lists_column( - 3, - fixed_width_column_wrapper{0, 4, 7, 7}.release(), - src_child.release(), - 0, - {} - ); - - auto target_list_column = lists_column_wrapper{ - {0,0}, {1,1}, {2,2}, {3,3}, {4,4}, {5,5}, {6,6} - }; - - auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column->view()}), - scatter_map, - cudf::table_view({target_list_column})); - - auto expected_child_ints = fixed_width_column_wrapper { - {8,8,8, 1,1, 9,9,9,9, 3,3, 4,4, 6,6 }, - {1,1,1, 1,1, 1,1,1,0, 1,1, 1,1, 1,1 } - }; - auto expected_lists_column = cudf::make_lists_column( - 7, - fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), - expected_child_ints.release(), - 0, - {} - ); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - expected_lists_column->view(), - ret->get_column(0) - ); + using namespace cudf::test; + using T = TypeParam; + + auto src_child = + fixed_width_column_wrapper{{9, 9, 9, 9, 8, 8, 8}, {1, 1, 1, 0, 1, 1, 1}}; + + // One null list row, and one row with nulls. + auto src_list_column = + cudf::make_lists_column(3, + fixed_width_column_wrapper{0, 4, 7, 7}.release(), + src_child.release(), + 0, + {}); + + auto target_list_column = + lists_column_wrapper{{0, 0}, {1, 1}, {2, 2}, {3, 3}, {4, 4}, {5, 5}, {6, 6}}; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; + + auto ret = cudf::scatter(cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_child_ints = fixed_width_column_wrapper{ + {8, 8, 8, 1, 1, 9, 9, 9, 9, 3, 3, 4, 4, 6, 6}, {1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1}}; + auto expected_lists_column = cudf::make_lists_column( + 7, + fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), + expected_child_ints.release(), + 0, + {}); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_lists_column->view(), ret->get_column(0)); } TYPED_TEST(TypedScatterListsTest, NullableListsOfNullableFixedWidth) { - using namespace cudf::test; - using T = TypeParam; - - auto src_child = fixed_width_column_wrapper { - {9, 9, 9, 9, 8, 8, 8}, - {1, 1, 1, 0, 1, 1, 1} - }; - - auto src_list_validity = make_counting_transform_iterator(0, [](auto i) { return i != 2; }); - // One null list row, and one row with nulls. - auto src_list_column = cudf::make_lists_column( - 3, - fixed_width_column_wrapper{0, 4, 7, 7}.release(), - src_child.release(), - 1, - detail::make_null_mask(src_list_validity, src_list_validity + 3) - ); - - auto target_list_column = lists_column_wrapper{ - {0,0}, {1,1}, {2,2}, {3,3}, {4,4}, {5,5}, {6,6} - }; - - auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column->view()}), - scatter_map, - cudf::table_view({target_list_column})); - - auto expected_child_ints = fixed_width_column_wrapper { - {8,8,8, 1,1, 9,9,9,9, 3,3, 4,4, 6,6 }, - {1,1,1, 1,1, 1,1,1,0, 1,1, 1,1, 1,1 } - }; - - auto expected_validity = make_counting_transform_iterator(0, [](auto i) { return i != 5; }); - auto expected_lists_column = cudf::make_lists_column( - 7, - fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), - expected_child_ints.release(), - 1, - detail::make_null_mask(expected_validity, expected_validity + 7) - ); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - expected_lists_column->view(), - ret->get_column(0) - ); + using namespace cudf::test; + using T = TypeParam; + + auto src_child = + fixed_width_column_wrapper{{9, 9, 9, 9, 8, 8, 8}, {1, 1, 1, 0, 1, 1, 1}}; + + auto src_list_validity = make_counting_transform_iterator(0, [](auto i) { return i != 2; }); + // One null list row, and one row with nulls. + auto src_list_column = + cudf::make_lists_column(3, + fixed_width_column_wrapper{0, 4, 7, 7}.release(), + src_child.release(), + 1, + detail::make_null_mask(src_list_validity, src_list_validity + 3)); + + auto target_list_column = + lists_column_wrapper{{0, 0}, {1, 1}, {2, 2}, {3, 3}, {4, 4}, {5, 5}, {6, 6}}; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 5}; + + auto ret = cudf::scatter(cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_child_ints = fixed_width_column_wrapper{ + {8, 8, 8, 1, 1, 9, 9, 9, 9, 3, 3, 4, 4, 6, 6}, {1, 1, 1, 1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1}}; + + auto expected_validity = make_counting_transform_iterator(0, [](auto i) { return i != 5; }); + auto expected_lists_column = cudf::make_lists_column( + 7, + fixed_width_column_wrapper{0, 3, 5, 9, 11, 13, 13, 15}.release(), + expected_child_ints.release(), + 1, + detail::make_null_mask(expected_validity, expected_validity + 7)); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_lists_column->view(), ret->get_column(0)); } TEST_F(ScatterListsTest, ListsOfStrings) { - using namespace cudf::test; - - auto src_list_column = lists_column_wrapper { - {"all", "the", "leaves", "are", "brown"}, - {"california", "dreaming"} - }; - - auto target_list_column = lists_column_wrapper { - {"zero"}, - {"one", "one"}, - {"two", "two"}, - {"three", "three", "three"}, - {"four", "four", "four", "four"} - }; - - auto scatter_map = fixed_width_column_wrapper{2, 0}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column}), - scatter_map, - cudf::table_view({target_list_column}) - ); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - lists_column_wrapper{ - {"california", "dreaming"}, - {"one", "one"}, - {"all", "the", "leaves", "are", "brown"}, - {"three", "three", "three"}, - {"four", "four", "four", "four"} - }, - ret->get_column(0) - ); + using namespace cudf::test; + + auto src_list_column = lists_column_wrapper{ + {"all", "the", "leaves", "are", "brown"}, {"california", "dreaming"}}; + + auto target_list_column = + lists_column_wrapper{{"zero"}, + {"one", "one"}, + {"two", "two"}, + {"three", "three", "three"}, + {"four", "four", "four", "four"}}; + + auto scatter_map = fixed_width_column_wrapper{2, 0}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), scatter_map, cudf::table_view({target_list_column})); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + lists_column_wrapper{{"california", "dreaming"}, + {"one", "one"}, + {"all", "the", "leaves", "are", "brown"}, + {"three", "three", "three"}, + {"four", "four", "four", "four"}}, + ret->get_column(0)); } TEST_F(ScatterListsTest, ListsOfNullableStrings) { - using namespace cudf::test; - - auto src_strings_column = strings_column_wrapper{ - {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, - { 1, 1, 1, 0, 1, 0, 1} - }; - - auto src_list_column = cudf::make_lists_column( - 2, - fixed_width_column_wrapper{0, 5, 7}.release(), - src_strings_column.release(), - 0, - {} - ); - - auto target_list_column = lists_column_wrapper { - {"zero"}, - {"one", "one"}, - {"two", "two"}, - {"three", "three"}, - {"four", "four"}, - {"five", "five"} - }; - - auto scatter_map = fixed_width_column_wrapper{2, 0}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column->view()}), - scatter_map, - cudf::table_view({target_list_column}) - ); - - auto expected_strings = strings_column_wrapper { - {"california", "dreaming", "one", "one", "all", "the", "leaves", "are", "brown", - "three", "three", "four", "four", "five", "five"}, - make_counting_transform_iterator(0, [](auto i) {return i!=0 && i!=7;}) - }; - - auto expected_lists = cudf::make_lists_column( - 6, - fixed_width_column_wrapper{0,2,4,9,11,13,15}.release(), - expected_strings.release(), - 0, - {} - ); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - expected_lists->view(), - ret->get_column(0) - ); + using namespace cudf::test; + + auto src_strings_column = strings_column_wrapper{ + {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, {1, 1, 1, 0, 1, 0, 1}}; + + auto src_list_column = + cudf::make_lists_column(2, + fixed_width_column_wrapper{0, 5, 7}.release(), + src_strings_column.release(), + 0, + {}); + + auto target_list_column = lists_column_wrapper{{"zero"}, + {"one", "one"}, + {"two", "two"}, + {"three", "three"}, + {"four", "four"}, + {"five", "five"}}; + + auto scatter_map = fixed_width_column_wrapper{2, 0}; + + auto ret = cudf::scatter(cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_strings = strings_column_wrapper{ + {"california", + "dreaming", + "one", + "one", + "all", + "the", + "leaves", + "are", + "brown", + "three", + "three", + "four", + "four", + "five", + "five"}, + make_counting_transform_iterator(0, [](auto i) { return i != 0 && i != 7; })}; + + auto expected_lists = cudf::make_lists_column( + 6, + fixed_width_column_wrapper{0, 2, 4, 9, 11, 13, 15}.release(), + expected_strings.release(), + 0, + {}); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_lists->view(), ret->get_column(0)); } TEST_F(ScatterListsTest, EmptyListsOfNullableStrings) { - using namespace cudf::test; - - auto src_strings_column = strings_column_wrapper{ - {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, - { 1, 1, 1, 0, 1, 0, 1} - }; - - auto src_list_column = cudf::make_lists_column( - 3, - fixed_width_column_wrapper{0, 5, 5, 7}.release(), - src_strings_column.release(), - 0, - {} - ); - - auto target_list_column = lists_column_wrapper { - {"zero"}, - {"one", "one"}, - {"two", "two"}, - {"three", "three"}, - {"four", "four"}, - {"five", "five"} - }; - - auto scatter_map = fixed_width_column_wrapper{2, 4, 0}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column->view()}), - scatter_map, - cudf::table_view({target_list_column}) - ); - - auto expected_strings = strings_column_wrapper { - {"california", "dreaming", - "one", "one", - "all", "the", "leaves", "are", "brown", - "three", "three", - "five", "five"}, - make_counting_transform_iterator(0, [](auto i) {return i!=0 && i!=7;}) - }; - - auto expected_lists = cudf::make_lists_column( - 6, - fixed_width_column_wrapper{0,2,4,9,11,11,13}.release(), - expected_strings.release(), - 0, - {} - ); - - std::cout << "Expected: " << std::endl; print(expected_lists->view()); - std::cout << "Received: " << std::endl; print(ret->get_column(0)); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - expected_lists->view(), - ret->get_column(0) - ); + using namespace cudf::test; + + auto src_strings_column = strings_column_wrapper{ + {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, {1, 1, 1, 0, 1, 0, 1}}; + + auto src_list_column = + cudf::make_lists_column(3, + fixed_width_column_wrapper{0, 5, 5, 7}.release(), + src_strings_column.release(), + 0, + {}); + + auto target_list_column = lists_column_wrapper{{"zero"}, + {"one", "one"}, + {"two", "two"}, + {"three", "three"}, + {"four", "four"}, + {"five", "five"}}; + + auto scatter_map = fixed_width_column_wrapper{2, 4, 0}; + + auto ret = cudf::scatter(cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_strings = strings_column_wrapper{ + {"california", + "dreaming", + "one", + "one", + "all", + "the", + "leaves", + "are", + "brown", + "three", + "three", + "five", + "five"}, + make_counting_transform_iterator(0, [](auto i) { return i != 0 && i != 7; })}; + + auto expected_lists = cudf::make_lists_column( + 6, + fixed_width_column_wrapper{0, 2, 4, 9, 11, 11, 13}.release(), + expected_strings.release(), + 0, + {}); + + std::cout << "Expected: " << std::endl; + print(expected_lists->view()); + std::cout << "Received: " << std::endl; + print(ret->get_column(0)); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_lists->view(), ret->get_column(0)); } TEST_F(ScatterListsTest, NullableListsOfNullableStrings) { - using namespace cudf::test; - - auto src_strings_column = strings_column_wrapper{ - {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, - { 1, 1, 1, 0, 1, 0, 1} - }; - - auto src_validity = make_counting_transform_iterator(0, [](auto i) { return i != 1;}); - auto src_list_column = cudf::make_lists_column( - 3, - fixed_width_column_wrapper{0, 5, 5, 7}.release(), - src_strings_column.release(), - 1, - detail::make_null_mask(src_validity, src_validity + 3) - ); - - auto target_list_column = lists_column_wrapper { - {"zero"}, - {"one", "one"}, - {"two", "two"}, - {"three", "three"}, - {"four", "four"}, - {"five", "five"} - }; - - auto scatter_map = fixed_width_column_wrapper{2, 4, 0}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column->view()}), - scatter_map, - cudf::table_view({target_list_column}) - ); - - auto expected_strings = strings_column_wrapper { - {"california", "dreaming", - "one", "one", - "all", "the", "leaves", "are", "brown", - "three", "three", - "five", "five"}, - make_counting_transform_iterator(0, [](auto i) {return i!=0 && i!=7;}) - }; - - auto expected_validity = make_counting_transform_iterator(0, [](auto i) { return i != 4; }); - auto expected_lists = cudf::make_lists_column( - 6, - fixed_width_column_wrapper{0,2,4,9,11,11,13}.release(), - expected_strings.release(), - 1, - detail::make_null_mask(expected_validity, expected_validity+6) - ); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - expected_lists->view(), - ret->get_column(0) - ); + using namespace cudf::test; + + auto src_strings_column = strings_column_wrapper{ + {"all", "the", "leaves", "are", "brown", "california", "dreaming"}, {1, 1, 1, 0, 1, 0, 1}}; + + auto src_validity = make_counting_transform_iterator(0, [](auto i) { return i != 1; }); + auto src_list_column = + cudf::make_lists_column(3, + fixed_width_column_wrapper{0, 5, 5, 7}.release(), + src_strings_column.release(), + 1, + detail::make_null_mask(src_validity, src_validity + 3)); + + auto target_list_column = lists_column_wrapper{{"zero"}, + {"one", "one"}, + {"two", "two"}, + {"three", "three"}, + {"four", "four"}, + {"five", "five"}}; + + auto scatter_map = fixed_width_column_wrapper{2, 4, 0}; + + auto ret = cudf::scatter(cudf::table_view({src_list_column->view()}), + scatter_map, + cudf::table_view({target_list_column})); + + auto expected_strings = strings_column_wrapper{ + {"california", + "dreaming", + "one", + "one", + "all", + "the", + "leaves", + "are", + "brown", + "three", + "three", + "five", + "five"}, + make_counting_transform_iterator(0, [](auto i) { return i != 0 && i != 7; })}; + + auto expected_validity = make_counting_transform_iterator(0, [](auto i) { return i != 4; }); + auto expected_lists = cudf::make_lists_column( + 6, + fixed_width_column_wrapper{0, 2, 4, 9, 11, 11, 13}.release(), + expected_strings.release(), + 1, + detail::make_null_mask(expected_validity, expected_validity + 6)); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected_lists->view(), ret->get_column(0)); } TYPED_TEST(TypedScatterListsTest, ListsOfLists) { - using namespace cudf::test; - using T = TypeParam; - - auto src_list_column = lists_column_wrapper { - { {1,1,1,1}, {2,2,2,2} }, - { {3,3,3,3}, {4,4,4,4} } - }; - - auto target_list_column = lists_column_wrapper { - { {9,9,9}, {8,8,8}, {7,7,7} }, - { {6,6,6}, {5,5,5}, {4,4,4} }, - { {3,3,3}, {2,2,2}, {1,1,1} }, - { {9,9}, {8,8}, {7,7} }, - { {6,6}, {5,5}, {4,4} }, - { {3,3}, {2,2}, {1,1} } - }; - - auto scatter_map = fixed_width_column_wrapper{2, 0}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column}), - scatter_map, - cudf::table_view({target_list_column}) - ); - - auto expected = lists_column_wrapper { - { {3,3,3,3}, {4,4,4,4} }, - { {6,6,6}, {5,5,5}, {4,4,4} }, - { {1,1,1,1}, {2,2,2,2} }, - { {9,9}, {8,8}, {7,7} }, - { {6,6}, {5,5}, {4,4} }, - { {3,3}, {2,2}, {1,1} } - }; - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - expected, - ret->get_column(0) - ); + using namespace cudf::test; + using T = TypeParam; + + auto src_list_column = + lists_column_wrapper{{{1, 1, 1, 1}, {2, 2, 2, 2}}, {{3, 3, 3, 3}, {4, 4, 4, 4}}}; + + auto target_list_column = lists_column_wrapper{{{9, 9, 9}, {8, 8, 8}, {7, 7, 7}}, + {{6, 6, 6}, {5, 5, 5}, {4, 4, 4}}, + {{3, 3, 3}, {2, 2, 2}, {1, 1, 1}}, + {{9, 9}, {8, 8}, {7, 7}}, + {{6, 6}, {5, 5}, {4, 4}}, + {{3, 3}, {2, 2}, {1, 1}}}; + + auto scatter_map = fixed_width_column_wrapper{2, 0}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), scatter_map, cudf::table_view({target_list_column})); + + auto expected = lists_column_wrapper{{{3, 3, 3, 3}, {4, 4, 4, 4}}, + {{6, 6, 6}, {5, 5, 5}, {4, 4, 4}}, + {{1, 1, 1, 1}, {2, 2, 2, 2}}, + {{9, 9}, {8, 8}, {7, 7}}, + {{6, 6}, {5, 5}, {4, 4}}, + {{3, 3}, {2, 2}, {1, 1}}}; + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT(expected, ret->get_column(0)); } TYPED_TEST(TypedScatterListsTest, EmptyListsOfLists) { - using namespace cudf::test; - using T = TypeParam; - - auto src_list_column = lists_column_wrapper { - { {1,1,1,1}, {2,2,2,2} }, - { {3,3,3,3}, {} }, - {} - }; - - auto target_list_column = lists_column_wrapper { - { {9,9,9}, {8,8,8}, {7,7,7} }, - { {6,6,6}, {5,5,5}, {4,4,4} }, - { {3,3,3}, {2,2,2}, {1,1,1} }, - { {9,9}, {8,8}, {7,7} }, - { {6,6}, {5,5}, {4,4} }, - { {3,3}, {2,2}, {1,1} } - }; - - auto scatter_map = fixed_width_column_wrapper{2, 0, 4}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column}), - scatter_map, - cudf::table_view({target_list_column}) - ); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - lists_column_wrapper { - { {3,3,3,3}, {} }, - { {6,6,6}, {5,5,5}, {4,4,4} }, - { {1,1,1,1}, {2,2,2,2} }, - { {9,9}, {8,8}, {7,7} }, - { }, - { {3,3}, {2,2}, {1,1} } - }, - ret->get_column(0) - ); + using namespace cudf::test; + using T = TypeParam; + + auto src_list_column = + lists_column_wrapper{{{1, 1, 1, 1}, {2, 2, 2, 2}}, {{3, 3, 3, 3}, {}}, {}}; + + auto target_list_column = lists_column_wrapper{{{9, 9, 9}, {8, 8, 8}, {7, 7, 7}}, + {{6, 6, 6}, {5, 5, 5}, {4, 4, 4}}, + {{3, 3, 3}, {2, 2, 2}, {1, 1, 1}}, + {{9, 9}, {8, 8}, {7, 7}}, + {{6, 6}, {5, 5}, {4, 4}}, + {{3, 3}, {2, 2}, {1, 1}}}; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 4}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), scatter_map, cudf::table_view({target_list_column})); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + lists_column_wrapper{{{3, 3, 3, 3}, {}}, + {{6, 6, 6}, {5, 5, 5}, {4, 4, 4}}, + {{1, 1, 1, 1}, {2, 2, 2, 2}}, + {{9, 9}, {8, 8}, {7, 7}}, + {}, + {{3, 3}, {2, 2}, {1, 1}}}, + ret->get_column(0)); } TYPED_TEST(TypedScatterListsTest, NullListsOfLists) { - using namespace cudf::test; - using T = TypeParam; - - auto src_list_column = lists_column_wrapper { - { - { {1,1,1,1}, {2,2,2,2} }, - { {3,3,3,3}, {} }, - {} - }, - make_counting_transform_iterator(0, [](auto i) { return i != 2; }) - }; - - auto target_list_column = lists_column_wrapper { - { {9,9,9}, {8,8,8}, {7,7,7} }, - { {6,6,6}, {5,5,5}, {4,4,4} }, - { {3,3,3}, {2,2,2}, {1,1,1} }, - { {9,9}, {8,8}, {7,7} }, - { {6,6}, {5,5}, {4,4} }, - { {3,3}, {2,2}, {1,1} } - }; - - auto scatter_map = fixed_width_column_wrapper{2, 0, 4}; - - auto ret = cudf::scatter( - cudf::table_view({src_list_column}), - scatter_map, - cudf::table_view({target_list_column}) - ); - - CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( - lists_column_wrapper { - { - { {3,3,3,3}, {} }, - { {6,6,6}, {5,5,5}, {4,4,4} }, - { {1,1,1,1}, {2,2,2,2} }, - { {9,9}, {8,8}, {7,7} }, - { }, - { {3,3}, {2,2}, {1,1} } - }, - make_counting_transform_iterator(0, [](auto i) { return i != 4; }) - }, - ret->get_column(0) - ); + using namespace cudf::test; + using T = TypeParam; + + auto src_list_column = lists_column_wrapper{ + {{{1, 1, 1, 1}, {2, 2, 2, 2}}, {{3, 3, 3, 3}, {}}, {}}, + make_counting_transform_iterator(0, [](auto i) { return i != 2; })}; + + auto target_list_column = lists_column_wrapper{{{9, 9, 9}, {8, 8, 8}, {7, 7, 7}}, + {{6, 6, 6}, {5, 5, 5}, {4, 4, 4}}, + {{3, 3, 3}, {2, 2, 2}, {1, 1, 1}}, + {{9, 9}, {8, 8}, {7, 7}}, + {{6, 6}, {5, 5}, {4, 4}}, + {{3, 3}, {2, 2}, {1, 1}}}; + + auto scatter_map = fixed_width_column_wrapper{2, 0, 4}; + + auto ret = cudf::scatter( + cudf::table_view({src_list_column}), scatter_map, cudf::table_view({target_list_column})); + + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + lists_column_wrapper{ + {{{3, 3, 3, 3}, {}}, + {{6, 6, 6}, {5, 5, 5}, {4, 4, 4}}, + {{1, 1, 1, 1}, {2, 2, 2, 2}}, + {{9, 9}, {8, 8}, {7, 7}}, + {}, + {{3, 3}, {2, 2}, {1, 1}}}, + make_counting_transform_iterator(0, [](auto i) { return i != 4; })}, + ret->get_column(0)); } From a7dc525df492026de7f7aea5bd29d9b8d2ad7619 Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Mon, 16 Nov 2020 11:18:52 -0800 Subject: [PATCH 04/10] [scatter] Switch to rmm::cuda_stream_view --- cpp/include/cudf/detail/scatter.cuh | 2 +- cpp/include/cudf/lists/detail/scatter.cuh | 86 +++++++++++------------ 2 files changed, 41 insertions(+), 47 deletions(-) diff --git a/cpp/include/cudf/detail/scatter.cuh b/cpp/include/cudf/detail/scatter.cuh index dc16780c05b..c59eb378019 100644 --- a/cpp/include/cudf/detail/scatter.cuh +++ b/cpp/include/cudf/detail/scatter.cuh @@ -129,7 +129,7 @@ struct column_scatterer_impl { cudaStream_t stream) const { return cudf::lists::detail::scatter( - source, scatter_map_begin, scatter_map_end, target, mr, stream); + source, scatter_map_begin, scatter_map_end, target, stream, mr); } }; diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 36ed1c2d7db..8316df61197 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -134,14 +134,14 @@ struct unbound_list_view { rmm::device_vector list_vector_from_column( unbound_list_view::label_t label, cudf::detail::lists_column_device_view const& lists_column, - cudaStream_t stream) + rmm::cuda_stream_view stream) { auto n_rows = lists_column.size(); auto vector = rmm::device_vector(n_rows); thrust::for_each_n( - rmm::exec_policy(stream)->on(stream), + rmm::exec_policy(stream)->on(stream.value()), thrust::make_counting_iterator(0), n_rows, [label, lists_column, output = vector.data().get()] __device__(size_type row_index) { @@ -161,7 +161,7 @@ rmm::device_vector list_vector_from_column( * @return int32_t The last element in the list_offsets column, indicating * the number of rows in the lists-column's child. */ -int32_t get_num_child_rows(cudf::column_view const& list_offsets, cudaStream_t stream) +int32_t get_num_child_rows(cudf::column_view const& list_offsets, rmm::cuda_stream_view stream) { // Number of rows in child-column == last offset value. int32_t num_child_rows{}; @@ -169,8 +169,8 @@ int32_t get_num_child_rows(cudf::column_view const& list_offsets, cudaStream_t s list_offsets.data() + list_offsets.size() - 1, sizeof(int32_t), cudaMemcpyDeviceToHost, - stream)); - CUDA_TRY(cudaStreamSynchronize(stream)); + stream.value())); + stream.synchronize(); return num_child_rows; } @@ -192,8 +192,8 @@ std::pair construct_child_nullmask( cudf::detail::lists_column_device_view const& source_lists, cudf::detail::lists_column_device_view const& target_lists, size_type num_child_rows, - rmm::mr::device_memory_resource* mr, - cudaStream_t stream) + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto is_valid_predicate = [d_list_vector = parent_list_vector.data().get(), d_offsets = parent_list_offsets.template data(), @@ -217,7 +217,7 @@ std::pair construct_child_nullmask( } #ifndef NDEBUG -void print(std::string const& msg, column_view const& col, cudaStream_t stream) +void print(std::string const& msg, column_view const& col, rmm::cuda_stream_view stream) { if (col.type().id() != type_id::INT32) { std::cout << "[Cannot print non-INT32 column.]" << std::endl; @@ -226,7 +226,7 @@ void print(std::string const& msg, column_view const& col, cudaStream_t stream) std::cout << msg << " = ["; thrust::for_each_n( - rmm::exec_policy(stream)->on(stream), + rmm::exec_policy(stream)->on(stream.value()), thrust::make_counting_iterator(0), col.size(), [c = col.template data()] __device__(auto const& i) { printf("%d,", c[i]); }); @@ -235,11 +235,11 @@ void print(std::string const& msg, column_view const& col, cudaStream_t stream) void print(std::string const& msg, rmm::device_vector const& scatter, - cudaStream_t stream) + rmm::cuda_stream_view stream) { std::cout << msg << " == ["; - thrust::for_each_n(rmm::exec_policy(stream)->on(stream), + thrust::for_each_n(rmm::exec_policy(stream)->on(stream.value()), thrust::make_counting_iterator(0), scatter.size(), [s = scatter.data().get()] __device__(auto const& i) { @@ -323,14 +323,8 @@ struct list_child_constructor { /** * @brief SFINAE catch-all, for unsupported child column types. */ - template - std::enable_if_t::value, std::unique_ptr> operator()( - rmm::device_vector const& list_vector, - cudf::column_view const& list_offsets, - cudf::lists_column_view const& source_list, - cudf::lists_column_view const& target_list, - rmm::mr::device_memory_resource* mr, - cudaStream_t stream) const + template + std::enable_if_t::value, std::unique_ptr> operator()(Args&&... args) { CUDF_FAIL("list_child_constructor unsupported!"); } @@ -344,8 +338,8 @@ struct list_child_constructor { cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, - rmm::mr::device_memory_resource* mr, - cudaStream_t stream) const + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const { auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); @@ -360,7 +354,7 @@ struct list_child_constructor { auto child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() ? construct_child_nullmask( - list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + list_vector, list_offsets, source_lists, target_lists, num_child_rows, stream, mr) : std::make_pair(rmm::device_buffer{}, 0); #ifndef NDEBUG @@ -377,7 +371,7 @@ struct list_child_constructor { num_child_rows, child_null_mask.first, child_null_mask.second, - stream, + stream.value(), mr); // Function to copy child-values for specified index of unbound_list_view @@ -424,7 +418,7 @@ struct list_child_constructor { }; // For each list-row, copy underlying elements to the child column. - thrust::for_each_n(rmm::exec_policy(stream)->on(stream), + thrust::for_each_n(rmm::exec_policy(stream)->on(stream.value()), thrust::make_counting_iterator(0), list_vector.size(), copy_child_values_for_list_index); @@ -441,8 +435,8 @@ struct list_child_constructor { cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, - rmm::mr::device_memory_resource* mr, - cudaStream_t stream) const + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const { auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); @@ -496,7 +490,7 @@ struct list_child_constructor { }); }; - thrust::for_each_n(rmm::exec_policy(stream)->on(stream), + thrust::for_each_n(rmm::exec_policy(stream)->on(stream.value()), thrust::make_counting_iterator(0), list_vector.size(), populate_string_views); @@ -504,13 +498,13 @@ struct list_child_constructor { // string_views should now have been populated with source and target references. auto string_offsets = - cudf::strings::detail::child_offsets_from_string_vector(string_views, mr, stream); + cudf::strings::detail::child_offsets_from_string_vector(string_views, mr, stream.value()); auto string_chars = cudf::strings::detail::child_chars_from_string_vector( - string_views, string_offsets->view().data(), 0, mr, stream); + string_views, string_offsets->view().data(), 0, mr, stream.value()); auto child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() ? construct_child_nullmask( - list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + list_vector, list_offsets, source_lists, target_lists, num_child_rows, stream, mr) : std::make_pair(rmm::device_buffer{}, 0); return cudf::make_strings_column(num_child_rows, @@ -518,7 +512,7 @@ struct list_child_constructor { std::move(string_chars), child_null_mask.second, // Null count. std::move(child_null_mask.first), // Null mask. - stream, + stream.value(), mr); } @@ -531,8 +525,8 @@ struct list_child_constructor { cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, - rmm::mr::device_memory_resource* mr, - cudaStream_t stream) const + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) const { auto source_column_device_view = column_device_view::create(source_lists_column_view.parent(), stream); @@ -586,7 +580,7 @@ struct list_child_constructor { }); }; - thrust::for_each_n(rmm::exec_policy(stream)->on(stream), + thrust::for_each_n(rmm::exec_policy(stream)->on(stream.value()), thrust::make_counting_iterator(0), list_vector.size(), populate_child_list_views); @@ -597,7 +591,7 @@ struct list_child_constructor { child_list_views.begin(), [] __device__(auto const& row) { return row.size(); }); auto child_offsets = cudf::strings::detail::make_offsets_child_column( - begin, begin + child_list_views.size(), mr, stream); + begin, begin + child_list_views.size(), mr, stream.value()); auto child_column = cudf::type_dispatcher(source_lists_column_view.child().child(1).type(), @@ -606,13 +600,13 @@ struct list_child_constructor { child_offsets->view(), cudf::lists_column_view(source_lists_column_view.child()), cudf::lists_column_view(target_lists_column_view.child()), - mr, - stream); + stream, + mr); auto child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() ? construct_child_nullmask( - list_vector, list_offsets, source_lists, target_lists, num_child_rows, mr, stream) + list_vector, list_offsets, source_lists, target_lists, num_child_rows, stream, mr) : std::make_pair(rmm::device_buffer{}, 0); return cudf::make_lists_column(num_child_rows, @@ -620,7 +614,7 @@ struct list_child_constructor { std::move(child_column), child_null_mask.second, // Null count std::move(child_null_mask.first), // Null mask - stream, + stream.value(), mr); } }; @@ -664,8 +658,8 @@ std::unique_ptr scatter( MapIterator scatter_map_begin, MapIterator scatter_map_end, column_view const& target, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource(), - cudaStream_t stream = 0) + rmm::cuda_stream_view stream = 0, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { auto num_rows = target.size(); @@ -690,7 +684,7 @@ std::unique_ptr scatter( unbound_list_view::TARGET, lists_column_device_view(*target_device_view), stream); // Scatter. - thrust::scatter(rmm::exec_policy(stream)->on(stream), + thrust::scatter(rmm::exec_policy(stream)->on(stream.value()), source_vector.begin(), source_vector.end(), scatter_map_begin, @@ -699,7 +693,7 @@ std::unique_ptr scatter( auto list_size_begin = thrust::make_transform_iterator( target_vector.begin(), [] __device__(unbound_list_view l) { return l.size(); }); auto offsets_column = cudf::strings::detail::make_offsets_child_column( - list_size_begin, list_size_begin + target.size(), mr, stream); + list_size_begin, list_size_begin + target.size(), mr, stream.value()); auto child_column = cudf::type_dispatcher(child_column_type, list_child_constructor{}, @@ -707,8 +701,8 @@ std::unique_ptr scatter( offsets_column->view(), source_lists_column_view, target_lists_column_view, - mr, - stream); + stream, + mr); rmm::device_buffer null_mask{0, stream, mr}; if (target.has_nulls()) { null_mask = copy_bitmask(target, stream, mr); } @@ -718,7 +712,7 @@ std::unique_ptr scatter( std::move(child_column), cudf::UNKNOWN_NULL_COUNT, std::move(null_mask), - stream, + stream.value(), mr); } From c764de21834d0b13123fcc1a13caecd567925fde Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Mon, 16 Nov 2020 11:31:46 -0800 Subject: [PATCH 05/10] [scatter] Move label_t to scoped label_type --- cpp/include/cudf/lists/detail/scatter.cuh | 30 +++++++++++++---------- 1 file changed, 17 insertions(+), 13 deletions(-) diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 8316df61197..d44d1b47f58 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -51,7 +51,7 @@ struct unbound_list_view { * @brief Flag type, indicating whether this list row originated from * the source or target column, in `scatter()`. */ - enum label_t : bool { SOURCE, TARGET }; + enum class label_type : bool { SOURCE, TARGET }; using lists_column_device_view = cudf::detail::lists_column_device_view; using list_device_view = cudf::list_device_view; @@ -69,7 +69,7 @@ struct unbound_list_view { * @param lists_column The actual source/target lists column * @param row_index Index of the row in lists_column that this instance represents */ - CUDA_DEVICE_CALLABLE unbound_list_view(label_t scatter_source_label, + CUDA_DEVICE_CALLABLE unbound_list_view(label_type scatter_source_label, cudf::detail::lists_column_device_view const& lists_column, size_type const& row_index) : _label{scatter_source_label}, _row_index{row_index} @@ -85,7 +85,7 @@ struct unbound_list_view { * @param row_index Index of the row that this instance represents in the source/target column * @param size The number of elements in this list row */ - CUDA_DEVICE_CALLABLE unbound_list_view(label_t scatter_source_label, + CUDA_DEVICE_CALLABLE unbound_list_view(label_type scatter_source_label, size_type const& row_index, size_type const& size) : _label{scatter_source_label}, _row_index{row_index}, _size{size} @@ -100,7 +100,7 @@ struct unbound_list_view { /** * @brief Returns whether this row came from the `scatter()` source or target */ - CUDA_DEVICE_CALLABLE label_t label() const { return _label; } + CUDA_DEVICE_CALLABLE label_type label() const { return _label; } /** * @brief Returns the index in the source/target column @@ -119,20 +119,22 @@ struct unbound_list_view { bind_to_column(lists_column_device_view const& scatter_source, lists_column_device_view const& scatter_target) const { - return list_device_view(_label == SOURCE ? scatter_source : scatter_target, _row_index); + return list_device_view(_label == label_type::SOURCE ? scatter_source : scatter_target, + _row_index); } private: // Note: Cannot store reference to list column, because of storage in device_vector. // Only keep track of whether this list row came from the source or target of scatter. - label_t _label{SOURCE}; // Whether this list row came from the scatter source or target. + label_type _label{ + label_type::SOURCE}; // Whether this list row came from the scatter source or target. size_type _row_index{}; // Row index in the Lists column. size_type _size{}; // Number of elements in *this* list row. }; rmm::device_vector list_vector_from_column( - unbound_list_view::label_t label, + unbound_list_view::label_type label, cudf::detail::lists_column_device_view const& lists_column, rmm::cuda_stream_view stream) { @@ -324,7 +326,8 @@ struct list_child_constructor { * @brief SFINAE catch-all, for unsupported child column types. */ template - std::enable_if_t::value, std::unique_ptr> operator()(Args&&... args) + std::enable_if_t::value, std::unique_ptr> operator()( + Args&&... args) { CUDF_FAIL("list_child_constructor unsupported!"); } @@ -386,7 +389,8 @@ struct list_child_constructor { auto unbound_list_row = d_scattered_lists[row_index]; auto actual_list_row = unbound_list_row.bind_to_column(source_lists, target_lists); auto const& bound_column = - (unbound_list_row.label() == unbound_list_view::SOURCE ? source_lists : target_lists); + (unbound_list_row.label() == unbound_list_view::label_type::SOURCE ? source_lists + : target_lists); auto list_begin_offset = bound_column.offsets().element(unbound_list_row.row_index()); auto list_end_offset = @@ -396,7 +400,7 @@ struct list_child_constructor { printf( "%d: Unbound == %s[%d](%d), Bound size == %d, calc_begin==%d, calc_end=%d, calc_size=%d\n", row_index, - (unbound_list_row.label() == unbound_list_view::SOURCE ? "S" : "T"), + (unbound_list_row.label() == unbound_list_view::label_type::SOURCE ? "S" : "T"), unbound_list_row.row_index(), unbound_list_row.size(), actual_list_row.size(), @@ -658,7 +662,7 @@ std::unique_ptr scatter( MapIterator scatter_map_begin, MapIterator scatter_map_end, column_view const& target, - rmm::cuda_stream_view stream = 0, + rmm::cuda_stream_view stream = 0, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { auto num_rows = target.size(); @@ -675,13 +679,13 @@ std::unique_ptr scatter( auto source_lists_column_view = lists_column_view(source); // Checks that this is a list column. auto source_device_view = column_device_view::create(source, stream); auto source_vector = list_vector_from_column( - unbound_list_view::SOURCE, lists_column_device_view(*source_device_view), stream); + unbound_list_view::label_type::SOURCE, lists_column_device_view(*source_device_view), stream); auto target_lists_column_view = lists_column_view(target); // Checks that target is a list column. auto target_device_view = column_device_view::create(target, stream); auto target_vector = list_vector_from_column( - unbound_list_view::TARGET, lists_column_device_view(*target_device_view), stream); + unbound_list_view::label_type::TARGET, lists_column_device_view(*target_device_view), stream); // Scatter. thrust::scatter(rmm::exec_policy(stream)->on(stream.value()), From b1caeef6847d7836304279f4d7ca43625c46a518 Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Mon, 16 Nov 2020 12:32:53 -0800 Subject: [PATCH 06/10] [scatter] Switch from device_vector to device_uvector. --- cpp/include/cudf/lists/detail/scatter.cuh | 63 ++++++++++++++--------- 1 file changed, 39 insertions(+), 24 deletions(-) diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index d44d1b47f58..7a455edf55d 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -27,6 +27,7 @@ #include #include #include +#include namespace cudf { namespace lists { @@ -63,7 +64,7 @@ struct unbound_list_view { unbound_list_view& operator=(unbound_list_view&&) = default; /** - * @brief (__device__) Constructor, for use from `scatter()`. + * @brief __device__ Constructor, for use from `scatter()`. * * @param scatter_source_label Whether the row came from source or target * @param lists_column The actual source/target lists column @@ -78,7 +79,7 @@ struct unbound_list_view { } /** - * @brief (__device__) Constructor, for use when constructing the child column + * @brief __device__ Constructor, for use when constructing the child column * of a scattered list column * * @param scatter_source_label Whether the row came from source or target @@ -93,7 +94,7 @@ struct unbound_list_view { } /** - * @brief Returns number of elements in this list-row. + * @brief Returns number of elements in this list row. */ CUDA_DEVICE_CALLABLE size_type size() const { return _size; } @@ -133,15 +134,25 @@ struct unbound_list_view { size_type _size{}; // Number of elements in *this* list row. }; -rmm::device_vector list_vector_from_column( +rmm::device_uvector list_vector_from_column( unbound_list_view::label_type label, cudf::detail::lists_column_device_view const& lists_column, - rmm::cuda_stream_view stream) + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) { auto n_rows = lists_column.size(); - auto vector = rmm::device_vector(n_rows); + auto vector = rmm::device_uvector(n_rows, stream, mr); + + thrust::transform(rmm::exec_policy(stream)->on(stream.value()), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(n_rows), + vector.begin(), + [label, lists_column] __device__(size_type row_index) { + return unbound_list_view{label, lists_column, row_index}; + }); + /* thrust::for_each_n( rmm::exec_policy(stream)->on(stream.value()), thrust::make_counting_iterator(0), @@ -149,6 +160,7 @@ rmm::device_vector list_vector_from_column( [label, lists_column, output = vector.data().get()] __device__(size_type row_index) { output[row_index] = unbound_list_view{label, lists_column, row_index}; }); + */ return vector; } @@ -189,7 +201,7 @@ int32_t get_num_child_rows(cudf::column_view const& list_offsets, rmm::cuda_stre * @return std::pair Child column's null mask and null row count */ std::pair construct_child_nullmask( - rmm::device_vector const& parent_list_vector, + rmm::device_uvector const& parent_list_vector, column_view const& parent_list_offsets, cudf::detail::lists_column_device_view const& source_lists, cudf::detail::lists_column_device_view const& target_lists, @@ -197,7 +209,7 @@ std::pair construct_child_nullmask( rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto is_valid_predicate = [d_list_vector = parent_list_vector.data().get(), + auto is_valid_predicate = [d_list_vector = parent_list_vector.begin(), d_offsets = parent_list_offsets.template data(), d_offsets_size = parent_list_offsets.size(), source_lists, @@ -236,7 +248,7 @@ void print(std::string const& msg, column_view const& col, rmm::cuda_stream_view } void print(std::string const& msg, - rmm::device_vector const& scatter, + rmm::device_uvector const& scatter, rmm::cuda_stream_view stream) { std::cout << msg << " == ["; @@ -337,7 +349,7 @@ struct list_child_constructor { */ template std::enable_if_t(), std::unique_ptr> operator()( - rmm::device_vector const& list_vector, + rmm::device_uvector const& list_vector, cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, @@ -380,7 +392,7 @@ struct list_child_constructor { // Function to copy child-values for specified index of unbound_list_view // to the child column. auto copy_child_values_for_list_index = [d_scattered_lists = - list_vector.data().get(), // unbound_list_view* + list_vector.begin(), // unbound_list_view* d_child_column = child_column->mutable_view().data(), d_offsets = list_offsets.template data(), @@ -435,7 +447,7 @@ struct list_child_constructor { */ template std::enable_if_t::value, std::unique_ptr> operator()( - rmm::device_vector const& list_vector, + rmm::device_uvector const& list_vector, cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, @@ -453,10 +465,9 @@ struct list_child_constructor { auto string_views = rmm::device_vector(num_child_rows); - auto populate_string_views = [d_scattered_lists = - list_vector.data().get(), // unbound_list_view* - d_list_offsets = list_offsets.template data(), - d_string_views = string_views.data().get(), + auto populate_string_views = [d_scattered_lists = list_vector.begin(), // unbound_list_view* + d_list_offsets = list_offsets.template data(), + d_string_views = string_views.data().get(), source_lists, target_lists] __device__(auto const& row_index) { auto unbound_list_view = d_scattered_lists[row_index]; @@ -525,7 +536,7 @@ struct list_child_constructor { */ template std::enable_if_t::value, std::unique_ptr> operator()( - rmm::device_vector const& list_vector, + rmm::device_uvector const& list_vector, cudf::column_view const& list_offsets, cudf::lists_column_view const& source_lists_column_view, cudf::lists_column_view const& target_lists_column_view, @@ -541,14 +552,14 @@ struct list_child_constructor { auto num_child_rows = get_num_child_rows(list_offsets, stream); - auto child_list_views = rmm::device_vector(num_child_rows); + auto child_list_views = rmm::device_uvector(num_child_rows, stream, mr); // Function to convert from parent list_device_view instances to child list_device_views. // For instance, if a parent list_device_view has 3 elements, it should have 3 corresponding // child list_device_view instances. - auto populate_child_list_views = [d_scattered_lists = list_vector.data().get(), + auto populate_child_list_views = [d_scattered_lists = list_vector.begin(), d_list_offsets = list_offsets.template data(), - d_child_list_views = child_list_views.data().get(), + d_child_list_views = child_list_views.begin(), source_lists, target_lists] __device__(auto const& row_index) { auto scattered_row = d_scattered_lists[row_index]; @@ -678,14 +689,18 @@ std::unique_ptr scatter( auto source_lists_column_view = lists_column_view(source); // Checks that this is a list column. auto source_device_view = column_device_view::create(source, stream); - auto source_vector = list_vector_from_column( - unbound_list_view::label_type::SOURCE, lists_column_device_view(*source_device_view), stream); + auto source_vector = list_vector_from_column(unbound_list_view::label_type::SOURCE, + lists_column_device_view(*source_device_view), + stream, + mr); auto target_lists_column_view = lists_column_view(target); // Checks that target is a list column. auto target_device_view = column_device_view::create(target, stream); - auto target_vector = list_vector_from_column( - unbound_list_view::label_type::TARGET, lists_column_device_view(*target_device_view), stream); + auto target_vector = list_vector_from_column(unbound_list_view::label_type::TARGET, + lists_column_device_view(*target_device_view), + stream, + mr); // Scatter. thrust::scatter(rmm::exec_policy(stream)->on(stream.value()), From 2da4f153471d0435733ac5e1ff02050cbcee1c21 Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Mon, 16 Nov 2020 13:30:43 -0800 Subject: [PATCH 07/10] [scatter] Fix documentation. Switch size types to size_type. --- cpp/include/cudf/lists/detail/scatter.cuh | 40 +++++++---------------- 1 file changed, 12 insertions(+), 28 deletions(-) diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 7a455edf55d..fc5d71da5a1 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -152,36 +152,24 @@ rmm::device_uvector list_vector_from_column( return unbound_list_view{label, lists_column, row_index}; }); - /* - thrust::for_each_n( - rmm::exec_policy(stream)->on(stream.value()), - thrust::make_counting_iterator(0), - n_rows, - [label, lists_column, output = vector.data().get()] __device__(size_type row_index) { - output[row_index] = unbound_list_view{label, lists_column, row_index}; - }); - */ - return vector; } /** - * @brief Utility function to fetch the number of rows in a lists column's - * child column, given its offsets column. - * (This is simply the last value in the offsets column.) + * @brief Fetch the number of rows in a lists column's child given its offsets column. * * @param list_offsets Offsets child of a lists column * @param stream The cuda-stream to synchronize on, when reading from device memory - * @return int32_t The last element in the list_offsets column, indicating + * @return cudf::size_type The last element in the list_offsets column, indicating * the number of rows in the lists-column's child. */ -int32_t get_num_child_rows(cudf::column_view const& list_offsets, rmm::cuda_stream_view stream) +cudf::size_type get_num_child_rows(cudf::column_view const& list_offsets, rmm::cuda_stream_view stream) { // Number of rows in child-column == last offset value. - int32_t num_child_rows{}; + cudf::size_type num_child_rows{}; CUDA_TRY(cudaMemcpyAsync(&num_child_rows, - list_offsets.data() + list_offsets.size() - 1, - sizeof(int32_t), + list_offsets.data() + list_offsets.size() - 1, + sizeof(cudf::size_type), cudaMemcpyDeviceToHost, stream.value())); stream.synchronize(); @@ -196,8 +184,8 @@ int32_t get_num_child_rows(cudf::column_view const& list_offsets, rmm::cuda_stre * @param source_lists Source lists column for scatter operation * @param target_lists Target lists column for scatter operation * @param num_child_rows Number of rows in child column - * @param mr Device memory resource used to allocate child column's null mask * @param stream CUDA stream used for device memory operations and kernel launches + * @param mr Device memory resource used to allocate child column's null mask * @return std::pair Child column's null mask and null row count */ std::pair construct_child_nullmask( @@ -303,9 +291,8 @@ void print(std::string const& msg, * a. Child: [8,8,8,2,2,9,9,9,9] <--- THIS * b. Offsets: [0, 3, 5, 9] * - * It is the Expected Child column above that list_child_constructor attempts - * to construct. - * + * `list_child_constructor` constructs the Expected Child column indicated above. + * * `list_child_constructor` expects to be called with the `Source`/`Target` * lists columns, along with the following: * @@ -322,8 +309,7 @@ void print(std::string const& msg, struct list_child_constructor { private: /** - * @brief Function to determine what types are supported as child column types, - * when scattering lists. + * @brief Determine whether the child column type is supported with scattering lists. * * @tparam T The data type of the child column of the list being scattered. */ @@ -334,9 +320,8 @@ struct list_child_constructor { }; public: - /** - * @brief SFINAE catch-all, for unsupported child column types. - */ + + // SFINAE catch-all, for unsupported child column types. template std::enable_if_t::value, std::unique_ptr> operator()( Args&&... args) @@ -363,7 +348,6 @@ struct list_child_constructor { auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view); auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view); - // Number of rows in child-column == last offset value. int32_t num_child_rows{get_num_child_rows(list_offsets, stream)}; auto child_null_mask = From 09325c7c80779f00d1300dced64562279ea12123 Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Mon, 16 Nov 2020 14:07:39 -0800 Subject: [PATCH 08/10] [scatter] Fix documentation. --- cpp/include/cudf/lists/detail/scatter.cuh | 25 ++++++++++------------- 1 file changed, 11 insertions(+), 14 deletions(-) diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index fc5d71da5a1..8075e56f910 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -163,7 +163,8 @@ rmm::device_uvector list_vector_from_column( * @return cudf::size_type The last element in the list_offsets column, indicating * the number of rows in the lists-column's child. */ -cudf::size_type get_num_child_rows(cudf::column_view const& list_offsets, rmm::cuda_stream_view stream) +cudf::size_type get_num_child_rows(cudf::column_view const& list_offsets, + rmm::cuda_stream_view stream) { // Number of rows in child-column == last offset value. cudf::size_type num_child_rows{}; @@ -292,7 +293,7 @@ void print(std::string const& msg, * b. Offsets: [0, 3, 5, 9] * * `list_child_constructor` constructs the Expected Child column indicated above. - * + * * `list_child_constructor` expects to be called with the `Source`/`Target` * lists columns, along with the following: * @@ -309,7 +310,7 @@ void print(std::string const& msg, struct list_child_constructor { private: /** - * @brief Determine whether the child column type is supported with scattering lists. + * @brief Determine whether the child column type is supported with scattering lists. * * @tparam T The data type of the child column of the list being scattered. */ @@ -320,8 +321,7 @@ struct list_child_constructor { }; public: - - // SFINAE catch-all, for unsupported child column types. + // SFINAE catch-all, for unsupported child column types. template std::enable_if_t::value, std::unique_ptr> operator()( Args&&... args) @@ -350,7 +350,7 @@ struct list_child_constructor { int32_t num_child_rows{get_num_child_rows(list_offsets, stream)}; - auto child_null_mask = + auto const child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() ? construct_child_nullmask( list_vector, list_offsets, source_lists, target_lists, num_child_rows, stream, mr) @@ -365,7 +365,6 @@ struct list_child_constructor { print("scatter_rows ", list_vector, stream); #endif // NDEBUG - // Init child-column. auto child_column = cudf::make_fixed_width_column(cudf::data_type{cudf::type_to_id()}, num_child_rows, child_null_mask.first, @@ -373,8 +372,6 @@ struct list_child_constructor { stream.value(), mr); - // Function to copy child-values for specified index of unbound_list_view - // to the child column. auto copy_child_values_for_list_index = [d_scattered_lists = list_vector.begin(), // unbound_list_view* d_child_column = @@ -382,14 +379,14 @@ struct list_child_constructor { d_offsets = list_offsets.template data(), source_lists, target_lists] __device__(auto const& row_index) { - auto unbound_list_row = d_scattered_lists[row_index]; - auto actual_list_row = unbound_list_row.bind_to_column(source_lists, target_lists); + auto const unbound_list_row = d_scattered_lists[row_index]; + auto const actual_list_row = unbound_list_row.bind_to_column(source_lists, target_lists); auto const& bound_column = (unbound_list_row.label() == unbound_list_view::label_type::SOURCE ? source_lists : target_lists); - auto list_begin_offset = + auto const list_begin_offset = bound_column.offsets().element(unbound_list_row.row_index()); - auto list_end_offset = + auto const list_end_offset = bound_column.offsets().element(unbound_list_row.row_index() + 1); #ifndef NDEBUG @@ -406,7 +403,7 @@ struct list_child_constructor { #endif // NDEBUG // Copy all elements in this list row, to "appropriate" offset in child-column. - auto destination_start_offset = d_offsets[row_index]; + auto const destination_start_offset = d_offsets[row_index]; thrust::for_each_n(thrust::seq, thrust::make_counting_iterator(0), actual_list_row.size(), From 4b6a54fb7df8ac7fa19a3c0f5d86e9877ac6dc15 Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Mon, 16 Nov 2020 15:18:30 -0800 Subject: [PATCH 09/10] [scatter] More documentation fixes, const correctness --- cpp/include/cudf/lists/detail/scatter.cuh | 33 +++++++++++---------- cpp/include/cudf/lists/list_device_view.cuh | 7 ++--- 2 files changed, 20 insertions(+), 20 deletions(-) diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index 8075e56f910..b9d8d150da9 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -477,7 +477,6 @@ struct list_child_constructor { d_string_offsets = string_offsets_column.template data(), d_string_chars = string_chars_column.template data()] __device__(auto const& string_idx) { - // auto string_offset = output_start_offset + string_idx; auto string_start_idx = d_string_offsets[input_list_start + string_idx]; auto string_end_idx = d_string_offsets[input_list_start + string_idx + 1]; @@ -644,8 +643,8 @@ void assert_same_data_type(column_view const& lhs, column_view const& rhs) * @tparam SourceIterator must produce list_view objects * @tparam MapIterator must produce index values within the target column. * - * @param mr Device memory resource used to allocate the returned column's device memory * @param stream CUDA stream used for device memory operations and kernel launches. + * @param mr Device memory resource used to allocate the returned column's device memory * @return New lists column. */ template @@ -657,28 +656,25 @@ std::unique_ptr scatter( rmm::cuda_stream_view stream = 0, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { - auto num_rows = target.size(); + auto const num_rows = target.size(); if (num_rows == 0) { return cudf::empty_like(target); } - auto child_column_type = lists_column_view(target).child().type(); + auto const child_column_type = lists_column_view(target).child().type(); assert_same_data_type(source, target); using lists_column_device_view = cudf::detail::lists_column_device_view; using unbound_list_view = cudf::lists::detail::unbound_list_view; - auto source_lists_column_view = lists_column_view(source); // Checks that this is a list column. - auto source_device_view = column_device_view::create(source, stream); - auto source_vector = list_vector_from_column(unbound_list_view::label_type::SOURCE, - lists_column_device_view(*source_device_view), - stream, - mr); + auto const source_device_view = column_device_view::create(source, stream); + auto const source_vector = list_vector_from_column(unbound_list_view::label_type::SOURCE, + lists_column_device_view(*source_device_view), + stream, + mr); - auto target_lists_column_view = - lists_column_view(target); // Checks that target is a list column. - auto target_device_view = column_device_view::create(target, stream); - auto target_vector = list_vector_from_column(unbound_list_view::label_type::TARGET, + auto const target_device_view = column_device_view::create(target, stream); + auto target_vector = list_vector_from_column(unbound_list_view::label_type::TARGET, lists_column_device_view(*target_device_view), stream, mr); @@ -690,6 +686,11 @@ std::unique_ptr scatter( scatter_map_begin, target_vector.begin()); + auto const source_lists_column_view = + lists_column_view(source); // Checks that this is a list column. + auto const target_lists_column_view = + lists_column_view(target); // Checks that target is a list column. + auto list_size_begin = thrust::make_transform_iterator( target_vector.begin(), [] __device__(unbound_list_view l) { return l.size(); }); auto offsets_column = cudf::strings::detail::make_offsets_child_column( @@ -704,8 +705,8 @@ std::unique_ptr scatter( stream, mr); - rmm::device_buffer null_mask{0, stream, mr}; - if (target.has_nulls()) { null_mask = copy_bitmask(target, stream, mr); } + auto null_mask = + target.has_nulls() ? copy_bitmask(target, stream, mr) : rmm::device_buffer{0, stream, mr}; return cudf::make_lists_column(num_rows, std::move(offsets_column), diff --git a/cpp/include/cudf/lists/list_device_view.cuh b/cpp/include/cudf/lists/list_device_view.cuh index 5e904c49cf3..be6bf88da30 100644 --- a/cpp/include/cudf/lists/list_device_view.cuh +++ b/cpp/include/cudf/lists/list_device_view.cuh @@ -36,10 +36,9 @@ class list_device_view { size_type const& row_index) : lists_column(lists_column), _row_index(row_index) { - release_assert(row_index >= 0 && row_index < lists_column.size() && "row_index out of bounds"); - column_device_view const& offsets = lists_column.offsets(); - release_assert(row_index < offsets.size() && "row_index should not have exceeded offset size"); + release_assert(row_index >= 0 && row_index < lists_column.size() && + row_index < offsets.size() && "row_index out of bounds"); begin_offset = offsets.element(row_index); release_assert(begin_offset >= 0 && begin_offset <= lists_column.child().size() && @@ -77,7 +76,7 @@ class list_device_view { } /** - * @brief Fetches the element at the specified index, within the list row. + * @brief Fetches the element at the specified index within the list row. * * @tparam The type of the list's element. * @param The index into the list row From 6931ac8cb1a73b3c7ac9baf114d69eb931f10d9b Mon Sep 17 00:00:00 2001 From: Mithun RK Date: Tue, 17 Nov 2020 14:28:57 -0800 Subject: [PATCH 10/10] [scatter] Optimize for parallel child column construction: WIP: Attempting to replace O(N**2) with a single thrust::for_each_n(). Borked, because of empty lists. Will need closer look. --- cpp/include/cudf/lists/detail/scatter.cuh | 85 +++++++++++++++++++++-- 1 file changed, 80 insertions(+), 5 deletions(-) diff --git a/cpp/include/cudf/lists/detail/scatter.cuh b/cpp/include/cudf/lists/detail/scatter.cuh index b9d8d150da9..8f110523654 100644 --- a/cpp/include/cudf/lists/detail/scatter.cuh +++ b/cpp/include/cudf/lists/detail/scatter.cuh @@ -245,17 +245,63 @@ void print(std::string const& msg, thrust::for_each_n(rmm::exec_policy(stream)->on(stream.value()), thrust::make_counting_iterator(0), scatter.size(), - [s = scatter.data().get()] __device__(auto const& i) { + [s = scatter.begin()] __device__(auto const& i) { auto si = s[i]; printf("%s[%d](%d), ", - (si.label() == unbound_list_view::SOURCE ? "S" : "T"), + (si.label() == unbound_list_view::label_type::SOURCE ? "S" : "T"), si.row_index(), si.size()); }); std::cout << "]" << std::endl; } + +void print(std::string const& msg, + rmm::device_vector const& scatter, + rmm::cuda_stream_view stream) +{ + std::cout << msg << " == ["; + + thrust::for_each_n(rmm::exec_policy(stream)->on(stream.value()), + thrust::make_counting_iterator(0), + scatter.size(), + [s = scatter.data().get()] __device__(auto const& i) { + auto si = s[i]; + printf("%d, ", si); + }); + std::cout << "]" << std::endl; +} #endif // NDEBUG +// Helper to generate mapping between each child row and which list it belongs to. +rmm::device_vector get_child_row_to_list_map(cudf::size_type num_child_rows, + column_view const& list_offsets, + rmm::cuda_stream_view stream) +{ + CUDF_EXPECTS(list_offsets.size() >= 2, "Invalid list offsets."); + + auto scatter_map = cudf::slice(list_offsets, {1, list_offsets.size()-1})[0]; + auto d_scatter_map = scatter_map.data(); + auto ret = rmm::device_vector(static_cast(num_child_rows), 0); + auto scatter_1 = thrust::make_constant_iterator(1); + + thrust::scatter( + rmm::exec_policy(stream)->on(stream.value()), + scatter_1, + scatter_1 + scatter_map.size(), + d_scatter_map, + ret.begin() + ); + + thrust::inclusive_scan( + rmm::exec_policy(stream)->on(stream.value()), + ret.begin(), + ret.end(), + ret.begin() + ); + + return ret; +} + /** * @brief (type_dispatch endpoint) Functor that constructs the child column result * of `scatter()`ing a list column. @@ -348,7 +394,7 @@ struct list_child_constructor { auto source_lists = cudf::detail::lists_column_device_view(*source_column_device_view); auto target_lists = cudf::detail::lists_column_device_view(*target_column_device_view); - int32_t num_child_rows{get_num_child_rows(list_offsets, stream)}; + auto const num_child_rows{get_num_child_rows(list_offsets, stream)}; auto const child_null_mask = source_lists_column_view.child().nullable() || target_lists_column_view.child().nullable() @@ -356,14 +402,17 @@ struct list_child_constructor { list_vector, list_offsets, source_lists, target_lists, num_child_rows, stream, mr) : std::make_pair(rmm::device_buffer{}, 0); -#ifndef NDEBUG + auto const child_row_to_list_mapping = get_child_row_to_list_map(num_child_rows, list_offsets, stream); + +// #ifndef NDEBUG print("list_offsets ", list_offsets, stream); print("source_lists.child() ", source_lists_column_view.child(), stream); print("source_lists.offsets() ", source_lists_column_view.offsets(), stream); print("target_lists.child() ", target_lists_column_view.child(), stream); print("target_lists.offsets() ", target_lists_column_view.offsets(), stream); print("scatter_rows ", list_vector, stream); -#endif // NDEBUG + print("child_row_to_list_mapping ", child_row_to_list_mapping, stream); +// #endif // NDEBUG auto child_column = cudf::make_fixed_width_column(cudf::data_type{cudf::type_to_id()}, num_child_rows, @@ -372,6 +421,30 @@ struct list_child_constructor { stream.value(), mr); + thrust::for_each_n( + rmm::exec_policy(stream)->on(stream.value()), + thrust::make_counting_iterator(0), + num_child_rows, + [ + d_scattered_lists = list_vector.begin(), + d_child_row_to_list_map = child_row_to_list_mapping.data().get(), + d_offsets = list_offsets.template data(), + d_child_column = child_column->mutable_view().data(), + source_lists, + target_lists + ] + __device__ (auto const& child_row_index) + { + auto const list_row_index = d_child_row_to_list_map[child_row_index]; + auto const unbound_list_row = d_scattered_lists[list_row_index]; + auto const bound_list_row = unbound_list_row.bind_to_column(source_lists, target_lists); + if (bound_list_row.size() > 0) { + d_child_column[child_row_index] = bound_list_row.template element(child_row_index - d_offsets[list_row_index]); + } + } + ); + +/* auto copy_child_values_for_list_index = [d_scattered_lists = list_vector.begin(), // unbound_list_view* d_child_column = @@ -420,6 +493,8 @@ struct list_child_constructor { list_vector.size(), copy_child_values_for_list_index); + */ + return std::make_unique(child_column->view()); }