20 #include <cudf/detail/utilities/assert.cuh>
21 #include <cudf/hashing/detail/hash_functions.cuh>
22 #include <cudf/hashing/detail/hashing.hpp>
28 #include <thrust/equal.h>
29 #include <thrust/execution_policy.h>
30 #include <thrust/iterator/counting_iterator.h>
31 #include <thrust/swap.h>
32 #include <thrust/transform_reduce.h>
61 template <
typename Element>
62 __device__
weak_ordering compare_elements(Element lhs, Element rhs)
65 return weak_ordering::LESS;
66 }
else if (rhs < lhs) {
69 return weak_ordering::EQUIVALENT;
82 template <
typename Element, std::enable_if_t<std::is_
floating_po
int_v<Element>>* =
nullptr>
85 if (isnan(lhs) and isnan(rhs)) {
86 return weak_ordering::EQUIVALENT;
87 }
else if (isnan(rhs)) {
88 return weak_ordering::LESS;
89 }
else if (isnan(lhs)) {
93 return detail::compare_elements(lhs, rhs);
106 if (lhs_is_null and rhs_is_null) {
107 return weak_ordering::EQUIVALENT;
108 }
else if (lhs_is_null) {
110 }
else if (rhs_is_null) {
113 return weak_ordering::EQUIVALENT;
124 template <
typename Element, std::enable_if_t<not std::is_
floating_po
int_v<Element>>* =
nullptr>
127 return detail::compare_elements(lhs, rhs);
138 template <
typename Element, std::enable_if_t<std::is_
floating_po
int_v<Element>>* =
nullptr>
141 if (isnan(lhs) and isnan(rhs)) {
return true; }
153 template <
typename Element, std::enable_if_t<not std::is_
floating_po
int_v<Element>>* =
nullptr>
164 template <
typename Nullate>
183 : lhs{lhs}, rhs{rhs}, nulls{
has_nulls}, nulls_are_equal{nulls_are_equal}
194 template <
typename Element,
195 std::enable_if_t<cudf::is_equality_comparable<Element, Element>()>* =
nullptr>
197 size_type rhs_element_index)
const noexcept
200 bool const lhs_is_null{lhs.is_null(lhs_element_index)};
201 bool const rhs_is_null{rhs.is_null(rhs_element_index)};
202 if (lhs_is_null and rhs_is_null) {
203 return nulls_are_equal == null_equality::EQUAL;
204 }
else if (lhs_is_null != rhs_is_null) {
210 rhs.element<Element>(rhs_element_index));
214 template <
typename Element,
215 std::enable_if_t<not cudf::is_equality_comparable<Element, Element>()>* =
nullptr>
218 CUDF_UNREACHABLE(
"Attempted to compare elements of uncomparable types.");
223 column_device_view lhs;
224 column_device_view rhs;
234 template <
typename Nullate>
249 : lhs{lhs}, rhs{rhs}, nulls{
has_nulls}, nulls_are_equal{nulls_are_equal}
251 CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(),
"Mismatched number of columns.");
270 return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements);
285 template <
typename Nullate>
303 : lhs{lhs}, rhs{rhs}, nulls{
has_nulls}, null_precedence{null_precedence}
330 template <
typename Element,
331 std::enable_if_t<cudf::is_relationally_comparable<Element, Element>()>* =
nullptr>
333 size_type rhs_element_index)
const noexcept
336 bool const lhs_is_null{lhs.is_null(lhs_element_index)};
337 bool const rhs_is_null{rhs.is_null(rhs_element_index)};
339 if (lhs_is_null or rhs_is_null) {
340 return null_compare(lhs_is_null, rhs_is_null, null_precedence);
345 rhs.element<Element>(rhs_element_index));
349 template <
typename Element,
350 std::enable_if_t<not cudf::is_relationally_comparable<Element, Element>()>* =
nullptr>
353 CUDF_UNREACHABLE(
"Attempted to compare elements of uncomparable types.");
358 column_device_view lhs;
359 column_device_view rhs;
379 template <
typename Nullate>
404 order const* column_order =
nullptr,
409 _column_order{column_order},
410 _null_precedence{null_precedence}
427 bool ascending = (_column_order ==
nullptr) or (_column_order[i] == order::ASCENDING);
438 if (state == weak_ordering::EQUIVALENT) {
continue; }
450 order const* _column_order{};
459 template <
template <
typename>
class hash_function,
typename Nullate>
470 template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
473 if (
has_nulls && col.
is_null(row_index)) {
return std::numeric_limits<hash_value_type>::max(); }
474 return hash_function<T>{}(col.
element<T>(row_index));
485 template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
488 CUDF_UNREACHABLE(
"Unsupported type in hash.");
500 template <
template <
typename>
class hash_function,
typename Nullate>
522 : _seed{seed}, _null_hash{null_hash}, _has_nulls{
has_nulls}
534 template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
537 if (_has_nulls && col.
is_null(row_index)) {
return _null_hash; }
538 return hash_function<T>{_seed}(col.
element<T>(row_index));
549 template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
552 CUDF_UNREACHABLE(
"Unsupported type in hash.");
556 uint32_t _seed{DEFAULT_HASH_SEED};
557 hash_value_type _null_hash{std::numeric_limits<hash_value_type>::max()};
567 template <
template <
typename>
class hash_function,
typename Nullate>
590 : _table{t}, _seed(seed), _has_nulls{
has_nulls}
603 auto const initial_hash = cudf::hashing::detail::hash_combine(
605 type_dispatcher<dispatch_storage_type>(
612 auto hasher = [=](
size_type column_index) {
613 return cudf::type_dispatcher<dispatch_storage_type>(
616 _table.
column(column_index),
621 return thrust::transform_reduce(
624 thrust::make_counting_iterator(1),
625 thrust::make_counting_iterator(_table.
num_columns()),
629 return cudf::hashing::detail::hash_combine(lhs, rhs);
636 uint32_t _seed{DEFAULT_HASH_SEED};