From 9b3d166c87d1f41fcc201db9f57f6b701db24f50 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Tue, 14 Jul 2026 18:47:53 +0000
Subject: [PATCH 01/12] Use iterative seeding for multi-column row hashing

---
 .../cudf/detail/row_operator/hashing.cuh      | 22 +++++++++++++------
 .../row_operator/preprocessed_table.cuh       |  9 +++++---
 cpp/src/row_operator/row_operators.cu         |  9 +++++---
 cpp/tests/hashing/murmurhash3_x86_32_test.cpp | 14 ++++++++++++
 4 files changed, 41 insertions(+), 13 deletions(-)

diff --git a/cpp/include/cudf/detail/row_operator/hashing.cuh b/cpp/include/cudf/detail/row_operator/hashing.cuh
index e71fc5213483..0f2c170a5c81 100644
--- a/cpp/include/cudf/detail/row_operator/hashing.cuh
+++ b/cpp/include/cudf/detail/row_operator/hashing.cuh
@@ -22,9 +22,9 @@
 #include <cudf/utilities/traits.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
-#include <thrust/iterator/transform_iterator.h>
 
 #include <memory>
 
@@ -117,16 +117,22 @@ class device_row_hasher {
    */
   __device__ result_type operator()(size_type row_index) const noexcept
   {
-    auto const hasher = [row_index, this](auto const& column) {
+    auto const hasher = [row_index, this](auto seed, auto const& column) {
       return cudf::type_dispatcher<dispatch_storage_type>(
-        column.type(), element_hasher_adapter{_check_nulls, _seed}, column, row_index);
+        column.type(), element_hasher_adapter{_check_nulls, seed}, column, row_index);
     };
 
+    if (_num_input_columns > 1) {
+      return detail::accumulate(_table.begin(), _table.end(), _seed, hasher);
+    }
+
     auto const has_columns = _table.num_columns() > 0;
-    auto const init        = has_columns ? hasher(_table.column(0)) : _seed;
+    auto const init        = has_columns ? hasher(_seed, _table.column(0)) : _seed;
     auto const start_col   = static_cast<size_type>(has_columns);
 
-    auto it = thrust::make_transform_iterator(_table.begin() + start_col, hasher);
+    auto it = cuda::transform_iterator{
+      _table.begin() + start_col,
+      [hasher, seed = _seed](auto const& column) { return hasher(seed, column); }};
     return detail::accumulate(
       it, it + (_table.num_columns() - start_col), init, [](auto hash, auto h) {
         return cudf::hashing::detail::hash_combine(hash, h);
@@ -217,13 +223,15 @@ class device_row_hasher {
 
   CUDF_HOST_DEVICE device_row_hasher(Nullate check_nulls,
                                      table_device_view t,
+                                     size_type num_input_columns,
                                      result_type seed = DEFAULT_HASH_SEED) noexcept
-    : _check_nulls{check_nulls}, _table{t}, _seed(seed)
+    : _check_nulls{check_nulls}, _table{t}, _num_input_columns{num_input_columns}, _seed(seed)
   {
   }
 
   Nullate const _check_nulls;
   table_device_view const _table;
+  size_type const _num_input_columns;
   // Assumes seeds are the same as the result type of the hash function
   result_type const _seed;
 };
@@ -278,7 +286,7 @@ class row_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);
+    return DeviceRowHasher<hash_function, Nullate>(nullate, *d_t, d_t->_num_input_columns, seed);
   }
 
  private:
diff --git a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh
index 3490b9b826e8..c004891afd1e 100644
--- a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh
+++ b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -73,16 +73,19 @@ struct preprocessed_table {
 
   preprocessed_table(table_device_view_owner&& table,
                      std::vector<rmm::device_buffer>&& null_buffers,
-                     std::vector<std::unique_ptr<column>>&& tmp_columns)
+                     std::vector<std::unique_ptr<column>>&& tmp_columns,
+                     size_type num_input_columns)
     : _t(std::move(table)),
       _null_buffers(std::move(null_buffers)),
-      _tmp_columns(std::move(tmp_columns))
+      _tmp_columns(std::move(tmp_columns)),
+      _num_input_columns(num_input_columns)
   {
   }
 
   table_device_view_owner _t;
   std::vector<rmm::device_buffer> _null_buffers;
   std::vector<std::unique_ptr<column>> _tmp_columns;
+  size_type _num_input_columns;
 };
 
 }  // namespace equality
diff --git a/cpp/src/row_operator/row_operators.cu b/cpp/src/row_operator/row_operators.cu
index 4b21d9c980cb..4400bdfa1eb2 100644
--- a/cpp/src/row_operator/row_operators.cu
+++ b/cpp/src/row_operator/row_operators.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -854,8 +854,11 @@ std::shared_ptr<preprocessed_table> preprocessed_table::create(table_view const&
     std::get<0>(decompose_structs(struct_offset_removed_table, decompose_lists_column::YES));
 
   auto d_t = table_device_view_owner(table_device_view::create(verticalized_t, stream));
-  return std::shared_ptr<preprocessed_table>(new preprocessed_table(
-    std::move(d_t), std::move(nullable_data.new_null_masks), std::move(nullable_data.new_columns)));
+  return std::shared_ptr<preprocessed_table>(
+    new preprocessed_table(std::move(d_t),
+                           std::move(nullable_data.new_null_masks),
+                           std::move(nullable_data.new_columns),
+                           t.num_columns()));
 }
 
 two_table_comparator::two_table_comparator(table_view const& left,
diff --git a/cpp/tests/hashing/murmurhash3_x86_32_test.cpp b/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
index bb0b5cc41efa..897b586008ec 100644
--- a/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
+++ b/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
@@ -50,6 +50,20 @@ TEST_F(MurmurHashTest, MultiValue)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view());
 }
 
+TEST_F(MurmurHashTest, SparkCompatibleIterativeSeeding)
+{
+  cudf::test::fixed_width_column_wrapper<int32_t> const first({0, 1, -1, 42, 123456789});
+  cudf::test::fixed_width_column_wrapper<int32_t> const second({10, 20, 30, -40, -987654321});
+  auto const input = cudf::table_view({first, second});
+
+  auto const output = cudf::hashing::murmurhash3_x86_32(input, 42);
+
+  // Generated with Apache Spark's Murmur3Hash using seed 42.
+  cudf::test::fixed_width_column_wrapper<uint32_t> const expected(
+    {2573243963u, 1151116018u, 1549484878u, 3007216400u, 2314233967u});
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected);
+}
+
 TEST_F(MurmurHashTest, MultiValueNulls)
 {
   // Nulls with different values should be equal

From 51cd9a21c377bbb3ec2ee319db8dd19e3aa525d4 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Tue, 14 Jul 2026 20:12:16 +0000
Subject: [PATCH 02/12] Preserve custom row hasher compatibility

---
 .../cudf/detail/row_operator/hashing.cuh      |  8 ++++++-
 cpp/tests/row_operator/row_operator_tests.cu  | 24 ++++++++++++++++++-
 2 files changed, 30 insertions(+), 2 deletions(-)

diff --git a/cpp/include/cudf/detail/row_operator/hashing.cuh b/cpp/include/cudf/detail/row_operator/hashing.cuh
index 0f2c170a5c81..b2eb62a0db04 100644
--- a/cpp/include/cudf/detail/row_operator/hashing.cuh
+++ b/cpp/include/cudf/detail/row_operator/hashing.cuh
@@ -286,7 +286,13 @@ class row_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, d_t->_num_input_columns, seed);
+    using device_hasher_type = DeviceRowHasher<hash_function, Nullate>;
+    if constexpr (cuda::std::is_same_v<device_hasher_type,
+                                       device_row_hasher<hash_function, Nullate>>) {
+      return device_hasher_type(nullate, *d_t, d_t->_num_input_columns, seed);
+    } else {
+      return device_hasher_type(nullate, *d_t, seed);
+    }
   }
 
  private:
diff --git a/cpp/tests/row_operator/row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu
index 08046c5e3128..8d0e64a8b677 100644
--- a/cpp/tests/row_operator/row_operator_tests.cu
+++ b/cpp/tests/row_operator/row_operator_tests.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -289,6 +289,28 @@ TYPED_TEST(NaNTableViewTest, TestEqualityComparatorTwoTableNaNCase)
 
 struct RowOperatorTest : public cudf::test::BaseFixture {};
 
+template <template <typename> class hash_function, typename Nullate>
+struct legacy_device_row_hasher {
+  using result_type = cuda::std::invoke_result_t<hash_function<int32_t>, int32_t>;
+
+  CUDF_HOST_DEVICE legacy_device_row_hasher(Nullate,
+                                            cudf::table_device_view,
+                                            result_type = cudf::DEFAULT_HASH_SEED) noexcept
+  {
+  }
+};
+
+TEST_F(RowOperatorTest, TestCustomRowHasherConstructorCompatibility)
+{
+  auto const col = cudf::test::fixed_width_column_wrapper<int32_t>{{1, 2}};
+  auto const row_hasher =
+    cudf::detail::row::hash::row_hasher(cudf::table_view{{col}}, cudf::get_default_stream());
+
+  [[maybe_unused]] auto const custom_hasher =
+    row_hasher.device_hasher<cudf::hashing::detail::XXHash_64, legacy_device_row_hasher>(
+      cudf::nullate::NO{});
+}
+
 TEST_F(RowOperatorTest, TestTwoTableComparatorColumnCountCheck)
 {
   rmm::cuda_stream_view stream{cudf::get_default_stream()};

From 6cabfd01b8dcb49ba38ab2f3dddc54bc891ba22f Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Thu, 16 Jul 2026 00:55:46 +0000
Subject: [PATCH 03/12] Isolate iterative seeding to MurmurHash3

---
 .../cudf/detail/row_operator/hashing.cuh      | 29 +++++++--------
 .../row_operator/preprocessed_table.cuh       |  9 ++---
 cpp/src/hash/murmurhash3_x86_32.cu            | 37 +++++++++++++------
 cpp/src/row_operator/row_operators.cu         |  9 ++---
 cpp/tests/hashing/murmurhash3_x86_32_test.cpp |  4 +-
 cpp/tests/row_operator/row_operator_tests.cu  | 24 +-----------
 python/cudf/cudf/core/indexed_frame.py        |  6 +--
 .../dataframe/methods/test_hash_values.py     | 17 ++++++++-
 8 files changed, 68 insertions(+), 67 deletions(-)

diff --git a/cpp/include/cudf/detail/row_operator/hashing.cuh b/cpp/include/cudf/detail/row_operator/hashing.cuh
index b2eb62a0db04..d3fc765fd026 100644
--- a/cpp/include/cudf/detail/row_operator/hashing.cuh
+++ b/cpp/include/cudf/detail/row_operator/hashing.cuh
@@ -22,9 +22,9 @@
 #include <cudf/utilities/traits.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 
-#include <cuda/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
+#include <thrust/iterator/transform_iterator.h>
 
 #include <memory>
 
@@ -101,8 +101,11 @@ class element_hasher {
  *
  * @tparam hash_function Hash functor to use for hashing elements.
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+ * @tparam use_iterative_seeding Whether to use each column hash as the seed for the next column
  */
-template <template <typename> class hash_function, typename Nullate>
+template <template <typename> class hash_function,
+          typename Nullate,
+          bool use_iterative_seeding = false>
 class device_row_hasher {
   friend class row_hasher;
 
@@ -122,7 +125,7 @@ class device_row_hasher {
         column.type(), element_hasher_adapter{_check_nulls, seed}, column, row_index);
     };
 
-    if (_num_input_columns > 1) {
+    if constexpr (use_iterative_seeding) {
       return detail::accumulate(_table.begin(), _table.end(), _seed, hasher);
     }
 
@@ -130,9 +133,9 @@ class device_row_hasher {
     auto const init        = has_columns ? hasher(_seed, _table.column(0)) : _seed;
     auto const start_col   = static_cast<size_type>(has_columns);
 
-    auto it = cuda::transform_iterator{
+    auto it = thrust::make_transform_iterator(
       _table.begin() + start_col,
-      [hasher, seed = _seed](auto const& column) { return hasher(seed, column); }};
+      [hasher, seed = _seed](auto const& column) { return hasher(seed, column); });
     return detail::accumulate(
       it, it + (_table.num_columns() - start_col), init, [](auto hash, auto h) {
         return cudf::hashing::detail::hash_combine(hash, h);
@@ -223,19 +226,21 @@ class device_row_hasher {
 
   CUDF_HOST_DEVICE device_row_hasher(Nullate check_nulls,
                                      table_device_view t,
-                                     size_type num_input_columns,
                                      result_type seed = DEFAULT_HASH_SEED) noexcept
-    : _check_nulls{check_nulls}, _table{t}, _num_input_columns{num_input_columns}, _seed(seed)
+    : _check_nulls{check_nulls}, _table{t}, _seed(seed)
   {
   }
 
   Nullate const _check_nulls;
   table_device_view const _table;
-  size_type const _num_input_columns;
   // Assumes seeds are the same as the result type of the hash function
   result_type const _seed;
 };
 
+/** @brief A device row hasher that uses each column hash as the next column seed. */
+template <template <typename> class hash_function, typename Nullate>
+using device_row_hasher_iterative = device_row_hasher<hash_function, Nullate, true>;
+
 /**
  * @brief Computes the hash value of a row in the given table.
  *
@@ -286,13 +291,7 @@ class row_hasher {
     Nullate nullate                                                  = {},
     cuda::std::invoke_result_t<hash_function<int32_t>, int32_t> seed = DEFAULT_HASH_SEED) const
   {
-    using device_hasher_type = DeviceRowHasher<hash_function, Nullate>;
-    if constexpr (cuda::std::is_same_v<device_hasher_type,
-                                       device_row_hasher<hash_function, Nullate>>) {
-      return device_hasher_type(nullate, *d_t, d_t->_num_input_columns, seed);
-    } else {
-      return device_hasher_type(nullate, *d_t, seed);
-    }
+    return DeviceRowHasher<hash_function, Nullate>(nullate, *d_t, seed);
   }
 
  private:
diff --git a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh
index c004891afd1e..3490b9b826e8 100644
--- a/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh
+++ b/cpp/include/cudf/detail/row_operator/preprocessed_table.cuh
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -73,19 +73,16 @@ struct preprocessed_table {
 
   preprocessed_table(table_device_view_owner&& table,
                      std::vector<rmm::device_buffer>&& null_buffers,
-                     std::vector<std::unique_ptr<column>>&& tmp_columns,
-                     size_type num_input_columns)
+                     std::vector<std::unique_ptr<column>>&& tmp_columns)
     : _t(std::move(table)),
       _null_buffers(std::move(null_buffers)),
-      _tmp_columns(std::move(tmp_columns)),
-      _num_input_columns(num_input_columns)
+      _tmp_columns(std::move(tmp_columns))
   {
   }
 
   table_device_view_owner _t;
   std::vector<rmm::device_buffer> _null_buffers;
   std::vector<std::unique_ptr<column>> _tmp_columns;
-  size_type _num_input_columns;
 };
 
 }  // namespace equality
diff --git a/cpp/src/hash/murmurhash3_x86_32.cu b/cpp/src/hash/murmurhash3_x86_32.cu
index 3a28c1a15553..0bd88fb3f59f 100644
--- a/cpp/src/hash/murmurhash3_x86_32.cu
+++ b/cpp/src/hash/murmurhash3_x86_32.cu
@@ -17,6 +17,22 @@ namespace cudf {
 namespace hashing {
 namespace detail {
 
+template <template <template <typename> class, typename> class DeviceRowHasher>
+void hash_rows(cudf::detail::row::hash::row_hasher const& row_hasher,
+               size_type num_rows,
+               hash_value_type* output_begin,
+               bool nullable,
+               uint32_t seed,
+               rmm::cuda_stream_view stream)
+{
+  auto const hasher = row_hasher.device_hasher<MurmurHash3_x86_32, DeviceRowHasher>(nullable, seed);
+  // thrust::tabulate is slow here, see NVIDIA/cccl#9070
+  CUDF_CUDA_TRY(cub::DeviceFor::Bulk(
+    num_rows,
+    [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); },
+    stream.value()));
+}
+
 std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
                                            uint32_t seed,
                                            rmm::cuda_stream_view stream,
@@ -30,18 +46,17 @@ std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
 
   if (input.num_rows() == 0) { return output; }
 
-  bool const nullable   = has_nulls(input);
-  auto const row_hasher = cudf::detail::row::hash::row_hasher(input, stream);
-  auto output_view      = output->mutable_view();
+  bool const nullable     = has_nulls(input);
+  auto const row_hasher   = cudf::detail::row::hash::row_hasher(input, stream);
+  auto const output_begin = output->mutable_view().begin<hash_value_type>();
 
-  // Compute the hash value for each row
-  auto const output_begin = output_view.begin<hash_value_type>();
-  auto const hasher       = row_hasher.device_hasher<MurmurHash3_x86_32>(nullable, seed);
-  // thrust::tabulate is slow here, see NVIDIA/cccl#9070
-  CUDF_CUDA_TRY(cub::DeviceFor::Bulk(
-    input.num_rows(),
-    [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); },
-    stream.value()));
+  if (input.num_columns() > 1) {
+    hash_rows<cudf::detail::row::hash::device_row_hasher_iterative>(
+      row_hasher, input.num_rows(), output_begin, nullable, seed, stream);
+  } else {
+    hash_rows<cudf::detail::row::hash::device_row_hasher>(
+      row_hasher, input.num_rows(), output_begin, nullable, seed, stream);
+  }
 
   return output;
 }
diff --git a/cpp/src/row_operator/row_operators.cu b/cpp/src/row_operator/row_operators.cu
index 4400bdfa1eb2..4b21d9c980cb 100644
--- a/cpp/src/row_operator/row_operators.cu
+++ b/cpp/src/row_operator/row_operators.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -854,11 +854,8 @@ std::shared_ptr<preprocessed_table> preprocessed_table::create(table_view const&
     std::get<0>(decompose_structs(struct_offset_removed_table, decompose_lists_column::YES));
 
   auto d_t = table_device_view_owner(table_device_view::create(verticalized_t, stream));
-  return std::shared_ptr<preprocessed_table>(
-    new preprocessed_table(std::move(d_t),
-                           std::move(nullable_data.new_null_masks),
-                           std::move(nullable_data.new_columns),
-                           t.num_columns()));
+  return std::shared_ptr<preprocessed_table>(new preprocessed_table(
+    std::move(d_t), std::move(nullable_data.new_null_masks), std::move(nullable_data.new_columns)));
 }
 
 two_table_comparator::two_table_comparator(table_view const& left,
diff --git a/cpp/tests/hashing/murmurhash3_x86_32_test.cpp b/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
index 897b586008ec..e83756ebbb6d 100644
--- a/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
+++ b/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
@@ -50,7 +50,7 @@ TEST_F(MurmurHashTest, MultiValue)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view());
 }
 
-TEST_F(MurmurHashTest, SparkCompatibleIterativeSeeding)
+TEST_F(MurmurHashTest, IterativeSeeding)
 {
   cudf::test::fixed_width_column_wrapper<int32_t> const first({0, 1, -1, 42, 123456789});
   cudf::test::fixed_width_column_wrapper<int32_t> const second({10, 20, 30, -40, -987654321});
@@ -58,7 +58,7 @@ TEST_F(MurmurHashTest, SparkCompatibleIterativeSeeding)
 
   auto const output = cudf::hashing::murmurhash3_x86_32(input, 42);
 
-  // Generated with Apache Spark's Murmur3Hash using seed 42.
+  // Matches Apache Spark for non-null fixed-width inputs with seed 42.
   cudf::test::fixed_width_column_wrapper<uint32_t> const expected(
     {2573243963u, 1151116018u, 1549484878u, 3007216400u, 2314233967u});
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected);
diff --git a/cpp/tests/row_operator/row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu
index 8d0e64a8b677..08046c5e3128 100644
--- a/cpp/tests/row_operator/row_operator_tests.cu
+++ b/cpp/tests/row_operator/row_operator_tests.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -289,28 +289,6 @@ TYPED_TEST(NaNTableViewTest, TestEqualityComparatorTwoTableNaNCase)
 
 struct RowOperatorTest : public cudf::test::BaseFixture {};
 
-template <template <typename> class hash_function, typename Nullate>
-struct legacy_device_row_hasher {
-  using result_type = cuda::std::invoke_result_t<hash_function<int32_t>, int32_t>;
-
-  CUDF_HOST_DEVICE legacy_device_row_hasher(Nullate,
-                                            cudf::table_device_view,
-                                            result_type = cudf::DEFAULT_HASH_SEED) noexcept
-  {
-  }
-};
-
-TEST_F(RowOperatorTest, TestCustomRowHasherConstructorCompatibility)
-{
-  auto const col = cudf::test::fixed_width_column_wrapper<int32_t>{{1, 2}};
-  auto const row_hasher =
-    cudf::detail::row::hash::row_hasher(cudf::table_view{{col}}, cudf::get_default_stream());
-
-  [[maybe_unused]] auto const custom_hasher =
-    row_hasher.device_hasher<cudf::hashing::detail::XXHash_64, legacy_device_row_hasher>(
-      cudf::nullate::NO{});
-}
-
 TEST_F(RowOperatorTest, TestTwoTableComparatorColumnCountCheck)
 {
   rmm::cuda_stream_view stream{cudf::get_default_stream()};
diff --git a/python/cudf/cudf/core/indexed_frame.py b/python/cudf/cudf/core/indexed_frame.py
index 57cae22e6572..212e96fe5697 100644
--- a/python/cudf/cudf/core/indexed_frame.py
+++ b/python/cudf/cudf/core/indexed_frame.py
@@ -2973,9 +2973,9 @@ def hash_values(
         1  120  0.25
         2   30  0.50
         >>> df.hash_values(method="murmur3")
-        0      92443192
-        1     116637743
-        2    2602001908
+        0    3826000917
+        1     657073880
+        2     339499243
         dtype: uint32
         >>> df.hash_values(method="md5")
         0    57ce879751b5169c525907d5c563fae1
diff --git a/python/cudf/cudf/tests/dataframe/methods/test_hash_values.py b/python/cudf/cudf/tests/dataframe/methods/test_hash_values.py
index 5ca7761c4bc1..92a64b6258d5 100644
--- a/python/cudf/cudf/tests/dataframe/methods/test_hash_values.py
+++ b/python/cudf/cudf/tests/dataframe/methods/test_hash_values.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
+# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: Apache-2.0
 
 import contextlib
@@ -86,6 +86,21 @@ def test_dataframe_hash_values_seed(method):
     assert_neq(out_one, out_two)
 
 
+def test_dataframe_hash_values_murmur3_iterative_seeding():
+    gdf = cudf.DataFrame(
+        {
+            "a": cudf.Series([0, 1, -1, 42, 123456789], dtype=np.int32),
+            "b": cudf.Series([10, 20, 30, -40, -987654321], dtype=np.int32),
+        }
+    )
+    expected = cudf.Series(
+        [2573243963, 1151116018, 1549484878, 3007216400, 2314233967],
+        dtype=np.uint32,
+    )
+
+    assert_eq(gdf.hash_values(method="murmur3", seed=42), expected)
+
+
 def test_dataframe_hash_values_xxhash32():
     # xxhash32 has no built-in implementation in Python and we don't want to
     # add a testing dependency, so we use regression tests against known good

From c47238e66ad6fd6e8b5e3b55f6d502d95c8d16a4 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Thu, 16 Jul 2026 01:00:58 +0000
Subject: [PATCH 04/12] Use cuda::transform_iterator in row hasher

---
 cpp/include/cudf/detail/row_operator/hashing.cuh | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/cpp/include/cudf/detail/row_operator/hashing.cuh b/cpp/include/cudf/detail/row_operator/hashing.cuh
index d3fc765fd026..1fa22e95b279 100644
--- a/cpp/include/cudf/detail/row_operator/hashing.cuh
+++ b/cpp/include/cudf/detail/row_operator/hashing.cuh
@@ -22,9 +22,9 @@
 #include <cudf/utilities/traits.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 
+#include <cuda/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
-#include <thrust/iterator/transform_iterator.h>
 
 #include <memory>
 
@@ -133,9 +133,9 @@ class device_row_hasher {
     auto const init        = has_columns ? hasher(_seed, _table.column(0)) : _seed;
     auto const start_col   = static_cast<size_type>(has_columns);
 
-    auto it = thrust::make_transform_iterator(
+    auto it = cuda::transform_iterator{
       _table.begin() + start_col,
-      [hasher, seed = _seed](auto const& column) { return hasher(seed, column); });
+      [hasher, seed = _seed](auto const& column) { return hasher(seed, column); }};
     return detail::accumulate(
       it, it + (_table.num_columns() - start_col), init, [](auto hash, auto h) {
         return cudf::hashing::detail::hash_combine(hash, h);

From e2bd413d612c5c03b353f997583ee7c9e62c41ac Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Thu, 16 Jul 2026 01:02:20 +0000
Subject: [PATCH 05/12] Document iterative row hashing semantics

---
 cpp/include/cudf/detail/row_operator/hashing.cuh | 13 ++++++++++++-
 1 file changed, 12 insertions(+), 1 deletion(-)

diff --git a/cpp/include/cudf/detail/row_operator/hashing.cuh b/cpp/include/cudf/detail/row_operator/hashing.cuh
index 1fa22e95b279..d736206f2275 100644
--- a/cpp/include/cudf/detail/row_operator/hashing.cuh
+++ b/cpp/include/cudf/detail/row_operator/hashing.cuh
@@ -237,7 +237,18 @@ class device_row_hasher {
   result_type const _seed;
 };
 
-/** @brief A device row hasher that uses each column hash as the next column seed. */
+/**
+ * @brief A device row hasher that uses each column hash as the seed for the next column.
+ *
+ * Spark computes multi-column MurmurHash3 values by using each column hash as the seed for the
+ * following column. This iterative seed propagation is therefore required for Spark-compatible
+ * results when the element-level hashing semantics are otherwise compatible.
+ *
+ * @note This hasher provides Spark-compatible seed propagation only.
+ *
+ * @tparam hash_function Hash functor to use for hashing elements
+ * @tparam Nullate A cudf::nullate type describing whether to check for nulls
+ */
 template <template <typename> class hash_function, typename Nullate>
 using device_row_hasher_iterative = device_row_hasher<hash_function, Nullate, true>;
 

From 13491786f4f584c15a88cd38e62c12bf0db9c9ef Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Thu, 6 Aug 2026 20:25:57 +0000
Subject: [PATCH 06/12] Add Spark-compatible MurmurHash to libcudf

---
 cpp/CMakeLists.txt                            |   1 +
 .../cudf/detail/row_operator/hashing.cuh      |  36 +-
 cpp/include/cudf/hashing.hpp                  |  28 +-
 .../cudf/hashing/detail/spark_murmurhash3.cuh | 382 ++++++++++++
 cpp/include/cudf/partitioning.hpp             |   3 +-
 cpp/src/hash/murmurhash3_x86_32.cu            |  37 +-
 cpp/src/hash/spark_murmurhash3_x86_32.cu      |  89 +++
 cpp/src/partitioning/partitioning.cu          |  25 +-
 cpp/tests/CMakeLists.txt                      |   1 +
 cpp/tests/hashing/murmurhash3_x86_32_test.cpp |  14 -
 .../hashing/spark_murmurhash3_x86_32_test.cpp | 579 ++++++++++++++++++
 .../partitioning/hash_partition_test.cpp      |  20 +-
 python/cudf/cudf/core/indexed_frame.py        |   6 +-
 .../dataframe/methods/test_hash_values.py     |  17 +-
 14 files changed, 1138 insertions(+), 100 deletions(-)
 create mode 100644 cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
 create mode 100644 cpp/src/hash/spark_murmurhash3_x86_32.cu
 create mode 100644 cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 4ba016d0da6c..1ef26de9f0af 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -716,6 +716,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
diff --git a/cpp/include/cudf/detail/row_operator/hashing.cuh b/cpp/include/cudf/detail/row_operator/hashing.cuh
index d736206f2275..e71fc5213483 100644
--- a/cpp/include/cudf/detail/row_operator/hashing.cuh
+++ b/cpp/include/cudf/detail/row_operator/hashing.cuh
@@ -22,9 +22,9 @@
 #include <cudf/utilities/traits.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 
-#include <cuda/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/type_traits>
+#include <thrust/iterator/transform_iterator.h>
 
 #include <memory>
 
@@ -101,11 +101,8 @@ class element_hasher {
  *
  * @tparam hash_function Hash functor to use for hashing elements.
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
- * @tparam use_iterative_seeding Whether to use each column hash as the seed for the next column
  */
-template <template <typename> class hash_function,
-          typename Nullate,
-          bool use_iterative_seeding = false>
+template <template <typename> class hash_function, typename Nullate>
 class device_row_hasher {
   friend class row_hasher;
 
@@ -120,22 +117,16 @@ class device_row_hasher {
    */
   __device__ result_type operator()(size_type row_index) const noexcept
   {
-    auto const hasher = [row_index, this](auto seed, auto const& column) {
+    auto const hasher = [row_index, this](auto const& column) {
       return cudf::type_dispatcher<dispatch_storage_type>(
-        column.type(), element_hasher_adapter{_check_nulls, seed}, column, row_index);
+        column.type(), element_hasher_adapter{_check_nulls, _seed}, column, row_index);
     };
 
-    if constexpr (use_iterative_seeding) {
-      return detail::accumulate(_table.begin(), _table.end(), _seed, hasher);
-    }
-
     auto const has_columns = _table.num_columns() > 0;
-    auto const init        = has_columns ? hasher(_seed, _table.column(0)) : _seed;
+    auto const init        = has_columns ? hasher(_table.column(0)) : _seed;
     auto const start_col   = static_cast<size_type>(has_columns);
 
-    auto it = cuda::transform_iterator{
-      _table.begin() + start_col,
-      [hasher, seed = _seed](auto const& column) { return hasher(seed, column); }};
+    auto it = thrust::make_transform_iterator(_table.begin() + start_col, hasher);
     return detail::accumulate(
       it, it + (_table.num_columns() - start_col), init, [](auto hash, auto h) {
         return cudf::hashing::detail::hash_combine(hash, h);
@@ -237,21 +228,6 @@ class device_row_hasher {
   result_type const _seed;
 };
 
-/**
- * @brief A device row hasher that uses each column hash as the seed for the next column.
- *
- * Spark computes multi-column MurmurHash3 values by using each column hash as the seed for the
- * following column. This iterative seed propagation is therefore required for Spark-compatible
- * results when the element-level hashing semantics are otherwise compatible.
- *
- * @note This hasher provides Spark-compatible seed propagation only.
- *
- * @tparam hash_function Hash functor to use for hashing elements
- * @tparam Nullate A cudf::nullate type describing whether to check for nulls
- */
-template <template <typename> class hash_function, typename Nullate>
-using device_row_hasher_iterative = device_row_hasher<hash_function, Nullate, true>;
-
 /**
  * @brief Computes the hash value of a row in the given table.
  *
diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp
index 669f4444751c..bf447a927061 100644
--- a/cpp/include/cudf/hashing.hpp
+++ b/cpp/include/cudf/hashing.hpp
@@ -39,9 +39,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 +55,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, 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,
+  rmm::cuda_stream_view 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
  *
diff --git a/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
new file mode 100644
index 000000000000..2caa0cfe9810
--- /dev/null
+++ b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
@@ -0,0 +1,382 @@
+/*
+ * 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/fixed_point/fixed_point.hpp>
+#include <cudf/hashing/detail/hash_functions.cuh>
+#include <cudf/strings/string_view.hpp>
+#include <cudf/utilities/traits.hpp>
+#include <cudf/utilities/type_dispatcher.hpp>
+
+#include <cuda/functional>
+#include <cuda/std/algorithm>
+#include <cuda/std/cstddef>
+#include <cuda/std/iterator>
+#include <cuda/std/type_traits>
+#include <thrust/find.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/reverse.h>
+
+namespace cudf::hashing::detail {
+
+using spark_hash_value_type = int32_t;
+
+template <typename Key, CUDF_ENABLE_IF(not cudf::is_nested<Key>())>
+struct Spark_MurmurHash3_x86_32 {
+  using result_type = spark_hash_value_type;
+
+  CUDF_HOST_DEVICE constexpr Spark_MurmurHash3_x86_32() = delete;
+  CUDF_HOST_DEVICE constexpr Spark_MurmurHash3_x86_32(result_type seed)
+    : m_seed(static_cast<uint32_t>(seed))
+  {
+  }
+
+  [[nodiscard]] __device__ inline uint32_t fmix32(uint32_t h) const
+  {
+    h ^= h >> 16;
+    h *= 0x85ebca6b;
+    h ^= h >> 13;
+    h *= 0xc2b2ae35;
+    h ^= h >> 16;
+    return h;
+  }
+
+  [[nodiscard]] __device__ inline uint32_t getblock32(cuda::std::byte const* data,
+                                                      cudf::size_type offset) const
+  {
+    // Read a 4-byte value from the data pointer as individual bytes for safe
+    // unaligned access (very likely for string types).
+    auto block = reinterpret_cast<uint8_t const*>(data + offset);
+    return block[0] | (block[1] << 8) | (block[2] << 16) | (block[3] << 24);
+  }
+
+  [[nodiscard]] result_type __device__ inline operator()(Key const& key) const
+  {
+    return compute(key);
+  }
+
+  template <typename T>
+  result_type __device__ inline compute(T const& key) const
+  {
+    return compute_bytes(reinterpret_cast<cuda::std::byte const*>(&key), sizeof(T));
+  }
+
+  result_type __device__ inline compute_remaining_bytes(cuda::std::byte const* data,
+                                                        cudf::size_type len,
+                                                        cudf::size_type tail_offset,
+                                                        result_type h) const
+  {
+    // Process remaining bytes that do not fill a four-byte chunk using Spark's approach
+    // (does not conform to normal MurmurHash3).
+    for (auto i = tail_offset; i < len; i++) {
+      // We require a two-step cast to get the k1 value from the byte. First,
+      // we must cast to a signed int8_t. Then, the sign bit is preserved when
+      // casting to uint32_t under 2's complement. Java preserves the sign when
+      // casting byte-to-int, but C++ does not.
+      uint32_t k1 = static_cast<uint32_t>(cuda::std::to_integer<int8_t>(data[i]));
+      k1 *= c1;
+      k1 = rotate_bits_left(k1, rot_c1);
+      k1 *= c2;
+      h ^= k1;
+      h = rotate_bits_left(static_cast<uint32_t>(h), rot_c2);
+      h = h * 5 + c3;
+    }
+    return h;
+  }
+
+  result_type __device__ compute_bytes(cuda::std::byte const* data, cudf::size_type const len) const
+  {
+    constexpr cudf::size_type BLOCK_SIZE = 4;
+    cudf::size_type const nblocks        = len / BLOCK_SIZE;
+    cudf::size_type const tail_offset    = nblocks * BLOCK_SIZE;
+    result_type h                        = m_seed;
+
+    // Process all four-byte chunks.
+    for (cudf::size_type i = 0; i < nblocks; i++) {
+      uint32_t k1 = getblock32(data, i * BLOCK_SIZE);
+      k1 *= c1;
+      k1 = rotate_bits_left(k1, rot_c1);
+      k1 *= c2;
+      h ^= k1;
+      h = rotate_bits_left(static_cast<uint32_t>(h), rot_c2);
+      h = h * 5 + c3;
+    }
+
+    h = compute_remaining_bytes(data, len, tail_offset, h);
+
+    // Finalize hash.
+    h ^= len;
+    h = fmix32(h);
+    return h;
+  }
+
+ private:
+  uint32_t m_seed{cudf::DEFAULT_HASH_SEED};
+  static constexpr uint32_t c1     = 0xcc9e2d51;
+  static constexpr uint32_t c2     = 0x1b873593;
+  static constexpr uint32_t c3     = 0xe6546b64;
+  static constexpr uint32_t rot_c1 = 15;
+  static constexpr uint32_t rot_c2 = 13;
+};
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<bool>::operator()(
+  bool const& key) const
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<int8_t>::operator()(
+  int8_t const& key) const
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<uint8_t>::operator()(
+  uint8_t const& key) const
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<int16_t>::operator()(
+  int16_t const& key) const
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<uint16_t>::operator()(
+  uint16_t const& key) const
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<float>::operator()(
+  float const& key) const
+{
+  return compute<float>(normalize_nans(key));
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<double>::operator()(
+  double const& key) const
+{
+  return compute<double>(normalize_nans(key));
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<cudf::string_view>::operator()(
+  cudf::string_view const& key) const
+{
+  auto const data = reinterpret_cast<cuda::std::byte const*>(key.data());
+  auto const len  = key.size_bytes();
+  return compute_bytes(data, len);
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<numeric::decimal32>::operator()(
+  numeric::decimal32 const& key) const
+{
+  return compute<uint64_t>(key.value());
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<numeric::decimal64>::operator()(
+  numeric::decimal64 const& key) const
+{
+  return compute<uint64_t>(key.value());
+}
+
+template <>
+spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<numeric::decimal128>::operator()(
+  numeric::decimal128 const& key) const
+{
+  // Generates the Spark MurmurHash3 hash value, mimicking the conversion:
+  // java.math.BigDecimal.valueOf(unscaled_value, _scale).unscaledValue().toByteArray()
+  // https://github.com/apache/spark/blob/master/sql/catalyst/src/main/scala/org/apache/spark/sql/catalyst/expressions/hash.scala#L381
+  __int128_t const val               = key.value();
+  constexpr cudf::size_type key_size = sizeof(__int128_t);
+  cuda::std::byte const* data        = reinterpret_cast<cuda::std::byte const*>(&val);
+
+  // Small negative values start with 0xff..., small positive values start with 0x00...
+  bool const is_negative           = val < 0;
+  cuda::std::byte const zero_value = is_negative ? cuda::std::byte{0xff} : cuda::std::byte{0x00};
+
+  // If the value can be represented with a shorter than 16-byte integer, the
+  // leading bytes of the little-endian value are truncated and are not hashed.
+  auto const reverse_begin = cuda::std::reverse_iterator(data + key_size);
+  auto const reverse_end   = cuda::std::reverse_iterator(data);
+  auto const first_nonzero_byte =
+    thrust::find_if_not(thrust::seq,
+                        reverse_begin,
+                        reverse_end,
+                        [zero_value](cuda::std::byte const& v) { return v == zero_value; })
+      .base();
+  // Max handles special case of 0 and -1 which would shorten to 0 length otherwise
+  cudf::size_type length =
+    cuda::std::max(1, static_cast<cudf::size_type>(cuda::std::distance(data, first_nonzero_byte)));
+
+  // Preserve the 2's complement sign bit by adding a byte back on if necessary.
+  // e.g. 0x0000ff would shorten to 0x00ff. The 0x00 byte is retained to
+  // preserve the sign bit, rather than leaving an "f" at the front which would
+  // change the sign bit. However, 0x00007f would shorten to 0x7f. No extra byte
+  // is needed because the leftmost bit matches the sign bit. Similarly for
+  // negative values: 0xffff00 --> 0xff00 and 0xffff80 --> 0x80.
+  if ((length < key_size) && (is_negative ^ bool(data[length - 1] & cuda::std::byte{0x80}))) {
+    ++length;
+  }
+
+  // Convert to big endian by reversing the range of nonzero bytes. Only those bytes are hashed.
+  __int128_t big_endian_value = 0;
+  auto big_endian_data        = reinterpret_cast<cuda::std::byte*>(&big_endian_value);
+  thrust::reverse_copy(thrust::seq, data, data + length, big_endian_data);
+  return compute_bytes(big_endian_data, length);
+}
+
+/**
+ * @brief Computes the hash value of a row in the given table.
+ *
+ * This functor uses Spark conventions for Murmur hashing, which differs from
+ * the Murmur implementation used in the rest of libcudf. These differences
+ * include:
+ * - Serially using the output hash as an input seed for the next item
+ * - Ignorance of null values
+ *
+ * The serial use of hashes as seeds means that data of different nested types
+ * can exhibit hash collisions. For example, a row of an integer column
+ * containing a 1 will have the same hash as a lists column of integers
+ * containing a list of [1] and a struct column of a single integer column
+ * containing a struct of {1}.
+ *
+ * As a consequence of ignoring null values, inputs like [1], [1, null], and
+ * [null, 1] have the same hash (an expected hash collision). This kind of
+ * collision can also occur across a table of nullable columns and with nulls
+ * in structs ({1, null} and {null, 1} have the same hash). The seed value (the
+ * previous element's hash value) is returned as the hash if an element is
+ * null.
+ *
+ * For additional differences such as special tail processing and decimal type
+ * handling, refer to the Spark_MurmurHash3_x86_32 functor.
+ *
+ * @tparam hash_function Hash functor to use for hashing elements. Must be Spark_MurmurHash3_x86_32.
+ * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+ */
+template <template <typename> class hash_function, typename Nullate>
+class spark_murmur_device_row_hasher {
+  friend class cudf::detail::row::hash::row_hasher;  ///< Allow row_hasher to access private
+                                                     ///< members.
+
+ public:
+  /**
+   * @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
+  {
+    return cudf::detail::accumulate(
+      _table.begin(),
+      _table.end(),
+      _seed,
+      cuda::proclaim_return_type<spark_hash_value_type>(
+        [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
+          return cudf::type_dispatcher(
+            column.type(), element_hasher_adapter<hash_function>{nulls, hash}, column, row_index);
+        }));
+  }
+
+ private:
+  /**
+   * @brief Computes the hash value of an element in the given column.
+   *
+   * When the column is non-nested, this is a simple wrapper around the element_hasher.
+   * When the column is nested, this uses a seed value to serially compute each
+   * nested element, with the output hash becoming the seed for the next value.
+   * This requires constructing a new hash functor for each nested element,
+   * using the new seed from the previous element's hash. The hash of a null
+   * element is the input seed (the previous element's hash).
+   */
+  template <template <typename> class hash_fn>
+  class element_hasher_adapter {
+   public:
+    using hash_functor = cudf::detail::row::hash::element_hasher<hash_fn, Nullate>;
+    using result_type  = typename hash_functor::result_type;
+
+    __device__ element_hasher_adapter(Nullate check_nulls, result_type seed) noexcept
+      : _check_nulls(check_nulls), _seed(seed)
+    {
+    }
+
+    template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
+    __device__ spark_hash_value_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, CUDF_ENABLE_IF(cudf::is_nested<T>())>
+    __device__ spark_hash_value_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 assumed to be decomposed and contain only one child
+          curr_col = cudf::structs_column_device_view(curr_col).get_sliced_child(0);
+        } else if (curr_col.type().id() == type_id::LIST) {
+          curr_col = cudf::lists_column_device_view(curr_col).get_sliced_child();
+        }
+      }
+
+      return cudf::detail::accumulate(
+        thrust::counting_iterator(0),
+        thrust::counting_iterator(curr_col.size()),
+        _seed,
+        [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);
+        });
+    }
+
+    Nullate const _check_nulls;  ///< Whether to check for nulls
+    result_type const _seed;     ///< The seed to use for hashing, also returned for null elements
+  };
+
+  using result_type = typename element_hasher_adapter<hash_function>::result_type;
+
+  CUDF_HOST_DEVICE spark_murmur_device_row_hasher(Nullate check_nulls,
+                                                  table_device_view t,
+                                                  result_type seed = DEFAULT_HASH_SEED) noexcept
+    : _check_nulls{check_nulls}, _table{t}, _seed(seed)
+  {
+    // Error out if passed an unsupported hash_function
+    static_assert(
+      cuda::std::is_same_v<Spark_MurmurHash3_x86_32<int>, hash_function<int>>,
+      "spark_murmur_device_row_hasher only supports the Spark_MurmurHash3_x86_32 hash function");
+  }
+
+  Nullate const _check_nulls;
+  table_device_view const _table;
+  result_type const _seed;
+};
+
+/**
+ * @brief Throws if `input` contains a nested shape unsupported by Spark MurmurHash3.
+ *
+ * @param input Table to validate
+ */
+void check_spark_murmurhash3_compatibility(table_view const& input);
+
+}  // namespace cudf::hashing::detail
diff --git a/cpp/include/cudf/partitioning.hpp b/cpp/include/cudf/partitioning.hpp
index 67cffce3e131..e261ff7613b8 100644
--- a/cpp/include/cudf/partitioning.hpp
+++ b/cpp/include/cudf/partitioning.hpp
@@ -31,7 +31,8 @@ namespace CUDF_EXPORT cudf {
  */
 enum class hash_id : int32_t {
   HASH_IDENTITY = 0,  ///< Identity hash function that simply returns the key to be hashed
-  HASH_MURMUR3        ///< Murmur3 hash function
+  HASH_MURMUR3,       ///< Murmur3 hash function
+  SPARK_MURMUR3       ///< Apache Spark-compatible Murmur3 hash function
 };
 
 /**
diff --git a/cpp/src/hash/murmurhash3_x86_32.cu b/cpp/src/hash/murmurhash3_x86_32.cu
index 0bd88fb3f59f..3a28c1a15553 100644
--- a/cpp/src/hash/murmurhash3_x86_32.cu
+++ b/cpp/src/hash/murmurhash3_x86_32.cu
@@ -17,22 +17,6 @@ namespace cudf {
 namespace hashing {
 namespace detail {
 
-template <template <template <typename> class, typename> class DeviceRowHasher>
-void hash_rows(cudf::detail::row::hash::row_hasher const& row_hasher,
-               size_type num_rows,
-               hash_value_type* output_begin,
-               bool nullable,
-               uint32_t seed,
-               rmm::cuda_stream_view stream)
-{
-  auto const hasher = row_hasher.device_hasher<MurmurHash3_x86_32, DeviceRowHasher>(nullable, seed);
-  // thrust::tabulate is slow here, see NVIDIA/cccl#9070
-  CUDF_CUDA_TRY(cub::DeviceFor::Bulk(
-    num_rows,
-    [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); },
-    stream.value()));
-}
-
 std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
                                            uint32_t seed,
                                            rmm::cuda_stream_view stream,
@@ -46,17 +30,18 @@ std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
 
   if (input.num_rows() == 0) { return output; }
 
-  bool const nullable     = has_nulls(input);
-  auto const row_hasher   = cudf::detail::row::hash::row_hasher(input, stream);
-  auto const output_begin = output->mutable_view().begin<hash_value_type>();
+  bool const nullable   = has_nulls(input);
+  auto const row_hasher = cudf::detail::row::hash::row_hasher(input, stream);
+  auto output_view      = output->mutable_view();
 
-  if (input.num_columns() > 1) {
-    hash_rows<cudf::detail::row::hash::device_row_hasher_iterative>(
-      row_hasher, input.num_rows(), output_begin, nullable, seed, stream);
-  } else {
-    hash_rows<cudf::detail::row::hash::device_row_hasher>(
-      row_hasher, input.num_rows(), output_begin, nullable, seed, stream);
-  }
+  // Compute the hash value for each row
+  auto const output_begin = output_view.begin<hash_value_type>();
+  auto const hasher       = row_hasher.device_hasher<MurmurHash3_x86_32>(nullable, seed);
+  // thrust::tabulate is slow here, see NVIDIA/cccl#9070
+  CUDF_CUDA_TRY(cub::DeviceFor::Bulk(
+    input.num_rows(),
+    [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); },
+    stream.value()));
 
   return output;
 }
diff --git a/cpp/src/hash/spark_murmurhash3_x86_32.cu b/cpp/src/hash/spark_murmurhash3_x86_32.cu
new file mode 100644
index 000000000000..8625b1e49112
--- /dev/null
+++ b/cpp/src/hash/spark_murmurhash3_x86_32.cu
@@ -0,0 +1,89 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+#include <cudf/column/column_factories.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/hashing.hpp>
+#include <cudf/hashing/detail/spark_murmurhash3.cuh>
+#include <cudf/lists/lists_column_view.hpp>
+#include <cudf/utilities/error.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+
+#include <cub/device/device_for.cuh>
+
+#include <functional>
+
+namespace cudf {
+namespace hashing {
+namespace detail {
+
+void check_spark_murmurhash3_compatibility(table_view const& input)
+{
+  using column_checker_fn_t = std::function<void(column_view const&)>;
+
+  column_checker_fn_t check_column = [&](column_view const& c) {
+    if (c.type().id() == type_id::LIST) {
+      auto const& list_col = lists_column_view(c);
+      CUDF_EXPECTS(list_col.child().type().id() != type_id::STRUCT,
+                   "Cannot compute hash of a table with a LIST of STRUCT columns.");
+      check_column(list_col.child());
+    } else if (c.type().id() == type_id::STRUCT) {
+      for (auto child = c.child_begin(); child != c.child_end(); ++child) {
+        check_column(*child);
+      }
+    }
+  };
+
+  for (column_view const& c : input) {
+    check_column(c);
+  }
+}
+
+std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
+                                                 uint32_t seed,
+                                                 rmm::cuda_stream_view stream,
+                                                 rmm::device_async_resource_ref mr)
+{
+  auto output = make_numeric_column(data_type(type_to_id<spark_hash_value_type>()),
+                                    input.num_rows(),
+                                    mask_state::UNALLOCATED,
+                                    stream,
+                                    mr);
+
+  // Return early if there's nothing to hash
+  if (input.num_rows() == 0) { return output; }
+
+  // Lists of structs are not supported
+  check_spark_murmurhash3_compatibility(input);
+
+  bool const nullable     = has_nested_nulls(input);
+  auto const row_hasher   = cudf::detail::row::hash::row_hasher(input, stream);
+  auto const output_begin = output->mutable_view().begin<spark_hash_value_type>();
+
+  // Compute the hash value for each row
+  auto const hasher =
+    row_hasher.device_hasher<Spark_MurmurHash3_x86_32, spark_murmur_device_row_hasher>(nullable,
+                                                                                       seed);
+  CUDF_CUDA_TRY(cub::DeviceFor::Bulk(
+    input.num_rows(),
+    [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); },
+    stream.value()));
+
+  return output;
+}
+
+}  // namespace detail
+
+std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
+                                                 uint32_t seed,
+                                                 rmm::cuda_stream_view stream,
+                                                 rmm::device_async_resource_ref mr)
+{
+  CUDF_FUNC_RANGE();
+  return detail::spark_murmurhash3_x86_32(input, seed, stream, mr);
+}
+
+}  // namespace hashing
+}  // namespace cudf
diff --git a/cpp/src/partitioning/partitioning.cu b/cpp/src/partitioning/partitioning.cu
index d9bf0d99864f..2ac3d673bb70 100644
--- a/cpp/src/partitioning/partitioning.cu
+++ b/cpp/src/partitioning/partitioning.cu
@@ -15,6 +15,7 @@
 #include <cudf/detail/utilities/integer_utils.hpp>
 #include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/hashing/detail/murmurhash3_x86_32.cuh>
+#include <cudf/hashing/detail/spark_murmurhash3.cuh>
 #include <cudf/partitioning.hpp>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/utilities/default_stream.hpp>
@@ -565,7 +566,10 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table_g
 }
 
 // NOTE hash_has_nulls must be true if table_to_hash has nulls
-template <template <typename> class hash_function, bool hash_has_nulls>
+template <template <typename> class hash_function,
+          bool hash_has_nulls,
+          template <template <typename> class, typename> class DeviceRowHasher =
+            detail::row::hash::device_row_hasher>
 std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
   table_view const& input,
   table_view const& table_to_hash,
@@ -577,8 +581,8 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
   auto const num_rows = table_to_hash.num_rows();
 
   auto const row_hasher = detail::row::hash::row_hasher(table_to_hash, stream);
-  auto const hasher =
-    row_hasher.device_hasher<hash_function>(nullate::DYNAMIC{hash_has_nulls}, seed);
+  auto const hasher     = row_hasher.device_hasher<hash_function, DeviceRowHasher>(
+    nullate::DYNAMIC{hash_has_nulls}, seed);
 
   // Check whether the per-block shared memory histograms fit in shared memory
   int dev;
@@ -871,7 +875,9 @@ struct IdentityHash {
   }
 };
 
-template <template <typename> class hash_function>
+template <template <typename> class hash_function,
+          template <template <typename> class, typename> class DeviceRowHasher =
+            detail::row::hash::device_row_hasher>
 std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition(
   table_view const& input,
   table_view const& table_to_hash,
@@ -891,10 +897,10 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition(
     }
   }
   if (has_nested_nulls(table_to_hash)) {
-    return hash_partition_table<hash_function, true>(
+    return hash_partition_table<hash_function, true, DeviceRowHasher>(
       input, table_to_hash, num_partitions, seed, stream, mr);
   } else {
-    return hash_partition_table<hash_function, false>(
+    return hash_partition_table<hash_function, false, DeviceRowHasher>(
       input, table_to_hash, num_partitions, seed, stream, mr);
   }
 }
@@ -939,6 +945,13 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition(
     case (hash_id::HASH_MURMUR3):
       return hash_partition<cudf::hashing::detail::MurmurHash3_x86_32>(
         input, keys, num_partitions, seed, stream, mr);
+    case (hash_id::SPARK_MURMUR3):
+      if (keys.num_rows() > 0 && keys.num_columns() > 0) {
+        cudf::hashing::detail::check_spark_murmurhash3_compatibility(keys);
+      }
+      return hash_partition<cudf::hashing::detail::Spark_MurmurHash3_x86_32,
+                            cudf::hashing::detail::spark_murmur_device_row_hasher>(
+        input, keys, num_partitions, seed, stream, mr);
     default: CUDF_FAIL("Unsupported hash function in hash_partition");
   }
 }
diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt
index ff214d35562a..0f2337ca52a7 100644
--- a/cpp/tests/CMakeLists.txt
+++ b/cpp/tests/CMakeLists.txt
@@ -198,6 +198,7 @@ ConfigureTest(
   hashing/md5_test.cpp
   hashing/murmurhash3_x86_32_test.cpp
   hashing/murmurhash3_x64_128_test.cpp
+  hashing/spark_murmurhash3_x86_32_test.cpp
   hashing/sha1_test.cpp
   hashing/sha224_test.cpp
   hashing/sha256_test.cpp
diff --git a/cpp/tests/hashing/murmurhash3_x86_32_test.cpp b/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
index e83756ebbb6d..bb0b5cc41efa 100644
--- a/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
+++ b/cpp/tests/hashing/murmurhash3_x86_32_test.cpp
@@ -50,20 +50,6 @@ TEST_F(MurmurHashTest, MultiValue)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(output1->view(), output2->view());
 }
 
-TEST_F(MurmurHashTest, IterativeSeeding)
-{
-  cudf::test::fixed_width_column_wrapper<int32_t> const first({0, 1, -1, 42, 123456789});
-  cudf::test::fixed_width_column_wrapper<int32_t> const second({10, 20, 30, -40, -987654321});
-  auto const input = cudf::table_view({first, second});
-
-  auto const output = cudf::hashing::murmurhash3_x86_32(input, 42);
-
-  // Matches Apache Spark for non-null fixed-width inputs with seed 42.
-  cudf::test::fixed_width_column_wrapper<uint32_t> const expected(
-    {2573243963u, 1151116018u, 1549484878u, 3007216400u, 2314233967u});
-  CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected);
-}
-
 TEST_F(MurmurHashTest, MultiValueNulls)
 {
   // Nulls with different values should be equal
diff --git a/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
new file mode 100644
index 000000000000..056fa28f374d
--- /dev/null
+++ b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
@@ -0,0 +1,579 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf_test/base_fixture.hpp>
+#include <cudf_test/column_utilities.hpp>
+#include <cudf_test/column_wrapper.hpp>
+#include <cudf_test/iterator_utilities.hpp>
+#include <cudf_test/type_lists.hpp>
+
+#include <cudf/detail/iterator.cuh>
+#include <cudf/fixed_point/fixed_point.hpp>
+#include <cudf/hashing.hpp>
+
+constexpr cudf::test::debug_output_level verbosity{cudf::test::debug_output_level::ALL_ERRORS};
+
+template <typename T>
+class SparkMurmurHashTestTyped : public cudf::test::BaseFixture {};
+
+TYPED_TEST_SUITE(SparkMurmurHashTestTyped, cudf::test::FixedWidthTypes);
+
+TYPED_TEST(SparkMurmurHashTestTyped, Equality)
+{
+  cudf::test::fixed_width_column_wrapper<TypeParam, int32_t> const col{0, 127, 1, 2, 8};
+  auto const input = cudf::table_view({col});
+
+  // Hash of same input should be equal
+  auto const spark_output1 = cudf::hashing::spark_murmurhash3_x86_32(input, 0);
+  auto const spark_output2 = cudf::hashing::spark_murmurhash3_x86_32(input);
+
+  EXPECT_EQ(input.num_rows(), spark_output1->size());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(spark_output1->view(), spark_output2->view());
+}
+
+TYPED_TEST(SparkMurmurHashTestTyped, EqualityNulls)
+{
+  using T = TypeParam;
+
+  // Nulls with different values should be equal
+  cudf::test::fixed_width_column_wrapper<T, int32_t> const col1({0, 127, 1, 2, 8}, {0, 1, 1, 1, 1});
+  cudf::test::fixed_width_column_wrapper<T, int32_t> const col2({1, 127, 1, 2, 8}, {0, 1, 1, 1, 1});
+
+  auto const input1 = cudf::table_view({col1});
+  auto const input2 = cudf::table_view({col2});
+
+  auto const spark_output1 = cudf::hashing::spark_murmurhash3_x86_32(input1, 0);
+  auto const spark_output2 = cudf::hashing::spark_murmurhash3_x86_32(input2);
+
+  EXPECT_EQ(input1.num_rows(), spark_output1->size());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(spark_output1->view(), spark_output2->view());
+}
+
+template <typename T>
+class SparkMurmurHashTestFloatTyped : public cudf::test::BaseFixture {};
+
+TYPED_TEST_SUITE(SparkMurmurHashTestFloatTyped, cudf::test::FloatingPointTypes);
+
+TYPED_TEST(SparkMurmurHashTestFloatTyped, TestExtremes)
+{
+  using T = TypeParam;
+  T min   = std::numeric_limits<T>::min();
+  T max   = std::numeric_limits<T>::max();
+  T nan   = std::numeric_limits<T>::quiet_NaN();
+  T inf   = std::numeric_limits<T>::infinity();
+
+  cudf::test::fixed_width_column_wrapper<T> const col(
+    {T(0.0), T(100.0), T(-100.0), min, max, nan, inf, -inf});
+  cudf::test::fixed_width_column_wrapper<T> const col_neg_zero(
+    {T(-0.0), T(100.0), T(-100.0), min, max, nan, inf, -inf});
+  cudf::test::fixed_width_column_wrapper<T> const col_neg_nan(
+    {T(0.0), T(100.0), T(-100.0), min, max, -nan, inf, -inf});
+
+  auto const table_col          = cudf::table_view({col});
+  auto const table_col_neg_zero = cudf::table_view({col_neg_zero});
+  auto const table_col_neg_nan  = cudf::table_view({col_neg_nan});
+
+  // Spark hash is sensitive to 0 and -0
+  auto const spark_col         = cudf::hashing::spark_murmurhash3_x86_32(table_col, 0);
+  auto const spark_col_neg_nan = cudf::hashing::spark_murmurhash3_x86_32(table_col_neg_nan);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*spark_col, *spark_col_neg_nan);
+}
+
+class SparkMurmurHashTest : public cudf::test::BaseFixture {};
+
+TEST_F(SparkMurmurHashTest, IterativeSeeding)
+{
+  cudf::test::fixed_width_column_wrapper<int32_t> const first({0, 1, -1, 42, 123456789});
+  cudf::test::fixed_width_column_wrapper<int32_t> const second({10, 20, 30, -40, -987654321});
+  auto const input = cudf::table_view({first, second});
+
+  auto const output = cudf::hashing::spark_murmurhash3_x86_32(input, 42);
+
+  // Values produced by Apache Spark for the same input and seed.
+  cudf::test::fixed_width_column_wrapper<int32_t> const expected(
+    {-1721723333, 1151116018, 1549484878, -1287750896, -1980733329});
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(output->view(), expected);
+}
+
+TEST_F(SparkMurmurHashTest, MultiValueNulls)
+{
+  // Nulls with different values should be equal
+  cudf::test::strings_column_wrapper const strings_col1(
+    {"",
+     "The quick brown fox",
+     "jumps over the lazy dog.",
+     "All work and no play makes Jack a dull boy",
+     R"(!"#$%&'()*+,-./0123456789:;<=>?@[\]^_`{|}~)"},
+    {0, 1, 1, 0, 1});
+  cudf::test::strings_column_wrapper const strings_col2(
+    {"different but null",
+     "The quick brown fox",
+     "jumps over the lazy dog.",
+     "I am Jack's complete lack of null value",
+     R"(!"#$%&'()*+,-./0123456789:;<=>?@[\]^_`{|}~)"},
+    {0, 1, 1, 0, 1});
+
+  // Nulls with different values should be equal
+  using limits = std::numeric_limits<int32_t>;
+  cudf::test::fixed_width_column_wrapper<int32_t> const ints_col1(
+    {0, 100, -100, limits::min(), limits::max()}, {1, 0, 0, 1, 1});
+  cudf::test::fixed_width_column_wrapper<int32_t> const ints_col2(
+    {0, -200, 200, limits::min(), limits::max()}, {1, 0, 0, 1, 1});
+
+  // Nulls with different values should be equal
+  // Different truth values should be equal
+  cudf::test::fixed_width_column_wrapper<bool> const bools_col1({0, 1, 0, 1, 1}, {1, 1, 0, 0, 1});
+  cudf::test::fixed_width_column_wrapper<bool> const bools_col2({0, 2, 1, 0, 255}, {1, 1, 0, 0, 1});
+
+  // Nulls with different values should be equal
+  using ts = cudf::timestamp_s;
+  cudf::test::fixed_width_column_wrapper<ts, ts::duration> const secs_col1(
+    {ts::duration::zero(),
+     static_cast<ts::duration>(100),
+     static_cast<ts::duration>(-100),
+     ts::duration::min(),
+     ts::duration::max()},
+    {1, 0, 0, 1, 1});
+  cudf::test::fixed_width_column_wrapper<ts, ts::duration> const secs_col2(
+    {ts::duration::zero(),
+     static_cast<ts::duration>(-200),
+     static_cast<ts::duration>(200),
+     ts::duration::min(),
+     ts::duration::max()},
+    {1, 0, 0, 1, 1});
+
+  auto const input1        = cudf::table_view({strings_col1, ints_col1, bools_col1, secs_col1});
+  auto const input2        = cudf::table_view({strings_col2, ints_col2, bools_col2, secs_col2});
+  auto const spark_output1 = cudf::hashing::spark_murmurhash3_x86_32(input1, 0);
+  auto const spark_output2 = cudf::hashing::spark_murmurhash3_x86_32(input2);
+
+  EXPECT_EQ(input1.num_rows(), spark_output1->size());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(spark_output1->view(), spark_output2->view());
+}
+
+TEST_F(SparkMurmurHashTest, MultiValueWithSeeds)
+{
+  // The hash values were determined by running the following Scala code in Apache Spark.
+  // Note that Spark >= 3.2 normalizes the float/double value of -0. to +0. and both values hash
+  // to the same result. This is normalized in the calling code (Spark RAPIDS plugin) for Spark
+  // >= 3.2. However, the reference values for -0. below must be obtained with Spark < 3.2 and
+  // libcudf will continue to implement the Spark < 3.2 behavior until Spark >= 3.2 is required and
+  // the workaround in the calling code is removed. This also affects the combined hash values.
+
+  /*
+  import org.apache.spark.sql.functions._
+  import org.apache.spark.sql.types._
+  import org.apache.spark.sql.Row
+  import org.apache.spark.sql.catalyst.util.DateTimeUtils
+
+  val schema = new StructType()
+      .add("structs", new StructType()
+          .add("a", IntegerType)
+          .add("b", StringType)
+          .add("c", new StructType()
+              .add("x", FloatType)
+              .add("y", LongType)))
+      .add("strings", StringType)
+      .add("doubles", DoubleType)
+      .add("timestamps", TimestampType)
+      .add("decimal64", DecimalType(18, 7))
+      .add("longs", LongType)
+      .add("floats", FloatType)
+      .add("dates", DateType)
+      .add("decimal32", DecimalType(9, 3))
+      .add("ints", IntegerType)
+      .add("shorts", ShortType)
+      .add("bytes", ByteType)
+      .add("bools", BooleanType)
+      .add("decimal128", DecimalType(38, 11))
+
+  val data = Seq(
+      Row(Row(0, "a", Row(0f, 0L)), "", 0.toDouble,
+          DateTimeUtils.toJavaTimestamp(0), BigDecimal(0), 0.toLong, 0.toFloat,
+          DateTimeUtils.toJavaDate(0), BigDecimal(0), 0, 0.toShort, 0.toByte,
+          false, BigDecimal(0)),
+      Row(Row(100, "bc", Row(100f, 100L)), "The quick brown fox", -(0.toDouble),
+          DateTimeUtils.toJavaTimestamp(100), BigDecimal("0.00001"), 100.toLong, -(0.toFloat),
+          DateTimeUtils.toJavaDate(100), BigDecimal("0.1"), 100, 100.toShort, 100.toByte,
+          true, BigDecimal("0.000000001")),
+      Row(Row(-100, "def", Row(-100f, -100L)), "jumps over the lazy dog.", -Double.NaN,
+          DateTimeUtils.toJavaTimestamp(-100), BigDecimal("-0.00001"), -100.toLong, -Float.NaN,
+          DateTimeUtils.toJavaDate(-100), BigDecimal("-0.1"), -100, -100.toShort, -100.toByte,
+          true, BigDecimal("-0.00000000001")),
+      Row(Row(0x12345678, "ghij", Row(Float.PositiveInfinity, 0x123456789abcdefL)),
+          "All work and no play makes Jack a dull boy", Double.MinValue,
+          DateTimeUtils.toJavaTimestamp(Long.MinValue/1000000), BigDecimal("-99999999999.9999999"),
+          Long.MinValue, Float.MinValue, DateTimeUtils.toJavaDate(Int.MinValue/100),
+          BigDecimal("-999999.999"), Int.MinValue, Short.MinValue, Byte.MinValue, true,
+          BigDecimal("-9999999999999999.99999999999")),
+      Row(Row(-0x76543210, "klmno", Row(Float.NegativeInfinity, -0x123456789abcdefL)),
+          "!\"#$%&\'()*+,-./:;<=>?@[\\]^_`{|}~\ud720\ud721", Double.MaxValue,
+          DateTimeUtils.toJavaTimestamp(Long.MaxValue/1000000), BigDecimal("99999999999.9999999"),
+          Long.MaxValue, Float.MaxValue, DateTimeUtils.toJavaDate(Int.MaxValue/100),
+          BigDecimal("999999.999"), Int.MaxValue, Short.MaxValue, Byte.MaxValue, false,
+          BigDecimal("99999999999999999999999999.99999999999")))
+
+  val df = spark.createDataFrame(sc.parallelize(data), schema)
+  df.columns.foreach(c => println(s"$c => ${df.select(hash(col(c))).collect.mkString(",")}"))
+  println(s"combined => ${df.select(hash(col("*"))).collect.mkString(",")}")
+  */
+
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_structs_expected(
+    {-105406170, 90479889, -678041645, 1667387937, 301478567});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_strings_expected(
+    {142593372, 1217302703, -715697185, -2061143941, -111635966});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_doubles_expected(
+    {-1670924195, -853646085, -1281358385, 1897734433, -508695674});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_timestamps_expected(
+    {-1670924195, 1114849490, 904948192, -1832979433, 1752430209});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_decimal64_expected(
+    {-1670924195, 1114849490, 904948192, 1962370902, -1795328666});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_longs_expected(
+    {-1670924195, 1114849490, 904948192, -853646085, -1604625029});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_floats_expected(
+    {933211791, 723455942, -349261430, -1225560532, -338752985});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_dates_expected(
+    {933211791, 751823303, -1080202046, -1906567553, -1503850410});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_decimal32_expected(
+    {-1670924195, 1114849490, 904948192, -1454351396, -193774131});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_ints_expected(
+    {933211791, 751823303, -1080202046, 723455942, 133916647});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_shorts_expected(
+    {933211791, 751823303, -1080202046, -1871935946, 1249274084});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_bytes_expected(
+    {933211791, 751823303, -1080202046, 1110053733, 1135925485});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_bools_expected(
+    {933211791, -559580957, -559580957, -559580957, 933211791});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_decimal128_expected(
+    {-783713497, -295670906, 1398487324, -52622807, -1359749815});
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_combined_expected(
+    {401603227, 588162166, 552160517, 1132537411, -326043017});
+
+  using double_limits = std::numeric_limits<double>;
+  using long_limits   = std::numeric_limits<int64_t>;
+  using float_limits  = std::numeric_limits<float>;
+  using int_limits    = std::numeric_limits<int32_t>;
+  cudf::test::fixed_width_column_wrapper<int32_t> a_col{0, 100, -100, 0x1234'5678, -0x7654'3210};
+  cudf::test::strings_column_wrapper b_col{"a", "bc", "def", "ghij", "klmno"};
+  cudf::test::fixed_width_column_wrapper<float> x_col{
+    0.f, 100.f, -100.f, float_limits::infinity(), -float_limits::infinity()};
+  cudf::test::fixed_width_column_wrapper<int64_t> y_col{
+    0L, 100L, -100L, 0x0123'4567'89ab'cdefL, -0x0123'4567'89ab'cdefL};
+  cudf::test::structs_column_wrapper c_col{{x_col, y_col}};
+  cudf::test::structs_column_wrapper const structs_col{{a_col, b_col, c_col}};
+
+  cudf::test::strings_column_wrapper const strings_col(
+    {"",
+     "The quick brown fox",
+     "jumps over the lazy dog.",
+     "All work and no play makes Jack a dull boy",
+     "!\"#$%&\'()*+,-./:;<=>?@[\\]^_`{|}~\ud720\ud721"});
+  cudf::test::fixed_width_column_wrapper<double> const doubles_col(
+    {0., -0., -double_limits::quiet_NaN(), double_limits::lowest(), double_limits::max()});
+  cudf::test::fixed_width_column_wrapper<cudf::timestamp_ms, cudf::timestamp_ms::rep> const
+    timestamps_col({0L, 100L, -100L, long_limits::min() / 1000000, long_limits::max() / 1000000});
+  cudf::test::fixed_point_column_wrapper<int64_t> const decimal64_col(
+    {0L, 100L, -100L, -999999999999999999L, 999999999999999999L}, numeric::scale_type{-7});
+  cudf::test::fixed_width_column_wrapper<int64_t> const longs_col(
+    {0L, 100L, -100L, long_limits::min(), long_limits::max()});
+  cudf::test::fixed_width_column_wrapper<float> const floats_col(
+    {0.f, -0.f, -float_limits::quiet_NaN(), float_limits::lowest(), float_limits::max()});
+  cudf::test::fixed_width_column_wrapper<cudf::timestamp_D, cudf::timestamp_D::rep> dates_col(
+    {0, 100, -100, int_limits::min() / 100, int_limits::max() / 100});
+  cudf::test::fixed_point_column_wrapper<int32_t> const decimal32_col(
+    {0, 100, -100, -999999999, 999999999}, numeric::scale_type{-3});
+  cudf::test::fixed_width_column_wrapper<int32_t> const ints_col(
+    {0, 100, -100, int_limits::min(), int_limits::max()});
+  cudf::test::fixed_width_column_wrapper<int16_t> const shorts_col({0, 100, -100, -32768, 32767});
+  cudf::test::fixed_width_column_wrapper<int8_t> const bytes_col({0, 100, -100, -128, 127});
+  cudf::test::fixed_width_column_wrapper<bool> const bools_col1({0, 1, 1, 1, 0});
+  cudf::test::fixed_width_column_wrapper<bool> const bools_col2({0, 1, 2, 255, 0});
+  cudf::test::fixed_point_column_wrapper<__int128_t> const decimal128_col(
+    {static_cast<__int128>(0),
+     static_cast<__int128>(100),
+     static_cast<__int128>(-1),
+     (static_cast<__int128>(0xFFFF'FFFF'FCC4'D1C3u) << 64 | 0x602F'7FC3'1800'0001u),
+     (static_cast<__int128>(0x0785'EE10'D5DA'46D9u) << 64 | 0x00F4'369F'FFFF'FFFFu)},
+    numeric::scale_type{-11});
+
+  auto const hash_structs =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({structs_col}), 42);
+  auto const hash_strings =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({strings_col}), 42);
+  auto const hash_doubles =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({doubles_col}), 42);
+  auto const hash_timestamps =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({timestamps_col}), 42);
+  auto const hash_decimal64 =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({decimal64_col}), 42);
+  auto const hash_longs =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({longs_col}), 42);
+  auto const hash_floats =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({floats_col}), 42);
+  auto const hash_dates =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({dates_col}), 42);
+  auto const hash_decimal32 =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({decimal32_col}), 42);
+  auto const hash_ints = cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({ints_col}), 42);
+  auto const hash_shorts =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({shorts_col}), 42);
+  auto const hash_bytes =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({bytes_col}), 42);
+  auto const hash_bools1 =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({bools_col1}), 42);
+  auto const hash_bools2 =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({bools_col2}), 42);
+  auto const hash_decimal128 =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({decimal128_col}), 42);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_structs, hash_structs_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_strings, hash_strings_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_doubles, hash_doubles_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_timestamps, hash_timestamps_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_decimal64, hash_decimal64_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_longs, hash_longs_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_floats, hash_floats_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_dates, hash_dates_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_decimal32, hash_decimal32_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_ints, hash_ints_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_shorts, hash_shorts_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_bytes, hash_bytes_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_bools1, hash_bools_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_bools2, hash_bools_expected, verbosity);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_decimal128, hash_decimal128_expected, verbosity);
+
+  auto const combined_table = cudf::table_view({structs_col,
+                                                strings_col,
+                                                doubles_col,
+                                                timestamps_col,
+                                                decimal64_col,
+                                                longs_col,
+                                                floats_col,
+                                                dates_col,
+                                                decimal32_col,
+                                                ints_col,
+                                                shorts_col,
+                                                bytes_col,
+                                                bools_col2,
+                                                decimal128_col});
+  auto const hash_combined  = cudf::hashing::spark_murmurhash3_x86_32(combined_table, 42);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_combined, hash_combined_expected, verbosity);
+}
+
+TEST_F(SparkMurmurHashTest, StringsWithSeed)
+{
+  // The hash values were determined by running the following Scala code in Apache Spark:
+  // val strs = Seq("", "The quick brown fox",
+  //              "jumps over the lazy dog.",
+  //              "All work and no play makes Jack a dull boy",
+  //              "!\"#$%&\'()*+,-./:;<=>?@[\\]^_`{|}~\ud720\ud721")
+  // println(strs.map(org.apache.spark.unsafe.types.UTF8String.fromString)
+  //   .map(org.apache.spark.sql.catalyst.expressions.Murmur3HashFunction.hash(
+  //     _, org.apache.spark.sql.types.StringType, 314)))
+
+  cudf::test::fixed_width_column_wrapper<int32_t> const hash_strings_expected_seed_314(
+    {1467149710, 723257560, -1620282500, -2001858707, 1588473657});
+
+  cudf::test::strings_column_wrapper const strings_col(
+    {"",
+     "The quick brown fox",
+     "jumps over the lazy dog.",
+     "All work and no play makes Jack a dull boy",
+     "!\"#$%&\'()*+,-./:;<=>?@[\\]^_`{|}~\ud720\ud721"});
+
+  auto const hash_strings =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({strings_col}), 314);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_strings, hash_strings_expected_seed_314, verbosity);
+}
+
+TEST_F(SparkMurmurHashTest, ListValues)
+{
+  /*
+  import org.apache.spark.sql.functions._
+  import org.apache.spark.sql.types.{ArrayType, IntegerType, StructType}
+  import org.apache.spark.sql.Row
+
+  val schema = new StructType()
+    .add("lists",ArrayType(ArrayType(IntegerType)))
+
+  val data = Seq(
+    Row(null),
+    Row(List(null)),
+    Row(List(List())),
+    Row(List(List(1))),
+    Row(List(List(1, 2))),
+    Row(List(List(1, 2, 3))),
+    Row(List(List(1, 2), List(3))),
+    Row(List(List(1), List(2, 3))),
+    Row(List(List(1), List(null, 2, 3))),
+    Row(List(List(1, 2), List(3), List(null))),
+    Row(List(List(1, 2), null, List(3))),
+  )
+
+  val df = spark.createDataFrame(
+    spark.sparkContext.parallelize(data), schema)
+
+  val df2 = df.selectExpr("lists", "hash(lists) as hash")
+  df2.printSchema()
+  df2.show(false)
+  */
+
+  auto const null = -1;
+  auto nested_list =
+    cudf::test::lists_column_wrapper<int>({{},
+                                           {1},
+                                           {1, 2},
+                                           {1, 2, 3},
+                                           {1, 2},
+                                           {3},
+                                           {1},
+                                           {2, 3},
+                                           {1},
+                                           {{null, 2, 3}, cudf::test::iterators::nulls_at({0})},
+                                           {1, 2},
+                                           {3},
+                                           {{null}, cudf::test::iterators::nulls_at({0})},
+                                           {1, 2},
+                                           {},
+                                           {3}},
+                                          cudf::test::iterators::nulls_at({0, 14}));
+  auto offsets =
+    cudf::test::fixed_width_column_wrapper<cudf::size_type>{0, 0, 0, 1, 2, 3, 4, 6, 8, 10, 13, 16};
+  auto list_validity = cudf::test::iterators::nulls_at({0});
+  auto [null_mask, null_count] =
+    cudf::test::detail::make_null_mask(list_validity, list_validity + 11);
+  auto list_column = cudf::make_lists_column(
+    11, offsets.release(), nested_list.release(), null_count, std::move(null_mask));
+
+  auto expect = cudf::test::fixed_width_column_wrapper<int32_t>{42,
+                                                                42,
+                                                                42,
+                                                                -559580957,
+                                                                -222940379,
+                                                                -912918097,
+                                                                -912918097,
+                                                                -912918097,
+                                                                -912918097,
+                                                                -912918097,
+                                                                -912918097};
+
+  auto output = cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({*list_column}), 42);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
+}
+
+TEST_F(SparkMurmurHashTest, StructOfListValues)
+{
+  /*
+  import org.apache.spark.sql.functions._
+  import org.apache.spark.sql.types.{ArrayType, IntegerType, StructType}
+  import org.apache.spark.sql.Row
+
+  val schema = new StructType()
+    .add("structs", new StructType()
+        .add("a", ArrayType(IntegerType))
+        .add("b", ArrayType(IntegerType)))
+
+  val data = Seq(
+    Row(Row(List(), List())),
+    Row(Row(List(0), List(0))),
+    Row(Row(List(1, null), null)),
+    Row(Row(List(1, null), List())),
+    Row(Row(List(), List(null, 1))),
+    Row(Row(null, List(1))),
+    Row(Row(List(2, 3), List(4, 5))),
+  )
+
+  val df = spark.createDataFrame(
+    spark.sparkContext.parallelize(data), schema)
+
+  val df2 = df.selectExpr("lists", "hash(lists) as hash")
+  df2.printSchema()
+  df2.show(false)
+  */
+
+  auto const null = -1;
+  auto col1 =
+    cudf::test::lists_column_wrapper<int>({{},
+                                           {0},
+                                           {{1, null}, cudf::test::iterators::nulls_at({1})},
+                                           {{1, null}, cudf::test::iterators::nulls_at({1})},
+                                           {},
+                                           {} /*NULL*/,
+                                           {2, 3}},
+                                          cudf::test::iterators::nulls_at({5}));
+  auto col2 = cudf::test::lists_column_wrapper<int>(
+    {{}, {0}, {} /*NULL*/, {}, {{null, 1}, cudf::test::iterators::nulls_at({0})}, {1}, {4, 5}},
+    cudf::test::iterators::nulls_at({2}));
+  auto struct_column = cudf::test::structs_column_wrapper{{col1, col2}};
+
+  auto expect = cudf::test::fixed_width_column_wrapper<int32_t>{
+    42, 59727262, -559580957, -559580957, -559580957, -559580957, 170038658};
+
+  auto output = cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({struct_column}), 42);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
+}
+
+TEST_F(SparkMurmurHashTest, ListOfStructValues)
+{
+  /*
+  import org.apache.spark.sql.functions._
+  import org.apache.spark.sql.types.{ArrayType, IntegerType, StructType}
+  import org.apache.spark.sql.Row
+
+  val schema = new StructType()
+    .add("lists", ArrayType(new StructType()
+      .add("a", IntegerType)
+      .add("b", IntegerType)))
+
+  val data = Seq(
+    Row(List(Row(0, 0))),
+    Row(List(null)),
+    Row(List(Row(null, null))),
+    Row(List(Row(1, null))),
+    Row(List(Row(null, 1))),
+    Row(List(Row(null, 1), Row(2, 3))),
+    Row(List(Row(2, 3), null)),
+    Row(List(Row(2, 3), Row(4, 5))),
+  )
+
+  val df = spark.createDataFrame(
+    spark.sparkContext.parallelize(data), schema)
+
+  val df2 = df.selectExpr("lists", "hash(lists) as hash")
+  df2.printSchema()
+  df2.show(false)
+  */
+
+  auto const null = -1;
+  auto col1       = cudf::test::fixed_width_column_wrapper<int32_t>(
+    {0, null, null, 1, null, null, 2, 2, null, 2, 4},
+    cudf::test::iterators::nulls_at({1, 2, 4, 5, 8}));
+  auto col2 = cudf::test::fixed_width_column_wrapper<int32_t>(
+    {0, null, null, null, 1, 1, 3, 3, null, 3, 5}, cudf::test::iterators::nulls_at({1, 2, 3, 8}));
+  auto struct_column =
+    cudf::test::structs_column_wrapper{{col1, col2}, {1, 0, 1, 1, 1, 1, 1, 1, 0, 1, 1}};
+  auto offsets =
+    cudf::test::fixed_width_column_wrapper<cudf::size_type>{0, 1, 2, 3, 4, 5, 7, 9, 11};
+  auto list_nullmask = std::vector<bool>(1, 8);
+  auto [null_mask, null_count] =
+    cudf::test::detail::make_null_mask(list_nullmask.begin(), list_nullmask.end());
+  auto list_column = cudf::make_lists_column(
+    8, offsets.release(), struct_column.release(), null_count, std::move(null_mask));
+
+  // TODO: Lists of structs are not yet supported. Once support is added,
+  // remove this EXPECT_THROW and uncomment the rest of this test.
+  EXPECT_THROW(cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({*list_column}), 42),
+               cudf::logic_error);
+
+  /*
+  auto expect = cudf::test::fixed_width_column_wrapper<int32_t>{
+    59727262, 42, 42, -559580957, -559580957, -912918097, 1092624418, 170038658};
+
+  auto output = cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({*list_column}), 42);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
+  */
+}
diff --git a/cpp/tests/partitioning/hash_partition_test.cpp b/cpp/tests/partitioning/hash_partition_test.cpp
index 0a9d4f4992fc..4469325ab599 100644
--- a/cpp/tests/partitioning/hash_partition_test.cpp
+++ b/cpp/tests/partitioning/hash_partition_test.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #include <cudf_test/base_fixture.hpp>
@@ -366,6 +366,24 @@ TEST_F(HashPartition, CustomSeedValue)
   CUDF_TEST_EXPECT_TABLES_EQUAL(output1->view(), output2->view());
 }
 
+TEST_F(HashPartition, SparkMurmur3)
+{
+  fixed_width_column_wrapper<int32_t> const first({0, 1, -1, 42, 123456789});
+  fixed_width_column_wrapper<int32_t> const second({10, 20, 30, -40, -987654321});
+  auto const input = cudf::table_view({first, second});
+
+  auto [output, offsets] = cudf::hash_partition(
+    input, std::vector<cudf::size_type>{0, 1}, 8, cudf::hash_id::SPARK_MURMUR3, 42);
+
+  fixed_width_column_wrapper<int32_t> const expected_first({42, 1, 0, -1, 123456789});
+  fixed_width_column_wrapper<int32_t> const expected_second({-40, 20, 10, 30, -987654321});
+  auto const expected = cudf::table_view({expected_first, expected_second});
+  std::vector<cudf::size_type> const expected_offsets({0, 1, 1, 2, 3, 3, 3, 4, 5});
+
+  EXPECT_EQ(offsets, expected_offsets);
+  CUDF_TEST_EXPECT_TABLES_EQUAL(output->view(), expected);
+}
+
 template <typename T>
 class HashPartitionFixedWidth : public cudf::test::BaseFixture {};
 
diff --git a/python/cudf/cudf/core/indexed_frame.py b/python/cudf/cudf/core/indexed_frame.py
index 699657c8a38d..8a1db59fa71a 100644
--- a/python/cudf/cudf/core/indexed_frame.py
+++ b/python/cudf/cudf/core/indexed_frame.py
@@ -3021,9 +3021,9 @@ def hash_values(
         1  120  0.25
         2   30  0.50
         >>> df.hash_values(method="murmur3")
-        0    3826000917
-        1     657073880
-        2     339499243
+        0      92443192
+        1     116637743
+        2    2602001908
         dtype: uint32
         >>> df.hash_values(method="md5")
         0    57ce879751b5169c525907d5c563fae1
diff --git a/python/cudf/cudf/tests/dataframe/methods/test_hash_values.py b/python/cudf/cudf/tests/dataframe/methods/test_hash_values.py
index 92a64b6258d5..5ca7761c4bc1 100644
--- a/python/cudf/cudf/tests/dataframe/methods/test_hash_values.py
+++ b/python/cudf/cudf/tests/dataframe/methods/test_hash_values.py
@@ -1,4 +1,4 @@
-# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION.
 # SPDX-License-Identifier: Apache-2.0
 
 import contextlib
@@ -86,21 +86,6 @@ def test_dataframe_hash_values_seed(method):
     assert_neq(out_one, out_two)
 
 
-def test_dataframe_hash_values_murmur3_iterative_seeding():
-    gdf = cudf.DataFrame(
-        {
-            "a": cudf.Series([0, 1, -1, 42, 123456789], dtype=np.int32),
-            "b": cudf.Series([10, 20, 30, -40, -987654321], dtype=np.int32),
-        }
-    )
-    expected = cudf.Series(
-        [2573243963, 1151116018, 1549484878, 3007216400, 2314233967],
-        dtype=np.uint32,
-    )
-
-    assert_eq(gdf.hash_values(method="murmur3", seed=42), expected)
-
-
 def test_dataframe_hash_values_xxhash32():
     # xxhash32 has no built-in implementation in Python and we don't want to
     # add a testing dependency, so we use regression tests against known good

From 8cdf078a13d33b8c8f4c8f1550e7581a82d4e72e Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Thu, 6 Aug 2026 21:07:19 +0000
Subject: [PATCH 07/12] Remove Spark hash partitioning support

---
 cpp/include/cudf/partitioning.hpp             |  3 +-
 cpp/src/partitioning/partitioning.cu          | 25 ++++------------
 .../partitioning/hash_partition_test.cpp      | 20 +------------
 cpp/tests/row_operator/row_operator_tests.cu  | 29 ++++++++++++++++++-
 4 files changed, 36 insertions(+), 41 deletions(-)

diff --git a/cpp/include/cudf/partitioning.hpp b/cpp/include/cudf/partitioning.hpp
index e261ff7613b8..67cffce3e131 100644
--- a/cpp/include/cudf/partitioning.hpp
+++ b/cpp/include/cudf/partitioning.hpp
@@ -31,8 +31,7 @@ namespace CUDF_EXPORT cudf {
  */
 enum class hash_id : int32_t {
   HASH_IDENTITY = 0,  ///< Identity hash function that simply returns the key to be hashed
-  HASH_MURMUR3,       ///< Murmur3 hash function
-  SPARK_MURMUR3       ///< Apache Spark-compatible Murmur3 hash function
+  HASH_MURMUR3        ///< Murmur3 hash function
 };
 
 /**
diff --git a/cpp/src/partitioning/partitioning.cu b/cpp/src/partitioning/partitioning.cu
index 2ac3d673bb70..d9bf0d99864f 100644
--- a/cpp/src/partitioning/partitioning.cu
+++ b/cpp/src/partitioning/partitioning.cu
@@ -15,7 +15,6 @@
 #include <cudf/detail/utilities/integer_utils.hpp>
 #include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/hashing/detail/murmurhash3_x86_32.cuh>
-#include <cudf/hashing/detail/spark_murmurhash3.cuh>
 #include <cudf/partitioning.hpp>
 #include <cudf/table/table_device_view.cuh>
 #include <cudf/utilities/default_stream.hpp>
@@ -566,10 +565,7 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table_g
 }
 
 // NOTE hash_has_nulls must be true if table_to_hash has nulls
-template <template <typename> class hash_function,
-          bool hash_has_nulls,
-          template <template <typename> class, typename> class DeviceRowHasher =
-            detail::row::hash::device_row_hasher>
+template <template <typename> class hash_function, bool hash_has_nulls>
 std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
   table_view const& input,
   table_view const& table_to_hash,
@@ -581,8 +577,8 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
   auto const num_rows = table_to_hash.num_rows();
 
   auto const row_hasher = detail::row::hash::row_hasher(table_to_hash, stream);
-  auto const hasher     = row_hasher.device_hasher<hash_function, DeviceRowHasher>(
-    nullate::DYNAMIC{hash_has_nulls}, seed);
+  auto const hasher =
+    row_hasher.device_hasher<hash_function>(nullate::DYNAMIC{hash_has_nulls}, seed);
 
   // Check whether the per-block shared memory histograms fit in shared memory
   int dev;
@@ -875,9 +871,7 @@ struct IdentityHash {
   }
 };
 
-template <template <typename> class hash_function,
-          template <template <typename> class, typename> class DeviceRowHasher =
-            detail::row::hash::device_row_hasher>
+template <template <typename> class hash_function>
 std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition(
   table_view const& input,
   table_view const& table_to_hash,
@@ -897,10 +891,10 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition(
     }
   }
   if (has_nested_nulls(table_to_hash)) {
-    return hash_partition_table<hash_function, true, DeviceRowHasher>(
+    return hash_partition_table<hash_function, true>(
       input, table_to_hash, num_partitions, seed, stream, mr);
   } else {
-    return hash_partition_table<hash_function, false, DeviceRowHasher>(
+    return hash_partition_table<hash_function, false>(
       input, table_to_hash, num_partitions, seed, stream, mr);
   }
 }
@@ -945,13 +939,6 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition(
     case (hash_id::HASH_MURMUR3):
       return hash_partition<cudf::hashing::detail::MurmurHash3_x86_32>(
         input, keys, num_partitions, seed, stream, mr);
-    case (hash_id::SPARK_MURMUR3):
-      if (keys.num_rows() > 0 && keys.num_columns() > 0) {
-        cudf::hashing::detail::check_spark_murmurhash3_compatibility(keys);
-      }
-      return hash_partition<cudf::hashing::detail::Spark_MurmurHash3_x86_32,
-                            cudf::hashing::detail::spark_murmur_device_row_hasher>(
-        input, keys, num_partitions, seed, stream, mr);
     default: CUDF_FAIL("Unsupported hash function in hash_partition");
   }
 }
diff --git a/cpp/tests/partitioning/hash_partition_test.cpp b/cpp/tests/partitioning/hash_partition_test.cpp
index 4469325ab599..0a9d4f4992fc 100644
--- a/cpp/tests/partitioning/hash_partition_test.cpp
+++ b/cpp/tests/partitioning/hash_partition_test.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 #include <cudf_test/base_fixture.hpp>
@@ -366,24 +366,6 @@ TEST_F(HashPartition, CustomSeedValue)
   CUDF_TEST_EXPECT_TABLES_EQUAL(output1->view(), output2->view());
 }
 
-TEST_F(HashPartition, SparkMurmur3)
-{
-  fixed_width_column_wrapper<int32_t> const first({0, 1, -1, 42, 123456789});
-  fixed_width_column_wrapper<int32_t> const second({10, 20, 30, -40, -987654321});
-  auto const input = cudf::table_view({first, second});
-
-  auto [output, offsets] = cudf::hash_partition(
-    input, std::vector<cudf::size_type>{0, 1}, 8, cudf::hash_id::SPARK_MURMUR3, 42);
-
-  fixed_width_column_wrapper<int32_t> const expected_first({42, 1, 0, -1, 123456789});
-  fixed_width_column_wrapper<int32_t> const expected_second({-40, 20, 10, 30, -987654321});
-  auto const expected = cudf::table_view({expected_first, expected_second});
-  std::vector<cudf::size_type> const expected_offsets({0, 1, 1, 2, 3, 3, 3, 4, 5});
-
-  EXPECT_EQ(offsets, expected_offsets);
-  CUDF_TEST_EXPECT_TABLES_EQUAL(output->view(), expected);
-}
-
 template <typename T>
 class HashPartitionFixedWidth : public cudf::test::BaseFixture {};
 
diff --git a/cpp/tests/row_operator/row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu
index 08046c5e3128..7a98590af5cc 100644
--- a/cpp/tests/row_operator/row_operator_tests.cu
+++ b/cpp/tests/row_operator/row_operator_tests.cu
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -14,6 +14,7 @@
 #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/spark_murmurhash3.cuh>
 #include <cudf/hashing/detail/xxhash_64.cuh>
 #include <cudf/strings/strings_column_view.hpp>
 
@@ -365,6 +366,32 @@ TEST_F(RowOperatorTest, TestRowHasher64BitHash)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(results, expected);
 }
 
+TEST_F(RowOperatorTest, TestSparkMurmurRowHasher)
+{
+  auto const first  = cudf::test::fixed_width_column_wrapper<int32_t>{0, 1, -1, 42, 123456789};
+  auto const second = cudf::test::fixed_width_column_wrapper<int32_t>{10, 20, 30, -40, -987654321};
+  auto const input  = cudf::table_view{{first, second}};
+
+  auto const stream     = cudf::get_default_stream();
+  auto const row_hasher = cudf::detail::row::hash::row_hasher{input, stream};
+  auto const hasher =
+    row_hasher.device_hasher<cudf::hashing::detail::Spark_MurmurHash3_x86_32,
+                             cudf::hashing::detail::spark_murmur_device_row_hasher>(
+      cudf::nullate::DYNAMIC{false}, 42);
+
+  auto results = cudf::test::fixed_width_column_wrapper<int32_t>{0, 0, 0, 0, 0};
+  thrust::transform(rmm::exec_policy_nosync(stream),
+                    cuda::counting_iterator<cudf::size_type>{0},
+                    cuda::counting_iterator<cudf::size_type>{input.num_rows()},
+                    cudf::mutable_column_view{results}.begin<int32_t>(),
+                    hasher);
+
+  // Values produced by Apache Spark for the same input and seed.
+  auto const expected = cudf::test::fixed_width_column_wrapper<int32_t>{
+    -1721723333, 1151116018, 1549484878, -1287750896, -1980733329};
+  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}};

From 353bef9c964c773ba1c127ff94e3edf14d12028b Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Thu, 6 Aug 2026 23:00:36 +0000
Subject: [PATCH 08/12] Generalize Spark row hashing policy

---
 .../detail/row_operator/spark_hashing.cuh     | 129 ++++++++++++
 .../cudf/hashing/detail/spark_murmurhash3.cuh | 191 +++---------------
 cpp/src/hash/spark_murmurhash3_x86_32.cu      |  21 +-
 cpp/tests/row_operator/row_operator_tests.cu  |   8 +-
 4 files changed, 169 insertions(+), 180 deletions(-)
 create mode 100644 cpp/include/cudf/detail/row_operator/spark_hashing.cuh

diff --git a/cpp/include/cudf/detail/row_operator/spark_hashing.cuh b/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
new file mode 100644
index 000000000000..411e8c37131f
--- /dev/null
+++ b/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
@@ -0,0 +1,129 @@
+/*
+ * 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.
+ *
+ * @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, CUDF_ENABLE_IF(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, CUDF_ENABLE_IF(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,
+        [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
diff --git a/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
index 2caa0cfe9810..ee209ded2583 100644
--- a/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
+++ b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
@@ -5,29 +5,22 @@
 
 #pragma once
 
-#include <cudf/detail/row_operator/hashing.cuh>
 #include <cudf/fixed_point/fixed_point.hpp>
 #include <cudf/hashing/detail/hash_functions.cuh>
 #include <cudf/strings/string_view.hpp>
 #include <cudf/utilities/traits.hpp>
-#include <cudf/utilities/type_dispatcher.hpp>
 
-#include <cuda/functional>
 #include <cuda/std/algorithm>
 #include <cuda/std/cstddef>
 #include <cuda/std/iterator>
-#include <cuda/std/type_traits>
 #include <thrust/find.h>
-#include <thrust/iterator/counting_iterator.h>
 #include <thrust/reverse.h>
 
 namespace cudf::hashing::detail {
 
-using spark_hash_value_type = int32_t;
-
 template <typename Key, CUDF_ENABLE_IF(not cudf::is_nested<Key>())>
 struct Spark_MurmurHash3_x86_32 {
-  using result_type = spark_hash_value_type;
+  using result_type = int32_t;
 
   CUDF_HOST_DEVICE constexpr Spark_MurmurHash3_x86_32() = delete;
   CUDF_HOST_DEVICE constexpr Spark_MurmurHash3_x86_32(result_type seed)
@@ -124,57 +117,57 @@ struct Spark_MurmurHash3_x86_32 {
 };
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<bool>::operator()(
-  bool const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<bool>::operator()(bool const& key) const
+  -> result_type
 {
   return compute<uint32_t>(key);
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<int8_t>::operator()(
-  int8_t const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<int8_t>::operator()(int8_t const& key) const
+  -> result_type
 {
   return compute<uint32_t>(key);
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<uint8_t>::operator()(
-  uint8_t const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<uint8_t>::operator()(uint8_t const& key) const
+  -> result_type
 {
   return compute<uint32_t>(key);
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<int16_t>::operator()(
-  int16_t const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<int16_t>::operator()(int16_t const& key) const
+  -> result_type
 {
   return compute<uint32_t>(key);
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<uint16_t>::operator()(
-  uint16_t const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<uint16_t>::operator()(uint16_t const& key) const
+  -> result_type
 {
   return compute<uint32_t>(key);
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<float>::operator()(
-  float const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<float>::operator()(float const& key) const
+  -> result_type
 {
   return compute<float>(normalize_nans(key));
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<double>::operator()(
-  double const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<double>::operator()(double const& key) const
+  -> result_type
 {
   return compute<double>(normalize_nans(key));
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<cudf::string_view>::operator()(
-  cudf::string_view const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<cudf::string_view>::operator()(
+  cudf::string_view const& key) const -> result_type
 {
   auto const data = reinterpret_cast<cuda::std::byte const*>(key.data());
   auto const len  = key.size_bytes();
@@ -182,22 +175,22 @@ spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<cudf::string_vi
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<numeric::decimal32>::operator()(
-  numeric::decimal32 const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<numeric::decimal32>::operator()(
+  numeric::decimal32 const& key) const -> result_type
 {
   return compute<uint64_t>(key.value());
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<numeric::decimal64>::operator()(
-  numeric::decimal64 const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<numeric::decimal64>::operator()(
+  numeric::decimal64 const& key) const -> result_type
 {
   return compute<uint64_t>(key.value());
 }
 
 template <>
-spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<numeric::decimal128>::operator()(
-  numeric::decimal128 const& key) const
+__device__ inline auto Spark_MurmurHash3_x86_32<numeric::decimal128>::operator()(
+  numeric::decimal128 const& key) const -> result_type
 {
   // Generates the Spark MurmurHash3 hash value, mimicking the conversion:
   // java.math.BigDecimal.valueOf(unscaled_value, _scale).unscaledValue().toByteArray()
@@ -241,142 +234,4 @@ spark_hash_value_type __device__ inline Spark_MurmurHash3_x86_32<numeric::decima
   return compute_bytes(big_endian_data, length);
 }
 
-/**
- * @brief Computes the hash value of a row in the given table.
- *
- * This functor uses Spark conventions for Murmur hashing, which differs from
- * the Murmur implementation used in the rest of libcudf. These differences
- * include:
- * - Serially using the output hash as an input seed for the next item
- * - Ignorance of null values
- *
- * The serial use of hashes as seeds means that data of different nested types
- * can exhibit hash collisions. For example, a row of an integer column
- * containing a 1 will have the same hash as a lists column of integers
- * containing a list of [1] and a struct column of a single integer column
- * containing a struct of {1}.
- *
- * As a consequence of ignoring null values, inputs like [1], [1, null], and
- * [null, 1] have the same hash (an expected hash collision). This kind of
- * collision can also occur across a table of nullable columns and with nulls
- * in structs ({1, null} and {null, 1} have the same hash). The seed value (the
- * previous element's hash value) is returned as the hash if an element is
- * null.
- *
- * For additional differences such as special tail processing and decimal type
- * handling, refer to the Spark_MurmurHash3_x86_32 functor.
- *
- * @tparam hash_function Hash functor to use for hashing elements. Must be Spark_MurmurHash3_x86_32.
- * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
- */
-template <template <typename> class hash_function, typename Nullate>
-class spark_murmur_device_row_hasher {
-  friend class cudf::detail::row::hash::row_hasher;  ///< Allow row_hasher to access private
-                                                     ///< members.
-
- public:
-  /**
-   * @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
-  {
-    return cudf::detail::accumulate(
-      _table.begin(),
-      _table.end(),
-      _seed,
-      cuda::proclaim_return_type<spark_hash_value_type>(
-        [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
-          return cudf::type_dispatcher(
-            column.type(), element_hasher_adapter<hash_function>{nulls, hash}, column, row_index);
-        }));
-  }
-
- private:
-  /**
-   * @brief Computes the hash value of an element in the given column.
-   *
-   * When the column is non-nested, this is a simple wrapper around the element_hasher.
-   * When the column is nested, this uses a seed value to serially compute each
-   * nested element, with the output hash becoming the seed for the next value.
-   * This requires constructing a new hash functor for each nested element,
-   * using the new seed from the previous element's hash. The hash of a null
-   * element is the input seed (the previous element's hash).
-   */
-  template <template <typename> class hash_fn>
-  class element_hasher_adapter {
-   public:
-    using hash_functor = cudf::detail::row::hash::element_hasher<hash_fn, Nullate>;
-    using result_type  = typename hash_functor::result_type;
-
-    __device__ element_hasher_adapter(Nullate check_nulls, result_type seed) noexcept
-      : _check_nulls(check_nulls), _seed(seed)
-    {
-    }
-
-    template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
-    __device__ spark_hash_value_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, CUDF_ENABLE_IF(cudf::is_nested<T>())>
-    __device__ spark_hash_value_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 assumed to be decomposed and contain only one child
-          curr_col = cudf::structs_column_device_view(curr_col).get_sliced_child(0);
-        } else if (curr_col.type().id() == type_id::LIST) {
-          curr_col = cudf::lists_column_device_view(curr_col).get_sliced_child();
-        }
-      }
-
-      return cudf::detail::accumulate(
-        thrust::counting_iterator(0),
-        thrust::counting_iterator(curr_col.size()),
-        _seed,
-        [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);
-        });
-    }
-
-    Nullate const _check_nulls;  ///< Whether to check for nulls
-    result_type const _seed;     ///< The seed to use for hashing, also returned for null elements
-  };
-
-  using result_type = typename element_hasher_adapter<hash_function>::result_type;
-
-  CUDF_HOST_DEVICE spark_murmur_device_row_hasher(Nullate check_nulls,
-                                                  table_device_view t,
-                                                  result_type seed = DEFAULT_HASH_SEED) noexcept
-    : _check_nulls{check_nulls}, _table{t}, _seed(seed)
-  {
-    // Error out if passed an unsupported hash_function
-    static_assert(
-      cuda::std::is_same_v<Spark_MurmurHash3_x86_32<int>, hash_function<int>>,
-      "spark_murmur_device_row_hasher only supports the Spark_MurmurHash3_x86_32 hash function");
-  }
-
-  Nullate const _check_nulls;
-  table_device_view const _table;
-  result_type const _seed;
-};
-
-/**
- * @brief Throws if `input` contains a nested shape unsupported by Spark MurmurHash3.
- *
- * @param input Table to validate
- */
-void check_spark_murmurhash3_compatibility(table_view const& input);
-
 }  // namespace cudf::hashing::detail
diff --git a/cpp/src/hash/spark_murmurhash3_x86_32.cu b/cpp/src/hash/spark_murmurhash3_x86_32.cu
index 8625b1e49112..9428060dd963 100644
--- a/cpp/src/hash/spark_murmurhash3_x86_32.cu
+++ b/cpp/src/hash/spark_murmurhash3_x86_32.cu
@@ -4,6 +4,7 @@
  */
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/detail/row_operator/spark_hashing.cuh>
 #include <cudf/hashing.hpp>
 #include <cudf/hashing/detail/spark_murmurhash3.cuh>
 #include <cudf/lists/lists_column_view.hpp>
@@ -19,6 +20,8 @@ namespace cudf {
 namespace hashing {
 namespace detail {
 
+namespace {
+
 void check_spark_murmurhash3_compatibility(table_view const& input)
 {
   using column_checker_fn_t = std::function<void(column_view const&)>;
@@ -41,16 +44,17 @@ void check_spark_murmurhash3_compatibility(table_view const& input)
   }
 }
 
+}  // namespace
+
 std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
                                                  uint32_t seed,
                                                  rmm::cuda_stream_view stream,
                                                  rmm::device_async_resource_ref mr)
 {
-  auto output = make_numeric_column(data_type(type_to_id<spark_hash_value_type>()),
-                                    input.num_rows(),
-                                    mask_state::UNALLOCATED,
-                                    stream,
-                                    mr);
+  using result_type = Spark_MurmurHash3_x86_32<int32_t>::result_type;
+
+  auto output = make_numeric_column(
+    data_type(type_to_id<result_type>()), input.num_rows(), mask_state::UNALLOCATED, stream, mr);
 
   // Return early if there's nothing to hash
   if (input.num_rows() == 0) { return output; }
@@ -60,12 +64,13 @@ std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
 
   bool const nullable     = has_nested_nulls(input);
   auto const row_hasher   = cudf::detail::row::hash::row_hasher(input, stream);
-  auto const output_begin = output->mutable_view().begin<spark_hash_value_type>();
+  auto const output_begin = output->mutable_view().begin<result_type>();
 
   // Compute the hash value for each row
   auto const hasher =
-    row_hasher.device_hasher<Spark_MurmurHash3_x86_32, spark_murmur_device_row_hasher>(nullable,
-                                                                                       seed);
+    row_hasher
+      .device_hasher<Spark_MurmurHash3_x86_32, cudf::detail::row::hash::spark_device_row_hasher>(
+        nullable, seed);
   CUDF_CUDA_TRY(cub::DeviceFor::Bulk(
     input.num_rows(),
     [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); },
diff --git a/cpp/tests/row_operator/row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu
index 7a98590af5cc..807eb16e8387 100644
--- a/cpp/tests/row_operator/row_operator_tests.cu
+++ b/cpp/tests/row_operator/row_operator_tests.cu
@@ -14,6 +14,7 @@
 #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/detail/row_operator/spark_hashing.cuh>
 #include <cudf/hashing/detail/spark_murmurhash3.cuh>
 #include <cudf/hashing/detail/xxhash_64.cuh>
 #include <cudf/strings/strings_column_view.hpp>
@@ -374,10 +375,9 @@ TEST_F(RowOperatorTest, TestSparkMurmurRowHasher)
 
   auto const stream     = cudf::get_default_stream();
   auto const row_hasher = cudf::detail::row::hash::row_hasher{input, stream};
-  auto const hasher =
-    row_hasher.device_hasher<cudf::hashing::detail::Spark_MurmurHash3_x86_32,
-                             cudf::hashing::detail::spark_murmur_device_row_hasher>(
-      cudf::nullate::DYNAMIC{false}, 42);
+  auto const hasher     = row_hasher.device_hasher<cudf::hashing::detail::Spark_MurmurHash3_x86_32,
+                                                   cudf::detail::row::hash::spark_device_row_hasher>(
+    cudf::nullate::DYNAMIC{false}, 42);
 
   auto results = cudf::test::fixed_width_column_wrapper<int32_t>{0, 0, 0, 0, 0};
   thrust::transform(rmm::exec_policy_nosync(stream),

From 6b72548636684a61c00ea68324f6ed848301c41a Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Fri, 7 Aug 2026 18:33:29 +0000
Subject: [PATCH 09/12] Fix Spark MurmurHash3 signed-zero handling

Signed-off-by: Yunsong Wang <yunsongw@nvidia.com>
---
 .../detail/row_operator/spark_hashing.cuh     | 14 +++---
 cpp/include/cudf/hashing.hpp                  |  2 +-
 cpp/include/cudf/hashing/detail/hashing.hpp   |  5 +++
 .../cudf/hashing/detail/spark_murmurhash3.cuh |  9 ++--
 cpp/src/hash/spark_murmurhash3_x86_32.cu      |  3 +-
 .../hashing/spark_murmurhash3_x86_32_test.cpp | 44 ++++++++++++++++---
 6 files changed, 61 insertions(+), 16 deletions(-)

diff --git a/cpp/include/cudf/detail/row_operator/spark_hashing.cuh b/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
index 411e8c37131f..a5d294617630 100644
--- a/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
+++ b/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
@@ -29,6 +29,9 @@ namespace detail::row::hash {
  * 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
  */
@@ -101,11 +104,12 @@ class spark_device_row_hasher {
         thrust::counting_iterator(0),
         thrust::counting_iterator(curr_col.size()),
         _seed,
-        [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);
-        });
+        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:
diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp
index bf447a927061..5792f17ba9a4 100644
--- a/cpp/include/cudf/hashing.hpp
+++ b/cpp/include/cudf/hashing.hpp
@@ -61,7 +61,7 @@ std::unique_ptr<column> murmurhash3_x86_32(
  * 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, fixed-point values, lists, and structs.
+ * NaNs and signed zeros, fixed-point values, lists, and structs.
  *
  * This function does not support LIST columns whose child is a STRUCT.
  *
diff --git a/cpp/include/cudf/hashing/detail/hashing.hpp b/cpp/include/cudf/hashing/detail/hashing.hpp
index 567bb864628a..05e2bfa31418 100644
--- a/cpp/include/cudf/hashing/detail/hashing.hpp
+++ b/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,
+                                                 rmm::cuda_stream_view,
+                                                 rmm::device_async_resource_ref mr);
+
 std::unique_ptr<table> murmurhash3_x64_128(table_view const& input,
                                            uint64_t seed,
                                            rmm::cuda_stream_view,
diff --git a/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
index ee209ded2583..cc84fc55daa3 100644
--- a/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
+++ b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
@@ -13,6 +13,7 @@
 #include <cuda/std/algorithm>
 #include <cuda/std/cstddef>
 #include <cuda/std/iterator>
+#include <thrust/execution_policy.h>
 #include <thrust/find.h>
 #include <thrust/reverse.h>
 
@@ -108,7 +109,7 @@ struct Spark_MurmurHash3_x86_32 {
   }
 
  private:
-  uint32_t m_seed{cudf::DEFAULT_HASH_SEED};
+  uint32_t m_seed;
   static constexpr uint32_t c1     = 0xcc9e2d51;
   static constexpr uint32_t c2     = 0x1b873593;
   static constexpr uint32_t c3     = 0xe6546b64;
@@ -155,14 +156,14 @@ template <>
 __device__ inline auto Spark_MurmurHash3_x86_32<float>::operator()(float const& key) const
   -> result_type
 {
-  return compute<float>(normalize_nans(key));
+  return compute<float>(normalize_nans_and_zeros(key));
 }
 
 template <>
 __device__ inline auto Spark_MurmurHash3_x86_32<double>::operator()(double const& key) const
   -> result_type
 {
-  return compute<double>(normalize_nans(key));
+  return compute<double>(normalize_nans_and_zeros(key));
 }
 
 template <>
@@ -194,7 +195,7 @@ __device__ inline auto Spark_MurmurHash3_x86_32<numeric::decimal128>::operator()
 {
   // Generates the Spark MurmurHash3 hash value, mimicking the conversion:
   // java.math.BigDecimal.valueOf(unscaled_value, _scale).unscaledValue().toByteArray()
-  // https://github.com/apache/spark/blob/master/sql/catalyst/src/main/scala/org/apache/spark/sql/catalyst/expressions/hash.scala#L381
+  // https://github.com/apache/spark/blob/ce5ddad990373636e94071e7cef2f31021add07b/sql/catalyst/src/main/scala/org/apache/spark/sql/catalyst/expressions/hash.scala#L391-L396
   __int128_t const val               = key.value();
   constexpr cudf::size_type key_size = sizeof(__int128_t);
   cuda::std::byte const* data        = reinterpret_cast<cuda::std::byte const*>(&val);
diff --git a/cpp/src/hash/spark_murmurhash3_x86_32.cu b/cpp/src/hash/spark_murmurhash3_x86_32.cu
index 9428060dd963..4c1bd2baf88e 100644
--- a/cpp/src/hash/spark_murmurhash3_x86_32.cu
+++ b/cpp/src/hash/spark_murmurhash3_x86_32.cu
@@ -6,6 +6,7 @@
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/row_operator/spark_hashing.cuh>
 #include <cudf/hashing.hpp>
+#include <cudf/hashing/detail/hashing.hpp>
 #include <cudf/hashing/detail/spark_murmurhash3.cuh>
 #include <cudf/lists/lists_column_view.hpp>
 #include <cudf/utilities/error.hpp>
@@ -62,7 +63,7 @@ std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
   // Lists of structs are not supported
   check_spark_murmurhash3_compatibility(input);
 
-  bool const nullable     = has_nested_nulls(input);
+  auto const nullable     = nullate::DYNAMIC{has_nested_nulls(input)};
   auto const row_hasher   = cudf::detail::row::hash::row_hasher(input, stream);
   auto const output_begin = output->mutable_view().begin<result_type>();
 
diff --git a/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
index 056fa28f374d..2083898b5707 100644
--- a/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
+++ b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
@@ -9,6 +9,7 @@
 #include <cudf_test/iterator_utilities.hpp>
 #include <cudf_test/type_lists.hpp>
 
+#include <cudf/copying.hpp>
 #include <cudf/detail/iterator.cuh>
 #include <cudf/fixed_point/fixed_point.hpp>
 #include <cudf/hashing.hpp>
@@ -75,15 +76,48 @@ TYPED_TEST(SparkMurmurHashTestFloatTyped, TestExtremes)
   auto const table_col_neg_zero = cudf::table_view({col_neg_zero});
   auto const table_col_neg_nan  = cudf::table_view({col_neg_nan});
 
-  // Spark hash is sensitive to 0 and -0
-  auto const spark_col         = cudf::hashing::spark_murmurhash3_x86_32(table_col, 0);
-  auto const spark_col_neg_nan = cudf::hashing::spark_murmurhash3_x86_32(table_col_neg_nan);
+  // Spark 3.2+ normalizes signed zero before invoking the JNI hasher.
+  // https://github.com/NVIDIA/spark-rapids/blob/aef688e232b9e2c710696d380b2539230897450a/sql-plugin/src/main/scala/com/nvidia/spark/rapids/shims/HashUtils.scala#L21-L44
+  // https://github.com/NVIDIA/spark-rapids-jni/blob/09bae9c7dccf050b5db53a70282c3001cba5a015/src/main/cpp/src/hash/murmur_hash.cuh#L163-L172
+  auto const spark_col          = cudf::hashing::spark_murmurhash3_x86_32(table_col, 0);
+  auto const spark_col_neg_zero = cudf::hashing::spark_murmurhash3_x86_32(table_col_neg_zero, 0);
+  auto const spark_col_neg_nan  = cudf::hashing::spark_murmurhash3_x86_32(table_col_neg_nan, 0);
 
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*spark_col, *spark_col_neg_zero);
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(*spark_col, *spark_col_neg_nan);
 }
 
 class SparkMurmurHashTest : public cudf::test::BaseFixture {};
 
+TEST_F(SparkMurmurHashTest, EmptyInput)
+{
+  cudf::test::fixed_width_column_wrapper<int32_t> const empty{};
+  auto const expected = cudf::test::fixed_width_column_wrapper<int32_t>{};
+
+  auto const empty_table_output = cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view{}, 42);
+  auto const empty_column_output =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view{{empty}}, 42);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, empty_table_output->view());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, empty_column_output->view());
+}
+
+TEST_F(SparkMurmurHashTest, SlicedInput)
+{
+  cudf::test::fixed_width_column_wrapper<int32_t> const first({999, 0, 1, -1, 42, 123456789, -999});
+  cudf::test::fixed_width_column_wrapper<int32_t> const second(
+    {999, 10, 20, 30, -40, -987654321, -999});
+
+  auto const first_slice  = cudf::slice(first, {1, 6}).front();
+  auto const second_slice = cudf::slice(second, {1, 6}).front();
+  auto const output =
+    cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view{{first_slice, second_slice}}, 42);
+
+  cudf::test::fixed_width_column_wrapper<int32_t> const expected(
+    {-1721723333, 1151116018, 1549484878, -1287750896, -1980733329});
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expected, output->view());
+}
+
 TEST_F(SparkMurmurHashTest, IterativeSeeding)
 {
   cudf::test::fixed_width_column_wrapper<int32_t> const first({0, 1, -1, 42, 123456789});
@@ -490,7 +524,7 @@ TEST_F(SparkMurmurHashTest, StructOfListValues)
   val df = spark.createDataFrame(
     spark.sparkContext.parallelize(data), schema)
 
-  val df2 = df.selectExpr("lists", "hash(lists) as hash")
+  val df2 = df.selectExpr("structs", "hash(structs) as hash")
   df2.printSchema()
   df2.show(false)
   */
@@ -558,7 +592,7 @@ TEST_F(SparkMurmurHashTest, ListOfStructValues)
     cudf::test::structs_column_wrapper{{col1, col2}, {1, 0, 1, 1, 1, 1, 1, 1, 0, 1, 1}};
   auto offsets =
     cudf::test::fixed_width_column_wrapper<cudf::size_type>{0, 1, 2, 3, 4, 5, 7, 9, 11};
-  auto list_nullmask = std::vector<bool>(1, 8);
+  auto list_nullmask = std::vector<bool>(8, true);
   auto [null_mask, null_count] =
     cudf::test::detail::make_null_mask(list_nullmask.begin(), list_nullmask.end());
   auto list_column = cudf::make_lists_column(

From 319948beab34908cb9bf7cf42f8dc0c2935333a4 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Fri, 7 Aug 2026 20:05:39 +0000
Subject: [PATCH 10/12] Fix Spark signed-zero test vectors

Signed-off-by: Yunsong Wang <yunsongw@nvidia.com>
---
 cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
index 2083898b5707..601c79d4b1c7 100644
--- a/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
+++ b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
@@ -260,7 +260,7 @@ TEST_F(SparkMurmurHashTest, MultiValueWithSeeds)
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_strings_expected(
     {142593372, 1217302703, -715697185, -2061143941, -111635966});
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_doubles_expected(
-    {-1670924195, -853646085, -1281358385, 1897734433, -508695674});
+    {-1670924195, -1670924195, -1281358385, 1897734433, -508695674});
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_timestamps_expected(
     {-1670924195, 1114849490, 904948192, -1832979433, 1752430209});
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_decimal64_expected(
@@ -268,7 +268,7 @@ TEST_F(SparkMurmurHashTest, MultiValueWithSeeds)
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_longs_expected(
     {-1670924195, 1114849490, 904948192, -853646085, -1604625029});
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_floats_expected(
-    {933211791, 723455942, -349261430, -1225560532, -338752985});
+    {933211791, 933211791, -349261430, -1225560532, -338752985});
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_dates_expected(
     {933211791, 751823303, -1080202046, -1906567553, -1503850410});
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_decimal32_expected(
@@ -284,7 +284,7 @@ TEST_F(SparkMurmurHashTest, MultiValueWithSeeds)
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_decimal128_expected(
     {-783713497, -295670906, 1398487324, -52622807, -1359749815});
   cudf::test::fixed_width_column_wrapper<int32_t> const hash_combined_expected(
-    {401603227, 588162166, 552160517, 1132537411, -326043017});
+    {401603227, 1728350528, 552160517, 1132537411, -326043017});
 
   using double_limits = std::numeric_limits<double>;
   using long_limits   = std::numeric_limits<int64_t>;

From 9e882f3eb3bb3216966d57660f311412fe279826 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Fri, 21 Aug 2026 21:40:29 +0000
Subject: [PATCH 11/12] Address review: fix signed overflow, use requires and
 cuda::stream_ref

---
 .../detail/row_operator/spark_hashing.cuh     |  6 +++--
 cpp/include/cudf/hashing.hpp                  |  4 +++-
 cpp/include/cudf/hashing/detail/hashing.hpp   |  2 +-
 .../cudf/hashing/detail/spark_murmurhash3.cuh | 21 +++++++++--------
 cpp/src/hash/spark_murmurhash3_x86_32.cu      |  7 +++---
 .../hashing/spark_murmurhash3_x86_32_test.cpp | 23 +++++++------------
 6 files changed, 31 insertions(+), 32 deletions(-)

diff --git a/cpp/include/cudf/detail/row_operator/spark_hashing.cuh b/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
index a5d294617630..54c74c15a6fb 100644
--- a/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
+++ b/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
@@ -77,7 +77,8 @@ class spark_device_row_hasher {
     {
     }
 
-    template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
+    template <typename T>
+      requires(not cudf::is_nested<T>())
     __device__ result_type operator()(column_device_view const& col,
                                       size_type row_index) const noexcept
     {
@@ -85,7 +86,8 @@ class spark_device_row_hasher {
       return hasher.template operator()<T>(col, row_index);
     }
 
-    template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
+    template <typename T>
+      requires(cudf::is_nested<T>())
     __device__ result_type operator()(column_device_view const& col,
                                       size_type row_index) const noexcept
     {
diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp
index 5792f17ba9a4..79faee82d447 100644
--- a/cpp/include/cudf/hashing.hpp
+++ b/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.
@@ -75,7 +77,7 @@ std::unique_ptr<column> murmurhash3_x86_32(
 std::unique_ptr<column> spark_murmurhash3_x86_32(
   table_view const& input,
   uint32_t seed                     = DEFAULT_HASH_SEED,
-  rmm::cuda_stream_view stream      = cudf::get_default_stream(),
+  cuda::stream_ref stream           = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
 /**
diff --git a/cpp/include/cudf/hashing/detail/hashing.hpp b/cpp/include/cudf/hashing/detail/hashing.hpp
index 05e2bfa31418..18ade03e6d7b 100644
--- a/cpp/include/cudf/hashing/detail/hashing.hpp
+++ b/cpp/include/cudf/hashing/detail/hashing.hpp
@@ -24,7 +24,7 @@ std::unique_ptr<column> murmurhash3_x86_32(table_view const& input,
 
 std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
                                                  uint32_t seed,
-                                                 rmm::cuda_stream_view,
+                                                 cuda::stream_ref,
                                                  rmm::device_async_resource_ref mr);
 
 std::unique_ptr<table> murmurhash3_x64_128(table_view const& input,
diff --git a/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
index cc84fc55daa3..63fffa882862 100644
--- a/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
+++ b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
@@ -19,7 +19,8 @@
 
 namespace cudf::hashing::detail {
 
-template <typename Key, CUDF_ENABLE_IF(not cudf::is_nested<Key>())>
+template <typename Key>
+  requires(not cudf::is_nested<Key>())
 struct Spark_MurmurHash3_x86_32 {
   using result_type = int32_t;
 
@@ -59,10 +60,10 @@ struct Spark_MurmurHash3_x86_32 {
     return compute_bytes(reinterpret_cast<cuda::std::byte const*>(&key), sizeof(T));
   }
 
-  result_type __device__ inline compute_remaining_bytes(cuda::std::byte const* data,
-                                                        cudf::size_type len,
-                                                        cudf::size_type tail_offset,
-                                                        result_type h) const
+  uint32_t __device__ inline compute_remaining_bytes(cuda::std::byte const* data,
+                                                     cudf::size_type len,
+                                                     cudf::size_type tail_offset,
+                                                     uint32_t h) const
   {
     // Process remaining bytes that do not fill a four-byte chunk using Spark's approach
     // (does not conform to normal MurmurHash3).
@@ -76,7 +77,7 @@ struct Spark_MurmurHash3_x86_32 {
       k1 = rotate_bits_left(k1, rot_c1);
       k1 *= c2;
       h ^= k1;
-      h = rotate_bits_left(static_cast<uint32_t>(h), rot_c2);
+      h = rotate_bits_left(h, rot_c2);
       h = h * 5 + c3;
     }
     return h;
@@ -87,7 +88,7 @@ struct Spark_MurmurHash3_x86_32 {
     constexpr cudf::size_type BLOCK_SIZE = 4;
     cudf::size_type const nblocks        = len / BLOCK_SIZE;
     cudf::size_type const tail_offset    = nblocks * BLOCK_SIZE;
-    result_type h                        = m_seed;
+    uint32_t h                           = m_seed;
 
     // Process all four-byte chunks.
     for (cudf::size_type i = 0; i < nblocks; i++) {
@@ -96,16 +97,16 @@ struct Spark_MurmurHash3_x86_32 {
       k1 = rotate_bits_left(k1, rot_c1);
       k1 *= c2;
       h ^= k1;
-      h = rotate_bits_left(static_cast<uint32_t>(h), rot_c2);
+      h = rotate_bits_left(h, rot_c2);
       h = h * 5 + c3;
     }
 
     h = compute_remaining_bytes(data, len, tail_offset, h);
 
     // Finalize hash.
-    h ^= len;
+    h ^= static_cast<uint32_t>(len);
     h = fmix32(h);
-    return h;
+    return static_cast<result_type>(h);
   }
 
  private:
diff --git a/cpp/src/hash/spark_murmurhash3_x86_32.cu b/cpp/src/hash/spark_murmurhash3_x86_32.cu
index 4c1bd2baf88e..2eb4ed182cec 100644
--- a/cpp/src/hash/spark_murmurhash3_x86_32.cu
+++ b/cpp/src/hash/spark_murmurhash3_x86_32.cu
@@ -14,6 +14,7 @@
 #include <rmm/cuda_stream_view.hpp>
 
 #include <cub/device/device_for.cuh>
+#include <cuda/stream_ref>
 
 #include <functional>
 
@@ -49,7 +50,7 @@ void check_spark_murmurhash3_compatibility(table_view const& input)
 
 std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
                                                  uint32_t seed,
-                                                 rmm::cuda_stream_view stream,
+                                                 cuda::stream_ref stream,
                                                  rmm::device_async_resource_ref mr)
 {
   using result_type = Spark_MurmurHash3_x86_32<int32_t>::result_type;
@@ -75,7 +76,7 @@ std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
   CUDF_CUDA_TRY(cub::DeviceFor::Bulk(
     input.num_rows(),
     [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); },
-    stream.value()));
+    stream.get()));
 
   return output;
 }
@@ -84,7 +85,7 @@ std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
 
 std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
                                                  uint32_t seed,
-                                                 rmm::cuda_stream_view stream,
+                                                 cuda::stream_ref stream,
                                                  rmm::device_async_resource_ref mr)
 {
   CUDF_FUNC_RANGE();
diff --git a/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
index 601c79d4b1c7..e465db85851c 100644
--- a/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
+++ b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
@@ -191,11 +191,11 @@ TEST_F(SparkMurmurHashTest, MultiValueNulls)
 TEST_F(SparkMurmurHashTest, MultiValueWithSeeds)
 {
   // The hash values were determined by running the following Scala code in Apache Spark.
-  // Note that Spark >= 3.2 normalizes the float/double value of -0. to +0. and both values hash
-  // to the same result. This is normalized in the calling code (Spark RAPIDS plugin) for Spark
-  // >= 3.2. However, the reference values for -0. below must be obtained with Spark < 3.2 and
-  // libcudf will continue to implement the Spark < 3.2 behavior until Spark >= 3.2 is required and
-  // the workaround in the calling code is removed. This also affects the combined hash values.
+  // Note that Spark >= 3.2 normalizes the float/double value of -0. to +0. so that both values
+  // hash to the same result, and libcudf normalizes signed zero to match. The reference values
+  // for -0. below are therefore the Spark >= 3.2 values, which are the ones for 0., and the
+  // combined hash values follow from them. The Spark RAPIDS plugin applies the same
+  // normalization in its calling code, so normalizing here is idempotent for that caller.
 
   /*
   import org.apache.spark.sql.functions._
@@ -598,16 +598,9 @@ TEST_F(SparkMurmurHashTest, ListOfStructValues)
   auto list_column = cudf::make_lists_column(
     8, offsets.release(), struct_column.release(), null_count, std::move(null_mask));
 
-  // TODO: Lists of structs are not yet supported. Once support is added,
-  // remove this EXPECT_THROW and uncomment the rest of this test.
+  // Lists of structs are rejected, matching the Spark RAPIDS JNI implementation, which raises the
+  // same error from check_hash_compatibility.
+  // https://github.com/NVIDIA/spark-rapids-jni/blob/09bae9c7dccf050b5db53a70282c3001cba5a015/src/main/cpp/src/hash/murmur_hash.cu#L167-L182
   EXPECT_THROW(cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({*list_column}), 42),
                cudf::logic_error);
-
-  /*
-  auto expect = cudf::test::fixed_width_column_wrapper<int32_t>{
-    59727262, 42, 42, -559580957, -559580957, -912918097, 1092624418, 170038658};
-
-  auto output = cudf::hashing::spark_murmurhash3_x86_32(cudf::table_view({*list_column}), 42);
-  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
-  */
 }

From d2dac9a84fe3d75b927bd1f86adca927181ec798 Mon Sep 17 00:00:00 2001
From: Yunsong Wang <yunsongw@nvidia.com>
Date: Fri, 21 Aug 2026 22:12:25 +0000
Subject: [PATCH 12/12] Add stream test for spark_murmurhash3_x86_32

---
 cpp/tests/streams/hash_test.cpp | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/cpp/tests/streams/hash_test.cpp b/cpp/tests/streams/hash_test.cpp
index fd49e85d38ea..f0f6728097eb 100644
--- a/cpp/tests/streams/hash_test.cpp
+++ b/cpp/tests/streams/hash_test.cpp
@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -41,6 +41,9 @@ TEST_F(HashTest, MultiValue)
 
   auto const output1 = cudf::hashing::murmurhash3_x86_32(
     input1, cudf::DEFAULT_HASH_SEED, cudf::test::get_default_stream());
+
+  auto const output2 = cudf::hashing::spark_murmurhash3_x86_32(
+    input1, cudf::DEFAULT_HASH_SEED, cudf::test::get_default_stream());
 }
 
 CUDF_TEST_PROGRAM_MAIN()