From ebf079a7f637f93f6ca6cdde3588cc82d6e57254 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 28 Jun 2022 14:01:12 -0700
Subject: [PATCH 01/40] Add test for Spark list-of-list hashing.

---
 cpp/tests/hashing/hash_test.cpp | 97 +++++++++++++++++++++++++++++++++
 1 file changed, 97 insertions(+)

diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index 1f3e6567253..43bd70a24f2 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -736,6 +736,103 @@ TEST_F(SparkMurmurHash3Test, StringsWithSeed)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(*hash_strings, hash_strings_expected_seed_314, verbosity);
 }
 
+TEST_F(SparkMurmurHash3Test, 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)
+  */
+
+  /*
+  child data: 1, 1, 2, 1, 2, 3, 1, 2, 3, 1, 2, 3, 1, null, 2, 3, 1, 2, 3, null, 1, 2, 3
+  23 items
+  validity: i != 13, i != 19
+
+  parent validity: i != 0 && i != 15
+  offsets: 0, 0, 1, 3, 6, 8, 9, 10, 12, 13, 16, 18, 19, 20, 22, 23
+  16 items
+
+  row validity: i != 0
+  row offsets: 0, 0, 0, 1, 2, 3, 4, 6, 8, 10, 13, 16
+  11 items
+  */
+
+  auto const null = -1;
+  auto child_validity =
+    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; });
+  auto parent_validity =
+    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0 && i != 15; });
+  auto row_validity =
+    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; });
+  auto list1 = cudf::test::lists_column_wrapper<int>({{},
+                                                      {1},
+                                                      {1, 2},
+                                                      {1, 2, 3},
+                                                      {1, 2},
+                                                      {3},
+                                                      {1},
+                                                      {2, 3},
+                                                      {1},
+                                                      {{null, 2, 3}, inner_validity},
+                                                      {1, 2},
+                                                      {3},
+                                                      {{null}, inner_validity},
+                                                      {1, 2},
+                                                      {},
+                                                      {3}},
+                                                     parent_validity);
+  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 = std::vector<bool>{0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1};
+  auto list_validity_buffer =
+    cudf::test::detail::make_null_mask(list_validity.begin(), list_validity.end());
+  auto list_column = cudf::make_lists_column(11,
+                                             offsets.release(),
+                                             list1.release(),
+                                             cudf::UNKNOWN_NULL_COUNT,
+                                             std::move(list_validity_buffer));
+
+  auto expect = cudf::test::fixed_width_column_wrapper<uint32_t>{42,
+                                                                 42,
+                                                                 42,
+                                                                 -559580957,
+                                                                 -222940379,
+                                                                 -912918097,
+                                                                 -912918097,
+                                                                 -912918097,
+                                                                 -912918097,
+                                                                 -912918097,
+                                                                 -912918097};
+
+  auto output = cudf::hash(cudf::table_view({*list_column}));
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
+}
+
 TEST_F(SparkMurmurHash3Test, ListThrows)
 {
   lists_column_wrapper<cudf::string_view> strings_list_col({{""}, {"abc"}, {"123"}});

From 5632e7cbb1d190e0112df394e07b105582a20707 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Thu, 7 Jul 2022 08:53:40 -0700
Subject: [PATCH 02/40] Improve test.

---
 cpp/tests/hashing/hash_test.cpp | 52 ++++++++++++++++-----------------
 1 file changed, 25 insertions(+), 27 deletions(-)

diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index 43bd70a24f2..1c8d7ae2df5 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -773,7 +773,7 @@ TEST_F(SparkMurmurHash3Test, ListValues)
   23 items
   validity: i != 13, i != 19
 
-  parent validity: i != 0 && i != 15
+  nested validity: i != 0 && i != 15
   offsets: 0, 0, 1, 3, 6, 8, 9, 10, 12, 13, 16, 18, 19, 20, 22, 23
   16 items
 
@@ -785,35 +785,33 @@ TEST_F(SparkMurmurHash3Test, ListValues)
   auto const null = -1;
   auto child_validity =
     cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; });
-  auto parent_validity =
+  auto nested_validity =
     cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0 && i != 15; });
-  auto row_validity =
-    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; });
-  auto list1 = cudf::test::lists_column_wrapper<int>({{},
-                                                      {1},
-                                                      {1, 2},
-                                                      {1, 2, 3},
-                                                      {1, 2},
-                                                      {3},
-                                                      {1},
-                                                      {2, 3},
-                                                      {1},
-                                                      {{null, 2, 3}, inner_validity},
-                                                      {1, 2},
-                                                      {3},
-                                                      {{null}, inner_validity},
-                                                      {1, 2},
-                                                      {},
-                                                      {3}},
-                                                     parent_validity);
+  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}, child_validity},
+                                                            {1, 2},
+                                                            {3},
+                                                            {{null}, child_validity},
+                                                            {1, 2},
+                                                            {},
+                                                            {3}},
+                                                           nested_validity);
   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 = std::vector<bool>{0, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1};
-  auto list_validity_buffer =
-    cudf::test::detail::make_null_mask(list_validity.begin(), list_validity.end());
-  auto list_column = cudf::make_lists_column(11,
+  auto list_validity =
+    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; });
+  auto list_validity_buffer = cudf::test::detail::make_null_mask(list_validity, list_validity + 11);
+  auto list_column          = cudf::make_lists_column(11,
                                              offsets.release(),
-                                             list1.release(),
+                                             nested_list.release(),
                                              cudf::UNKNOWN_NULL_COUNT,
                                              std::move(list_validity_buffer));
 
@@ -829,7 +827,7 @@ TEST_F(SparkMurmurHash3Test, ListValues)
                                                                  -912918097,
                                                                  -912918097};
 
-  auto output = cudf::hash(cudf::table_view({*list_column}));
+  auto output = cudf::hash(cudf::table_view({*list_column}), cudf::hash_id::HASH_SPARK_MURMUR3);
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
 }
 

From 6e55198ed18a70d14f478855512af156e09d0c46 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Mon, 18 Jul 2022 22:04:43 -0700
Subject: [PATCH 03/40] Copy experimental row hasher for modification.

---
 cpp/CMakeLists.txt                  |   1 +
 cpp/include/cudf/detail/hashing.hpp |   6 +
 cpp/src/hash/hashing.cu             |   3 +-
 cpp/src/hash/spark_murmur_hash.cu   | 287 ++++++++++++++++++++++++++++
 4 files changed, 295 insertions(+), 2 deletions(-)
 create mode 100644 cpp/src/hash/spark_murmur_hash.cu

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 86bfdc1444b..55ef34d07fc 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -299,6 +299,7 @@ add_library(
   src/hash/hashing.cu
   src/hash/md5_hash.cu
   src/hash/murmur_hash.cu
+  src/hash/spark_murmur_hash.cu
   src/interop/dlpack.cpp
   src/interop/from_arrow.cu
   src/interop/to_arrow.cu
diff --git a/cpp/include/cudf/detail/hashing.hpp b/cpp/include/cudf/detail/hashing.hpp
index 29522764dad..8b0ea42658f 100644
--- a/cpp/include/cudf/detail/hashing.hpp
+++ b/cpp/include/cudf/detail/hashing.hpp
@@ -44,6 +44,12 @@ std::unique_ptr<column> murmur_hash3_32(
   rmm::cuda_stream_view stream        = cudf::default_stream_value,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
+std::unique_ptr<column> spark_murmur_hash3_32(
+  table_view const& input,
+  uint32_t seed                       = cudf::DEFAULT_HASH_SEED,
+  rmm::cuda_stream_view stream        = cudf::default_stream_value,
+  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,
diff --git a/cpp/src/hash/hashing.cu b/cpp/src/hash/hashing.cu
index a3bfa99b000..af26042426f 100644
--- a/cpp/src/hash/hashing.cu
+++ b/cpp/src/hash/hashing.cu
@@ -99,8 +99,7 @@ std::unique_ptr<column> hash(table_view const& input,
     case (hash_id::HASH_MURMUR3): return 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);
+    case (hash_id::HASH_SPARK_MURMUR3): return spark_murmur_hash3_32(input, seed, stream, mr);
     case (hash_id::HASH_MD5): return md5_hash(input, stream, mr);
     default: CUDF_FAIL("Unsupported hash function.");
   }
diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
new file mode 100644
index 00000000000..b45775f9cc8
--- /dev/null
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -0,0 +1,287 @@
+/*
+ * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#include <cudf/column/column_factories.hpp>
+#include <cudf/detail/hashing.hpp>
+#include <cudf/detail/utilities/hash_functions.cuh>
+#include <cudf/detail/utilities/vector_factories.hpp>
+#include <cudf/table/experimental/row_operators.cuh>
+#include <cudf/table/table_device_view.cuh>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/tabulate.h>
+
+namespace cudf {
+namespace detail {
+
+namespace {
+
+/**
+ * @brief Computes the hash value of an element in the given column.
+ *
+ * @tparam hash_function Hash functor to use for hashing elements.
+ * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+ */
+template <template <typename> class hash_function, typename Nullate>
+class element_hasher {
+ public:
+  /**
+   * @brief Constructs an element_hasher object.
+   *
+   * @param nulls Indicates whether to check for nulls
+   * @param seed  The seed to use for the hash function
+   * @param null_hash The hash value to use for nulls
+   */
+  __device__ element_hasher(
+    Nullate nulls,
+    uint32_t seed             = DEFAULT_HASH_SEED,
+    hash_value_type null_hash = std::numeric_limits<hash_value_type>::max()) noexcept
+    : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
+  {
+  }
+
+  /**
+   * @brief Returns the hash value of the given element.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
+  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
+  __device__ hash_value_type operator()(column_device_view const& col,
+                                        size_type row_index) const noexcept
+  {
+    if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
+    return hash_function<T>{_seed}(col.element<T>(row_index));
+  }
+
+  /**
+   * @brief Returns the hash value of the given element.
+   *
+   * @tparam T The type of the element to hash
+   * @param col The column to hash
+   * @param row_index The index of the row to hash
+   * @return The hash value of the given element
+   */
+  template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
+  __device__ hash_value_type operator()(column_device_view const&, size_type) const noexcept
+  {
+    CUDF_UNREACHABLE("Unsupported type in hash.");
+  }
+
+  uint32_t _seed;              ///< The seed to use for hashing
+  hash_value_type _null_hash;  ///< Hash value to use for null elements
+  Nullate _check_nulls;        ///< Whether to check for nulls
+};
+
+/**
+ * @brief Computes the hash value of a row in the given table.
+ *
+ * @tparam hash_function Hash functor to use for hashing elements.
+ * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+ */
+template <template <typename> class hash_function, typename Nullate>
+class device_spark_row_hasher {
+  friend class spark_row_hasher;  ///< Allow spark_row_hasher to access private members.
+
+ public:
+  device_spark_row_hasher() = delete;
+
+  /**
+   * @brief Return the hash value of a row in the given table.
+   *
+   * @param row_index The row index to compute the hash value of
+   * @return The hash value of the row
+   */
+  __device__ auto operator()(size_type row_index) const noexcept
+  {
+    auto it = thrust::make_transform_iterator(_table.begin(), [=](auto const& column) {
+      return cudf::type_dispatcher<dispatch_storage_type>(
+        column.type(), element_hasher_adapter<hash_function>{_check_nulls}, column, row_index);
+    });
+
+    // Hash each element and combine all the hash values together
+    return detail::accumulate(it, it + _table.num_columns(), _seed, [](auto hash, auto h) {
+      return cudf::detail::hash_combine(hash, h);
+    });
+  }
+
+ private:
+  /**
+   * @brief Computes the hash value of an element in the given column.
+   *
+   * When the column is non-nested, this is a simple wrapper around the element_hasher.
+   * When the column is nested, this uses the element_hasher to hash the shape and values of the
+   * column.
+   */
+  template <template <typename> class hash_fn>
+  class element_hasher_adapter {
+    static constexpr hash_value_type NULL_HASH     = std::numeric_limits<hash_value_type>::max();
+    static constexpr hash_value_type NON_NULL_HASH = 0;
+
+   public:
+    __device__ element_hasher_adapter(Nullate check_nulls) noexcept
+      : _element_hasher(check_nulls), _check_nulls(check_nulls)
+    {
+    }
+
+    template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
+    __device__ hash_value_type operator()(column_device_view const& col,
+                                          size_type row_index) const noexcept
+    {
+      return _element_hasher.template operator()<T>(col, row_index);
+    }
+
+    template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
+    __device__ hash_value_type operator()(column_device_view const& col,
+                                          size_type row_index) const noexcept
+    {
+      auto hash                   = hash_value_type{0};
+      column_device_view curr_col = col.slice(row_index, 1);
+      while (is_nested(curr_col.type())) {
+        if (_check_nulls) {
+          auto validity_it = detail::make_validity_iterator<true>(curr_col);
+          hash             = detail::accumulate(
+            validity_it, validity_it + curr_col.size(), hash, [](auto hash, auto is_valid) {
+              return cudf::detail::hash_combine(hash, is_valid ? NON_NULL_HASH : NULL_HASH);
+            });
+        }
+        if (curr_col.type().id() == type_id::STRUCT) {
+          if (curr_col.num_child_columns() == 0) { return hash; }
+          // Non-empty structs are assumed to be decomposed and contain only one child
+          curr_col = detail::structs_column_device_view(curr_col).get_sliced_child(0);
+        } else if (curr_col.type().id() == type_id::LIST) {
+          auto list_col   = detail::lists_column_device_view(curr_col);
+          auto list_sizes = make_list_size_iterator(list_col);
+          hash            = detail::accumulate(
+            list_sizes, list_sizes + list_col.size(), hash, [](auto hash, auto size) {
+              return cudf::detail::hash_combine(hash, hash_fn<size_type>{}(size));
+            });
+          curr_col = list_col.get_sliced_child();
+        }
+      }
+      for (int i = 0; i < curr_col.size(); ++i) {
+        hash =
+          cudf::detail::hash_combine(hash,
+                                     type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
+                                       curr_col.type(), _element_hasher, curr_col, i));
+      }
+      return hash;
+    }
+
+    element_hasher<hash_fn, Nullate> const _element_hasher;
+    Nullate const _check_nulls;
+  };
+
+  CUDF_HOST_DEVICE device_spark_row_hasher(Nullate check_nulls,
+                                           table_device_view t,
+                                           uint32_t seed = DEFAULT_HASH_SEED) noexcept
+    : _table{t}, _seed(seed), _check_nulls{check_nulls}
+  {
+  }
+
+  table_device_view const _table;
+  Nullate const _check_nulls;
+  uint32_t const _seed;
+};
+
+using preprocessed_table = cudf::experimental::row::hash::preprocessed_table;
+
+/**
+ * @brief Computes the hash value of a row in the given table.
+ *
+ */
+class spark_row_hasher {
+ public:
+  /**
+   * @brief Construct an owning object for hashing the rows of a table
+   *
+   * @param t The table containing rows to hash
+   * @param stream The stream to construct this object on. Not the stream that will be used for
+   * comparisons using this object.
+   */
+  spark_row_hasher(table_view const& t, rmm::cuda_stream_view stream)
+    : d_t(preprocessed_table::create(t, stream))
+  {
+  }
+
+  /**
+   * @brief Construct an owning object for hashing the rows of a table from an existing
+   * preprocessed_table
+   *
+   * This constructor allows independently constructing a `preprocessed_table` and sharing it among
+   * multiple `spark_row_hasher` and `equality::self_comparator` objects.
+   *
+   * @param t A table preprocessed for hashing or equality.
+   */
+  spark_row_hasher(std::shared_ptr<preprocessed_table> t) : d_t{std::move(t)} {}
+
+  /**
+   * @brief Get the hash operator to use on the device
+   *
+   * Returns a unary callable, `F`, with signature `hash_function::hash_value_type F(size_type)`.
+   *
+   * `F(i)` returns the hash of row i.
+   *
+   * @tparam Nullate A cudf::nullate type describing whether to check for nulls
+   * @param nullate Indicates if any input column contains nulls
+   * @param seed The seed to use for the hash function
+   * @return A hash operator to use on the device
+   */
+  template <template <typename> class hash_function = detail::default_hash, typename Nullate>
+  device_spark_row_hasher<hash_function, Nullate> device_hasher(
+    Nullate nullate = {}, uint32_t seed = DEFAULT_HASH_SEED) const
+  {
+    return device_spark_row_hasher<hash_function, Nullate>(nullate, *d_t, seed);
+  }
+
+ private:
+  std::shared_ptr<preprocessed_table> d_t;
+};
+
+}  // namespace
+
+std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
+                                              uint32_t seed,
+                                              rmm::cuda_stream_view stream,
+                                              rmm::mr::device_memory_resource* mr)
+{
+  auto output = make_numeric_column(data_type(type_to_id<hash_value_type>()),
+                                    input.num_rows(),
+                                    mask_state::UNALLOCATED,
+                                    stream,
+                                    mr);
+
+  // Return early if there's nothing to hash
+  if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
+
+  bool const nullable   = has_nulls(input);
+  auto const row_hasher = spark_row_hasher(input, stream);
+  auto output_view      = output->mutable_view();
+
+  // Compute the hash value for each row
+  thrust::tabulate(rmm::exec_policy(stream),
+                   output_view.begin<hash_value_type>(),
+                   output_view.end<hash_value_type>(),
+                   row_hasher.device_hasher<SparkMurmurHash3_32>(nullable, seed));
+
+  return output;
+}
+
+}  // namespace detail
+}  // namespace cudf

From 95f25cb0e2030576d3079ffeb44e45e6cdb860e2 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 19 Jul 2022 13:04:55 -0700
Subject: [PATCH 04/40] Make preprocessed table methods public.

---
 .../cudf/table/experimental/row_operators.cuh       | 13 ++++++++-----
 1 file changed, 8 insertions(+), 5 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index e3acde6b90b..fc30462c00e 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -1132,14 +1132,16 @@ struct preprocessed_table {
   static std::shared_ptr<preprocessed_table> create(table_view const& table,
                                                     rmm::cuda_stream_view stream);
 
- private:
-  friend class self_comparator;       ///< Allow self_comparator to access private members
-  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
-  friend class hash::row_hasher;      ///< Allow row_hasher to access private members
-
   using table_device_view_owner =
     std::invoke_result_t<decltype(table_device_view::create), table_view, rmm::cuda_stream_view>;
+  ///< Type returned by table creation.
 
+  /**
+   * @brief Construct a preprocessed table.
+   *
+   * @param table The owning table device view
+   * @param null_buffers Null masks superimposed from parent columns
+   */
   preprocessed_table(table_device_view_owner&& table,
                      std::vector<rmm::device_buffer>&& null_buffers)
     : _t(std::move(table)), _null_buffers(std::move(null_buffers))
@@ -1153,6 +1155,7 @@ struct preprocessed_table {
    */
   operator table_device_view() { return *_t; }
 
+ private:
   table_device_view_owner _t;
   std::vector<rmm::device_buffer> _null_buffers;
 };

From ac21705d303b7094dffa39765863f9e48784ef6a Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 19 Jul 2022 13:06:10 -0700
Subject: [PATCH 05/40] Use structured binding to make names map more clearly
 to the objects.

---
 cpp/src/table/row_operators.cu | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/cpp/src/table/row_operators.cu b/cpp/src/table/row_operators.cu
index f6ca9a04e6d..0afc36b0379 100644
--- a/cpp/src/table/row_operators.cu
+++ b/cpp/src/table/row_operators.cu
@@ -358,13 +358,13 @@ std::shared_ptr<preprocessed_table> preprocessed_table::create(table_view const&
 {
   check_eq_compatibility(t);
 
-  auto null_pushed_table              = structs::detail::superimpose_parent_nulls(t, stream);
-  auto struct_offset_removed_table    = remove_struct_child_offsets(std::get<0>(null_pushed_table));
-  auto [verticalized_lhs, _, __, ___] = decompose_structs(struct_offset_removed_table);
+  auto [null_pushed_table, null_masks] = structs::detail::superimpose_parent_nulls(t, stream);
+  auto struct_offset_removed_table     = remove_struct_child_offsets(null_pushed_table);
+  auto [verticalized_lhs, _, __, ___]  = decompose_structs(struct_offset_removed_table);
 
   auto d_t = table_device_view_owner(table_device_view::create(verticalized_lhs, stream));
   return std::shared_ptr<preprocessed_table>(
-    new preprocessed_table(std::move(d_t), std::move(std::get<1>(null_pushed_table))));
+    new preprocessed_table(std::move(d_t), std::move(null_masks)));
 }
 
 two_table_comparator::two_table_comparator(table_view const& left,

From 4106aef5190f48e71a8662e0df64e8f2d8ce501d Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 19 Jul 2022 13:25:35 -0700
Subject: [PATCH 06/40] Add note about hash types, fix tests.

---
 cpp/src/hash/spark_murmur_hash.cu |  5 +++++
 cpp/tests/hashing/hash_test.cpp   | 30 +++++++++++-------------------
 2 files changed, 16 insertions(+), 19 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index b45775f9cc8..f0524282473 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -261,6 +261,11 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
                                               rmm::cuda_stream_view stream,
                                               rmm::mr::device_memory_resource* mr)
 {
+  // TODO: Spark uses int32_t hash values, but libcudf defines hash_value_type
+  // as uint32_t elsewhere. This should be investigated and unified. I suspect
+  // we should use int32_t everywhere. --bdice
+  using hash_value_type = int32_t;
+
   auto output = make_numeric_column(data_type(type_to_id<hash_value_type>()),
                                     input.num_rows(),
                                     mask_state::UNALLOCATED,
diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index 1c8d7ae2df5..b78c2dfd1b2 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -815,30 +815,22 @@ TEST_F(SparkMurmurHash3Test, ListValues)
                                              cudf::UNKNOWN_NULL_COUNT,
                                              std::move(list_validity_buffer));
 
-  auto expect = cudf::test::fixed_width_column_wrapper<uint32_t>{42,
-                                                                 42,
-                                                                 42,
-                                                                 -559580957,
-                                                                 -222940379,
-                                                                 -912918097,
-                                                                 -912918097,
-                                                                 -912918097,
-                                                                 -912918097,
-                                                                 -912918097,
-                                                                 -912918097};
+  auto expect = cudf::test::fixed_width_column_wrapper<int32_t>{42,
+                                                                42,
+                                                                42,
+                                                                -559580957,
+                                                                -222940379,
+                                                                -912918097,
+                                                                -912918097,
+                                                                -912918097,
+                                                                -912918097,
+                                                                -912918097,
+                                                                -912918097};
 
   auto output = cudf::hash(cudf::table_view({*list_column}), cudf::hash_id::HASH_SPARK_MURMUR3);
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
 }
 
-TEST_F(SparkMurmurHash3Test, ListThrows)
-{
-  lists_column_wrapper<cudf::string_view> strings_list_col({{""}, {"abc"}, {"123"}});
-  EXPECT_THROW(
-    cudf::hash(cudf::table_view({strings_list_col}), cudf::hash_id::HASH_SPARK_MURMUR3, {}),
-    cudf::logic_error);
-}
-
 class MD5HashTest : public cudf::test::BaseFixture {
 };
 

From ca8557d23d361b6415997145a465272c7a29876a Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 19 Jul 2022 14:01:37 -0700
Subject: [PATCH 07/40] Align hashing behavior closer to Spark. Some tests
 passing, but not all.

---
 cpp/src/hash/spark_murmur_hash.cu | 150 +++++++++++++++++-------------
 1 file changed, 85 insertions(+), 65 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index f0524282473..53a3b1ae025 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -30,64 +30,63 @@ namespace detail {
 
 namespace {
 
-/**
- * @brief Computes the hash value of an element in the given column.
- *
- * @tparam hash_function Hash functor to use for hashing elements.
- * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
- */
-template <template <typename> class hash_function, typename Nullate>
-class element_hasher {
- public:
-  /**
-   * @brief Constructs an element_hasher object.
-   *
-   * @param nulls Indicates whether to check for nulls
-   * @param seed  The seed to use for the hash function
-   * @param null_hash The hash value to use for nulls
-   */
-  __device__ element_hasher(
-    Nullate nulls,
-    uint32_t seed             = DEFAULT_HASH_SEED,
-    hash_value_type null_hash = std::numeric_limits<hash_value_type>::max()) noexcept
-    : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
-  {
-  }
-
-  /**
-   * @brief Returns the hash value of the given element.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(column_device_view const& col,
-                                        size_type row_index) const noexcept
-  {
-    if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
-    return hash_function<T>{_seed}(col.element<T>(row_index));
-  }
-
-  /**
-   * @brief Returns the hash value of the given element.
-   *
-   * @tparam T The type of the element to hash
-   * @param col The column to hash
-   * @param row_index The index of the row to hash
-   * @return The hash value of the given element
-   */
-  template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
-  __device__ hash_value_type operator()(column_device_view const&, size_type) const noexcept
-  {
-    CUDF_UNREACHABLE("Unsupported type in hash.");
-  }
-
-  uint32_t _seed;              ///< The seed to use for hashing
-  hash_value_type _null_hash;  ///< Hash value to use for null elements
-  Nullate _check_nulls;        ///< Whether to check for nulls
-};
+// /**
+//  * @brief Computes the hash value of an element in the given column.
+//  *
+//  * @tparam hash_function Hash functor to use for hashing elements.
+//  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
+//  */
+// class element_hasher {
+//  public:
+// /**
+//  * @brief Constructs an element_hasher object.
+//  *
+//  * @param nulls Indicates whether to check for nulls
+//  * @param seed  The seed to use for the hash function
+//  * @param null_hash The hash value to use for nulls
+//  */
+// __device__ element_hasher(
+// Nullate nulls,
+// uint32_t seed             = DEFAULT_HASH_SEED,
+// hash_value_type null_hash = std::numeric_limits<hash_value_type>::max()) noexcept
+// : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
+// {
+// }
+//
+// /**
+//  * @brief Returns the hash value of the given element.
+//  *
+//  * @tparam T The type of the element to hash
+//  * @param col The column to hash
+//  * @param row_index The index of the row to hash
+//  * @return The hash value of the given element
+//  */
+// template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
+// __device__ hash_value_type operator()(column_device_view const& col,
+// size_type row_index) const noexcept
+// {
+// if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
+// return hash_function<T>{_seed}(col.element<T>(row_index));
+// }
+//
+// /**
+//  * @brief Returns the hash value of the given element.
+//  *
+//  * @tparam T The type of the element to hash
+//  * @param col The column to hash
+//  * @param row_index The index of the row to hash
+//  * @return The hash value of the given element
+//  */
+// template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
+// __device__ hash_value_type operator()(column_device_view const&, size_type) const noexcept
+// {
+// CUDF_UNREACHABLE("Unsupported type in hash.");
+// }
+//
+// uint32_t _seed;              ///< The seed to use for hashing
+// hash_value_type _null_hash;  ///< Hash value to use for null elements
+// Nullate _check_nulls;        ///< Whether to check for nulls
+// };
 
 /**
  * @brief Computes the hash value of a row in the given table.
@@ -110,6 +109,7 @@ class device_spark_row_hasher {
    */
   __device__ auto operator()(size_type row_index) const noexcept
   {
+    /*
     auto it = thrust::make_transform_iterator(_table.begin(), [=](auto const& column) {
       return cudf::type_dispatcher<dispatch_storage_type>(
         column.type(), element_hasher_adapter<hash_function>{_check_nulls}, column, row_index);
@@ -119,6 +119,19 @@ class device_spark_row_hasher {
     return detail::accumulate(it, it + _table.num_columns(), _seed, [](auto hash, auto h) {
       return cudf::detail::hash_combine(hash, h);
     });
+    */
+
+    return detail::accumulate(
+      _table.begin(),
+      _table.end(),
+      _seed,
+      [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
+        return cudf::type_dispatcher<dispatch_storage_type>(
+          column.type(),
+          element_hasher_adapter<hash_function>{nulls, hash, hash},
+          column,
+          row_index);
+      });
   }
 
  private:
@@ -131,12 +144,15 @@ class device_spark_row_hasher {
    */
   template <template <typename> class hash_fn>
   class element_hasher_adapter {
-    static constexpr hash_value_type NULL_HASH     = std::numeric_limits<hash_value_type>::max();
-    static constexpr hash_value_type NON_NULL_HASH = 0;
+    // static constexpr hash_value_type NULL_HASH     = std::numeric_limits<hash_value_type>::max();
+    // static constexpr hash_value_type NON_NULL_HASH = 0;
 
    public:
-    __device__ element_hasher_adapter(Nullate check_nulls) noexcept
-      : _element_hasher(check_nulls), _check_nulls(check_nulls)
+    __device__ element_hasher_adapter(
+      Nullate check_nulls,
+      uint32_t seed             = DEFAULT_HASH_SEED,
+      hash_value_type null_hash = std::numeric_limits<hash_value_type>::max()) noexcept
+      : _element_hasher(check_nulls, seed, null_hash), _check_nulls(check_nulls)
     {
     }
 
@@ -154,6 +170,7 @@ class device_spark_row_hasher {
       auto hash                   = hash_value_type{0};
       column_device_view curr_col = col.slice(row_index, 1);
       while (is_nested(curr_col.type())) {
+        /*
         if (_check_nulls) {
           auto validity_it = detail::make_validity_iterator<true>(curr_col);
           hash             = detail::accumulate(
@@ -161,17 +178,20 @@ class device_spark_row_hasher {
               return cudf::detail::hash_combine(hash, is_valid ? NON_NULL_HASH : NULL_HASH);
             });
         }
+        */
         if (curr_col.type().id() == type_id::STRUCT) {
           if (curr_col.num_child_columns() == 0) { return hash; }
           // Non-empty structs are assumed to be decomposed and contain only one child
           curr_col = detail::structs_column_device_view(curr_col).get_sliced_child(0);
         } else if (curr_col.type().id() == type_id::LIST) {
-          auto list_col   = detail::lists_column_device_view(curr_col);
+          auto list_col = detail::lists_column_device_view(curr_col);
+          /*
           auto list_sizes = make_list_size_iterator(list_col);
           hash            = detail::accumulate(
             list_sizes, list_sizes + list_col.size(), hash, [](auto hash, auto size) {
               return cudf::detail::hash_combine(hash, hash_fn<size_type>{}(size));
             });
+          */
           curr_col = list_col.get_sliced_child();
         }
       }
@@ -184,7 +204,7 @@ class device_spark_row_hasher {
       return hash;
     }
 
-    element_hasher<hash_fn, Nullate> const _element_hasher;
+    cudf::experimental::row::hash::element_hasher<hash_fn, Nullate> const _element_hasher;
     Nullate const _check_nulls;
   };
 
@@ -263,7 +283,7 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
 {
   // TODO: Spark uses int32_t hash values, but libcudf defines hash_value_type
   // as uint32_t elsewhere. This should be investigated and unified. I suspect
-  // we should use int32_t everywhere. --bdice
+  // we should use int32_t everywhere. Also check this for hash seeds. --bdice
   using hash_value_type = int32_t;
 
   auto output = make_numeric_column(data_type(type_to_id<hash_value_type>()),

From 703956f76f754056bfb666ab19d68ca4db517821 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 19 Jul 2022 14:03:45 -0700
Subject: [PATCH 08/40] Remove commented code, fix dispatch for decimal32 and
 decimal128.

---
 cpp/src/hash/spark_murmur_hash.cu | 100 ++----------------------------
 1 file changed, 5 insertions(+), 95 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 53a3b1ae025..16c0bbf6507 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -30,64 +30,6 @@ namespace detail {
 
 namespace {
 
-// /**
-//  * @brief Computes the hash value of an element in the given column.
-//  *
-//  * @tparam hash_function Hash functor to use for hashing elements.
-//  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
-//  */
-// class element_hasher {
-//  public:
-// /**
-//  * @brief Constructs an element_hasher object.
-//  *
-//  * @param nulls Indicates whether to check for nulls
-//  * @param seed  The seed to use for the hash function
-//  * @param null_hash The hash value to use for nulls
-//  */
-// __device__ element_hasher(
-// Nullate nulls,
-// uint32_t seed             = DEFAULT_HASH_SEED,
-// hash_value_type null_hash = std::numeric_limits<hash_value_type>::max()) noexcept
-// : _check_nulls(nulls), _seed(seed), _null_hash(null_hash)
-// {
-// }
-//
-// /**
-//  * @brief Returns the hash value of the given element.
-//  *
-//  * @tparam T The type of the element to hash
-//  * @param col The column to hash
-//  * @param row_index The index of the row to hash
-//  * @return The hash value of the given element
-//  */
-// template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-// __device__ hash_value_type operator()(column_device_view const& col,
-// size_type row_index) const noexcept
-// {
-// if (_check_nulls && col.is_null(row_index)) { return _null_hash; }
-// return hash_function<T>{_seed}(col.element<T>(row_index));
-// }
-//
-// /**
-//  * @brief Returns the hash value of the given element.
-//  *
-//  * @tparam T The type of the element to hash
-//  * @param col The column to hash
-//  * @param row_index The index of the row to hash
-//  * @return The hash value of the given element
-//  */
-// template <typename T, CUDF_ENABLE_IF(not column_device_view::has_element_accessor<T>())>
-// __device__ hash_value_type operator()(column_device_view const&, size_type) const noexcept
-// {
-// CUDF_UNREACHABLE("Unsupported type in hash.");
-// }
-//
-// uint32_t _seed;              ///< The seed to use for hashing
-// hash_value_type _null_hash;  ///< Hash value to use for null elements
-// Nullate _check_nulls;        ///< Whether to check for nulls
-// };
-
 /**
  * @brief Computes the hash value of a row in the given table.
  *
@@ -109,28 +51,15 @@ class device_spark_row_hasher {
    */
   __device__ auto operator()(size_type row_index) const noexcept
   {
-    /*
-    auto it = thrust::make_transform_iterator(_table.begin(), [=](auto const& column) {
-      return cudf::type_dispatcher<dispatch_storage_type>(
-        column.type(), element_hasher_adapter<hash_function>{_check_nulls}, column, row_index);
-    });
-
-    // Hash each element and combine all the hash values together
-    return detail::accumulate(it, it + _table.num_columns(), _seed, [](auto hash, auto h) {
-      return cudf::detail::hash_combine(hash, h);
-    });
-    */
-
     return detail::accumulate(
       _table.begin(),
       _table.end(),
       _seed,
       [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
-        return cudf::type_dispatcher<dispatch_storage_type>(
-          column.type(),
-          element_hasher_adapter<hash_function>{nulls, hash, hash},
-          column,
-          row_index);
+        return cudf::type_dispatcher(column.type(),
+                                     element_hasher_adapter<hash_function>{nulls, hash, hash},
+                                     column,
+                                     row_index);
       });
   }
 
@@ -144,9 +73,6 @@ class device_spark_row_hasher {
    */
   template <template <typename> class hash_fn>
   class element_hasher_adapter {
-    // static constexpr hash_value_type NULL_HASH     = std::numeric_limits<hash_value_type>::max();
-    // static constexpr hash_value_type NON_NULL_HASH = 0;
-
    public:
     __device__ element_hasher_adapter(
       Nullate check_nulls,
@@ -170,29 +96,13 @@ class device_spark_row_hasher {
       auto hash                   = hash_value_type{0};
       column_device_view curr_col = col.slice(row_index, 1);
       while (is_nested(curr_col.type())) {
-        /*
-        if (_check_nulls) {
-          auto validity_it = detail::make_validity_iterator<true>(curr_col);
-          hash             = detail::accumulate(
-            validity_it, validity_it + curr_col.size(), hash, [](auto hash, auto is_valid) {
-              return cudf::detail::hash_combine(hash, is_valid ? NON_NULL_HASH : NULL_HASH);
-            });
-        }
-        */
         if (curr_col.type().id() == type_id::STRUCT) {
           if (curr_col.num_child_columns() == 0) { return hash; }
           // Non-empty structs are assumed to be decomposed and contain only one child
           curr_col = detail::structs_column_device_view(curr_col).get_sliced_child(0);
         } else if (curr_col.type().id() == type_id::LIST) {
           auto list_col = detail::lists_column_device_view(curr_col);
-          /*
-          auto list_sizes = make_list_size_iterator(list_col);
-          hash            = detail::accumulate(
-            list_sizes, list_sizes + list_col.size(), hash, [](auto hash, auto size) {
-              return cudf::detail::hash_combine(hash, hash_fn<size_type>{}(size));
-            });
-          */
-          curr_col = list_col.get_sliced_child();
+          curr_col      = list_col.get_sliced_child();
         }
       }
       for (int i = 0; i < curr_col.size(); ++i) {

From a98c1300d5209c773682d374a77ad3577b65ef76 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 19 Jul 2022 14:19:45 -0700
Subject: [PATCH 09/40] Reorder members to match constructor signature.

---
 cpp/include/cudf/table/experimental/row_operators.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index fc30462c00e..868c79a996f 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -1381,9 +1381,9 @@ class element_hasher {
     CUDF_UNREACHABLE("Unsupported type in hash.");
   }
 
+  Nullate _check_nulls;        ///< Whether to check for nulls
   uint32_t _seed;              ///< The seed to use for hashing
   hash_value_type _null_hash;  ///< Hash value to use for null elements
-  Nullate _check_nulls;        ///< Whether to check for nulls
 };
 
 /**

From 659bca9f3891c2227472d0a283f5c8890dbbac0c Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 19 Jul 2022 14:28:37 -0700
Subject: [PATCH 10/40] Update nested type hashing to act more like the
 previous serial hash code.

---
 cpp/src/hash/spark_murmur_hash.cu | 39 +++++++++++++++++--------------
 1 file changed, 22 insertions(+), 17 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 16c0bbf6507..e94e2c8d679 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -74,11 +74,10 @@ class device_spark_row_hasher {
   template <template <typename> class hash_fn>
   class element_hasher_adapter {
    public:
-    __device__ element_hasher_adapter(
-      Nullate check_nulls,
-      uint32_t seed             = DEFAULT_HASH_SEED,
-      hash_value_type null_hash = std::numeric_limits<hash_value_type>::max()) noexcept
-      : _element_hasher(check_nulls, seed, null_hash), _check_nulls(check_nulls)
+    __device__ element_hasher_adapter(Nullate check_nulls,
+                                      uint32_t seed,
+                                      hash_value_type null_hash) noexcept
+      : _check_nulls(check_nulls), _seed(seed), _null_hash(null_hash)
     {
     }
 
@@ -86,18 +85,19 @@ class device_spark_row_hasher {
     __device__ hash_value_type operator()(column_device_view const& col,
                                           size_type row_index) const noexcept
     {
-      return _element_hasher.template operator()<T>(col, row_index);
+      auto const hasher = cudf::experimental::row::hash::element_hasher<hash_fn, Nullate>(
+        _check_nulls, _seed, _null_hash);
+      return hasher.template operator()<T>(col, row_index);
     }
 
     template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
     __device__ hash_value_type operator()(column_device_view const& col,
                                           size_type row_index) const noexcept
     {
-      auto hash                   = hash_value_type{0};
       column_device_view curr_col = col.slice(row_index, 1);
       while (is_nested(curr_col.type())) {
         if (curr_col.type().id() == type_id::STRUCT) {
-          if (curr_col.num_child_columns() == 0) { return hash; }
+          if (curr_col.num_child_columns() == 0) { return _seed; }
           // Non-empty structs are assumed to be decomposed and contain only one child
           curr_col = detail::structs_column_device_view(curr_col).get_sliced_child(0);
         } else if (curr_col.type().id() == type_id::LIST) {
@@ -105,17 +105,22 @@ class device_spark_row_hasher {
           curr_col      = list_col.get_sliced_child();
         }
       }
-      for (int i = 0; i < curr_col.size(); ++i) {
-        hash =
-          cudf::detail::hash_combine(hash,
-                                     type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
-                                       curr_col.type(), _element_hasher, curr_col, i));
-      }
-      return hash;
+
+      return detail::accumulate(
+        thrust::make_counting_iterator(0),
+        thrust::make_counting_iterator(curr_col.size()),
+        _seed,
+        [curr_col, nulls = this->_check_nulls] __device__(auto hash, auto element_index) {
+          auto const hasher =
+            cudf::experimental::row::hash::element_hasher<hash_fn, Nullate>(nulls, hash, hash);
+          return cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
+            curr_col.type(), hasher, curr_col, element_index);
+        });
     }
 
-    cudf::experimental::row::hash::element_hasher<hash_fn, Nullate> const _element_hasher;
-    Nullate const _check_nulls;
+    Nullate const _check_nulls;        ///< Whether to check for nulls
+    uint32_t const _seed;              ///< The seed to use for hashing
+    hash_value_type const _null_hash;  ///< Hash value to use for null elements
   };
 
   CUDF_HOST_DEVICE device_spark_row_hasher(Nullate check_nulls,

From fcca18a9861df79c6fa732a617e2cbda868a6ac2 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 19 Jul 2022 14:28:52 -0700
Subject: [PATCH 11/40] Use seed 42 for consistency with Spark.

---
 cpp/tests/hashing/hash_test.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index b78c2dfd1b2..6d909c56239 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -827,7 +827,7 @@ TEST_F(SparkMurmurHash3Test, ListValues)
                                                                 -912918097,
                                                                 -912918097};
 
-  auto output = cudf::hash(cudf::table_view({*list_column}), cudf::hash_id::HASH_SPARK_MURMUR3);
+  auto output = cudf::hash(cudf::table_view({*list_column}), cudf::hash_id::HASH_SPARK_MURMUR3, 42);
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
 }
 

From 189f3bfb692e0f517c8142667b2b17fed383b83b Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 20 Jul 2022 07:31:26 -0700
Subject: [PATCH 12/40] Clean up.

---
 cpp/src/hash/spark_murmur_hash.cu | 7 +++----
 1 file changed, 3 insertions(+), 4 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index e94e2c8d679..3500fc0af1f 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -101,14 +101,13 @@ class device_spark_row_hasher {
           // Non-empty structs are assumed to be decomposed and contain only one child
           curr_col = detail::structs_column_device_view(curr_col).get_sliced_child(0);
         } else if (curr_col.type().id() == type_id::LIST) {
-          auto list_col = detail::lists_column_device_view(curr_col);
-          curr_col      = list_col.get_sliced_child();
+          curr_col = detail::lists_column_device_view(curr_col).get_sliced_child();
         }
       }
 
       return detail::accumulate(
-        thrust::make_counting_iterator(0),
-        thrust::make_counting_iterator(curr_col.size()),
+        thrust::counting_iterator(0),
+        thrust::counting_iterator(curr_col.size()),
         _seed,
         [curr_col, nulls = this->_check_nulls] __device__(auto hash, auto element_index) {
           auto const hasher =

From ee754a7a1c77aa12db55536d7a2231435535e678 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 20 Jul 2022 07:31:38 -0700
Subject: [PATCH 13/40] Fix bug in test (wrong index).

---
 cpp/tests/hashing/hash_test.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index 6d909c56239..94d819be2e9 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -786,7 +786,7 @@ TEST_F(SparkMurmurHash3Test, ListValues)
   auto child_validity =
     cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; });
   auto nested_validity =
-    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0 && i != 15; });
+    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0 && i != 14; });
   auto nested_list = cudf::test::lists_column_wrapper<int>({{},
                                                             {1},
                                                             {1, 2},

From b163b4965cefa0501a76d1d04a528e3f472cb559 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 20 Jul 2022 09:32:41 -0700
Subject: [PATCH 14/40] Template row_hasher on device_row_hasher class.

---
 .../cudf/table/experimental/row_operators.cuh | 18 +++--
 cpp/src/hash/spark_murmur_hash.cu             | 68 +++----------------
 2 files changed, 21 insertions(+), 65 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index 868c79a996f..c3b6acfb357 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -824,8 +824,15 @@ class two_table_comparator {
 }  // namespace lexicographic
 
 namespace hash {
+
+template <template <typename> class hash_function, typename Nullate>
+class device_row_hasher;
+
+template <template <template <typename> class hash_function, typename> class device_row_hasher =
+            device_row_hasher>
 class row_hasher;
-}
+
+}  // namespace hash
 
 namespace equality {
 
@@ -1134,7 +1141,7 @@ struct preprocessed_table {
 
   using table_device_view_owner =
     std::invoke_result_t<decltype(table_device_view::create), table_view, rmm::cuda_stream_view>;
-  ///< Type returned by table creation.
+  ///< Owning type returned by table device view creation.
 
   /**
    * @brief Construct a preprocessed table.
@@ -1394,7 +1401,7 @@ class element_hasher {
  */
 template <template <typename> class hash_function, typename Nullate>
 class device_row_hasher {
-  friend class row_hasher;  ///< Allow row_hasher to access private members.
+  friend class row_hasher<device_row_hasher>;  ///< Allow row_hasher to access private members.
 
  public:
   device_row_hasher() = delete;
@@ -1487,12 +1494,12 @@ class device_row_hasher {
   CUDF_HOST_DEVICE device_row_hasher(Nullate check_nulls,
                                      table_device_view t,
                                      uint32_t seed = DEFAULT_HASH_SEED) noexcept
-    : _table{t}, _seed(seed), _check_nulls{check_nulls}
+    : _check_nulls{check_nulls}, _table{t}, _seed(seed)
   {
   }
 
-  table_device_view const _table;
   Nullate const _check_nulls;
+  table_device_view const _table;
   uint32_t const _seed;
 };
 
@@ -1505,6 +1512,7 @@ using preprocessed_table = row::equality::preprocessed_table;
  * @brief Computes the hash value of a row in the given table.
  *
  */
+template <template <template <typename> class, typename> class device_row_hasher>
 class row_hasher {
  public:
   /**
diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 3500fc0af1f..68f50169791 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -38,7 +38,8 @@ namespace {
  */
 template <template <typename> class hash_function, typename Nullate>
 class device_spark_row_hasher {
-  friend class spark_row_hasher;  ///< Allow spark_row_hasher to access private members.
+  friend class cudf::experimental::row::hash::row_hasher<
+    device_spark_row_hasher>;  ///< Allow row_hasher to access private members.
 
  public:
   device_spark_row_hasher() = delete;
@@ -125,69 +126,15 @@ class device_spark_row_hasher {
   CUDF_HOST_DEVICE device_spark_row_hasher(Nullate check_nulls,
                                            table_device_view t,
                                            uint32_t seed = DEFAULT_HASH_SEED) noexcept
-    : _table{t}, _seed(seed), _check_nulls{check_nulls}
+    : _check_nulls{check_nulls}, _table{t}, _seed(seed)
   {
   }
 
-  table_device_view const _table;
   Nullate const _check_nulls;
+  table_device_view const _table;
   uint32_t const _seed;
 };
 
-using preprocessed_table = cudf::experimental::row::hash::preprocessed_table;
-
-/**
- * @brief Computes the hash value of a row in the given table.
- *
- */
-class spark_row_hasher {
- public:
-  /**
-   * @brief Construct an owning object for hashing the rows of a table
-   *
-   * @param t The table containing rows to hash
-   * @param stream The stream to construct this object on. Not the stream that will be used for
-   * comparisons using this object.
-   */
-  spark_row_hasher(table_view const& t, rmm::cuda_stream_view stream)
-    : d_t(preprocessed_table::create(t, stream))
-  {
-  }
-
-  /**
-   * @brief Construct an owning object for hashing the rows of a table from an existing
-   * preprocessed_table
-   *
-   * This constructor allows independently constructing a `preprocessed_table` and sharing it among
-   * multiple `spark_row_hasher` and `equality::self_comparator` objects.
-   *
-   * @param t A table preprocessed for hashing or equality.
-   */
-  spark_row_hasher(std::shared_ptr<preprocessed_table> t) : d_t{std::move(t)} {}
-
-  /**
-   * @brief Get the hash operator to use on the device
-   *
-   * Returns a unary callable, `F`, with signature `hash_function::hash_value_type F(size_type)`.
-   *
-   * `F(i)` returns the hash of row i.
-   *
-   * @tparam Nullate A cudf::nullate type describing whether to check for nulls
-   * @param nullate Indicates if any input column contains nulls
-   * @param seed The seed to use for the hash function
-   * @return A hash operator to use on the device
-   */
-  template <template <typename> class hash_function = detail::default_hash, typename Nullate>
-  device_spark_row_hasher<hash_function, Nullate> device_hasher(
-    Nullate nullate = {}, uint32_t seed = DEFAULT_HASH_SEED) const
-  {
-    return device_spark_row_hasher<hash_function, Nullate>(nullate, *d_t, seed);
-  }
-
- private:
-  std::shared_ptr<preprocessed_table> d_t;
-};
-
 }  // namespace
 
 std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
@@ -209,9 +156,10 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
   // Return early if there's nothing to hash
   if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
 
-  bool const nullable   = has_nulls(input);
-  auto const row_hasher = spark_row_hasher(input, stream);
-  auto output_view      = output->mutable_view();
+  bool const nullable = has_nulls(input);
+  auto const row_hasher =
+    cudf::experimental::row::hash::row_hasher<device_spark_row_hasher>(input, stream);
+  auto output_view = output->mutable_view();
 
   // Compute the hash value for each row
   thrust::tabulate(rmm::exec_policy(stream),

From 387390af4ec763aca706b2a8b1ce575ca7df3288 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 20 Jul 2022 09:44:48 -0700
Subject: [PATCH 15/40] Fix up friends / private methods.

---
 .../cudf/table/experimental/row_operators.cuh     | 15 +++++++--------
 1 file changed, 7 insertions(+), 8 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index c3b6acfb357..ac37c33f531 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -1139,16 +1139,16 @@ struct preprocessed_table {
   static std::shared_ptr<preprocessed_table> create(table_view const& table,
                                                     rmm::cuda_stream_view stream);
 
+ private:
+  friend class self_comparator;       ///< Allow self_comparator to access private members
+  friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
+
+  template <template <template <typename> class hash_function, typename> class device_row_hasher>
+  friend class hash::row_hasher;  ///< Allow row_hasher to access private members
+
   using table_device_view_owner =
     std::invoke_result_t<decltype(table_device_view::create), table_view, rmm::cuda_stream_view>;
-  ///< Owning type returned by table device view creation.
 
-  /**
-   * @brief Construct a preprocessed table.
-   *
-   * @param table The owning table device view
-   * @param null_buffers Null masks superimposed from parent columns
-   */
   preprocessed_table(table_device_view_owner&& table,
                      std::vector<rmm::device_buffer>&& null_buffers)
     : _t(std::move(table)), _null_buffers(std::move(null_buffers))
@@ -1162,7 +1162,6 @@ struct preprocessed_table {
    */
   operator table_device_view() { return *_t; }
 
- private:
   table_device_view_owner _t;
   std::vector<rmm::device_buffer> _null_buffers;
 };

From 2218338d6e987f5441c06a77adc31ffdf5658ec6 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Fri, 22 Jul 2022 13:41:10 -0500
Subject: [PATCH 16/40] Enable list types in spark32BitMurmurHash3.

---
 java/src/main/java/ai/rapids/cudf/ColumnVector.java | 1 -
 1 file changed, 1 deletion(-)

diff --git a/java/src/main/java/ai/rapids/cudf/ColumnVector.java b/java/src/main/java/ai/rapids/cudf/ColumnVector.java
index aab8e7dd475..c8e9cc313a0 100644
--- a/java/src/main/java/ai/rapids/cudf/ColumnVector.java
+++ b/java/src/main/java/ai/rapids/cudf/ColumnVector.java
@@ -778,7 +778,6 @@ public static ColumnVector spark32BitMurmurHash3(int seed, ColumnView columns[])
       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().equals(DType.LIST) : "List columns are not supported";
       columnViews[i] = columns[i].getNativeView();
     }
     return new ColumnVector(hash(columnViews, HashType.HASH_SPARK_MURMUR3.getNativeId(), seed));

From bad0b5eb1a267ba9195331978b4de5d9349bc188 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Fri, 22 Jul 2022 13:41:28 -0500
Subject: [PATCH 17/40] Remove template, always use Spark hash functor.

---
 cpp/src/hash/spark_murmur_hash.cu | 14 ++++++--------
 1 file changed, 6 insertions(+), 8 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 68f50169791..b713d1e5340 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -33,10 +33,9 @@ namespace {
 /**
  * @brief Computes the hash value of a row in the given table.
  *
- * @tparam hash_function Hash functor to use for hashing elements.
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
  */
-template <template <typename> class hash_function, typename Nullate>
+template <typename Nullate>
 class device_spark_row_hasher {
   friend class cudf::experimental::row::hash::row_hasher<
     device_spark_row_hasher>;  ///< Allow row_hasher to access private members.
@@ -58,7 +57,7 @@ class device_spark_row_hasher {
       _seed,
       [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
         return cudf::type_dispatcher(column.type(),
-                                     element_hasher_adapter<hash_function>{nulls, hash, hash},
+                                     element_hasher_adapter{nulls, hash, hash},
                                      column,
                                      row_index);
       });
@@ -72,7 +71,6 @@ class device_spark_row_hasher {
    * When the column is nested, this uses the element_hasher to hash the shape and values of the
    * column.
    */
-  template <template <typename> class hash_fn>
   class element_hasher_adapter {
    public:
     __device__ element_hasher_adapter(Nullate check_nulls,
@@ -82,12 +80,13 @@ class device_spark_row_hasher {
     {
     }
 
+    using hash_functor = cudf::experimental::row::hash::element_hasher<SparkMurmurHash3_32, Nullate>;
+
     template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
     __device__ hash_value_type operator()(column_device_view const& col,
                                           size_type row_index) const noexcept
     {
-      auto const hasher = cudf::experimental::row::hash::element_hasher<hash_fn, Nullate>(
-        _check_nulls, _seed, _null_hash);
+      auto const hasher = hash_functor(_check_nulls, _seed, _null_hash);
       return hasher.template operator()<T>(col, row_index);
     }
 
@@ -111,8 +110,7 @@ class device_spark_row_hasher {
         thrust::counting_iterator(curr_col.size()),
         _seed,
         [curr_col, nulls = this->_check_nulls] __device__(auto hash, auto element_index) {
-          auto const hasher =
-            cudf::experimental::row::hash::element_hasher<hash_fn, Nullate>(nulls, hash, hash);
+          auto const hasher = hash_functor(nulls, hash, hash);
           return cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
             curr_col.type(), hasher, curr_col, element_index);
         });

From 2d30d2c6b7432cd95e49f450ad721f71fc4cc89f Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Fri, 22 Jul 2022 13:45:24 -0500
Subject: [PATCH 18/40] Add summary of differences in Spark hashing.

---
 cpp/src/hash/spark_murmur_hash.cu | 9 +++++++++
 1 file changed, 9 insertions(+)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index b713d1e5340..ca3755a010f 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -33,6 +33,15 @@ namespace {
 /**
  * @brief Computes the hash value of a row in the given table.
  *
+ * This functor uses Spark conventions for MurmurHash3_32 hashing, which
+ * differs from the standard MurmurHash3_32 implementation. These differences
+ * include:
+ * - specialized tail processing in the hash functor
+ * - specializations for decimal types
+ * - special behavior for serially using the output hash as an input seed for
+ *   the next item
+ * - ignorance of null values: [1], [1, null], and [null, 1] have the same hash
+ *
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
  */
 template <typename Nullate>

From 4fee0f8c73fae65fee71fca8c0258a0efd793a5a Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Fri, 22 Jul 2022 13:46:09 -0500
Subject: [PATCH 19/40] clang-format

---
 cpp/src/hash/spark_murmur_hash.cu | 9 ++++-----
 1 file changed, 4 insertions(+), 5 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index ca3755a010f..56c5e04a41c 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -65,10 +65,8 @@ class device_spark_row_hasher {
       _table.end(),
       _seed,
       [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
-        return cudf::type_dispatcher(column.type(),
-                                     element_hasher_adapter{nulls, hash, hash},
-                                     column,
-                                     row_index);
+        return cudf::type_dispatcher(
+          column.type(), element_hasher_adapter{nulls, hash, hash}, column, row_index);
       });
   }
 
@@ -89,7 +87,8 @@ class device_spark_row_hasher {
     {
     }
 
-    using hash_functor = cudf::experimental::row::hash::element_hasher<SparkMurmurHash3_32, Nullate>;
+    using hash_functor =
+      cudf::experimental::row::hash::element_hasher<SparkMurmurHash3_32, Nullate>;
 
     template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
     __device__ hash_value_type operator()(column_device_view const& col,

From 11947231877aa3170b47e1df5dd53b9df6bd5333 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Fri, 22 Jul 2022 11:55:51 -0700
Subject: [PATCH 20/40] Switch back to template parameter so the device row
 hasher matches the template template parameter needed upstream.

---
 cpp/src/hash/spark_murmur_hash.cu | 13 ++++++++-----
 1 file changed, 8 insertions(+), 5 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 56c5e04a41c..b2b3c474ee5 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -42,9 +42,10 @@ namespace {
  *   the next item
  * - ignorance of null values: [1], [1, null], and [null, 1] have the same hash
  *
+ * @tparam hash_function Hash functor to use for hashing elements.
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
  */
-template <typename Nullate>
+template <template <typename> class hash_function, typename Nullate>
 class device_spark_row_hasher {
   friend class cudf::experimental::row::hash::row_hasher<
     device_spark_row_hasher>;  ///< Allow row_hasher to access private members.
@@ -65,8 +66,10 @@ class device_spark_row_hasher {
       _table.end(),
       _seed,
       [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
-        return cudf::type_dispatcher(
-          column.type(), element_hasher_adapter{nulls, hash, hash}, column, row_index);
+        return cudf::type_dispatcher(column.type(),
+                                     element_hasher_adapter<hash_function>{nulls, hash, hash},
+                                     column,
+                                     row_index);
       });
   }
 
@@ -78,6 +81,7 @@ class device_spark_row_hasher {
    * When the column is nested, this uses the element_hasher to hash the shape and values of the
    * column.
    */
+  template <template <typename> class hash_fn>
   class element_hasher_adapter {
    public:
     __device__ element_hasher_adapter(Nullate check_nulls,
@@ -87,8 +91,7 @@ class device_spark_row_hasher {
     {
     }
 
-    using hash_functor =
-      cudf::experimental::row::hash::element_hasher<SparkMurmurHash3_32, Nullate>;
+    using hash_functor = cudf::experimental::row::hash::element_hasher<hash_fn, Nullate>;
 
     template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
     __device__ hash_value_type operator()(column_device_view const& col,

From 873452c7d8392d29cc9a81d135efaefb9a901592 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Mon, 25 Jul 2022 15:45:38 -0500
Subject: [PATCH 21/40] Use has_nested_nulls instead of has_nulls.

---
 cpp/src/hash/spark_murmur_hash.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index b2b3c474ee5..1760851d200 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -165,7 +165,7 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
   // Return early if there's nothing to hash
   if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
 
-  bool const nullable = has_nulls(input);
+  bool const nullable = has_nested_nulls(input);
   auto const row_hasher =
     cudf::experimental::row::hash::row_hasher<device_spark_row_hasher>(input, stream);
   auto output_view = output->mutable_view();

From 28266550746c6bed5b20449c71dec91bd5375f9c Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Mon, 25 Jul 2022 17:20:14 -0500
Subject: [PATCH 22/40] Rename to spark_device_row_hasher.

---
 cpp/src/hash/spark_murmur_hash.cu | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 1760851d200..b1c41617784 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -46,12 +46,12 @@ namespace {
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
  */
 template <template <typename> class hash_function, typename Nullate>
-class device_spark_row_hasher {
+class spark_device_row_hasher {
   friend class cudf::experimental::row::hash::row_hasher<
-    device_spark_row_hasher>;  ///< Allow row_hasher to access private members.
+    spark_device_row_hasher>;  ///< Allow row_hasher to access private members.
 
  public:
-  device_spark_row_hasher() = delete;
+  spark_device_row_hasher() = delete;
 
   /**
    * @brief Return the hash value of a row in the given table.
@@ -132,7 +132,7 @@ class device_spark_row_hasher {
     hash_value_type const _null_hash;  ///< Hash value to use for null elements
   };
 
-  CUDF_HOST_DEVICE device_spark_row_hasher(Nullate check_nulls,
+  CUDF_HOST_DEVICE spark_device_row_hasher(Nullate check_nulls,
                                            table_device_view t,
                                            uint32_t seed = DEFAULT_HASH_SEED) noexcept
     : _check_nulls{check_nulls}, _table{t}, _seed(seed)
@@ -167,7 +167,7 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
 
   bool const nullable = has_nested_nulls(input);
   auto const row_hasher =
-    cudf::experimental::row::hash::row_hasher<device_spark_row_hasher>(input, stream);
+    cudf::experimental::row::hash::row_hasher<spark_device_row_hasher>(input, stream);
   auto output_view = output->mutable_view();
 
   // Compute the hash value for each row

From 3e70ac590643e6d3ff4de93d085db4bb73a4e035 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Mon, 25 Jul 2022 19:27:42 -0500
Subject: [PATCH 23/40] Fix copyright.

Co-authored-by: Vyas Ramasubramani <vyas.ramasubramani@gmail.com>
---
 cpp/src/hash/spark_murmur_hash.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index b1c41617784..540cc8cffa3 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.

From c4586f35cdc425006aa703b40a74770ca1768c29 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 11:54:04 -0500
Subject: [PATCH 24/40] Add test for structs of lists.

---
 cpp/tests/hashing/hash_test.cpp | 77 +++++++++++++++++++++++----------
 1 file changed, 53 insertions(+), 24 deletions(-)

diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index 94d819be2e9..3ed2cc78577 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -21,11 +21,13 @@
 #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>
 
 using cudf::test::fixed_width_column_wrapper;
 using cudf::test::strings_column_wrapper;
 using namespace cudf::test;
+using namespace cudf::test::iterators;
 
 constexpr debug_output_level verbosity{debug_output_level::ALL_ERRORS};
 
@@ -768,25 +770,7 @@ TEST_F(SparkMurmurHash3Test, ListValues)
   df2.show(false)
   */
 
-  /*
-  child data: 1, 1, 2, 1, 2, 3, 1, 2, 3, 1, 2, 3, 1, null, 2, 3, 1, 2, 3, null, 1, 2, 3
-  23 items
-  validity: i != 13, i != 19
-
-  nested validity: i != 0 && i != 15
-  offsets: 0, 0, 1, 3, 6, 8, 9, 10, 12, 13, 16, 18, 19, 20, 22, 23
-  16 items
-
-  row validity: i != 0
-  row offsets: 0, 0, 0, 1, 2, 3, 4, 6, 8, 10, 13, 16
-  11 items
-  */
-
-  auto const null = -1;
-  auto child_validity =
-    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; });
-  auto nested_validity =
-    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0 && i != 14; });
+  auto const null  = -1;
   auto nested_list = cudf::test::lists_column_wrapper<int>({{},
                                                             {1},
                                                             {1, 2},
@@ -796,18 +780,17 @@ TEST_F(SparkMurmurHash3Test, ListValues)
                                                             {1},
                                                             {2, 3},
                                                             {1},
-                                                            {{null, 2, 3}, child_validity},
+                                                            {{null, 2, 3}, nulls_at({0})},
                                                             {1, 2},
                                                             {3},
-                                                            {{null}, child_validity},
+                                                            {{null}, nulls_at({0})},
                                                             {1, 2},
                                                             {},
                                                             {3}},
-                                                           nested_validity);
+                                                           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::detail::make_counting_transform_iterator(0, [](auto i) { return i != 0; });
+  auto list_validity        = nulls_at({0});
   auto list_validity_buffer = cudf::test::detail::make_null_mask(list_validity, list_validity + 11);
   auto list_column          = cudf::make_lists_column(11,
                                              offsets.release(),
@@ -831,6 +814,52 @@ TEST_F(SparkMurmurHash3Test, ListValues)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
 }
 
+TEST_F(SparkMurmurHash3Test, StructOfListValues)
+{
+  /*
+  import org.apache.spark.sql.functions._
+  import org.apache.spark.sql.types.{ArrayType, IntegerType, StructType}
+  import org.apache.spark.sql.Row
+
+  val schema = new StructType()
+    .add("structs", new StructType()
+        .add("a", ArrayType(IntegerType))
+        .add("b", ArrayType(IntegerType)))
+
+  val data = Seq(
+    Row(Row(List(), List())),
+    Row(Row(List(0), List(0))),
+    Row(Row(List(1, null), null)),
+    Row(Row(List(1, null), List())),
+    Row(Row(List(), List(null, 1))),
+    Row(Row(null, List(1))),
+    Row(Row(List(2, 3), List(4, 5))),
+  )
+
+  val df = spark.createDataFrame(
+    spark.sparkContext.parallelize(data), schema)
+
+  val df2 = df.selectExpr("lists", "hash(lists) as hash")
+  df2.printSchema()
+  df2.show(false)
+  */
+
+  auto const null = -1;
+  auto col1       = cudf::test::lists_column_wrapper<int>(
+    {{}, {0}, {{1, null}, nulls_at({1})}, {{1, null}, nulls_at({1})}, {}, {} /*NULL*/, {2, 3}},
+    nulls_at({5}));
+  auto col2 = cudf::test::lists_column_wrapper<int>(
+    {{}, {0}, {} /*NULL*/, {}, {{null, 1}, nulls_at({0})}, {1}, {4, 5}}, 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::hash(cudf::table_view({struct_column}), cudf::hash_id::HASH_SPARK_MURMUR3, 42);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
+}
+
 class MD5HashTest : public cudf::test::BaseFixture {
 };
 

From 1f75acf83f2bbbf9cbc068120e314497b6358993 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 13:23:04 -0500
Subject: [PATCH 25/40] Add failing test for lists of structs.

---
 cpp/tests/hashing/hash_test.cpp | 54 +++++++++++++++++++++++++++++++++
 1 file changed, 54 insertions(+)

diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index 3ed2cc78577..354ed653df8 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -860,6 +860,60 @@ TEST_F(SparkMurmurHash3Test, StructOfListValues)
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
 }
 
+TEST_F(SparkMurmurHash3Test, 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 = fixed_width_column_wrapper<int32_t>({0, null, null, 1, null, null, 2, 2, null, 2, 4},
+                                                  nulls_at({1, 2, 4, 5, 8}));
+  auto col2 = fixed_width_column_wrapper<int32_t>({0, null, null, null, 1, 1, 3, 3, null, 3, 5},
+                                                  nulls_at({1, 2, 3, 8}));
+  auto struct_column = structs_column_wrapper{{col1, col2}, {1, 0, 1, 1, 1, 1, 1, 1, 0, 1, 1}};
+  auto offsets       = fixed_width_column_wrapper<cudf::size_type>{0, 1, 2, 3, 4, 5, 7, 9, 11};
+  auto list_nullmask = std::vector<bool>(1, 8);
+  auto list_validity_buffer =
+    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(),
+                                             cudf::UNKNOWN_NULL_COUNT,
+                                             std::move(list_validity_buffer));
+
+  auto expect = cudf::test::fixed_width_column_wrapper<int32_t>{
+    59727262, 42, 42, -559580957, -559580957, -912918097, 1092624418, 170038658};
+
+  auto output = cudf::hash(cudf::table_view({*list_column}), cudf::hash_id::HASH_SPARK_MURMUR3, 42);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
+}
+
 class MD5HashTest : public cudf::test::BaseFixture {
 };
 

From e179a4c55d82d49771a452049c6e12bad3d277f3 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 13:35:17 -0500
Subject: [PATCH 26/40] Fail if a table with lists of structs is passed.

---
 cpp/src/hash/spark_murmur_hash.cu | 25 +++++++++++++++++++++++++
 cpp/tests/hashing/hash_test.cpp   |  7 +++++++
 2 files changed, 32 insertions(+)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 540cc8cffa3..657bf0a20d1 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -144,6 +144,28 @@ class spark_device_row_hasher {
   uint32_t const _seed;
 };
 
+void check_hash_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_murmur_hash3_32(table_view const& input,
@@ -165,6 +187,9 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
   // Return early if there's nothing to hash
   if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
 
+  // Lists of structs are not supported
+  check_hash_compatibility(input);
+
   bool const nullable = has_nested_nulls(input);
   auto const row_hasher =
     cudf::experimental::row::hash::row_hasher<spark_device_row_hasher>(input, stream);
diff --git a/cpp/tests/hashing/hash_test.cpp b/cpp/tests/hashing/hash_test.cpp
index 354ed653df8..863c186fe44 100644
--- a/cpp/tests/hashing/hash_test.cpp
+++ b/cpp/tests/hashing/hash_test.cpp
@@ -907,11 +907,18 @@ TEST_F(SparkMurmurHash3Test, ListOfStructValues)
                                              cudf::UNKNOWN_NULL_COUNT,
                                              std::move(list_validity_buffer));
 
+  // TODO: Lists of structs are not yet supported. Once support is added,
+  // remove this EXPECT_THROW and uncomment the rest of this test.
+  EXPECT_THROW(cudf::hash(cudf::table_view({*list_column}), cudf::hash_id::HASH_SPARK_MURMUR3, 42),
+               cudf::logic_error);
+
+  /*
   auto expect = cudf::test::fixed_width_column_wrapper<int32_t>{
     59727262, 42, 42, -559580957, -559580957, -912918097, 1092624418, 170038658};
 
   auto output = cudf::hash(cudf::table_view({*list_column}), cudf::hash_id::HASH_SPARK_MURMUR3, 42);
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(expect, output->view(), verbosity);
+  */
 }
 
 class MD5HashTest : public cudf::test::BaseFixture {

From f057bc6efcc5fb805500a2ac82d04238252ddf89 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 13:57:03 -0500
Subject: [PATCH 27/40] Use spark_hash_value_type.

---
 cpp/src/hash/spark_murmur_hash.cu | 9 +++++----
 1 file changed, 5 insertions(+), 4 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 657bf0a20d1..fb8f554178f 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -174,11 +174,12 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
                                               rmm::mr::device_memory_resource* mr)
 {
   // TODO: Spark uses int32_t hash values, but libcudf defines hash_value_type
-  // as uint32_t elsewhere. This should be investigated and unified. I suspect
-  // we should use int32_t everywhere. Also check this for hash seeds. --bdice
-  using hash_value_type = int32_t;
+  // as uint32_t elsewhere. I plan to move the SparkMurmurHash3_32 functor into
+  // this file (since it is only used here), and replace its use of
+  // hash_value_type with spark_hash_value_type. --bdice
+  using spark_hash_value_type = int32_t;
 
-  auto output = make_numeric_column(data_type(type_to_id<hash_value_type>()),
+  auto output = make_numeric_column(data_type(type_to_id<spark_hash_value_type>()),
                                     input.num_rows(),
                                     mask_state::UNALLOCATED,
                                     stream,

From 929c5890fff814a45ff7de520ec59f99779d7032 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 14:04:41 -0500
Subject: [PATCH 28/40] Change template name to DeviceRowHasher.

---
 .../cudf/table/experimental/row_operators.cuh        | 12 ++++++------
 1 file changed, 6 insertions(+), 6 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index ac37c33f531..e699a8c32db 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -828,7 +828,7 @@ namespace hash {
 template <template <typename> class hash_function, typename Nullate>
 class device_row_hasher;
 
-template <template <template <typename> class hash_function, typename> class device_row_hasher =
+template <template <template <typename> class hash_function, typename> class DeviceRowHasher =
             device_row_hasher>
 class row_hasher;
 
@@ -1143,7 +1143,7 @@ struct preprocessed_table {
   friend class self_comparator;       ///< Allow self_comparator to access private members
   friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
 
-  template <template <template <typename> class hash_function, typename> class device_row_hasher>
+  template <template <template <typename> class hash_function, typename> class DeviceRowHasher>
   friend class hash::row_hasher;  ///< Allow row_hasher to access private members
 
   using table_device_view_owner =
@@ -1511,7 +1511,7 @@ using preprocessed_table = row::equality::preprocessed_table;
  * @brief Computes the hash value of a row in the given table.
  *
  */
-template <template <template <typename> class, typename> class device_row_hasher>
+template <template <template <typename> class, typename> class DeviceRowHasher>
 class row_hasher {
  public:
   /**
@@ -1550,10 +1550,10 @@ class row_hasher {
    * @return A hash operator to use on the device
    */
   template <template <typename> class hash_function = detail::default_hash, typename Nullate>
-  device_row_hasher<hash_function, Nullate> device_hasher(Nullate nullate = {},
-                                                          uint32_t seed   = DEFAULT_HASH_SEED) const
+  DeviceRowHasher<hash_function, Nullate> device_hasher(Nullate nullate = {},
+                                                        uint32_t seed   = DEFAULT_HASH_SEED) const
   {
-    return device_row_hasher<hash_function, Nullate>(nullate, *d_t, seed);
+    return DeviceRowHasher<hash_function, Nullate>(nullate, *d_t, seed);
   }
 
  private:

From 248cfaafe75b785e6ce5aa343c53904935d03d5a Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 17:42:38 -0500
Subject: [PATCH 29/40] Improve comment.

---
 cpp/src/hash/spark_murmur_hash.cu | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index fb8f554178f..f9831a1baf7 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -78,8 +78,10 @@ class spark_device_row_hasher {
    * @brief Computes the hash value of an element in the given column.
    *
    * When the column is non-nested, this is a simple wrapper around the element_hasher.
-   * When the column is nested, this uses the element_hasher to hash the shape and values of the
-   * column.
+   * When the column is nested, this uses a seed value to serially compute each
+   * nested element, with the output hash becoming the seed for the next value.
+   * The hash of a null element is the input seed (the previous element's
+   * hash).
    */
   template <template <typename> class hash_fn>
   class element_hasher_adapter {

From f90af13564d7ef2536f537bfa9aaed350e1627ab Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 17:43:03 -0500
Subject: [PATCH 30/40] Rename to spark_murmur_device_row_hasher.

---
 cpp/src/hash/spark_murmur_hash.cu | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index f9831a1baf7..2ca3eed7a3c 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -46,12 +46,12 @@ namespace {
  * @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 {
+class spark_murmur_device_row_hasher {
   friend class cudf::experimental::row::hash::row_hasher<
-    spark_device_row_hasher>;  ///< Allow row_hasher to access private members.
+    spark_murmur_device_row_hasher>;  ///< Allow row_hasher to access private members.
 
  public:
-  spark_device_row_hasher() = delete;
+  spark_murmur_device_row_hasher() = delete;
 
   /**
    * @brief Return the hash value of a row in the given table.
@@ -134,9 +134,9 @@ class spark_device_row_hasher {
     hash_value_type const _null_hash;  ///< Hash value to use for null elements
   };
 
-  CUDF_HOST_DEVICE spark_device_row_hasher(Nullate check_nulls,
-                                           table_device_view t,
-                                           uint32_t seed = DEFAULT_HASH_SEED) noexcept
+  CUDF_HOST_DEVICE spark_murmur_device_row_hasher(Nullate check_nulls,
+                                                  table_device_view t,
+                                                  uint32_t seed = DEFAULT_HASH_SEED) noexcept
     : _check_nulls{check_nulls}, _table{t}, _seed(seed)
   {
   }
@@ -195,7 +195,7 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
 
   bool const nullable = has_nested_nulls(input);
   auto const row_hasher =
-    cudf::experimental::row::hash::row_hasher<spark_device_row_hasher>(input, stream);
+    cudf::experimental::row::hash::row_hasher<spark_murmur_device_row_hasher>(input, stream);
   auto output_view = output->mutable_view();
 
   // Compute the hash value for each row

From 4a808760fe140b63e86c52749da92e5b7ce52ae9 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 17:59:45 -0500
Subject: [PATCH 31/40] Use spark_hash_value_type more consistently, and
 enforce that null hash values are always the seed.

---
 cpp/src/hash/spark_murmur_hash.cu | 34 ++++++++++++++-----------------
 1 file changed, 15 insertions(+), 19 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 2ca3eed7a3c..7ff143433ba 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -30,6 +30,8 @@ namespace detail {
 
 namespace {
 
+using spark_hash_value_type = int32_t;
+
 /**
  * @brief Computes the hash value of a row in the given table.
  *
@@ -66,10 +68,8 @@ class spark_murmur_device_row_hasher {
       _table.end(),
       _seed,
       [row_index, nulls = this->_check_nulls] __device__(auto hash, auto column) {
-        return cudf::type_dispatcher(column.type(),
-                                     element_hasher_adapter<hash_function>{nulls, hash, hash},
-                                     column,
-                                     row_index);
+        return cudf::type_dispatcher(
+          column.type(), element_hasher_adapter<hash_function>{nulls, hash}, column, row_index);
       });
   }
 
@@ -86,26 +86,24 @@ class spark_murmur_device_row_hasher {
   template <template <typename> class hash_fn>
   class element_hasher_adapter {
    public:
-    __device__ element_hasher_adapter(Nullate check_nulls,
-                                      uint32_t seed,
-                                      hash_value_type null_hash) noexcept
-      : _check_nulls(check_nulls), _seed(seed), _null_hash(null_hash)
+    __device__ element_hasher_adapter(Nullate check_nulls, uint32_t seed) noexcept
+      : _check_nulls(check_nulls), _seed(seed)
     {
     }
 
     using hash_functor = cudf::experimental::row::hash::element_hasher<hash_fn, Nullate>;
 
     template <typename T, CUDF_ENABLE_IF(not cudf::is_nested<T>())>
-    __device__ hash_value_type operator()(column_device_view const& col,
-                                          size_type row_index) const noexcept
+    __device__ spark_hash_value_type operator()(column_device_view const& col,
+                                                size_type row_index) const noexcept
     {
-      auto const hasher = hash_functor(_check_nulls, _seed, _null_hash);
+      auto const hasher = hash_functor(_check_nulls, _seed, _seed);
       return hasher.template operator()<T>(col, row_index);
     }
 
     template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
-    __device__ hash_value_type operator()(column_device_view const& col,
-                                          size_type row_index) const noexcept
+    __device__ spark_hash_value_type operator()(column_device_view const& col,
+                                                size_type row_index) const noexcept
     {
       column_device_view curr_col = col.slice(row_index, 1);
       while (is_nested(curr_col.type())) {
@@ -129,9 +127,8 @@ class spark_murmur_device_row_hasher {
         });
     }
 
-    Nullate const _check_nulls;        ///< Whether to check for nulls
-    uint32_t const _seed;              ///< The seed to use for hashing
-    hash_value_type const _null_hash;  ///< Hash value to use for null elements
+    Nullate const _check_nulls;  ///< Whether to check for nulls
+    uint32_t const _seed;        ///< The seed to use for hashing, also returned for null elements
   };
 
   CUDF_HOST_DEVICE spark_murmur_device_row_hasher(Nullate check_nulls,
@@ -179,7 +176,6 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
   // as uint32_t elsewhere. I plan to move the SparkMurmurHash3_32 functor into
   // this file (since it is only used here), and replace its use of
   // hash_value_type with spark_hash_value_type. --bdice
-  using spark_hash_value_type = int32_t;
 
   auto output = make_numeric_column(data_type(type_to_id<spark_hash_value_type>()),
                                     input.num_rows(),
@@ -200,8 +196,8 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
 
   // Compute the hash value for each row
   thrust::tabulate(rmm::exec_policy(stream),
-                   output_view.begin<hash_value_type>(),
-                   output_view.end<hash_value_type>(),
+                   output_view.begin<spark_hash_value_type>(),
+                   output_view.end<spark_hash_value_type>(),
                    row_hasher.device_hasher<SparkMurmurHash3_32>(nullable, seed));
 
   return output;

From 5be3ea09977044f3d5111ebb226ca0c156ecb430 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 18:27:27 -0500
Subject: [PATCH 32/40] Added developer documentation.

---
 cpp/src/hash/spark_murmur_hash.cu | 43 ++++++++++++++++++++-----------
 1 file changed, 28 insertions(+), 15 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 7ff143433ba..9b01fdc5a1c 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -30,21 +30,38 @@ namespace detail {
 
 namespace {
 
+// TODO: Spark uses int32_t hash values, but libcudf defines hash_value_type as
+// uint32_t elsewhere. I plan to move the SparkMurmurHash3_32 functor into this
+// file (since it is only used here), and replace its use of hash_value_type
+// with spark_hash_value_type. --bdice
 using spark_hash_value_type = int32_t;
 
 /**
  * @brief Computes the hash value of a row in the given table.
  *
- * This functor uses Spark conventions for MurmurHash3_32 hashing, which
- * differs from the standard MurmurHash3_32 implementation. These differences
+ * This functor uses Spark conventions for Murmur hashing, which differs from
+ * the Murmur implementation used in the rest of libcudf. These differences
  * include:
- * - specialized tail processing in the hash functor
- * - specializations for decimal types
- * - special behavior for serially using the output hash as an input seed for
- *   the next item
- * - ignorance of null values: [1], [1, null], and [null, 1] have the same hash
+ * - Serially using the output hash as an input seed for the next item
+ * - Ignorance of null values
  *
- * @tparam hash_function Hash functor to use for hashing elements.
+ * The serial use of hashes as seeds means that data of different nested types
+ * can exhibit hash collisions. For example, a row of an integer column
+ * containing a 1 will have the same hash as a lists column of integers
+ * containing a list of [1] and a struct column of a single integer column
+ * containing a struct of {1}.
+ *
+ * As a consequence of ignoring null values, inputs like [1], [1, null], and
+ * [null, 1] have the same hash (an expected hash collision). This kind of
+ * collision can also occur across a table of nullable columns and with nulls
+ * in structs ({1, null} and {null, 1} have the same hash). The seed value (the
+ * previous element's hash value) is returned as the hash if an element is
+ * null.
+ *
+ * For additional differences such as special tail processing and decimal type
+ * handling, refer to the SparkMurmurHash3_32 functor.
+ *
+ * @tparam hash_function Hash functor to use for hashing elements. Must be SparkMurmurHash3_32.
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
  */
 template <template <typename> class hash_function, typename Nullate>
@@ -80,8 +97,9 @@ class spark_murmur_device_row_hasher {
    * When the column is non-nested, this is a simple wrapper around the element_hasher.
    * When the column is nested, this uses a seed value to serially compute each
    * nested element, with the output hash becoming the seed for the next value.
-   * The hash of a null element is the input seed (the previous element's
-   * hash).
+   * This requires constructing a new hash functor for each nested element,
+   * using the new seed from the previous element's hash. The hash of a null
+   * element is the input seed (the previous element's hash).
    */
   template <template <typename> class hash_fn>
   class element_hasher_adapter {
@@ -172,11 +190,6 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
                                               rmm::cuda_stream_view stream,
                                               rmm::mr::device_memory_resource* mr)
 {
-  // TODO: Spark uses int32_t hash values, but libcudf defines hash_value_type
-  // as uint32_t elsewhere. I plan to move the SparkMurmurHash3_32 functor into
-  // this file (since it is only used here), and replace its use of
-  // hash_value_type with spark_hash_value_type. --bdice
-
   auto output = make_numeric_column(data_type(type_to_id<spark_hash_value_type>()),
                                     input.num_rows(),
                                     mask_state::UNALLOCATED,

From cf6f6cdc2ca8f0666f6165780fefbb21226c891d Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 26 Jul 2022 18:30:21 -0500
Subject: [PATCH 33/40] Use list initialization.

---
 cpp/src/hash/spark_murmur_hash.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 9b01fdc5a1c..b3d11c7f3ed 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -115,7 +115,7 @@ class spark_murmur_device_row_hasher {
     __device__ spark_hash_value_type operator()(column_device_view const& col,
                                                 size_type row_index) const noexcept
     {
-      auto const hasher = hash_functor(_check_nulls, _seed, _seed);
+      auto const hasher = hash_functor{_check_nulls, _seed, _seed};
       return hasher.template operator()<T>(col, row_index);
     }
 
@@ -139,7 +139,7 @@ class spark_murmur_device_row_hasher {
         thrust::counting_iterator(curr_col.size()),
         _seed,
         [curr_col, nulls = this->_check_nulls] __device__(auto hash, auto element_index) {
-          auto const hasher = hash_functor(nulls, hash, hash);
+          auto const hasher = hash_functor{nulls, hash, hash};
           return cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
             curr_col.type(), hasher, curr_col, element_index);
         });

From 6475517f1c5bb501be1b9c5972e4eb817d7d4832 Mon Sep 17 00:00:00 2001
From: Ryan Lee <ryanlee@nvidia.com>
Date: Tue, 26 Jul 2022 14:34:23 -0700
Subject: [PATCH 34/40] add list test

---
 .../java/ai/rapids/cudf/ColumnVectorTest.java | 49 +++++++++++++++++++
 1 file changed, 49 insertions(+)

diff --git a/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java b/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
index 05abe4958e2..5b2b90cbf69 100644
--- a/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
+++ b/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
@@ -764,6 +764,55 @@ void testSpark32BitMurmur3HashStruct() {
     }
   }
 
+  @Test
+  void testSpark32BitMurmur3HashListsAndNestedLists() {
+    try (ColumnVector stringListCV = ColumnVector.fromLists(
+             new ListType(true, new BasicType(true, DType.STRING)),
+             Arrays.asList(null, "a"),
+             Arrays.asList("B\n", ""),
+             Arrays.asList("dE\"\u0100\t\u0101", " \ud720\ud721"),
+             Collections.singletonList("A very long (greater than 128 bytes/char string) to test a multi" +
+             " hash-step data point in the Murmur3 hash function. This string needed to be longer."),
+             Collections.singletonList(""),
+             null);
+         ColumnVector strings1 = ColumnVector.fromStrings(
+             "a", "B\n", "dE\"\u0100\t\u0101",
+             "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
+             "in the Murmur3 hash function. This string needed to be longer.", null, null);
+         ColumnVector strings2 = ColumnVector.fromStrings(
+             null, "", " \ud720\ud721", null, "", null);
+         ColumnView stringStruct = ColumnView.makeStructView(strings1, strings2);
+         ColumnVector stringExpected = ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{stringStruct});
+         ColumnVector stringResult = ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{stringListCV});
+         ColumnVector intListCV = ColumnVector.fromLists(
+             new ListType(true, new BasicType(true, DType.INT32)),
+             null,
+             Arrays.asList(0, -2, 3),
+             Collections.singletonList(Integer.MAX_VALUE),
+             Arrays.asList(5, -6, null),
+             Collections.singletonList(Integer.MIN_VALUE),
+             null);
+         ColumnVector integers1 = ColumnVector.fromBoxedInts(null, 0, null, 5, Integer.MIN_VALUE, null);
+         ColumnVector integers2 = ColumnVector.fromBoxedInts(null, -2, Integer.MAX_VALUE, null, null, null);
+         ColumnVector integers3 = ColumnVector.fromBoxedInts(null, 3, null, -6, null, null);
+         ColumnVector intExpected =
+             ColumnVector.spark32BitMurmurHash3(1868, new ColumnVector[]{integers1, integers2, integers3});
+         ColumnVector intResult = ColumnVector.spark32BitMurmurHash3(1868, new ColumnVector[]{intListCV});
+         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);
+         ColumnView structCV = ColumnView.makeStructView(intListCV, stringListCV, doubles, floats);
+         ColumnVector nestedExpected =
+             ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{intListCV, strings1, strings2, doubles, floats});
+         ColumnVector nestedResult =
+             ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{structCV})) {
+      assertColumnsAreEqual(stringExpected, stringResult);
+      assertColumnsAreEqual(intExpected, intResult);
+      assertColumnsAreEqual(nestedExpected, nestedResult);
+    }
+  }
+
   @Test
   void testAndNullReconfigureNulls() {
     try (ColumnVector v0 = ColumnVector.fromBoxedInts(0, 100, null, null, Integer.MIN_VALUE, null);

From 028b0bc54acaf719156d264079c89301f0de922a Mon Sep 17 00:00:00 2001
From: Ryan Lee <ryanlee@nvidia.com>
Date: Tue, 26 Jul 2022 18:37:35 -0700
Subject: [PATCH 35/40] nested struct java test

---
 .../java/ai/rapids/cudf/ColumnVectorTest.java | 23 +++++++++++++++++++
 1 file changed, 23 insertions(+)

diff --git a/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java b/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
index 5b2b90cbf69..a0eb1c69c4c 100644
--- a/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
+++ b/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
@@ -764,6 +764,29 @@ void testSpark32BitMurmur3HashStruct() {
     }
   }
 
+  @Test
+  void testSpark32BitMurmur3HashNestedStruct() {
+    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);
+         ColumnView structs1 = ColumnView.makeStructView(strings, integers);
+         ColumnView structs2 = ColumnView.makeStructView(structs1, doubles);
+         ColumnView structs3 = ColumnView.makeStructView(bools);
+         ColumnView structs = ColumnView.makeStructView(structs2, floats, structs3);
+         ColumnVector expected = ColumnVector.spark32BitMurmurHash3(1868, new ColumnVector[]{strings, integers, doubles, floats, bools})) {
+         ColumnVector result = ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{structs});
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
   @Test
   void testSpark32BitMurmur3HashListsAndNestedLists() {
     try (ColumnVector stringListCV = ColumnVector.fromLists(

From 5d05984cf33da2c316db5d700b77b5562e0d7183 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 27 Jul 2022 17:02:44 -0500
Subject: [PATCH 36/40] Remove deleted constructor.

---
 cpp/src/hash/spark_murmur_hash.cu | 2 --
 1 file changed, 2 deletions(-)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index b3d11c7f3ed..f27137334e5 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -70,8 +70,6 @@ class spark_murmur_device_row_hasher {
     spark_murmur_device_row_hasher>;  ///< Allow row_hasher to access private members.
 
  public:
-  spark_murmur_device_row_hasher() = delete;
-
   /**
    * @brief Return the hash value of a row in the given table.
    *

From 9281aeedbc8c24e522f025857cd9a4cf4fe1ae06 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 27 Jul 2022 17:06:20 -0500
Subject: [PATCH 37/40] Require SparkMurmurHash3_32 in constructor.

---
 cpp/src/hash/spark_murmur_hash.cu | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index f27137334e5..7bc2fb0c21a 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -152,6 +152,10 @@ class spark_murmur_device_row_hasher {
                                                   uint32_t seed = DEFAULT_HASH_SEED) noexcept
     : _check_nulls{check_nulls}, _table{t}, _seed(seed)
   {
+    // Error out if passed an unsupported hash_function
+    static_assert(
+      std::is_base_of_v<SparkMurmurHash3_32<int>, hash_function<int>>,
+      "spark_murmur_device_row_hasher only supports the SparkMurmurHash3_32 hash function");
   }
 
   Nullate const _check_nulls;

From fa50d2c2902d12c58ddcc52f485917a7d9e41cae Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 27 Jul 2022 17:07:33 -0500
Subject: [PATCH 38/40] Remove deleted public constructor because the
 constructor is declared to be private.

---
 cpp/include/cudf/table/experimental/row_operators.cuh | 2 --
 1 file changed, 2 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index e699a8c32db..d6f8dd9bdca 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -1403,8 +1403,6 @@ class device_row_hasher {
   friend class row_hasher<device_row_hasher>;  ///< Allow row_hasher to access private members.
 
  public:
-  device_row_hasher() = delete;
-
   /**
    * @brief Return the hash value of a row in the given table.
    *

From ef0ad3bc2c6ea0d3093ec6561a00befbdd434c41 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 27 Jul 2022 17:28:08 -0500
Subject: [PATCH 39/40] Template only the device_hasher method.

---
 .../cudf/table/experimental/row_operators.cuh | 19 ++++++------------
 cpp/src/hash/spark_murmur_hash.cu             | 20 +++++++++----------
 2 files changed, 16 insertions(+), 23 deletions(-)

diff --git a/cpp/include/cudf/table/experimental/row_operators.cuh b/cpp/include/cudf/table/experimental/row_operators.cuh
index d6f8dd9bdca..fcdc7948667 100644
--- a/cpp/include/cudf/table/experimental/row_operators.cuh
+++ b/cpp/include/cudf/table/experimental/row_operators.cuh
@@ -824,14 +824,7 @@ class two_table_comparator {
 }  // namespace lexicographic
 
 namespace hash {
-
-template <template <typename> class hash_function, typename Nullate>
-class device_row_hasher;
-
-template <template <template <typename> class hash_function, typename> class DeviceRowHasher =
-            device_row_hasher>
 class row_hasher;
-
 }  // namespace hash
 
 namespace equality {
@@ -1142,9 +1135,7 @@ struct preprocessed_table {
  private:
   friend class self_comparator;       ///< Allow self_comparator to access private members
   friend class two_table_comparator;  ///< Allow two_table_comparator to access private members
-
-  template <template <template <typename> class hash_function, typename> class DeviceRowHasher>
-  friend class hash::row_hasher;  ///< Allow row_hasher to access private members
+  friend class hash::row_hasher;      ///< Allow row_hasher to access private members
 
   using table_device_view_owner =
     std::invoke_result_t<decltype(table_device_view::create), table_view, rmm::cuda_stream_view>;
@@ -1400,7 +1391,7 @@ class element_hasher {
  */
 template <template <typename> class hash_function, typename Nullate>
 class device_row_hasher {
-  friend class row_hasher<device_row_hasher>;  ///< Allow row_hasher to access private members.
+  friend class row_hasher;  ///< Allow row_hasher to access private members.
 
  public:
   /**
@@ -1509,7 +1500,6 @@ using preprocessed_table = row::equality::preprocessed_table;
  * @brief Computes the hash value of a row in the given table.
  *
  */
-template <template <template <typename> class, typename> class DeviceRowHasher>
 class row_hasher {
  public:
   /**
@@ -1547,7 +1537,10 @@ class row_hasher {
    * @param seed The seed to use for the hash function
    * @return A hash operator to use on the device
    */
-  template <template <typename> class hash_function = detail::default_hash, typename Nullate>
+  template <template <typename> class hash_function = detail::default_hash,
+            template <template <typename> class, typename>
+            class DeviceRowHasher = device_row_hasher,
+            typename Nullate>
   DeviceRowHasher<hash_function, Nullate> device_hasher(Nullate nullate = {},
                                                         uint32_t seed   = DEFAULT_HASH_SEED) const
   {
diff --git a/cpp/src/hash/spark_murmur_hash.cu b/cpp/src/hash/spark_murmur_hash.cu
index 7bc2fb0c21a..613def643bf 100644
--- a/cpp/src/hash/spark_murmur_hash.cu
+++ b/cpp/src/hash/spark_murmur_hash.cu
@@ -66,8 +66,8 @@ using spark_hash_value_type = int32_t;
  */
 template <template <typename> class hash_function, typename Nullate>
 class spark_murmur_device_row_hasher {
-  friend class cudf::experimental::row::hash::row_hasher<
-    spark_murmur_device_row_hasher>;  ///< Allow row_hasher to access private members.
+  friend class cudf::experimental::row::hash::row_hasher;  ///< Allow row_hasher to access private
+                                                           ///< members.
 
  public:
   /**
@@ -204,16 +204,16 @@ std::unique_ptr<column> spark_murmur_hash3_32(table_view const& input,
   // Lists of structs are not supported
   check_hash_compatibility(input);
 
-  bool const nullable = has_nested_nulls(input);
-  auto const row_hasher =
-    cudf::experimental::row::hash::row_hasher<spark_murmur_device_row_hasher>(input, stream);
-  auto output_view = output->mutable_view();
+  bool const nullable   = has_nested_nulls(input);
+  auto const row_hasher = cudf::experimental::row::hash::row_hasher(input, stream);
+  auto output_view      = output->mutable_view();
 
   // Compute the hash value for each row
-  thrust::tabulate(rmm::exec_policy(stream),
-                   output_view.begin<spark_hash_value_type>(),
-                   output_view.end<spark_hash_value_type>(),
-                   row_hasher.device_hasher<SparkMurmurHash3_32>(nullable, seed));
+  thrust::tabulate(
+    rmm::exec_policy(stream),
+    output_view.begin<spark_hash_value_type>(),
+    output_view.end<spark_hash_value_type>(),
+    row_hasher.device_hasher<SparkMurmurHash3_32, spark_murmur_device_row_hasher>(nullable, seed));
 
   return output;
 }

From aaadac2d49e753e421e942a0cc198c8a8541625e Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Wed, 27 Jul 2022 22:11:00 -0500
Subject: [PATCH 40/40] Temporarily remove JNI tests due to an outstanding
 allocation bug.

---
 .../java/ai/rapids/cudf/ColumnVectorTest.java | 72 -------------------
 1 file changed, 72 deletions(-)

diff --git a/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java b/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
index a0eb1c69c4c..05abe4958e2 100644
--- a/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
+++ b/java/src/test/java/ai/rapids/cudf/ColumnVectorTest.java
@@ -764,78 +764,6 @@ void testSpark32BitMurmur3HashStruct() {
     }
   }
 
-  @Test
-  void testSpark32BitMurmur3HashNestedStruct() {
-    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);
-         ColumnView structs1 = ColumnView.makeStructView(strings, integers);
-         ColumnView structs2 = ColumnView.makeStructView(structs1, doubles);
-         ColumnView structs3 = ColumnView.makeStructView(bools);
-         ColumnView structs = ColumnView.makeStructView(structs2, floats, structs3);
-         ColumnVector expected = ColumnVector.spark32BitMurmurHash3(1868, new ColumnVector[]{strings, integers, doubles, floats, bools})) {
-         ColumnVector result = ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{structs});
-      assertColumnsAreEqual(expected, result);
-    }
-  }
-
-  @Test
-  void testSpark32BitMurmur3HashListsAndNestedLists() {
-    try (ColumnVector stringListCV = ColumnVector.fromLists(
-             new ListType(true, new BasicType(true, DType.STRING)),
-             Arrays.asList(null, "a"),
-             Arrays.asList("B\n", ""),
-             Arrays.asList("dE\"\u0100\t\u0101", " \ud720\ud721"),
-             Collections.singletonList("A very long (greater than 128 bytes/char string) to test a multi" +
-             " hash-step data point in the Murmur3 hash function. This string needed to be longer."),
-             Collections.singletonList(""),
-             null);
-         ColumnVector strings1 = ColumnVector.fromStrings(
-             "a", "B\n", "dE\"\u0100\t\u0101",
-             "A very long (greater than 128 bytes/char string) to test a multi hash-step data point " +
-             "in the Murmur3 hash function. This string needed to be longer.", null, null);
-         ColumnVector strings2 = ColumnVector.fromStrings(
-             null, "", " \ud720\ud721", null, "", null);
-         ColumnView stringStruct = ColumnView.makeStructView(strings1, strings2);
-         ColumnVector stringExpected = ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{stringStruct});
-         ColumnVector stringResult = ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{stringListCV});
-         ColumnVector intListCV = ColumnVector.fromLists(
-             new ListType(true, new BasicType(true, DType.INT32)),
-             null,
-             Arrays.asList(0, -2, 3),
-             Collections.singletonList(Integer.MAX_VALUE),
-             Arrays.asList(5, -6, null),
-             Collections.singletonList(Integer.MIN_VALUE),
-             null);
-         ColumnVector integers1 = ColumnVector.fromBoxedInts(null, 0, null, 5, Integer.MIN_VALUE, null);
-         ColumnVector integers2 = ColumnVector.fromBoxedInts(null, -2, Integer.MAX_VALUE, null, null, null);
-         ColumnVector integers3 = ColumnVector.fromBoxedInts(null, 3, null, -6, null, null);
-         ColumnVector intExpected =
-             ColumnVector.spark32BitMurmurHash3(1868, new ColumnVector[]{integers1, integers2, integers3});
-         ColumnVector intResult = ColumnVector.spark32BitMurmurHash3(1868, new ColumnVector[]{intListCV});
-         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);
-         ColumnView structCV = ColumnView.makeStructView(intListCV, stringListCV, doubles, floats);
-         ColumnVector nestedExpected =
-             ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{intListCV, strings1, strings2, doubles, floats});
-         ColumnVector nestedResult =
-             ColumnVector.spark32BitMurmurHash3(1868, new ColumnView[]{structCV})) {
-      assertColumnsAreEqual(stringExpected, stringResult);
-      assertColumnsAreEqual(intExpected, intResult);
-      assertColumnsAreEqual(nestedExpected, nestedResult);
-    }
-  }
-
   @Test
   void testAndNullReconfigureNulls() {
     try (ColumnVector v0 = ColumnVector.fromBoxedInts(0, 100, null, null, Integer.MIN_VALUE, null);