diff --git a/cpp/include/cudf/detail/hashing.hpp b/cpp/include/cudf/detail/hashing.hpp
index 18f97f00f21..06f523c2320 100644
--- a/cpp/include/cudf/detail/hashing.hpp
+++ b/cpp/include/cudf/detail/hashing.hpp
@@ -46,6 +46,7 @@ std::unique_ptr<column> md5_hash(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+template <template <typename> class hash_function>
 std::unique_ptr<column> serial_murmur_hash3_32(
   table_view const& input,
   uint32_t seed                       = 0,
diff --git a/cpp/include/cudf/detail/utilities/hash_functions.cuh b/cpp/include/cudf/detail/utilities/hash_functions.cuh
index 08d88aee137..e9d66d125dd 100644
--- a/cpp/include/cudf/detail/utilities/hash_functions.cuh
+++ b/cpp/include/cudf/detail/utilities/hash_functions.cuh
@@ -111,7 +111,7 @@ void CUDA_DEVICE_CALLABLE md5_process(TKey const& key, md5_intermediate_data* ha
 /**
  * Normalization of floating point NANs and zeros helper
  */
-template <typename T, typename std::enable_if_t<std::is_floating_point<T>::value>* = nullptr>
+template <typename T, std::enable_if_t<std::is_floating_point<T>::value>* = nullptr>
 T CUDA_DEVICE_CALLABLE normalize_nans_and_zeros_helper(T key)
 {
   if (isnan(key)) {
@@ -148,7 +148,7 @@ void CUDA_DEVICE_CALLABLE uint32ToLowercaseHexString(uint32_t num, char* destina
 }
 
 struct MD5ListHasher {
-  template <typename T, typename std::enable_if_t<is_chrono<T>()>* = nullptr>
+  template <typename T, std::enable_if_t<is_chrono<T>()>* = nullptr>
   void __device__ operator()(column_device_view data_col,
                              size_type offset_begin,
                              size_type offset_end,
@@ -157,7 +157,7 @@ struct MD5ListHasher {
     release_assert(false && "MD5 Unsupported chrono type column");
   }
 
-  template <typename T, typename std::enable_if_t<!is_fixed_width<T>()>* = nullptr>
+  template <typename T, std::enable_if_t<!is_fixed_width<T>()>* = nullptr>
   void __device__ operator()(column_device_view data_col,
                              size_type offset_begin,
                              size_type offset_end,
@@ -166,7 +166,7 @@ struct MD5ListHasher {
     release_assert(false && "MD5 Unsupported non-fixed-width type column");
   }
 
-  template <typename T, typename std::enable_if_t<is_floating_point<T>()>* = nullptr>
+  template <typename T, std::enable_if_t<is_floating_point<T>()>* = nullptr>
   void __device__ operator()(column_device_view data_col,
                              size_type offset_begin,
                              size_type offset_end,
@@ -179,9 +179,9 @@ struct MD5ListHasher {
     }
   }
 
-  template <typename T,
-            typename std::enable_if_t<is_fixed_width<T>() && !is_floating_point<T>() &&
-                                      !is_chrono<T>()>* = nullptr>
+  template <
+    typename T,
+    std::enable_if_t<is_fixed_width<T>() && !is_floating_point<T>() && !is_chrono<T>()>* = nullptr>
   void CUDA_DEVICE_CALLABLE operator()(column_device_view data_col,
                                        size_type offset_begin,
                                        size_type offset_end,
@@ -268,7 +268,7 @@ struct MD5Hash {
       uint32ToLowercaseHexString(hash_state->hash_value[i], result_location + (8 * i));
   }
 
-  template <typename T, typename std::enable_if_t<is_chrono<T>()>* = nullptr>
+  template <typename T, std::enable_if_t<is_chrono<T>()>* = nullptr>
   void __device__ operator()(column_device_view col,
                              size_type row_index,
                              md5_intermediate_data* hash_state) const
@@ -276,7 +276,7 @@ struct MD5Hash {
     release_assert(false && "MD5 Unsupported chrono type column");
   }
 
-  template <typename T, typename std::enable_if_t<!is_fixed_width<T>()>* = nullptr>
+  template <typename T, std::enable_if_t<!is_fixed_width<T>()>* = nullptr>
   void __device__ operator()(column_device_view col,
                              size_type row_index,
                              md5_intermediate_data* hash_state) const
@@ -284,7 +284,7 @@ struct MD5Hash {
     release_assert(false && "MD5 Unsupported non-fixed-width type column");
   }
 
-  template <typename T, typename std::enable_if_t<is_floating_point<T>()>* = nullptr>
+  template <typename T, std::enable_if_t<is_floating_point<T>()>* = nullptr>
   void __device__ operator()(column_device_view col,
                              size_type row_index,
                              md5_intermediate_data* hash_state) const
@@ -292,9 +292,9 @@ struct MD5Hash {
     md5_process(normalize_nans_and_zeros_helper<T>(col.element<T>(row_index)), hash_state);
   }
 
-  template <typename T,
-            typename std::enable_if_t<is_fixed_width<T>() && !is_floating_point<T>() &&
-                                      !is_chrono<T>()>* = nullptr>
+  template <
+    typename T,
+    std::enable_if_t<is_fixed_width<T>() && !is_floating_point<T>() && !is_chrono<T>()>* = nullptr>
   void CUDA_DEVICE_CALLABLE operator()(column_device_view col,
                                        size_type row_index,
                                        md5_intermediate_data* hash_state) const
@@ -370,15 +370,15 @@ struct MurmurHash3_32 {
   using argument_type = Key;
   using result_type   = hash_value_type;
 
-  CUDA_HOST_DEVICE_CALLABLE MurmurHash3_32() : m_seed(0) {}
+  MurmurHash3_32() = default;
   CUDA_HOST_DEVICE_CALLABLE MurmurHash3_32(uint32_t seed) : m_seed(seed) {}
 
-  CUDA_HOST_DEVICE_CALLABLE uint32_t rotl32(uint32_t x, int8_t r) const
+  CUDA_DEVICE_CALLABLE uint32_t rotl32(uint32_t x, int8_t r) const
   {
     return (x << r) | (x >> (32 - r));
   }
 
-  CUDA_HOST_DEVICE_CALLABLE uint32_t fmix32(uint32_t h) const
+  CUDA_DEVICE_CALLABLE uint32_t fmix32(uint32_t h) const
   {
     h ^= h >> 16;
     h *= 0x85ebca6b;
@@ -399,7 +399,7 @@ struct MurmurHash3_32 {
    *
    * @returns A hash value that intelligently combines the lhs and rhs hash values
    */
-  CUDA_HOST_DEVICE_CALLABLE result_type hash_combine(result_type lhs, result_type rhs)
+  CUDA_DEVICE_CALLABLE result_type hash_combine(result_type lhs, result_type rhs)
   {
     result_type combined{lhs};
 
@@ -408,11 +408,11 @@ struct MurmurHash3_32 {
     return combined;
   }
 
-  result_type CUDA_HOST_DEVICE_CALLABLE operator()(Key const& key) const { return compute(key); }
+  result_type CUDA_DEVICE_CALLABLE operator()(Key const& key) const { return compute(key); }
 
   // compute wrapper for floating point types
-  template <typename T, typename std::enable_if_t<std::is_floating_point<T>::value>* = nullptr>
-  hash_value_type CUDA_HOST_DEVICE_CALLABLE compute_floating_point(T const& key) const
+  template <typename T, std::enable_if_t<std::is_floating_point<T>::value>* = nullptr>
+  hash_value_type CUDA_DEVICE_CALLABLE compute_floating_point(T const& key) const
   {
     if (key == T{0.0}) {
       return compute(T{0.0});
@@ -425,10 +425,10 @@ struct MurmurHash3_32 {
   }
 
   template <typename TKey>
-  result_type CUDA_HOST_DEVICE_CALLABLE compute(TKey const& key) const
+  result_type CUDA_DEVICE_CALLABLE compute(TKey const& key) const
   {
     constexpr int len         = sizeof(argument_type);
-    const uint8_t* const data = (const uint8_t*)&key;
+    uint8_t const* const data = reinterpret_cast<uint8_t const*>(&key);
     constexpr int nblocks     = len / 4;
 
     uint32_t h1           = m_seed;
@@ -436,7 +436,7 @@ struct MurmurHash3_32 {
     constexpr uint32_t c2 = 0x1b873593;
     //----------
     // body
-    const uint32_t* const blocks = (const uint32_t*)(data + nblocks * 4);
+    uint32_t const* const blocks = reinterpret_cast<uint32_t const*>(data + nblocks * 4);
     for (int i = -nblocks; i; i++) {
       uint32_t k1 = blocks[i];  // getblock32(blocks,i);
       k1 *= c1;
@@ -448,7 +448,7 @@ struct MurmurHash3_32 {
     }
     //----------
     // tail
-    const uint8_t* tail = (const uint8_t*)(data + nblocks * 4);
+    uint8_t const* tail = reinterpret_cast<uint8_t const*>(data + nblocks * 4);
     uint32_t k1         = 0;
     switch (len & 3) {
       case 3: k1 ^= tail[2] << 16;
@@ -468,11 +468,11 @@ struct MurmurHash3_32 {
   }
 
  private:
-  uint32_t m_seed;
+  uint32_t m_seed{0};
 };
 
 template <>
-hash_value_type CUDA_HOST_DEVICE_CALLABLE MurmurHash3_32<bool>::operator()(bool const& key) const
+hash_value_type CUDA_DEVICE_CALLABLE MurmurHash3_32<bool>::operator()(bool const& key) const
 {
   return this->compute(static_cast<uint8_t>(key));
 }
@@ -481,28 +481,24 @@ hash_value_type CUDA_HOST_DEVICE_CALLABLE MurmurHash3_32<bool>::operator()(bool
  * @brief Specialization of MurmurHash3_32 operator for strings.
  */
 template <>
-hash_value_type CUDA_HOST_DEVICE_CALLABLE
+hash_value_type CUDA_DEVICE_CALLABLE
 MurmurHash3_32<cudf::string_view>::operator()(cudf::string_view const& key) const
 {
-  const int len         = (int)key.size_bytes();
-  const uint8_t* data   = (const uint8_t*)key.data();
-  const int nblocks     = len / 4;
+  auto const len        = key.size_bytes();
+  uint8_t const* data   = reinterpret_cast<uint8_t const*>(key.data());
+  int const nblocks     = len / 4;
   result_type h1        = m_seed;
   constexpr uint32_t c1 = 0xcc9e2d51;
   constexpr uint32_t c2 = 0x1b873593;
-  auto getblock32       = [] __host__ __device__(const uint32_t* p, int i) -> uint32_t {
-  // Individual byte reads for unaligned accesses (very likely)
-#ifndef __CUDA_ARCH__
-    CUDF_FAIL("Hashing a string in host code is not supported.");
-#else
-    auto q = (const uint8_t*)(p + i);
+  auto getblock32       = [] __device__(uint32_t const* p, int i) -> uint32_t {
+    // Individual byte reads for unaligned accesses (very likely)
+    auto q = (uint8_t const*)(p + i);
     return q[0] | (q[1] << 8) | (q[2] << 16) | (q[3] << 24);
-#endif
   };
 
   //----------
   // body
-  const uint32_t* const blocks = (const uint32_t*)(data + nblocks * 4);
+  uint32_t const* const blocks = reinterpret_cast<uint32_t const*>(data + nblocks * 4);
   for (int i = -nblocks; i; i++) {
     uint32_t k1 = getblock32(blocks, i);
     k1 *= c1;
@@ -514,7 +510,7 @@ MurmurHash3_32<cudf::string_view>::operator()(cudf::string_view const& key) cons
   }
   //----------
   // tail
-  const uint8_t* tail = (const uint8_t*)(data + nblocks * 4);
+  uint8_t const* tail = reinterpret_cast<uint8_t const*>(data + nblocks * 4);
   uint32_t k1         = 0;
   switch (len & 3) {
     case 3: k1 ^= tail[2] << 16;
@@ -534,14 +530,164 @@ MurmurHash3_32<cudf::string_view>::operator()(cudf::string_view const& key) cons
 }
 
 template <>
-hash_value_type CUDA_HOST_DEVICE_CALLABLE MurmurHash3_32<float>::operator()(float const& key) const
+hash_value_type CUDA_DEVICE_CALLABLE MurmurHash3_32<float>::operator()(float const& key) const
+{
+  return this->compute_floating_point(key);
+}
+
+template <>
+hash_value_type CUDA_DEVICE_CALLABLE MurmurHash3_32<double>::operator()(double const& key) const
+{
+  return this->compute_floating_point(key);
+}
+
+template <typename Key>
+struct SparkMurmurHash3_32 {
+  using argument_type = Key;
+  using result_type   = hash_value_type;
+
+  SparkMurmurHash3_32() = default;
+  CUDA_HOST_DEVICE_CALLABLE SparkMurmurHash3_32(uint32_t seed) : m_seed(seed) {}
+
+  CUDA_DEVICE_CALLABLE uint32_t rotl32(uint32_t x, int8_t r) const
+  {
+    return (x << r) | (x >> (32 - r));
+  }
+
+  CUDA_DEVICE_CALLABLE uint32_t fmix32(uint32_t h) const
+  {
+    h ^= h >> 16;
+    h *= 0x85ebca6b;
+    h ^= h >> 13;
+    h *= 0xc2b2ae35;
+    h ^= h >> 16;
+    return h;
+  }
+
+  result_type CUDA_DEVICE_CALLABLE operator()(Key const& key) const { return compute(key); }
+
+  // compute wrapper for floating point types
+  template <typename T, std::enable_if_t<std::is_floating_point<T>::value>* = nullptr>
+  hash_value_type CUDA_DEVICE_CALLABLE compute_floating_point(T const& key) const
+  {
+    if (key == T{0.0}) {
+      return compute(T{0.0});
+    } else if (isnan(key)) {
+      T nan = std::numeric_limits<T>::quiet_NaN();
+      return compute(nan);
+    } else {
+      return compute(key);
+    }
+  }
+
+  template <typename TKey>
+  result_type CUDA_DEVICE_CALLABLE compute(TKey const& key) const
+  {
+    constexpr int len        = sizeof(TKey);
+    int8_t const* const data = reinterpret_cast<int8_t const*>(&key);
+    constexpr int nblocks    = len / 4;
+
+    uint32_t h1           = m_seed;
+    constexpr uint32_t c1 = 0xcc9e2d51;
+    constexpr uint32_t c2 = 0x1b873593;
+    //----------
+    // body
+    uint32_t const* const blocks = reinterpret_cast<uint32_t const*>(data + nblocks * 4);
+    for (int i = -nblocks; i; i++) {
+      uint32_t k1 = blocks[i];
+      k1 *= c1;
+      k1 = rotl32(k1, 15);
+      k1 *= c2;
+      h1 ^= k1;
+      h1 = rotl32(h1, 13);
+      h1 = h1 * 5 + 0xe6546b64;
+    }
+    //----------
+    // byte by byte tail processing
+    for (int i = nblocks * 4; i < len; i++) {
+      int32_t k1 = data[i];
+      k1 *= c1;
+      k1 = rotl32(k1, 15);
+      k1 *= c2;
+      h1 ^= k1;
+      h1 = rotl32(h1, 13);
+      h1 = h1 * 5 + 0xe6546b64;
+    }
+    //----------
+    // finalization
+    h1 ^= len;
+    h1 = fmix32(h1);
+    return h1;
+  }
+
+ private:
+  uint32_t m_seed{0};
+};
+
+template <>
+hash_value_type CUDA_DEVICE_CALLABLE SparkMurmurHash3_32<bool>::operator()(bool const& key) const
+{
+  return this->compute<uint32_t>(key);
+}
+
+/**
+ * @brief Specialization of MurmurHash3_32 operator for strings.
+ */
+template <>
+hash_value_type CUDA_DEVICE_CALLABLE
+SparkMurmurHash3_32<cudf::string_view>::operator()(cudf::string_view const& key) const
+{
+  auto const len        = key.size_bytes();
+  int8_t const* data    = reinterpret_cast<int8_t const*>(key.data());
+  int const nblocks     = len / 4;
+  result_type h1        = m_seed;
+  constexpr uint32_t c1 = 0xcc9e2d51;
+  constexpr uint32_t c2 = 0x1b873593;
+  auto getblock32       = [] __device__(uint32_t const* p, int i) -> uint32_t {
+    // Individual byte reads for unaligned accesses (very likely)
+    auto q = (const uint8_t*)(p + i);
+    return q[0] | (q[1] << 8) | (q[2] << 16) | (q[3] << 24);
+  };
+
+  //----------
+  // body
+  uint32_t const* const blocks = reinterpret_cast<uint32_t const*>(data + nblocks * 4);
+  for (int i = -nblocks; i; i++) {
+    uint32_t k1 = getblock32(blocks, i);
+    k1 *= c1;
+    k1 = rotl32(k1, 15);
+    k1 *= c2;
+    h1 ^= k1;
+    h1 = rotl32(h1, 13);
+    h1 = h1 * 5 + 0xe6546b64;
+  }
+  //----------
+  // Spark's byte by byte tail processing
+  for (int i = nblocks * 4; i < len; i++) {
+    int32_t k1 = data[i];
+    k1 *= c1;
+    k1 = rotl32(k1, 15);
+    k1 *= c2;
+    h1 ^= k1;
+    h1 = rotl32(h1, 13);
+    h1 = h1 * 5 + 0xe6546b64;
+  }
+  //----------
+  // finalization
+  h1 ^= len;
+  h1 = fmix32(h1);
+  return h1;
+}
+
+template <>
+hash_value_type CUDA_DEVICE_CALLABLE SparkMurmurHash3_32<float>::operator()(float const& key) const
 {
   return this->compute_floating_point(key);
 }
 
 template <>
-hash_value_type CUDA_HOST_DEVICE_CALLABLE
-MurmurHash3_32<double>::operator()(double const& key) const
+hash_value_type CUDA_DEVICE_CALLABLE
+SparkMurmurHash3_32<double>::operator()(double const& key) const
 {
   return this->compute_floating_point(key);
 }
@@ -576,7 +722,7 @@ struct IdentityHash {
 
   template <typename return_type = result_type>
   CUDA_HOST_DEVICE_CALLABLE std::enable_if_t<!std::is_arithmetic<Key>::value, return_type>
-  operator()(const Key& key) const
+  operator()(Key const& key) const
   {
     release_assert(false && "IdentityHash does not support this data type");
     return 0;
@@ -584,7 +730,7 @@ struct IdentityHash {
 
   template <typename return_type = result_type>
   CUDA_HOST_DEVICE_CALLABLE std::enable_if_t<std::is_arithmetic<Key>::value, return_type>
-  operator()(const Key& key) const
+  operator()(Key const& key) const
   {
     return static_cast<result_type>(key);
   }
diff --git a/cpp/include/cudf/types.hpp b/cpp/include/cudf/types.hpp
index a08ae366f68..48e5d9543b8 100644
--- a/cpp/include/cudf/types.hpp
+++ b/cpp/include/cudf/types.hpp
@@ -324,10 +324,11 @@ std::size_t size_of(data_type t);
  *  @brief Identifies the hash function to be used
  */
 enum class hash_id {
-  HASH_IDENTITY = 0,   ///< Identity hash function that simply returns the key to be hashed
-  HASH_MURMUR3,        ///< Murmur3 hash function
-  HASH_MD5,            ///< MD5 hash function
-  HASH_SERIAL_MURMUR3  ///< Serial Murmur3 hash function
+  HASH_IDENTITY = 0,    ///< Identity hash function that simply returns the key to be hashed
+  HASH_MURMUR3,         ///< Murmur3 hash function
+  HASH_MD5,             ///< MD5 hash function
+  HASH_SERIAL_MURMUR3,  ///< Serial Murmur3 hash function
+  HASH_SPARK_MURMUR3    ///< Spark Murmur3 hash function
 };
 
 /** @} */
diff --git a/cpp/src/hash/hashing.cu b/cpp/src/hash/hashing.cu
index 13ee340ba27..16efb666b3e 100644
--- a/cpp/src/hash/hashing.cu
+++ b/cpp/src/hash/hashing.cu
@@ -52,7 +52,10 @@ std::unique_ptr<column> hash(table_view const& input,
   switch (hash_function) {
     case (hash_id::HASH_MURMUR3): return murmur_hash3_32(input, initial_hash, stream, mr);
     case (hash_id::HASH_MD5): return md5_hash(input, stream, mr);
-    case (hash_id::HASH_SERIAL_MURMUR3): return serial_murmur_hash3_32(input, seed, stream, mr);
+    case (hash_id::HASH_SERIAL_MURMUR3):
+      return serial_murmur_hash3_32<MurmurHash3_32>(input, seed, stream, mr);
+    case (hash_id::HASH_SPARK_MURMUR3):
+      return serial_murmur_hash3_32<SparkMurmurHash3_32>(input, seed, stream, mr);
     default: return nullptr;
   }
 }
@@ -119,6 +122,7 @@ std::unique_ptr<column> md5_hash(table_view const& input,
                              mr);
 }
 
+template <template <typename> class hash_function>
 std::unique_ptr<column> serial_murmur_hash3_32(table_view const& input,
                                                uint32_t seed,
                                                rmm::cuda_stream_view stream,
@@ -145,7 +149,7 @@ std::unique_ptr<column> serial_murmur_hash3_32(table_view const& input,
                          [rindex = row_index] __device__(auto hash, auto column) {
                            return cudf::type_dispatcher(
                              column.type(),
-                             element_hasher_with_seed<MurmurHash3_32, true>{hash, hash},
+                             element_hasher_with_seed<hash_function, true>{hash, hash},
                              column,
                              rindex);
                          });
@@ -163,7 +167,7 @@ std::unique_ptr<column> serial_murmur_hash3_32(table_view const& input,
                          [rindex = row_index] __device__(auto hash, auto column) {
                            return cudf::type_dispatcher(
                              column.type(),
-                             element_hasher_with_seed<MurmurHash3_32, false>{hash, hash},
+                             element_hasher_with_seed<hash_function, false>{hash, hash},
                              column,
                              rindex);
                          });
diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index 12cea8bc0de..88415d820d6 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -115,6 +115,12 @@ TEST_F(HashTest, MultiValueNulls)
 
   EXPECT_EQ(input1.num_rows(), serial_output1->size());
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(serial_output1->view(), serial_output2->view());
+
+  auto const spark_output1 = cudf::hash(input1, cudf::hash_id::HASH_SPARK_MURMUR3, {}, 0);
+  auto const spark_output2 = cudf::hash(input2, cudf::hash_id::HASH_SPARK_MURMUR3);
+
+  EXPECT_EQ(input1.num_rows(), spark_output1->size());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(spark_output1->view(), spark_output2->view());
 }
 
 template <typename T>
@@ -140,6 +146,12 @@ TYPED_TEST(HashTestTyped, Equality)
 
   EXPECT_EQ(input.num_rows(), serial_output1->size());
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(serial_output1->view(), serial_output2->view());
+
+  auto const spark_output1 = cudf::hash(input, cudf::hash_id::HASH_SPARK_MURMUR3, {}, 0);
+  auto const spark_output2 = cudf::hash(input, cudf::hash_id::HASH_SPARK_MURMUR3);
+
+  EXPECT_EQ(input.num_rows(), spark_output1->size());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(spark_output1->view(), spark_output2->view());
 }
 
 TYPED_TEST(HashTestTyped, EqualityNulls)
@@ -164,6 +176,12 @@ TYPED_TEST(HashTestTyped, EqualityNulls)
 
   EXPECT_EQ(input1.num_rows(), serial_output1->size());
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(serial_output1->view(), serial_output2->view());
+
+  auto const spark_output1 = cudf::hash(input1, cudf::hash_id::HASH_SPARK_MURMUR3, {}, 0);
+  auto const spark_output2 = cudf::hash(input2, cudf::hash_id::HASH_SPARK_MURMUR3);
+
+  EXPECT_EQ(input1.num_rows(), spark_output1->size());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(spark_output1->view(), spark_output2->view());
 }
 
 template <typename T>
@@ -196,6 +214,11 @@ TYPED_TEST(HashTestFloatTyped, TestExtremes)
   auto const serial_output2 = cudf::hash(input2, cudf::hash_id::HASH_SERIAL_MURMUR3);
 
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(serial_output1->view(), serial_output2->view());
+
+  auto const spark_output1 = cudf::hash(input1, cudf::hash_id::HASH_SPARK_MURMUR3, {}, 0);
+  auto const spark_output2 = cudf::hash(input2, cudf::hash_id::HASH_SPARK_MURMUR3);
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(spark_output1->view(), spark_output2->view());
 }
 
 class SerialMurmurHash3Test : public cudf::test::BaseFixture {
@@ -204,7 +227,7 @@ class SerialMurmurHash3Test : public cudf::test::BaseFixture {
 TEST_F(SerialMurmurHash3Test, MultiValueWithSeeds)
 {
   fixed_width_column_wrapper<int32_t> const strings_col_result(
-    {1467149710, -680899318, -1620282500, 91106683, -1165223012});
+    {1467149710, -680899318, -1620282500, 91106683, -1564993834});
   fixed_width_column_wrapper<int32_t> const ints_col_result(
     {933211791, 751823303, -1080202046, 723455942, 133916647});
 
@@ -212,7 +235,7 @@ TEST_F(SerialMurmurHash3Test, MultiValueWithSeeds)
                                             "The quick brown fox",
                                             "jumps over the lazy dog.",
                                             "All work and no play makes Jack a dull boy",
-                                            "!\"#$%&\'()*+,-./0123456789:;<=>?@[\\]^_`{|}~"});
+                                            "!\"#$%&\'()*+,-./]:;<=>?@[\\]^_`{|}~\ud720\ud721"});
 
   using limits = std::numeric_limits<int32_t>;
   fixed_width_column_wrapper<int32_t> const ints_col({0, 100, -100, limits::min(), limits::max()});
@@ -236,6 +259,44 @@ TEST_F(SerialMurmurHash3Test, MultiValueWithSeeds)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(hashed_output3->view(), hashed_output4->view(), true);
 }
 
+class SparkMurmurHash3Test : public cudf::test::BaseFixture {
+};
+
+TEST_F(SparkMurmurHash3Test, MultiValueWithSeeds)
+{
+  fixed_width_column_wrapper<int32_t> const strings_col_result(
+    {1467149710, 723257560, -1620282500, -2001858707, 1588473657});
+  fixed_width_column_wrapper<int32_t> const ints_col_result(
+    {933211791, 751823303, -1080202046, 723455942, 133916647});
+
+  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"});
+
+  using limits = std::numeric_limits<int32_t>;
+  fixed_width_column_wrapper<int32_t> const ints_col({0, 100, -100, limits::min(), limits::max()});
+
+  fixed_width_column_wrapper<bool> const bools_col1({0, 1, 1, 1, 0});
+  fixed_width_column_wrapper<bool> const bools_col2({0, 1, 2, 255, 0});
+
+  auto const input1 = cudf::table_view({strings_col});
+  auto const input2 = cudf::table_view({ints_col});
+  auto const input3 = cudf::table_view({strings_col, ints_col, bools_col1});
+  auto const input4 = cudf::table_view({strings_col, ints_col, bools_col2});
+
+  auto const hashed_output1 = cudf::hash(input1, cudf::hash_id::HASH_SPARK_MURMUR3, {}, 314);
+  auto const hashed_output2 = cudf::hash(input2, cudf::hash_id::HASH_SPARK_MURMUR3, {}, 42);
+  auto const hashed_output3 = cudf::hash(input3, cudf::hash_id::HASH_SPARK_MURMUR3, {});
+  auto const hashed_output4 = cudf::hash(input4, cudf::hash_id::HASH_SPARK_MURMUR3, {});
+
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(hashed_output1->view(), strings_col_result, true);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(hashed_output2->view(), ints_col_result, true);
+  EXPECT_EQ(input3.num_rows(), hashed_output3->size());
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(hashed_output3->view(), hashed_output4->view(), true);
+}
+
 class MD5HashTest : public cudf::test::BaseFixture {
 };
 
diff --git a/java/src/main/java/ai/rapids/cudf/ColumnVector.java b/java/src/main/java/ai/rapids/cudf/ColumnVector.java
index 926fba30418..34a2e2fe374 100644
--- a/java/src/main/java/ai/rapids/cudf/ColumnVector.java
+++ b/java/src/main/java/ai/rapids/cudf/ColumnVector.java
@@ -435,15 +435,15 @@ public static ColumnVector md5Hash(ColumnView... columns) {
   }
 
   /**
-   * Create a new vector containing the MD5 hash of each row in the table.
+   * Create a new vector containing the murmur3 hash of each row in the table.
    *
    * @param seed integer seed for the murmur3 hash function
    * @param columns array of columns to hash, must have identical number of rows.
-   * @return the new ColumnVector of 32 character hex strings representing each row's hash value.
+   * @return the new ColumnVector of 32-bit values representing each row's hash value.
    */
   public static ColumnVector serial32BitMurmurHash3(int seed, ColumnView columns[]) {
     if (columns.length < 1) {
-      throw new IllegalArgumentException("MD5 hashing requires at least 1 column of input");
+      throw new IllegalArgumentException("Murmur3 hashing requires at least 1 column of input");
     }
     long[] columnViews = new long[columns.length];
     long size = columns[0].getRowCount();
@@ -460,15 +460,52 @@ public static ColumnVector serial32BitMurmurHash3(int seed, ColumnView columns[]
   }
 
   /**
-   * Create a new vector containing the MD5 hash of each row in the table, seed defaulted to 0.
+   * Create a new vector containing the murmur3 hash of each row in the table, seed defaulted to 0.
    *
    * @param columns array of columns to hash, must have identical number of rows.
-   * @return the new ColumnVector of 32 character hex strings representing each row's hash value.
+   * @return the new ColumnVector of 32-bit values representing each row's hash value.
    */
   public static ColumnVector serial32BitMurmurHash3(ColumnView columns[]) {
     return serial32BitMurmurHash3(0, columns);
   }
 
+  /**
+   * Create a new vector containing spark's 32-bit murmur3 hash of each row in the table.
+   * Spark's murmur3 hash uses a different tail processing algorithm.
+   *
+   * @param seed integer seed for the murmur3 hash function
+   * @param columns array of columns to hash, must have identical number of rows.
+   * @return the new ColumnVector of 32-bit values representing each row's hash value.
+   */
+  public static ColumnVector spark32BitMurmurHash3(int seed, ColumnView columns[]) {
+    if (columns.length < 1) {
+      throw new IllegalArgumentException("Murmur3 hashing requires at least 1 column of input");
+    }
+    long[] columnViews = new long[columns.length];
+    long size = columns[0].getRowCount();
+
+    for(int i = 0; i < columns.length; i++) {
+      assert columns[i] != null : "Column vectors passed may not be null";
+      assert columns[i].getRowCount() == size : "Row count mismatch, all columns must be the same size";
+      assert !columns[i].getType().isDurationType() : "Unsupported column type Duration";
+      assert !columns[i].getType().isTimestampType() : "Unsupported column type Timestamp";
+      assert !columns[i].getType().isNestedType() : "Unsupported column of nested type";
+      columnViews[i] = columns[i].getNativeView();
+    }
+    return new ColumnVector(hash(columnViews, HashType.HASH_SPARK_MURMUR3.getNativeId(), new int[0], seed));
+  }
+
+  /**
+   * Create a new vector containing spark's 32-bit murmur3 hash of each row in the table with the
+   * seed set to 0. Spark's murmur3 hash uses a different tail processing algorithm.
+   *
+   * @param columns array of columns to hash, must have identical number of rows.
+   * @return the new ColumnVector of 32-bit values representing each row's hash value.
+   */
+  public static ColumnVector spark32BitMurmurHash3(ColumnView columns[]) {
+    return spark32BitMurmurHash3(0, columns);
+  }
+
   /**
    * Generic method to cast ColumnVector
    * When casting from a Date, Timestamp, or Boolean to a numerical type the underlying numerical
diff --git a/java/src/main/java/ai/rapids/cudf/HashType.java b/java/src/main/java/ai/rapids/cudf/HashType.java
index 9266545bd94..b521bc5c42c 100644
--- a/java/src/main/java/ai/rapids/cudf/HashType.java
+++ b/java/src/main/java/ai/rapids/cudf/HashType.java
@@ -25,7 +25,8 @@ public enum HashType {
   // TODO IDENTITY(0),
   // TODO MURMUR3(1),
   HASH_MD5(2),
-  HASH_SERIAL_MURMUR3(3);
+  HASH_SERIAL_MURMUR3(3),
+  HASH_SPARK_MURMUR3(4);
 
   private static final HashType[] HASH_TYPES = HashType.values();
   final int nativeId;
diff --git a/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java b/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
index e01b0b8ef1b..d3cde7f77cc 100644
--- a/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
+++ b/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
@@ -284,7 +284,7 @@ void testNormalizeNANsAndZeros() {
   @Test
   void testMD5HashStrings() {
     try (ColumnVector v0 = ColumnVector.fromStrings(
-          "a", "B\n", "dE\"\u0100\t\u0101 \u0500\u0501",
+          "a", "B\n", "dE\"\u0100\t\u0101 \ud720\ud721",
           "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
           "in the MD5 hash function. This string needed to be longer.",
           null, null);
@@ -293,7 +293,7 @@ void testMD5HashStrings() {
            "A 60 character string to test MD5's message padding algorithm",
            "hiJ\ud720\ud721\ud720\ud721", null);
          ColumnVector v2 = ColumnVector.fromStrings(
-           "a", "B\nc",  "dE\"\u0100\t\u0101 \u0500\u0501\\Fg2\'",
+           "a", "B\nc",  "dE\"\u0100\t\u0101 \ud720\ud721\\Fg2\'",
            "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
            "in the MD5 hash function. This string needed to be longer.A 60 character string to " +
            "test MD5's message padding algorithm",
@@ -302,7 +302,7 @@ void testMD5HashStrings() {
          ColumnVector result2 = ColumnVector.md5Hash(v2);
          ColumnVector expected = ColumnVector.fromStrings(
           "0cc175b9c0f1b6a831c399e269772661", "f5112705c2f6dc7d3fc6bd496df6c2e8",
-          "675c30ce6d1b27dcb5009b01be42e9bd", "8fa29148f63c1fe9248fdc4644e3a193",
+          "d49db62680847e0e7107e0937d29668e", "8fa29148f63c1fe9248fdc4644e3a193",
           "1bc221b25e6c4825929e884092f4044f", "d41d8cd98f00b204e9800998ecf8427e")) {
       assertColumnsAreEqual(result01, expected);
       assertColumnsAreEqual(result2, expected);
@@ -318,7 +318,7 @@ void testMD5HashInts() {
           "7dea362b3fac8e00956a4952a3d4f474", "cdc824bf721df654130ed7447fb878ac",
           "7fb89558e395330c6a10ab98915fcafb", "d41d8cd98f00b204e9800998ecf8427e",
           "152afd7bf4dbe26f85032eee0269201a", "ed0622e1179e101cf7edc0b952cc5262")) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -337,7 +337,7 @@ void testMD5HashDoubles() {
           "bf26d90b64827fdbc58da0aa195156fe", "bf26d90b64827fdbc58da0aa195156fe",
           "bf26d90b64827fdbc58da0aa195156fe", "bf26d90b64827fdbc58da0aa195156fe",
           "73c82437c94e197f7e35e14f0140497a", "740660a5f71e7a264fca45330c34da31")) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -356,7 +356,7 @@ void testMD5HashFloats() {
           "d6fd2bac25776d9a7269ca0e24b21b36", "d6fd2bac25776d9a7269ca0e24b21b36",
           "d6fd2bac25776d9a7269ca0e24b21b36", "d6fd2bac25776d9a7269ca0e24b21b36",
           "55e3a4db046ad9065bd7d64243de408f", "33b552ad34a825b275f5f2b59778b5c3")){
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -369,14 +369,14 @@ void testMD5HashBools() {
           "d41d8cd98f00b204e9800998ecf8427e", "249ba6277758050695e8f5909bacd6d3",
           "c4103f122d27677c9db144cae1394a66", "55a54008ad1ba589aa210d2629c1df41",
           "93b885adfe0da089cdf634904fd59f71", "441077cc9e57554dd476bdfb8b8b8102")) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
   @Test
   void testMD5HashMixed() {
     try (ColumnVector strings = ColumnVector.fromStrings(
-          "a", "B\n", "dE\"\u0100\t\u0101 \u0500\u0501",
+          "a", "B\n", "dE\"\u0100\t\u0101 \ud720\ud721",
           "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
           "in the MD5 hash function. This string needed to be longer.",
           null, null);
@@ -389,9 +389,9 @@ void testMD5HashMixed() {
          ColumnVector result = ColumnVector.md5Hash(strings, integers, doubles, floats, bools);
          ColumnVector expected = ColumnVector.fromStrings(
           "c12c8638819fdd8377bbf537a4ebf0b4", "abad86357c1ae28eeb89f4b59700946a",
-          "8c176e687e677a752868fe9d319d9d5c", "2f64d6a1d5b730fd97115924cf9aa486",
+          "7e376255c6354716cd63418208dc7b90", "2f64d6a1d5b730fd97115924cf9aa486",
           "9f9d26bb5d25d56453a91f0558370fa4", "d41d8cd98f00b204e9800998ecf8427e")) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -411,23 +411,22 @@ void testMD5HashLists() {
       assertColumnsAreEqual(expected, result);
     }
   }
-
   @Test
-  void test32BitMurmur3HashStrings() {
+  void testSerial32BitMurmur3HashStrings() {
     try (ColumnVector v0 = ColumnVector.fromStrings(
-           "a", "B\nc",  "dE\"\u0100\t\u0101 \u0500\u0501\\Fg2\'",
+           "a", "B\nc",  "dE\"\u0100\t\u0101 \ud720\ud721\\Fg2\'",
            "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
            "in the MD5 hash function. This string needed to be longer.A 60 character string to " +
            "test MD5's message padding algorithm",
            "hiJ\ud720\ud721\ud720\ud721", null);
          ColumnVector result = ColumnVector.serial32BitMurmurHash3(42, new ColumnVector[]{v0});
-         ColumnVector expected = ColumnVector.fromBoxedInts(-1293573533, 1163854319, 1247767417, 1504480835, 1249086584, 42)) {
+         ColumnVector expected = ColumnVector.fromBoxedInts(-1293573533, 1163854319, 1423943036, 1504480835, 1249086584, 42)) {
       assertColumnsAreEqual(expected, result);
     }
   }
 
   @Test
-  void test32BitMurmur3HashInts() {
+  void testSerial32BitMurmur3HashInts() {
     try (ColumnVector v0 = ColumnVector.fromBoxedInts(0, 100, null, null, Integer.MIN_VALUE, null);
          ColumnVector v1 = ColumnVector.fromBoxedInts(0, null, -100, null, null, Integer.MAX_VALUE);
          ColumnVector result = ColumnVector.serial32BitMurmurHash3(42, new ColumnVector[]{v0, v1});
@@ -437,7 +436,7 @@ void test32BitMurmur3HashInts() {
   }
 
   @Test
-  void test32BitMurmur3HashDoubles() {
+  void testSerial32BitMurmur3HashDoubles() {
     try (ColumnVector v = ColumnVector.fromBoxedDoubles(
           0.0, null, 100.0, -100.0, Double.MIN_NORMAL, Double.MAX_VALUE,
           POSITIVE_DOUBLE_NAN_UPPER_RANGE, POSITIVE_DOUBLE_NAN_LOWER_RANGE,
@@ -450,7 +449,7 @@ void test32BitMurmur3HashDoubles() {
   }
 
   @Test
-  void test32BitMurmur3HashFloats() {
+  void testSerialBitMurmur3HashFloats() {
     try (ColumnVector v = ColumnVector.fromBoxedFloats(
           0f, 100f, -100f, Float.MIN_NORMAL, Float.MAX_VALUE, null,
           POSITIVE_FLOAT_NAN_LOWER_RANGE, POSITIVE_FLOAT_NAN_UPPER_RANGE,
@@ -463,7 +462,7 @@ void test32BitMurmur3HashFloats() {
   }
 
   @Test
-  void test32BitMurmur3HashBools() {
+  void testSerial32BitMurmur3HashBools() {
     try (ColumnVector v0 = ColumnVector.fromBoxedBooleans(null, true, false, true, null, false);
          ColumnVector v1 = ColumnVector.fromBoxedBooleans(null, true, false, null, false, true);
          ColumnVector result = ColumnVector.serial32BitMurmurHash3(0, new ColumnVector[]{v0, v1});
@@ -473,9 +472,9 @@ void test32BitMurmur3HashBools() {
   }
 
   @Test
-  void test32BitMurmur3HashMixed() {
+  void testSerial32BitMurmur3HashMixed() {
     try (ColumnVector strings = ColumnVector.fromStrings(
-          "a", "B\n", "dE\"\u0100\t\u0101 \u0500\u0501",
+          "a", "B\n", "dE\"\u0100\t\u0101 \ud720\ud721",
           "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
           "in the MD5 hash function. This string needed to be longer.",
           null, null);
@@ -486,7 +485,87 @@ void test32BitMurmur3HashMixed() {
           0f, 100f, -100f, NEGATIVE_FLOAT_NAN_LOWER_RANGE, NEGATIVE_FLOAT_NAN_UPPER_RANGE, null);
          ColumnVector bools = ColumnVector.fromBoxedBooleans(true, false, null, false, true, null);
          ColumnVector result = ColumnVector.serial32BitMurmurHash3(1868, new ColumnVector[]{strings, integers, doubles, floats, bools});
-         ColumnVector expected = ColumnVector.fromBoxedInts(387200465, 1988790727, 1843539474, 814731646, -1073686048, 1868)) {
+         ColumnVector expected = ColumnVector.fromBoxedInts(387200465, 1988790727, 774895031, 814731646, -1073686048, 1868)) {
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
+  @Test
+  void testSpark32BitMurmur3HashStrings() {
+    try (ColumnVector v0 = ColumnVector.fromStrings(
+           "a", "B\nc",  "dE\"\u0100\t\u0101 \ud720\ud721\\Fg2\'",
+           "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
+           "in the MD5 hash function. This string needed to be longer.A 60 character string to " +
+           "test MD5's message padding algorithm",
+           "hiJ\ud720\ud721\ud720\ud721", null);
+         ColumnVector result = ColumnVector.spark32BitMurmurHash3(42, new ColumnVector[]{v0});
+         ColumnVector expected = ColumnVector.fromBoxedInts(1485273170, 1709559900, 1423943036, 176121990, 1199621434, 42)) {
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
+  @Test
+  void testSpark32BitMurmur3HashInts() {
+    try (ColumnVector v0 = ColumnVector.fromBoxedInts(0, 100, null, null, Integer.MIN_VALUE, null);
+         ColumnVector v1 = ColumnVector.fromBoxedInts(0, null, -100, null, null, Integer.MAX_VALUE);
+         ColumnVector result = ColumnVector.spark32BitMurmurHash3(42, new ColumnVector[]{v0, v1});
+         ColumnVector expected = ColumnVector.fromBoxedInts(59727262, 751823303, -1080202046, 42, 723455942, 133916647)) {
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
+  @Test
+  void testSpark32BitMurmur3HashDoubles() {
+    try (ColumnVector v = ColumnVector.fromBoxedDoubles(
+          0.0, null, 100.0, -100.0, Double.MIN_NORMAL, Double.MAX_VALUE,
+          POSITIVE_DOUBLE_NAN_UPPER_RANGE, POSITIVE_DOUBLE_NAN_LOWER_RANGE,
+          NEGATIVE_DOUBLE_NAN_UPPER_RANGE, NEGATIVE_DOUBLE_NAN_LOWER_RANGE,
+          Double.POSITIVE_INFINITY, Double.NEGATIVE_INFINITY);
+         ColumnVector result = ColumnVector.spark32BitMurmurHash3(new ColumnVector[]{v});
+         ColumnVector expected = ColumnVector.fromBoxedInts(1669671676, 0, -544903190, -1831674681, 150502665, 474144502, 1428788237, 1428788237, 1428788237, 1428788237, 420913893, 1915664072)) {
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
+  @Test
+  void testSpark32BitMurmur3HashFloats() {
+    try (ColumnVector v = ColumnVector.fromBoxedFloats(
+          0f, 100f, -100f, Float.MIN_NORMAL, Float.MAX_VALUE, null,
+          POSITIVE_FLOAT_NAN_LOWER_RANGE, POSITIVE_FLOAT_NAN_UPPER_RANGE,
+          NEGATIVE_FLOAT_NAN_LOWER_RANGE, NEGATIVE_FLOAT_NAN_UPPER_RANGE,
+          Float.POSITIVE_INFINITY, Float.NEGATIVE_INFINITY);
+         ColumnVector result = ColumnVector.spark32BitMurmurHash3(411, new ColumnVector[]{v});
+         ColumnVector expected = ColumnVector.fromBoxedInts(-235179434, 1812056886, 2028471189, 1775092689, -1531511762, 411, -1053523253, -1053523253, -1053523253, -1053523253, -1526256646, 930080402)){
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
+  @Test
+  void testSpark32BitMurmur3HashBools() {
+    try (ColumnVector v0 = ColumnVector.fromBoxedBooleans(null, true, false, true, null, false);
+         ColumnVector v1 = ColumnVector.fromBoxedBooleans(null, true, false, null, false, true);
+         ColumnVector result = ColumnVector.spark32BitMurmurHash3(0, new ColumnVector[]{v0, v1});
+         ColumnVector expected = ColumnVector.fromBoxedInts(0, -1589400010, -239939054, -68075478, 593689054, -1194558265)) {
+      assertColumnsAreEqual(expected, result);
+
+    }
+  }
+
+  @Test
+  void testSpark32BitMurmur3HashMixed() {
+    try (ColumnVector strings = ColumnVector.fromStrings(
+          "a", "B\n", "dE\"\u0100\t\u0101 \ud720\ud721",
+          "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
+          "in the MD5 hash function. This string needed to be longer.",
+          null, null);
+         ColumnVector integers = ColumnVector.fromBoxedInts(0, 100, -100, Integer.MIN_VALUE, Integer.MAX_VALUE, null);
+         ColumnVector doubles = ColumnVector.fromBoxedDoubles(
+          0.0, 100.0, -100.0, POSITIVE_DOUBLE_NAN_LOWER_RANGE, POSITIVE_DOUBLE_NAN_UPPER_RANGE, null);
+         ColumnVector floats = ColumnVector.fromBoxedFloats(
+          0f, 100f, -100f, NEGATIVE_FLOAT_NAN_LOWER_RANGE, NEGATIVE_FLOAT_NAN_UPPER_RANGE, null);
+         ColumnVector bools = ColumnVector.fromBoxedBooleans(true, false, null, false, true, null);
+         ColumnVector result = ColumnVector.spark32BitMurmurHash3(1868, new ColumnVector[]{strings, integers, doubles, floats, bools});
+         ColumnVector expected = ColumnVector.fromBoxedInts(1936985022, 720652989, 339312041, 1400354989, 769988643, 1868)) {
       assertColumnsAreEqual(expected, result);
     }
   }
@@ -511,7 +590,7 @@ void isNotNullTestEmptyColumn() {
     try (ColumnVector v = ColumnVector.fromBoxedInts();
          ColumnVector expected = ColumnVector.fromBoxedBooleans();
          ColumnVector result = v.isNotNull()) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -520,7 +599,7 @@ void isNotNullTest() {
     try (ColumnVector v = ColumnVector.fromBoxedInts(1, 2, null, 4, null, 6);
          ColumnVector expected = ColumnVector.fromBoxedBooleans(true, true, false, true, false, true);
          ColumnVector result = v.isNotNull()) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -529,7 +608,7 @@ void isNotNullTestAllNulls() {
     try (ColumnVector v = ColumnVector.fromBoxedInts(null, null, null, null, null, null);
          ColumnVector expected = ColumnVector.fromBoxedBooleans(false, false, false, false, false, false);
          ColumnVector result = v.isNotNull()) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -1232,7 +1311,7 @@ void testReplaceEmptyColumn() {
          ColumnVector expected = ColumnVector.fromBoxedBooleans();
          Scalar s = Scalar.fromBool(false);
          ColumnVector result = input.replaceNulls(s)) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -1242,7 +1321,7 @@ void testReplaceNullBoolsWithAllNulls() {
          ColumnVector expected = ColumnVector.fromBoxedBooleans(false, false, false, false);
          Scalar s = Scalar.fromBool(false);
          ColumnVector result = input.replaceNulls(s)) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -1252,7 +1331,7 @@ void testReplaceSomeNullBools() {
          ColumnVector expected = ColumnVector.fromBoxedBooleans(false, true, true, false);
          Scalar s = Scalar.fromBool(true);
          ColumnVector result = input.replaceNulls(s)) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -1262,7 +1341,7 @@ void testReplaceNullIntegersWithAllNulls() {
          ColumnVector expected = ColumnVector.fromBoxedInts(0, 0, 0, 0);
          Scalar s = Scalar.fromInt(0);
          ColumnVector result = input.replaceNulls(s)) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
@@ -1272,7 +1351,7 @@ void testReplaceSomeNullIntegers() {
          ColumnVector expected = ColumnVector.fromBoxedInts(1, 2, 999, 4, 999);
          Scalar s = Scalar.fromInt(999);
          ColumnVector result = input.replaceNulls(s)) {
-      assertColumnsAreEqual(result, expected);
+      assertColumnsAreEqual(expected, result);
     }
   }
 
diff --git a/python/cudf/cudf/_lib/cpp/types.pxd b/python/cudf/cudf/_lib/cpp/types.pxd
index 86f4e03741d..cf86076f8d6 100644
--- a/python/cudf/cudf/_lib/cpp/types.pxd
+++ b/python/cudf/cudf/_lib/cpp/types.pxd
@@ -79,6 +79,7 @@ cdef extern from "cudf/types.hpp" namespace "cudf" nogil:
         HASH_MURMUR3 "cudf::hash_id::HASH_MURMUR3"
         HASH_MD5 "cudf::hash_id::HASH_MD5"
         HASH_SERIAL_MURMUR3 "cudf::hash_id::HASH_SERIAL_MURMUR3"
+        HASH_SPARK_MURMUR3 "cudf::hash_id::HASH_SPARK_MURMUR3"
 
     cdef cppclass data_type:
         data_type() except +