From 933c974ddc0a29d002251e9558cdd41c9ff3cd2c Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Thu, 26 Aug 2021 19:28:51 +0530 Subject: [PATCH 001/333] First commit Many iterations already happened. I just realized late that I should commit --- cpp/benchmarks/CMakeLists.txt | 2 + .../compare/comparator_benchmark.cu | 92 ++++ cpp/include/cudf/sort2.cuh | 77 ++++ cpp/include/cudf/table/row_operator2.cuh | 420 ++++++++++++++++++ cpp/tests/CMakeLists.txt | 4 + cpp/tests/sort/sort2_test.cu | 66 +++ 6 files changed, 661 insertions(+) create mode 100644 cpp/benchmarks/compare/comparator_benchmark.cu create mode 100644 cpp/include/cudf/sort2.cuh create mode 100644 cpp/include/cudf/table/row_operator2.cuh create mode 100644 cpp/tests/sort/sort2_test.cu diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index 56f17dc7090..9ec90bd6afa 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -113,6 +113,8 @@ ConfigureBench(ITERATOR_BENCH iterator/iterator_benchmark.cu) # - search benchmark ------------------------------------------------------------------------------ ConfigureBench(SEARCH_BENCH search/search_benchmark.cpp) +ConfigureBench(COMPARE_BENCH compare/comparator_benchmark.cu) + ################################################################################################### # - sort benchmark -------------------------------------------------------------------------------- ConfigureBench(SORT_BENCH diff --git a/cpp/benchmarks/compare/comparator_benchmark.cu b/cpp/benchmarks/compare/comparator_benchmark.cu new file mode 100644 index 00000000000..f012456ac39 --- /dev/null +++ b/cpp/benchmarks/compare/comparator_benchmark.cu @@ -0,0 +1,92 @@ +/* + * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * 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 + +template +class Sort : public cudf::benchmark { +}; + +template +static void BM_sort(benchmark::State& state, bool nulls) +{ + using Type = int; + using column_wrapper = cudf::test::fixed_width_column_wrapper; + std::default_random_engine generator; + std::uniform_int_distribution distribution(0, 100); + + const cudf::size_type n_rows{(cudf::size_type)state.range(0)}; + const cudf::size_type n_cols{1}; + + // Create columns with values in the range [0,100) + std::vector columns; + columns.reserve(n_cols); + std::generate_n(std::back_inserter(columns), n_cols, [&, n_rows]() { + auto elements = cudf::detail::make_counting_transform_iterator( + 0, [&](auto row) { return distribution(generator); }); + if (!nulls) return column_wrapper(elements, elements + n_rows); + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 100 == 0 ? false : true; }); + return column_wrapper(elements, elements + n_rows, valids); + }); + + std::vector> cols; + std::transform(columns.begin(), columns.end(), std::back_inserter(cols), [](column_wrapper& col) { + return col.release(); + }); + + // Lets add some nulls + std::vector struct_validity; + std::uniform_int_distribution bool_distribution(0, 1000); + std::generate_n(std::back_inserter(struct_validity), cols[0]->size(), [&]() { + return bool_distribution(generator); + }); + cudf::test::structs_column_wrapper struct_col(std::move(cols), struct_validity); + + // // Create table view + auto input = cudf::table_view({struct_col}); + // auto input = cudf::table_view({cols[0]->view()}); + + for (auto _ : state) { + cuda_event_timer raii(state, true, rmm::cuda_stream_default); + + // auto result = cudf::sorted_order(input); + auto result = cudf::detail::sorted_order2(input); + } +} + +#define SORT_BENCHMARK_DEFINE(name, stable, nulls) \ + BENCHMARK_TEMPLATE_DEFINE_F(Sort, name, stable) \ + (::benchmark::State & st) { BM_sort(st, nulls); } \ + BENCHMARK_REGISTER_F(Sort, name) \ + ->RangeMultiplier(8) \ + ->Ranges({{1 << 10, 1 << 26}, {1, 1}}) \ + ->UseManualTime() \ + ->Unit(benchmark::kMillisecond); + +SORT_BENCHMARK_DEFINE(unstable, false, true) diff --git a/cpp/include/cudf/sort2.cuh b/cpp/include/cudf/sort2.cuh new file mode 100644 index 00000000000..96006550716 --- /dev/null +++ b/cpp/include/cudf/sort2.cuh @@ -0,0 +1,77 @@ +/* + * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * 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 + +#include +#include + +namespace cudf { +namespace detail { + +/** + * @copydoc + * sorted_order(table_view&,std::vector,std::vector,rmm::mr::device_memory_resource*) + * + * @param stream CUDA stream used for device memory operations and kernel launches + */ +template +std::unique_ptr sorted_order2( + table_view input, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) +{ + if (input.num_rows() == 0 or input.num_columns() == 0) { + return cudf::make_numeric_column(data_type(type_to_id()), 0); + } + + std::unique_ptr sorted_indices = cudf::make_numeric_column( + data_type(type_to_id()), input.num_rows(), mask_state::UNALLOCATED, stream, mr); + mutable_column_view mutable_indices_view = sorted_indices->mutable_view(); + thrust::sequence(rmm::exec_policy(stream), + mutable_indices_view.begin(), + mutable_indices_view.end(), + 0); + + auto device_table = table_device_view::create(input, stream); + auto const comparator = row_lexicographic_comparator2(*device_table, *device_table); + + thrust::sort(rmm::exec_policy(stream), + mutable_indices_view.begin(), + mutable_indices_view.end(), + comparator); + // protection for temporary d_column_order and d_null_precedence + stream.synchronize(); + + return sorted_indices; +} + +} // namespace detail +} // namespace cudf diff --git a/cpp/include/cudf/table/row_operator2.cuh b/cpp/include/cudf/table/row_operator2.cuh new file mode 100644 index 00000000000..9edb60cc112 --- /dev/null +++ b/cpp/include/cudf/table/row_operator2.cuh @@ -0,0 +1,420 @@ +/* + * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * 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 { + +/** + * @brief Result type of the `element_relational_comparator2` function object. + * + * Indicates how two elements `a` and `b` compare with one and another. + * + * Equivalence is defined as `not (a +__device__ weak_ordering2 compare_elements2(Element lhs, Element rhs) +{ + if (lhs < rhs) { + return weak_ordering2::LESS; + } else if (rhs < lhs) { + return weak_ordering2::GREATER; + } + return weak_ordering2::EQUIVALENT; +} +} // namespace detail + +/* + * @brief A specialization for floating-point `Element` type relational comparison + * to derive the order of the elements with respect to `lhs`. Specialization is to + * handle `nan` in the order shown below. + * `[-Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN, null] (for null_order::AFTER)` + * `[null, -Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN] (for null_order::BEFORE)` + * + * @param[in] lhs first element + * @param[in] rhs second element + * @return weak_ordering2 Indicates the relationship between the elements in + * the `lhs` and `rhs` columns. + */ +template ::value>* = nullptr> +__device__ weak_ordering2 relational_compare2(Element lhs, Element rhs) +{ + if (isnan(lhs) and isnan(rhs)) { + return weak_ordering2::EQUIVALENT; + } else if (isnan(rhs)) { + return weak_ordering2::LESS; + } else if (isnan(lhs)) { + return weak_ordering2::GREATER; + } + + return detail::compare_elements2(lhs, rhs); +} + +/** + * @brief Compare the nulls according to null order. + * + * @param lhs_is_null boolean representing if lhs is null + * @param rhs_is_null boolean representing if lhs is null + * @param null_precedence null order + * @return Indicates the relationship between null in lhs and rhs columns. + */ +inline __device__ auto null_compare2(bool lhs_is_null, bool rhs_is_null, null_order null_precedence) +{ + if (lhs_is_null and rhs_is_null) { // null ::value>* = nullptr> +__device__ weak_ordering2 relational_compare2(Element lhs, Element rhs) +{ + return detail::compare_elements2(lhs, rhs); +} + +/** + * @brief A specialization for floating-point `Element` type to check if + * `lhs` is equivalent to `rhs`. `nan == nan`. + * + * @param[in] lhs first element + * @param[in] rhs second element + * @return bool `true` if `lhs` == `rhs` else `false`. + */ +template ::value>* = nullptr> +__device__ bool equality_compare2(Element lhs, Element rhs) +{ + if (isnan(lhs) and isnan(rhs)) { return true; } + return lhs == rhs; +} + +/** + * @brief A specialization for non-floating-point `Element` type to check if + * `lhs` is equivalent to `rhs`. + * + * @param[in] lhs first element + * @param[in] rhs second element + * @return bool `true` if `lhs` == `rhs` else `false`. + */ +template ::value>* = nullptr> +__device__ bool equality_compare2(Element const lhs, Element const rhs) +{ + return lhs == rhs; +} + +/** + * @brief Performs an equality comparison between two elements in two columns. + * + * @tparam has_nulls Indicates the potential for null values in either column. + */ +template +class element_equality_comparator2 { + public: + /** + * @brief Construct type-dispatched function object for comparing equality + * between two elements. + * + * @note `lhs` and `rhs` may be the same. + * + * @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 + */ + __host__ __device__ element_equality_comparator2(column_device_view lhs, + column_device_view rhs, + bool nulls_are_equal = true) + : lhs{lhs}, rhs{rhs}, 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 + * + */ + template ()>* = nullptr> + __device__ bool operator()(size_type lhs_element_index, + size_type rhs_element_index) const noexcept + { + if (has_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; + } else if (lhs_is_null != rhs_is_null) { + return false; + } + } + + return equality_compare2(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); + } + + template ()>* = nullptr> + __device__ bool operator()(size_type lhs_element_index, size_type rhs_element_index) + { + cudf_assert(false && "Attempted to compare elements of uncomparable types."); + return false; + } + + private: + column_device_view lhs; + column_device_view rhs; + bool nulls_are_equal; +}; + +template +class row_equality_comparator2 { + public: + row_equality_comparator2(table_device_view lhs, + table_device_view rhs, + bool nulls_are_equal = true) + : lhs{lhs}, rhs{rhs}, nulls_are_equal{nulls_are_equal} + { + CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(), "Mismatched number of columns."); + } + + __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept + { + auto equal_elements = [=](column_device_view l, column_device_view r) { + return cudf::type_dispatcher(l.type(), + element_equality_comparator2{l, r, nulls_are_equal}, + lhs_row_index, + rhs_row_index); + }; + + return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements); + } + + private: + table_device_view lhs; + table_device_view rhs; + bool nulls_are_equal; +}; + +/** + * @brief Performs a relational comparison between two elements in two columns. + * + * @tparam has_nulls Indicates the potential for null values in either column. + */ +template +class element_relational_comparator2 { + public: + /** + * @brief Construct type-dispatched function object for performing a + * relational comparison between two elements. + * + * @note `lhs` and `rhs` may be the same. + * + * @param lhs The column containing the first element + * @param rhs The column containing the second element (may be the same as lhs) + * @param null_precedence Indicates how null values are ordered with other + * values + */ + __host__ __device__ element_relational_comparator2(column_device_view lhs, + column_device_view rhs, + null_order null_precedence) + : lhs{lhs}, rhs{rhs}, null_precedence{null_precedence} + { + } + + /** + * @brief Performs a relational comparison between the specified elements + * + * @param lhs_element_index The index of the first element + * @param rhs_element_index The index of the second element + * @param null_precedence Indicates how null values are ordered with other + * values + * @return weak_ordering2 Indicates the relationship between the elements in + * the `lhs` and `rhs` columns. + */ + template ()>* = nullptr> + __device__ weak_ordering2 operator()(size_type lhs_element_index, + size_type rhs_element_index) const noexcept + { + if (has_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 or rhs_is_null) { // atleast one is null + return null_compare2(lhs_is_null, rhs_is_null, null_precedence); + } + } + + return relational_compare2(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); + } + + template ()>* = nullptr> + __device__ weak_ordering2 operator()(size_type lhs_element_index, size_type rhs_element_index) + { + cudf_assert(false && "Attempted to compare elements of uncomparable types."); + return weak_ordering2::LESS; + } + + private: + column_device_view lhs; + column_device_view rhs; + null_order null_precedence; +}; + +/** + * @brief Computes whether one row is lexicographically *less* than another row. + * + * Lexicographic ordering is determined by: + * - Two rows are compared element by element. + * - The first mismatching element defines which row is lexicographically less + * or greater than the other. + * + * Lexicographic ordering is exactly equivalent to doing an alphabetical sort of + * two words, for example, `aac` would be *less* than (or precede) `abb`. The + * second letter in both words is the first non-equal letter, and `a < b`, thus + * `aac < abb`. + * + * @tparam has_nulls Indicates the potential for null values in either row. + */ +template +class row_lexicographic_comparator2 { + public: + /** + * @brief Construct a function object for performing a lexicographic + * comparison between the rows of two tables. + * + * @throws cudf::logic_error if `lhs.num_columns() != rhs.num_columns()` + * @throws cudf::logic_error if column types of `lhs` and `rhs` are not comparable. + * + * @param lhs The first table + * @param rhs The second table (may be the same table as `lhs`) + * @param column_order Optional, device array the same length as a row that + * indicates the desired ascending/descending order of each column in a row. + * If `nullptr`, it is assumed all columns are sorted in ascending order. + * @param null_precedence Optional, device array the same length as a row + * and indicates how null values compare to all other for every column. If + * it is nullptr, then null precedence would be `null_order::BEFORE` for all + * columns. + */ + row_lexicographic_comparator2(table_device_view lhs, + table_device_view rhs, + order const* column_order = nullptr, + null_order const* null_precedence = nullptr) + : _lhs{lhs}, _rhs{rhs}, _column_order{column_order}, _null_precedence{null_precedence} + { + CUDF_EXPECTS(_lhs.num_columns() == _rhs.num_columns(), "Mismatched number of columns."); + // CUDF_EXPECTS(detail::is_relationally_comparable(_lhs, _rhs), + // "Attempted to compare elements of uncomparable types."); + } + + /** + * @brief Checks whether the row at `lhs_index` in the `lhs` table compares + * lexicographically less than 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 compares less than row in the + * `rhs` table + */ + __device__ bool operator()(size_type lhs_index, size_type rhs_index) const noexcept + { + for (size_type i = 0; i < _lhs.num_columns(); ++i) { + bool ascending = (_column_order == nullptr) or (_column_order[i] == order::ASCENDING); + + weak_ordering2 state{weak_ordering2::EQUIVALENT}; + null_order null_precedence = + _null_precedence == nullptr ? null_order::BEFORE : _null_precedence[i]; + + column_device_view lcol = _lhs.column(i); + column_device_view rcol = _rhs.column(i); + while (lcol.type().id() == type_id::STRUCT) { + bool const lhs_is_null{lcol.is_null(lhs_index)}; + bool const rhs_is_null{rcol.is_null(rhs_index)}; + + if (lhs_is_null or rhs_is_null) { // atleast one is null + state = null_compare2(lhs_is_null, rhs_is_null, null_precedence); + if (state != weak_ordering2::EQUIVALENT) break; + } + + lcol = lcol.children()[0]; + rcol = rcol.children()[0]; + } + + if (state == weak_ordering2::EQUIVALENT) { + auto comparator = element_relational_comparator2{lcol, rcol, null_precedence}; + state = cudf::type_dispatcher(lcol.type(), comparator, lhs_index, rhs_index); + } + + if (state == weak_ordering2::EQUIVALENT) { continue; } + + return state == (ascending ? weak_ordering2::LESS : weak_ordering2::GREATER); + } + return false; + } + + private: + table_device_view _lhs; + table_device_view _rhs; + null_order const* _null_precedence{}; + order const* _column_order{}; +}; // class row_lexicographic_comparator2 + +} // namespace cudf diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index c82826b8c60..de690a4af79 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -206,6 +206,10 @@ ConfigureTest(SORT_TEST sort/sort_test.cpp sort/rank_test.cpp) +ConfigureTest(SORT2_TEST + sort/sort2_test.cu +) + ################################################################################################### # - copying tests --------------------------------------------------------------------------------- ConfigureTest(COPYING_TEST diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu new file mode 100644 index 00000000000..3ab67b0d0df --- /dev/null +++ b/cpp/tests/sort/sort2_test.cu @@ -0,0 +1,66 @@ +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +#include + +struct NewRowOp : public cudf::test::BaseFixture { +}; + +#include + +TEST_F(NewRowOp, BasicTest) +{ + using Type = int; + using column_wrapper = cudf::test::fixed_width_column_wrapper; + std::default_random_engine generator; + std::uniform_int_distribution distribution(0, 100); + + const cudf::size_type n_rows{1 << 10}; + const cudf::size_type n_cols{1}; + + // Create columns with values in the range [0,100) + std::vector columns; + columns.reserve(n_cols); + std::generate_n(std::back_inserter(columns), n_cols, [&, n_rows]() { + auto elements = cudf::detail::make_counting_transform_iterator( + 0, [&](auto row) { return distribution(generator); }); + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 3 == 0 ? false : true; }); + return column_wrapper(elements, elements + n_rows, valids); + }); + + std::vector> cols; + std::transform(columns.begin(), columns.end(), std::back_inserter(cols), [](column_wrapper& col) { + return col.release(); + }); + + // Lets add some nulls + std::vector struct_validity; + std::uniform_int_distribution bool_distribution(0, 10); + std::generate_n(std::back_inserter(struct_validity), cols[0]->size(), [&]() { + return bool_distribution(generator); + }); + cudf::test::structs_column_wrapper struct_col(std::move(cols), struct_validity); + + // cudf::test::print(struct_col); + + // // Create table view + auto input = cudf::table_view({struct_col}); + + auto result1 = cudf::sorted_order(input); + // cudf::test::print(result1->view()); + auto result2 = cudf::detail::sorted_order2(input); + // cudf::test::print(result2->view()); + cudf::test::expect_columns_equal(result1->view(), result2->view()); +} From a1636e52a925460d27e2b4bb7ec2bac8c3710f75 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Sat, 28 Aug 2021 05:17:22 +0530 Subject: [PATCH 002/333] testing and profiling deep single hierarchy struct --- .../compare/comparator_benchmark.cu | 25 +++++++++++------- cpp/tests/sort/sort2_test.cu | 26 ++++++++++++------- 2 files changed, 31 insertions(+), 20 deletions(-) diff --git a/cpp/benchmarks/compare/comparator_benchmark.cu b/cpp/benchmarks/compare/comparator_benchmark.cu index f012456ac39..e4f690102b0 100644 --- a/cpp/benchmarks/compare/comparator_benchmark.cu +++ b/cpp/benchmarks/compare/comparator_benchmark.cu @@ -41,6 +41,7 @@ static void BM_sort(benchmark::State& state, bool nulls) std::uniform_int_distribution distribution(0, 100); const cudf::size_type n_rows{(cudf::size_type)state.range(0)}; + const cudf::size_type depth{(cudf::size_type)state.range(1)}; const cudf::size_type n_cols{1}; // Create columns with values in the range [0,100) @@ -51,7 +52,7 @@ static void BM_sort(benchmark::State& state, bool nulls) 0, [&](auto row) { return distribution(generator); }); if (!nulls) return column_wrapper(elements, elements + n_rows); auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 100 == 0 ? false : true; }); + 0, [](auto i) { return i % 10 == 0 ? false : true; }); return column_wrapper(elements, elements + n_rows, valids); }); @@ -60,16 +61,20 @@ static void BM_sort(benchmark::State& state, bool nulls) return col.release(); }); - // Lets add some nulls - std::vector struct_validity; - std::uniform_int_distribution bool_distribution(0, 1000); - std::generate_n(std::back_inserter(struct_validity), cols[0]->size(), [&]() { - return bool_distribution(generator); - }); - cudf::test::structs_column_wrapper struct_col(std::move(cols), struct_validity); + std::vector> child_cols = std::move(cols); + // Lets add some layers + for (int i = 0; i < depth; i++) { + std::vector struct_validity; + std::uniform_int_distribution bool_distribution(0, 100 * (i + 1)); + std::generate_n( + std::back_inserter(struct_validity), n_rows, [&]() { return bool_distribution(generator); }); + cudf::test::structs_column_wrapper struct_col(std::move(child_cols), struct_validity); + child_cols = std::vector>{}; + child_cols.push_back(struct_col.release()); + } // // Create table view - auto input = cudf::table_view({struct_col}); + auto input = cudf::table(std::move(child_cols)); // auto input = cudf::table_view({cols[0]->view()}); for (auto _ : state) { @@ -85,7 +90,7 @@ static void BM_sort(benchmark::State& state, bool nulls) (::benchmark::State & st) { BM_sort(st, nulls); } \ BENCHMARK_REGISTER_F(Sort, name) \ ->RangeMultiplier(8) \ - ->Ranges({{1 << 10, 1 << 26}, {1, 1}}) \ + ->Ranges({{1 << 10, 1 << 26}, {1, 8}}) \ ->UseManualTime() \ ->Unit(benchmark::kMillisecond); diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu index 3ab67b0d0df..44abf545a71 100644 --- a/cpp/tests/sort/sort2_test.cu +++ b/cpp/tests/sort/sort2_test.cu @@ -26,8 +26,9 @@ TEST_F(NewRowOp, BasicTest) std::default_random_engine generator; std::uniform_int_distribution distribution(0, 100); - const cudf::size_type n_rows{1 << 10}; + const cudf::size_type n_rows{1 << 6}; const cudf::size_type n_cols{1}; + const cudf::size_type depth{8}; // Create columns with values in the range [0,100) std::vector columns; @@ -45,18 +46,23 @@ TEST_F(NewRowOp, BasicTest) return col.release(); }); - // Lets add some nulls - std::vector struct_validity; - std::uniform_int_distribution bool_distribution(0, 10); - std::generate_n(std::back_inserter(struct_validity), cols[0]->size(), [&]() { - return bool_distribution(generator); - }); - cudf::test::structs_column_wrapper struct_col(std::move(cols), struct_validity); + std::vector> child_cols = std::move(cols); + // Lets add some layers + for (int i = 0; i < depth; i++) { + std::vector struct_validity; + std::uniform_int_distribution bool_distribution(0, 10 * (i + 1)); + std::generate_n( + std::back_inserter(struct_validity), n_rows, [&]() { return bool_distribution(generator); }); + cudf::test::structs_column_wrapper struct_col(std::move(child_cols), struct_validity); + child_cols = std::vector>{}; + child_cols.push_back(struct_col.release()); + } - // cudf::test::print(struct_col); + cudf::test::print(child_cols[0]->view()); // // Create table view - auto input = cudf::table_view({struct_col}); + // auto input = cudf::table_view({struct_col}); + auto input = cudf::table(std::move(child_cols)); auto result1 = cudf::sorted_order(input); // cudf::test::print(result1->view()); From 3d21daf53efe83e82f088c61ac4ff19d68d8a431 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Fri, 14 Jan 2022 20:18:32 +0530 Subject: [PATCH 003/333] Make the sandboxed test compile again --- cpp/include/cudf/sort2.cuh | 2 +- cpp/tests/CMakeLists.txt | 4 +--- cpp/tests/sort/sort2_test.cu | 2 ++ 3 files changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/include/cudf/sort2.cuh b/cpp/include/cudf/sort2.cuh index 96006550716..6732e5a5db1 100644 --- a/cpp/include/cudf/sort2.cuh +++ b/cpp/include/cudf/sort2.cuh @@ -24,7 +24,7 @@ #include #include -#include +#include #include #include diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index f9f157dbcb4..23a158af4c2 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -205,9 +205,7 @@ endif() # * sort tests ------------------------------------------------------------------------------------ ConfigureTest(SORT_TEST sort/segmented_sort_tests.cpp sort/sort_test.cpp sort/rank_test.cpp) -ConfigureTest(SORT2_TEST - sort/sort2_test.cu -) +ConfigureTest(SORT2_TEST sort/sort2_test.cu) # ################################################################################################## # * copying tests --------------------------------------------------------------------------------- diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu index 44abf545a71..649951d1d5b 100644 --- a/cpp/tests/sort/sort2_test.cu +++ b/cpp/tests/sort/sort2_test.cu @@ -70,3 +70,5 @@ TEST_F(NewRowOp, BasicTest) // cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); } + +CUDF_TEST_PROGRAM_MAIN() From 9f32e6b1eafe5db74ae6157cbec057eb1259b58d Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Sat, 15 Jan 2022 05:37:18 +0530 Subject: [PATCH 004/333] Update my row_comparator with nullate --- cpp/include/cudf/sort2.cuh | 7 +- cpp/include/cudf/table/row_operator3.cuh | 441 +++++++++++++++++++++++ 2 files changed, 446 insertions(+), 2 deletions(-) create mode 100644 cpp/include/cudf/table/row_operator3.cuh diff --git a/cpp/include/cudf/sort2.cuh b/cpp/include/cudf/sort2.cuh index 6732e5a5db1..1d1a5300787 100644 --- a/cpp/include/cudf/sort2.cuh +++ b/cpp/include/cudf/sort2.cuh @@ -20,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -60,8 +61,10 @@ std::unique_ptr sorted_order2( mutable_indices_view.end(), 0); - auto device_table = table_device_view::create(input, stream); - auto const comparator = row_lexicographic_comparator2(*device_table, *device_table); + auto device_table = table_device_view::create(input, stream); + // auto const comparator = row_lexicographic_comparator2(*device_table, *device_table); + auto const comparator = + row_lexicographic_comparator3(nullate::DYNAMIC{true}, *device_table, *device_table); thrust::sort(rmm::exec_policy(stream), mutable_indices_view.begin(), diff --git a/cpp/include/cudf/table/row_operator3.cuh b/cpp/include/cudf/table/row_operator3.cuh new file mode 100644 index 00000000000..1e3c2f7d75b --- /dev/null +++ b/cpp/include/cudf/table/row_operator3.cuh @@ -0,0 +1,441 @@ +/* + * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * 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 { + +/** + * @brief Result type of the `element_relational_comparator2` function object. + * + * Indicates how two elements `a` and `b` compare with one and another. + * + * Equivalence is defined as `not (a +__device__ weak_ordering3 compare_elements3(Element lhs, Element rhs) +{ + if (lhs < rhs) { + return weak_ordering3::LESS; + } else if (rhs < lhs) { + return weak_ordering3::GREATER; + } + return weak_ordering3::EQUIVALENT; +} +} // namespace detail + +/** + * @brief A specialization for floating-point `Element` type relational comparison + * to derive the order of the elements with respect to `lhs`. + * + * This Specialization handles `nan` in the following order: + * `[-Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN, null] (for null_order::AFTER)` + * `[null, -Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN] (for null_order::BEFORE)` + * + * @param lhs first element + * @param rhs second element + * @return Indicates the relationship between the elements in + * the `lhs` and `rhs` columns. + */ +template ::value>* = nullptr> +__device__ weak_ordering3 relational_compare3(Element lhs, Element rhs) +{ + if (isnan(lhs) and isnan(rhs)) { + return weak_ordering3::EQUIVALENT; + } else if (isnan(rhs)) { + return weak_ordering3::LESS; + } else if (isnan(lhs)) { + return weak_ordering3::GREATER; + } + + return detail::compare_elements3(lhs, rhs); +} + +/** + * @brief Compare the nulls according to null order. + * + * @param lhs_is_null boolean representing if lhs is null + * @param rhs_is_null boolean representing if lhs is null + * @param null_precedence null order + * @return Indicates the relationship between null in lhs and rhs columns. + */ +inline __device__ auto null_compare3(bool lhs_is_null, bool rhs_is_null, null_order null_precedence) +{ + if (lhs_is_null and rhs_is_null) { // null ::value>* = nullptr> +__device__ weak_ordering3 relational_compare3(Element lhs, Element rhs) +{ + return detail::compare_elements3(lhs, rhs); +} + +/** + * @brief A specialization for floating-point `Element` type to check if + * `lhs` is equivalent to `rhs`. `nan == nan`. + * + * @param lhs first element + * @param rhs second element + * @return `true` if `lhs` == `rhs` else `false`. + */ +template ::value>* = nullptr> +__device__ bool equality_compare3(Element lhs, Element rhs) +{ + if (isnan(lhs) and isnan(rhs)) { return true; } + return lhs == rhs; +} + +/** + * @brief A specialization for non-floating-point `Element` type to check if + * `lhs` is equivalent to `rhs`. + * + * @param lhs first element + * @param rhs second element + * @return `true` if `lhs` == `rhs` else `false`. + */ +template ::value>* = nullptr> +__device__ bool equality_compare3(Element const lhs, Element const rhs) +{ + return lhs == rhs; +} + +/** + * @brief Performs an equality comparison between two elements in two columns. + * + * @tparam Nullate A cudf::nullate type describing how to check for nulls. + */ +template +class element_equality_comparator3 { + 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 + */ + __host__ __device__ + element_equality_comparator3(Nullate has_nulls, + column_device_view lhs, + column_device_view rhs, + null_equality nulls_are_equal = null_equality::EQUAL) + : 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 both lhs and rhs element are both nulls and `nulls_are_equal` is true, or equal + */ + template ()>* = nullptr> + __device__ bool operator()(size_type lhs_element_index, + size_type 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_compare3(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); + } + + template ()>* = nullptr> + __device__ bool operator()(size_type lhs_element_index, size_type rhs_element_index) + { + cudf_assert(false && "Attempted to compare elements of uncomparable types."); + return false; + } + + private: + column_device_view lhs; + column_device_view rhs; + Nullate nulls; + null_equality nulls_are_equal; +}; + +template +class row_equality_comparator3 { + public: + row_equality_comparator3(Nullate has_nulls, + table_device_view lhs, + table_device_view rhs, + null_equality nulls_are_equal = true) + : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal} + { + CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(), "Mismatched number of columns."); + } + + __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept + { + auto equal_elements = [=](column_device_view l, column_device_view r) { + return cudf::type_dispatcher(l.type(), + element_equality_comparator3{nulls, l, r, nulls_are_equal}, + lhs_row_index, + rhs_row_index); + }; + + return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements); + } + + private: + table_device_view lhs; + table_device_view rhs; + Nullate nulls; + null_equality nulls_are_equal; +}; + +/** + * @brief Performs a relational comparison between two elements in two columns. + * + * @tparam Nullate A cudf::nullate type describing how to check for nulls. + */ +template +class element_relational_comparator3 { + public: + /** + * @brief Construct type-dispatched function object for performing a + * relational comparison between two elements. + * + * @note `lhs` and `rhs` may be the same. + * + * @param lhs The column containing the first element + * @param rhs The column containing the second element (may be the same as lhs) + * @param has_nulls Indicates if either input column contains nulls. + * @param null_precedence Indicates how null values are ordered with other values + */ + __host__ __device__ element_relational_comparator3(Nullate has_nulls, + column_device_view lhs, + column_device_view rhs, + null_order null_precedence) + : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, null_precedence{null_precedence} + { + } + + __host__ __device__ element_relational_comparator3(Nullate has_nulls, + column_device_view lhs, + column_device_view rhs) + : lhs{lhs}, rhs{rhs}, nulls{has_nulls} + { + } + + /** + * @brief Performs a relational comparison between the specified elements + * + * @param lhs_element_index The index of the first element + * @param rhs_element_index The index of the second element + * @return Indicates the relationship between the elements in + * the `lhs` and `rhs` columns. + */ + template ()>* = nullptr> + __device__ weak_ordering3 operator()(size_type lhs_element_index, + size_type 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 or rhs_is_null) { // at least one is null + return null_compare3(lhs_is_null, rhs_is_null, null_precedence); + } + } + + return relational_compare3(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); + } + + template ()>* = nullptr> + __device__ weak_ordering3 operator()(size_type lhs_element_index, size_type rhs_element_index) + { + cudf_assert(false && "Attempted to compare elements of uncomparable types."); + return weak_ordering3::LESS; + } + + private: + column_device_view lhs; + column_device_view rhs; + Nullate nulls; + null_order null_precedence{}; +}; + +/** + * @brief Computes whether one row is lexicographically *less* than another row. + * + * Lexicographic ordering is determined by: + * - Two rows are compared element by element. + * - The first mismatching element defines which row is lexicographically less + * or greater than the other. + * + * Lexicographic ordering is exactly equivalent to doing an alphabetical sort of + * two words, for example, `aac` would be *less* than (or precede) `abb`. The + * second letter in both words is the first non-equal letter, and `a < b`, thus + * `aac < abb`. + * + * @tparam Nullate A cudf::nullate type describing how to check for nulls. + */ +template +class row_lexicographic_comparator3 { + public: + /** + * @brief Construct a function object for performing a lexicographic + * comparison between the rows of two tables. + * + * @throws cudf::logic_error if `lhs.num_columns() != rhs.num_columns()` + * @throws cudf::logic_error if column types of `lhs` and `rhs` are not comparable. + * + * @param lhs The first table + * @param rhs The second table (may be the same table as `lhs`) + * @param has_nulls Indicates if either input table contains columns with nulls. + * @param column_order Optional, device array the same length as a row that + * indicates the desired ascending/descending order of each column in a row. + * If `nullptr`, it is assumed all columns are sorted in ascending order. + * @param null_precedence Optional, device array the same length as a row + * and indicates how null values compare to all other for every column. If + * it is nullptr, then null precedence would be `null_order::BEFORE` for all + * columns. + */ + row_lexicographic_comparator3(Nullate has_nulls, + table_device_view lhs, + table_device_view rhs, + order const* column_order = nullptr, + null_order const* null_precedence = nullptr) + : _lhs{lhs}, + _rhs{rhs}, + _nulls{has_nulls}, + _column_order{column_order}, + _null_precedence{null_precedence} + { + CUDF_EXPECTS(_lhs.num_columns() == _rhs.num_columns(), "Mismatched number of columns."); + // CUDF_EXPECTS(detail::is_relationally_comparable(_lhs, _rhs), + // "Attempted to compare elements of uncomparable types."); + } + + /** + * @brief Checks whether the row at `lhs_index` in the `lhs` table compares + * lexicographically less than 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 compares less than row in the + * `rhs` table + */ + __device__ bool operator()(size_type lhs_index, size_type rhs_index) const noexcept + { + for (size_type i = 0; i < _lhs.num_columns(); ++i) { + bool ascending = (_column_order == nullptr) or (_column_order[i] == order::ASCENDING); + + weak_ordering3 state{weak_ordering3::EQUIVALENT}; + null_order null_precedence = + _null_precedence == nullptr ? null_order::BEFORE : _null_precedence[i]; + + column_device_view lcol = _lhs.column(i); + column_device_view rcol = _rhs.column(i); + while (lcol.type().id() == type_id::STRUCT) { + bool const lhs_is_null{lcol.is_null(lhs_index)}; + bool const rhs_is_null{rcol.is_null(rhs_index)}; + + if (lhs_is_null or rhs_is_null) { // atleast one is null + state = null_compare3(lhs_is_null, rhs_is_null, null_precedence); + if (state != weak_ordering3::EQUIVALENT) break; + } + + lcol = lcol.children()[0]; + rcol = rcol.children()[0]; + } + + if (state == weak_ordering3::EQUIVALENT) { + auto comparator = element_relational_comparator3{_nulls, lcol, rcol, null_precedence}; + state = cudf::type_dispatcher(lcol.type(), comparator, lhs_index, rhs_index); + } + + if (state == weak_ordering3::EQUIVALENT) { continue; } + + return state == (ascending ? weak_ordering3::LESS : weak_ordering3::GREATER); + } + return false; + } + + private: + table_device_view _lhs; + table_device_view _rhs; + Nullate _nulls{}; + null_order const* _null_precedence{}; + order const* _column_order{}; +}; // class row_lexicographic_comparator3 + +} // namespace cudf From 022e2a49d0cc4b6cc65da972595e5b3c7f2d7e9a Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Tue, 25 Jan 2022 02:35:22 +0530 Subject: [PATCH 005/333] Basic verticalization utility and experimental namespace --- cpp/include/cudf/detail/structs/utilities.hpp | 5 + cpp/include/cudf/sort2.cuh | 8 +- cpp/include/cudf/table/row_operator3.cuh | 124 +++++++++--------- cpp/src/structs/utilities.cpp | 52 ++++++++ cpp/tests/sort/sort2_test.cu | 105 +++++++++++++++ 5 files changed, 230 insertions(+), 64 deletions(-) diff --git a/cpp/include/cudf/detail/structs/utilities.hpp b/cpp/include/cudf/detail/structs/utilities.hpp index 751b7c00e8a..23601548cbf 100644 --- a/cpp/include/cudf/detail/structs/utilities.hpp +++ b/cpp/include/cudf/detail/structs/utilities.hpp @@ -150,6 +150,11 @@ flattened_table flatten_nested_columns( std::vector const& null_precedence, column_nullability nullability = column_nullability::MATCH_INCOMING); +namespace experimental { +std::tuple> verticalize_nested_columns( + table_view input); +} + /** * @brief Unflatten columns flattened as by `flatten_nested_columns()`, * based on the provided `blueprint`. diff --git a/cpp/include/cudf/sort2.cuh b/cpp/include/cudf/sort2.cuh index 1d1a5300787..125d696d055 100644 --- a/cpp/include/cudf/sort2.cuh +++ b/cpp/include/cudf/sort2.cuh @@ -61,10 +61,12 @@ std::unique_ptr sorted_order2( mutable_indices_view.end(), 0); - auto device_table = table_device_view::create(input, stream); + auto [vertical_table, nullmasks] = + cudf::structs::detail::experimental::verticalize_nested_columns(input); + auto device_table = table_device_view::create(vertical_table, stream); // auto const comparator = row_lexicographic_comparator2(*device_table, *device_table); - auto const comparator = - row_lexicographic_comparator3(nullate::DYNAMIC{true}, *device_table, *device_table); + auto const comparator = experimental::row_lexicographic_comparator( + nullate::DYNAMIC{true}, *device_table, *device_table); thrust::sort(rmm::exec_policy(stream), mutable_indices_view.begin(), diff --git a/cpp/include/cudf/table/row_operator3.cuh b/cpp/include/cudf/table/row_operator3.cuh index 1e3c2f7d75b..71bbb59339e 100644 --- a/cpp/include/cudf/table/row_operator3.cuh +++ b/cpp/include/cudf/table/row_operator3.cuh @@ -31,6 +31,7 @@ #include namespace cudf { +namespace experimental { /** * @brief Result type of the `element_relational_comparator2` function object. @@ -40,7 +41,7 @@ namespace cudf { * Equivalence is defined as `not (a -__device__ weak_ordering3 compare_elements3(Element lhs, Element rhs) +__device__ weak_ordering compare_elements(Element lhs, Element rhs) { if (lhs < rhs) { - return weak_ordering3::LESS; + return weak_ordering::LESS; } else if (rhs < lhs) { - return weak_ordering3::GREATER; + return weak_ordering::GREATER; } - return weak_ordering3::EQUIVALENT; + return weak_ordering::EQUIVALENT; } } // namespace detail @@ -81,17 +82,17 @@ __device__ weak_ordering3 compare_elements3(Element lhs, Element rhs) * the `lhs` and `rhs` columns. */ template ::value>* = nullptr> -__device__ weak_ordering3 relational_compare3(Element lhs, Element rhs) +__device__ weak_ordering relational_compare(Element lhs, Element rhs) { if (isnan(lhs) and isnan(rhs)) { - return weak_ordering3::EQUIVALENT; + return weak_ordering::EQUIVALENT; } else if (isnan(rhs)) { - return weak_ordering3::LESS; + return weak_ordering::LESS; } else if (isnan(lhs)) { - return weak_ordering3::GREATER; + return weak_ordering::GREATER; } - return detail::compare_elements3(lhs, rhs); + return detail::compare_elements(lhs, rhs); } /** @@ -102,16 +103,16 @@ __device__ weak_ordering3 relational_compare3(Element lhs, Element rhs) * @param null_precedence null order * @return Indicates the relationship between null in lhs and rhs columns. */ -inline __device__ auto null_compare3(bool lhs_is_null, bool rhs_is_null, null_order null_precedence) +inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_order null_precedence) { if (lhs_is_null and rhs_is_null) { // null ::value>* = nullptr> -__device__ weak_ordering3 relational_compare3(Element lhs, Element rhs) +__device__ weak_ordering relational_compare(Element lhs, Element rhs) { - return detail::compare_elements3(lhs, rhs); + return detail::compare_elements(lhs, rhs); } /** @@ -138,7 +139,7 @@ __device__ weak_ordering3 relational_compare3(Element lhs, Element rhs) * @return `true` if `lhs` == `rhs` else `false`. */ template ::value>* = nullptr> -__device__ bool equality_compare3(Element lhs, Element rhs) +__device__ bool equality_compare(Element lhs, Element rhs) { if (isnan(lhs) and isnan(rhs)) { return true; } return lhs == rhs; @@ -153,7 +154,7 @@ __device__ bool equality_compare3(Element lhs, Element rhs) * @return `true` if `lhs` == `rhs` else `false`. */ template ::value>* = nullptr> -__device__ bool equality_compare3(Element const lhs, Element const rhs) +__device__ bool equality_compare(Element const lhs, Element const rhs) { return lhs == rhs; } @@ -164,7 +165,7 @@ __device__ bool equality_compare3(Element const lhs, Element const rhs) * @tparam Nullate A cudf::nullate type describing how to check for nulls. */ template -class element_equality_comparator3 { +class element_equality_comparator { public: /** * @brief Construct type-dispatched function object for comparing equality @@ -178,10 +179,10 @@ class element_equality_comparator3 { * @param nulls_are_equal Indicates if two null elements are treated as equivalent */ __host__ __device__ - element_equality_comparator3(Nullate has_nulls, - column_device_view lhs, - column_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL) + element_equality_comparator(Nullate has_nulls, + column_device_view lhs, + column_device_view rhs, + null_equality nulls_are_equal = null_equality::EQUAL) : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal} { } @@ -208,8 +209,8 @@ class element_equality_comparator3 { } } - return equality_compare3(lhs.element(lhs_element_index), - rhs.element(rhs_element_index)); + return equality_compare(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); } template -class row_equality_comparator3 { +class row_equality_comparator { public: - row_equality_comparator3(Nullate has_nulls, - table_device_view lhs, - table_device_view rhs, - null_equality nulls_are_equal = true) + row_equality_comparator(Nullate has_nulls, + table_device_view lhs, + table_device_view rhs, + null_equality nulls_are_equal = true) : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal} { CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(), "Mismatched number of columns."); @@ -243,7 +244,7 @@ class row_equality_comparator3 { { auto equal_elements = [=](column_device_view l, column_device_view r) { return cudf::type_dispatcher(l.type(), - element_equality_comparator3{nulls, l, r, nulls_are_equal}, + element_equality_comparator{nulls, l, r, nulls_are_equal}, lhs_row_index, rhs_row_index); }; @@ -264,7 +265,7 @@ class row_equality_comparator3 { * @tparam Nullate A cudf::nullate type describing how to check for nulls. */ template -class element_relational_comparator3 { +class element_relational_comparator { public: /** * @brief Construct type-dispatched function object for performing a @@ -277,17 +278,17 @@ class element_relational_comparator3 { * @param has_nulls Indicates if either input column contains nulls. * @param null_precedence Indicates how null values are ordered with other values */ - __host__ __device__ element_relational_comparator3(Nullate has_nulls, - column_device_view lhs, - column_device_view rhs, - null_order null_precedence) + __host__ __device__ element_relational_comparator(Nullate has_nulls, + column_device_view lhs, + column_device_view rhs, + null_order null_precedence) : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, null_precedence{null_precedence} { } - __host__ __device__ element_relational_comparator3(Nullate has_nulls, - column_device_view lhs, - column_device_view rhs) + __host__ __device__ element_relational_comparator(Nullate has_nulls, + column_device_view lhs, + column_device_view rhs) : lhs{lhs}, rhs{rhs}, nulls{has_nulls} { } @@ -302,28 +303,28 @@ class element_relational_comparator3 { */ template ()>* = nullptr> - __device__ weak_ordering3 operator()(size_type lhs_element_index, - size_type rhs_element_index) const noexcept + __device__ weak_ordering operator()(size_type lhs_element_index, + size_type 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 or rhs_is_null) { // at least one is null - return null_compare3(lhs_is_null, rhs_is_null, null_precedence); + return null_compare(lhs_is_null, rhs_is_null, null_precedence); } } - return relational_compare3(lhs.element(lhs_element_index), - rhs.element(rhs_element_index)); + return relational_compare(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); } template ()>* = nullptr> - __device__ weak_ordering3 operator()(size_type lhs_element_index, size_type rhs_element_index) + __device__ weak_ordering operator()(size_type lhs_element_index, size_type rhs_element_index) { cudf_assert(false && "Attempted to compare elements of uncomparable types."); - return weak_ordering3::LESS; + return weak_ordering::LESS; } private: @@ -349,7 +350,7 @@ class element_relational_comparator3 { * @tparam Nullate A cudf::nullate type describing how to check for nulls. */ template -class row_lexicographic_comparator3 { +class row_lexicographic_comparator { public: /** * @brief Construct a function object for performing a lexicographic @@ -369,11 +370,11 @@ class row_lexicographic_comparator3 { * it is nullptr, then null precedence would be `null_order::BEFORE` for all * columns. */ - row_lexicographic_comparator3(Nullate has_nulls, - table_device_view lhs, - table_device_view rhs, - order const* column_order = nullptr, - null_order const* null_precedence = nullptr) + row_lexicographic_comparator(Nullate has_nulls, + table_device_view lhs, + table_device_view rhs, + order const* column_order = nullptr, + null_order const* null_precedence = nullptr) : _lhs{lhs}, _rhs{rhs}, _nulls{has_nulls}, @@ -399,7 +400,7 @@ class row_lexicographic_comparator3 { for (size_type i = 0; i < _lhs.num_columns(); ++i) { bool ascending = (_column_order == nullptr) or (_column_order[i] == order::ASCENDING); - weak_ordering3 state{weak_ordering3::EQUIVALENT}; + weak_ordering state{weak_ordering::EQUIVALENT}; null_order null_precedence = _null_precedence == nullptr ? null_order::BEFORE : _null_precedence[i]; @@ -410,22 +411,22 @@ class row_lexicographic_comparator3 { bool const rhs_is_null{rcol.is_null(rhs_index)}; if (lhs_is_null or rhs_is_null) { // atleast one is null - state = null_compare3(lhs_is_null, rhs_is_null, null_precedence); - if (state != weak_ordering3::EQUIVALENT) break; + state = null_compare(lhs_is_null, rhs_is_null, null_precedence); + if (state != weak_ordering::EQUIVALENT) break; } lcol = lcol.children()[0]; rcol = rcol.children()[0]; } - if (state == weak_ordering3::EQUIVALENT) { - auto comparator = element_relational_comparator3{_nulls, lcol, rcol, null_precedence}; + if (state == weak_ordering::EQUIVALENT) { + auto comparator = element_relational_comparator{_nulls, lcol, rcol, null_precedence}; state = cudf::type_dispatcher(lcol.type(), comparator, lhs_index, rhs_index); } - if (state == weak_ordering3::EQUIVALENT) { continue; } + if (state == weak_ordering::EQUIVALENT) { continue; } - return state == (ascending ? weak_ordering3::LESS : weak_ordering3::GREATER); + return state == (ascending ? weak_ordering::LESS : weak_ordering::GREATER); } return false; } @@ -436,6 +437,7 @@ class row_lexicographic_comparator3 { Nullate _nulls{}; null_order const* _null_precedence{}; order const* _column_order{}; -}; // class row_lexicographic_comparator3 +}; // class row_lexicographic_comparator +} // namespace experimental } // namespace cudf diff --git a/cpp/src/structs/utilities.cpp b/cpp/src/structs/utilities.cpp index afea8a55b16..1fc0c4457f8 100644 --- a/cpp/src/structs/utilities.cpp +++ b/cpp/src/structs/utilities.cpp @@ -208,6 +208,58 @@ flattened_table flatten_nested_columns(table_view const& input, return table_flattener{input, column_order, null_precedence, nullability}(); } +namespace experimental { + +std::tuple> verticalize_nested_columns( + table_view input) +{ + auto [table, null_masks] = superimpose_parent_nulls(input); + std::vector verticalized_columns; + for (auto const& col : table) { + if (is_nested(col.type())) { + // convert and insert + std::vector r_verticalized_columns; + std::vector flattened; + // TODO: Here I added a bogus leaf column at the beginning to help in the while loop below. + // Refactor the while loop so that it can handle the last case. + flattened.push_back(make_empty_column(type_id::INT32)->view()); + std::function recursive_child = [&](column_view const& c) { + flattened.push_back(c); + for (int child_idx = 0; child_idx < c.num_children(); ++child_idx) { + recursive_child(c.child(child_idx)); + } + }; + recursive_child(col); + int curr_col_idx = flattened.size() - 1; + column_view curr_col = flattened[curr_col_idx]; + while (curr_col_idx > 0) { + auto const& prev_col = flattened[curr_col_idx - 1]; + if (not is_nested(prev_col.type())) { + // We hit a column that's a leaf so seal this hierarchy + r_verticalized_columns.push_back(curr_col); + curr_col = prev_col; + } else { + curr_col = column_view(prev_col.type(), + prev_col.size(), + nullptr, + prev_col.null_mask(), + UNKNOWN_NULL_COUNT, + prev_col.offset(), + {curr_col}); + } + --curr_col_idx; + } + verticalized_columns.insert( + verticalized_columns.end(), r_verticalized_columns.rbegin(), r_verticalized_columns.rend()); + } else { + verticalized_columns.push_back(col); + } + } + return std::make_tuple(table_view(verticalized_columns), std::move(null_masks)); +} + +} // namespace experimental + namespace { using vector_of_columns = std::vector>; using column_index_t = typename vector_of_columns::size_type; diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu index 649951d1d5b..346a4983560 100644 --- a/cpp/tests/sort/sort2_test.cu +++ b/cpp/tests/sort/sort2_test.cu @@ -71,4 +71,109 @@ TEST_F(NewRowOp, BasicTest) cudf::test::expect_columns_equal(result1->view(), result2->view()); } +TEST_F(NewRowOp, StructTwoChildTest) +{ + using Type = int; + using column_wrapper = cudf::test::fixed_width_column_wrapper; + std::default_random_engine generator; + std::uniform_int_distribution distribution(0, 100); + + const cudf::size_type n_rows{1 << 2}; + const cudf::size_type n_cols{2}; + + // Create columns with values in the range [0,100) + std::vector columns; + columns.reserve(n_cols); + std::generate_n(std::back_inserter(columns), n_cols, [&]() { + auto elements = cudf::detail::make_counting_transform_iterator( + 0, [&](auto row) { return distribution(generator); }); + return column_wrapper(elements, elements + n_rows); + }); + + std::vector> cols; + std::transform(columns.begin(), columns.end(), std::back_inserter(cols), [](column_wrapper& col) { + return col.release(); + }); + + auto make_struct = [&](std::vector> child_cols) { + cudf::test::structs_column_wrapper struct_col(std::move(child_cols)); + return struct_col.release(); + }; + + std::vector> s2_children; + s2_children.push_back(std::move(cols[0])); + s2_children.push_back(std::move(cols[1])); + auto s2 = make_struct(std::move(s2_children)); + + cudf::test::print(s2->view()); + + // // Create table view + // auto input = cudf::table_view({struct_col}); + auto input = cudf::table_view({s2->view()}); + + auto result1 = cudf::sorted_order(input); + cudf::test::print(result1->view()); + auto result2 = cudf::detail::sorted_order2(input); + cudf::test::print(result2->view()); + cudf::test::expect_columns_equal(result1->view(), result2->view()); +} + +TEST_F(NewRowOp, SampleStructTest) +{ + using Type = int; + using column_wrapper = cudf::test::fixed_width_column_wrapper; + std::default_random_engine generator; + std::uniform_int_distribution distribution(0, 100); + + const cudf::size_type n_rows{1 << 6}; + const cudf::size_type n_cols{3}; + + // Create columns with values in the range [0,100) + std::vector columns; + columns.reserve(n_cols); + std::generate_n(std::back_inserter(columns), n_cols, [&]() { + auto elements = cudf::detail::make_counting_transform_iterator( + 0, [&](auto row) { return distribution(generator); }); + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 3 == 0 ? false : true; }); + return column_wrapper(elements, elements + n_rows, valids); + }); + + std::vector> cols; + std::transform(columns.begin(), columns.end(), std::back_inserter(cols), [](column_wrapper& col) { + return col.release(); + }); + + auto make_struct = [&](std::vector> child_cols, int nullfreq) { + std::vector struct_validity; + std::uniform_int_distribution bool_distribution(0, 10 * (nullfreq)); + std::generate_n( + std::back_inserter(struct_validity), n_rows, [&]() { return bool_distribution(generator); }); + cudf::test::structs_column_wrapper struct_col(std::move(child_cols), struct_validity); + return struct_col.release(); + }; + + std::vector> s2_children; + s2_children.push_back(std::move(cols[0])); + s2_children.push_back(std::move(cols[1])); + auto s2 = make_struct(std::move(s2_children), 1); + + std::vector> s1_children; + s1_children.push_back(std::move(s2)); + s1_children.push_back(std::move(cols[2])); + auto s1 = make_struct(std::move(s1_children), 2); + + cudf::test::print(s1->view()); + + // // Create table view + // auto input = cudf::table_view({struct_col}); + auto input = cudf::table_view({s1->view()}); + + auto result1 = cudf::sorted_order(input); + cudf::test::print(result1->view()); + auto result2 = cudf::detail::sorted_order2(input); + cudf::test::print(result2->view()); + cudf::test::expect_columns_equal(result1->view(), result2->view()); +} + CUDF_TEST_PROGRAM_MAIN() From 7fef64335d5581ebf36a490d9efcf4808c577b2c Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Thu, 27 Jan 2022 05:16:10 +0530 Subject: [PATCH 006/333] clean up most of row operators that I didn't change. --- cpp/include/cudf/table/row_operator3.cuh | 302 +---------------------- cpp/tests/sort/sort2_test.cu | 4 +- 2 files changed, 3 insertions(+), 303 deletions(-) diff --git a/cpp/include/cudf/table/row_operator3.cuh b/cpp/include/cudf/table/row_operator3.cuh index 71bbb59339e..47ab65ff0e1 100644 --- a/cpp/include/cudf/table/row_operator3.cuh +++ b/cpp/include/cudf/table/row_operator3.cuh @@ -20,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -33,307 +34,6 @@ namespace cudf { namespace experimental { -/** - * @brief Result type of the `element_relational_comparator2` function object. - * - * Indicates how two elements `a` and `b` compare with one and another. - * - * Equivalence is defined as `not (a -__device__ weak_ordering compare_elements(Element lhs, Element rhs) -{ - if (lhs < rhs) { - return weak_ordering::LESS; - } else if (rhs < lhs) { - return weak_ordering::GREATER; - } - return weak_ordering::EQUIVALENT; -} -} // namespace detail - -/** - * @brief A specialization for floating-point `Element` type relational comparison - * to derive the order of the elements with respect to `lhs`. - * - * This Specialization handles `nan` in the following order: - * `[-Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN, null] (for null_order::AFTER)` - * `[null, -Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN] (for null_order::BEFORE)` - * - * @param lhs first element - * @param rhs second element - * @return Indicates the relationship between the elements in - * the `lhs` and `rhs` columns. - */ -template ::value>* = nullptr> -__device__ weak_ordering relational_compare(Element lhs, Element rhs) -{ - if (isnan(lhs) and isnan(rhs)) { - return weak_ordering::EQUIVALENT; - } else if (isnan(rhs)) { - return weak_ordering::LESS; - } else if (isnan(lhs)) { - return weak_ordering::GREATER; - } - - return detail::compare_elements(lhs, rhs); -} - -/** - * @brief Compare the nulls according to null order. - * - * @param lhs_is_null boolean representing if lhs is null - * @param rhs_is_null boolean representing if lhs is null - * @param null_precedence null order - * @return Indicates the relationship between null in lhs and rhs columns. - */ -inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_order null_precedence) -{ - if (lhs_is_null and rhs_is_null) { // null ::value>* = nullptr> -__device__ weak_ordering relational_compare(Element lhs, Element rhs) -{ - return detail::compare_elements(lhs, rhs); -} - -/** - * @brief A specialization for floating-point `Element` type to check if - * `lhs` is equivalent to `rhs`. `nan == nan`. - * - * @param lhs first element - * @param rhs second element - * @return `true` if `lhs` == `rhs` else `false`. - */ -template ::value>* = nullptr> -__device__ bool equality_compare(Element lhs, Element rhs) -{ - if (isnan(lhs) and isnan(rhs)) { return true; } - return lhs == rhs; -} - -/** - * @brief A specialization for non-floating-point `Element` type to check if - * `lhs` is equivalent to `rhs`. - * - * @param lhs first element - * @param rhs second element - * @return `true` if `lhs` == `rhs` else `false`. - */ -template ::value>* = nullptr> -__device__ bool equality_compare(Element const lhs, Element const rhs) -{ - return lhs == rhs; -} - -/** - * @brief Performs an equality comparison between two elements in two columns. - * - * @tparam Nullate A cudf::nullate type describing how to check for nulls. - */ -template -class element_equality_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 - */ - __host__ __device__ - element_equality_comparator(Nullate has_nulls, - column_device_view lhs, - column_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL) - : 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 both lhs and rhs element are both nulls and `nulls_are_equal` is true, or equal - */ - template ()>* = nullptr> - __device__ bool operator()(size_type lhs_element_index, - size_type 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 ()>* = nullptr> - __device__ bool operator()(size_type lhs_element_index, size_type rhs_element_index) - { - cudf_assert(false && "Attempted to compare elements of uncomparable types."); - return false; - } - - private: - column_device_view lhs; - column_device_view rhs; - Nullate nulls; - null_equality nulls_are_equal; -}; - -template -class row_equality_comparator { - public: - row_equality_comparator(Nullate has_nulls, - table_device_view lhs, - table_device_view rhs, - null_equality nulls_are_equal = true) - : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal} - { - CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(), "Mismatched number of columns."); - } - - __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept - { - auto equal_elements = [=](column_device_view l, column_device_view r) { - return cudf::type_dispatcher(l.type(), - element_equality_comparator{nulls, l, r, nulls_are_equal}, - lhs_row_index, - rhs_row_index); - }; - - return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements); - } - - private: - table_device_view lhs; - table_device_view rhs; - Nullate nulls; - null_equality nulls_are_equal; -}; - -/** - * @brief Performs a relational comparison between two elements in two columns. - * - * @tparam Nullate A cudf::nullate type describing how to check for nulls. - */ -template -class element_relational_comparator { - public: - /** - * @brief Construct type-dispatched function object for performing a - * relational comparison between two elements. - * - * @note `lhs` and `rhs` may be the same. - * - * @param lhs The column containing the first element - * @param rhs The column containing the second element (may be the same as lhs) - * @param has_nulls Indicates if either input column contains nulls. - * @param null_precedence Indicates how null values are ordered with other values - */ - __host__ __device__ element_relational_comparator(Nullate has_nulls, - column_device_view lhs, - column_device_view rhs, - null_order null_precedence) - : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, null_precedence{null_precedence} - { - } - - __host__ __device__ element_relational_comparator(Nullate has_nulls, - column_device_view lhs, - column_device_view rhs) - : lhs{lhs}, rhs{rhs}, nulls{has_nulls} - { - } - - /** - * @brief Performs a relational comparison between the specified elements - * - * @param lhs_element_index The index of the first element - * @param rhs_element_index The index of the second element - * @return Indicates the relationship between the elements in - * the `lhs` and `rhs` columns. - */ - template ()>* = nullptr> - __device__ weak_ordering operator()(size_type lhs_element_index, - size_type 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 or rhs_is_null) { // at least one is null - return null_compare(lhs_is_null, rhs_is_null, null_precedence); - } - } - - return relational_compare(lhs.element(lhs_element_index), - rhs.element(rhs_element_index)); - } - - template ()>* = nullptr> - __device__ weak_ordering operator()(size_type lhs_element_index, size_type rhs_element_index) - { - cudf_assert(false && "Attempted to compare elements of uncomparable types."); - return weak_ordering::LESS; - } - - private: - column_device_view lhs; - column_device_view rhs; - Nullate nulls; - null_order null_precedence{}; -}; - /** * @brief Computes whether one row is lexicographically *less* than another row. * diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu index 346a4983560..fb78e4a4629 100644 --- a/cpp/tests/sort/sort2_test.cu +++ b/cpp/tests/sort/sort2_test.cu @@ -123,7 +123,7 @@ TEST_F(NewRowOp, SampleStructTest) using Type = int; using column_wrapper = cudf::test::fixed_width_column_wrapper; std::default_random_engine generator; - std::uniform_int_distribution distribution(0, 100); + std::uniform_int_distribution distribution(0, 20); const cudf::size_type n_rows{1 << 6}; const cudf::size_type n_cols{3}; @@ -135,7 +135,7 @@ TEST_F(NewRowOp, SampleStructTest) auto elements = cudf::detail::make_counting_transform_iterator( 0, [&](auto row) { return distribution(generator); }); auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 3 == 0 ? false : true; }); + 0, [](auto i) { return i % 5 == 0 ? false : true; }); return column_wrapper(elements, elements + n_rows, valids); }); From 930d8de8ca9b06322c9ef15849d0ec9b968fafcd Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Fri, 28 Jan 2022 04:03:56 +0530 Subject: [PATCH 007/333] Sliced column test --- cpp/tests/sort/sort2_test.cu | 38 ++++++++++++++++++++++++------------ 1 file changed, 26 insertions(+), 12 deletions(-) diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu index fb78e4a4629..bd24befc88e 100644 --- a/cpp/tests/sort/sort2_test.cu +++ b/cpp/tests/sort/sort2_test.cu @@ -14,12 +14,12 @@ #include -struct NewRowOp : public cudf::test::BaseFixture { +struct NewRowOpTest : public cudf::test::BaseFixture { }; #include -TEST_F(NewRowOp, BasicTest) +TEST_F(NewRowOpTest, DeepStruct) { using Type = int; using column_wrapper = cudf::test::fixed_width_column_wrapper; @@ -64,14 +64,16 @@ TEST_F(NewRowOp, BasicTest) // auto input = cudf::table_view({struct_col}); auto input = cudf::table(std::move(child_cols)); - auto result1 = cudf::sorted_order(input); - // cudf::test::print(result1->view()); - auto result2 = cudf::detail::sorted_order2(input); - // cudf::test::print(result2->view()); + auto sliced_input = cudf::slice(input, {7, input.num_rows() - 12}); + + auto result1 = cudf::sorted_order(sliced_input); + cudf::test::print(result1->view()); + auto result2 = cudf::detail::sorted_order2(sliced_input); + cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); } -TEST_F(NewRowOp, StructTwoChildTest) +TEST_F(NewRowOpTest, StructTwoChildTest) { using Type = int; using column_wrapper = cudf::test::fixed_width_column_wrapper; @@ -118,15 +120,15 @@ TEST_F(NewRowOp, StructTwoChildTest) cudf::test::expect_columns_equal(result1->view(), result2->view()); } -TEST_F(NewRowOp, SampleStructTest) +TEST_F(NewRowOpTest, SampleStructTest) { using Type = int; using column_wrapper = cudf::test::fixed_width_column_wrapper; std::default_random_engine generator; - std::uniform_int_distribution distribution(0, 20); + std::uniform_int_distribution distribution(0, 10); const cudf::size_type n_rows{1 << 6}; - const cudf::size_type n_cols{3}; + const cudf::size_type n_cols{6}; // Create columns with values in the range [0,100) std::vector columns; @@ -135,7 +137,7 @@ TEST_F(NewRowOp, SampleStructTest) auto elements = cudf::detail::make_counting_transform_iterator( 0, [&](auto row) { return distribution(generator); }); auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 5 == 0 ? false : true; }); + 0, [](auto i) { return i % 7 == 0 ? false : true; }); return column_wrapper(elements, elements + n_rows, valids); }); @@ -165,9 +167,21 @@ TEST_F(NewRowOp, SampleStructTest) cudf::test::print(s1->view()); + std::vector> s22_children; + s22_children.push_back(std::move(cols[3])); + s22_children.push_back(std::move(cols[4])); + auto s22 = make_struct(std::move(s22_children), 1); + + std::vector> s12_children; + s12_children.push_back(std::move(cols[5])); + s12_children.push_back(std::move(s22)); + auto s12 = make_struct(std::move(s12_children), 2); + + cudf::test::print(s1->view()); + // // Create table view // auto input = cudf::table_view({struct_col}); - auto input = cudf::table_view({s1->view()}); + auto input = cudf::table_view({s1->view(), s12->view()}); auto result1 = cudf::sorted_order(input); cudf::test::print(result1->view()); From 0ecc4f87c20df2a03c1f1d4b869632058843ad58 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Sat, 29 Jan 2022 00:37:59 +0530 Subject: [PATCH 008/333] column order and null precendence support --- cpp/include/cudf/detail/structs/utilities.hpp | 5 +- cpp/include/cudf/sort2.cuh | 22 +++-- cpp/src/structs/utilities.cpp | 24 ++++- cpp/tests/sort/sort2_test.cu | 94 ++++++++++--------- 4 files changed, 89 insertions(+), 56 deletions(-) diff --git a/cpp/include/cudf/detail/structs/utilities.hpp b/cpp/include/cudf/detail/structs/utilities.hpp index 23601548cbf..1f29a247121 100644 --- a/cpp/include/cudf/detail/structs/utilities.hpp +++ b/cpp/include/cudf/detail/structs/utilities.hpp @@ -151,8 +151,9 @@ flattened_table flatten_nested_columns( column_nullability nullability = column_nullability::MATCH_INCOMING); namespace experimental { -std::tuple> verticalize_nested_columns( - table_view input); +flattened_table verticalize_nested_columns(table_view input, + std::vector const& column_order, + std::vector const& null_precedence); } /** diff --git a/cpp/include/cudf/sort2.cuh b/cpp/include/cudf/sort2.cuh index 125d696d055..6710615aa35 100644 --- a/cpp/include/cudf/sort2.cuh +++ b/cpp/include/cudf/sort2.cuh @@ -46,8 +46,10 @@ namespace detail { template std::unique_ptr sorted_order2( table_view input, - rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + std::vector const& column_order = {}, + std::vector const& null_precedence = {}, + rmm::cuda_stream_view stream = rmm::cuda_stream_default, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { if (input.num_rows() == 0 or input.num_columns() == 0) { return cudf::make_numeric_column(data_type(type_to_id()), 0); @@ -61,12 +63,18 @@ std::unique_ptr sorted_order2( mutable_indices_view.end(), 0); - auto [vertical_table, nullmasks] = - cudf::structs::detail::experimental::verticalize_nested_columns(input); - auto device_table = table_device_view::create(vertical_table, stream); + auto verticalized = cudf::structs::detail::experimental::verticalize_nested_columns( + input, column_order, null_precedence); + auto device_table = table_device_view::create(verticalized.flattened_columns(), stream); + auto const d_column_order = make_device_uvector_async(verticalized.orders(), stream); + auto const d_null_precedence = make_device_uvector_async(verticalized.null_orders(), stream); + // auto const comparator = row_lexicographic_comparator2(*device_table, *device_table); - auto const comparator = experimental::row_lexicographic_comparator( - nullate::DYNAMIC{true}, *device_table, *device_table); + auto const comparator = experimental::row_lexicographic_comparator(nullate::DYNAMIC{true}, + *device_table, + *device_table, + d_column_order.data(), + d_null_precedence.data()); thrust::sort(rmm::exec_policy(stream), mutable_indices_view.begin(), diff --git a/cpp/src/structs/utilities.cpp b/cpp/src/structs/utilities.cpp index 1fc0c4457f8..42c167466c1 100644 --- a/cpp/src/structs/utilities.cpp +++ b/cpp/src/structs/utilities.cpp @@ -210,12 +210,16 @@ flattened_table flatten_nested_columns(table_view const& input, namespace experimental { -std::tuple> verticalize_nested_columns( - table_view input) +flattened_table verticalize_nested_columns(table_view input, + std::vector const& column_order, + std::vector const& null_precedence) { auto [table, null_masks] = superimpose_parent_nulls(input); std::vector verticalized_columns; - for (auto const& col : table) { + std::vector new_column_order; + std::vector new_null_precedence; + 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())) { // convert and insert std::vector r_verticalized_columns; @@ -251,11 +255,23 @@ std::tuple> verticalize_nested } verticalized_columns.insert( verticalized_columns.end(), r_verticalized_columns.rbegin(), r_verticalized_columns.rend()); + if (not column_order.empty()) { + new_column_order.insert( + new_column_order.end(), r_verticalized_columns.size(), column_order[col_idx]); + } + if (not null_precedence.empty()) { + new_null_precedence.insert( + new_null_precedence.end(), r_verticalized_columns.size(), null_precedence[col_idx]); + } } else { verticalized_columns.push_back(col); } } - return std::make_tuple(table_view(verticalized_columns), std::move(null_masks)); + return flattened_table(table_view(verticalized_columns), + new_column_order, + new_null_precedence, + {}, + std::move(null_masks)); } } // namespace experimental diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu index bd24befc88e..98bd28fb181 100644 --- a/cpp/tests/sort/sort2_test.cu +++ b/cpp/tests/sort/sort2_test.cu @@ -19,26 +19,23 @@ struct NewRowOpTest : public cudf::test::BaseFixture { #include -TEST_F(NewRowOpTest, DeepStruct) +TEST_F(NewRowOpTest, BasicStructTwoChild) { using Type = int; using column_wrapper = cudf::test::fixed_width_column_wrapper; std::default_random_engine generator; std::uniform_int_distribution distribution(0, 100); - const cudf::size_type n_rows{1 << 6}; - const cudf::size_type n_cols{1}; - const cudf::size_type depth{8}; + const cudf::size_type n_rows{1 << 2}; + const cudf::size_type n_cols{2}; // Create columns with values in the range [0,100) std::vector columns; columns.reserve(n_cols); - std::generate_n(std::back_inserter(columns), n_cols, [&, n_rows]() { + std::generate_n(std::back_inserter(columns), n_cols, [&]() { auto elements = cudf::detail::make_counting_transform_iterator( 0, [&](auto row) { return distribution(generator); }); - auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 3 == 0 ? false : true; }); - return column_wrapper(elements, elements + n_rows, valids); + return column_wrapper(elements, elements + n_rows); }); std::vector> cols; @@ -46,50 +43,49 @@ TEST_F(NewRowOpTest, DeepStruct) return col.release(); }); - std::vector> child_cols = std::move(cols); - // Lets add some layers - for (int i = 0; i < depth; i++) { - std::vector struct_validity; - std::uniform_int_distribution bool_distribution(0, 10 * (i + 1)); - std::generate_n( - std::back_inserter(struct_validity), n_rows, [&]() { return bool_distribution(generator); }); - cudf::test::structs_column_wrapper struct_col(std::move(child_cols), struct_validity); - child_cols = std::vector>{}; - child_cols.push_back(struct_col.release()); - } + auto make_struct = [&](std::vector> child_cols) { + cudf::test::structs_column_wrapper struct_col(std::move(child_cols)); + return struct_col.release(); + }; - cudf::test::print(child_cols[0]->view()); + std::vector> s2_children; + s2_children.push_back(std::move(cols[0])); + s2_children.push_back(std::move(cols[1])); + auto s2 = make_struct(std::move(s2_children)); + + cudf::test::print(s2->view()); // // Create table view // auto input = cudf::table_view({struct_col}); - auto input = cudf::table(std::move(child_cols)); - - auto sliced_input = cudf::slice(input, {7, input.num_rows() - 12}); + auto input = cudf::table_view({s2->view()}); - auto result1 = cudf::sorted_order(sliced_input); + auto result1 = cudf::sorted_order(input); cudf::test::print(result1->view()); - auto result2 = cudf::detail::sorted_order2(sliced_input); + auto result2 = cudf::detail::sorted_order2(input); cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); } -TEST_F(NewRowOpTest, StructTwoChildTest) +TEST_F(NewRowOpTest, DeepStruct) { using Type = int; using column_wrapper = cudf::test::fixed_width_column_wrapper; std::default_random_engine generator; std::uniform_int_distribution distribution(0, 100); - const cudf::size_type n_rows{1 << 2}; - const cudf::size_type n_cols{2}; + const cudf::size_type n_rows{1 << 6}; + const cudf::size_type n_cols{1}; + const cudf::size_type depth{8}; // Create columns with values in the range [0,100) std::vector columns; columns.reserve(n_cols); - std::generate_n(std::back_inserter(columns), n_cols, [&]() { + std::generate_n(std::back_inserter(columns), n_cols, [&, n_rows]() { auto elements = cudf::detail::make_counting_transform_iterator( 0, [&](auto row) { return distribution(generator); }); - return column_wrapper(elements, elements + n_rows); + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 3 == 0 ? false : true; }); + return column_wrapper(elements, elements + n_rows, valids); }); std::vector> cols; @@ -97,25 +93,29 @@ TEST_F(NewRowOpTest, StructTwoChildTest) return col.release(); }); - auto make_struct = [&](std::vector> child_cols) { - cudf::test::structs_column_wrapper struct_col(std::move(child_cols)); - return struct_col.release(); - }; - - std::vector> s2_children; - s2_children.push_back(std::move(cols[0])); - s2_children.push_back(std::move(cols[1])); - auto s2 = make_struct(std::move(s2_children)); + std::vector> child_cols = std::move(cols); + // Lets add some layers + for (int i = 0; i < depth; i++) { + std::vector struct_validity; + std::uniform_int_distribution bool_distribution(0, 10 * (i + 1)); + std::generate_n( + std::back_inserter(struct_validity), n_rows, [&]() { return bool_distribution(generator); }); + cudf::test::structs_column_wrapper struct_col(std::move(child_cols), struct_validity); + child_cols = std::vector>{}; + child_cols.push_back(struct_col.release()); + } - cudf::test::print(s2->view()); + cudf::test::print(child_cols[0]->view()); // // Create table view // auto input = cudf::table_view({struct_col}); - auto input = cudf::table_view({s2->view()}); + auto input = cudf::table(std::move(child_cols)); - auto result1 = cudf::sorted_order(input); + auto sliced_input = cudf::slice(input, {7, input.num_rows() - 12}); + + auto result1 = cudf::sorted_order(sliced_input); cudf::test::print(result1->view()); - auto result2 = cudf::detail::sorted_order2(input); + auto result2 = cudf::detail::sorted_order2(sliced_input); cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); } @@ -188,6 +188,14 @@ TEST_F(NewRowOpTest, SampleStructTest) auto result2 = cudf::detail::sorted_order2(input); cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); + + std::vector col_order = {cudf::order::DESCENDING, cudf::order::ASCENDING}; + std::vector null_order = {cudf::null_order::BEFORE, cudf::null_order::AFTER}; + result1 = cudf::sorted_order(input, col_order, null_order); + result2 = cudf::detail::sorted_order2(input, col_order, null_order); + cudf::test::print(result1->view()); + cudf::test::print(result2->view()); + cudf::test::expect_columns_equal(result1->view(), result2->view()); } CUDF_TEST_PROGRAM_MAIN() From ff36d2dbe3e535625ee011c2ff8656e7a78b0bed Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Sat, 29 Jan 2022 04:37:46 +0530 Subject: [PATCH 009/333] Manually managed stack --- cpp/include/cudf/table/row_operator2.cuh | 79 ++++++++++++++++++++---- 1 file changed, 68 insertions(+), 11 deletions(-) diff --git a/cpp/include/cudf/table/row_operator2.cuh b/cpp/include/cudf/table/row_operator2.cuh index 9edb60cc112..1288cef70a9 100644 --- a/cpp/include/cudf/table/row_operator2.cuh +++ b/cpp/include/cudf/table/row_operator2.cuh @@ -319,6 +319,35 @@ class element_relational_comparator2 { null_order null_precedence; }; +template +struct device_stack { + __device__ device_stack(T* stack_storage, int capacity) + : stack(stack_storage), capacity(capacity), size(0) + { + } + __device__ void push(T const& val) + { + cudf_assert(size < capacity and "Stack overflow"); + stack[size++] = val; + } + __device__ T pop() + { + cudf_assert(size > 0 and "Stack underflow"); + return stack[--size]; + } + __device__ T top() + { + cudf_assert(size > 0 and "Stack underflow"); + return stack[size - 1]; + } + __device__ bool empty() { return size == 0; } + + private: + T* stack; + int capacity; + int size; +}; + /** * @brief Computes whether one row is lexicographically *less* than another row. * @@ -376,31 +405,59 @@ class row_lexicographic_comparator2 { */ __device__ bool operator()(size_type lhs_index, size_type rhs_index) const noexcept { + using stack_value_type = + thrust::tuple; + stack_value_type stack_storage[10]; + for (size_type i = 0; i < _lhs.num_columns(); ++i) { + device_stack stack(stack_storage, 9); bool ascending = (_column_order == nullptr) or (_column_order[i] == order::ASCENDING); weak_ordering2 state{weak_ordering2::EQUIVALENT}; null_order null_precedence = _null_precedence == nullptr ? null_order::BEFORE : _null_precedence[i]; - column_device_view lcol = _lhs.column(i); - column_device_view rcol = _rhs.column(i); - while (lcol.type().id() == type_id::STRUCT) { - bool const lhs_is_null{lcol.is_null(lhs_index)}; - bool const rhs_is_null{rcol.is_null(rhs_index)}; + column_device_view const* lcol = _lhs.begin() + i; + column_device_view const* rcol = _rhs.begin() + i; + size_t curr_child = 0; + + while (true) { + bool const lhs_is_null{lcol->is_null(lhs_index)}; + bool const rhs_is_null{rcol->is_null(rhs_index)}; if (lhs_is_null or rhs_is_null) { // atleast one is null state = null_compare2(lhs_is_null, rhs_is_null, null_precedence); if (state != weak_ordering2::EQUIVALENT) break; + } else if (lcol->type().id() != type_id::STRUCT) { + auto comparator = + element_relational_comparator2{*lcol, *rcol, null_precedence}; + state = cudf::type_dispatcher(lcol->type(), comparator, lhs_index, rhs_index); + if (state != weak_ordering2::EQUIVALENT) break; } - lcol = lcol.children()[0]; - rcol = rcol.children()[0]; - } + // Reaching here means the nullability was same and we need to continue comparing + if (lcol->type().id() == type_id::STRUCT) { + stack.push({lcol, rcol, 0}); + } else { + // unwind stack until we reach a struct level with children still left to compare + bool completed_comparison = false; + do { + if (stack.empty()) { + completed_comparison = true; + break; + } + thrust::tie(lcol, rcol, curr_child) = stack.pop(); + } while (lcol->num_child_columns() <= curr_child + 1); + if (completed_comparison) { break; } + stack.push({lcol, rcol, curr_child + 1}); + // break; + } + + // The top of the stack now is where we have to continue comparing from + thrust::tie(lcol, rcol, curr_child) = stack.top(); - if (state == weak_ordering2::EQUIVALENT) { - auto comparator = element_relational_comparator2{lcol, rcol, null_precedence}; - state = cudf::type_dispatcher(lcol.type(), comparator, lhs_index, rhs_index); + lcol = &lcol->children()[curr_child]; + rcol = &rcol->children()[curr_child]; } if (state == weak_ordering2::EQUIVALENT) { continue; } From cd0f93885efb17f643e269e4d73e9f52c216b30c Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Thu, 3 Feb 2022 01:01:19 +0530 Subject: [PATCH 010/333] New depth based method to avoid superimpose nulls sliced no longer works --- cpp/include/cudf/detail/structs/utilities.hpp | 7 +-- cpp/include/cudf/sort2.cuh | 6 ++- cpp/include/cudf/table/row_operator3.cuh | 19 +++++++- cpp/src/structs/utilities.cpp | 34 +++++++++----- cpp/tests/sort/sort2_test.cu | 45 +++++++++++++------ 5 files changed, 80 insertions(+), 31 deletions(-) diff --git a/cpp/include/cudf/detail/structs/utilities.hpp b/cpp/include/cudf/detail/structs/utilities.hpp index 1f29a247121..6b3a297cbaa 100644 --- a/cpp/include/cudf/detail/structs/utilities.hpp +++ b/cpp/include/cudf/detail/structs/utilities.hpp @@ -151,9 +151,10 @@ flattened_table flatten_nested_columns( column_nullability nullability = column_nullability::MATCH_INCOMING); namespace experimental { -flattened_table verticalize_nested_columns(table_view input, - std::vector const& column_order, - std::vector const& null_precedence); +std::tuple> verticalize_nested_columns( + table_view input, + std::vector const& column_order, + std::vector const& null_precedence); } /** diff --git a/cpp/include/cudf/sort2.cuh b/cpp/include/cudf/sort2.cuh index 6710615aa35..e612a5340d9 100644 --- a/cpp/include/cudf/sort2.cuh +++ b/cpp/include/cudf/sort2.cuh @@ -18,6 +18,7 @@ #include #include +#include #include #include #include @@ -51,6 +52,7 @@ std::unique_ptr sorted_order2( rmm::cuda_stream_view stream = rmm::cuda_stream_default, rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) { + CUDF_FUNC_RANGE(); if (input.num_rows() == 0 or input.num_columns() == 0) { return cudf::make_numeric_column(data_type(type_to_id()), 0); } @@ -63,16 +65,18 @@ std::unique_ptr sorted_order2( mutable_indices_view.end(), 0); - auto verticalized = cudf::structs::detail::experimental::verticalize_nested_columns( + auto [verticalized, depths] = cudf::structs::detail::experimental::verticalize_nested_columns( input, column_order, null_precedence); auto device_table = table_device_view::create(verticalized.flattened_columns(), stream); auto const d_column_order = make_device_uvector_async(verticalized.orders(), stream); auto const d_null_precedence = make_device_uvector_async(verticalized.null_orders(), stream); + auto const d_depths = make_device_uvector_async(depths, stream); // auto const comparator = row_lexicographic_comparator2(*device_table, *device_table); auto const comparator = experimental::row_lexicographic_comparator(nullate::DYNAMIC{true}, *device_table, *device_table, + d_depths.data(), d_column_order.data(), d_null_precedence.data()); diff --git a/cpp/include/cudf/table/row_operator3.cuh b/cpp/include/cudf/table/row_operator3.cuh index 47ab65ff0e1..6f40ff0c414 100644 --- a/cpp/include/cudf/table/row_operator3.cuh +++ b/cpp/include/cudf/table/row_operator3.cuh @@ -73,11 +73,13 @@ class row_lexicographic_comparator { row_lexicographic_comparator(Nullate has_nulls, table_device_view lhs, table_device_view rhs, + int const* depth = nullptr, order const* column_order = nullptr, null_order const* null_precedence = nullptr) : _lhs{lhs}, _rhs{rhs}, _nulls{has_nulls}, + _depth{depth}, _column_order{column_order}, _null_precedence{null_precedence} { @@ -97,7 +99,15 @@ class row_lexicographic_comparator { */ __device__ bool operator()(size_type lhs_index, size_type rhs_index) const noexcept { + int last_null_depth = std::numeric_limits::max(); for (size_type i = 0; i < _lhs.num_columns(); ++i) { + if (_depth[i] > last_null_depth) { + continue; + } else { + last_null_depth = std::numeric_limits::max(); + } + + bool continue_to_next_col = false; bool ascending = (_column_order == nullptr) or (_column_order[i] == order::ASCENDING); weak_ordering state{weak_ordering::EQUIVALENT}; @@ -106,19 +116,25 @@ class row_lexicographic_comparator { column_device_view lcol = _lhs.column(i); column_device_view rcol = _rhs.column(i); + int depth = _depth[i]; while (lcol.type().id() == type_id::STRUCT) { bool const lhs_is_null{lcol.is_null(lhs_index)}; bool const rhs_is_null{rcol.is_null(rhs_index)}; if (lhs_is_null or rhs_is_null) { // atleast one is null state = null_compare(lhs_is_null, rhs_is_null, null_precedence); - if (state != weak_ordering::EQUIVALENT) break; + if (state == weak_ordering::EQUIVALENT) { continue_to_next_col = true; } + last_null_depth = depth; + break; } lcol = lcol.children()[0]; rcol = rcol.children()[0]; + ++depth; } + if (continue_to_next_col) { continue; } + if (state == weak_ordering::EQUIVALENT) { auto comparator = element_relational_comparator{_nulls, lcol, rcol, null_precedence}; state = cudf::type_dispatcher(lcol.type(), comparator, lhs_index, rhs_index); @@ -137,6 +153,7 @@ class row_lexicographic_comparator { Nullate _nulls{}; null_order const* _null_precedence{}; order const* _column_order{}; + int const* _depth; }; // class row_lexicographic_comparator } // namespace experimental diff --git a/cpp/src/structs/utilities.cpp b/cpp/src/structs/utilities.cpp index 42c167466c1..3ad9917db56 100644 --- a/cpp/src/structs/utilities.cpp +++ b/cpp/src/structs/utilities.cpp @@ -210,30 +210,37 @@ flattened_table flatten_nested_columns(table_view const& input, namespace experimental { -flattened_table verticalize_nested_columns(table_view input, - std::vector const& column_order, - std::vector const& null_precedence) +std::tuple> verticalize_nested_columns( + table_view input, + std::vector const& column_order, + std::vector const& null_precedence) { - auto [table, null_masks] = superimpose_parent_nulls(input); + // auto [table, null_masks] = superimpose_parent_nulls(input); + auto table = input; 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())) { // convert and insert std::vector r_verticalized_columns; + std::vector r_verticalized_col_depths; std::vector flattened; + std::vector depths; // TODO: Here I added a bogus leaf column at the beginning to help in the while loop below. // Refactor the while loop so that it can handle the last case. flattened.push_back(make_empty_column(type_id::INT32)->view()); - std::function recursive_child = [&](column_view const& c) { + std::function recursive_child = [&](column_view const& c, + int depth) { flattened.push_back(c); + depths.push_back(depth); for (int child_idx = 0; child_idx < c.num_children(); ++child_idx) { - recursive_child(c.child(child_idx)); + recursive_child(c.child(child_idx), depth + 1); } }; - recursive_child(col); + recursive_child(col, 0); int curr_col_idx = flattened.size() - 1; column_view curr_col = flattened[curr_col_idx]; while (curr_col_idx > 0) { @@ -241,6 +248,7 @@ flattened_table verticalize_nested_columns(table_view input, if (not is_nested(prev_col.type())) { // We hit a column that's a leaf so seal this hierarchy r_verticalized_columns.push_back(curr_col); + r_verticalized_col_depths.push_back(depths[curr_col_idx - 1]); curr_col = prev_col; } else { curr_col = column_view(prev_col.type(), @@ -255,6 +263,9 @@ flattened_table verticalize_nested_columns(table_view input, } verticalized_columns.insert( verticalized_columns.end(), r_verticalized_columns.rbegin(), r_verticalized_columns.rend()); + verticalized_col_depths.insert(verticalized_col_depths.end(), + r_verticalized_col_depths.rbegin(), + r_verticalized_col_depths.rend()); if (not column_order.empty()) { new_column_order.insert( new_column_order.end(), r_verticalized_columns.size(), column_order[col_idx]); @@ -267,11 +278,10 @@ flattened_table verticalize_nested_columns(table_view input, verticalized_columns.push_back(col); } } - return flattened_table(table_view(verticalized_columns), - new_column_order, - new_null_precedence, - {}, - std::move(null_masks)); + return std::make_tuple( + flattened_table( + table_view(verticalized_columns), new_column_order, new_null_precedence, {}, {}), + std::move(verticalized_col_depths)); } } // namespace experimental diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu index 98bd28fb181..56ef562be39 100644 --- a/cpp/tests/sort/sort2_test.cu +++ b/cpp/tests/sort/sort2_test.cu @@ -1,6 +1,7 @@ #include #include #include +#include #include #include @@ -26,7 +27,7 @@ TEST_F(NewRowOpTest, BasicStructTwoChild) std::default_random_engine generator; std::uniform_int_distribution distribution(0, 100); - const cudf::size_type n_rows{1 << 2}; + const cudf::size_type n_rows{1 << 4}; const cudf::size_type n_cols{2}; // Create columns with values in the range [0,100) @@ -35,7 +36,9 @@ TEST_F(NewRowOpTest, BasicStructTwoChild) std::generate_n(std::back_inserter(columns), n_cols, [&]() { auto elements = cudf::detail::make_counting_transform_iterator( 0, [&](auto row) { return distribution(generator); }); - return column_wrapper(elements, elements + n_rows); + auto valids = cudf::detail::make_counting_transform_iterator( + 0, [](auto i) { return i % 4 == 0 ? false : true; }); + return column_wrapper(elements, elements + n_rows, valids); }); std::vector> cols; @@ -43,15 +46,25 @@ TEST_F(NewRowOpTest, BasicStructTwoChild) return col.release(); }); - auto make_struct = [&](std::vector> child_cols) { + auto make_struct = [&](std::vector> child_cols, int nullfreq) { + // std::vector struct_validity; + std::uniform_int_distribution bool_distribution(0, 10 * (nullfreq)); + // std::generate_n( + // std::back_inserter(struct_validity), n_rows, [&]() { return bool_distribution(generator); + // }); + auto null_iter = cudf::detail::make_counting_transform_iterator( + 0, [&](int i) { return bool_distribution(generator); }); + cudf::test::structs_column_wrapper struct_col(std::move(child_cols)); - return struct_col.release(); + auto struct_ = struct_col.release(); + struct_->set_null_mask(cudf::test::detail::make_null_mask(null_iter, null_iter + n_rows)); + return struct_; }; std::vector> s2_children; s2_children.push_back(std::move(cols[0])); s2_children.push_back(std::move(cols[1])); - auto s2 = make_struct(std::move(s2_children)); + auto s2 = make_struct(std::move(s2_children), 1); cudf::test::print(s2->view()); @@ -111,11 +124,11 @@ TEST_F(NewRowOpTest, DeepStruct) // auto input = cudf::table_view({struct_col}); auto input = cudf::table(std::move(child_cols)); - auto sliced_input = cudf::slice(input, {7, input.num_rows() - 12}); + // auto sliced_input = cudf::slice(input, {7, input.num_rows() - 12}); - auto result1 = cudf::sorted_order(sliced_input); + auto result1 = cudf::sorted_order(input); cudf::test::print(result1->view()); - auto result2 = cudf::detail::sorted_order2(sliced_input); + auto result2 = cudf::detail::sorted_order2(input); cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); } @@ -136,8 +149,9 @@ TEST_F(NewRowOpTest, SampleStructTest) std::generate_n(std::back_inserter(columns), n_cols, [&]() { auto elements = cudf::detail::make_counting_transform_iterator( 0, [&](auto row) { return distribution(generator); }); + int start = distribution(generator); auto valids = cudf::detail::make_counting_transform_iterator( - 0, [](auto i) { return i % 7 == 0 ? false : true; }); + 0, [&](auto i) { return (i + start) % 7 == 0 ? false : true; }); return column_wrapper(elements, elements + n_rows, valids); }); @@ -149,10 +163,13 @@ TEST_F(NewRowOpTest, SampleStructTest) auto make_struct = [&](std::vector> child_cols, int nullfreq) { std::vector struct_validity; std::uniform_int_distribution bool_distribution(0, 10 * (nullfreq)); - std::generate_n( - std::back_inserter(struct_validity), n_rows, [&]() { return bool_distribution(generator); }); - cudf::test::structs_column_wrapper struct_col(std::move(child_cols), struct_validity); - return struct_col.release(); + auto null_iter = cudf::detail::make_counting_transform_iterator( + 0, [&](int i) { return bool_distribution(generator); }); + + cudf::test::structs_column_wrapper struct_col(std::move(child_cols)); + auto struct_ = struct_col.release(); + struct_->set_null_mask(cudf::test::detail::make_null_mask(null_iter, null_iter + n_rows)); + return struct_; }; std::vector> s2_children; @@ -177,7 +194,7 @@ TEST_F(NewRowOpTest, SampleStructTest) s12_children.push_back(std::move(s22)); auto s12 = make_struct(std::move(s12_children), 2); - cudf::test::print(s1->view()); + cudf::test::print(s12->view()); // // Create table view // auto input = cudf::table_view({struct_col}); From 7b8e060ab1bd646408154882c349de1b0f2502b0 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Thu, 3 Feb 2022 01:06:06 +0530 Subject: [PATCH 011/333] Put sort2 impl in separate TU --- cpp/CMakeLists.txt | 1 + .../compare/comparator_benchmark.cu | 2 +- cpp/include/cudf/sort2.cuh | 43 +-------- cpp/src/sort/sort2.cu | 93 +++++++++++++++++++ cpp/tests/sort/sort2_test.cu | 8 +- 5 files changed, 102 insertions(+), 45 deletions(-) create mode 100644 cpp/src/sort/sort2.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 4db9f6de4d5..50d664a0661 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -391,6 +391,7 @@ add_library( src/sort/segmented_sort.cu src/sort/sort_column.cu src/sort/sort.cu + src/sort/sort2.cu src/sort/stable_sort_column.cu src/sort/stable_sort.cu src/stream_compaction/apply_boolean_mask.cu diff --git a/cpp/benchmarks/compare/comparator_benchmark.cu b/cpp/benchmarks/compare/comparator_benchmark.cu index e4f690102b0..8c9accc0a00 100644 --- a/cpp/benchmarks/compare/comparator_benchmark.cu +++ b/cpp/benchmarks/compare/comparator_benchmark.cu @@ -81,7 +81,7 @@ static void BM_sort(benchmark::State& state, bool nulls) cuda_event_timer raii(state, true, rmm::cuda_stream_default); // auto result = cudf::sorted_order(input); - auto result = cudf::detail::sorted_order2(input); + auto result = cudf::detail::experimental::sorted_order2(input); } } diff --git a/cpp/include/cudf/sort2.cuh b/cpp/include/cudf/sort2.cuh index e612a5340d9..f9feab8246b 100644 --- a/cpp/include/cudf/sort2.cuh +++ b/cpp/include/cudf/sort2.cuh @@ -37,6 +37,7 @@ namespace cudf { namespace detail { +namespace experimental { /** * @copydoc @@ -44,51 +45,13 @@ namespace detail { * * @param stream CUDA stream used for device memory operations and kernel launches */ -template std::unique_ptr sorted_order2( table_view input, std::vector const& column_order = {}, std::vector const& null_precedence = {}, rmm::cuda_stream_view stream = rmm::cuda_stream_default, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) -{ - CUDF_FUNC_RANGE(); - if (input.num_rows() == 0 or input.num_columns() == 0) { - return cudf::make_numeric_column(data_type(type_to_id()), 0); - } - - std::unique_ptr sorted_indices = cudf::make_numeric_column( - data_type(type_to_id()), input.num_rows(), mask_state::UNALLOCATED, stream, mr); - mutable_column_view mutable_indices_view = sorted_indices->mutable_view(); - thrust::sequence(rmm::exec_policy(stream), - mutable_indices_view.begin(), - mutable_indices_view.end(), - 0); - - auto [verticalized, depths] = cudf::structs::detail::experimental::verticalize_nested_columns( - input, column_order, null_precedence); - auto device_table = table_device_view::create(verticalized.flattened_columns(), stream); - auto const d_column_order = make_device_uvector_async(verticalized.orders(), stream); - auto const d_null_precedence = make_device_uvector_async(verticalized.null_orders(), stream); - auto const d_depths = make_device_uvector_async(depths, stream); - - // auto const comparator = row_lexicographic_comparator2(*device_table, *device_table); - auto const comparator = experimental::row_lexicographic_comparator(nullate::DYNAMIC{true}, - *device_table, - *device_table, - d_depths.data(), - d_column_order.data(), - d_null_precedence.data()); - - thrust::sort(rmm::exec_policy(stream), - mutable_indices_view.begin(), - mutable_indices_view.end(), - comparator); - // protection for temporary d_column_order and d_null_precedence - stream.synchronize(); - - return sorted_indices; -} + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()); +} // namespace experimental } // namespace detail } // namespace cudf diff --git a/cpp/src/sort/sort2.cu b/cpp/src/sort/sort2.cu new file mode 100644 index 00000000000..81deb1789fa --- /dev/null +++ b/cpp/src/sort/sort2.cu @@ -0,0 +1,93 @@ +/* + * Copyright (c) 2020-2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * 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 +#include + +namespace cudf { +namespace detail { +namespace experimental { + +/** + * @copydoc + * sorted_order(table_view&,std::vector,std::vector,rmm::mr::device_memory_resource*) + * + * @param stream CUDA stream used for device memory operations and kernel launches + */ +std::unique_ptr sorted_order2(table_view input, + std::vector const& column_order, + std::vector const& null_precedence, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + CUDF_FUNC_RANGE(); + if (input.num_rows() == 0 or input.num_columns() == 0) { + return cudf::make_numeric_column(data_type(type_to_id()), 0); + } + + std::unique_ptr sorted_indices = cudf::make_numeric_column( + data_type(type_to_id()), input.num_rows(), mask_state::UNALLOCATED, stream, mr); + mutable_column_view mutable_indices_view = sorted_indices->mutable_view(); + thrust::sequence(rmm::exec_policy(stream), + mutable_indices_view.begin(), + mutable_indices_view.end(), + 0); + + auto [verticalized, depths] = cudf::structs::detail::experimental::verticalize_nested_columns( + input, column_order, null_precedence); + auto device_table = table_device_view::create(verticalized.flattened_columns(), stream); + auto const d_column_order = make_device_uvector_async(verticalized.orders(), stream); + auto const d_null_precedence = make_device_uvector_async(verticalized.null_orders(), stream); + auto const d_depths = make_device_uvector_async(depths, stream); + + // auto const comparator = row_lexicographic_comparator2(*device_table, *device_table); + auto const comparator = + cudf::experimental::row_lexicographic_comparator(nullate::DYNAMIC{true}, + *device_table, + *device_table, + d_depths.data(), + d_column_order.data(), + d_null_precedence.data()); + + thrust::sort(rmm::exec_policy(stream), + mutable_indices_view.begin(), + mutable_indices_view.end(), + comparator); + // protection for temporary d_column_order and d_null_precedence + stream.synchronize(); + + return sorted_indices; +} + +} // namespace experimental +} // namespace detail +} // namespace cudf diff --git a/cpp/tests/sort/sort2_test.cu b/cpp/tests/sort/sort2_test.cu index 56ef562be39..2c173e0a4e2 100644 --- a/cpp/tests/sort/sort2_test.cu +++ b/cpp/tests/sort/sort2_test.cu @@ -74,7 +74,7 @@ TEST_F(NewRowOpTest, BasicStructTwoChild) auto result1 = cudf::sorted_order(input); cudf::test::print(result1->view()); - auto result2 = cudf::detail::sorted_order2(input); + auto result2 = cudf::detail::experimental::sorted_order2(input); cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); } @@ -128,7 +128,7 @@ TEST_F(NewRowOpTest, DeepStruct) auto result1 = cudf::sorted_order(input); cudf::test::print(result1->view()); - auto result2 = cudf::detail::sorted_order2(input); + auto result2 = cudf::detail::experimental::sorted_order2(input); cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); } @@ -202,14 +202,14 @@ TEST_F(NewRowOpTest, SampleStructTest) auto result1 = cudf::sorted_order(input); cudf::test::print(result1->view()); - auto result2 = cudf::detail::sorted_order2(input); + auto result2 = cudf::detail::experimental::sorted_order2(input); cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); std::vector col_order = {cudf::order::DESCENDING, cudf::order::ASCENDING}; std::vector null_order = {cudf::null_order::BEFORE, cudf::null_order::AFTER}; result1 = cudf::sorted_order(input, col_order, null_order); - result2 = cudf::detail::sorted_order2(input, col_order, null_order); + result2 = cudf::detail::experimental::sorted_order2(input, col_order, null_order); cudf::test::print(result1->view()); cudf::test::print(result2->view()); cudf::test::expect_columns_equal(result1->view(), result2->view()); From c8e527e8f9fd9b12f5108cf5499e751fcb89f411 Mon Sep 17 00:00:00 2001 From: Devavret Makkar Date: Fri, 4 Feb 2022 23:40:24 +0530 Subject: [PATCH 012/333] Basic working list == comp --- cpp/include/cudf/table/row_operator_list.cuh | 291 +++++++++++++++++++ cpp/src/reductions/scan/rank_scan.cu | 6 +- cpp/tests/CMakeLists.txt | 2 + cpp/tests/reductions/list_rank_test.cpp | 51 ++++ 4 files changed, 346 insertions(+), 4 deletions(-) create mode 100644 cpp/include/cudf/table/row_operator_list.cuh create mode 100644 cpp/tests/reductions/list_rank_test.cpp diff --git a/cpp/include/cudf/table/row_operator_list.cuh b/cpp/include/cudf/table/row_operator_list.cuh new file mode 100644 index 00000000000..2f473c837c8 --- /dev/null +++ b/cpp/include/cudf/table/row_operator_list.cuh @@ -0,0 +1,291 @@ +/* + * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * 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 +#include + +#include + +namespace cudf { +namespace experimental { + +#pragma nv_exec_check_disable +template