Add XXHash_64 hash function to cudf

cpp/CMakeLists.txt

@@ -346,6 +346,7 @@ add_library(
   src/hash/md5_hash.cu
   src/hash/murmurhash3_x86_32.cu
   src/hash/spark_murmurhash3_x86_32.cu
+  src/hash/xxhash64.cu
   src/interop/dlpack.cpp
   src/interop/from_arrow.cu
   src/interop/to_arrow.cu

cpp/include/cudf/hashing.hpp

@@ -70,6 +70,7 @@ std::unique_ptr<column> hash(
   rmm::cuda_stream_view stream        = cudf::get_default_stream(),
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+//! Hash APIs
 namespace hashing {
 
 /**
@@ -125,6 +126,24 @@ std::unique_ptr<column> md5(
   rmm::cuda_stream_view stream        = cudf::get_default_stream(),
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+/**
+ * @brief Computes the hash value of each row in the input set of columns
+ *
+ * This function takes a 64-bit seed value and returns a column of type UINT64.
+ *
+ * @param input The table of columns to hash
+ * @param seed Optional seed value to use for the hash function
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ * @param mr Device memory resource used to allocate the returned column's device memory
+ *
+ * @returns A column where each row is the hash of a row from the input
+ */
+std::unique_ptr<column> xxhash64(
+  table_view const& input,
+  uint64_t seed                       = DEFAULT_HASH_SEED,
+  rmm::cuda_stream_view stream        = cudf::get_default_stream(),
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 }  // namespace hashing
 
 /** @} */  // end of group

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

@@ -52,6 +52,11 @@ __device__ inline uint32_t rotate_bits_left(uint32_t x, uint32_t r)
   return __funnelshift_l(x, x, r);
 }
 
+__device__ inline uint64_t rotate_bits_left(uint64_t x, uint32_t r)
+{
+  return ((x << r) | (x >> (64 - r)));
+}
+
 __device__ inline uint32_t rotate_bits_right(uint32_t x, uint32_t r)
 {
   // This function is equivalent to (x >> r) | (x << (32 - r))

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

@@ -41,6 +41,11 @@ std::unique_ptr<column> md5(table_view const& input,
                             rmm::cuda_stream_view stream,
                             rmm::mr::device_memory_resource* mr);
 
+std::unique_ptr<column> xxhash64(table_view const& input,
+                                 uint64_t seed,
+                                 rmm::cuda_stream_view,
+                                 rmm::mr::device_memory_resource* mr);
+
 /* Copyright 2005-2014 Daniel James.
  *
  * Use, modification and distribution is subject to the Boost Software

cpp/src/hash/spark_murmurhash3_x86_32.cu

@@ -89,7 +89,7 @@ struct Spark_MurmurHash3_x86_32 {
       k1 = rotate_bits_left(k1, rot_c1);
       k1 *= c2;
       h ^= k1;
-      h = rotate_bits_left(h, rot_c2);
+      h = rotate_bits_left(static_cast<uint32_t>(h), rot_c2);
       h = h * 5 + c3;
     }
     return h;
@@ -109,7 +109,7 @@ struct Spark_MurmurHash3_x86_32 {
       k1 = rotate_bits_left(k1, rot_c1);
       k1 *= c2;
       h ^= k1;
-      h = rotate_bits_left(h, rot_c2);
+      h = rotate_bits_left(static_cast<uint32_t>(h), rot_c2);
       h = h * 5 + c3;
     }