From eaffdeac4b5e6568d7c68cd9624448895f88da27 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Mon, 16 May 2022 17:10:48 -0700 Subject: [PATCH 01/31] NaN handling in device_row_comparators --- .../cudf/table/experimental/row_operators.cuh | 87 ++++++++++++++----- cpp/include/cudf/table/row_operators.cuh | 45 ++++++++-- 2 files changed, 102 insertions(+), 30 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 43f09ff55f0..de0b7304ae4 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -87,8 +87,10 @@ namespace lexicographic { * `aac < abb`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. + * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized IEEE + * 754 compliant nan handling */ -template +template class device_row_comparator { friend class self_comparator; /** @@ -140,17 +142,20 @@ class device_row_comparator { * @param null_precedence Indicates how null values are ordered with other values * @param depth The depth of the column if part of a nested column @see * preprocessed_table::depths + * @param nan_result Specifies what value should be returned if either element is `nan` */ __device__ element_comparator(Nullate check_nulls, column_device_view lhs, column_device_view rhs, null_order null_precedence = null_order::BEFORE, - int depth = 0) + int depth = 0, + weak_ordering nan_result = weak_ordering::EQUIVALENT) : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, _null_precedence{null_precedence}, - _depth{depth} + _depth{depth}, + _nan_result{nan_result} { } @@ -176,8 +181,12 @@ class device_row_comparator { } } - return cuda::std::pair(relational_compare(_lhs.element(lhs_element_index), - _rhs.element(rhs_element_index)), + return cuda::std::pair(NanConfig + ? relational_compare(_lhs.element(lhs_element_index), + _rhs.element(rhs_element_index), + _nan_result) + : relational_compare(_lhs.element(lhs_element_index), + _rhs.element(rhs_element_index)), std::numeric_limits::max()); } @@ -216,7 +225,8 @@ class device_row_comparator { ++depth; } - auto const comparator = element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth}; + auto const comparator = + element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth, _nan_result}; return cudf::type_dispatcher( lcol.type(), comparator, lhs_element_index, rhs_element_index); } @@ -227,6 +237,7 @@ class device_row_comparator { Nullate const _check_nulls; null_order const _null_precedence; int const _depth; + weak_ordering const _nan_result; }; public: @@ -253,7 +264,18 @@ class device_row_comparator { _null_precedence.has_value() ? (*_null_precedence)[i] : null_order::BEFORE; auto const comparator = - element_comparator{_check_nulls, _lhs.column(i), _rhs.column(i), null_precedence, depth}; + NanConfig ? element_comparator{_check_nulls, + _lhs.column(i), + _rhs.column(i), + null_precedence, + depth, + ascending ? weak_ordering::GREATER : weak_ordering::LESS} + : element_comparator{_check_nulls, + _lhs.column(i), + _rhs.column(i), + null_precedence, + depth, + weak_ordering::EQUIVALENT}; weak_ordering state; cuda::std::tie(state, last_null_depth) = cudf::type_dispatcher(_lhs.column(i).type(), comparator, lhs_index, rhs_index); @@ -301,15 +323,18 @@ struct weak_ordering_comparator_impl { * weak_ordering::LESS meaning one row is lexicographically *less* than another row. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. + * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized IEEE + * 754 compliant nan handling */ -template +template using less_comparator = - weak_ordering_comparator_impl, weak_ordering::LESS>; + weak_ordering_comparator_impl, weak_ordering::LESS>; -template -using less_equivalent_comparator = weak_ordering_comparator_impl, - weak_ordering::LESS, - weak_ordering::EQUIVALENT>; +template +using less_equivalent_comparator = + weak_ordering_comparator_impl, + weak_ordering::LESS, + weak_ordering::EQUIVALENT>; struct preprocessed_table { using table_device_view_owner = @@ -457,11 +482,13 @@ class self_comparator { * `F(i,j)` returns true if and only if row `i` compares lexicographically less than row `j`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. + * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized + * IEEE 754 compliant nan handling */ - template - less_comparator device_comparator(Nullate nullate = {}) const + template + less_comparator device_comparator(Nullate nullate = {}) const { - return less_comparator{device_row_comparator( + return less_comparator{device_row_comparator( nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence())}; } @@ -494,7 +521,10 @@ class device_row_comparator { { auto equal_elements = [=](column_device_view l, column_device_view r) { return cudf::type_dispatcher( - l.type(), element_comparator{check_nulls, l, r, nulls_are_equal}, lhs_index, rhs_index); + l.type(), + element_comparator{check_nulls, l, r, nulls_are_equal, nans_are_equal}, + lhs_index, + rhs_index); }; return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements); @@ -513,8 +543,13 @@ class device_row_comparator { device_row_comparator(Nullate check_nulls, table_device_view lhs, table_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL) noexcept - : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal} + null_equality nulls_are_equal = null_equality::EQUAL, + nan_equality nans_are_equal = nan_equality::ALL_EQUAL) noexcept + : lhs{lhs}, + rhs{rhs}, + check_nulls{check_nulls}, + nulls_are_equal{nulls_are_equal}, + nans_are_equal{nans_are_equal} { } @@ -537,8 +572,13 @@ class device_row_comparator { __device__ element_comparator(Nullate check_nulls, column_device_view lhs, column_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL) noexcept - : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal} + null_equality nulls_are_equal = null_equality::EQUAL, + nan_equality nans_are_equal = nan_equality::ALL_EQUAL) noexcept + : lhs{lhs}, + rhs{rhs}, + check_nulls{check_nulls}, + nulls_are_equal{nulls_are_equal}, + nans_are_equal{nans_are_equal} { } @@ -565,7 +605,8 @@ class device_row_comparator { } return equality_compare(lhs.element(lhs_element_index), - rhs.element(rhs_element_index)); + rhs.element(rhs_element_index), + nans_are_equal); } template ::value>* = nullptr> +__device__ weak_ordering relational_compare(Element lhs, Element rhs, weak_ordering nan_result) +{ + return isnan(lhs) or isnan(rhs) ? nan_result : detail::compare_elements(lhs, rhs); +} + /** * @brief Compare the nulls according to null order. * @@ -123,11 +143,14 @@ inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_ord * * @param[in] lhs first element * @param[in] rhs second element + * @param nan_result ignored for non-floating point operation * @return Indicates the relationship between the elements in * the `lhs` and `rhs` columns. */ template >* = nullptr> -__device__ weak_ordering relational_compare(Element lhs, Element rhs) +__device__ weak_ordering relational_compare(Element lhs, + Element rhs, + weak_ordering const nan_result = weak_ordering::GREATER) { return detail::compare_elements(lhs, rhs); } @@ -138,12 +161,15 @@ __device__ weak_ordering relational_compare(Element lhs, Element rhs) * * @param lhs first element * @param rhs second element + * @param nan_result specifies what value should be returned if either element is `nan` * @return `true` if `lhs` == `rhs` else `false`. */ template >* = nullptr> -__device__ bool equality_compare(Element lhs, Element rhs) +__device__ bool equality_compare(Element lhs, + Element rhs, + nan_equality const nan_result = nan_equality::ALL_EQUAL) { - if (isnan(lhs) and isnan(rhs)) { return true; } + if (isnan(lhs) and isnan(rhs)) { return nan_result == nan_equality::ALL_EQUAL; } return lhs == rhs; } @@ -153,10 +179,13 @@ __device__ bool equality_compare(Element lhs, Element rhs) * * @param lhs first element * @param rhs second element + * @param nan_result ignored for non-floating point operation * @return `true` if `lhs` == `rhs` else `false`. */ template >* = nullptr> -__device__ bool equality_compare(Element const lhs, Element const rhs) +__device__ bool equality_compare(Element const lhs, + Element const rhs, + nan_equality const nan_result = nan_equality::ALL_EQUAL) { return lhs == rhs; } @@ -225,8 +254,8 @@ class element_equality_comparator { private: column_device_view lhs; column_device_view rhs; - Nullate nulls; - null_equality nulls_are_equal; + Nullate const nulls; + null_equality const nulls_are_equal; }; template @@ -256,8 +285,8 @@ class row_equality_comparator { private: table_device_view lhs; table_device_view rhs; - Nullate nulls; - null_equality nulls_are_equal; + Nullate const nulls; + null_equality const nulls_are_equal; }; /** From e518668fcf0f9f175b9ff7e1e1b477c46cf310b7 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Fri, 20 May 2022 14:33:22 -0700 Subject: [PATCH 02/31] template the comparator --- .../cudf/table/experimental/row_operators.cuh | 116 +++++++++++++----- 1 file changed, 84 insertions(+), 32 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index b182a03dfae..4b6e35f9d63 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -115,6 +115,51 @@ using rhs_iterator = strong_index_iterator; namespace lexicographic { +struct relational_comparator { + template >* = nullptr> + __device__ weak_ordering relational_comparator()(Element const lhs, Element const rhs) + { + return detail::compare_elements(lhs, rhs); + } + + template >* = nullptr> + __device__ weak_ordering operator()(Element const lhs, Element const 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); + } +}; + +template +struct IEE740_relational_comparator : relational_comparator { + /** + * @brief A specialization for floating-point `Element` type relational comparison + * to derive the order of the elements with respect to `lhs`. Returns specified weak_ordering if + * either value is `nan`, enabling IEEE 754 compliant comparison. + * + * This specialization allows `nan` values to be evaluated as not equal to any other value, while + * also not evaluating as greater or less than + * + * @param lhs first element + * @param rhs second element + * @param nan_result specifies what value should be returned if either element is `nan` + * @return Indicates the relationship between the elements in + * the `lhs` and `rhs` columns. + */ + template >* = nullptr> + __device__ weak_ordering operator()(Element const lhs, Element const rhs) + { + return isnan(lhs) or isnan(rhs) ? nan_result : detail::compare_elements(lhs, rhs); + } +}; + /** * @brief Computes the lexicographic comparison between 2 rows. * @@ -134,7 +179,9 @@ namespace lexicographic { * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized IEEE * 754 compliant nan handling */ -template +template class device_row_comparator { friend class self_comparator; friend class two_table_comparator; @@ -174,6 +221,7 @@ class device_row_comparator { /** * @brief Performs a relational comparison between two elements in two columns. */ + template class element_comparator { public: /** @@ -194,14 +242,12 @@ class device_row_comparator { column_device_view lhs, column_device_view rhs, null_order null_precedence = null_order::BEFORE, - int depth = 0, - weak_ordering nan_result = weak_ordering::EQUIVALENT) + int depth = 0) : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, _null_precedence{null_precedence}, - _depth{depth}, - _nan_result{nan_result} + _depth{depth} { } @@ -227,12 +273,8 @@ class device_row_comparator { } } - return cuda::std::pair(NanConfig - ? relational_compare(_lhs.element(lhs_element_index), - _rhs.element(rhs_element_index), - _nan_result) - : relational_compare(_lhs.element(lhs_element_index), - _rhs.element(rhs_element_index)), + return cuda::std::pair(ElementComparator2{}(_lhs.element(lhs_element_index), + _rhs.element(rhs_element_index)), std::numeric_limits::max()); } @@ -271,8 +313,7 @@ class device_row_comparator { ++depth; } - auto const comparator = - element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth, _nan_result}; + auto const comparator = element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth}; return cudf::type_dispatcher( lcol.type(), comparator, lhs_element_index, rhs_element_index); } @@ -283,7 +324,6 @@ class device_row_comparator { Nullate const _check_nulls; null_order const _null_precedence; int const _depth; - weak_ordering const _nan_result; }; public: @@ -310,22 +350,28 @@ class device_row_comparator { null_order const null_precedence = _null_precedence.has_value() ? (*_null_precedence)[i] : null_order::BEFORE; - auto const comparator = - NanConfig ? element_comparator{_check_nulls, - _lhs.column(i), - _rhs.column(i), - null_precedence, - depth, - ascending ? weak_ordering::GREATER : weak_ordering::LESS} - : element_comparator{_check_nulls, - _lhs.column(i), - _rhs.column(i), - null_precedence, - depth, - weak_ordering::EQUIVALENT}; weak_ordering state; cuda::std::tie(state, last_null_depth) = - cudf::type_dispatcher(_lhs.column(i).type(), comparator, lhs_index, rhs_index); + ascending ? cudf::type_dispatcher( + _lhs.column(i).type(), + element_comparator{// AscendingElementComparator{ + _check_nulls, + _lhs.column(i), + _rhs.column(i), + null_precedence, + depth}, + lhs_index, + rhs_index) + : cudf::type_dispatcher(_lhs.column(i).type(), + element_comparator{ + // DescendingElementComparator{ + _check_nulls, + _lhs.column(i), + _rhs.column(i), + null_precedence, + depth}, + lhs_index, + rhs_index); if (state == weak_ordering::EQUIVALENT) { continue; } @@ -532,11 +578,17 @@ class self_comparator { * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized * IEEE 754 compliant nan handling */ - template - less_comparator> device_comparator(Nullate nullate = {}) const + template + less_comparator< + device_row_comparator> + device_comparator(Nullate nullate = {}) const { - return less_comparator>{device_row_comparator( - nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence())}; + return less_comparator< + device_row_comparator>{ + device_row_comparator( + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence())}; } private: From c99e3c51bd88799eda7008f8e97ad27c258574cc Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Mon, 23 May 2022 16:59:26 -0700 Subject: [PATCH 03/31] partial fix to performance regression --- .../cudf/table/experimental/row_operators.cuh | 277 ++++++++---------- cpp/include/cudf/table/row_operators.cuh | 14 +- 2 files changed, 135 insertions(+), 156 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 4b6e35f9d63..ca72d7ae103 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -114,16 +114,115 @@ using lhs_iterator = strong_index_iterator; using rhs_iterator = strong_index_iterator; namespace lexicographic { +/** + * @brief Performs a relational comparison between two elements in two columns. + */ +template +struct element_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 check_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 null_precedence Indicates how null values are ordered with other values + * @param depth The depth of the column if part of a nested column @see + * preprocessed_table::depths + * @param nan_result Specifies what value should be returned if either element is `nan` + */ + __device__ element_comparator(Nullate check_nulls, + column_device_view lhs, + column_device_view rhs, + null_order null_precedence = null_order::BEFORE, + int depth = 0, + bool ascending = true) + : _lhs{lhs}, + _rhs{rhs}, + _check_nulls{check_nulls}, + _null_precedence{null_precedence}, + _depth{depth}, + _ascending{ascending} + { + } + + /** + * @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, along + * with the depth at which a null value was encountered. + */ + template ())> + __device__ cuda::std::pair operator()( + size_type const lhs_element_index, size_type const rhs_element_index) const noexcept + { + if (_check_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 cuda::std::pair(null_compare(lhs_is_null, rhs_is_null, _null_precedence), _depth); + } + } + + return cuda::std::pair(relational_comparator(_lhs.element(lhs_element_index), + _rhs.element(rhs_element_index)), + std::numeric_limits::max()); + } + + template () and + not std::is_same_v)> + __device__ cuda::std::pair operator()(size_type const, + size_type const) const noexcept + { + CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types."); + } + + template )> + __device__ cuda::std::pair operator()( + size_type const lhs_element_index, size_type const rhs_element_index) const noexcept + { + column_device_view lcol = _lhs; + column_device_view rcol = _rhs; + int depth = _depth; + while (lcol.type().id() == type_id::STRUCT) { + bool const lhs_is_null{lcol.is_null(lhs_element_index)}; + bool const rhs_is_null{rcol.is_null(rhs_element_index)}; + + if (lhs_is_null or rhs_is_null) { // at least one is null + weak_ordering state = null_compare(lhs_is_null, rhs_is_null, _null_precedence); + return cuda::std::pair(state, depth); + } + + if (lcol.num_child_columns() == 0) { + return cuda::std::pair(weak_ordering::EQUIVALENT, depth); + } + + // Non-empty structs have been modified to only have 1 child when using this. + lcol = detail::structs_column_device_view(lcol).get_sliced_child(0); + rcol = detail::structs_column_device_view(rcol).get_sliced_child(0); + ++depth; + } + + auto const comparator = element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth}; + return cudf::type_dispatcher( + lcol.type(), comparator, lhs_element_index, rhs_element_index); + } -struct relational_comparator { template >* = nullptr> - __device__ weak_ordering relational_comparator()(Element const lhs, Element const rhs) + __device__ weak_ordering relational_comparator(Element lhs, Element rhs) const noexcept { return detail::compare_elements(lhs, rhs); } template >* = nullptr> - __device__ weak_ordering operator()(Element const lhs, Element const rhs) + __device__ weak_ordering relational_comparator(Element lhs, Element rhs) const noexcept { if (isnan(lhs) and isnan(rhs)) { return weak_ordering::EQUIVALENT; @@ -135,10 +234,19 @@ struct relational_comparator { return detail::compare_elements(lhs, rhs); } + + bool const _ascending; + + private: + column_device_view const _lhs; + column_device_view const _rhs; + Nullate const _check_nulls; + null_order const _null_precedence; + int const _depth; }; -template -struct IEE740_relational_comparator : relational_comparator { +template +struct IEE740_element_comparator : element_comparator { /** * @brief A specialization for floating-point `Element` type relational comparison * to derive the order of the elements with respect to `lhs`. Returns specified weak_ordering if @@ -154,9 +262,11 @@ struct IEE740_relational_comparator : relational_comparator { * the `lhs` and `rhs` columns. */ template >* = nullptr> - __device__ weak_ordering operator()(Element const lhs, Element const rhs) + __device__ weak_ordering relational_comparator(Element lhs, Element rhs) const noexcept { - return isnan(lhs) or isnan(rhs) ? nan_result : detail::compare_elements(lhs, rhs); + return isnan(lhs) or isnan(rhs) + ? (this->_ascending ? weak_ordering::GREATER : weak_ordering::LESS) + : detail::compare_elements(lhs, rhs); } }; @@ -179,9 +289,7 @@ struct IEE740_relational_comparator : relational_comparator { * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized IEEE * 754 compliant nan handling */ -template +template > class device_row_comparator { friend class self_comparator; friend class two_table_comparator; @@ -218,114 +326,6 @@ class device_row_comparator { { } - /** - * @brief Performs a relational comparison between two elements in two columns. - */ - template - class element_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 check_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 null_precedence Indicates how null values are ordered with other values - * @param depth The depth of the column if part of a nested column @see - * preprocessed_table::depths - * @param nan_result Specifies what value should be returned if either element is `nan` - */ - __device__ element_comparator(Nullate check_nulls, - column_device_view lhs, - column_device_view rhs, - null_order null_precedence = null_order::BEFORE, - int depth = 0) - : _lhs{lhs}, - _rhs{rhs}, - _check_nulls{check_nulls}, - _null_precedence{null_precedence}, - _depth{depth} - { - } - - /** - * @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, along - * with the depth at which a null value was encountered. - */ - template ())> - __device__ cuda::std::pair operator()( - size_type const lhs_element_index, size_type const rhs_element_index) const noexcept - { - if (_check_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 cuda::std::pair(null_compare(lhs_is_null, rhs_is_null, _null_precedence), _depth); - } - } - - return cuda::std::pair(ElementComparator2{}(_lhs.element(lhs_element_index), - _rhs.element(rhs_element_index)), - std::numeric_limits::max()); - } - - template () and - not std::is_same_v)> - __device__ cuda::std::pair operator()(size_type const, - size_type const) const noexcept - { - CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types."); - } - - template )> - __device__ cuda::std::pair operator()( - size_type const lhs_element_index, size_type const rhs_element_index) const noexcept - { - column_device_view lcol = _lhs; - column_device_view rcol = _rhs; - int depth = _depth; - while (lcol.type().id() == type_id::STRUCT) { - bool const lhs_is_null{lcol.is_null(lhs_element_index)}; - bool const rhs_is_null{rcol.is_null(rhs_element_index)}; - - if (lhs_is_null or rhs_is_null) { // at least one is null - weak_ordering state = null_compare(lhs_is_null, rhs_is_null, _null_precedence); - return cuda::std::pair(state, depth); - } - - if (lcol.num_child_columns() == 0) { - return cuda::std::pair(weak_ordering::EQUIVALENT, depth); - } - - // Non-empty structs have been modified to only have 1 child when using this. - lcol = detail::structs_column_device_view(lcol).get_sliced_child(0); - rcol = detail::structs_column_device_view(rcol).get_sliced_child(0); - ++depth; - } - - auto const comparator = element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth}; - return cudf::type_dispatcher( - lcol.type(), comparator, lhs_element_index, rhs_element_index); - } - - private: - column_device_view const _lhs; - column_device_view const _rhs; - Nullate const _check_nulls; - null_order const _null_precedence; - int const _depth; - }; - public: /** * @brief Checks whether the row at `lhs_index` in the `lhs` table compares @@ -351,27 +351,12 @@ class device_row_comparator { _null_precedence.has_value() ? (*_null_precedence)[i] : null_order::BEFORE; weak_ordering state; - cuda::std::tie(state, last_null_depth) = - ascending ? cudf::type_dispatcher( - _lhs.column(i).type(), - element_comparator{// AscendingElementComparator{ - _check_nulls, - _lhs.column(i), - _rhs.column(i), - null_precedence, - depth}, - lhs_index, - rhs_index) - : cudf::type_dispatcher(_lhs.column(i).type(), - element_comparator{ - // DescendingElementComparator{ - _check_nulls, - _lhs.column(i), - _rhs.column(i), - null_precedence, - depth}, - lhs_index, - rhs_index); + cuda::std::tie(state, last_null_depth) = cudf::type_dispatcher( + _lhs.column(i).type(), + ElementComparator{ + _check_nulls, _lhs.column(i), _rhs.column(i), null_precedence, depth, ascending}, + lhs_index, + rhs_index); if (state == weak_ordering::EQUIVALENT) { continue; } @@ -578,16 +563,12 @@ class self_comparator { * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized * IEEE 754 compliant nan handling */ - template - less_comparator< - device_row_comparator> - device_comparator(Nullate nullate = {}) const + template + less_comparator>> device_comparator( + Nullate nullate = {}) const { - return less_comparator< - device_row_comparator>{ - device_row_comparator( + return less_comparator>>{ + device_row_comparator>( nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence())}; } diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh index 9fe51c88160..a4f672daf1f 100644 --- a/cpp/include/cudf/table/row_operators.cuh +++ b/cpp/include/cudf/table/row_operators.cuh @@ -111,11 +111,11 @@ __device__ weak_ordering relational_compare(Element lhs, Element rhs) * @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, weak_ordering nan_result) -{ - return isnan(lhs) or isnan(rhs) ? nan_result : detail::compare_elements(lhs, rhs); -} +// template ::value>* = nullptr> +// __device__ weak_ordering relational_compare(Element lhs, Element rhs, weak_ordering nan_result) +// { +// return isnan(lhs) or isnan(rhs) ? nan_result : detail::compare_elements(lhs, rhs); +// } /** * @brief Compare the nulls according to null order. @@ -148,9 +148,7 @@ inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_ord * the `lhs` and `rhs` columns. */ template >* = nullptr> -__device__ weak_ordering relational_compare(Element lhs, - Element rhs, - weak_ordering const nan_result = weak_ordering::GREATER) +__device__ weak_ordering relational_compare(Element lhs, Element rhs) { return detail::compare_elements(lhs, rhs); } From 798d6c5a47da063c0e99bc088edee7fba9a15fdc Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Tue, 31 May 2022 16:21:32 -0700 Subject: [PATCH 04/31] Template NaN config lexicographic and equality --- .../cudf/table/experimental/row_operators.cuh | 444 ++++++++++-------- cpp/include/cudf/table/row_operators.cuh | 41 +- 2 files changed, 268 insertions(+), 217 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index f9c515c37ac..157f72e9c84 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -134,115 +134,57 @@ using lhs_iterator = strong_index_iterator; using rhs_iterator = strong_index_iterator; namespace lexicographic { + /** - * @brief Performs a relational comparison between two elements in two columns. + * @brief Relational comparator functor that evaluates `NaN` as not less, not equivalent, and not greater + * than all other values. This is IEE 754 compliant. */ -template -struct element_comparator { - public: +struct physical_element_comparator { /** - * @brief Construct type-dispatched function object for performing a - * relational comparison between two elements. - * - * @note `lhs` and `rhs` may be the same. + * @brief Shared operator for all `Element` types relational comparisons. * - * @param check_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 null_precedence Indicates how null values are ordered with other values - * @param depth The depth of the column if part of a nested column @see - * preprocessed_table::depths - * @param nan_result Specifies what value should be returned if either element is `nan` + * @param[in] lhs first element + * @param[in] rhs second element + * @return Indicates the relationship between the elements in + * the `lhs` and `rhs` columns. */ - __device__ element_comparator(Nullate check_nulls, - column_device_view lhs, - column_device_view rhs, - null_order null_precedence = null_order::BEFORE, - int depth = 0, - bool ascending = true) - : _lhs{lhs}, - _rhs{rhs}, - _check_nulls{check_nulls}, - _null_precedence{null_precedence}, - _depth{depth}, - _ascending{ascending} + template + __device__ weak_ordering operator()(Element lhs, Element rhs) { + return detail::compare_elements(lhs, rhs); } +}; +/** + * @brief Relational comparator functor that evaluates `NaN` as equivalent to other `NaN`s and + * greater than all other values. + */ +struct sorting_physical_element_comparator : physical_element_comparator { /** - * @brief Performs a relational comparison between the specified elements + * @brief A specialization for non-floating-point `Element` type relational + * comparison to derive the order of the elements with respect to `lhs`. * - * @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, along - * with the depth at which a null value was encountered. + * @param[in] lhs first element + * @param[in] rhs second element + * @return Indicates the relationship between the elements in + * the `lhs` and `rhs` columns. */ - template ())> - __device__ cuda::std::pair operator()( - size_type const lhs_element_index, size_type const rhs_element_index) const noexcept - { - if (_check_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 cuda::std::pair(null_compare(lhs_is_null, rhs_is_null, _null_precedence), _depth); - } - } - - return cuda::std::pair(relational_comparator(_lhs.element(lhs_element_index), - _rhs.element(rhs_element_index)), - std::numeric_limits::max()); - } - - template () and - not std::is_same_v)> - __device__ cuda::std::pair operator()(size_type const, - size_type const) const noexcept - { - CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types."); - } - - template )> - __device__ cuda::std::pair operator()( - size_type const lhs_element_index, size_type const rhs_element_index) const noexcept - { - column_device_view lcol = _lhs; - column_device_view rcol = _rhs; - int depth = _depth; - while (lcol.type().id() == type_id::STRUCT) { - bool const lhs_is_null{lcol.is_null(lhs_element_index)}; - bool const rhs_is_null{rcol.is_null(rhs_element_index)}; - - if (lhs_is_null or rhs_is_null) { // at least one is null - weak_ordering state = null_compare(lhs_is_null, rhs_is_null, _null_precedence); - return cuda::std::pair(state, depth); - } - - if (lcol.num_child_columns() == 0) { - return cuda::std::pair(weak_ordering::EQUIVALENT, depth); - } - - // Non-empty structs have been modified to only have 1 child when using this. - lcol = detail::structs_column_device_view(lcol).get_sliced_child(0); - rcol = detail::structs_column_device_view(rcol).get_sliced_child(0); - ++depth; - } - - auto const comparator = element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth}; - return cudf::type_dispatcher( - lcol.type(), comparator, lhs_element_index, rhs_element_index); - } - template >* = nullptr> - __device__ weak_ordering relational_comparator(Element lhs, Element rhs) const noexcept + __device__ weak_ordering operator()(Element lhs, Element rhs) { return detail::compare_elements(lhs, rhs); } + /** + * @brief A specialization for floating-point `Element` type relational comparison that compares + * 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 + * @return Indicates the relationship between null in lhs and rhs columns. + */ template >* = nullptr> - __device__ weak_ordering relational_comparator(Element lhs, Element rhs) const noexcept + __device__ weak_ordering operator()(Element lhs, Element rhs) { if (isnan(lhs) and isnan(rhs)) { return weak_ordering::EQUIVALENT; @@ -254,40 +196,6 @@ struct element_comparator { return detail::compare_elements(lhs, rhs); } - - bool const _ascending; - - private: - column_device_view const _lhs; - column_device_view const _rhs; - Nullate const _check_nulls; - null_order const _null_precedence; - int const _depth; -}; - -template -struct IEE740_element_comparator : element_comparator { - /** - * @brief A specialization for floating-point `Element` type relational comparison - * to derive the order of the elements with respect to `lhs`. Returns specified weak_ordering if - * either value is `nan`, enabling IEEE 754 compliant comparison. - * - * This specialization allows `nan` values to be evaluated as not equal to any other value, while - * also not evaluating as greater or less than - * - * @param lhs first element - * @param rhs second element - * @param nan_result specifies what value should be returned if either element is `nan` - * @return Indicates the relationship between the elements in - * the `lhs` and `rhs` columns. - */ - template >* = nullptr> - __device__ weak_ordering relational_comparator(Element lhs, Element rhs) const noexcept - { - return isnan(lhs) or isnan(rhs) - ? (this->_ascending ? weak_ordering::GREATER : weak_ordering::LESS) - : detail::compare_elements(lhs, rhs); - } }; /** @@ -306,10 +214,11 @@ struct IEE740_element_comparator : element_comparator { * `aac < abb`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized IEEE - * 754 compliant nan handling + * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware + * relational comparator that evaluates `NaN` as greater than standard values. */ -template > +template class device_row_comparator { friend class self_comparator; friend class two_table_comparator; @@ -346,6 +255,112 @@ class device_row_comparator { { } + /** + * @brief Performs a relational comparison between two elements in two columns. + */ + class element_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 check_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 null_precedence Indicates how null values are ordered with other values + * @param depth The depth of the column if part of a nested column @see + * preprocessed_table::depths + */ + __device__ element_comparator(Nullate check_nulls, + column_device_view lhs, + column_device_view rhs, + null_order null_precedence = null_order::BEFORE, + int depth = 0) + : _lhs{lhs}, + _rhs{rhs}, + _check_nulls{check_nulls}, + _null_precedence{null_precedence}, + _depth{depth} + { + } + + /** + * @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, along + * with the depth at which a null value was encountered. + */ + template ())> + __device__ cuda::std::pair operator()( + size_type const lhs_element_index, size_type const rhs_element_index) const noexcept + { + if (_check_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 cuda::std::pair(null_compare(lhs_is_null, rhs_is_null, _null_precedence), _depth); + } + } + + return cuda::std::pair(PhysicalElementComparator{}(_lhs.element(lhs_element_index), + _rhs.element(rhs_element_index)), + std::numeric_limits::max()); + } + + template () and + not std::is_same_v)> + __device__ cuda::std::pair operator()(size_type const, + size_type const) const noexcept + { + CUDF_UNREACHABLE("Attempted to compare elements of uncomparable types."); + } + + template )> + __device__ cuda::std::pair operator()( + size_type const lhs_element_index, size_type const rhs_element_index) const noexcept + { + column_device_view lcol = _lhs; + column_device_view rcol = _rhs; + int depth = _depth; + while (lcol.type().id() == type_id::STRUCT) { + bool const lhs_is_null{lcol.is_null(lhs_element_index)}; + bool const rhs_is_null{rcol.is_null(rhs_element_index)}; + + if (lhs_is_null or rhs_is_null) { // at least one is null + weak_ordering state = null_compare(lhs_is_null, rhs_is_null, _null_precedence); + return cuda::std::pair(state, depth); + } + + if (lcol.num_child_columns() == 0) { + return cuda::std::pair(weak_ordering::EQUIVALENT, depth); + } + + // Non-empty structs have been modified to only have 1 child when using this. + lcol = detail::structs_column_device_view(lcol).get_sliced_child(0); + rcol = detail::structs_column_device_view(rcol).get_sliced_child(0); + ++depth; + } + + auto const comparator = element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth}; + return cudf::type_dispatcher( + lcol.type(), comparator, lhs_element_index, rhs_element_index); + } + + private: + column_device_view const _lhs; + column_device_view const _rhs; + Nullate const _check_nulls; + null_order const _null_precedence; + int const _depth; + }; + public: /** * @brief Checks whether the row at `lhs_index` in the `lhs` table compares @@ -370,13 +385,11 @@ class device_row_comparator { null_order const null_precedence = _null_precedence.has_value() ? (*_null_precedence)[i] : null_order::BEFORE; + auto const comparator = + element_comparator{_check_nulls, _lhs.column(i), _rhs.column(i), null_precedence, depth}; weak_ordering state; - cuda::std::tie(state, last_null_depth) = cudf::type_dispatcher( - _lhs.column(i).type(), - ElementComparator{ - _check_nulls, _lhs.column(i), _rhs.column(i), null_precedence, depth, ascending}, - lhs_index, - rhs_index); + cuda::std::tie(state, last_null_depth) = + cudf::type_dispatcher(_lhs.column(i).type(), comparator, lhs_index, rhs_index); if (state == weak_ordering::EQUIVALENT) { continue; } @@ -423,8 +436,6 @@ struct weak_ordering_comparator_impl { * weak_ordering::LESS meaning one row is lexicographically *less* than another row. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized IEEE - * 754 compliant nan handling */ template using less_comparator = weak_ordering_comparator_impl; @@ -580,15 +591,16 @@ class self_comparator { * `F(i,j)` returns true if and only if row `i` compares lexicographically less than row `j`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam NanConfig default configuration nans are equal, if set to true triggers specialized - * IEEE 754 compliant nan handling + * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware + * relational comparator that evaluates `NaN` as greater than standard values. */ - template - less_comparator>> device_comparator( + template + less_comparator> device_comparator( Nullate nullate = {}) const { - return less_comparator>>{ - device_row_comparator>( + return less_comparator>{ + device_row_comparator( nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence())}; } @@ -695,18 +707,23 @@ class two_table_comparator { * `j` of the left table. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. + * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware + * relational comparator that evaluates `NaN` as greater than standard values. */ - template - less_comparator>> + template + less_comparator< + strong_index_comparator_adapter>> device_comparator(Nullate nullate = {}) const { - return less_comparator>>{ - device_row_comparator(nullate, - *d_left_table, - *d_right_table, - d_left_table->depths(), - d_left_table->column_order(), - d_left_table->null_precedence())}; + return less_comparator< + strong_index_comparator_adapter>>{ + device_row_comparator(nullate, + *d_left_table, + *d_right_table, + d_left_table->depths(), + d_left_table->column_order(), + d_left_table->null_precedence())}; } private: @@ -722,7 +739,75 @@ class row_hasher; namespace equality { -template +/** + * @brief Equality comparator functor that evaluates `NaN` not equal to all other values for IEE 754 + * compliance. + */ +struct physical_equality_comparator { + /** + * @brief Shared operator for all `Element` types 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 + __device__ bool operator()(Element const lhs, Element const rhs) + { + return lhs == rhs; + } +}; + +/** + * @brief Equality comparator functor that evaluates `NaN` as equivalent to other `NaN`s . + */ +struct sorting_physical_equality_comparator : physical_equality_comparator { + /** + * @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 >* = nullptr> + __device__ bool operator()(Element const lhs, Element const rhs) + { + return 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 >* = nullptr> + __device__ bool operator()(Element lhs, + Element rhs, + nan_equality const nan_result = nan_equality::ALL_EQUAL) + { + return isnan(lhs) and isnan(rhs) ? true : lhs == rhs; + } +}; + +/** + * @brief Computes the equality comparison between 2 rows. + * + * Equality is determined by: + * - Two rows are compared element by element. + * - The first mismatching element returns false representing unequal rows + * - If the rows are compared without mismatched elements, the rows are equivalent + * + * @tparam Nullate A cudf::nullate type describing whether to check for nulls. + * @tparam PhysicalElementComparator A equality comparator functor, defaults to a `NaN == NaN` + * equality comparator. + */ +template class device_row_comparator { friend class self_comparator; friend class two_table_comparator; @@ -740,10 +825,7 @@ class device_row_comparator { { auto equal_elements = [=](column_device_view l, column_device_view r) { return cudf::type_dispatcher( - l.type(), - element_comparator{check_nulls, l, r, nulls_are_equal, nans_are_equal}, - lhs_index, - rhs_index); + l.type(), element_comparator{check_nulls, l, r, nulls_are_equal}, lhs_index, rhs_index); }; return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements); @@ -762,13 +844,8 @@ class device_row_comparator { device_row_comparator(Nullate check_nulls, table_device_view lhs, table_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL, - nan_equality nans_are_equal = nan_equality::ALL_EQUAL) noexcept - : lhs{lhs}, - rhs{rhs}, - check_nulls{check_nulls}, - nulls_are_equal{nulls_are_equal}, - nans_are_equal{nans_are_equal} + null_equality nulls_are_equal = null_equality::EQUAL) noexcept + : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal} { } @@ -791,13 +868,8 @@ class device_row_comparator { __device__ element_comparator(Nullate check_nulls, column_device_view lhs, column_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL, - nan_equality nans_are_equal = nan_equality::ALL_EQUAL) noexcept - : lhs{lhs}, - rhs{rhs}, - check_nulls{check_nulls}, - nulls_are_equal{nulls_are_equal}, - nans_are_equal{nans_are_equal} + null_equality nulls_are_equal = null_equality::EQUAL) noexcept + : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal} { } @@ -823,9 +895,8 @@ class device_row_comparator { } } - return equality_compare(lhs.element(lhs_element_index), - rhs.element(rhs_element_index), - nans_are_equal); + return PhysicalEqualityComparator{}(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); } template - device_row_comparator device_comparator( + template + device_row_comparator device_comparator( Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL) const { - return device_row_comparator(nullate, *d_t, *d_t, nulls_are_equal); + return device_row_comparator( + nullate, *d_t, *d_t, nulls_are_equal); } private: @@ -1098,13 +1171,18 @@ class two_table_comparator { * right table compares equal to row `j` of the left table. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. + * @tparam PhysicalElementComparator A equality comparator functor, defaults to a `NaN == NaN` + * equality comparator. */ - template + template auto device_comparator(Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL) const { - return strong_index_comparator_adapter>{ - device_row_comparator(nullate, *d_left_table, *d_right_table, nulls_are_equal)}; + return strong_index_comparator_adapter< + device_row_comparator>{ + device_row_comparator( + nullate, *d_left_table, *d_right_table, nulls_are_equal)}; } private: diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh index a4f672daf1f..a181e9bae63 100644 --- a/cpp/include/cudf/table/row_operators.cuh +++ b/cpp/include/cudf/table/row_operators.cuh @@ -97,26 +97,6 @@ __device__ weak_ordering relational_compare(Element lhs, Element rhs) return detail::compare_elements(lhs, rhs); } -/** - * @brief A specialization for floating-point `Element` type relational comparison - * to derive the order of the elements with respect to `lhs`. Returns specified weak_ordering if - * either value is `nan`, enabling IEEE 754 compliant comparison. - * - * This specialization allows `nan` values to be evaluated as not equal to any other value, while - * also not evaluating as greater or less than - * - * @param lhs first element - * @param rhs second element - * @param nan_result specifies what value should be returned if either element is `nan` - * @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, weak_ordering nan_result) -// { -// return isnan(lhs) or isnan(rhs) ? nan_result : detail::compare_elements(lhs, rhs); -// } - /** * @brief Compare the nulls according to null order. * @@ -143,7 +123,6 @@ inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_ord * * @param[in] lhs first element * @param[in] rhs second element - * @param nan_result ignored for non-floating point operation * @return Indicates the relationship between the elements in * the `lhs` and `rhs` columns. */ @@ -159,15 +138,12 @@ __device__ weak_ordering relational_compare(Element lhs, Element rhs) * * @param lhs first element * @param rhs second element - * @param nan_result specifies what value should be returned if either element is `nan` * @return `true` if `lhs` == `rhs` else `false`. */ template >* = nullptr> -__device__ bool equality_compare(Element lhs, - Element rhs, - nan_equality const nan_result = nan_equality::ALL_EQUAL) +__device__ bool equality_compare(Element lhs, Element rhs) { - if (isnan(lhs) and isnan(rhs)) { return nan_result == nan_equality::ALL_EQUAL; } + if (isnan(lhs) and isnan(rhs)) { return true; } return lhs == rhs; } @@ -177,13 +153,10 @@ __device__ bool equality_compare(Element lhs, * * @param lhs first element * @param rhs second element - * @param nan_result ignored for non-floating point operation * @return `true` if `lhs` == `rhs` else `false`. */ template >* = nullptr> -__device__ bool equality_compare(Element const lhs, - Element const rhs, - nan_equality const nan_result = nan_equality::ALL_EQUAL) +__device__ bool equality_compare(Element const lhs, Element const rhs) { return lhs == rhs; } @@ -252,8 +225,8 @@ class element_equality_comparator { private: column_device_view lhs; column_device_view rhs; - Nullate const nulls; - null_equality const nulls_are_equal; + Nullate nulls; + null_equality nulls_are_equal; }; template @@ -283,8 +256,8 @@ class row_equality_comparator { private: table_device_view lhs; table_device_view rhs; - Nullate const nulls; - null_equality const nulls_are_equal; + Nullate nulls; + null_equality nulls_are_equal; }; /** From c5f996140f651f7a9f9cd961239a2c8003c05ac4 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Tue, 31 May 2022 16:45:25 -0700 Subject: [PATCH 05/31] Add experimental row operator tests --- .../cudf/table/experimental/row_operators.cuh | 4 +- cpp/tests/CMakeLists.txt | 6 +- .../table/experimental_table_view_tests.cu | 214 ++++++++++++++++++ 3 files changed, 221 insertions(+), 3 deletions(-) create mode 100644 cpp/tests/table/experimental_table_view_tests.cu diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 157f72e9c84..a1621661060 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -136,8 +136,8 @@ using rhs_iterator = strong_index_iterator; namespace lexicographic { /** - * @brief Relational comparator functor that evaluates `NaN` as not less, not equivalent, and not greater - * than all other values. This is IEE 754 compliant. + * @brief Relational comparator functor that evaluates `NaN` as not less, not equivalent, and not + * greater than all other values. This is IEE 754 compliant. */ struct physical_element_comparator { /** diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index eadcd985de3..9273ad53097 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -309,7 +309,11 @@ ConfigureTest(TRANSPOSE_TEST transpose/transpose_test.cpp) # ################################################################################################## # * table tests ----------------------------------------------------------------------------------- ConfigureTest( - TABLE_TEST table/table_tests.cpp table/table_view_tests.cu table/row_operators_tests.cpp + TABLE_TEST + table/table_tests.cpp + table/table_view_tests.cu + table/experimental_table_view_tests.cu + table/row_operators_tests.cpp ) # ################################################################################################## diff --git a/cpp/tests/table/experimental_table_view_tests.cu b/cpp/tests/table/experimental_table_view_tests.cu new file mode 100644 index 00000000000..5603995bad6 --- /dev/null +++ b/cpp/tests/table/experimental_table_view_tests.cu @@ -0,0 +1,214 @@ +/* + * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include +#include + +#include +#include + +#include + +using namespace cudf::test; +using namespace cudf::experimental::row::lexicographic; + +// Compares two tables row by row, if table1 row is less than table2, then corresponding row value +// in `output` would be `true`/1 else `false`/0. +template +struct TypedTableViewTest : public cudf::test::BaseFixture { +}; + +using NumericTypesNotBool = Concat; +TYPED_TEST_SUITE(TypedTableViewTest, NumericTypesNotBool); + +template +void row_comparison(cudf::table_view input1, + cudf::table_view input2, + cudf::mutable_column_view output, + std::vector const& column_order) +{ + rmm::cuda_stream_view stream{}; + + auto table_comparator = two_table_comparator{input1, input2, column_order, {}, stream}; + auto comparator = + table_comparator.device_comparator(cudf::nullate::NO{}); + auto const lhs_it = cudf::experimental::row::lhs_iterator(0); + auto const rhs_it = cudf::experimental::row::rhs_iterator(0); + thrust::transform(rmm::exec_policy(stream), + lhs_it, + lhs_it + input1.num_rows(), + rhs_it, + output.data(), + comparator); +} + +template +void self_comparison(cudf::table_view input, + cudf::mutable_column_view output, + std::vector const& column_order) +{ + rmm::cuda_stream_view stream{}; + + auto table_comparator = self_comparator{input, column_order, {}, stream}; + auto comparator = + table_comparator.device_comparator(cudf::nullate::NO{}); + thrust::transform(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(input.num_rows()), + thrust::make_counting_iterator(0), + output.data(), + comparator); +} + +TYPED_TEST(TypedTableViewTest, EmptyColumnedTable) +{ + std::vector cols{}; + + cudf::table_view input(cols); + cudf::size_type expected = 0; + + EXPECT_EQ(input.num_columns(), expected); +} + +TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTableCase) +{ + using T = TypeParam; + + fixed_width_column_wrapper col1{{1, 2, 3, 4}}; + fixed_width_column_wrapper col2{{0, 1, 4, 3}}; + std::vector column_order{cudf::order::DESCENDING}; + + cudf::table_view input_table_1{{col1}}; + cudf::table_view input_table_2{{col2}}; + + auto got = cudf::make_numeric_column( + cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); + fixed_width_column_wrapper expected{{1, 1, 0, 1}}; + + row_comparison( + input_table_1, input_table_2, got->mutable_view(), column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + + row_comparison( + input_table_1, input_table_2, got->mutable_view(), column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); +} + +TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorSameTable) +{ + using T = TypeParam; + + fixed_width_column_wrapper col1{{1, 2, 3, 4}}; + std::vector column_order{cudf::order::DESCENDING}; + + cudf::table_view input_table_1{{col1}}; + + auto got = cudf::make_numeric_column( + cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); + fixed_width_column_wrapper expected{{0, 0, 0, 0}}; + + self_comparison(input_table_1, got->mutable_view(), column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + + self_comparison( + input_table_1, got->mutable_view(), column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); +} + +TYPED_TEST(TypedTableViewTest, Select) +{ + using T = TypeParam; + + fixed_width_column_wrapper col1{{1, 2, 3, 4}}; + fixed_width_column_wrapper col2{{1, 2, 3, 4}}; + fixed_width_column_wrapper col3{{4, 5, 6, 7}}; + fixed_width_column_wrapper col4{{4, 5, 6, 7}}; + cudf::table_view t{{col1, col2, col3, col4}}; + + cudf::table_view selected = t.select({2, 3}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(t.column(2), selected.column(0)); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(t.column(3), selected.column(1)); +} + +TYPED_TEST(TypedTableViewTest, SelectOutOfBounds) +{ + using T = TypeParam; + + fixed_width_column_wrapper col1{{1, 2, 3, 4}}; + fixed_width_column_wrapper col2{{1, 2, 3, 4}}; + fixed_width_column_wrapper col3{{4, 5, 6, 7}}; + fixed_width_column_wrapper col4{{4, 5, 6, 7}}; + cudf::table_view t{{col1, col2}}; + + EXPECT_THROW((void)t.select({2, 3, 4}), std::out_of_range); +} + +TYPED_TEST(TypedTableViewTest, SelectNoColumns) +{ + using T = TypeParam; + + fixed_width_column_wrapper col1{{1, 2, 3, 4}}; + fixed_width_column_wrapper col2{{1, 2, 3, 4}}; + fixed_width_column_wrapper col3{{4, 5, 6, 7}}; + fixed_width_column_wrapper col4{{4, 5, 6, 7}}; + cudf::table_view t{{col1, col2, col3, col4}}; + + cudf::table_view selected = t.select({}); + EXPECT_EQ(selected.num_columns(), 0); +} + +template +struct NaNTableViewTest : public cudf::test::BaseFixture { +}; + +TYPED_TEST_SUITE(NaNTableViewTest, FloatingPointTypes); + +TYPED_TEST(NaNTableViewTest, TestLexicographicalComparatorTwoTableNaNCase) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper col1{{T(NAN), T(NAN), T(1), T(1)}}; + cudf::test::fixed_width_column_wrapper col2{{T(NAN), T(1), T(NAN), T(1)}}; + std::vector column_order{cudf::order::DESCENDING}; + + cudf::table_view input_table_1{{col1}}; + cudf::table_view input_table_2{{col2}}; + + auto got = cudf::make_numeric_column( + cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); + + cudf::test::fixed_width_column_wrapper expected{{0, 0, 0, 0}}; + row_comparison( + input_table_1, input_table_2, got->mutable_view(), column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + + cudf::test::fixed_width_column_wrapper sorting_expected{{0, 1, 0, 0}}; + row_comparison( + input_table_1, input_table_2, got->mutable_view(), column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, got->view()); +} From 0d4e7988d1886eee9e5e4dc4fc78b84117395c1a Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Tue, 31 May 2022 17:06:05 -0700 Subject: [PATCH 06/31] switch to CUDF_ENABLE_IF --- cpp/include/cudf/table/experimental/row_operators.cuh | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index a1621661060..12ee0ef5462 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -169,7 +169,7 @@ struct sorting_physical_element_comparator : physical_element_comparator { * @return Indicates the relationship between the elements in * the `lhs` and `rhs` columns. */ - template >* = nullptr> + template )> __device__ weak_ordering operator()(Element lhs, Element rhs) { return detail::compare_elements(lhs, rhs); @@ -183,7 +183,7 @@ struct sorting_physical_element_comparator : physical_element_comparator { * @param rhs_is_null boolean representing if lhs is null * @return Indicates the relationship between null in lhs and rhs columns. */ - template >* = nullptr> + template )> __device__ weak_ordering operator()(Element lhs, Element rhs) { if (isnan(lhs) and isnan(rhs)) { @@ -771,7 +771,7 @@ struct sorting_physical_equality_comparator : physical_equality_comparator { * @param rhs second element * @return `true` if `lhs` == `rhs` else `false`. */ - template >* = nullptr> + template )> __device__ bool operator()(Element const lhs, Element const rhs) { return lhs == rhs; @@ -785,7 +785,7 @@ struct sorting_physical_equality_comparator : physical_equality_comparator { * @param rhs second element * @return `true` if `lhs` == `rhs` else `false`. */ - template >* = nullptr> + template )> __device__ bool operator()(Element lhs, Element rhs, nan_equality const nan_result = nan_equality::ALL_EQUAL) From f35d9e39377029bcdfe91adb96b57761b325a453 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Wed, 1 Jun 2022 01:35:53 -0700 Subject: [PATCH 07/31] Naming and add equality tests --- .../cudf/table/experimental/row_operators.cuh | 39 +++++---- .../table/experimental_table_view_tests.cu | 84 +++++++++++++++++-- 2 files changed, 94 insertions(+), 29 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 12ee0ef5462..06b8363da26 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -149,7 +149,7 @@ struct physical_element_comparator { * the `lhs` and `rhs` columns. */ template - __device__ weak_ordering operator()(Element lhs, Element rhs) + __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept { return detail::compare_elements(lhs, rhs); } @@ -159,7 +159,7 @@ struct physical_element_comparator { * @brief Relational comparator functor that evaluates `NaN` as equivalent to other `NaN`s and * greater than all other values. */ -struct sorting_physical_element_comparator : physical_element_comparator { +struct sorting_physical_element_comparator { /** * @brief A specialization for non-floating-point `Element` type relational * comparison to derive the order of the elements with respect to `lhs`. @@ -170,7 +170,7 @@ struct sorting_physical_element_comparator : physical_element_comparator { * the `lhs` and `rhs` columns. */ template )> - __device__ weak_ordering operator()(Element lhs, Element rhs) + __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept { return detail::compare_elements(lhs, rhs); } @@ -184,7 +184,7 @@ struct sorting_physical_element_comparator : physical_element_comparator { * @return Indicates the relationship between null in lhs and rhs columns. */ template )> - __device__ weak_ordering operator()(Element lhs, Element rhs) + __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept { if (isnan(lhs) and isnan(rhs)) { return weak_ordering::EQUIVALENT; @@ -371,8 +371,8 @@ class device_row_comparator { * @return weak ordering comparison of the row in the `lhs` table relative to the row in the `rhs` * table */ - __device__ weak_ordering operator()(size_type const lhs_index, - size_type const rhs_index) const noexcept + __device__ constexpr weak_ordering operator()(size_type const lhs_index, + size_type const rhs_index) const noexcept { int last_null_depth = std::numeric_limits::max(); for (size_type i = 0; i < _lhs.num_columns(); ++i) { @@ -597,7 +597,7 @@ class self_comparator { template less_comparator> device_comparator( - Nullate nullate = {}) const + Nullate nullate = {}) const noexcept { return less_comparator>{ device_row_comparator( @@ -714,7 +714,7 @@ class two_table_comparator { typename PhysicalElementComparator = sorting_physical_element_comparator> less_comparator< strong_index_comparator_adapter>> - device_comparator(Nullate nullate = {}) const + device_comparator(Nullate nullate = {}) const noexcept { return less_comparator< strong_index_comparator_adapter>>{ @@ -753,7 +753,7 @@ struct physical_equality_comparator { * @return `true` if `lhs` == `rhs` else `false`. */ template - __device__ bool operator()(Element const lhs, Element const rhs) + __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept { return lhs == rhs; } @@ -762,7 +762,7 @@ struct physical_equality_comparator { /** * @brief Equality comparator functor that evaluates `NaN` as equivalent to other `NaN`s . */ -struct sorting_physical_equality_comparator : physical_equality_comparator { +struct nan_equal_physical_equality_comparator { /** * @brief A specialization for non-floating-point `Element` type to check if * `lhs` is equivalent to `rhs`. @@ -772,7 +772,7 @@ struct sorting_physical_equality_comparator : physical_equality_comparator { * @return `true` if `lhs` == `rhs` else `false`. */ template )> - __device__ bool operator()(Element const lhs, Element const rhs) + __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept { return lhs == rhs; } @@ -786,9 +786,7 @@ struct sorting_physical_equality_comparator : physical_equality_comparator { * @return `true` if `lhs` == `rhs` else `false`. */ template )> - __device__ bool operator()(Element lhs, - Element rhs, - nan_equality const nan_result = nan_equality::ALL_EQUAL) + __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept { return isnan(lhs) and isnan(rhs) ? true : lhs == rhs; } @@ -807,7 +805,7 @@ struct sorting_physical_equality_comparator : physical_equality_comparator { * equality comparator. */ template + typename PhysicalEqualityComparator = nan_equal_physical_equality_comparator> class device_row_comparator { friend class self_comparator; friend class two_table_comparator; @@ -821,7 +819,8 @@ class device_row_comparator { * @param rhs_index The index of the row in the `rhs` table to examine * @return `true` if row from the `lhs` table is equal to the row in the `rhs` table */ - __device__ bool operator()(size_type const lhs_index, size_type const rhs_index) const noexcept + __device__ constexpr bool operator()(size_type const lhs_index, + size_type const rhs_index) const noexcept { auto equal_elements = [=](column_device_view l, column_device_view r) { return cudf::type_dispatcher( @@ -1081,9 +1080,9 @@ class self_comparator { * equality comparator. */ template + typename PhysicalEqualityComparator = nan_equal_physical_equality_comparator> device_row_comparator device_comparator( - Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL) const + Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL) const noexcept { return device_row_comparator( nullate, *d_t, *d_t, nulls_are_equal); @@ -1175,9 +1174,9 @@ class two_table_comparator { * equality comparator. */ template + typename PhysicalEqualityComparator = nan_equal_physical_equality_comparator> auto device_comparator(Nullate nullate = {}, - null_equality nulls_are_equal = null_equality::EQUAL) const + null_equality nulls_are_equal = null_equality::EQUAL) const noexcept { return strong_index_comparator_adapter< device_row_comparator>{ diff --git a/cpp/tests/table/experimental_table_view_tests.cu b/cpp/tests/table/experimental_table_view_tests.cu index 5603995bad6..d847a8272be 100644 --- a/cpp/tests/table/experimental_table_view_tests.cu +++ b/cpp/tests/table/experimental_table_view_tests.cu @@ -35,7 +35,7 @@ #include using namespace cudf::test; -using namespace cudf::experimental::row::lexicographic; +using namespace cudf::experimental::row; // Compares two tables row by row, if table1 row is less than table2, then corresponding row value // in `output` would be `true`/1 else `false`/0. @@ -54,7 +54,8 @@ void row_comparison(cudf::table_view input1, { rmm::cuda_stream_view stream{}; - auto table_comparator = two_table_comparator{input1, input2, column_order, {}, stream}; + auto table_comparator = + lexicographic::two_table_comparator{input1, input2, column_order, {}, stream}; auto comparator = table_comparator.device_comparator(cudf::nullate::NO{}); auto const lhs_it = cudf::experimental::row::lhs_iterator(0); @@ -74,7 +75,46 @@ void self_comparison(cudf::table_view input, { rmm::cuda_stream_view stream{}; - auto table_comparator = self_comparator{input, column_order, {}, stream}; + auto table_comparator = lexicographic::self_comparator{input, column_order, {}, stream}; + auto comparator = + table_comparator.device_comparator(cudf::nullate::NO{}); + thrust::transform(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(input.num_rows()), + thrust::make_counting_iterator(0), + output.data(), + comparator); +} + +template +void row_equality(cudf::table_view input1, + cudf::table_view input2, + cudf::mutable_column_view output, + std::vector const& column_order) +{ + rmm::cuda_stream_view stream{}; + + auto table_comparator = equality::two_table_comparator{input1, input2, stream}; + auto comparator = + table_comparator.device_comparator(cudf::nullate::NO{}); + auto const lhs_it = cudf::experimental::row::lhs_iterator(0); + auto const rhs_it = cudf::experimental::row::rhs_iterator(0); + thrust::transform(rmm::exec_policy(stream), + lhs_it, + lhs_it + input1.num_rows(), + rhs_it, + output.data(), + comparator); +} + +template +void self_equality(cudf::table_view input, + cudf::mutable_column_view output, + std::vector const& column_order) +{ + rmm::cuda_stream_view stream{}; + + auto table_comparator = equality::self_comparator{input, stream}; auto comparator = table_comparator.device_comparator(cudf::nullate::NO{}); thrust::transform(rmm::exec_policy(stream), @@ -110,11 +150,11 @@ TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTableCase) cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); fixed_width_column_wrapper expected{{1, 1, 0, 1}}; - row_comparison( + row_comparison( input_table_1, input_table_2, got->mutable_view(), column_order); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - row_comparison( + row_comparison( input_table_1, input_table_2, got->mutable_view(), column_order); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); } @@ -132,10 +172,11 @@ TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorSameTable) cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); fixed_width_column_wrapper expected{{0, 0, 0, 0}}; - self_comparison(input_table_1, got->mutable_view(), column_order); + self_comparison( + input_table_1, got->mutable_view(), column_order); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - self_comparison( + self_comparison( input_table_1, got->mutable_view(), column_order); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); } @@ -203,12 +244,37 @@ TYPED_TEST(NaNTableViewTest, TestLexicographicalComparatorTwoTableNaNCase) cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); cudf::test::fixed_width_column_wrapper expected{{0, 0, 0, 0}}; - row_comparison( + row_comparison( input_table_1, input_table_2, got->mutable_view(), column_order); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); cudf::test::fixed_width_column_wrapper sorting_expected{{0, 1, 0, 0}}; - row_comparison( + row_comparison( + input_table_1, input_table_2, got->mutable_view(), column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, got->view()); +} + +TYPED_TEST(NaNTableViewTest, TestEqualityComparatorTwoTableNaNCase) +{ + using T = TypeParam; + + cudf::test::fixed_width_column_wrapper col1{{T(NAN), T(NAN), T(1), T(1)}}; + cudf::test::fixed_width_column_wrapper col2{{T(NAN), T(1), T(NAN), T(1)}}; + std::vector column_order{cudf::order::DESCENDING}; + + cudf::table_view input_table_1{{col1}}; + cudf::table_view input_table_2{{col2}}; + + auto got = cudf::make_numeric_column( + cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); + + cudf::test::fixed_width_column_wrapper expected{{0, 0, 0, 1}}; + row_equality( + input_table_1, input_table_2, got->mutable_view(), column_order); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + + cudf::test::fixed_width_column_wrapper sorting_expected{{1, 0, 0, 1}}; + row_equality( input_table_1, input_table_2, got->mutable_view(), column_order); CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, got->view()); } From e4cee9542d408bb2857dcd36e3dcdecba4606f25 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Wed, 1 Jun 2022 15:38:21 -0700 Subject: [PATCH 08/31] reorder cmake test file --- cpp/tests/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 9273ad53097..fef8f003a03 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -310,10 +310,10 @@ ConfigureTest(TRANSPOSE_TEST transpose/transpose_test.cpp) # * table tests ----------------------------------------------------------------------------------- ConfigureTest( TABLE_TEST - table/table_tests.cpp - table/table_view_tests.cu table/experimental_table_view_tests.cu table/row_operators_tests.cpp + table/table_tests.cpp + table/table_view_tests.cu ) # ################################################################################################## From 58d26635712158743dbde003dc525a341564c8b8 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Wed, 1 Jun 2022 16:02:16 -0700 Subject: [PATCH 09/31] fix cmake formatting --- cpp/tests/CMakeLists.txt | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index fef8f003a03..aeb1d822717 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -309,11 +309,8 @@ ConfigureTest(TRANSPOSE_TEST transpose/transpose_test.cpp) # ################################################################################################## # * table tests ----------------------------------------------------------------------------------- ConfigureTest( - TABLE_TEST - table/experimental_table_view_tests.cu + TABLE_TEST table/table_tests.cpp table/table_view_tests.cu table/experimental_table_view_tests.cu table/row_operators_tests.cpp - table/table_tests.cpp - table/table_view_tests.cu ) # ################################################################################################## From 49be087cd1c24681e446df77b2d15afe962859a7 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 09:53:17 -0700 Subject: [PATCH 10/31] Apply suggestions from code review Co-authored-by: Bradley Dice --- .../cudf/table/experimental/row_operators.cuh | 76 +++++++++---------- 1 file changed, 37 insertions(+), 39 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 06b8363da26..48666518fef 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -136,15 +136,15 @@ using rhs_iterator = strong_index_iterator; namespace lexicographic { /** - * @brief Relational comparator functor that evaluates `NaN` as not less, not equivalent, and not - * greater than all other values. This is IEE 754 compliant. + * @brief Relational comparator functor that evaluates `NaN` as not less than, equal to, or + * greater than other values. This is IEEE-754 compliant. */ struct physical_element_comparator { /** - * @brief Shared operator for all `Element` types relational comparisons. + * @brief Operator for relational comparisons. * - * @param[in] lhs first element - * @param[in] rhs second element + * @param lhs First element + * @param rhs Second element * @return Indicates the relationship between the elements in * the `lhs` and `rhs` columns. */ @@ -161,13 +161,11 @@ struct physical_element_comparator { */ struct sorting_physical_element_comparator { /** - * @brief A specialization for non-floating-point `Element` type relational - * comparison to derive the order of the elements with respect to `lhs`. + * @brief Operator for relational comparison of non-floating point values. * - * @param[in] lhs first element - * @param[in] rhs second element - * @return Indicates the relationship between the elements in - * the `lhs` and `rhs` columns. + * @param lhs First element + * @param rhs Second element + * @return Relation between elements */ template )> __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept @@ -176,12 +174,11 @@ struct sorting_physical_element_comparator { } /** - * @brief A specialization for floating-point `Element` type relational comparison that compares - * the nulls according to null order. + * @brief Operator for relational comparison of floating point values. * - * @param lhs_is_null boolean representing if lhs is null - * @param rhs_is_null boolean representing if lhs is null - * @return Indicates the relationship between null in lhs and rhs columns. + * @param lhs First element + * @param rhs Second element + * @return Relation between elements */ template )> __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept @@ -592,7 +589,7 @@ class self_comparator { * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware - * relational comparator that evaluates `NaN` as greater than standard values. + * relational comparator that evaluates `NaN` as greater than all other values. */ template @@ -708,7 +705,7 @@ class two_table_comparator { * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware - * relational comparator that evaluates `NaN` as greater than standard values. + * relational comparator that evaluates `NaN` as greater than all other values. */ template @@ -740,17 +737,18 @@ class row_hasher; namespace equality { /** - * @brief Equality comparator functor that evaluates `NaN` not equal to all other values for IEE 754 + * @brief Equality comparator functor that evaluates `NaN` not equal to all other values for IEEE-754 * compliance. */ struct physical_equality_comparator { /** - * @brief Shared operator for all `Element` types to check if `lhs` is equivalent to `rhs`. `nan - * != nan`. + * @brief Operator for equality comparisons. + * + * Note that `NaN != NaN`, following IEEE-754. * - * @param lhs first element - * @param rhs second element - * @return `true` if `lhs` == `rhs` else `false`. + * @param lhs First element + * @param rhs Second element + * @return `true` if `lhs == rhs` else `false` */ template __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept @@ -760,16 +758,15 @@ struct physical_equality_comparator { }; /** - * @brief Equality comparator functor that evaluates `NaN` as equivalent to other `NaN`s . + * @brief Equality comparator functor that evaluates `NaN` as equal to other `NaN`s. */ struct nan_equal_physical_equality_comparator { /** - * @brief A specialization for non-floating-point `Element` type to check if - * `lhs` is equivalent to `rhs`. + * @brief Operator for equality comparison of non-floating point values. * - * @param lhs first element - * @param rhs second element - * @return `true` if `lhs` == `rhs` else `false`. + * @param lhs First element + * @param rhs Second element + * @return `true` if `lhs == rhs` else `false` */ template )> __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept @@ -778,12 +775,13 @@ struct nan_equal_physical_equality_comparator { } /** - * @brief A specialization for floating-point `Element` type to check if - * `lhs` is equivalent to `rhs`. `nan == nan`. + * @brief Operator for equality comparison of floating point values. * - * @param lhs first element - * @param rhs second element - * @return `true` if `lhs` == `rhs` else `false`. + * Note that `NaN == NaN`. + * + * @param lhs First element + * @param rhs Second element + * @return `true` if `lhs` == `rhs` else `false` */ template )> __device__ constexpr bool operator()(Element const lhs, Element const rhs) const noexcept @@ -801,8 +799,8 @@ struct nan_equal_physical_equality_comparator { * - If the rows are compared without mismatched elements, the rows are equivalent * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A equality comparator functor, defaults to a `NaN == NaN` - * equality comparator. + * @tparam PhysicalElementComparator A equality comparator functor, defaults to a comparator for + * which `NaN == NaN`. */ template @@ -1076,8 +1074,8 @@ class self_comparator { * `F(i,j)` returns true if and only if row `i` compares equal to row `j`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A equality comparator functor, defaults to a `NaN == NaN` - * equality comparator. + * @tparam PhysicalElementComparator A equality comparator functor, defaults to a comparator for + * which `NaN == NaN`. */ template From d0f64f2ef0cf52bd40cd884b1f09ab1ff26768a5 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 10:18:33 -0700 Subject: [PATCH 11/31] ctad refactor + split device_comaparator call --- .../cudf/table/experimental/row_operators.cuh | 67 ++++++++++++++----- cpp/src/search/search_ordered.cu | 2 +- cpp/src/sort/sort_impl.cuh | 2 +- .../table/experimental_table_view_tests.cu | 4 +- 4 files changed, 54 insertions(+), 21 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 06b8363da26..d2bc3a62634 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -421,6 +421,10 @@ class device_row_comparator { */ template struct weak_ordering_comparator_impl { + static_assert(not((weak_ordering::EQUIVALENT == values) && ...), + "weak_ordering_comparator should not be used for pure equality comparisons. The " + "`row_equality_comparator` should be used instead"); + template __device__ constexpr bool operator()(LhsType const lhs_index, RhsType const rhs_index) const noexcept @@ -438,11 +442,21 @@ struct weak_ordering_comparator_impl { * @tparam Nullate A cudf::nullate type describing whether to check for nulls. */ template -using less_comparator = weak_ordering_comparator_impl; +struct less_comparator : weak_ordering_comparator_impl { + less_comparator(Comparator const& c) + : weak_ordering_comparator_impl{c} + { + } +}; template -using less_equivalent_comparator = - weak_ordering_comparator_impl; +struct less_equivalent_comparator + : weak_ordering_comparator_impl { + less_equivalent_comparator(Comparator const& c) + : weak_ordering_comparator_impl{c} + { + } +}; struct preprocessed_table { using table_device_view_owner = @@ -596,12 +610,18 @@ class self_comparator { */ template - less_comparator> device_comparator( - Nullate nullate = {}) const noexcept + auto device_less_comparator(Nullate nullate = {}) const noexcept + { + return less_comparator{device_row_comparator{ + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence()}}; + } + + template + auto device_less_equivalent_comparator(Nullate nullate = {}) const noexcept { - return less_comparator>{ - device_row_comparator( - nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence())}; + return less_equivalent_comparator{device_row_comparator( + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence())}; } private: @@ -610,6 +630,8 @@ class self_comparator { template struct strong_index_comparator_adapter { + strong_index_comparator_adapter(Comparator const& c) : comparator{c} {} + __device__ constexpr weak_ordering operator()(lhs_index_type const lhs_index, rhs_index_type const rhs_index) const noexcept { @@ -712,18 +734,28 @@ class two_table_comparator { */ template - less_comparator< - strong_index_comparator_adapter>> - device_comparator(Nullate nullate = {}) const noexcept + auto device_less_comparator(Nullate nullate = {}) const noexcept { - return less_comparator< - strong_index_comparator_adapter>>{ - device_row_comparator(nullate, + return less_comparator{strong_index_comparator_adapter{ + device_row_comparator{nullate, *d_left_table, *d_right_table, d_left_table->depths(), d_left_table->column_order(), - d_left_table->null_precedence())}; + d_left_table->null_precedence()}}}; + } + + template + auto device_less_equivalent_comparator(Nullate nullate = {}) const noexcept + { + return less_equivalent_comparator{strong_index_comparator_adapter{ + device_row_comparator{nullate, + *d_left_table, + *d_right_table, + d_left_table->depths(), + d_left_table->column_order(), + d_left_table->null_precedence()}}}; } private: @@ -1094,6 +1126,8 @@ class self_comparator { template struct strong_index_comparator_adapter { + strong_index_comparator_adapter(Comparator const& c) : comparator{c} {} + __device__ constexpr bool operator()(lhs_index_type const lhs_index, rhs_index_type const rhs_index) const noexcept { @@ -1178,8 +1212,7 @@ class two_table_comparator { auto device_comparator(Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL) const noexcept { - return strong_index_comparator_adapter< - device_row_comparator>{ + return strong_index_comparator_adapter{ device_row_comparator( nullate, *d_left_table, *d_right_table, nulls_are_equal)}; } diff --git a/cpp/src/search/search_ordered.cu b/cpp/src/search/search_ordered.cu index 2ae776420e2..5486d79111d 100644 --- a/cpp/src/search/search_ordered.cu +++ b/cpp/src/search/search_ordered.cu @@ -68,7 +68,7 @@ std::unique_ptr search_ordered(table_view const& haystack, auto const comparator = cudf::experimental::row::lexicographic::two_table_comparator( matched_haystack, matched_needles, column_order, null_precedence, stream); auto const has_nulls = has_nested_nulls(matched_haystack) or has_nested_nulls(matched_needles); - auto const d_comparator = comparator.device_comparator(nullate::DYNAMIC{has_nulls}); + auto const d_comparator = comparator.device_less_comparator(nullate::DYNAMIC{has_nulls}); auto const haystack_it = cudf::experimental::row::lhs_iterator(0); auto const needles_it = cudf::experimental::row::rhs_iterator(0); diff --git a/cpp/src/sort/sort_impl.cuh b/cpp/src/sort/sort_impl.cuh index 7f84c49a417..52e2578e809 100644 --- a/cpp/src/sort/sort_impl.cuh +++ b/cpp/src/sort/sort_impl.cuh @@ -127,7 +127,7 @@ std::unique_ptr sorted_order(table_view input, auto comp = experimental::row::lexicographic::self_comparator(input, column_order, null_precedence, stream); - auto comparator = comp.device_comparator(nullate::DYNAMIC{has_nested_nulls(input)}); + auto comparator = comp.device_less_comparator(nullate::DYNAMIC{has_nested_nulls(input)}); if (stable) { thrust::stable_sort(rmm::exec_policy(stream), diff --git a/cpp/tests/table/experimental_table_view_tests.cu b/cpp/tests/table/experimental_table_view_tests.cu index d847a8272be..2bf9b79996f 100644 --- a/cpp/tests/table/experimental_table_view_tests.cu +++ b/cpp/tests/table/experimental_table_view_tests.cu @@ -57,7 +57,7 @@ void row_comparison(cudf::table_view input1, auto table_comparator = lexicographic::two_table_comparator{input1, input2, column_order, {}, stream}; auto comparator = - table_comparator.device_comparator(cudf::nullate::NO{}); + table_comparator.device_less_comparator(cudf::nullate::NO{}); auto const lhs_it = cudf::experimental::row::lhs_iterator(0); auto const rhs_it = cudf::experimental::row::rhs_iterator(0); thrust::transform(rmm::exec_policy(stream), @@ -77,7 +77,7 @@ void self_comparison(cudf::table_view input, auto table_comparator = lexicographic::self_comparator{input, column_order, {}, stream}; auto comparator = - table_comparator.device_comparator(cudf::nullate::NO{}); + table_comparator.device_less_comparator(cudf::nullate::NO{}); thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.num_rows()), From 1bd540560a4bac47fd75345c0da91f50c219e420 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 10:31:22 -0700 Subject: [PATCH 12/31] rename experimental op test file --- cpp/tests/CMakeLists.txt | 4 ++-- ...table_view_tests.cu => experimental_row_operator_tests.cu} | 0 2 files changed, 2 insertions(+), 2 deletions(-) rename cpp/tests/table/{experimental_table_view_tests.cu => experimental_row_operator_tests.cu} (100%) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index c8ab920e57e..816c5a1c59c 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -309,8 +309,8 @@ ConfigureTest(TRANSPOSE_TEST transpose/transpose_test.cpp) # ################################################################################################## # * table tests ----------------------------------------------------------------------------------- ConfigureTest( - TABLE_TEST table/table_tests.cpp table/table_view_tests.cu table/experimental_table_view_tests.cu - table/row_operators_tests.cpp + TABLE_TEST table/table_tests.cpp table/table_view_tests.cu table/row_operators_tests.cpp + table/experimental_row_operator_tests.cu ) # ################################################################################################## diff --git a/cpp/tests/table/experimental_table_view_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu similarity index 100% rename from cpp/tests/table/experimental_table_view_tests.cu rename to cpp/tests/table/experimental_row_operator_tests.cu From 70da3b45a9fb7dcee2ed980d871457e2dbe68b8c Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 11:04:34 -0700 Subject: [PATCH 13/31] comment cleanup and pr fixes --- .../cudf/table/experimental/row_operators.cuh | 17 +++++++---------- .../table/experimental_row_operator_tests.cu | 4 ++-- 2 files changed, 9 insertions(+), 12 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index e402dc48245..9b6eacd07bc 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -137,7 +137,7 @@ namespace lexicographic { /** * @brief Relational comparator functor that evaluates `NaN` as not less than, equal to, or - * greater than other values. This is IEEE-754 compliant. + * greater than other values. This is IEEE-754 compliant. */ struct physical_element_comparator { /** @@ -145,8 +145,7 @@ struct physical_element_comparator { * * @param lhs First element * @param rhs Second element - * @return Indicates the relationship between the elements in - * the `lhs` and `rhs` columns. + * @return Relation between elements */ template __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept @@ -157,7 +156,7 @@ struct physical_element_comparator { /** * @brief Relational comparator functor that evaluates `NaN` as equivalent to other `NaN`s and - * greater than all other values. + * greater than all other values. */ struct sorting_physical_element_comparator { /** @@ -183,12 +182,10 @@ struct sorting_physical_element_comparator { template )> __device__ constexpr weak_ordering operator()(Element const lhs, Element const rhs) const noexcept { - if (isnan(lhs) and isnan(rhs)) { - return weak_ordering::EQUIVALENT; + if (isnan(lhs)) { + return isnan(rhs) ? weak_ordering::EQUIVALENT : weak_ordering::GREATER; } else if (isnan(rhs)) { return weak_ordering::LESS; - } else if (isnan(lhs)) { - return weak_ordering::GREATER; } return detail::compare_elements(lhs, rhs); @@ -212,7 +209,7 @@ struct sorting_physical_element_comparator { * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware - * relational comparator that evaluates `NaN` as greater than standard values. + * relational comparator that evaluates `NaN` as greater than all non-null values. */ template @@ -1031,7 +1028,7 @@ class device_row_comparator { table_device_view const rhs; Nullate const check_nulls; null_equality const nulls_are_equal; -}; // class device_row_comparator +}; struct preprocessed_table { /** diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 2bf9b79996f..5af41c1adda 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -103,7 +103,7 @@ void row_equality(cudf::table_view input1, lhs_it, lhs_it + input1.num_rows(), rhs_it, - output.data(), + output.data(), comparator); } @@ -121,7 +121,7 @@ void self_equality(cudf::table_view input, thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.num_rows()), thrust::make_counting_iterator(0), - output.data(), + output.data(), comparator); } From e4a70297b6de0cf84e5db706ddb385491c12ad94 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 12:17:46 -0700 Subject: [PATCH 14/31] physical comparator clarification docs --- .../cudf/table/experimental/row_operators.cuh | 43 +++++++++++-------- 1 file changed, 25 insertions(+), 18 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 9b6eacd07bc..eebab275b60 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -136,8 +136,9 @@ using rhs_iterator = strong_index_iterator; namespace lexicographic { /** - * @brief Relational comparator functor that evaluates `NaN` as not less than, equal to, or - * greater than other values. This is IEEE-754 compliant. + * @brief Relational comparator functor that compares individual physical values rather than logical + * elements like lists, strings, or structs. It evaluates `NaN` as not less than, equal to, or + * greater than other values and is IEEE-754 compliant. */ struct physical_element_comparator { /** @@ -155,7 +156,8 @@ struct physical_element_comparator { }; /** - * @brief Relational comparator functor that evaluates `NaN` as equivalent to other `NaN`s and + * @brief Relational comparator functor that compares individual physical values rather than logical + * elements like lists, strings, or structs. It evaluates `NaN` as equivalent to other `NaN`s and * greater than all other values. */ struct sorting_physical_element_comparator { @@ -208,8 +210,9 @@ struct sorting_physical_element_comparator { * `aac < abb`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware - * relational comparator that evaluates `NaN` as greater than all non-null values. + * @tparam PhysicalElementComparator A relational comparator functor that compares individual values + * rather than logical elements, defaults to `NaN` aware relational comparator that evaluates `NaN` + * as greater than all non-null values. */ template @@ -599,8 +602,9 @@ class self_comparator { * `F(i,j)` returns true if and only if row `i` compares lexicographically less than row `j`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware - * relational comparator that evaluates `NaN` as greater than all other values. + * @tparam PhysicalElementComparator A relational comparator functor that compares individual + * values rather than logical elements, defaults to `NaN` aware relational comparator that + * evaluates `NaN` as greater than all other values. */ template @@ -723,8 +727,9 @@ class two_table_comparator { * `j` of the left table. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A relational comparator functor, defaults to `NaN` aware - * relational comparator that evaluates `NaN` as greater than all other values. + * @tparam PhysicalElementComparator A relational comparator functor that compares individual + * values rather than logical elements, defaults to `NaN` aware relational comparator that + * evaluates `NaN` as greater than all other values. */ template @@ -766,8 +771,9 @@ class row_hasher; namespace equality { /** - * @brief Equality comparator functor that evaluates `NaN` not equal to all other values for IEEE-754 - * compliance. + * @brief Equality comparator functor that compares individual physical values rather than logical + * elements like lists, strings, or structs. It evaluates `NaN` not equal to all other values for + * IEEE-754. compliance. */ struct physical_equality_comparator { /** @@ -787,7 +793,8 @@ struct physical_equality_comparator { }; /** - * @brief Equality comparator functor that evaluates `NaN` as equal to other `NaN`s. + * @brief Equality comparator functor that compares individual physical values rather than logical + * elements like lists, strings, or structs. It evaluates `NaN` as equal to other `NaN`s. */ struct nan_equal_physical_equality_comparator { /** @@ -828,8 +835,8 @@ struct nan_equal_physical_equality_comparator { * - If the rows are compared without mismatched elements, the rows are equivalent * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A equality comparator functor, defaults to a comparator for - * which `NaN == NaN`. + * @tparam PhysicalElementComparator A equality comparator functor that compares individual values + * rather than logical elements, defaults to a comparator for which `NaN == NaN`. */ template @@ -1103,8 +1110,8 @@ class self_comparator { * `F(i,j)` returns true if and only if row `i` compares equal to row `j`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A equality comparator functor, defaults to a comparator for - * which `NaN == NaN`. + * @tparam PhysicalElementComparator A equality comparator functor that compares individual values + * rather than logical elements, defaults to a comparator for which `NaN == NaN`. */ template @@ -1199,8 +1206,8 @@ class two_table_comparator { * right table compares equal to row `j` of the left table. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A equality comparator functor, defaults to a `NaN == NaN` - * equality comparator. + * @tparam PhysicalElementComparator A equality comparator functor that compares individual values + * rather than logical elements, defaults to a `NaN == NaN` equality comparator. */ template From 1240c85192f59b026e28bdee56ad581be9698dcd Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 12:24:14 -0700 Subject: [PATCH 15/31] fix whitespace --- .../cudf/table/experimental/row_operators.cuh | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index eebab275b60..4db6bc2f491 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -110,7 +110,7 @@ struct strong_index_iterator : public thrust::iterator_facade const& other) const noexcept { return begin == other.begin; - } + }q __device__ constexpr Index dereference() const noexcept { return static_cast(begin); } @@ -136,9 +136,9 @@ using rhs_iterator = strong_index_iterator; namespace lexicographic { /** - * @brief Relational comparator functor that compares individual physical values rather than logical - * elements like lists, strings, or structs. It evaluates `NaN` as not less than, equal to, or - * greater than other values and is IEEE-754 compliant. + * @brief Relational comparator functor that compares physical values rather than logical + * elements like lists, strings, or structs. It evaluates `NaN` as not less than, equal to, or greater + * than other values and is IEEE-754 compliant. */ struct physical_element_comparator { /** @@ -156,9 +156,9 @@ struct physical_element_comparator { }; /** - * @brief Relational comparator functor that compares individual physical values rather than logical + * @brief Relational comparator functor that compares physical values rather than logical * elements like lists, strings, or structs. It evaluates `NaN` as equivalent to other `NaN`s and - * greater than all other values. + * greater than all other values. */ struct sorting_physical_element_comparator { /** @@ -771,7 +771,7 @@ class row_hasher; namespace equality { /** - * @brief Equality comparator functor that compares individual physical values rather than logical + * @brief Equality comparator functor that compares physical values rather than logical * elements like lists, strings, or structs. It evaluates `NaN` not equal to all other values for * IEEE-754. compliance. */ @@ -793,7 +793,7 @@ struct physical_equality_comparator { }; /** - * @brief Equality comparator functor that compares individual physical values rather than logical + * @brief Equality comparator functor that compares physical values rather than logical * elements like lists, strings, or structs. It evaluates `NaN` as equal to other `NaN`s. */ struct nan_equal_physical_equality_comparator { From 33046152f1e89169bdffd139cadc443e4d89077d Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 13:18:30 -0700 Subject: [PATCH 16/31] fix formatting --- cpp/include/cudf/table/experimental/row_operators.cuh | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 4db6bc2f491..da89c813583 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -110,7 +110,7 @@ struct strong_index_iterator : public thrust::iterator_facade const& other) const noexcept { return begin == other.begin; - }q + } __device__ constexpr Index dereference() const noexcept { return static_cast(begin); } @@ -137,8 +137,8 @@ namespace lexicographic { /** * @brief Relational comparator functor that compares physical values rather than logical - * elements like lists, strings, or structs. It evaluates `NaN` as not less than, equal to, or greater - * than other values and is IEEE-754 compliant. + * elements like lists, strings, or structs. It evaluates `NaN` as not less than, equal to, or + * greater than other values and is IEEE-754 compliant. */ struct physical_element_comparator { /** From 9ca2d2912a27df489ba55b610b550c5f6e2c5c98 Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 13:44:16 -0700 Subject: [PATCH 17/31] fix copyright --- cpp/tests/table/experimental_row_operator_tests.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 5af41c1adda..393b81379e2 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2022, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. From 51938fdc4b9b4ec8bdf831ac95306c1e972b420f Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 16:23:32 -0700 Subject: [PATCH 18/31] Functors as arguments, enabling CTAD --- .../cudf/table/experimental/row_operators.cuh | 166 +++++++++++------- .../table/experimental_row_operator_tests.cu | 74 +++++--- 2 files changed, 150 insertions(+), 90 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index bf89039d06d..9783846e6a2 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -242,20 +242,22 @@ class device_row_comparator { * @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 `nullopt`, then null precedence would be * `null_order::BEFORE` for all columns. + * @param comparator Physical element relational comparison functor. */ - device_row_comparator( - Nullate check_nulls, - table_device_view lhs, - table_device_view rhs, - std::optional> depth = std::nullopt, - std::optional> column_order = std::nullopt, - std::optional> null_precedence = std::nullopt) noexcept + device_row_comparator(Nullate check_nulls, + table_device_view lhs, + table_device_view rhs, + std::optional> depth = std::nullopt, + std::optional> column_order = std::nullopt, + std::optional> null_precedence = std::nullopt, + PhysicalElementComparator comparator = {}) noexcept : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, _depth{depth}, _column_order{column_order}, - _null_precedence{null_precedence} + _null_precedence{null_precedence}, + _comparator{comparator} { } @@ -276,17 +278,20 @@ class device_row_comparator { * @param null_precedence Indicates how null values are ordered with other values * @param depth The depth of the column if part of a nested column @see * preprocessed_table::depths + * @param comparator Physical element relational comparison functor. */ __device__ element_comparator(Nullate check_nulls, column_device_view lhs, column_device_view rhs, - null_order null_precedence = null_order::BEFORE, - int depth = 0) + null_order null_precedence = null_order::BEFORE, + int depth = 0, + PhysicalElementComparator comparator = {}) : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, _null_precedence{null_precedence}, - _depth{depth} + _depth{depth}, + _comparator{comparator} { } @@ -312,8 +317,8 @@ class device_row_comparator { } } - return cuda::std::pair(PhysicalElementComparator{}(_lhs.element(lhs_element_index), - _rhs.element(rhs_element_index)), + return cuda::std::pair(_comparator(_lhs.element(lhs_element_index), + _rhs.element(rhs_element_index)), std::numeric_limits::max()); } @@ -352,15 +357,18 @@ class device_row_comparator { ++depth; } - auto const comparator = element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth}; return cudf::type_dispatcher( - lcol.type(), comparator, lhs_element_index, rhs_element_index); + lcol.type(), + element_comparator{_check_nulls, lcol, rcol, _null_precedence, depth, _comparator}, + lhs_element_index, + rhs_element_index); } private: column_device_view const _lhs; column_device_view const _rhs; Nullate const _check_nulls; + PhysicalElementComparator const _comparator; null_order const _null_precedence; int const _depth; }; @@ -389,11 +397,13 @@ class device_row_comparator { null_order const null_precedence = _null_precedence.has_value() ? (*_null_precedence)[i] : null_order::BEFORE; - auto const comparator = - element_comparator{_check_nulls, _lhs.column(i), _rhs.column(i), null_precedence, depth}; weak_ordering state; - cuda::std::tie(state, last_null_depth) = - cudf::type_dispatcher(_lhs.column(i).type(), comparator, lhs_index, rhs_index); + cuda::std::tie(state, last_null_depth) = cudf::type_dispatcher( + _lhs.column(i).type(), + element_comparator{ + _check_nulls, _lhs.column(i), _rhs.column(i), null_precedence, depth, _comparator}, + lhs_index, + rhs_index); if (state == weak_ordering::EQUIVALENT) { continue; } @@ -407,7 +417,8 @@ class device_row_comparator { private: table_device_view const _lhs; table_device_view const _rhs; - Nullate const _check_nulls{}; + Nullate const _check_nulls; + PhysicalElementComparator const _comparator; std::optional> const _depth; std::optional> const _column_order; std::optional> const _null_precedence; @@ -619,22 +630,25 @@ class self_comparator { * values rather than logical elements, defaults to `NaN` aware relational comparator that * evaluates `NaN` as greater than all other values. * @param nullate Indicates if either input column contains nulls. + * @param comparator Physical element relational comparison functor. * @return A binary callable object. */ template - auto device_less_comparator(Nullate nullate = {}) const noexcept + auto device_less_comparator(Nullate nullate = {}, + PhysicalElementComparator comparator = {}) const noexcept { - return less_comparator{device_row_comparator{ - nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence()}}; + return less_comparator{device_row_comparator{ + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), comparator}}; } template - auto device_less_equivalent_comparator(Nullate nullate = {}) const noexcept + auto device_less_equivalent_comparator(Nullate nullate = {}, + PhysicalElementComparator comparator = {}) const noexcept { - return less_equivalent_comparator{device_row_comparator( - nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence())}; + return less_equivalent_comparator{device_row_comparator{ + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), comparator}}; } private: @@ -748,32 +762,37 @@ class two_table_comparator { * values rather than logical elements, defaults to `NaN` aware relational comparator that * evaluates `NaN` as greater than all other values. * @param nullate Indicates if either input column contains nulls. + * @param comparator Physical element relational comparison functor. * @return A binary callable object. */ template - auto device_less_comparator(Nullate nullate = {}) const noexcept + auto device_less_comparator(Nullate nullate = {}, + PhysicalElementComparator comparator = {}) const noexcept { - return less_comparator{strong_index_comparator_adapter{ - device_row_comparator{nullate, - *d_left_table, - *d_right_table, - d_left_table->depths(), - d_left_table->column_order(), - d_left_table->null_precedence()}}}; + return less_comparator{ + strong_index_comparator_adapter{device_row_comparator{nullate, + *d_left_table, + *d_right_table, + d_left_table->depths(), + d_left_table->column_order(), + d_left_table->null_precedence(), + comparator}}}; } template - auto device_less_equivalent_comparator(Nullate nullate = {}) const noexcept + auto device_less_equivalent_comparator(Nullate nullate = {}, + PhysicalElementComparator comparator = {}) const noexcept { - return less_equivalent_comparator{strong_index_comparator_adapter{ - device_row_comparator{nullate, - *d_left_table, - *d_right_table, - d_left_table->depths(), - d_left_table->column_order(), - d_left_table->null_precedence()}}}; + return less_equivalent_comparator{ + strong_index_comparator_adapter{device_row_comparator{nullate, + *d_left_table, + *d_right_table, + d_left_table->depths(), + d_left_table->column_order(), + d_left_table->null_precedence(), + comparator}}}; } private: @@ -854,7 +873,7 @@ struct nan_equal_physical_equality_comparator { * - If the rows are compared without mismatched elements, the rows are equivalent * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A equality comparator functor that compares individual values + * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual values * rather than logical elements, defaults to a comparator for which `NaN == NaN`. */ template (lhs_element_index), - rhs.element(rhs_element_index)); + return comparator(lhs.element(lhs_element_index), + rhs.element(rhs_element_index)); } template (lcol.type(), comp); } @@ -1048,12 +1082,14 @@ class device_row_comparator { column_device_view const rhs; Nullate const check_nulls; null_equality const nulls_are_equal; + PhysicalEqualityComparator const comparator; }; table_device_view const lhs; table_device_view const rhs; Nullate const check_nulls; null_equality const nulls_are_equal; + PhysicalEqualityComparator const comparator; }; /** @@ -1139,19 +1175,20 @@ class self_comparator { * `F(i,j)` returns true if and only if row `i` compares equal to row `j`. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A equality comparator functor that compares individual values - * rather than logical elements, defaults to a comparator for which `NaN == NaN`. + * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual + * values rather than logical elements, defaults to a comparator for which `NaN == NaN`. * @param nullate Indicates if either input column contains nulls. * @param nulls_are_equal Indicates if nulls are equal. + * @param comparator Physical element equality comparison functor. * @return A binary callable object */ template - device_row_comparator device_comparator( - Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL) const noexcept + auto device_comparator(Nullate nullate = {}, + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator comparator = {}) const noexcept { - return device_row_comparator( - nullate, *d_t, *d_t, nulls_are_equal); + return device_row_comparator{nullate, *d_t, *d_t, nulls_are_equal, comparator}; } private: @@ -1240,20 +1277,21 @@ class two_table_comparator { * right table compares equal to row `j` of the left table. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. - * @tparam PhysicalElementComparator A equality comparator functor that compares individual values - * rather than logical elements, defaults to a `NaN == NaN` equality comparator. + * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual + * values rather than logical elements, defaults to a `NaN == NaN` equality comparator. * @param nullate Indicates if either input column contains nulls. * @param nulls_are_equal Indicates if nulls are equal. + * @param comparator Physical element equality comparison functor. * @return A binary callable object */ template - auto device_comparator(Nullate nullate = {}, - null_equality nulls_are_equal = null_equality::EQUAL) const noexcept + auto device_comparator(Nullate nullate = {}, + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator comparator = {}) const noexcept { return strong_index_comparator_adapter{ - device_row_comparator( - nullate, *d_left_table, *d_right_table, nulls_are_equal)}; + device_row_comparator(nullate, *d_left_table, *d_right_table, nulls_are_equal, comparator)}; } private: diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 393b81379e2..d308c580799 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -50,14 +50,14 @@ template void row_comparison(cudf::table_view input1, cudf::table_view input2, cudf::mutable_column_view output, - std::vector const& column_order) + std::vector const& column_order, + Comparator c) { rmm::cuda_stream_view stream{}; auto table_comparator = lexicographic::two_table_comparator{input1, input2, column_order, {}, stream}; - auto comparator = - table_comparator.device_less_comparator(cudf::nullate::NO{}); + auto comparator = table_comparator.device_less_comparator(cudf::nullate::NO{}, c); auto const lhs_it = cudf::experimental::row::lhs_iterator(0); auto const rhs_it = cudf::experimental::row::rhs_iterator(0); thrust::transform(rmm::exec_policy(stream), @@ -71,13 +71,13 @@ void row_comparison(cudf::table_view input1, template void self_comparison(cudf::table_view input, cudf::mutable_column_view output, - std::vector const& column_order) + std::vector const& column_order, + Comparator c) { rmm::cuda_stream_view stream{}; auto table_comparator = lexicographic::self_comparator{input, column_order, {}, stream}; - auto comparator = - table_comparator.device_less_comparator(cudf::nullate::NO{}); + auto comparator = table_comparator.device_less_comparator(cudf::nullate::NO{}, c); thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.num_rows()), @@ -90,13 +90,14 @@ template void row_equality(cudf::table_view input1, cudf::table_view input2, cudf::mutable_column_view output, - std::vector const& column_order) + std::vector const& column_order, + Comparator c) { rmm::cuda_stream_view stream{}; auto table_comparator = equality::two_table_comparator{input1, input2, stream}; auto comparator = - table_comparator.device_comparator(cudf::nullate::NO{}); + table_comparator.device_comparator(cudf::nullate::NO{}, cudf::null_equality::EQUAL, c); auto const lhs_it = cudf::experimental::row::lhs_iterator(0); auto const rhs_it = cudf::experimental::row::rhs_iterator(0); thrust::transform(rmm::exec_policy(stream), @@ -110,13 +111,14 @@ void row_equality(cudf::table_view input1, template void self_equality(cudf::table_view input, cudf::mutable_column_view output, - std::vector const& column_order) + std::vector const& column_order, + Comparator c) { rmm::cuda_stream_view stream{}; auto table_comparator = equality::self_comparator{input, stream}; auto comparator = - table_comparator.device_comparator(cudf::nullate::NO{}); + table_comparator.device_comparator(cudf::nullate::NO{}, cudf::null_equality::EQUAL, c); thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.num_rows()), @@ -150,12 +152,18 @@ TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTableCase) cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); fixed_width_column_wrapper expected{{1, 1, 0, 1}}; - row_comparison( - input_table_1, input_table_2, got->mutable_view(), column_order); + row_comparison(input_table_1, + input_table_2, + got->mutable_view(), + column_order, + lexicographic::physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - row_comparison( - input_table_1, input_table_2, got->mutable_view(), column_order); + row_comparison(input_table_1, + input_table_2, + got->mutable_view(), + column_order, + lexicographic::sorting_physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); } @@ -172,12 +180,14 @@ TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorSameTable) cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); fixed_width_column_wrapper expected{{0, 0, 0, 0}}; - self_comparison( - input_table_1, got->mutable_view(), column_order); + self_comparison( + input_table_1, got->mutable_view(), column_order, lexicographic::physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - self_comparison( - input_table_1, got->mutable_view(), column_order); + self_comparison(input_table_1, + got->mutable_view(), + column_order, + lexicographic::sorting_physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); } @@ -244,13 +254,19 @@ TYPED_TEST(NaNTableViewTest, TestLexicographicalComparatorTwoTableNaNCase) cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); cudf::test::fixed_width_column_wrapper expected{{0, 0, 0, 0}}; - row_comparison( - input_table_1, input_table_2, got->mutable_view(), column_order); + row_comparison(input_table_1, + input_table_2, + got->mutable_view(), + column_order, + lexicographic::physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); cudf::test::fixed_width_column_wrapper sorting_expected{{0, 1, 0, 0}}; - row_comparison( - input_table_1, input_table_2, got->mutable_view(), column_order); + row_comparison(input_table_1, + input_table_2, + got->mutable_view(), + column_order, + lexicographic::sorting_physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, got->view()); } @@ -269,12 +285,18 @@ TYPED_TEST(NaNTableViewTest, TestEqualityComparatorTwoTableNaNCase) cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); cudf::test::fixed_width_column_wrapper expected{{0, 0, 0, 1}}; - row_equality( - input_table_1, input_table_2, got->mutable_view(), column_order); + row_equality(input_table_1, + input_table_2, + got->mutable_view(), + column_order, + equality::physical_equality_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); cudf::test::fixed_width_column_wrapper sorting_expected{{1, 0, 0, 1}}; - row_equality( - input_table_1, input_table_2, got->mutable_view(), column_order); + row_equality(input_table_1, + input_table_2, + got->mutable_view(), + column_order, + equality::nan_equal_physical_equality_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, got->view()); } From 3f545fef5306850e084499f263d0b86335c05f4d Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 17:25:00 -0700 Subject: [PATCH 19/31] Docs, cleanup, and renaming --- .../cudf/table/experimental/row_operators.cuh | 97 +++++++++---------- cpp/src/search/search_ordered.cu | 2 +- cpp/src/sort/sort_impl.cuh | 2 +- .../table/experimental_row_operator_tests.cu | 4 +- 4 files changed, 49 insertions(+), 56 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 9783846e6a2..decf0819a2e 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -143,9 +143,11 @@ using rhs_iterator = strong_index_iterator; namespace lexicographic { /** - * @brief Relational comparator functor that compares physical values rather than logical - * elements like lists, strings, or structs. It evaluates `NaN` as not less than, equal to, or - * greater than other values and is IEEE-754 compliant. + * @brief Computes a weak ordering of two values with special sorting behavior. + * + * This relational comparator functor compares physical values rather than logical + * elements like lists, strings, or structs. It evaluates `NaN` as equivalent to other `NaN`s and + * greater than all other values. */ struct physical_element_comparator { /** @@ -219,7 +221,7 @@ struct sorting_physical_element_comparator { * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalElementComparator A relational comparator functor that compares individual values * rather than logical elements, defaults to `NaN` aware relational comparator that evaluates `NaN` - * as greater than all non-null values. + * as greater than all other values. */ template @@ -231,7 +233,7 @@ class device_row_comparator { * @brief Construct a function object for performing a lexicographic * comparison between the rows of two tables. * - * @param check_nulls Indicates if either input table contains columns with nulls. + * @param check_nulls Indicates if any input column contains columns with nulls. * @param lhs The first table * @param rhs The second table (may be the same table as `lhs`) * @param depth Optional, device array the same length as a row that contains starting depths of @@ -250,14 +252,14 @@ class device_row_comparator { std::optional> depth = std::nullopt, std::optional> column_order = std::nullopt, std::optional> null_precedence = std::nullopt, - PhysicalElementComparator comparator = {}) noexcept + PhysicalElementComparator c = {}) noexcept : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, _depth{depth}, _column_order{column_order}, _null_precedence{null_precedence}, - _comparator{comparator} + _comparator{c} { } @@ -283,15 +285,15 @@ class device_row_comparator { __device__ element_comparator(Nullate check_nulls, column_device_view lhs, column_device_view rhs, - null_order null_precedence = null_order::BEFORE, - int depth = 0, - PhysicalElementComparator comparator = {}) + null_order null_precedence = null_order::BEFORE, + int depth = 0, + PhysicalElementComparator c = {}) : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, _null_precedence{null_precedence}, _depth{depth}, - _comparator{comparator} + _comparator{c} { } @@ -629,26 +631,24 @@ class self_comparator { * @tparam PhysicalElementComparator A relational comparator functor that compares individual * values rather than logical elements, defaults to `NaN` aware relational comparator that * evaluates `NaN` as greater than all other values. - * @param nullate Indicates if either input column contains nulls. + * @param nullate Indicates if any input column contains nulls. * @param comparator Physical element relational comparison functor. * @return A binary callable object. */ template - auto device_less_comparator(Nullate nullate = {}, - PhysicalElementComparator comparator = {}) const noexcept + auto less(Nullate nullate = {}, PhysicalElementComparator c = {}) const noexcept { return less_comparator{device_row_comparator{ - nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), comparator}}; + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), c}}; } template - auto device_less_equivalent_comparator(Nullate nullate = {}, - PhysicalElementComparator comparator = {}) const noexcept + auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator c = {}) const noexcept { return less_equivalent_comparator{device_row_comparator{ - nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), comparator}}; + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), c}}; } private: @@ -761,14 +761,13 @@ class two_table_comparator { * @tparam PhysicalElementComparator A relational comparator functor that compares individual * values rather than logical elements, defaults to `NaN` aware relational comparator that * evaluates `NaN` as greater than all other values. - * @param nullate Indicates if either input column contains nulls. + * @param nullate Indicates if any input column contains nulls. * @param comparator Physical element relational comparison functor. * @return A binary callable object. */ template - auto device_less_comparator(Nullate nullate = {}, - PhysicalElementComparator comparator = {}) const noexcept + auto less(Nullate nullate = {}, PhysicalElementComparator c = {}) const noexcept { return less_comparator{ strong_index_comparator_adapter{device_row_comparator{nullate, @@ -777,13 +776,12 @@ class two_table_comparator { d_left_table->depths(), d_left_table->column_order(), d_left_table->null_precedence(), - comparator}}}; + c}}}; } template - auto device_less_equivalent_comparator(Nullate nullate = {}, - PhysicalElementComparator comparator = {}) const noexcept + auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator c = {}) const noexcept { return less_equivalent_comparator{ strong_index_comparator_adapter{device_row_comparator{nullate, @@ -792,7 +790,7 @@ class two_table_comparator { d_left_table->depths(), d_left_table->column_order(), d_left_table->null_precedence(), - comparator}}}; + c}}}; } private: @@ -811,7 +809,7 @@ namespace equality { /** * @brief Equality comparator functor that compares physical values rather than logical * elements like lists, strings, or structs. It evaluates `NaN` not equal to all other values for - * IEEE-754. compliance. + * IEEE-754 compliance. */ struct physical_equality_comparator { /** @@ -867,10 +865,9 @@ struct nan_equal_physical_equality_comparator { /** * @brief Computes the equality comparison between 2 rows. * - * Equality is determined by: - * - Two rows are compared element by element. - * - The first mismatching element returns false representing unequal rows - * - If the rows are compared without mismatched elements, the rows are equivalent + * Equality is determined by comparing rows element by element. The first mismatching element + * returns false, representing unequal rows. If the rows are compared without mismatched elements, + * the rows are equal. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual values @@ -910,7 +907,7 @@ class device_row_comparator { * @brief Construct a function object for performing equality comparison between the rows of two * tables. * - * @param check_nulls Indicates if either input table contains columns with nulls. + * @param check_nulls Indicates if any input column contains columns with nulls. * @param lhs The first table * @param rhs The second table (may be the same table as `lhs`) * @param nulls_are_equal Indicates if two null elements are treated as equivalent @@ -919,13 +916,9 @@ class device_row_comparator { device_row_comparator(Nullate check_nulls, table_device_view lhs, table_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator comparator = {}) noexcept - : lhs{lhs}, - rhs{rhs}, - check_nulls{check_nulls}, - nulls_are_equal{nulls_are_equal}, - comparator{comparator} + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator c = {}) noexcept + : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal}, comparator{c} { } @@ -949,13 +942,13 @@ class device_row_comparator { __device__ element_comparator(Nullate check_nulls, column_device_view lhs, column_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator comparator = {}) noexcept + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator c = {}) noexcept : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal}, - comparator{comparator} + comparator{c} { } @@ -1177,18 +1170,18 @@ class self_comparator { * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual * values rather than logical elements, defaults to a comparator for which `NaN == NaN`. - * @param nullate Indicates if either input column contains nulls. + * @param nullate Indicates if any input column contains nulls. * @param nulls_are_equal Indicates if nulls are equal. * @param comparator Physical element equality comparison functor. * @return A binary callable object */ template - auto device_comparator(Nullate nullate = {}, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator comparator = {}) const noexcept + auto device_comparator(Nullate nullate = {}, + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator c = {}) const noexcept { - return device_row_comparator{nullate, *d_t, *d_t, nulls_are_equal, comparator}; + return device_row_comparator{nullate, *d_t, *d_t, nulls_are_equal, c}; } private: @@ -1279,19 +1272,19 @@ class two_table_comparator { * @tparam Nullate A cudf::nullate type describing whether to check for nulls. * @tparam PhysicalEqualityComparator A equality comparator functor that compares individual * values rather than logical elements, defaults to a `NaN == NaN` equality comparator. - * @param nullate Indicates if either input column contains nulls. + * @param nullate Indicates if any input column contains nulls. * @param nulls_are_equal Indicates if nulls are equal. * @param comparator Physical element equality comparison functor. * @return A binary callable object */ template - auto device_comparator(Nullate nullate = {}, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator comparator = {}) const noexcept + auto device_comparator(Nullate nullate = {}, + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator c = {}) const noexcept { return strong_index_comparator_adapter{ - device_row_comparator(nullate, *d_left_table, *d_right_table, nulls_are_equal, comparator)}; + device_row_comparator(nullate, *d_left_table, *d_right_table, nulls_are_equal, c)}; } private: @@ -1515,7 +1508,7 @@ class row_hasher { * `F(i)` returns the hash of row i. * * @tparam Nullate A cudf::nullate type describing whether to check for nulls - * @param nullate Indicates if either input column contains nulls + * @param nullate Indicates if any input column contains nulls * @param seed The seed to use for the hash function * @return A hash operator to use on the device */ diff --git a/cpp/src/search/search_ordered.cu b/cpp/src/search/search_ordered.cu index 5486d79111d..01b990facdc 100644 --- a/cpp/src/search/search_ordered.cu +++ b/cpp/src/search/search_ordered.cu @@ -68,7 +68,7 @@ std::unique_ptr search_ordered(table_view const& haystack, auto const comparator = cudf::experimental::row::lexicographic::two_table_comparator( matched_haystack, matched_needles, column_order, null_precedence, stream); auto const has_nulls = has_nested_nulls(matched_haystack) or has_nested_nulls(matched_needles); - auto const d_comparator = comparator.device_less_comparator(nullate::DYNAMIC{has_nulls}); + auto const d_comparator = comparator.less(nullate::DYNAMIC{has_nulls}); auto const haystack_it = cudf::experimental::row::lhs_iterator(0); auto const needles_it = cudf::experimental::row::rhs_iterator(0); diff --git a/cpp/src/sort/sort_impl.cuh b/cpp/src/sort/sort_impl.cuh index 52e2578e809..f98fda307b8 100644 --- a/cpp/src/sort/sort_impl.cuh +++ b/cpp/src/sort/sort_impl.cuh @@ -127,7 +127,7 @@ std::unique_ptr sorted_order(table_view input, auto comp = experimental::row::lexicographic::self_comparator(input, column_order, null_precedence, stream); - auto comparator = comp.device_less_comparator(nullate::DYNAMIC{has_nested_nulls(input)}); + auto comparator = comp.less(nullate::DYNAMIC{has_nested_nulls(input)}); if (stable) { thrust::stable_sort(rmm::exec_policy(stream), diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index d308c580799..7233538de6f 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -57,7 +57,7 @@ void row_comparison(cudf::table_view input1, auto table_comparator = lexicographic::two_table_comparator{input1, input2, column_order, {}, stream}; - auto comparator = table_comparator.device_less_comparator(cudf::nullate::NO{}, c); + auto comparator = table_comparator.less(cudf::nullate::NO{}, c); auto const lhs_it = cudf::experimental::row::lhs_iterator(0); auto const rhs_it = cudf::experimental::row::rhs_iterator(0); thrust::transform(rmm::exec_policy(stream), @@ -77,7 +77,7 @@ void self_comparison(cudf::table_view input, rmm::cuda_stream_view stream{}; auto table_comparator = lexicographic::self_comparator{input, column_order, {}, stream}; - auto comparator = table_comparator.device_less_comparator(cudf::nullate::NO{}, c); + auto comparator = table_comparator.less(cudf::nullate::NO{}, c); thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.num_rows()), From 2eb810338e84117deefbccdda8e317dad9a5fcac Mon Sep 17 00:00:00 2001 From: Ryan Lee Date: Thu, 2 Jun 2022 21:44:38 -0700 Subject: [PATCH 20/31] device_comparator --> equal_to --- .../cudf/table/experimental/row_operators.cuh | 12 ++++++------ cpp/src/groupby/hash/groupby.cu | 2 +- cpp/src/reductions/scan/rank_scan.cu | 2 +- cpp/src/search/contains_nested.cu | 2 +- cpp/src/stream_compaction/distinct.cu | 2 +- cpp/tests/groupby/lists_tests.cu | 2 +- cpp/tests/table/experimental_row_operator_tests.cu | 6 ++---- 7 files changed, 13 insertions(+), 15 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index decf0819a2e..c5b99880e24 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -1177,9 +1177,9 @@ class self_comparator { */ template - auto device_comparator(Nullate nullate = {}, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator c = {}) const noexcept + auto equal_to(Nullate nullate = {}, + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator c = {}) const noexcept { return device_row_comparator{nullate, *d_t, *d_t, nulls_are_equal, c}; } @@ -1279,9 +1279,9 @@ class two_table_comparator { */ template - auto device_comparator(Nullate nullate = {}, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator c = {}) const noexcept + auto equal_to(Nullate nullate = {}, + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator c = {}) const noexcept { return strong_index_comparator_adapter{ device_row_comparator(nullate, *d_left_table, *d_right_table, nulls_are_equal, c)}; diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu index ab8d0089347..c07833520ab 100644 --- a/cpp/src/groupby/hash/groupby.cu +++ b/cpp/src/groupby/hash/groupby.cu @@ -568,7 +568,7 @@ std::unique_ptr groupby(table_view const& keys, auto preprocessed_keys = cudf::experimental::row::hash::preprocessed_table::create(keys, stream); auto const comparator = cudf::experimental::row::equality::self_comparator{preprocessed_keys}; auto const row_hash = cudf::experimental::row::hash::row_hasher{std::move(preprocessed_keys)}; - auto const d_key_equal = comparator.device_comparator(has_null, null_keys_are_equal); + auto const d_key_equal = comparator.equal_to(has_null, null_keys_are_equal); auto const d_row_hash = row_hash.device_hasher(has_null); size_type constexpr unused_key{std::numeric_limits::max()}; diff --git a/cpp/src/reductions/scan/rank_scan.cu b/cpp/src/reductions/scan/rank_scan.cu index 0ababbf0a3d..c6909bfd601 100644 --- a/cpp/src/reductions/scan/rank_scan.cu +++ b/cpp/src/reductions/scan/rank_scan.cu @@ -53,7 +53,7 @@ std::unique_ptr rank_generator(column_view const& order_by, { auto comp = cudf::experimental::row::equality::self_comparator(table_view{{order_by}}, stream); auto const device_comparator = - comp.device_comparator(nullate::DYNAMIC{has_nested_nulls(table_view({order_by}))}); + comp.equal_to(nullate::DYNAMIC{has_nested_nulls(table_view({order_by}))}); auto ranks = make_fixed_width_column( data_type{type_to_id()}, order_by.size(), mask_state::UNALLOCATED, stream, mr); auto mutable_ranks = ranks->mutable_view(); diff --git a/cpp/src/search/contains_nested.cu b/cpp/src/search/contains_nested.cu index c3143b12e90..f4332efb23f 100644 --- a/cpp/src/search/contains_nested.cu +++ b/cpp/src/search/contains_nested.cu @@ -37,7 +37,7 @@ bool contains_nested_element(column_view const& haystack, auto const comparator = cudf::experimental::row::equality::two_table_comparator(haystack_tv, needle_tv, stream); - auto const d_comp = comparator.device_comparator(nullate::DYNAMIC{has_nulls}); + auto const d_comp = comparator.equal_to(nullate::DYNAMIC{has_nulls}); auto const begin = cudf::experimental::row::lhs_iterator(0); auto const end = begin + haystack.size(); diff --git a/cpp/src/stream_compaction/distinct.cu b/cpp/src/stream_compaction/distinct.cu index d3b31dccf77..91d8c85120f 100644 --- a/cpp/src/stream_compaction/distinct.cu +++ b/cpp/src/stream_compaction/distinct.cu @@ -73,7 +73,7 @@ std::unique_ptr
distinct(table_view const& input, experimental::compaction_hash hash_key(row_hash.device_hasher(has_null)); cudf::experimental::row::equality::self_comparator row_equal(preprocessed_keys); - auto key_equal = row_equal.device_comparator(has_null, nulls_equal); + auto key_equal = row_equal.equal_to(has_null, nulls_equal); auto iter = cudf::detail::make_counting_transform_iterator( 0, [] __device__(size_type i) { return cuco::make_pair(i, i); }); diff --git a/cpp/tests/groupby/lists_tests.cu b/cpp/tests/groupby/lists_tests.cu index 7c145271662..81322d87747 100644 --- a/cpp/tests/groupby/lists_tests.cu +++ b/cpp/tests/groupby/lists_tests.cu @@ -118,7 +118,7 @@ inline void test_hash_based_sum_agg(column_view const& keys, auto const null_keys_are_equal = include_null_keys == null_policy::INCLUDE ? null_equality::EQUAL : null_equality::UNEQUAL; - auto row_equal = comparator.device_comparator(nullate::DYNAMIC{true}, null_keys_are_equal); + auto row_equal = comparator.equal_to(nullate::DYNAMIC{true}, null_keys_are_equal); auto func = match_expected_fn{num_rows, row_equal}; // For each row in expected table `t[0, num_rows)`, there must be a match diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 7233538de6f..6f1f43229e3 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -96,8 +96,7 @@ void row_equality(cudf::table_view input1, rmm::cuda_stream_view stream{}; auto table_comparator = equality::two_table_comparator{input1, input2, stream}; - auto comparator = - table_comparator.device_comparator(cudf::nullate::NO{}, cudf::null_equality::EQUAL, c); + auto comparator = table_comparator.equal_to(cudf::nullate::NO{}, cudf::null_equality::EQUAL, c); auto const lhs_it = cudf::experimental::row::lhs_iterator(0); auto const rhs_it = cudf::experimental::row::rhs_iterator(0); thrust::transform(rmm::exec_policy(stream), @@ -117,8 +116,7 @@ void self_equality(cudf::table_view input, rmm::cuda_stream_view stream{}; auto table_comparator = equality::self_comparator{input, stream}; - auto comparator = - table_comparator.device_comparator(cudf::nullate::NO{}, cudf::null_equality::EQUAL, c); + auto comparator = table_comparator.equal_to(cudf::nullate::NO{}, cudf::null_equality::EQUAL, c); thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.num_rows()), From 7293b6178085a4e73f18891fc09a2830b90c09ff Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 07:51:46 -0500 Subject: [PATCH 21/31] Fix docstring. --- cpp/include/cudf/table/experimental/row_operators.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index c5b99880e24..e78cce5325c 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -146,8 +146,8 @@ namespace lexicographic { * @brief Computes a weak ordering of two values with special sorting behavior. * * This relational comparator functor compares physical values rather than logical - * elements like lists, strings, or structs. It evaluates `NaN` as equivalent to other `NaN`s and - * greater than all other values. + * elements like lists, strings, or structs. It evaluates `NaN` as not less than, equal to, or + * greater than other values and is IEEE-754 compliant. */ struct physical_element_comparator { /** From b42c8959ab20d03b657896442ff364cd60c630dd Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 07:52:47 -0500 Subject: [PATCH 22/31] Fix check_nulls parameter docstring. --- cpp/include/cudf/table/experimental/row_operators.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index e78cce5325c..f9eef432415 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -233,7 +233,7 @@ class device_row_comparator { * @brief Construct a function object for performing a lexicographic * comparison between the rows of two tables. * - * @param check_nulls Indicates if any input column contains columns with nulls. + * @param check_nulls Indicates if any input column contains nulls. * @param lhs The first table * @param rhs The second table (may be the same table as `lhs`) * @param depth Optional, device array the same length as a row that contains starting depths of @@ -907,7 +907,7 @@ class device_row_comparator { * @brief Construct a function object for performing equality comparison between the rows of two * tables. * - * @param check_nulls Indicates if any input column contains columns with nulls. + * @param check_nulls Indicates if any input column contains nulls. * @param lhs The first table * @param rhs The second table (may be the same table as `lhs`) * @param nulls_are_equal Indicates if two null elements are treated as equivalent From 25fcff2c5eed2c2c6d4ceac35770e8f3c1e285aa Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 07:55:01 -0500 Subject: [PATCH 23/31] Rename c to comparator. --- .../cudf/table/experimental/row_operators.cuh | 52 +++++++++---------- 1 file changed, 26 insertions(+), 26 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index f9eef432415..50d3a8d9fb0 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -252,14 +252,14 @@ class device_row_comparator { std::optional> depth = std::nullopt, std::optional> column_order = std::nullopt, std::optional> null_precedence = std::nullopt, - PhysicalElementComparator c = {}) noexcept + PhysicalElementComparator comparator = {}) noexcept : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, _depth{depth}, _column_order{column_order}, _null_precedence{null_precedence}, - _comparator{c} + _comparator{comparator} { } @@ -287,13 +287,13 @@ class device_row_comparator { column_device_view rhs, null_order null_precedence = null_order::BEFORE, int depth = 0, - PhysicalElementComparator c = {}) + PhysicalElementComparator comparator = {}) : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, _null_precedence{null_precedence}, _depth{depth}, - _comparator{c} + _comparator{comparator} { } @@ -460,8 +460,8 @@ struct weak_ordering_comparator_impl { */ template struct less_comparator : weak_ordering_comparator_impl { - less_comparator(Comparator const& c) - : weak_ordering_comparator_impl{c} + less_comparator(Comparator const& comparator) + : weak_ordering_comparator_impl{comparator} { } }; @@ -469,8 +469,8 @@ struct less_comparator : weak_ordering_comparator_impl struct less_equivalent_comparator : weak_ordering_comparator_impl { - less_equivalent_comparator(Comparator const& c) - : weak_ordering_comparator_impl{c} + less_equivalent_comparator(Comparator const& comparator) + : weak_ordering_comparator_impl{comparator} { } }; @@ -637,18 +637,18 @@ class self_comparator { */ template - auto less(Nullate nullate = {}, PhysicalElementComparator c = {}) const noexcept + auto less(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const noexcept { return less_comparator{device_row_comparator{ - nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), c}}; + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), comparator}}; } template - auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator c = {}) const noexcept + auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const noexcept { return less_equivalent_comparator{device_row_comparator{ - nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), c}}; + nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), comparator}}; } private: @@ -658,7 +658,7 @@ class self_comparator { // @cond template struct strong_index_comparator_adapter { - strong_index_comparator_adapter(Comparator const& c) : comparator{c} {} + strong_index_comparator_adapter(Comparator const& comparator) : comparator{comparator} {} __device__ constexpr weak_ordering operator()(lhs_index_type const lhs_index, rhs_index_type const rhs_index) const noexcept @@ -767,7 +767,7 @@ class two_table_comparator { */ template - auto less(Nullate nullate = {}, PhysicalElementComparator c = {}) const noexcept + auto less(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const noexcept { return less_comparator{ strong_index_comparator_adapter{device_row_comparator{nullate, @@ -776,12 +776,12 @@ class two_table_comparator { d_left_table->depths(), d_left_table->column_order(), d_left_table->null_precedence(), - c}}}; + comparator}}}; } template - auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator c = {}) const noexcept + auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const noexcept { return less_equivalent_comparator{ strong_index_comparator_adapter{device_row_comparator{nullate, @@ -790,7 +790,7 @@ class two_table_comparator { d_left_table->depths(), d_left_table->column_order(), d_left_table->null_precedence(), - c}}}; + comparator}}}; } private: @@ -917,8 +917,8 @@ class device_row_comparator { table_device_view lhs, table_device_view rhs, null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator c = {}) noexcept - : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal}, comparator{c} + PhysicalEqualityComparator comparator = {}) noexcept + : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal}, comparator{comparator} { } @@ -943,12 +943,12 @@ class device_row_comparator { column_device_view lhs, column_device_view rhs, null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator c = {}) noexcept + PhysicalEqualityComparator comparator = {}) noexcept : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal}, - comparator{c} + comparator{comparator} { } @@ -1179,9 +1179,9 @@ class self_comparator { typename PhysicalEqualityComparator = nan_equal_physical_equality_comparator> auto equal_to(Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator c = {}) const noexcept + PhysicalEqualityComparator comparator = {}) const noexcept { - return device_row_comparator{nullate, *d_t, *d_t, nulls_are_equal, c}; + return device_row_comparator{nullate, *d_t, *d_t, nulls_are_equal, comparator}; } private: @@ -1191,7 +1191,7 @@ class self_comparator { // @cond template struct strong_index_comparator_adapter { - strong_index_comparator_adapter(Comparator const& c) : comparator{c} {} + strong_index_comparator_adapter(Comparator const& comparator) : comparator{comparator} {} __device__ constexpr bool operator()(lhs_index_type const lhs_index, rhs_index_type const rhs_index) const noexcept @@ -1281,10 +1281,10 @@ class two_table_comparator { typename PhysicalEqualityComparator = nan_equal_physical_equality_comparator> auto equal_to(Nullate nullate = {}, null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator c = {}) const noexcept + PhysicalEqualityComparator comparator = {}) const noexcept { return strong_index_comparator_adapter{ - device_row_comparator(nullate, *d_left_table, *d_right_table, nulls_are_equal, c)}; + device_row_comparator(nullate, *d_left_table, *d_right_table, nulls_are_equal, comparator)}; } private: From 20e498a12b0a87e9894fd567a611977aaa48c8c1 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 07:57:05 -0500 Subject: [PATCH 24/31] List members in initialization order. --- cpp/include/cudf/table/experimental/row_operators.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 50d3a8d9fb0..1e704c2da57 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -370,9 +370,9 @@ class device_row_comparator { column_device_view const _lhs; column_device_view const _rhs; Nullate const _check_nulls; - PhysicalElementComparator const _comparator; null_order const _null_precedence; int const _depth; + PhysicalElementComparator const _comparator; }; public: @@ -420,10 +420,10 @@ class device_row_comparator { table_device_view const _lhs; table_device_view const _rhs; Nullate const _check_nulls; - PhysicalElementComparator const _comparator; std::optional> const _depth; std::optional> const _column_order; std::optional> const _null_precedence; + PhysicalElementComparator const _comparator; }; // class device_row_comparator /** From 1eb657ec9deea8f762fa39916529071e2a190ec6 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 08:08:05 -0500 Subject: [PATCH 25/31] Style. --- .../cudf/table/experimental/row_operators.cuh | 41 +++++++++++-------- 1 file changed, 24 insertions(+), 17 deletions(-) diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh index 1e704c2da57..2397f9bc6fb 100644 --- a/cpp/include/cudf/table/experimental/row_operators.cuh +++ b/cpp/include/cudf/table/experimental/row_operators.cuh @@ -252,7 +252,7 @@ class device_row_comparator { std::optional> depth = std::nullopt, std::optional> column_order = std::nullopt, std::optional> null_precedence = std::nullopt, - PhysicalElementComparator comparator = {}) noexcept + PhysicalElementComparator comparator = {}) noexcept : _lhs{lhs}, _rhs{rhs}, _check_nulls{check_nulls}, @@ -285,8 +285,8 @@ class device_row_comparator { __device__ element_comparator(Nullate check_nulls, column_device_view lhs, column_device_view rhs, - null_order null_precedence = null_order::BEFORE, - int depth = 0, + null_order null_precedence = null_order::BEFORE, + int depth = 0, PhysicalElementComparator comparator = {}) : _lhs{lhs}, _rhs{rhs}, @@ -470,7 +470,8 @@ template struct less_equivalent_comparator : weak_ordering_comparator_impl { less_equivalent_comparator(Comparator const& comparator) - : weak_ordering_comparator_impl{comparator} + : weak_ordering_comparator_impl{ + comparator} { } }; @@ -645,7 +646,8 @@ class self_comparator { template - auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const noexcept + auto less_equivalent(Nullate nullate = {}, + PhysicalElementComparator comparator = {}) const noexcept { return less_equivalent_comparator{device_row_comparator{ nullate, *d_t, *d_t, d_t->depths(), d_t->column_order(), d_t->null_precedence(), comparator}}; @@ -781,7 +783,8 @@ class two_table_comparator { template - auto less_equivalent(Nullate nullate = {}, PhysicalElementComparator comparator = {}) const noexcept + auto less_equivalent(Nullate nullate = {}, + PhysicalElementComparator comparator = {}) const noexcept { return less_equivalent_comparator{ strong_index_comparator_adapter{device_row_comparator{nullate, @@ -916,9 +919,13 @@ class device_row_comparator { device_row_comparator(Nullate check_nulls, table_device_view lhs, table_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator comparator = {}) noexcept - : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, nulls_are_equal{nulls_are_equal}, comparator{comparator} + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator comparator = {}) noexcept + : lhs{lhs}, + rhs{rhs}, + check_nulls{check_nulls}, + nulls_are_equal{nulls_are_equal}, + comparator{comparator} { } @@ -942,8 +949,8 @@ class device_row_comparator { __device__ element_comparator(Nullate check_nulls, column_device_view lhs, column_device_view rhs, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator comparator = {}) noexcept + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator comparator = {}) noexcept : lhs{lhs}, rhs{rhs}, check_nulls{check_nulls}, @@ -1177,9 +1184,9 @@ class self_comparator { */ template - auto equal_to(Nullate nullate = {}, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator comparator = {}) const noexcept + auto equal_to(Nullate nullate = {}, + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator comparator = {}) const noexcept { return device_row_comparator{nullate, *d_t, *d_t, nulls_are_equal, comparator}; } @@ -1279,9 +1286,9 @@ class two_table_comparator { */ template - auto equal_to(Nullate nullate = {}, - null_equality nulls_are_equal = null_equality::EQUAL, - PhysicalEqualityComparator comparator = {}) const noexcept + auto equal_to(Nullate nullate = {}, + null_equality nulls_are_equal = null_equality::EQUAL, + PhysicalEqualityComparator comparator = {}) const noexcept { return strong_index_comparator_adapter{ device_row_comparator(nullate, *d_left_table, *d_right_table, nulls_are_equal, comparator)}; From f80cdc34f28e45e7e2de6ecc53a9a016c1a1a40f Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 08:08:45 -0500 Subject: [PATCH 26/31] Rename to nan_equal_expected. --- cpp/tests/table/experimental_row_operator_tests.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 6f1f43229e3..5b1dd7e6d5a 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -290,11 +290,11 @@ TYPED_TEST(NaNTableViewTest, TestEqualityComparatorTwoTableNaNCase) equality::physical_equality_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - cudf::test::fixed_width_column_wrapper sorting_expected{{1, 0, 0, 1}}; + cudf::test::fixed_width_column_wrapper nan_equal_expected{{1, 0, 0, 1}}; row_equality(input_table_1, input_table_2, got->mutable_view(), column_order, equality::nan_equal_physical_equality_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, got->view()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(nan_equal_expected, got->view()); } From 5cdd94a1f9083895f4c14f25574dd0ac042485dc Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 08:09:17 -0500 Subject: [PATCH 27/31] Include . --- cpp/tests/table/experimental_row_operator_tests.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 5b1dd7e6d5a..2c13c1a9622 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -32,6 +32,7 @@ #include #include +#include #include using namespace cudf::test; From 3a7f118ff92528257ebd7415b49286558f0385d8 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 08:53:34 -0500 Subject: [PATCH 28/31] Clean up test file. --- .../table/experimental_row_operator_tests.cu | 291 +++++++----------- 1 file changed, 113 insertions(+), 178 deletions(-) diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 2c13c1a9622..0102b4247f7 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -38,8 +38,6 @@ using namespace cudf::test; using namespace cudf::experimental::row; -// Compares two tables row by row, if table1 row is less than table2, then corresponding row value -// in `output` would be `true`/1 else `false`/0. template struct TypedTableViewTest : public cudf::test::BaseFixture { }; @@ -47,189 +45,141 @@ struct TypedTableViewTest : public cudf::test::BaseFixture { using NumericTypesNotBool = Concat; TYPED_TEST_SUITE(TypedTableViewTest, NumericTypesNotBool); -template -void row_comparison(cudf::table_view input1, - cudf::table_view input2, - cudf::mutable_column_view output, - std::vector const& column_order, - Comparator c) +template +auto self_comparison(cudf::table_view input, + std::vector const& column_order, + PhysicalElementComparator comparator) { rmm::cuda_stream_view stream{}; - auto table_comparator = - lexicographic::two_table_comparator{input1, input2, column_order, {}, stream}; - auto comparator = table_comparator.less(cudf::nullate::NO{}, c); - auto const lhs_it = cudf::experimental::row::lhs_iterator(0); - auto const rhs_it = cudf::experimental::row::rhs_iterator(0); - thrust::transform(rmm::exec_policy(stream), - lhs_it, - lhs_it + input1.num_rows(), - rhs_it, - output.data(), - comparator); -} + auto const table_comparator = lexicographic::self_comparator{input, column_order, {}, stream}; + auto const less_comparator = table_comparator.less(cudf::nullate::NO{}, comparator); -template -void self_comparison(cudf::table_view input, - cudf::mutable_column_view output, - std::vector const& column_order, - Comparator c) -{ - rmm::cuda_stream_view stream{}; + auto output = cudf::make_numeric_column( + cudf::data_type(cudf::type_id::BOOL8), input.num_rows(), cudf::mask_state::UNALLOCATED); - auto table_comparator = lexicographic::self_comparator{input, column_order, {}, stream}; - auto comparator = table_comparator.less(cudf::nullate::NO{}, c); thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.num_rows()), thrust::make_counting_iterator(0), - output.data(), - comparator); + output->mutable_view().data(), + less_comparator); + return output; } -template -void row_equality(cudf::table_view input1, - cudf::table_view input2, - cudf::mutable_column_view output, - std::vector const& column_order, - Comparator c) +template +auto two_table_comparison(cudf::table_view lhs, + cudf::table_view rhs, + std::vector const& column_order, + PhysicalElementComparator comparator) { rmm::cuda_stream_view stream{}; - auto table_comparator = equality::two_table_comparator{input1, input2, stream}; - auto comparator = table_comparator.equal_to(cudf::nullate::NO{}, cudf::null_equality::EQUAL, c); - auto const lhs_it = cudf::experimental::row::lhs_iterator(0); - auto const rhs_it = cudf::experimental::row::rhs_iterator(0); + auto const table_comparator = + lexicographic::two_table_comparator{lhs, rhs, column_order, {}, stream}; + auto const less_comparator = table_comparator.less(cudf::nullate::NO{}, comparator); + auto const lhs_it = cudf::experimental::row::lhs_iterator(0); + auto const rhs_it = cudf::experimental::row::rhs_iterator(0); + + auto output = cudf::make_numeric_column( + cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED); + thrust::transform(rmm::exec_policy(stream), lhs_it, - lhs_it + input1.num_rows(), + lhs_it + lhs.num_rows(), rhs_it, - output.data(), - comparator); + output->mutable_view().data(), + less_comparator); + return output; } -template -void self_equality(cudf::table_view input, - cudf::mutable_column_view output, +template +auto self_equality(cudf::table_view input, std::vector const& column_order, - Comparator c) + PhysicalElementComparator comparator) { rmm::cuda_stream_view stream{}; - auto table_comparator = equality::self_comparator{input, stream}; - auto comparator = table_comparator.equal_to(cudf::nullate::NO{}, cudf::null_equality::EQUAL, c); + auto const table_comparator = equality::self_comparator{input, stream}; + auto const equal_comparator = + table_comparator.equal_to(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator); + + auto output = cudf::make_numeric_column( + cudf::data_type(cudf::type_id::BOOL8), input.num_rows(), cudf::mask_state::UNALLOCATED); + thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(input.num_rows()), thrust::make_counting_iterator(0), - output.data(), - comparator); + output->mutable_view().data(), + equal_comparator); + return output; } -TYPED_TEST(TypedTableViewTest, EmptyColumnedTable) +template +auto two_table_equality(cudf::table_view lhs, + cudf::table_view rhs, + std::vector const& column_order, + PhysicalElementComparator comparator) { - std::vector cols{}; + rmm::cuda_stream_view stream{}; + + auto const table_comparator = equality::two_table_comparator{lhs, rhs, stream}; + auto const equal_comparator = + table_comparator.equal_to(cudf::nullate::NO{}, cudf::null_equality::EQUAL, comparator); + auto const lhs_it = cudf::experimental::row::lhs_iterator(0); + auto const rhs_it = cudf::experimental::row::rhs_iterator(0); - cudf::table_view input(cols); - cudf::size_type expected = 0; + auto output = cudf::make_numeric_column( + cudf::data_type(cudf::type_id::BOOL8), lhs.num_rows(), cudf::mask_state::UNALLOCATED); - EXPECT_EQ(input.num_columns(), expected); + thrust::transform(rmm::exec_policy(stream), + lhs_it, + lhs_it + lhs.num_rows(), + rhs_it, + output->mutable_view().data(), + equal_comparator); + return output; } -TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTableCase) +TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTables) { using T = TypeParam; - fixed_width_column_wrapper col1{{1, 2, 3, 4}}; - fixed_width_column_wrapper col2{{0, 1, 4, 3}}; - std::vector column_order{cudf::order::DESCENDING}; - - cudf::table_view input_table_1{{col1}}; - cudf::table_view input_table_2{{col2}}; - - auto got = cudf::make_numeric_column( - cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); - fixed_width_column_wrapper expected{{1, 1, 0, 1}}; + auto const col1 = fixed_width_column_wrapper{{1, 2, 3, 4}}; + auto const col2 = fixed_width_column_wrapper{{0, 1, 4, 3}}; + auto const column_order = std::vector{cudf::order::DESCENDING}; + auto const lhs = cudf::table_view{{col1}}; + auto const rhs = cudf::table_view{{col2}}; - row_comparison(input_table_1, - input_table_2, - got->mutable_view(), - column_order, - lexicographic::physical_element_comparator{}); + auto const expected = fixed_width_column_wrapper{{1, 1, 0, 1}}; + auto const got = two_table_comparison( + lhs, rhs, column_order, lexicographic::physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - row_comparison(input_table_1, - input_table_2, - got->mutable_view(), - column_order, - lexicographic::sorting_physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); + auto const sorting_got = two_table_comparison(lhs, + rhs, + column_order, + lexicographic::sorting_physical_element_comparator{}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, sorting_got->view()); } TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorSameTable) { using T = TypeParam; - fixed_width_column_wrapper col1{{1, 2, 3, 4}}; - std::vector column_order{cudf::order::DESCENDING}; - - cudf::table_view input_table_1{{col1}}; - - auto got = cudf::make_numeric_column( - cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); - fixed_width_column_wrapper expected{{0, 0, 0, 0}}; + auto const col1 = fixed_width_column_wrapper{{1, 2, 3, 4}}; + auto const column_order = std::vector{cudf::order::DESCENDING}; + auto const input_table = cudf::table_view{{col1}}; - self_comparison( - input_table_1, got->mutable_view(), column_order, lexicographic::physical_element_comparator{}); + auto const expected = fixed_width_column_wrapper{{0, 0, 0, 0}}; + auto const got = + self_comparison(input_table, column_order, lexicographic::physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - self_comparison(input_table_1, - got->mutable_view(), - column_order, - lexicographic::sorting_physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); -} - -TYPED_TEST(TypedTableViewTest, Select) -{ - using T = TypeParam; - - fixed_width_column_wrapper col1{{1, 2, 3, 4}}; - fixed_width_column_wrapper col2{{1, 2, 3, 4}}; - fixed_width_column_wrapper col3{{4, 5, 6, 7}}; - fixed_width_column_wrapper col4{{4, 5, 6, 7}}; - cudf::table_view t{{col1, col2, col3, col4}}; - - cudf::table_view selected = t.select({2, 3}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(t.column(2), selected.column(0)); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(t.column(3), selected.column(1)); -} - -TYPED_TEST(TypedTableViewTest, SelectOutOfBounds) -{ - using T = TypeParam; - - fixed_width_column_wrapper col1{{1, 2, 3, 4}}; - fixed_width_column_wrapper col2{{1, 2, 3, 4}}; - fixed_width_column_wrapper col3{{4, 5, 6, 7}}; - fixed_width_column_wrapper col4{{4, 5, 6, 7}}; - cudf::table_view t{{col1, col2}}; - - EXPECT_THROW((void)t.select({2, 3, 4}), std::out_of_range); -} - -TYPED_TEST(TypedTableViewTest, SelectNoColumns) -{ - using T = TypeParam; - - fixed_width_column_wrapper col1{{1, 2, 3, 4}}; - fixed_width_column_wrapper col2{{1, 2, 3, 4}}; - fixed_width_column_wrapper col3{{4, 5, 6, 7}}; - fixed_width_column_wrapper col4{{4, 5, 6, 7}}; - cudf::table_view t{{col1, col2, col3, col4}}; - - cudf::table_view selected = t.select({}); - EXPECT_EQ(selected.num_columns(), 0); + auto const sorting_got = self_comparison( + input_table, column_order, lexicographic::sorting_physical_element_comparator{}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, sorting_got->view()); } template @@ -242,60 +192,45 @@ TYPED_TEST(NaNTableViewTest, TestLexicographicalComparatorTwoTableNaNCase) { using T = TypeParam; - cudf::test::fixed_width_column_wrapper col1{{T(NAN), T(NAN), T(1), T(1)}}; - cudf::test::fixed_width_column_wrapper col2{{T(NAN), T(1), T(NAN), T(1)}}; - std::vector column_order{cudf::order::DESCENDING}; - - cudf::table_view input_table_1{{col1}}; - cudf::table_view input_table_2{{col2}}; + auto const col1 = fixed_width_column_wrapper{{T(NAN), T(NAN), T(1), T(1)}}; + auto const col2 = fixed_width_column_wrapper{{T(NAN), T(1), T(NAN), T(1)}}; + auto const column_order = std::vector{cudf::order::DESCENDING}; - auto got = cudf::make_numeric_column( - cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); + auto const lhs = cudf::table_view{{col1}}; + auto const rhs = cudf::table_view{{col2}}; - cudf::test::fixed_width_column_wrapper expected{{0, 0, 0, 0}}; - row_comparison(input_table_1, - input_table_2, - got->mutable_view(), - column_order, - lexicographic::physical_element_comparator{}); + auto const expected = fixed_width_column_wrapper{{0, 0, 0, 0}}; + auto const got = two_table_comparison( + lhs, rhs, column_order, lexicographic::physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - cudf::test::fixed_width_column_wrapper sorting_expected{{0, 1, 0, 0}}; - row_comparison(input_table_1, - input_table_2, - got->mutable_view(), - column_order, - lexicographic::sorting_physical_element_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, got->view()); + auto const sorting_expected = fixed_width_column_wrapper{{0, 1, 0, 0}}; + auto const sorting_got = + two_table_comparison(lhs, + rhs, + column_order, + lexicographic::sorting_physical_element_comparator{}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, sorting_got->view()); } TYPED_TEST(NaNTableViewTest, TestEqualityComparatorTwoTableNaNCase) { using T = TypeParam; - cudf::test::fixed_width_column_wrapper col1{{T(NAN), T(NAN), T(1), T(1)}}; - cudf::test::fixed_width_column_wrapper col2{{T(NAN), T(1), T(NAN), T(1)}}; - std::vector column_order{cudf::order::DESCENDING}; - - cudf::table_view input_table_1{{col1}}; - cudf::table_view input_table_2{{col2}}; + auto const col1 = fixed_width_column_wrapper{{T(NAN), T(NAN), T(1), T(1)}}; + auto const col2 = fixed_width_column_wrapper{{T(NAN), T(1), T(NAN), T(1)}}; + auto const column_order = std::vector{cudf::order::DESCENDING}; - auto got = cudf::make_numeric_column( - cudf::data_type(cudf::type_id::INT8), input_table_1.num_rows(), cudf::mask_state::UNALLOCATED); + auto const lhs = cudf::table_view{{col1}}; + auto const rhs = cudf::table_view{{col2}}; - cudf::test::fixed_width_column_wrapper expected{{0, 0, 0, 1}}; - row_equality(input_table_1, - input_table_2, - got->mutable_view(), - column_order, - equality::physical_equality_comparator{}); + auto const expected = fixed_width_column_wrapper{{0, 0, 0, 1}}; + auto const got = + two_table_equality(lhs, rhs, column_order, equality::physical_equality_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - cudf::test::fixed_width_column_wrapper nan_equal_expected{{1, 0, 0, 1}}; - row_equality(input_table_1, - input_table_2, - got->mutable_view(), - column_order, - equality::nan_equal_physical_equality_comparator{}); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(nan_equal_expected, got->view()); + auto const nan_equal_expected = fixed_width_column_wrapper{{1, 0, 0, 1}}; + auto const nan_equal_got = two_table_equality( + lhs, rhs, column_order, equality::nan_equal_physical_equality_comparator{}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(nan_equal_expected, nan_equal_got->view()); } From a1391319402ea9e6be6545df58dc8dbce654345e Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 08:54:19 -0500 Subject: [PATCH 29/31] Move tests to cpp file. --- cpp/tests/CMakeLists.txt | 2 +- ...ow_operator_tests.cu => experimental_row_operator_tests.cpp} | 0 2 files changed, 1 insertion(+), 1 deletion(-) rename cpp/tests/table/{experimental_row_operator_tests.cu => experimental_row_operator_tests.cpp} (100%) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 816c5a1c59c..319cece357d 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -310,7 +310,7 @@ ConfigureTest(TRANSPOSE_TEST transpose/transpose_test.cpp) # * table tests ----------------------------------------------------------------------------------- ConfigureTest( TABLE_TEST table/table_tests.cpp table/table_view_tests.cu table/row_operators_tests.cpp - table/experimental_row_operator_tests.cu + table/experimental_row_operator_tests.cpp ) # ################################################################################################## diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cpp similarity index 100% rename from cpp/tests/table/experimental_row_operator_tests.cu rename to cpp/tests/table/experimental_row_operator_tests.cpp From ec519e9f4201621c13fa3780bf356d122e406fb0 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 07:12:19 -0700 Subject: [PATCH 30/31] Revert "Move tests to cpp file." This reverts commit a1391319402ea9e6be6545df58dc8dbce654345e. --- cpp/tests/CMakeLists.txt | 2 +- ...ow_operator_tests.cpp => experimental_row_operator_tests.cu} | 0 2 files changed, 1 insertion(+), 1 deletion(-) rename cpp/tests/table/{experimental_row_operator_tests.cpp => experimental_row_operator_tests.cu} (100%) diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 319cece357d..816c5a1c59c 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -310,7 +310,7 @@ ConfigureTest(TRANSPOSE_TEST transpose/transpose_test.cpp) # * table tests ----------------------------------------------------------------------------------- ConfigureTest( TABLE_TEST table/table_tests.cpp table/table_view_tests.cu table/row_operators_tests.cpp - table/experimental_row_operator_tests.cpp + table/experimental_row_operator_tests.cu ) # ################################################################################################## diff --git a/cpp/tests/table/experimental_row_operator_tests.cpp b/cpp/tests/table/experimental_row_operator_tests.cu similarity index 100% rename from cpp/tests/table/experimental_row_operator_tests.cpp rename to cpp/tests/table/experimental_row_operator_tests.cu From 6da4a4c01a26ad4252c4a23e5ae41cb098c3ae16 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 3 Jun 2022 07:15:23 -0700 Subject: [PATCH 31/31] Style. --- .../table/experimental_row_operator_tests.cu | 41 ++++++++----------- 1 file changed, 18 insertions(+), 23 deletions(-) diff --git a/cpp/tests/table/experimental_row_operator_tests.cu b/cpp/tests/table/experimental_row_operator_tests.cu index 0102b4247f7..6b392fe57fe 100644 --- a/cpp/tests/table/experimental_row_operator_tests.cu +++ b/cpp/tests/table/experimental_row_operator_tests.cu @@ -146,21 +146,19 @@ TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorTwoTables) { using T = TypeParam; - auto const col1 = fixed_width_column_wrapper{{1, 2, 3, 4}}; - auto const col2 = fixed_width_column_wrapper{{0, 1, 4, 3}}; + auto const col1 = fixed_width_column_wrapper{{1, 2, 3, 4}}; + auto const col2 = fixed_width_column_wrapper{{0, 1, 4, 3}}; auto const column_order = std::vector{cudf::order::DESCENDING}; - auto const lhs = cudf::table_view{{col1}}; - auto const rhs = cudf::table_view{{col2}}; + auto const lhs = cudf::table_view{{col1}}; + auto const rhs = cudf::table_view{{col2}}; auto const expected = fixed_width_column_wrapper{{1, 1, 0, 1}}; - auto const got = two_table_comparison( - lhs, rhs, column_order, lexicographic::physical_element_comparator{}); + auto const got = + two_table_comparison(lhs, rhs, column_order, lexicographic::physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); - auto const sorting_got = two_table_comparison(lhs, - rhs, - column_order, - lexicographic::sorting_physical_element_comparator{}); + auto const sorting_got = two_table_comparison( + lhs, rhs, column_order, lexicographic::sorting_physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, sorting_got->view()); } @@ -168,9 +166,9 @@ TYPED_TEST(TypedTableViewTest, TestLexicographicalComparatorSameTable) { using T = TypeParam; - auto const col1 = fixed_width_column_wrapper{{1, 2, 3, 4}}; + auto const col1 = fixed_width_column_wrapper{{1, 2, 3, 4}}; auto const column_order = std::vector{cudf::order::DESCENDING}; - auto const input_table = cudf::table_view{{col1}}; + auto const input_table = cudf::table_view{{col1}}; auto const expected = fixed_width_column_wrapper{{0, 0, 0, 0}}; auto const got = @@ -192,24 +190,21 @@ TYPED_TEST(NaNTableViewTest, TestLexicographicalComparatorTwoTableNaNCase) { using T = TypeParam; - auto const col1 = fixed_width_column_wrapper{{T(NAN), T(NAN), T(1), T(1)}}; - auto const col2 = fixed_width_column_wrapper{{T(NAN), T(1), T(NAN), T(1)}}; + auto const col1 = fixed_width_column_wrapper{{T(NAN), T(NAN), T(1), T(1)}}; + auto const col2 = fixed_width_column_wrapper{{T(NAN), T(1), T(NAN), T(1)}}; auto const column_order = std::vector{cudf::order::DESCENDING}; auto const lhs = cudf::table_view{{col1}}; auto const rhs = cudf::table_view{{col2}}; auto const expected = fixed_width_column_wrapper{{0, 0, 0, 0}}; - auto const got = two_table_comparison( - lhs, rhs, column_order, lexicographic::physical_element_comparator{}); + auto const got = + two_table_comparison(lhs, rhs, column_order, lexicographic::physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); auto const sorting_expected = fixed_width_column_wrapper{{0, 1, 0, 0}}; - auto const sorting_got = - two_table_comparison(lhs, - rhs, - column_order, - lexicographic::sorting_physical_element_comparator{}); + auto const sorting_got = two_table_comparison( + lhs, rhs, column_order, lexicographic::sorting_physical_element_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(sorting_expected, sorting_got->view()); } @@ -230,7 +225,7 @@ TYPED_TEST(NaNTableViewTest, TestEqualityComparatorTwoTableNaNCase) CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, got->view()); auto const nan_equal_expected = fixed_width_column_wrapper{{1, 0, 0, 1}}; - auto const nan_equal_got = two_table_equality( - lhs, rhs, column_order, equality::nan_equal_physical_equality_comparator{}); + auto const nan_equal_got = + two_table_equality(lhs, rhs, column_order, equality::nan_equal_physical_equality_comparator{}); CUDF_TEST_EXPECT_COLUMNS_EQUAL(nan_equal_expected, nan_equal_got->view()); }