From 0ea6f8ee649579618caa990c38515acdbf9d3775 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Wed, 13 Apr 2022 13:16:42 +0530 Subject: [PATCH] List element Equality comparator (#10289) This PR implements equality comparator for LIST columns. This only supports "self" comparison for now, meaning the two rows to be compared should belong to the same table. A comparator that works on rows of two different tables will be implemented in another PR. This works only on "sanitized" list columns. See #10291 for details. This will partially support #10186. Authors: - Devavret Makkar (https://github.com/devavret) Approvers: - Robert Maynard (https://github.com/robertmaynard) - Vyas Ramasubramani (https://github.com/vyasr) - Mike Wilson (https://github.com/hyperbolic2346) - Jake Hemstad (https://github.com/jrhemstad) - Jordan Jacobelli (https://github.com/Ethyling) URL: https://github.com/rapidsai/cudf/pull/10289 --- conda/recipes/libcudf/meta.yaml | 1 + cpp/benchmarks/CMakeLists.txt | 2 +- cpp/benchmarks/reduction/rank.cpp | 64 ++++ .../cudf/column/column_device_view.cuh | 92 ++++-- cpp/include/cudf/detail/iterator.cuh | 47 ++- cpp/include/cudf/detail/utilities/column.hpp | 84 +++++ cpp/include/cudf/lists/list_device_view.cuh | 33 +- .../cudf/lists/lists_column_device_view.cuh | 57 ++-- .../structs/structs_column_device_view.cuh | 68 ++++ .../cudf/table/experimental/row_operators.cuh | 299 +++++++++++++++++- cpp/src/io/parquet/writer_impl.cu | 79 +---- cpp/src/reductions/scan/rank_scan.cu | 25 +- cpp/src/table/row_operators.cu | 214 +++++++++++-- cpp/tests/CMakeLists.txt | 1 + cpp/tests/reductions/list_rank_test.cpp | 228 +++++++++++++ 15 files changed, 1106 insertions(+), 188 deletions(-) create mode 100644 cpp/benchmarks/reduction/rank.cpp create mode 100644 cpp/include/cudf/detail/utilities/column.hpp create mode 100644 cpp/include/cudf/structs/structs_column_device_view.cuh create mode 100644 cpp/tests/reductions/list_rank_test.cpp diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml index fdd9011ae34..0806bb964cf 100644 --- a/conda/recipes/libcudf/meta.yaml +++ b/conda/recipes/libcudf/meta.yaml @@ -110,6 +110,7 @@ outputs: - test -f $PREFIX/include/cudf/detail/transpose.hpp - test -f $PREFIX/include/cudf/detail/unary.hpp - test -f $PREFIX/include/cudf/detail/utilities/alignment.hpp + - test -f $PREFIX/include/cudf/detail/utilities/column.hpp - test -f $PREFIX/include/cudf/detail/utilities/integer_utils.hpp - test -f $PREFIX/include/cudf/detail/utilities/int_fastdiv.h - test -f $PREFIX/include/cudf/detail/utilities/vector_factories.hpp diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index d863e6e05a9..26bb10da69f 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -181,7 +181,7 @@ ConfigureBench( REDUCTION_BENCH reduction/anyall.cpp reduction/dictionary.cpp reduction/minmax.cpp reduction/reduce.cpp reduction/scan.cpp ) -ConfigureNVBench(REDUCTION_NVBENCH reduction/segment_reduce.cu) +ConfigureNVBench(REDUCTION_NVBENCH reduction/segment_reduce.cu reduction/rank.cpp) # ################################################################################################## # * reduction benchmark --------------------------------------------------------------------------- diff --git a/cpp/benchmarks/reduction/rank.cpp b/cpp/benchmarks/reduction/rank.cpp new file mode 100644 index 00000000000..5e2848d7f0b --- /dev/null +++ b/cpp/benchmarks/reduction/rank.cpp @@ -0,0 +1,64 @@ +/* + * Copyright (c) 2022, 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 + +template +static void nvbench_reduction_scan(nvbench::state& state, nvbench::type_list) +{ + cudf::rmm_pool_raii pool_raii; + + auto const dtype = cudf::type_to_id(); + + double const null_frequency = state.get_float64("null_frequency"); + size_t const size = state.get_int64("data_size"); + + data_profile table_data_profile; + table_data_profile.set_distribution_params(dtype, distribution_id::UNIFORM, 0, 5); + table_data_profile.set_null_frequency(null_frequency); + + auto const table = create_random_table({dtype}, table_size_bytes{size / 2}, table_data_profile); + + auto const new_tbl = cudf::repeat(table->view(), 2); + cudf::column_view input(new_tbl->view().column(0)); + + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { + rmm::cuda_stream_view stream_view{launch.get_stream()}; + auto result = cudf::detail::inclusive_dense_rank_scan( + input, stream_view, rmm::mr::get_current_device_resource()); + }); +} + +using data_type = nvbench::type_list; + +NVBENCH_BENCH_TYPES(nvbench_reduction_scan, NVBENCH_TYPE_AXES(data_type)) + .set_name("rank_scan") + .add_float64_axis("null_frequency", {0, 0.1, 0.5, 0.9}) + .add_int64_axis("data_size", + { + 10000, // 10k + 100000, // 100k + 1000000, // 1M + 10000000, // 10M + 100000000, // 100M + }); \ No newline at end of file diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh index ec3795238b0..070ca80858b 100644 --- a/cpp/include/cudf/column/column_device_view.cuh +++ b/cpp/include/cudf/column/column_device_view.cuh @@ -111,7 +111,7 @@ class alignas(16) column_device_view_base { */ template or is_rep_layout_compatible())> - __host__ __device__ T const* head() const noexcept + [[nodiscard]] CUDF_HOST_DEVICE T const* head() const noexcept { return static_cast(_data); } @@ -132,7 +132,7 @@ class alignas(16) column_device_view_base { * @return T const* Typed pointer to underlying data, including the offset */ template ())> - __host__ __device__ T const* data() const noexcept + [[nodiscard]] CUDF_HOST_DEVICE T const* data() const noexcept { return head() + _offset; } @@ -140,12 +140,12 @@ class alignas(16) column_device_view_base { /** * @brief Returns the number of elements in the column. */ - [[nodiscard]] __host__ __device__ size_type size() const noexcept { return _size; } + [[nodiscard]] CUDF_HOST_DEVICE size_type size() const noexcept { return _size; } /** * @brief Returns the element type */ - [[nodiscard]] __host__ __device__ data_type type() const noexcept { return _type; } + [[nodiscard]] CUDF_HOST_DEVICE data_type type() const noexcept { return _type; } /** * @brief Indicates whether the column can contain null elements, i.e., if it @@ -156,7 +156,7 @@ class alignas(16) column_device_view_base { * @return true The bitmask is allocated * @return false The bitmask is not allocated */ - [[nodiscard]] __host__ __device__ bool nullable() const noexcept { return nullptr != _null_mask; } + [[nodiscard]] CUDF_HOST_DEVICE bool nullable() const noexcept { return nullptr != _null_mask; } /** * @brief Returns raw pointer to the underlying bitmask allocation. @@ -165,7 +165,7 @@ class alignas(16) column_device_view_base { * * @note If `null_count() == 0`, this may return `nullptr`. */ - [[nodiscard]] __host__ __device__ bitmask_type const* null_mask() const noexcept + [[nodiscard]] CUDF_HOST_DEVICE bitmask_type const* null_mask() const noexcept { return _null_mask; } @@ -174,7 +174,7 @@ class alignas(16) column_device_view_base { * @brief Returns the index of the first element relative to the base memory * allocation, i.e., what is returned from `head()`. */ - [[nodiscard]] __host__ __device__ size_type offset() const noexcept { return _offset; } + [[nodiscard]] CUDF_HOST_DEVICE size_type offset() const noexcept { return _offset; } /** * @brief Returns whether the specified element holds a valid value (i.e., not @@ -269,11 +269,11 @@ class alignas(16) column_device_view_base { size_type _offset{}; ///< Index position of the first element. ///< Enables zero-copy slicing - column_device_view_base(data_type type, - size_type size, - void const* data, - bitmask_type const* null_mask, - size_type offset) + CUDF_HOST_DEVICE column_device_view_base(data_type type, + size_type size, + void const* data, + bitmask_type const* null_mask, + size_type offset) : _type{type}, _size{size}, _data{data}, _null_mask{null_mask}, _offset{offset} { } @@ -329,6 +329,33 @@ class alignas(16) column_device_view : public detail::column_device_view_base { */ column_device_view(column_view column, void* h_ptr, void* d_ptr); + /** + * @brief Get a new column_device_view which is a slice of this column. + * + * Example: + * @code{.cpp} + * // column = column_device_view([1, 2, 3, 4, 5, 6, 7]) + * auto c = column.slice(1, 3); + * // c = column_device_view([2, 3, 4]) + * auto c1 = column.slice(2, 3); + * // c1 = column_device_view([3, 4, 5]) + * @endcode + * + * @param offset The index of the first element in the slice + * @param size The number of elements in the slice + */ + [[nodiscard]] CUDF_HOST_DEVICE column_device_view slice(size_type offset, + size_type size) const noexcept + { + return column_device_view{this->type(), + size, + this->head(), + this->null_mask(), + this->offset() + offset, + d_children, + this->num_child_columns()}; + } + /** * @brief Returns reference to element at the specified index. * @@ -346,7 +373,7 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * @param element_index Position of the desired element */ template ())> - __device__ T element(size_type element_index) const noexcept + [[nodiscard]] __device__ T element(size_type element_index) const noexcept { return data()[element_index]; } @@ -365,9 +392,8 @@ class alignas(16) column_device_view : public detail::column_device_view_base { template )> __device__ T element(size_type element_index) const noexcept { - size_type index = element_index + offset(); // account for this view's _offset - const int32_t* d_offsets = - d_children[strings_column_view::offsets_column_index].data(); + size_type index = element_index + offset(); // account for this view's _offset + const auto* d_offsets = d_children[strings_column_view::offsets_column_index].data(); const char* d_strings = d_children[strings_column_view::chars_column_index].data(); size_type offset = d_offsets[index]; return string_view{d_strings + offset, d_offsets[index + 1] - offset}; @@ -763,11 +789,37 @@ class alignas(16) column_device_view : public detail::column_device_view_base { * * @return The number of child columns */ - [[nodiscard]] __host__ __device__ size_type num_child_columns() const noexcept + [[nodiscard]] CUDF_HOST_DEVICE size_type num_child_columns() const noexcept { return _num_children; } + private: + /** + * @brief Creates an instance of this class using pre-existing device memory pointers to data, + * nullmask, and offset. + * + * @param type The type of the column + * @param size The number of elements in the column + * @param data Pointer to the device memory containing the data + * @param null_mask Pointer to the device memory containing the null bitmask + * @param offset The index of the first element in the column + * @param children Pointer to the device memory containing child data + * @param num_children The number of child columns + */ + CUDF_HOST_DEVICE column_device_view(data_type type, + size_type size, + void const* data, + bitmask_type const* null_mask, + size_type offset, + column_device_view* children, + size_type num_children) + : column_device_view_base(type, size, data, null_mask, offset), + d_children(children), + _num_children(num_children) + { + } + protected: column_device_view* d_children{}; ///< Array of `column_device_view` ///< objects in device memory. @@ -852,7 +904,7 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view */ template or is_rep_layout_compatible())> - __host__ __device__ T* head() const noexcept + CUDF_HOST_DEVICE T* head() const noexcept { return const_cast(detail::column_device_view_base::head()); } @@ -870,7 +922,7 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view * @return T* Typed pointer to underlying data, including the offset */ template ())> - __host__ __device__ T* data() const noexcept + CUDF_HOST_DEVICE T* data() const noexcept { return const_cast(detail::column_device_view_base::data()); } @@ -912,7 +964,7 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view * * @note If `null_count() == 0`, this may return `nullptr`. */ - [[nodiscard]] __host__ __device__ bitmask_type* null_mask() const noexcept + [[nodiscard]] CUDF_HOST_DEVICE bitmask_type* null_mask() const noexcept { return const_cast(detail::column_device_view_base::null_mask()); } diff --git a/cpp/include/cudf/detail/iterator.cuh b/cpp/include/cudf/detail/iterator.cuh index 4442af8fab1..7a83298c72a 100644 --- a/cpp/include/cudf/detail/iterator.cuh +++ b/cpp/include/cudf/detail/iterator.cuh @@ -67,7 +67,8 @@ namespace detail { * @return A transform iterator that applies `f` to a counting iterator */ template -inline auto make_counting_transform_iterator(cudf::size_type start, UnaryFunction f) +CUDF_HOST_DEVICE inline auto make_counting_transform_iterator(cudf::size_type start, + UnaryFunction f) { return thrust::make_transform_iterator(thrust::make_counting_iterator(start), f); } @@ -117,26 +118,42 @@ struct null_replaced_value_accessor { /** * @brief validity accessor of column with null bitmask - * A unary functor returns validity at `id`. - * `operator() (cudf::size_type id)` computes validity flag at `id` - * This functor is only allowed for nullable columns. + * A unary functor that returns validity at index `i`. * - * @throws cudf::logic_error if the column is not nullable. + * @tparam safe If false, the accessor with throw logic_error if the column is not nullable. If + * true, the accessor checks for nullability and if col is not nullable, returns true. */ +template struct validity_accessor { column_device_view const col; /** * @brief constructor + * + * @throws cudf::logic_error if not safe and `col` does not have a validity bitmask + * * @param[in] _col column device view of cudf column */ - validity_accessor(column_device_view const& _col) : col{_col} + CUDF_HOST_DEVICE validity_accessor(column_device_view const& _col) : col{_col} { - // verify valid is non-null, otherwise, is_valid() will crash - CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); + if constexpr (not safe) { + // verify col is nullable, otherwise, is_valid_nocheck() will crash +#if defined(__CUDA_ARCH__) + cudf_assert(_col.nullable() && "Unexpected non-nullable column."); +#else + CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); +#endif + } } - __device__ inline bool operator()(cudf::size_type i) const { return col.is_valid_nocheck(i); } + __device__ inline bool operator()(cudf::size_type i) const + { + if constexpr (safe) { + return col.is_valid(i); + } else { + return col.is_valid_nocheck(i); + } + } }; /** @@ -289,16 +306,20 @@ auto make_pair_rep_iterator(column_device_view const& column) * * Dereferencing the returned iterator for element `i` will return the validity * of `column[i]` - * This iterator is only allowed for nullable columns. + * This iterator is only allowed for nullable columns if `safe` = false + * When safe = true, if the column is not nullable then the validity is always true. * - * @throws cudf::logic_error if the column is not nullable. + * @throws cudf::logic_error if the column is not nullable when safe = false * + * @tparam safe If false, the accessor with throw logic_error if the column is not nullable. If + * true, the accessor checks for nullability and if col is not nullable, returns true. * @param column The column to iterate * @return auto Iterator that returns validities of column elements. */ -auto inline make_validity_iterator(column_device_view const& column) +template +CUDF_HOST_DEVICE auto inline make_validity_iterator(column_device_view const& column) { - return make_counting_transform_iterator(cudf::size_type{0}, validity_accessor{column}); + return make_counting_transform_iterator(cudf::size_type{0}, validity_accessor{column}); } /** diff --git a/cpp/include/cudf/detail/utilities/column.hpp b/cpp/include/cudf/detail/utilities/column.hpp new file mode 100644 index 00000000000..7d22bbd60af --- /dev/null +++ b/cpp/include/cudf/detail/utilities/column.hpp @@ -0,0 +1,84 @@ +/* + * Copyright (c) 2022, 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 + +namespace cudf::detail { + +struct linked_column_view; + +using LinkedColPtr = std::shared_ptr; +using LinkedColVector = std::vector; + +/** + * @brief column_view with the added member pointer to the parent of this column. + * + */ +struct linked_column_view : public column_view_base { + linked_column_view(linked_column_view const&) = delete; + linked_column_view& operator=(linked_column_view const&) = delete; + + linked_column_view(column_view const& col) : linked_column_view(nullptr, col) {} + + linked_column_view(linked_column_view* parent, column_view const& col) + : column_view_base(col), parent(parent) + { + children.reserve(col.num_children()); + std::transform( + col.child_begin(), col.child_end(), std::back_inserter(children), [&](column_view const& c) { + return std::make_shared(this, c); + }); + } + + operator column_view() const + { + auto child_it = thrust::make_transform_iterator( + children.begin(), [](auto const& c) { return static_cast(*c); }); + return column_view(this->type(), + this->size(), + this->head(), + this->null_mask(), + UNKNOWN_NULL_COUNT, + this->offset(), + std::vector(child_it, child_it + children.size())); + } + + linked_column_view* parent; //!< Pointer to parent of this column. Nullptr if root + LinkedColVector children; +}; + +/** + * @brief Converts all column_views of a table into linked_column_views + * + * @param table table of columns to convert + * @return Vector of converted linked_column_views + */ +inline LinkedColVector table_to_linked_columns(table_view const& table) +{ + LinkedColVector result; + result.reserve(table.num_columns()); + std::transform(table.begin(), table.end(), std::back_inserter(result), [&](column_view const& c) { + return std::make_shared(c); + }); + + return result; +} + +} // namespace cudf::detail \ No newline at end of file diff --git a/cpp/include/cudf/lists/list_device_view.cuh b/cpp/include/cudf/lists/list_device_view.cuh index ae0a247f005..5cc1e3d166b 100644 --- a/cpp/include/cudf/lists/list_device_view.cuh +++ b/cpp/include/cudf/lists/list_device_view.cuh @@ -16,6 +16,7 @@ #pragma once #include +#include #include #include #include @@ -293,22 +294,34 @@ class list_device_view { * */ struct list_size_functor { - column_device_view const d_column; - CUDF_HOST_DEVICE inline list_size_functor(column_device_view const& d_col) : d_column(d_col) + detail::lists_column_device_view const d_column; + CUDF_HOST_DEVICE inline list_size_functor(detail::lists_column_device_view const& d_col) + : d_column(d_col) { -#if defined(__CUDA_ARCH__) - cudf_assert(d_col.type().id() == type_id::LIST && "Only list type column is supported"); -#else - CUDF_EXPECTS(d_col.type().id() == type_id::LIST, "Only list type column is supported"); -#endif } __device__ inline size_type operator()(size_type idx) { if (d_column.is_null(idx)) return size_type{0}; - auto d_offsets = - d_column.child(lists_column_view::offsets_column_index).data() + d_column.offset(); - return d_offsets[idx + 1] - d_offsets[idx]; + return d_column.offset_at(idx + 1) - d_column.offset_at(idx); } }; +/** + * @brief Makes an iterator that returns size of the list by row index + * + * Example: + * For a list_column_device_view with 3 rows, `l = {[1, 2, 3], [4, 5], [6, 7, 8, 9]}`, + * @code{.cpp} + * auto it = make_list_size_iterator(l); + * assert(it[0] == 3); + * assert(it[1] == 2); + * assert(it[2] == 4); + * @endcode + * + */ +CUDF_HOST_DEVICE auto inline make_list_size_iterator(detail::lists_column_device_view const& c) +{ + return detail::make_counting_transform_iterator(0, list_size_functor{c}); +} + } // 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 e48707ec298..06c20933a70 100644 --- a/cpp/include/cudf/lists/lists_column_device_view.cuh +++ b/cpp/include/cudf/lists/lists_column_device_view.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -25,67 +25,70 @@ namespace cudf { namespace detail { /** - * @brief Given a column-device-view, an instance of this class provides a + * @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 { +class lists_column_device_view : private column_device_view { public: + lists_column_device_view() = delete; ~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_HOST_DEVICE lists_column_device_view(column_device_view const& underlying_) + : column_device_view(underlying_) { +#ifdef __CUDA_ARCH__ + cudf_assert(underlying_.type().id() == type_id::LIST and + "lists_column_device_view only supports lists"); +#else CUDF_EXPECTS(underlying_.type().id() == type_id::LIST, "lists_column_device_view only supports lists"); +#endif } - /** - * @brief Fetches number of rows in the lists column - */ - [[nodiscard]] CUDF_HOST_DEVICE inline cudf::size_type size() const { return underlying.size(); } + using column_device_view::is_null; + using column_device_view::nullable; + using column_device_view::offset; + using column_device_view::size; /** * @brief Fetches the offsets column of the underlying list column. */ [[nodiscard]] __device__ inline column_device_view offsets() const { - return underlying.child(lists_column_view::offsets_column_index); + return column_device_view::child(lists_column_view::offsets_column_index); } /** - * @brief Fetches the child column of the underlying list column. + * @brief Fetches the list offset value at a given row index while taking column offset into + * account. */ - [[nodiscard]] __device__ inline column_device_view child() const + [[nodiscard]] __device__ inline size_type offset_at(size_type idx) const { - return underlying.child(lists_column_view::child_column_index); + return offsets().size() > 0 ? offsets().element(offset() + idx) : 0; } /** - * @brief Indicates whether the list column is nullable. - */ - [[nodiscard]] __device__ inline bool nullable() const { return underlying.nullable(); } - - /** - * @brief Indicates whether the row (i.e. list) at the specified - * index is null. + * @brief Fetches the child column of the underlying list column. */ - [[nodiscard]] __device__ inline bool is_null(size_type idx) const + [[nodiscard]] __device__ inline column_device_view child() const { - return underlying.is_null(idx); + return column_device_view::child(lists_column_view::child_column_index); } /** - * @brief Fetches the offset of the underlying column_device_view, - * in case it is a sliced/offset column. + * @brief Fetches the child column of the underlying list column with offset and size applied */ - [[nodiscard]] __device__ inline size_type offset() const { return underlying.offset(); } - - private: - column_device_view underlying; + [[nodiscard]] __device__ inline column_device_view sliced_child() const + { + auto start = offset_at(0); + auto end = offset_at(size()); + return child().slice(start, end - start); + } }; } // namespace detail diff --git a/cpp/include/cudf/structs/structs_column_device_view.cuh b/cpp/include/cudf/structs/structs_column_device_view.cuh new file mode 100644 index 00000000000..09bbb46a93c --- /dev/null +++ b/cpp/include/cudf/structs/structs_column_device_view.cuh @@ -0,0 +1,68 @@ +/* + * Copyright (c) 2022, 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 + +namespace cudf { + +namespace detail { + +/** + * @brief Given a column_device_view, an instance of this class provides a + * wrapper on this compound column for struct operations. + * Analogous to struct_column_view. + */ +class structs_column_device_view : private column_device_view { + public: + structs_column_device_view() = delete; + ~structs_column_device_view() = default; + structs_column_device_view(structs_column_device_view const&) = default; + structs_column_device_view(structs_column_device_view&&) = default; + structs_column_device_view& operator=(structs_column_device_view const&) = default; + structs_column_device_view& operator=(structs_column_device_view&&) = default; + + CUDF_HOST_DEVICE structs_column_device_view(column_device_view const& underlying_) + : column_device_view(underlying_) + { +#ifdef __CUDA_ARCH__ + cudf_assert(underlying_.type().id() == type_id::STRUCT and + "structs_column_device_view only supports structs"); +#else + CUDF_EXPECTS(underlying_.type().id() == type_id::STRUCT, + "structs_column_device_view only supports structs"); +#endif + } + + using column_device_view::child; + using column_device_view::is_null; + using column_device_view::nullable; + using column_device_view::offset; + using column_device_view::size; + + /** + * @brief Fetches the child column of the underlying struct column. + */ + [[nodiscard]] __device__ inline column_device_view sliced_child(size_type idx) const + { + return child(idx).slice(offset(), size()); + } +}; + +} // namespace detail + +} // namespace cudf diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 0fb1ad7ca68..88e31744fdf 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -17,15 +17,20 @@ #pragma once #include +#include #include #include +#include +#include #include +#include #include #include #include #include #include +#include #include #include @@ -172,13 +177,11 @@ class device_row_comparator { template () and - not std::is_same_v)> - __device__ cuda::std::pair operator()(size_type const lhs_element_index, - size_type const rhs_element_index) + not std::is_same_v), + typename... Args> + __device__ cuda::std::pair operator()(Args...) { - // TODO: make this CUDF_UNREACHABLE - cudf_assert(false && "Attempted to compare elements of uncomparable types."); - return cuda::std::make_pair(weak_ordering::LESS, std::numeric_limits::max()); + CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types."); } template )> @@ -424,6 +427,290 @@ class self_comparator { }; } // namespace lexicographic + +namespace equality { + +template +class device_row_comparator { + friend class self_comparator; + + public: + /** + * @brief Checks whether the row at `lhs_index` in the `lhs` table is equal to the row at + * `rhs_index` in the `rhs` table. + * + * @param lhs_index The index of row in the `lhs` table to examine + * @param rhs_index The index of the row in the `rhs` table to examine + * @return `true` if row from the `lhs` table is equal to the row in the `rhs` table + */ + __device__ bool operator()(size_type const lhs_index, size_type const rhs_index) const noexcept + { + auto equal_elements = [=](column_device_view l, column_device_view r) { + return cudf::type_dispatcher( + l.type(), element_comparator{nulls, l, r, nulls_are_equal}, lhs_index, rhs_index); + }; + + return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements); + } + + private: + /** + * @brief Construct a function object for performing equality comparison between the rows of two + * tables. + * + * @param has_nulls Indicates if either input table contains columns with nulls. + * @param lhs The first table + * @param rhs The second table (may be the same table as `lhs`) + * @param nulls_are_equal Indicates if two null elements are treated as equivalent + */ + device_row_comparator(Nullate has_nulls, + table_device_view lhs, + table_device_view rhs, + null_equality nulls_are_equal = null_equality::EQUAL) noexcept + : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal} + { + } + + /** + * @brief Performs an equality comparison between two elements in two columns. + * + * @tparam Nullate A cudf::nullate type describing how to check for nulls. + */ + class element_comparator { + public: + /** + * @brief Construct type-dispatched function object for comparing equality + * between two elements. + * + * @note `lhs` and `rhs` may be the same. + * + * @param has_nulls Indicates if either input column contains nulls. + * @param lhs The column containing the first element + * @param rhs The column containing the second element (may be the same as lhs) + * @param nulls_are_equal Indicates if two null elements are treated as equivalent + */ + __device__ element_comparator(Nullate has_nulls, + column_device_view lhs, + column_device_view rhs, + null_equality nulls_are_equal = null_equality::EQUAL) noexcept + : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal} + { + } + + /** + * @brief Compares the specified elements for equality. + * + * @param lhs_element_index The index of the first element + * @param rhs_element_index The index of the second element + * @return True if lhs and rhs are equal or if both lhs and rhs are null and nulls are + * configured to be considered equal (`nulls_are_equal` == `null_equality::EQUAL`) + */ + template ())> + __device__ bool operator()(size_type const lhs_element_index, + size_type const rhs_element_index) const noexcept + { + if (nulls) { + bool const lhs_is_null{lhs.is_null(lhs_element_index)}; + bool const rhs_is_null{rhs.is_null(rhs_element_index)}; + if (lhs_is_null and rhs_is_null) { + return nulls_are_equal == null_equality::EQUAL; + } else if (lhs_is_null != rhs_is_null) { + return false; + } + } + + return equality_compare(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); + } + + template () and + not cudf::is_nested()), + typename... Args> + __device__ bool operator()(Args...) + { + CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types."); + } + + template ())> + __device__ bool operator()(size_type const lhs_element_index, + size_type const rhs_element_index) const noexcept + { + column_device_view lcol = lhs.slice(lhs_element_index, 1); + column_device_view rcol = rhs.slice(rhs_element_index, 1); + while (is_nested(lcol.type())) { + if (nulls) { + auto lvalid = detail::make_validity_iterator(lcol); + auto rvalid = detail::make_validity_iterator(rcol); + if (nulls_are_equal == null_equality::UNEQUAL) { + if (thrust::any_of( + thrust::seq, lvalid, lvalid + lcol.size(), thrust::logical_not()) or + thrust::any_of( + thrust::seq, rvalid, rvalid + rcol.size(), thrust::logical_not())) { + return false; + } + } else { + if (not thrust::equal(thrust::seq, lvalid, lvalid + lcol.size(), rvalid)) { + return false; + } + } + } + if (lcol.type().id() == type_id::STRUCT) { + if (lcol.num_child_columns() == 0) { return true; } + lcol = detail::structs_column_device_view(lcol).sliced_child(0); + rcol = detail::structs_column_device_view(rcol).sliced_child(0); + } else if (lcol.type().id() == type_id::LIST) { + auto l_list_col = detail::lists_column_device_view(lcol); + auto r_list_col = detail::lists_column_device_view(rcol); + + auto lsizes = make_list_size_iterator(l_list_col); + auto rsizes = make_list_size_iterator(r_list_col); + if (not thrust::equal(thrust::seq, lsizes, lsizes + lcol.size(), rsizes)) { + return false; + } + + lcol = l_list_col.sliced_child(); + rcol = r_list_col.sliced_child(); + if (lcol.size() != rcol.size()) { return false; } + } + } + + auto comp = + column_comparator{element_comparator{nulls, lcol, rcol, nulls_are_equal}, lcol.size()}; + return type_dispatcher(lcol.type(), comp); + } + + private: + /** + * @brief Serially compare two columns for equality. + * + * When we want to get the equivalence of two columns by serially comparing all elements in a + * one column with the corresponding elements in the other column, this saves us from type + * dispatching for each individual element in the range + */ + struct column_comparator { + element_comparator const comp; + size_type const size; + + /** + * @brief Serially compare two columns for equality. + * + * @return True if ALL elements compare equal, false otherwise + */ + template ())> + __device__ bool operator()() const noexcept + { + return thrust::all_of(thrust::seq, + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(0) + size, + [=](auto i) { return comp.template operator()(i, i); }); + } + + template ()), + typename... Args> + __device__ bool operator()(Args...) const noexcept + { + CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types."); + } + }; + + column_device_view const lhs; + column_device_view const rhs; + Nullate const nulls; + null_equality const nulls_are_equal; + }; + + table_device_view const lhs; + table_device_view const rhs; + Nullate const nulls; + null_equality const nulls_are_equal; +}; + +struct preprocessed_table { + /** + * @brief Preprocess table for use with row equality comparison or row hashing + * + * Sets up the table for use with row equality comparison or row hashing. The resulting + * preprocessed table can be passed to the constructor of `equality::self_comparator` to + * avoid preprocessing again. + * + * @param table The table to preprocess + * @param stream The cuda stream to use while preprocessing. + */ + static std::shared_ptr create(table_view const& table, + rmm::cuda_stream_view stream); + + private: + friend class self_comparator; + + using table_device_view_owner = + std::invoke_result_t; + + preprocessed_table(table_device_view_owner&& table, + std::vector&& null_buffers) + : _t(std::move(table)), _null_buffers(std::move(null_buffers)) + { + } + + /** + * @brief Implicit conversion operator to a `table_device_view` of the preprocessed table. + * + * @return table_device_view + */ + operator table_device_view() { return *_t; } + + table_device_view_owner _t; + std::vector _null_buffers; +}; + +class self_comparator { + public: + /** + * @brief Construct an owning object for performing equality comparisons between two rows of the + * same table. + * + * @param t The table to compare + * @param stream The stream to construct this object on. Not the stream that will be used for + * comparisons using this object. + */ + self_comparator(table_view const& t, rmm::cuda_stream_view stream) + : d_t(preprocessed_table::create(t, stream)) + { + } + + /** + * @brief Construct an owning object for performing equality comparisons between two rows of the + * same table. + * + * This constructor allows independently constructing a `preprocessed_table` and sharing it among + * multiple comparators. + * + * @param t A table preprocessed for equality comparison + */ + self_comparator(std::shared_ptr t) : d_t{std::move(t)} {} + + /** + * @brief Get the comparison operator to use on the device + * + * Returns a binary callable, `F`, with signature `bool F(size_t, size_t)`. + * + * `F(i,j)` returns true if and only if row `i` compares equal to row `j`. + * + * @tparam Nullate Optional, A cudf::nullate type describing how to check for nulls. + */ + template + device_row_comparator device_comparator(Nullate nullate = {}) const + { + return device_row_comparator(nullate, *d_t, *d_t); + } + + private: + std::shared_ptr d_t; +}; + +} // namespace equality + } // namespace row } // namespace experimental } // namespace cudf diff --git a/cpp/src/io/parquet/writer_impl.cu b/cpp/src/io/parquet/writer_impl.cu index 70a594423c9..cb1acb4d9ec 100644 --- a/cpp/src/io/parquet/writer_impl.cu +++ b/cpp/src/io/parquet/writer_impl.cu @@ -30,6 +30,7 @@ #include #include +#include #include #include #include @@ -190,55 +191,6 @@ struct aggregate_writer_metadata { uint32_t column_order_listsize = 0; }; -struct linked_column_view; - -using LinkedColPtr = std::shared_ptr; -using LinkedColVector = std::vector; - -/** - * @brief column_view with the added member pointer to the parent of this column. - * - */ -struct linked_column_view : public column_view { - // TODO(cp): we are currently keeping all column_view children info multiple times - once for each - // copy of this object. Options: - // 1. Inherit from column_view_base. Only lose out on children vector. That is not needed. - // 2. Don't inherit at all. make linked_column_view keep a reference wrapper to its column_view - linked_column_view(column_view const& col) : column_view(col), parent(nullptr) - { - for (auto child_it = col.child_begin(); child_it < col.child_end(); ++child_it) { - children.push_back(std::make_shared(this, *child_it)); - } - } - - linked_column_view(linked_column_view* parent, column_view const& col) - : column_view(col), parent(parent) - { - for (auto child_it = col.child_begin(); child_it < col.child_end(); ++child_it) { - children.push_back(std::make_shared(this, *child_it)); - } - } - - linked_column_view* parent; //!< Pointer to parent of this column. Nullptr if root - LinkedColVector children; -}; - -/** - * @brief Converts all column_views of a table into linked_column_views - * - * @param table table of columns to convert - * @return Vector of converted linked_column_views - */ -LinkedColVector input_table_to_linked_columns(table_view const& table) -{ - LinkedColVector result; - for (column_view const& col : table) { - result.emplace_back(std::make_shared(col)); - } - - return result; -} - /** * @brief Extends SchemaElement to add members required in constructing parquet_column_view * @@ -250,7 +202,7 @@ LinkedColVector input_table_to_linked_columns(table_view const& table) * supported types */ struct schema_tree_node : public SchemaElement { - LinkedColPtr leaf_column; + cudf::detail::LinkedColPtr leaf_column; statistics_dtype stats_dtype; int32_t ts_scale; @@ -262,7 +214,7 @@ struct schema_tree_node : public SchemaElement { struct leaf_schema_fn { schema_tree_node& col_schema; - LinkedColPtr const& col; + cudf::detail::LinkedColPtr const& col; column_in_metadata const& col_meta; bool timestamp_is_int96; @@ -494,7 +446,7 @@ struct leaf_schema_fn { } }; -inline bool is_col_nullable(LinkedColPtr const& col, +inline bool is_col_nullable(cudf::detail::LinkedColPtr const& col, column_in_metadata const& col_meta, bool single_write_mode) { @@ -520,10 +472,11 @@ inline bool is_col_nullable(LinkedColPtr const& col, * Recursively traverses through linked_columns and corresponding metadata to construct schema tree. * The resulting schema tree is stored in a vector in pre-order traversal order. */ -std::vector construct_schema_tree(LinkedColVector const& linked_columns, - table_input_metadata& metadata, - bool single_write_mode, - bool int96_timestamps) +std::vector construct_schema_tree( + cudf::detail::LinkedColVector const& linked_columns, + table_input_metadata& metadata, + bool single_write_mode, + bool int96_timestamps) { std::vector schema; schema_tree_node root{}; @@ -534,8 +487,8 @@ std::vector construct_schema_tree(LinkedColVector const& linke root.parent_idx = -1; // root schema has no parent schema.push_back(std::move(root)); - std::function add_schema = - [&](LinkedColPtr const& col, column_in_metadata& col_meta, size_t parent_idx) { + std::function add_schema = + [&](cudf::detail::LinkedColPtr const& col, column_in_metadata& col_meta, size_t parent_idx) { bool col_nullable = is_col_nullable(col, col_meta, single_write_mode); if (col->type().id() == type_id::STRUCT) { @@ -545,7 +498,7 @@ std::vector construct_schema_tree(LinkedColVector const& linke col_nullable ? FieldRepetitionType::OPTIONAL : FieldRepetitionType::REQUIRED; struct_schema.name = (schema[parent_idx].name == "list") ? "element" : col_meta.get_name(); - struct_schema.num_children = col->num_children(); + struct_schema.num_children = col->children.size(); struct_schema.parent_idx = parent_idx; schema.push_back(std::move(struct_schema)); @@ -553,7 +506,7 @@ std::vector construct_schema_tree(LinkedColVector const& linke // for (auto child_it = col->children.begin(); child_it < col->children.end(); child_it++) { // add_schema(*child_it, struct_node_index); // } - CUDF_EXPECTS(col->num_children() == static_cast(col_meta.num_children()), + CUDF_EXPECTS(col->children.size() == static_cast(col_meta.num_children()), "Mismatch in number of child columns between input table and metadata"); for (size_t i = 0; i < col->children.size(); ++i) { add_schema(col->children[i], col_meta.child(i), struct_node_index); @@ -592,7 +545,7 @@ std::vector construct_schema_tree(LinkedColVector const& linke // "col_name" : { "key_value" : { "key", "value" } } // verify the List child structure is a struct - auto const& struct_col = col->child(lists_column_view::child_column_index); + column_view struct_col = *col->children[lists_column_view::child_column_index]; CUDF_EXPECTS(struct_col.type().id() == type_id::STRUCT, "Map should be a List of struct"); CUDF_EXPECTS(struct_col.num_children() == 2, "Map should be a List of struct with two children only but found " + @@ -740,7 +693,7 @@ parquet_column_view::parquet_column_view(schema_tree_node const& schema_node, // For list columns, we still need to retain the offset child column. auto children = (parent.type().id() == type_id::LIST) - ? std::vector{parent.child(lists_column_view::offsets_column_index), + ? std::vector{*parent.children[lists_column_view::offsets_column_index], single_inheritance_cudf_col} : std::vector{single_inheritance_cudf_col}; @@ -1221,7 +1174,7 @@ void writer::impl::write(table_view const& table, std::vector co add_default_name(table_meta->column_metadata[i], "_col" + std::to_string(i)); } - auto vec = input_table_to_linked_columns(table); + auto vec = table_to_linked_columns(table); auto schema_tree = construct_schema_tree(vec, *table_meta, single_write_mode, int96_timestamps); // Construct parquet_column_views from the schema tree leaf nodes. std::vector parquet_columns; diff --git a/cpp/src/reductions/scan/rank_scan.cu b/cpp/src/reductions/scan/rank_scan.cu index 67b4b594f2e..521f8e2d06f 100644 --- a/cpp/src/reductions/scan/rank_scan.cu +++ b/cpp/src/reductions/scan/rank_scan.cu @@ -19,7 +19,7 @@ #include #include #include -#include +#include #include #include @@ -38,7 +38,6 @@ namespace { * @tparam value_resolver flag value resolver with boolean first and row number arguments * @tparam scan_operator scan function ran on the flag values * @param order_by input column to generate ranks for - * @param has_nulls if the order_by column has nested nulls * @param resolver flag value resolver * @param scan_op scan operation ran on the flag results * @param stream CUDA stream used for device memory operations and kernel launches @@ -47,28 +46,22 @@ namespace { */ template std::unique_ptr rank_generator(column_view const& order_by, - bool has_nulls, value_resolver resolver, scan_operator scan_op, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto const flattened = cudf::structs::detail::flatten_nested_columns( - table_view{{order_by}}, {}, {}, structs::detail::column_nullability::MATCH_INCOMING); - auto const d_flat_order = table_device_view::create(flattened, stream); - row_equality_comparator comparator( - nullate::DYNAMIC{has_nulls}, *d_flat_order, *d_flat_order, null_equality::EQUAL); - auto ranks = make_fixed_width_column(data_type{type_to_id()}, - flattened.flattened_columns().num_rows(), - mask_state::UNALLOCATED, - stream, - mr); + auto comp = cudf::experimental::row::equality::self_comparator(table_view{{order_by}}, stream); + auto const device_comparator = + comp.device_comparator(nullate::DYNAMIC{has_nested_nulls(table_view({order_by}))}); + auto ranks = make_fixed_width_column( + data_type{type_to_id()}, order_by.size(), mask_state::UNALLOCATED, stream, mr); auto mutable_ranks = ranks->mutable_view(); thrust::tabulate(rmm::exec_policy(stream), mutable_ranks.begin(), mutable_ranks.end(), - [comparator, resolver] __device__(size_type row_index) { + [comparator = device_comparator, resolver] __device__(size_type row_index) { return resolver(row_index == 0 || !comparator(row_index, row_index - 1), row_index); }); @@ -87,11 +80,8 @@ std::unique_ptr inclusive_dense_rank_scan(column_view const& order_by, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - CUDF_EXPECTS(!cudf::structs::detail::is_or_has_nested_lists(order_by), - "Unsupported list type in dense_rank scan."); return rank_generator( order_by, - has_nested_nulls(table_view{{order_by}}), [] __device__(bool const unequal, size_type const) { return unequal ? 1 : 0; }, DeviceSum{}, stream, @@ -106,7 +96,6 @@ std::unique_ptr inclusive_rank_scan(column_view const& order_by, "Unsupported list type in rank scan."); return rank_generator( order_by, - has_nested_nulls(table_view{{order_by}}), [] __device__(bool unequal, auto row_index) { return unequal ? row_index + 1 : 0; }, DeviceMax{}, stream, diff --git a/cpp/src/table/row_operators.cu b/cpp/src/table/row_operators.cu index a0400133c68..408d4e51425 100644 --- a/cpp/src/table/row_operators.cu +++ b/cpp/src/table/row_operators.cu @@ -17,7 +17,9 @@ #include #include #include +#include #include +#include #include #include @@ -28,6 +30,59 @@ namespace experimental { namespace { +/** + * @brief Applies the offsets of struct column onto its children + * + * @param c The column whose children are to be sliced + * @return Children of `c` with offsets applied + */ +std::vector slice_children(column_view const& c) +{ + if (c.type().id() == type_id::STRUCT) { + std::vector sliced_children; + sliced_children.reserve(c.num_children()); + auto struct_col = structs_column_view(c); + for (size_type i = 0; i < struct_col.num_children(); ++i) { + auto sliced = struct_col.get_sliced_child(i); + // We cannot directly use the output of `structs_column_view::get_sliced_child` because we + // must first traverse its children recursively to push offsets all the way down to the leaf + // children. + sliced_children.emplace_back(sliced.type(), + sliced.size(), + sliced.head(), + sliced.null_mask(), + sliced.null_count(), + sliced.offset(), + slice_children(sliced)); + } + return sliced_children; + } + return {c.child_begin(), c.child_end()}; +}; + +/** + * @brief Applies the offsets of struct columns in a table onto their children. + * + * Given a table, this replaces any struct columns with similar struct columns that have their + * offsets applied to their children. Structs that are children of list columns are not affected. + * + */ +table_view pushdown_struct_offsets(table_view table) +{ + std::vector cols; + cols.reserve(table.num_columns()); + std::transform(table.begin(), table.end(), std::back_inserter(cols), [&](column_view const& c) { + return column_view(c.type(), + c.size(), + c.head(), + c.null_mask(), + c.null_count(), + c.offset(), + slice_children(c)); + }); + return table_view(cols); +} + /** * @brief Decompose all struct columns in a table * @@ -39,33 +94,60 @@ namespace { * non-decomposed table, which are pruned during decomposition. * * For example, if the original table has a column `Struct, decimal>`, + * * S1 * / \ * S2 d * / \ * i f + * * then after decomposition, we get three columns: * `Struct>`, `float`, and `decimal`. - * 0 2 1 <- depths - * S1 - * | - * S2 d - * | - * i f + * + * 0 2 1 <- depths + * S1 + * | + * S2 d + * | + * i f + * * The depth of the first column is 0 because it contains all its parent levels, while the depth * of the second column is 2 because two of its parent struct levels were pruned. * - * Similarly, a struct column of type Struct> is decomposed as follows + * Similarly, a struct column of type Struct> is decomposed as follows + * * S1 * / \ * i S2 * / \ * f d * - * 0 1 2 <- depths - * S1 S2 d - * | | - * i f + * 0 1 2 <- depths + * S1 S2 d + * | | + * i f + * + * When list columns are present, the decomposition is performed similarly to pure structs but list + * parent columns are NOT pruned + * + * For example, if the original table has a column `List>`, + * + * L + * | + * S + * / \ + * i f + * + * after decomposition, we get two columns + * + * L L + * | | + * S f + * | + * i + * + * The list parents are still needed to define the range of elements in the leaf that belong to the + * same row. * * @param table The table whose struct columns to decompose. * @param column_order The per-column order if using output with lexicographic comparison @@ -77,26 +159,34 @@ auto decompose_structs(table_view table, host_span column_order = {}, host_span null_precedence = {}) { + auto sliced = pushdown_struct_offsets(table); + auto linked_columns = detail::table_to_linked_columns(sliced); + std::vector verticalized_columns; std::vector new_column_order; std::vector new_null_precedence; std::vector verticalized_col_depths; - for (size_type col_idx = 0; col_idx < table.num_columns(); ++col_idx) { - auto const& col = table.column(col_idx); - if (is_nested(col.type())) { + for (size_t col_idx = 0; col_idx < linked_columns.size(); ++col_idx) { + detail::linked_column_view const* col = linked_columns[col_idx].get(); + if (is_nested(col->type())) { // convert and insert - std::vector> flattened; - std::function*, int)> recursive_child = - [&](column_view const& c, std::vector* branch, int depth) { + std::vector> flattened; + std::function*, int)> + recursive_child = [&](detail::linked_column_view const* c, + std::vector* branch, + int depth) { branch->push_back(c); - if (c.type().id() == type_id::STRUCT) { - for (int child_idx = 0; child_idx < c.num_children(); ++child_idx) { - auto scol = structs_column_view(c); + if (c->type().id() == type_id::LIST) { + recursive_child( + c->children[lists_column_view::child_column_index].get(), branch, depth + 1); + } else if (c->type().id() == type_id::STRUCT) { + for (size_t child_idx = 0; child_idx < c->children.size(); ++child_idx) { if (child_idx > 0) { verticalized_col_depths.push_back(depth + 1); branch = &flattened.emplace_back(); } - recursive_child(scol.get_sliced_child(child_idx), branch, depth + 1); + recursive_child(c->children[child_idx].get(), branch, depth + 1); } } }; @@ -105,17 +195,39 @@ auto decompose_structs(table_view table, recursive_child(col, &branch, 0); for (auto const& branch : flattened) { - column_view curr_col = branch.back(); + column_view temp_col = *branch.back(); for (auto it = branch.crbegin() + 1; it < branch.crend(); ++it) { - curr_col = column_view(it->type(), - it->size(), + auto const& prev_col = *(*it); + auto children = + (prev_col.type().id() == type_id::LIST) + ? std::vector{*prev_col + .children[lists_column_view::offsets_column_index], + temp_col} + : std::vector{temp_col}; + temp_col = column_view(prev_col.type(), + prev_col.size(), nullptr, - it->null_mask(), + prev_col.null_mask(), UNKNOWN_NULL_COUNT, - it->offset(), - {curr_col}); + prev_col.offset(), + std::move(children)); + } + // Traverse upward and include any list columns in the ancestors + for (detail::linked_column_view* parent = branch.front()->parent; parent; + parent = parent->parent) { + if (parent->type().id() == type_id::LIST) { + // Include this parent + temp_col = column_view( + parent->type(), + parent->size(), + nullptr, // list has no data of its own + nullptr, // If we're going through this then nullmask is already in another branch + UNKNOWN_NULL_COUNT, + parent->offset(), + {*parent->children[lists_column_view::offsets_column_index], temp_col}); + } } - verticalized_columns.push_back(curr_col); + verticalized_columns.push_back(temp_col); } if (not column_order.empty()) { new_column_order.insert(new_column_order.end(), flattened.size(), column_order[col_idx]); @@ -125,7 +237,7 @@ auto decompose_structs(table_view table, new_null_precedence.end(), flattened.size(), null_precedence[col_idx]); } } else { - verticalized_columns.push_back(col); + verticalized_columns.push_back(*col); verticalized_col_depths.push_back(0); if (not column_order.empty()) { new_column_order.push_back(column_order[col_idx]); } if (not null_precedence.empty()) { new_null_precedence.push_back(null_precedence[col_idx]); } @@ -137,6 +249,8 @@ auto decompose_structs(table_view table, std::move(verticalized_col_depths)); } +using column_checker_fn_t = std::function; + /** * @brief Check a table for compatibility with lexicographic comparison * @@ -145,7 +259,7 @@ auto decompose_structs(table_view table, void check_lex_compatibility(table_view const& input) { // Basically check if there's any LIST hiding anywhere in the table - std::function check_column = [&](column_view const& c) { + column_checker_fn_t check_column = [&](column_view const& c) { CUDF_EXPECTS(c.type().id() != type_id::LIST, "Cannot lexicographic compare a table with a LIST column"); if (not is_nested(c.type())) { @@ -162,6 +276,28 @@ void check_lex_compatibility(table_view const& input) } } +/** + * @brief Check a table for compatibility with equality comparison + * + * Checks whether a given table contains columns of non-equality comparable types. + */ +void check_eq_compatibility(table_view const& input) +{ + column_checker_fn_t check_column = [&](column_view const& c) { + if (not is_nested(c.type())) { + CUDF_EXPECTS(is_equality_comparable(c.type()), + "Cannot compare equality for a table with a column of type " + + jit::get_type_name(c.type())); + } + for (auto child = c.child_begin(); child < c.child_end(); ++child) { + check_column(*child); + } + }; + for (column_view const& c : input) { + check_column(c); + } +} + } // namespace namespace row { @@ -189,6 +325,24 @@ std::shared_ptr preprocessed_table::create( } } // namespace lexicographic + +namespace equality { + +std::shared_ptr preprocessed_table::create(table_view const& t, + rmm::cuda_stream_view stream) +{ + check_eq_compatibility(t); + + auto null_pushed_table = structs::detail::superimpose_parent_nulls(t, stream); + auto [verticalized_lhs, _, __, ___] = decompose_structs(std::get<0>(null_pushed_table)); + + auto d_t = table_device_view_owner(table_device_view::create(verticalized_lhs, stream)); + return std::shared_ptr( + new preprocessed_table(std::move(d_t), std::move(std::get<1>(null_pushed_table)))); +} + +} // namespace equality + } // namespace row } // namespace experimental } // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 1ed921d1f08..e016f47616b 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -156,6 +156,7 @@ ConfigureTest( reductions/reduction_tests.cpp reductions/scan_tests.cpp reductions/segmented_reduction_tests.cpp + reductions/list_rank_test.cpp reductions/tdigest_tests.cu ) diff --git a/cpp/tests/reductions/list_rank_test.cpp b/cpp/tests/reductions/list_rank_test.cpp new file mode 100644 index 00000000000..d263677f23b --- /dev/null +++ b/cpp/tests/reductions/list_rank_test.cpp @@ -0,0 +1,228 @@ +/* + * Copyright (c) 2022, 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 "benchmarks/common/generate_input.hpp" +#include +#include +#include + +struct ListRankScanTest : public cudf::test::BaseFixture { + inline void test_ungrouped_rank_scan(cudf::column_view const& input, + cudf::column_view const& expect_vals, + std::unique_ptr const& agg, + cudf::null_policy null_handling) + { + auto col_out = cudf::scan(input, agg, cudf::scan_type::INCLUSIVE, null_handling); + CUDF_TEST_EXPECT_COLUMNS_EQUIVALENT( + expect_vals, col_out->view(), cudf::test::debug_output_level::ALL_ERRORS); + } +}; + +TEST_F(ListRankScanTest, BasicList) +{ + using lcw = cudf::test::lists_column_wrapper; + auto const col = lcw{{}, {}, {1}, {1, 1}, {1}, {1, 2}, {2, 2}, {2}, {2}, {2, 1}, {2, 2}, {2, 2}}; + + auto const expected_dense_vals = + cudf::test::fixed_width_column_wrapper{1, 1, 2, 3, 4, 5, 6, 7, 7, 8, 9, 9}; + this->test_ungrouped_rank_scan(col, + expected_dense_vals, + cudf::make_dense_rank_aggregation(), + cudf::null_policy::INCLUDE); +} + +TEST_F(ListRankScanTest, DeepList) +{ + using lcw = cudf::test::lists_column_wrapper; + lcw col{ + {{1, 2, 3}, {}, {4, 5}, {}, {0, 6, 0}}, + {{1, 2, 3}, {}, {4, 5}, {}, {0, 6, 0}}, + {{1, 2, 3}, {}, {4, 5}, {0, 6, 0}}, + {{7, 8}, {}}, + lcw{lcw{}, lcw{}, lcw{}}, + lcw{lcw{}}, + lcw{lcw{}}, + lcw{lcw{}}, + lcw{lcw{}, lcw{}, lcw{}}, + lcw{lcw{}, lcw{}, lcw{}}, + {lcw{10}}, + {lcw{10}}, + {{13, 14}, {15}}, + {{13, 14}, {16}}, + lcw{}, + lcw{lcw{}}, + }; + + { // Non-sliced + auto const expected_dense_vals = cudf::test::fixed_width_column_wrapper{ + 1, 1, 2, 3, 4, 5, 5, 5, 6, 6, 7, 7, 8, 9, 10, 11}; + this->test_ungrouped_rank_scan(col, + expected_dense_vals, + cudf::make_dense_rank_aggregation(), + cudf::null_policy::INCLUDE); + } + + { // sliced + auto sliced_col = cudf::slice(col, {3, 12})[0]; + auto const expected_dense_vals = + cudf::test::fixed_width_column_wrapper{1, 2, 3, 3, 3, 4, 4, 5, 5}; + this->test_ungrouped_rank_scan(sliced_col, + expected_dense_vals, + cudf::make_dense_rank_aggregation(), + cudf::null_policy::INCLUDE); + } +} + +TEST_F(ListRankScanTest, ListOfStruct) +{ + // Constructing a list of struct of two elements + // 0. [] == + // 1. [] != + // 2. Null == + // 3. Null != + // 4. [Null, Null] != + // 5. [Null] == + // 6. [Null] == + // 7. [Null] != + // 8. [{Null, Null}] != + // 9. [{1,'a'}, {2,'b'}] != + // 10. [{0,'a'}, {2,'b'}] != + // 11. [{0,'a'}, {2,'c'}] == + // 12. [{0,'a'}, {2,'c'}] != + // 13. [{0,Null}] == + // 14. [{0,Null}] != + // 15. [{Null, 0}] == + // 16. [{Null, 0}] + + auto col1 = cudf::test::fixed_width_column_wrapper{ + {-1, -1, 0, 2, 2, 2, 1, 2, 0, 2, 0, 2, 0, 2, 0, 0, 1, 2}, + {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0}}; + auto col2 = cudf::test::strings_column_wrapper{ + {"x", "x", "a", "a", "b", "b", "a", "b", "a", "b", "a", "c", "a", "c", "a", "c", "b", "b"}, + {1, 1, 1, 1, 1, 0, 1, 1, 1, 1, 1, 1, 1, 1, 0, 0, 1, 1}}; + auto struc = cudf::test::structs_column_wrapper{ + {col1, col2}, {0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}}; + + auto offsets = cudf::test::fixed_width_column_wrapper{ + 0, 0, 0, 0, 0, 2, 3, 4, 5, 6, 8, 10, 12, 14, 15, 16, 17, 18}; + + auto list_nullmask = std::vector{1, 1, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1}; + auto nullmask_buf = + cudf::test::detail::make_null_mask(list_nullmask.begin(), list_nullmask.end()); + auto list_column = cudf::column_view(cudf::data_type(cudf::type_id::LIST), + 17, + nullptr, + static_cast(nullmask_buf.data()), + cudf::UNKNOWN_NULL_COUNT, + 0, + {offsets, struc}); + + { // Non-sliced + auto expect = cudf::test::fixed_width_column_wrapper{ + 1, 1, 2, 2, 3, 4, 4, 4, 5, 6, 7, 8, 8, 9, 9, 10, 10}; + + this->test_ungrouped_rank_scan(list_column, + expect, + cudf::make_dense_rank_aggregation(), + cudf::null_policy::INCLUDE); + } + + { // Sliced + auto sliced_col = cudf::slice(list_column, {3, 15})[0]; + auto expect = + cudf::test::fixed_width_column_wrapper{1, 2, 3, 3, 3, 4, 5, 6, 7, 7, 8, 8}; + + this->test_ungrouped_rank_scan(sliced_col, + expect, + cudf::make_dense_rank_aggregation(), + cudf::null_policy::INCLUDE); + } +} + +TEST_F(ListRankScanTest, ListOfEmptyStruct) +{ + // [] + // [] + // Null + // Null + // [Null, Null] + // [Null, Null] + // [Null, Null] + // [Null] + // [Null] + // [{}] + // [{}] + // [{}, {}] + // [{}, {}] + + auto struct_validity = std::vector{0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1}; + auto struct_validity_buffer = + cudf::test::detail::make_null_mask(struct_validity.begin(), struct_validity.end()); + auto struct_col = + cudf::make_structs_column(14, {}, cudf::UNKNOWN_NULL_COUNT, std::move(struct_validity_buffer)); + + auto offsets = cudf::test::fixed_width_column_wrapper{ + 0, 0, 0, 0, 0, 2, 4, 6, 7, 8, 9, 10, 12, 14}; + auto list_nullmask = std::vector{1, 1, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1, 1}; + auto list_validity_buffer = + cudf::test::detail::make_null_mask(list_nullmask.begin(), list_nullmask.end()); + auto list_column = cudf::make_lists_column(13, + offsets.release(), + std::move(struct_col), + cudf::UNKNOWN_NULL_COUNT, + std::move(list_validity_buffer)); + + auto expect = + cudf::test::fixed_width_column_wrapper{1, 1, 2, 2, 3, 3, 3, 4, 4, 5, 5, 6, 6}; + + this->test_ungrouped_rank_scan(*list_column, + expect, + cudf::make_dense_rank_aggregation(), + cudf::null_policy::INCLUDE); +} + +TEST_F(ListRankScanTest, EmptyDeepList) +{ + // List>, where all lists are empty + // [] + // [] + // Null + // Null + + // Internal empty list + auto list1 = cudf::test::lists_column_wrapper{}; + + auto offsets = cudf::test::fixed_width_column_wrapper{0, 0, 0, 0, 0}; + auto list_nullmask = std::vector{1, 1, 0, 0}; + auto list_validity_buffer = + cudf::test::detail::make_null_mask(list_nullmask.begin(), list_nullmask.end()); + auto list_column = cudf::make_lists_column(4, + offsets.release(), + list1.release(), + cudf::UNKNOWN_NULL_COUNT, + std::move(list_validity_buffer)); + + auto expect = cudf::test::fixed_width_column_wrapper{1, 1, 2, 2}; + + this->test_ungrouped_rank_scan(*list_column, + expect, + cudf::make_dense_rank_aggregation(), + cudf::null_policy::INCLUDE); +}