Make row hasher 64-bit hashing compatible (rapidsai#20777)

Closes: rapidsai#20757

This PR modifies the normal row hasher and the primitive row hasher to enable compatibility with 64-bit hashing.

Authors:
  - Yunsong Wang (https://github.com/PointKernel)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - David Wendt (https://github.com/davidwendt)
  - Muhammad Haseeb (https://github.com/mhaseeb123)

URL: rapidsai#20777

Author: PointKernel

cpp/include/cudf/detail/row_operator/hashing.cuh

@@ -22,6 +22,7 @@
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <cuda/std/limits>
+#include <cuda/std/type_traits>
 #include <thrust/iterator/transform_iterator.h>
 
 #include <memory>
@@ -38,6 +39,8 @@ namespace detail::row::hash {
 template <template <typename> class hash_function, typename Nullate>
 class element_hasher {
  public:
+  using result_type = cuda::std::invoke_result_t<hash_function<int32_t>, int32_t>;
+
   /**
    * @brief Constructs an element_hasher object.
    *
@@ -47,8 +50,8 @@ class element_hasher {
    */
   __device__ element_hasher(
     Nullate nulls,
-    uint32_t seed             = DEFAULT_HASH_SEED,
-    hash_value_type null_hash = cuda::std::numeric_limits<hash_value_type>::max()) noexcept
+    result_type seed      = DEFAULT_HASH_SEED,
+    result_type null_hash = cuda::std::numeric_limits<result_type>::max()) noexcept
     : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
   {
   }
@@ -62,8 +65,8 @@ class element_hasher {
    * @return The hash value of the given element
    */
   template <typename T>
-  __device__ hash_value_type operator()(column_device_view const& col,
-                                        size_type row_index) const noexcept
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
     requires(column_device_view::has_element_accessor<T>())
   {
     if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
@@ -79,16 +82,17 @@ class element_hasher {
    * @return The hash value of the given element
    */
   template <typename T>
-  __device__ hash_value_type operator()(column_device_view const& col,
-                                        size_type row_index) const noexcept
+  __device__ result_type operator()(column_device_view const& col,
+                                    size_type row_index) const noexcept
     requires(not column_device_view::has_element_accessor<T>())
   {
     CUDF_UNREACHABLE("Unsupported type in hash.");
   }
 
   Nullate _check_nulls;
-  uint32_t _seed;
-  hash_value_type _null_hash;
+  // Assumes seeds are the same as the result type of the hash function
+  result_type _seed;
+  result_type _null_hash;
 };
 
 /**
@@ -102,21 +106,20 @@ class device_row_hasher {
   friend class row_hasher;
 
  public:
+  using result_type = cuda::std::invoke_result_t<hash_function<int32_t>, int32_t>;
+
   /**
    * @brief Return the hash value of a row in the given table.
    *
    * @param row_index The row index to compute the hash value of
    * @return The hash value of the row
    */
-  __device__ auto operator()(size_type row_index) const noexcept
+  __device__ result_type operator()(size_type row_index) const noexcept
   {
     auto it =
       thrust::make_transform_iterator(_table.begin(), [row_index, this](auto const& column) {
         return cudf::type_dispatcher<dispatch_storage_type>(
-          column.type(),
-          element_hasher_adapter<hash_function>{_check_nulls, _seed},
-          column,
-          row_index);
+          column.type(), element_hasher_adapter{_check_nulls, _seed}, column, row_index);
       });
 
     return detail::accumulate(it, it + _table.num_columns(), _seed, [](auto hash, auto h) {
@@ -132,31 +135,30 @@ class device_row_hasher {
    * When the column is nested, this uses the element_hasher to hash the shape and values of the
    * column.
    */
-  template <template <typename> class hash_fn>
   class element_hasher_adapter {
-    static constexpr hash_value_type NULL_HASH = cuda::std::numeric_limits<hash_value_type>::max();
-    static constexpr hash_value_type NON_NULL_HASH = 0;
+    static constexpr result_type NULL_HASH     = cuda::std::numeric_limits<result_type>::max();
+    static constexpr result_type NON_NULL_HASH = 0;
 
    public:
-    __device__ element_hasher_adapter(Nullate check_nulls, uint32_t seed) noexcept
+    __device__ element_hasher_adapter(Nullate check_nulls, result_type seed) noexcept
       : _element_hasher(check_nulls, seed), _check_nulls(check_nulls)
     {
     }
 
     template <typename T>
-    __device__ hash_value_type operator()(column_device_view const& col,
-                                          size_type row_index) const noexcept
+    __device__ result_type operator()(column_device_view const& col,
+                                      size_type row_index) const noexcept
       requires(not cudf::is_nested<T>())
     {
       return _element_hasher.template operator()<T>(col, row_index);
     }
 
     template <typename T>
-    __device__ hash_value_type operator()(column_device_view const& col,
-                                          size_type row_index) const noexcept
+    __device__ result_type operator()(column_device_view const& col,
+                                      size_type row_index) const noexcept
       requires(cudf::is_nested<T>())
     {
-      auto hash                   = hash_value_type{0};
+      auto hash                   = result_type{0};
       column_device_view curr_col = col.slice(row_index, 1);
       while (curr_col.type().id() == type_id::STRUCT || curr_col.type().id() == type_id::LIST) {
         if (_check_nulls) {
@@ -175,7 +177,7 @@ class device_row_hasher {
           auto list_sizes = make_list_size_iterator(list_col);
           hash            = detail::accumulate(
             list_sizes, list_sizes + list_col.size(), hash, [](auto hash, auto size) {
-              return cudf::hashing::detail::hash_combine(hash, hash_fn<size_type>{}(size));
+              return cudf::hashing::detail::hash_combine(hash, hash_function<size_type>{}(size));
             });
           curr_col = list_col.get_sliced_child();
         }
@@ -188,20 +190,21 @@ class device_row_hasher {
       return hash;
     }
 
-    element_hasher<hash_fn, Nullate> const _element_hasher;
+    element_hasher<hash_function, Nullate> const _element_hasher;
     Nullate const _check_nulls;
   };
 
   CUDF_HOST_DEVICE device_row_hasher(Nullate check_nulls,
                                      table_device_view t,
-                                     uint32_t seed = DEFAULT_HASH_SEED) noexcept
+                                     result_type seed = DEFAULT_HASH_SEED) noexcept
     : _check_nulls{check_nulls}, _table{t}, _seed(seed)
   {
   }
 
   Nullate const _check_nulls;
   table_device_view const _table;
-  uint32_t const _seed;
+  // Assumes seeds are the same as the result type of the hash function
+  result_type const _seed;
 };
 
 /**
@@ -236,11 +239,12 @@ class row_hasher {
   /**
    * @brief Get the hash operator to use on the device
    *
-   * Returns a unary callable, `F`, with signature `hash_function::hash_value_type F(size_type)`.
-   *
-   * `F(i)` returns the hash of row i.
+   * Returns a unary callable, `F`, where `F(i)` returns the hash value of row i.
    *
+   * @tparam hash_function Hash functor to use for hashing elements
+   * @tparam DeviceRowHasher The device row hasher type to use
    * @tparam Nullate A cudf::nullate type describing whether to check for 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
@@ -249,8 +253,9 @@ class row_hasher {
     template <typename> class hash_function = cudf::hashing::detail::default_hash,
     template <template <typename> class, typename> class DeviceRowHasher = device_row_hasher,
     typename Nullate>
-  DeviceRowHasher<hash_function, Nullate> device_hasher(Nullate nullate = {},
-                                                        uint32_t seed   = DEFAULT_HASH_SEED) const
+  DeviceRowHasher<hash_function, Nullate> device_hasher(
+    Nullate nullate                                                  = {},
+    cuda::std::invoke_result_t<hash_function<int32_t>, int32_t> seed = DEFAULT_HASH_SEED) const
   {
     return DeviceRowHasher<hash_function, Nullate>(nullate, *d_t, seed);
   }

cpp/include/cudf/detail/row_operator/primitive_row_operators.cuh

@@ -170,6 +170,8 @@ class row_equality_comparator {
 template <template <typename> class Hash>
 class element_hasher {
  public:
+  using result_type = cuda::std::invoke_result_t<Hash<int32_t>, int32_t>;
+
   /**
    * @brief Returns the hash value of the given element in the given column.
    *
@@ -180,16 +182,16 @@ class element_hasher {
    * @return The hash value of the given element
    */
   template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(hash_value_type seed,
-                                        column_device_view const& col,
-                                        size_type row_index) const
+  __device__ result_type operator()(result_type seed,
+                                    column_device_view const& col,
+                                    size_type row_index) const
   {
     return Hash<T>{seed}(col.element<T>(row_index));
   }
 
   // @cond
   template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(hash_value_type, column_device_view const&, size_type) const
+  __device__ result_type operator()(result_type, column_device_view const&, size_type) const
   {
     CUDF_UNREACHABLE("Unsupported type in hash.");
   }
@@ -204,6 +206,8 @@ class element_hasher {
 template <template <typename> class Hash = cudf::hashing::detail::default_hash>
 class row_hasher {
  public:
+  using result_type = cuda::std::invoke_result_t<Hash<int32_t>, int32_t>;
+
   row_hasher() = delete;
 
   /**
@@ -215,7 +219,7 @@ class row_hasher {
    */
   row_hasher(cudf::nullate::DYNAMIC const& has_nulls,
              table_device_view t,
-             hash_value_type seed = DEFAULT_HASH_SEED)
+             result_type seed = DEFAULT_HASH_SEED)
     : _has_nulls{has_nulls}, _table{t}, _seed{seed}
   {
   }
@@ -229,7 +233,7 @@ class row_hasher {
    */
   row_hasher(cudf::nullate::DYNAMIC const& has_nulls,
              std::shared_ptr<cudf::detail::row::equality::preprocessed_table> t,
-             hash_value_type seed = DEFAULT_HASH_SEED)
+             result_type seed = DEFAULT_HASH_SEED)
     : _has_nulls{has_nulls}, _table{*t}, _seed{seed}
   {
   }
@@ -243,8 +247,7 @@ class row_hasher {
   __device__ auto operator()(size_type row_index) const
   {
     element_hasher<Hash> hasher;
-    // avoid hash combine call if there is only one column
-    auto hash = cuda::std::numeric_limits<hash_value_type>::max();
+    auto hash = cuda::std::numeric_limits<result_type>::max();
     if (!_has_nulls || !_table.column(0).is_null(row_index)) {
       hash = cudf::type_dispatcher<dispatch_primitive_type>(
         _table.column(0).type(), hasher, _seed, _table.column(0), row_index);
@@ -257,8 +260,8 @@ class row_hasher {
           cudf::type_dispatcher<dispatch_primitive_type>(
             _table.column(i).type(), hasher, _seed, _table.column(i), row_index));
       } else {
-        hash = cudf::hashing::detail::hash_combine(
-          hash, cuda::std::numeric_limits<hash_value_type>::max());
+        hash =
+          cudf::hashing::detail::hash_combine(hash, cuda::std::numeric_limits<result_type>::max());
       }
     }
     return hash;
@@ -267,7 +270,7 @@ class row_hasher {
  private:
   cudf::nullate::DYNAMIC _has_nulls;
   table_device_view _table;
-  hash_value_type _seed;
+  result_type _seed;
 };
 
 }  // namespace row::primitive

cpp/include/cudf/hashing/detail/hashing.hpp

@@ -11,6 +11,7 @@
 #include <rmm/cuda_stream_view.hpp>
 
 #include <cstddef>
+#include <cstdint>
 #include <functional>
 
 namespace CUDF_EXPORT cudf {
@@ -95,7 +96,7 @@ CUDF_HOST_DEVICE constexpr uint32_t hash_combine(uint32_t lhs, uint32_t rhs)
  * http://www.boost.org/LICENSE_1_0.txt)
  */
 /**
- * @brief Combines two hash values into a single hash value.
+ * @brief Combines two 64-bit hash values into a single hash value.
  *
  * Adapted from Boost hash_combine function and modified for 64-bit.
  * https://www.boost.org/doc/libs/1_35_0/doc/html/boost/hash_combine_id241013.html
@@ -104,7 +105,7 @@ CUDF_HOST_DEVICE constexpr uint32_t hash_combine(uint32_t lhs, uint32_t rhs)
  * @param rhs The second hash value
  * @return Combined hash value
  */
-constexpr std::size_t hash_combine(std::size_t lhs, std::size_t rhs)
+CUDF_HOST_DEVICE constexpr uint64_t hash_combine(uint64_t lhs, uint64_t rhs)
 {
   return lhs ^ (rhs + 0x9e37'79b9'7f4a'7c15 + (lhs << 6) + (lhs >> 2));
 }

cpp/tests/row_operator/row_operator_tests.cu

@@ -11,10 +11,17 @@
 #include <cudf_test/type_lists.hpp>
 
 #include <cudf/detail/row_operator/equality.cuh>
+#include <cudf/detail/row_operator/hashing.cuh>
 #include <cudf/detail/row_operator/lexicographic.cuh>
+#include <cudf/detail/row_operator/primitive_row_operators.cuh>
+#include <cudf/hashing/detail/xxhash_64.cuh>
 #include <cudf/strings/strings_column_view.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/transform.h>
 
 template <typename T>
 struct TypedTableViewTest : public cudf::test::BaseFixture {};
@@ -332,3 +339,51 @@ TEST_F(RowOperatorTest, TestCheckShapeCompatibility)
     cudf::detail::row::equality::two_table_comparator(string_table, numeric_table, stream),
     std::invalid_argument);
 }
+
+TEST_F(RowOperatorTest, TestRowHasher64BitHash)
+{
+  auto const col   = cudf::test::fixed_width_column_wrapper<int32_t>{{0, 42, 123456789}};
+  auto const input = cudf::table_view{{col}};
+
+  auto const stream       = cudf::get_default_stream();
+  auto const preprocessed = cudf::detail::row::hash::preprocessed_table::create(input, stream);
+  auto const row_hasher   = cudf::detail::row::hash::row_hasher{preprocessed};
+  auto const hasher =
+    row_hasher.device_hasher<cudf::hashing::detail::XXHash_64>(cudf::nullate::DYNAMIC{false});
+
+  auto results = cudf::test::fixed_width_column_wrapper<std::uint64_t>{{0, 0, 0}};
+  thrust::transform(rmm::exec_policy_nosync(stream),
+                    thrust::counting_iterator<cudf::size_type>{0},
+                    thrust::counting_iterator<cudf::size_type>{3},
+                    cudf::mutable_column_view{results}.begin<std::uint64_t>(),
+                    hasher);
+
+  // Expected values are hash_combine(seed=0, xxhash_64(element))
+  auto const expected = cudf::test::fixed_width_column_wrapper<std::uint64_t>{
+    {15647511400073222857ul, 8470797489250732038ul, 2416304890555758815ul}};
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(results, expected);
+}
+
+TEST_F(RowOperatorTest, TestPrimitiveRowHasher64BitHash)
+{
+  auto const col   = cudf::test::fixed_width_column_wrapper<int32_t>{{0, 42, 123456789}};
+  auto const input = cudf::table_view{{col}};
+
+  auto const stream  = cudf::get_default_stream();
+  auto const d_input = cudf::table_device_view::create(input, stream);
+
+  auto const hasher = cudf::detail::row::primitive::row_hasher<cudf::hashing::detail::XXHash_64>(
+    cudf::nullate::DYNAMIC{false}, *d_input, static_cast<std::uint64_t>(cudf::DEFAULT_HASH_SEED));
+
+  auto results = cudf::test::fixed_width_column_wrapper<std::uint64_t>{{0, 0, 0}};
+
+  thrust::transform(rmm::exec_policy_nosync(stream),
+                    thrust::counting_iterator<cudf::size_type>{0},
+                    thrust::counting_iterator<cudf::size_type>{3},
+                    cudf::mutable_column_view{results}.begin<std::uint64_t>(),
+                    hasher);
+
+  auto const expected = cudf::test::fixed_width_column_wrapper<std::uint64_t>{
+    {4246796580750024372ul, 15516826743637085169ul, 9462334144942111946ul}};
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(results, expected);
+}