20 #include <cudf/detail/hashing.hpp>
21 #include <cudf/detail/utilities/assert.cuh>
22 #include <cudf/detail/utilities/hash_functions.cuh>
28 #include <thrust/equal.h>
29 #include <thrust/swap.h>
30 #include <thrust/transform_reduce.h>
59 template <
typename Element>
60 __device__
weak_ordering compare_elements(Element lhs, Element rhs)
63 return weak_ordering::LESS;
64 }
else if (rhs < lhs) {
67 return weak_ordering::EQUIVALENT;
84 template <
typename Element, std::enable_if_t<std::is_
floating_po
int_v<Element>>* =
nullptr>
87 if (isnan(lhs) and isnan(rhs)) {
88 return weak_ordering::EQUIVALENT;
89 }
else if (isnan(rhs)) {
90 return weak_ordering::LESS;
91 }
else if (isnan(lhs)) {
95 return detail::compare_elements(lhs, rhs);
108 if (lhs_is_null and rhs_is_null) {
109 return weak_ordering::EQUIVALENT;
110 }
else if (lhs_is_null) {
112 }
else if (rhs_is_null) {
115 return weak_ordering::EQUIVALENT;
127 template <
typename Element, std::enable_if_t<not std::is_
floating_po
int_v<Element>>* =
nullptr>
130 return detail::compare_elements(lhs, rhs);
141 template <
typename Element, std::enable_if_t<std::is_
floating_po
int_v<Element>>* =
nullptr>
144 if (isnan(lhs) and isnan(rhs)) {
return true; }
156 template <
typename Element, std::enable_if_t<not std::is_
floating_po
int_v<Element>>* =
nullptr>
167 template <
typename Nullate>
186 : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal}
197 template <
typename Element,
198 std::enable_if_t<cudf::is_equality_comparable<Element, Element>()>* =
nullptr>
200 size_type rhs_element_index)
const noexcept
203 bool const lhs_is_null{lhs.is_null(lhs_element_index)};
204 bool const rhs_is_null{rhs.is_null(rhs_element_index)};
205 if (lhs_is_null and rhs_is_null) {
206 return nulls_are_equal == null_equality::EQUAL;
207 }
else if (lhs_is_null != rhs_is_null) {
213 rhs.element<Element>(rhs_element_index));
216 template <
typename Element,
217 std::enable_if_t<not cudf::is_equality_comparable<Element, Element>()>* =
nullptr>
218 __device__
bool operator()(size_type lhs_element_index, size_type rhs_element_index)
220 CUDF_UNREACHABLE(
"Attempted to compare elements of uncomparable types.");
224 column_device_view lhs;
225 column_device_view rhs;
230 template <
typename Nullate>
237 : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal}
239 CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(),
"Mismatched number of columns.");
242 __device__
bool operator()(size_type lhs_row_index, size_type rhs_row_index)
const noexcept
251 return thrust::equal(thrust::seq, lhs.begin(), lhs.end(), rhs.begin(), equal_elements);
266 template <
typename Nullate>
284 : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, null_precedence{null_precedence}
289 column_device_view lhs,
290 column_device_view rhs)
291 : lhs{lhs}, rhs{rhs}, nulls{has_nulls}
303 template <
typename Element,
304 std::enable_if_t<cudf::is_relationally_comparable<Element, Element>()>* =
nullptr>
306 size_type rhs_element_index)
const noexcept
309 bool const lhs_is_null{lhs.is_null(lhs_element_index)};
310 bool const rhs_is_null{rhs.is_null(rhs_element_index)};
312 if (lhs_is_null or rhs_is_null) {
313 return null_compare(lhs_is_null, rhs_is_null, null_precedence);
318 rhs.element<Element>(rhs_element_index));
321 template <
typename Element,
322 std::enable_if_t<not cudf::is_relationally_comparable<Element, Element>()>* =
nullptr>
325 CUDF_UNREACHABLE(
"Attempted to compare elements of uncomparable types.");
329 column_device_view lhs;
330 column_device_view rhs;
350 template <
typename Nullate>
375 order const* column_order =
nullptr,
380 _column_order{column_order},
381 _null_precedence{null_precedence}
383 CUDF_EXPECTS(_lhs.num_columns() == _rhs.num_columns(),
"Mismatched number of columns.");
395 __device__
bool operator()(size_type lhs_index, size_type rhs_index)
const noexcept
397 for (size_type i = 0; i < _lhs.num_columns(); ++i) {
398 bool ascending = (_column_order ==
nullptr) or (_column_order[i] == order::ASCENDING);
409 if (state == weak_ordering::EQUIVALENT) {
continue; }
421 order const* _column_order{};
430 template <
template <
typename>
class hash_function,
typename Nullate>
433 template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
434 __device__ hash_value_type
operator()(
column_device_view col, size_type row_index)
const
436 if (has_nulls && col.
is_null(row_index)) {
return std::numeric_limits<hash_value_type>::max(); }
437 return hash_function<T>{}(col.
element<T>(row_index));
440 template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
441 __device__ hash_value_type
operator()(
column_device_view col, size_type row_index)
const
443 CUDF_UNREACHABLE(
"Unsupported type in hash.");
449 template <
template <
typename>
class hash_function,
typename Nullate>
453 : _seed{seed}, _has_nulls{has_nulls}
458 : _seed{seed}, _null_hash{null_hash}, _has_nulls{has_nulls}
462 template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
463 __device__ hash_value_type
operator()(
column_device_view col, size_type row_index)
const
465 if (_has_nulls && col.
is_null(row_index)) {
return _null_hash; }
466 return hash_function<T>{_seed}(col.
element<T>(row_index));
469 template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
470 __device__ hash_value_type
operator()(
column_device_view col, size_type row_index)
const
472 CUDF_UNREACHABLE(
"Unsupported type in hash.");
476 uint32_t _seed{DEFAULT_HASH_SEED};
477 hash_value_type _null_hash{std::numeric_limits<hash_value_type>::max()};
487 template <
template <
typename>
class hash_function,
typename Nullate>
492 : _table{t}, _has_nulls{has_nulls}
496 : _table{t}, _seed(seed), _has_nulls{has_nulls}
500 __device__
auto operator()(size_type row_index)
const
503 auto const initial_hash = cudf::detail::hash_combine(
505 type_dispatcher<dispatch_storage_type>(
506 _table.column(0).
type(),
512 auto hasher = [=](size_type column_index) {
513 return cudf::type_dispatcher<dispatch_storage_type>(
514 _table.column(column_index).
type(),
516 _table.column(column_index),
521 return thrust::transform_reduce(
524 thrust::make_counting_iterator(1),
525 thrust::make_counting_iterator(_table.num_columns()),
528 [](hash_value_type lhs, hash_value_type rhs) {
529 return cudf::detail::hash_combine(lhs, rhs);
536 uint32_t _seed{DEFAULT_HASH_SEED};