Add Spark-compatible MurmurHash3

cpp/CMakeLists.txt

@@ -706,6 +706,7 @@ add_library(
   src/hash/md5_hash.cu
   src/hash/murmurhash3_x86_32.cu
   src/hash/murmurhash3_x64_128.cu
+  src/hash/spark_murmurhash3_x86_32.cu
   src/hash/sha1_hash.cu
   src/hash/sha224_hash.cu
   src/hash/sha256_hash.cu

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

@@ -0,0 +1,135 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include <cudf/detail/row_operator/hashing.cuh>
+#include <cudf/detail/utilities/accumulate.cuh>
+#include <cudf/lists/lists_column_device_view.cuh>
+#include <cudf/structs/structs_column_device_view.cuh>
+#include <cudf/utilities/type_dispatcher.hpp>
+
+#include <cuda/functional>
+#include <thrust/iterator/counting_iterator.h>
+
+namespace CUDF_EXPORT cudf {
+namespace detail::row::hash {
+
+/**
+ * @brief Computes row hashes using Spark's iterative seeding convention.
+ *
+ * Spark uses the hash of each value as the seed for the next value and ignores
+ * null values. Consequently, values of different nested shapes can collide.
+ * For example, the integer `1`, the list `[1]`, and a struct containing only
+ * `1` have the same hash. Likewise, `[1]`, `[1, null]`, and `[null, 1]` have
+ * the same hash. A null element returns its input seed unchanged.
+ *
+ * The element hash function is responsible for Spark-specific type encodings
+ * and hash algorithm behavior.
+ *
+ * LIST columns whose child is a STRUCT are unsupported and must be rejected
+ * before invoking this hasher.
+ *
+ * @tparam hash_function Seeded element hash functor with a `result_type` member
+ * @tparam Nullate A cudf::nullate type describing whether to check for nulls
+ */
+template <template <typename> class hash_function, typename Nullate>
+class spark_device_row_hasher {
+  friend class row_hasher;
+
+ public:
+  using result_type = typename hash_function<int32_t>::result_type;
+
+  /**
+   * @brief Returns the hash value of a row in the table.
+   *
+   * @param row_index The row index to hash
+   * @return The hash value of the row
+   */
+  __device__ result_type operator()(size_type row_index) const noexcept
+  {
+    return cudf::detail::accumulate(
+      _table.begin(),
+      _table.end(),
+      _seed,
+      cuda::proclaim_return_type<result_type>(
+        [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
+          return cudf::type_dispatcher(
+            column.type(), element_hasher_adapter{nulls, hash}, column, row_index);
+        }));
+  }
+
+ private:
+  /**
+   * @brief Computes the hash value of an element in a column.
+   *
+   * Nested values are flattened and hashed serially, with each output becoming
+   * the seed for the next value. A null element returns the input seed.
+   */
+  class element_hasher_adapter {
+    using hash_functor = element_hasher<hash_function, Nullate>;
+
+   public:
+    __device__ element_hasher_adapter(Nullate check_nulls, result_type seed) noexcept
+      : _check_nulls(check_nulls), _seed(seed)
+    {
+    }
+
+    template <typename T>
+      requires(not cudf::is_nested<T>())
+    __device__ result_type operator()(column_device_view const& col,
+                                      size_type row_index) const noexcept
+    {
+      auto const hasher = hash_functor{_check_nulls, _seed, _seed};
+      return hasher.template operator()<T>(col, row_index);
+    }
+
+    template <typename T>
+      requires(cudf::is_nested<T>())
+    __device__ result_type operator()(column_device_view const& col,
+                                      size_type row_index) const noexcept
+    {
+      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 (curr_col.type().id() == type_id::STRUCT) {
+          if (curr_col.num_child_columns() == 0) { return _seed; }
+          // Non-empty structs are decomposed and contain only one child.
+          curr_col = structs_column_device_view(curr_col).get_sliced_child(0);
+        } else {
+          curr_col = lists_column_device_view(curr_col).get_sliced_child();
+        }
+      }
+
+      return cudf::detail::accumulate(
+        thrust::counting_iterator(0),
+        thrust::counting_iterator(curr_col.size()),
+        _seed,
+        cuda::proclaim_return_type<result_type>(
+          [curr_col, nulls = this->_check_nulls] __device__(auto hash, auto element_index) {
+            auto const hasher = hash_functor{nulls, hash, hash};
+            return cudf::type_dispatcher<cudf::detail::dispatch_void_if_nested>(
+              curr_col.type(), hasher, curr_col, element_index);
+          }));
+    }
+
+   private:
+    Nullate const _check_nulls;
+    result_type const _seed;
+  };
+
+  CUDF_HOST_DEVICE spark_device_row_hasher(Nullate check_nulls,
+                                           table_device_view table,
+                                           result_type seed = DEFAULT_HASH_SEED) noexcept
+    : _check_nulls{check_nulls}, _table{table}, _seed(seed)
+  {
+  }
+
+  Nullate const _check_nulls;
+  table_device_view const _table;
+  result_type const _seed;
+};
+
+}  // namespace detail::row::hash
+}  // namespace CUDF_EXPORT cudf

cpp/include/cudf/hashing.hpp

@@ -9,6 +9,8 @@
 #include <cudf/utilities/export.hpp>
 #include <cudf/utilities/memory_resource.hpp>
 
+#include <cuda/stream_ref>
+
 /**
  * @file
  * @brief APIs for computing hash values of columns and tables using various hash algorithms.
@@ -39,9 +41,8 @@ namespace hashing {
 /**
  * @brief Computes the MurmurHash3 32-bit hash value of each row in the given table
  *
- * This function computes the hash of each column using the `seed` for the first column
- * and the resulting hash as a seed for the next column and so on.
- * The result is a uint32 value for each row.
+ * This function hashes each column using the same initial `seed`, then combines the column
+ * hashes into a single uint32 value for each row.
  *
  * @param input The table of columns to hash
  * @param seed Optional seed value to use for the hash function
@@ -56,6 +57,29 @@ std::unique_ptr<column> murmurhash3_x86_32(
   rmm::cuda_stream_view stream      = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+/**
+ * @brief Computes the Apache Spark-compatible MurmurHash3 32-bit hash of each row
+ *
+ * This function follows Apache Spark value hashing and row traversal semantics. Each non-null value
+ * is hashed using the preceding value hash as its seed. Null values leave the current hash
+ * unchanged. Spark-specific handling is applied to strings, narrow integral types, floating-point
+ * NaNs and signed zeros, fixed-point values, lists, and structs.
+ *
+ * This function does not support LIST columns whose child is a STRUCT.
+ *
+ * @param input The table of columns to hash
+ * @param seed Optional initial seed value
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ * @param mr Device memory resource used to allocate the returned column device memory
+ *
+ * @returns A non-nullable INT32 column containing one Spark-compatible hash per input row
+ */
+std::unique_ptr<column> spark_murmurhash3_x86_32(
+  table_view const& input,
+  uint32_t seed                     = DEFAULT_HASH_SEED,
+  cuda::stream_ref stream           = cudf::get_default_stream(),
+  rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+
 /**
  * @brief Computes the MurmurHash3 64-bit hash value of each row in the given table
  *

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

@@ -22,6 +22,11 @@ std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
                                            rmm::cuda_stream_view,
                                            rmm::device_async_resource_ref mr);
 
+std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
+                                                 uint32_t seed,
+                                                 cuda::stream_ref,
+                                                 rmm::device_async_resource_ref mr);
+
 std::unique_ptr<table> murmurhash3_x64_128(table_view const& input,
                                            uint64_t seed,
                                            rmm::cuda_stream_view,