diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 26b7ee68aad7..a402333db81e 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -706,6 +706,7 @@ add_library(
   src/hash/md5_hash.cu
   src/hash/murmurhash3_x86_32.cu
   src/hash/murmurhash3_x64_128.cu
+  src/hash/spark_murmurhash3_x86_32.cu
   src/hash/sha1_hash.cu
   src/hash/sha224_hash.cu
   src/hash/sha256_hash.cu
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..54c74c15a6fb
--- /dev/null
+++ b/cpp/include/cudf/detail/row_operator/spark_hashing.cuh
@@ -0,0 +1,135 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#include <cudf/detail/row_operator/hashing.cuh>
+#include <cudf/detail/utilities/accumulate.cuh>
+#include <cudf/lists/lists_column_device_view.cuh>
+#include <cudf/structs/structs_column_device_view.cuh>
+#include <cudf/utilities/type_dispatcher.hpp>
+
+#include <cuda/functional>
+#include <thrust/iterator/counting_iterator.h>
+
+namespace CUDF_EXPORT cudf {
+namespace detail::row::hash {
+
+/**
+ * @brief Computes row hashes using Spark's iterative seeding convention.
+ *
+ * Spark uses the hash of each value as the seed for the next value and ignores
+ * null values. Consequently, values of different nested shapes can collide.
+ * For example, the integer `1`, the list `[1]`, and a struct containing only
+ * `1` have the same hash. Likewise, `[1]`, `[1, null]`, and `[null, 1]` have
+ * the same hash. A null element returns its input seed unchanged.
+ *
+ * The element hash function is responsible for Spark-specific type encodings
+ * and hash algorithm behavior.
+ *
+ * LIST columns whose child is a STRUCT are unsupported and must be rejected
+ * before invoking this hasher.
+ *
+ * @tparam hash_function Seeded element hash functor with a `result_type` member
+ * @tparam Nullate A cudf::nullate type describing whether to check for nulls
+ */
+template <template <typename> class hash_function, typename Nullate>
+class spark_device_row_hasher {
+  friend class row_hasher;
+
+ public:
+  using result_type = typename hash_function<int32_t>::result_type;
+
+  /**
+   * @brief Returns the hash value of a row in the table.
+   *
+   * @param row_index The row index to hash
+   * @return The hash value of the row
+   */
+  __device__ result_type operator()(size_type row_index) const noexcept
+  {
+    return cudf::detail::accumulate(
+      _table.begin(),
+      _table.end(),
+      _seed,
+      cuda::proclaim_return_type<result_type>(
+        [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
+          return cudf::type_dispatcher(
+            column.type(), element_hasher_adapter{nulls, hash}, column, row_index);
+        }));
+  }
+
+ private:
+  /**
+   * @brief Computes the hash value of an element in a column.
+   *
+   * Nested values are flattened and hashed serially, with each output becoming
+   * the seed for the next value. A null element returns the input seed.
+   */
+  class element_hasher_adapter {
+    using hash_functor = element_hasher<hash_function, Nullate>;
+
+   public:
+    __device__ element_hasher_adapter(Nullate check_nulls, result_type seed) noexcept
+      : _check_nulls(check_nulls), _seed(seed)
+    {
+    }
+
+    template <typename T>
+      requires(not cudf::is_nested<T>())
+    __device__ result_type operator()(column_device_view const& col,
+                                      size_type row_index) const noexcept
+    {
+      auto const hasher = hash_functor{_check_nulls, _seed, _seed};
+      return hasher.template operator()<T>(col, row_index);
+    }
+
+    template <typename T>
+      requires(cudf::is_nested<T>())
+    __device__ result_type operator()(column_device_view const& col,
+                                      size_type row_index) const noexcept
+    {
+      column_device_view curr_col = col.slice(row_index, 1);
+      while (curr_col.type().id() == type_id::STRUCT || curr_col.type().id() == type_id::LIST) {
+        if (curr_col.type().id() == type_id::STRUCT) {
+          if (curr_col.num_child_columns() == 0) { return _seed; }
+          // Non-empty structs are decomposed and contain only one child.
+          curr_col = structs_column_device_view(curr_col).get_sliced_child(0);
+        } else {
+          curr_col = lists_column_device_view(curr_col).get_sliced_child();
+        }
+      }
+
+      return cudf::detail::accumulate(
+        thrust::counting_iterator(0),
+        thrust::counting_iterator(curr_col.size()),
+        _seed,
+        cuda::proclaim_return_type<result_type>(
+          [curr_col, nulls = this->_check_nulls] __device__(auto hash, auto element_index) {
+            auto const hasher = hash_functor{nulls, hash, hash};
+            return cudf::type_dispatcher<cudf::detail::dispatch_void_if_nested>(
+              curr_col.type(), hasher, curr_col, element_index);
+          }));
+    }
+
+   private:
+    Nullate const _check_nulls;
+    result_type const _seed;
+  };
+
+  CUDF_HOST_DEVICE spark_device_row_hasher(Nullate check_nulls,
+                                           table_device_view table,
+                                           result_type seed = DEFAULT_HASH_SEED) noexcept
+    : _check_nulls{check_nulls}, _table{table}, _seed(seed)
+  {
+  }
+
+  Nullate const _check_nulls;
+  table_device_view const _table;
+  result_type const _seed;
+};
+
+}  // namespace detail::row::hash
+}  // namespace CUDF_EXPORT cudf
diff --git a/cpp/include/cudf/hashing.hpp b/cpp/include/cudf/hashing.hpp
index 669f4444751c..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.
@@ -39,9 +41,8 @@ namespace hashing {
 /**
  * @brief Computes the MurmurHash3 32-bit hash value of each row in the given table
  *
- * This function computes the hash of each column using the `seed` for the first column
- * and the resulting hash as a seed for the next column and so on.
- * The result is a uint32 value for each row.
+ * This function hashes each column using the same initial `seed`, then combines the column
+ * hashes into a single uint32 value for each row.
  *
  * @param input The table of columns to hash
  * @param seed Optional seed value to use for the hash function
@@ -56,6 +57,29 @@ std::unique_ptr<column> murmurhash3_x86_32(
   rmm::cuda_stream_view stream      = cudf::get_default_stream(),
   rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
 
+/**
+ * @brief Computes the Apache Spark-compatible MurmurHash3 32-bit hash of each row
+ *
+ * This function follows Apache Spark value hashing and row traversal semantics. Each non-null value
+ * is hashed using the preceding value hash as its seed. Null values leave the current hash
+ * unchanged. Spark-specific handling is applied to strings, narrow integral types, floating-point
+ * NaNs and signed zeros, fixed-point values, lists, and structs.
+ *
+ * This function does not support LIST columns whose child is a STRUCT.
+ *
+ * @param input The table of columns to hash
+ * @param seed Optional initial seed value
+ * @param stream CUDA stream used for device memory operations and kernel launches
+ * @param mr Device memory resource used to allocate the returned column device memory
+ *
+ * @returns A non-nullable INT32 column containing one Spark-compatible hash per input row
+ */
+std::unique_ptr<column> spark_murmurhash3_x86_32(
+  table_view const& input,
+  uint32_t seed                     = DEFAULT_HASH_SEED,
+  cuda::stream_ref stream           = cudf::get_default_stream(),
+  rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
+
 /**
  * @brief Computes the MurmurHash3 64-bit hash value of each row in the given table
  *
diff --git a/cpp/include/cudf/hashing/detail/hashing.hpp b/cpp/include/cudf/hashing/detail/hashing.hpp
index 567bb864628a..18ade03e6d7b 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,
+                                                 cuda::stream_ref,
+                                                 rmm::device_async_resource_ref mr);
+
 std::unique_ptr<table> murmurhash3_x64_128(table_view const& input,
                                            uint64_t seed,
                                            rmm::cuda_stream_view,
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..63fffa882862
--- /dev/null
+++ b/cpp/include/cudf/hashing/detail/spark_murmurhash3.cuh
@@ -0,0 +1,239 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#pragma once
+
+#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 <cuda/std/algorithm>
+#include <cuda/std/cstddef>
+#include <cuda/std/iterator>
+#include <thrust/execution_policy.h>
+#include <thrust/find.h>
+#include <thrust/reverse.h>
+
+namespace cudf::hashing::detail {
+
+template <typename Key>
+  requires(not cudf::is_nested<Key>())
+struct Spark_MurmurHash3_x86_32 {
+  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)
+    : 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));
+  }
+
+  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).
+    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(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;
+    uint32_t 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(h, rot_c2);
+      h = h * 5 + c3;
+    }
+
+    h = compute_remaining_bytes(data, len, tail_offset, h);
+
+    // Finalize hash.
+    h ^= static_cast<uint32_t>(len);
+    h = fmix32(h);
+    return static_cast<result_type>(h);
+  }
+
+ private:
+  uint32_t m_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 <>
+__device__ inline auto Spark_MurmurHash3_x86_32<bool>::operator()(bool const& key) const
+  -> result_type
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+__device__ inline auto Spark_MurmurHash3_x86_32<int8_t>::operator()(int8_t const& key) const
+  -> result_type
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+__device__ inline auto Spark_MurmurHash3_x86_32<uint8_t>::operator()(uint8_t const& key) const
+  -> result_type
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+__device__ inline auto Spark_MurmurHash3_x86_32<int16_t>::operator()(int16_t const& key) const
+  -> result_type
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+__device__ inline auto Spark_MurmurHash3_x86_32<uint16_t>::operator()(uint16_t const& key) const
+  -> result_type
+{
+  return compute<uint32_t>(key);
+}
+
+template <>
+__device__ inline auto Spark_MurmurHash3_x86_32<float>::operator()(float const& key) const
+  -> result_type
+{
+  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_and_zeros(key));
+}
+
+template <>
+__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();
+  return compute_bytes(data, len);
+}
+
+template <>
+__device__ inline auto Spark_MurmurHash3_x86_32<numeric::decimal32>::operator()(
+  numeric::decimal32 const& key) const -> result_type
+{
+  return compute<uint64_t>(key.value());
+}
+
+template <>
+__device__ inline auto Spark_MurmurHash3_x86_32<numeric::decimal64>::operator()(
+  numeric::decimal64 const& key) const -> result_type
+{
+  return compute<uint64_t>(key.value());
+}
+
+template <>
+__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()
+  // 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);
+
+  // 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);
+}
+
+}  // namespace cudf::hashing::detail
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..2eb4ed182cec
--- /dev/null
+++ b/cpp/src/hash/spark_murmurhash3_x86_32.cu
@@ -0,0 +1,96 @@
+/*
+ * 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/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>
+
+#include <rmm/cuda_stream_view.hpp>
+
+#include <cub/device/device_for.cuh>
+#include <cuda/stream_ref>
+
+#include <functional>
+
+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&)>;
+
+  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);
+  }
+}
+
+}  // namespace
+
+std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
+                                                 uint32_t seed,
+                                                 cuda::stream_ref stream,
+                                                 rmm::device_async_resource_ref 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; }
+
+  // Lists of structs are not supported
+  check_spark_murmurhash3_compatibility(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>();
+
+  // Compute the hash value for each row
+  auto const hasher =
+    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); },
+    stream.get()));
+
+  return output;
+}
+
+}  // namespace detail
+
+std::unique_ptr<column> spark_murmurhash3_x86_32(table_view const& input,
+                                                 uint32_t seed,
+                                                 cuda::stream_ref 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/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt
index 6b27b65b18a2..4ee086555ed7 100644
--- a/cpp/tests/CMakeLists.txt
+++ b/cpp/tests/CMakeLists.txt
@@ -199,6 +199,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/spark_murmurhash3_x86_32_test.cpp b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
new file mode 100644
index 000000000000..e465db85851c
--- /dev/null
+++ b/cpp/tests/hashing/spark_murmurhash3_x86_32_test.cpp
@@ -0,0 +1,606 @@
+/*
+ * 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/copying.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 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});
+  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. 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._
+  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, -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(
+    {-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, 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(
+    {-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, 1728350528, 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("structs", "hash(structs) 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>(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(
+    8, offsets.release(), struct_column.release(), null_count, std::move(null_mask));
+
+  // 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);
+}
diff --git a/cpp/tests/row_operator/row_operator_tests.cu b/cpp/tests/row_operator/row_operator_tests.cu
index 073c2dd58989..cd3cde9f921b 100644
--- a/cpp/tests/row_operator/row_operator_tests.cu
+++ b/cpp/tests/row_operator/row_operator_tests.cu
@@ -14,6 +14,8 @@
 #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>
 
@@ -453,6 +455,31 @@ TEST_F(RowOperatorTest, TestRowHasher64BitHash)
     results, expected, cudf::test::debug_output_level::FIRST_ERROR, stream, mr);
 }
 
+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::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),
+                    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 stream = this->stream();
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()