From 12e83e7fd3ff9ed6aec8f6a30c06d803eba15972 Mon Sep 17 00:00:00 2001
From: ustcfy <fengyan_@mail.ustc.edu.cn>
Date: Tue, 29 Oct 2024 15:29:50 +0800
Subject: [PATCH 01/25] Add HiveHash support for nested types

Signed-off-by: ustcfy <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu                 | 269 +++++++++++++++++-
 .../com/nvidia/spark/rapids/jni/Hash.java     |   1 -
 .../com/nvidia/spark/rapids/jni/HashTest.java | 170 +++++++++++
 3 files changed, 429 insertions(+), 11 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 85598565a9..71b25f1b5e 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -157,7 +157,7 @@ hive_hash_value_t __device__ inline hive_hash_function<cudf::timestamp_us>::oper
  * @tparam hash_function Hash functor to use for hashing elements. Must be hive_hash_function.
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
  */
-template <template <typename> class hash_function, typename Nullate>
+template <template <typename> class hash_function, typename Nullate, int MAX_NESTED_DEPTH>
 class hive_device_row_hasher {
  public:
   CUDF_HOST_DEVICE hive_device_row_hasher(Nullate check_nulls, cudf::table_device_view t) noexcept
@@ -182,8 +182,8 @@ class hive_device_row_hasher {
       HIVE_INIT_HASH,
       cuda::proclaim_return_type<hive_hash_value_t>(
         [row_index, nulls = this->_check_nulls] __device__(auto hash, auto const& column) {
-          auto cur_hash = cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
-            column.type(), element_hasher_adapter{nulls}, column, row_index);
+          auto cur_hash =
+            cudf::type_dispatcher(column.type(), element_hasher_adapter{nulls}, column, row_index);
           return HIVE_HASH_FACTOR * hash + cur_hash;
         }));
   }
@@ -191,8 +191,6 @@ class hive_device_row_hasher {
  private:
   /**
    * @brief Computes the hash value of an element in the given column.
-   *
-   * Only supported non nested types now
    */
   class element_hasher_adapter {
    public:
@@ -210,11 +208,228 @@ class hive_device_row_hasher {
       return this->hash_functor.template operator()<T>(col, row_index);
     }
 
+    /**
+     * @brief A structure representing an element in the stack used for processing columns.
+     *
+     * This structure is used to keep track of the current column being processed, the index of the
+     * next child column to process, and current hash value for the column.
+     *
+     * @param column The current column being processed.
+     * @param child_idx The index of the next child column to process.
+     * @param cur_hash The current hash value for the column.
+     *
+     * @note The default constructor is deleted to prevent uninitialized usage.
+     *
+     * @constructor
+     * @param col The column device view to be processed.
+     */
+    class col_stack_element {
+     private:
+      cudf::column_device_view column;  // current column
+      hive_hash_value_t cur_hash;       // current hash value of the column
+      int child_idx;  // index of the child column to process next, initialized as 0
+
+     public:
+      __device__ col_stack_element() =
+        delete;  // Because the default constructor of `cudf::column_device_view` is deleted
+
+      __device__ col_stack_element(cudf::column_device_view col)
+        : column(col), child_idx(0), cur_hash(HIVE_INIT_HASH)
+      {
+      }
+
+      __device__ void update_cur_hash(hive_hash_value_t child_hash)
+      {
+        this->cur_hash = this->cur_hash * HIVE_HASH_FACTOR + child_hash;
+      }
+
+      __device__ hive_hash_value_t get_hash() { return this->cur_hash; }
+
+      __device__ int child_idx_inc_one() { return this->child_idx++; }
+
+      __device__ int cur_child_idx() { return this->child_idx; }
+
+      __device__ cudf::column_device_view get_column() { return this->column; }
+    };
+
+    typedef col_stack_element* col_stack_element_ptr;
+
+    /**
+     * @brief Functor to compute the hive hash value for a nested column.
+     *
+     * This functor produces the same result as "HiveHash" in Spark for nested types.
+     * The pseudocode of Spark's HiveHash function is as follows:
+     *
+     * hive_hash_value_t hive_hash(NestedType element) {
+     *    hive_hash_value_t hash = HIVE_INIT_HASH;
+     *    for (int i = 0; i < element.num_child(); i++) {
+     *        hash = hash * HIVE_HASH_FACTOR + hive_hash(element.get_child(i));
+     *    }
+     *    return hash;
+     * }
+     *
+     * This functor uses a stack to simulate the recursive process of the above pseudocode.
+     * When an element is popped from the stack, it means that the hash value of it has been
+     * computed. Therefore, we should update the parent's `cur_hash` upon popping the element.
+     *
+     * The algorithm is as follows:
+     *
+     * 1. Initialize the stack and push the root column into the stack.
+     * 2. While the stack is not empty:
+     *    a. Get the top element of the stack. Don't pop it until it is processed.
+     *    b. If the column is a nested column:
+     *        i.  If all child columns are processed, pop the element and update `cur_hash` of its
+     *            parent column.
+     *        ii. Otherwise, push the next child column into the stack.
+     *    c. If the column is a primitive column, compute the hash value, pop the element,
+     *       and update `cur_hash` of its parent column.
+     * 3. Return the hash value of the root column.
+     *
+     * For example, consider the following nested column: `Struct<Struct<int, float>, decimal>`
+     *
+     *      S1
+     *     / \
+     *    S2  d
+     *   / \
+     *  i   f
+     *
+     * The `pop` order of the stack is: i, f, S2, d, S1.
+     * - When i   is popped, S2's cur_hash is updated to `hash(i)`.
+     * - When f   is popped, S2's cur_hash is updated to `hash(i) * HIVE_HASH_FACTOR + hash(f)`,
+     *   which is the hash value of S2.
+     * - When S2  is popped, S1's cur_hash is updated to `hash(S2)`.
+     * - When d   is popped, S1's cur_hash is updated to `hash(S2) * HIVE_HASH_FACTOR + hash(d)`,
+     *   which is the hash value of S1.
+     * - When S1  is popped, the hash value of the root column is returned.
+     *
+     * As lists columns have a different interface from structs columns, we need to handle them
+     * separately.
+     *
+     * For example, consider the following nested column: `List<List<int>>`
+     * lists_column = {{1, 0}, null, {2, null}}
+     *
+     *     L1
+     *     |
+     *     L2
+     *     |
+     *     i
+     *
+     * List level L1:
+     * |Index|      List<list<int>>    |
+     * |-----|-------------------------|
+     * |0    |{{1, 0}, null, {2, null}}|
+     * length: 1
+     * Offsets: 0, 3
+     *
+     * List level L2:
+     * |Index|List<int>|
+     * |-----|---------|
+     * |0    |{1, 0}   |
+     * |1    |null     |
+     * |2    |{2, null}|
+     * length: 3
+     * Offsets: 0, 2, 2, 4
+     * null_mask: 101
+     *
+     * Int level i:
+     * |Index|Int |
+     * |-----|----|
+     * |0    |1   |
+     * |1    |0   |
+     * |2    |2   |
+     * |3    |null|
+     * length: 4
+     *
+     * Since the underlying data loses the null information of the top-level list column,
+     * It will produce different results than Spark to compute the hash value using the
+     * underlying data merely.
+     *
+     * For example, `List<List<int>>` column {{1, 0}, {2, null}} has the same underlying data as the
+     * above `List<List<int>>` column {{1, 0}, null, {2, null}}. However, they have different hive
+     * hash values.
+     *
+     * The computation process for lists columns in this solution is as follows:
+     *            L1              List<list<int>>
+     *            |
+     *            L2              List<int>
+     *      /     |     \
+     *    L2[0] L2[1] L2[2]       List<int>
+     *     |            |
+     *    i1           i2         Int
+     *  /    \        /   \
+     * i1[0] i1[1]  i2[0] i2[1]   Int
+     *
+     * Note: L2、i1、i2 are all temporary columns, which would not be pushed into the stack.
+     *
+     *  Int level i1
+     * |Index|Int |
+     * |-----|----|
+     * |0    |1   |
+     * |1    |0   |
+     * length: 2
+     *
+     * Int level i2
+     * |Index|Int |
+     * |-----|----|
+     * |0    |2   |
+     * |1    |null|
+     * length: 2
+     *
+     * @tparam T The type of the column.
+     * @param col The column device view.
+     * @param row_index The index of the row to compute the hash for.
+     * @return The computed hive hash value.
+     *
+     * @note This function is only enabled for nested column types.
+     */
     template <typename T, CUDF_ENABLE_IF(cudf::is_nested<T>())>
     __device__ hive_hash_value_t operator()(cudf::column_device_view const& col,
                                             cudf::size_type row_index) const noexcept
     {
-      CUDF_UNREACHABLE("Nested type is not supported");
+      cudf::column_device_view curr_col = col.slice(row_index, 1);
+      // The default constructor of `col_stack_element` is deleted, so it can not allocate an array
+      // of `col_stack_element` directly.
+      // Instead leverage the byte array to create the col_stack_element array.
+      uint8_t stack_wrapper[MAX_NESTED_DEPTH * sizeof(col_stack_element)];
+      col_stack_element_ptr col_stack = reinterpret_cast<col_stack_element_ptr>(stack_wrapper);
+      int stack_size                  = 0;
+
+      col_stack[stack_size++] = col_stack_element(curr_col);
+
+      while (stack_size > 0) {
+        col_stack_element& element = col_stack[stack_size - 1];
+        curr_col                   = element.get_column();
+        // Do not pop it until it is processed. The definition of `processed` is:
+        // - For nested types, it is when the children are processed.
+        // - For primitive types, it is when the hash value is computed.
+        if (curr_col.type().id() == cudf::type_id::STRUCT) {
+          // All child columns processed, pop the element and update `cur_hash` of its parent column
+          if (element.cur_child_idx() == curr_col.num_child_columns()) {
+            if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(element.get_hash()); }
+          } else {
+            // Push the next child column into the stack
+            col_stack[stack_size++] =
+              col_stack_element(cudf::detail::structs_column_device_view(curr_col).get_sliced_child(
+                element.child_idx_inc_one()));
+          }
+        } else if (curr_col.type().id() == cudf::type_id::LIST) {
+          // lists_column_device_view has a different interface from structs_column_device_view
+          curr_col = cudf::detail::lists_column_device_view(curr_col).get_sliced_child();
+          if (element.cur_child_idx() == curr_col.size()) {
+            if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(element.get_hash()); }
+          } else {
+            col_stack[stack_size++] =
+              col_stack_element(curr_col.slice(element.child_idx_inc_one(), 1));
+          }
+        } else {
+          // There is only one element in the column for primitive types
+          auto hash = cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
+            curr_col.type(), this->hash_functor, curr_col, 0);
+          element.update_cur_hash(hash);
+          if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(element.get_hash()); }
+        }
+      }
+      return col_stack[0].get_hash();
     }
 
    private:
@@ -224,6 +439,35 @@ class hive_device_row_hasher {
   Nullate const _check_nulls;
   cudf::table_device_view const _table;
 };
+
+void check_nested_depth(cudf::table_view const& input, int max_nested_depth)
+{
+  using column_checker_fn_t = std::function<int(cudf::column_view const&)>;
+
+  column_checker_fn_t get_nested_depth = [&](cudf::column_view const& col) {
+    if (col.type().id() == cudf::type_id::LIST) {
+      return 1 + get_nested_depth(cudf::lists_column_view(col).child());
+    } else if (col.type().id() == cudf::type_id::STRUCT) {
+      int max_child_depth = 0;
+      for (auto child = col.child_begin(); child != col.child_end(); ++child) {
+        max_child_depth = std::max(max_child_depth, get_nested_depth(*child));
+      }
+      return 1 + max_child_depth;
+    } else {  // Primitive type
+      return 1;
+    }
+  };
+
+  for (auto i = 0; i < input.num_columns(); i++) {
+    cudf::column_view const& col = input.column(i);
+    CUDF_EXPECTS(get_nested_depth(col) <= max_nested_depth,
+                 "The " + std::to_string(i) +
+                   "-th column exceeds the maximum allowed nested depth. " +
+                   "Current depth: " + std::to_string(get_nested_depth(col)) + ", " +
+                   "Maximum allowed depth: " + std::to_string(max_nested_depth));
+  }
+}
+
 }  // namespace
 
 std::unique_ptr<cudf::column> hive_hash(cudf::table_view const& input,
@@ -239,15 +483,20 @@ std::unique_ptr<cudf::column> hive_hash(cudf::table_view const& input,
   // Return early if there's nothing to hash
   if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
 
+  // Nested depth cannot exceed 8
+  constexpr int max_nested_depth = 8;
+  check_nested_depth(input, max_nested_depth);
+
   bool const nullable   = has_nested_nulls(input);
   auto const input_view = cudf::table_device_view::create(input, stream);
   auto output_view      = output->mutable_view();
 
   // Compute the hash value for each row
-  thrust::tabulate(rmm::exec_policy(stream),
-                   output_view.begin<hive_hash_value_t>(),
-                   output_view.end<hive_hash_value_t>(),
-                   hive_device_row_hasher<hive_hash_function, bool>(nullable, *input_view));
+  thrust::tabulate(
+    rmm::exec_policy(stream),
+    output_view.begin<hive_hash_value_t>(),
+    output_view.end<hive_hash_value_t>(),
+    hive_device_row_hasher<hive_hash_function, bool, max_nested_depth>(nullable, *input_view));
 
   return output;
 }
diff --git a/src/main/java/com/nvidia/spark/rapids/jni/Hash.java b/src/main/java/com/nvidia/spark/rapids/jni/Hash.java
index a25fead0fd..16971c5bdb 100644
--- a/src/main/java/com/nvidia/spark/rapids/jni/Hash.java
+++ b/src/main/java/com/nvidia/spark/rapids/jni/Hash.java
@@ -96,7 +96,6 @@ public static ColumnVector hiveHash(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().isNestedType() : "Unsupported column type Nested";
       columnViews[i] = columns[i].getNativeView();
     }
     return new ColumnVector(hiveHash(columnViews));
diff --git a/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java b/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
index d35f20fe2c..3a396d8366 100644
--- a/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
+++ b/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
@@ -18,6 +18,7 @@
 
 import ai.rapids.cudf.ColumnVector;
 import ai.rapids.cudf.ColumnView;
+import ai.rapids.cudf.CudfException;
 import ai.rapids.cudf.DType;
 import ai.rapids.cudf.HostColumnVector.*;
 import org.junit.jupiter.api.Test;
@@ -26,6 +27,7 @@
 import java.util.Collections;
 
 import static ai.rapids.cudf.AssertUtils.*;
+import static org.junit.jupiter.api.Assertions.assertThrows;
 
 public class HashTest {
 // IEEE 754 NaN values
@@ -510,4 +512,172 @@ void testHiveHashMixed() {
       assertColumnsAreEqual(expected, result);
     }
   }
+
+  @Test
+  void testHiveHashStruct() {
+    try (ColumnVector strings = ColumnVector.fromStrings(
+          "a", "B\n", "dE\"\u0100\t\u0101 \ud720\ud721",
+          "This is a long string (greater than 128 bytes/char string) case to test this " +
+          "hash function. Just want an abnormal case here to see if any error may happen when" +
+          "doing the hive hashing",
+          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 structs = ColumnView.makeStructView(strings, integers, doubles, floats, bools);
+         ColumnVector result = Hash.hiveHash(new ColumnView[]{structs});
+         ColumnVector expected = ColumnVector.fromInts(89581538, 363542820, 413439036,
+          1272817854, 1513589666, 0)) {
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
+  @Test
+  void testHiveHashNestedStruct() {
+    try (ColumnVector strings = ColumnVector.fromStrings(
+          "a", "B\n", "dE\"\u0100\t\u0101 \ud720\ud721",
+          "This is a long string (greater than 128 bytes/char string) case to test this " +
+          "hash function. Just want an abnormal case here to see if any error may happen when" +
+          "doing the hive hashing",
+          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 result = Hash.hiveHash(new ColumnView[]{structs});
+         ColumnVector expected = ColumnVector.fromInts(89581538, 363542820, 413439036,
+          1272817854, 1513589666, 0)) {
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
+  @Test
+  void testHiveHashLists() {
+    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("This is a long string (greater than 128 bytes/char string) case to test this " +
+            "hash function. Just want an abnormal case here to see if any error may happen when" +
+            "doing the hive hashing"),
+             Collections.singletonList(""),
+             null);
+         ColumnVector stringResult = Hash.hiveHash(new ColumnView[]{stringListCV});
+         ColumnVector stringExpected = ColumnVector.fromInts(97, 63736, -96263528, 2112075710, 0, 0);
+         ColumnVector intListCV = ColumnVector.fromLists(
+             new ListType(true, new BasicType(true, DType.INT32)),
+             Collections.singletonList(null),
+             Arrays.asList(0, -2, 3),
+             Collections.singletonList(Integer.MAX_VALUE),
+             Arrays.asList(5, -6, null),
+             Collections.singletonList(Integer.MIN_VALUE),
+             null);
+         ColumnVector intExpected = ColumnVector.fromInts(0, -59, 2147483647, 4619, -2147483648, 0);
+         ColumnVector intResult = Hash.hiveHash(new ColumnVector[]{intListCV});) {
+      assertColumnsAreEqual(stringExpected, stringResult);
+      assertColumnsAreEqual(intExpected, intResult);
+    }
+  }
+
+  @Test
+  void testHiveHashNestedLists() {
+    try (ColumnVector nestedStringListCV = ColumnVector.fromLists(
+            new ListType(true, new ListType(true, new BasicType(true, DType.STRING))),
+            Arrays.asList(null, Arrays.asList("a", null)),
+            Arrays.asList(Arrays.asList("B\n", "")),
+            Arrays.asList(Collections.singletonList("dE\"\u0100\t\u0101"), Collections.singletonList(" \ud720\ud721")),
+            Arrays.asList(Collections.singletonList("This is a long string (greater than 128 bytes/char string) case to test this " +
+              "hash function. Just want an abnormal case here to see if any error may happen when" +
+              "doing the hive hashing")),
+            Arrays.asList(Collections.singletonList(""), null),
+            null);
+         ColumnVector stringExpected = ColumnVector.fromInts(3007, 63736, -96263528, 2112075710, 0, 0);
+         ColumnVector stringResult = Hash.hiveHash(new ColumnView[]{nestedStringListCV});
+         ColumnVector nestedIntListCV = ColumnVector.fromLists(
+             new ListType(true, new ListType(true, new BasicType(true, DType.INT32))),
+             Arrays.asList(Arrays.asList(null, null), null),
+             Arrays.asList(Collections.singletonList(0), Collections.singletonList(-2), Collections.singletonList(3)),
+             Arrays.asList(null, Collections.singletonList(Integer.MAX_VALUE)),
+             Arrays.asList(Collections.singletonList(5), Arrays.asList(-6, null)),
+             Arrays.asList(Collections.singletonList(Integer.MIN_VALUE), null),
+             null);
+         ColumnVector intExpected = ColumnVector.fromInts(0, -59, 2147483647, -31, -2147483648, 0);
+         ColumnVector intResult = Hash.hiveHash(new ColumnVector[]{nestedIntListCV});) {
+      assertColumnsAreEqual(stringExpected, stringResult);
+      assertColumnsAreEqual(intExpected, intResult);
+    }
+  }
+
+  @Test
+  void testHiveHashNestedType() {
+    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("This is a long string (greater than 128 bytes/char string) case to test this " +
+            "hash function. Just want an abnormal case here to see if any error may happen when" +
+            "doing the hive hashing"),
+             Collections.singletonList(""),
+             null);
+         ColumnVector intListCV = ColumnVector.fromLists(
+             new ListType(true, new BasicType(true, DType.INT32)),
+             Collections.singletonList(null),
+             Arrays.asList(0, -2, 3),
+             Collections.singletonList(Integer.MAX_VALUE),
+             Arrays.asList(5, -6, null),
+             Collections.singletonList(Integer.MIN_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);
+         ColumnView structCV = ColumnView.makeStructView(intListCV, stringListCV, doubles, floats);
+         ColumnVector nestedExpected = ColumnVector.fromInts(93217, 286968083, 59992121, -1697616301, 2127036416, 0);
+         ColumnVector nestedResult = Hash.hiveHash(new ColumnView[]{structCV})) {
+      assertColumnsAreEqual(nestedExpected, nestedResult);
+    }
+  }
+
+  @Test
+  void testHiveHashNestedDepthExceedsLimit() {
+    try (ColumnVector nestedIntListCV = ColumnVector.fromLists(
+            new ListType(true, new ListType(true, new BasicType(true, DType.INT32))),
+            Arrays.asList(Arrays.asList(null, null), null),
+            Arrays.asList(Collections.singletonList(0), Collections.singletonList(-2), Collections.singletonList(3)),
+            Arrays.asList(null, Collections.singletonList(Integer.MAX_VALUE)),
+            Arrays.asList(Collections.singletonList(5), Arrays.asList(-6, null)),
+            Arrays.asList(Collections.singletonList(Integer.MIN_VALUE), 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(nestedIntListCV, integers);
+         ColumnView structs2 = ColumnView.makeStructView(structs1, doubles);
+         ColumnView structs3 = ColumnView.makeStructView(structs2, bools);
+         ColumnView structs4 = ColumnView.makeStructView(structs3);
+         ColumnView structs5 = ColumnView.makeStructView(structs4, floats);
+         ColumnView nestedResult = ColumnView.makeStructView(structs5);) {
+      assertThrows(CudfException.class, () -> Hash.hiveHash(new ColumnView[]{nestedResult}));
+    }
+  }
 }

From a70742536615e606f442d544feaa60d0191eac05 Mon Sep 17 00:00:00 2001
From: ustcfy <fengyan_@mail.ustc.edu.cn>
Date: Thu, 7 Nov 2024 12:53:50 +0800
Subject: [PATCH 02/25] Fix minor details

Signed-off-by: ustcfy <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu                 | 28 +++++++++---------
 .../com/nvidia/spark/rapids/jni/HashTest.java | 29 +++++++++++++++----
 2 files changed, 37 insertions(+), 20 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 71b25f1b5e..d26acae42b 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -240,16 +240,16 @@ class hive_device_row_hasher {
 
       __device__ void update_cur_hash(hive_hash_value_t child_hash)
       {
-        this->cur_hash = this->cur_hash * HIVE_HASH_FACTOR + child_hash;
+        cur_hash = cur_hash * HIVE_HASH_FACTOR + child_hash;
       }
 
-      __device__ hive_hash_value_t get_hash() { return this->cur_hash; }
+      __device__ hive_hash_value_t get_hash() { return cur_hash; }
 
-      __device__ int child_idx_inc_one() { return this->child_idx++; }
+      __device__ int get_and_inc_child_idx() { return child_idx++; }
 
-      __device__ int cur_child_idx() { return this->child_idx; }
+      __device__ int cur_child_idx() { return child_idx; }
 
-      __device__ cudf::column_device_view get_column() { return this->column; }
+      __device__ cudf::column_device_view get_column() { return column; }
     };
 
     typedef col_stack_element* col_stack_element_ptr;
@@ -306,7 +306,8 @@ class hive_device_row_hasher {
      * separately.
      *
      * For example, consider the following nested column: `List<List<int>>`
-     * lists_column = {{1, 0}, null, {2, null}}
+     * lists_column = [ [[1, 0], null, [2, null]] ], which has only one item
+     * [[1, 0], null, [2, null]
      *
      *     L1
      *     |
@@ -317,16 +318,16 @@ class hive_device_row_hasher {
      * List level L1:
      * |Index|      List<list<int>>    |
      * |-----|-------------------------|
-     * |0    |{{1, 0}, null, {2, null}}|
+     * |0    |[[1, 0], null, [2, null]]|
      * length: 1
      * Offsets: 0, 3
      *
      * List level L2:
      * |Index|List<int>|
      * |-----|---------|
-     * |0    |{1, 0}   |
+     * |0    |[1, 0]   |
      * |1    |null     |
-     * |2    |{2, null}|
+     * |2    |[2, null]|
      * length: 3
      * Offsets: 0, 2, 2, 4
      * null_mask: 101
@@ -344,9 +345,8 @@ class hive_device_row_hasher {
      * It will produce different results than Spark to compute the hash value using the
      * underlying data merely.
      *
-     * For example, `List<List<int>>` column {{1, 0}, {2, null}} has the same underlying data as the
-     * above `List<List<int>>` column {{1, 0}, null, {2, null}}. However, they have different hive
-     * hash values.
+     * For example, [[1, 0], [2, null]] has the same underlying data as [[1, 0], null, [2, null]].
+     * However, they have different hive hash values.
      *
      * The computation process for lists columns in this solution is as follows:
      *            L1              List<list<int>>
@@ -410,7 +410,7 @@ class hive_device_row_hasher {
             // Push the next child column into the stack
             col_stack[stack_size++] =
               col_stack_element(cudf::detail::structs_column_device_view(curr_col).get_sliced_child(
-                element.child_idx_inc_one()));
+                element.get_and_inc_child_idx()));
           }
         } else if (curr_col.type().id() == cudf::type_id::LIST) {
           // lists_column_device_view has a different interface from structs_column_device_view
@@ -419,7 +419,7 @@ class hive_device_row_hasher {
             if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(element.get_hash()); }
           } else {
             col_stack[stack_size++] =
-              col_stack_element(curr_col.slice(element.child_idx_inc_one(), 1));
+              col_stack_element(curr_col.slice(element.get_and_inc_child_idx(), 1));
           }
         } else {
           // There is only one element in the column for primitive types
diff --git a/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java b/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
index 3a396d8366..20fba6f383 100644
--- a/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
+++ b/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
@@ -572,15 +572,15 @@ void testHiveHashLists() {
              Arrays.asList("B\n", ""),
              Arrays.asList("dE\"\u0100\t\u0101", " \ud720\ud721"),
              Collections.singletonList("This is a long string (greater than 128 bytes/char string) case to test this " +
-            "hash function. Just want an abnormal case here to see if any error may happen when" +
-            "doing the hive hashing"),
+             "hash function. Just want an abnormal case here to see if any error may happen when" +
+             "doing the hive hashing"),
              Collections.singletonList(""),
              null);
          ColumnVector stringResult = Hash.hiveHash(new ColumnView[]{stringListCV});
          ColumnVector stringExpected = ColumnVector.fromInts(97, 63736, -96263528, 2112075710, 0, 0);
          ColumnVector intListCV = ColumnVector.fromLists(
              new ListType(true, new BasicType(true, DType.INT32)),
-             Collections.singletonList(null),
+             Collections.emptyList(),
              Arrays.asList(0, -2, 3),
              Collections.singletonList(Integer.MAX_VALUE),
              Arrays.asList(5, -6, null),
@@ -623,15 +623,15 @@ void testHiveHashNestedLists() {
   }
 
   @Test
-  void testHiveHashNestedType() {
+  void testHiveHashStructOfList() {
     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("This is a long string (greater than 128 bytes/char string) case to test this " +
-            "hash function. Just want an abnormal case here to see if any error may happen when" +
-            "doing the hive hashing"),
+             "hash function. Just want an abnormal case here to see if any error may happen when" +
+             "doing the hive hashing"),
              Collections.singletonList(""),
              null);
          ColumnVector intListCV = ColumnVector.fromLists(
@@ -653,6 +653,23 @@ void testHiveHashNestedType() {
     }
   }
 
+  @Test
+  void testHiveHashListOfStruct() {
+    try (ColumnVector structListCV = ColumnVector.fromLists(new ListType(true, new StructType(true,
+              new BasicType(true, DType.STRING), new BasicType(true, DType.INT32), new BasicType(true, DType.FLOAT64), new BasicType(true, DType.FLOAT32), new BasicType(true, DType.BOOL8))),
+             Collections.emptyList(),
+             Collections.singletonList(new StructData("a", 0, 0.0, 0f, true)),
+             Arrays.asList(new StructData("B\n", 100, 100.0, 100f, false), new StructData("dE\"\u0100\t\u0101 \ud720\ud721", -100, -100.0, -100f, null)),
+             Collections.singletonList(new StructData("This is a long string (greater than 128 bytes/char string) case to test this " +
+             "hash function. Just want an abnormal case here to see if any error may happen when" + "doing the hive hashing", Integer.MIN_VALUE, POSITIVE_DOUBLE_NAN_LOWER_RANGE, NEGATIVE_FLOAT_NAN_LOWER_RANGE, false)),
+             Arrays.asList(new StructData(null, Integer.MAX_VALUE, POSITIVE_DOUBLE_NAN_UPPER_RANGE, NEGATIVE_FLOAT_NAN_UPPER_RANGE, true), new StructData(null, null, null, null, null)),
+             null);
+         ColumnVector result = Hash.hiveHash(new ColumnView[]{structListCV});
+         ColumnVector expected = ColumnVector.fromInts(0, 89581538, -1201635432, 1272817854, -323360610, 0)) {
+      assertColumnsAreEqual(expected, result);
+    }
+  }
+
   @Test
   void testHiveHashNestedDepthExceedsLimit() {
     try (ColumnVector nestedIntListCV = ColumnVector.fromLists(

From 215d0fba80eeb4250cf6ef481fcb06525f5712c0 Mon Sep 17 00:00:00 2001
From: ustcfy <fengyan_@mail.ustc.edu.cn>
Date: Fri, 8 Nov 2024 00:37:48 +0800
Subject: [PATCH 03/25] Rename class member variables

Signed-off-by: ustcfy <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu | 24 ++++++++++++------------
 1 file changed, 12 insertions(+), 12 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index d26acae42b..810f84d15a 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -214,9 +214,9 @@ class hive_device_row_hasher {
      * This structure is used to keep track of the current column being processed, the index of the
      * next child column to process, and current hash value for the column.
      *
-     * @param column The current column being processed.
-     * @param child_idx The index of the next child column to process.
-     * @param cur_hash The current hash value for the column.
+     * @param _column The current column being processed.
+     * @param _child_idx The index of the next child column to process.
+     * @param _cur_hash The current hash value for the column.
      *
      * @note The default constructor is deleted to prevent uninitialized usage.
      *
@@ -225,31 +225,31 @@ class hive_device_row_hasher {
      */
     class col_stack_element {
      private:
-      cudf::column_device_view column;  // current column
-      hive_hash_value_t cur_hash;       // current hash value of the column
-      int child_idx;  // index of the child column to process next, initialized as 0
+      cudf::column_device_view _column;  // current column
+      hive_hash_value_t _cur_hash;       // current hash value of the column
+      int _child_idx;  // index of the child column to process next, initialized as 0
 
      public:
       __device__ col_stack_element() =
         delete;  // Because the default constructor of `cudf::column_device_view` is deleted
 
       __device__ col_stack_element(cudf::column_device_view col)
-        : column(col), child_idx(0), cur_hash(HIVE_INIT_HASH)
+        : _column(col), _child_idx(0), _cur_hash(HIVE_INIT_HASH)
       {
       }
 
       __device__ void update_cur_hash(hive_hash_value_t child_hash)
       {
-        cur_hash = cur_hash * HIVE_HASH_FACTOR + child_hash;
+        _cur_hash = _cur_hash * HIVE_HASH_FACTOR + child_hash;
       }
 
-      __device__ hive_hash_value_t get_hash() { return cur_hash; }
+      __device__ hive_hash_value_t get_hash() { return _cur_hash; }
 
-      __device__ int get_and_inc_child_idx() { return child_idx++; }
+      __device__ int get_and_inc_child_idx() { return _child_idx++; }
 
-      __device__ int cur_child_idx() { return child_idx; }
+      __device__ int cur_child_idx() { return _child_idx; }
 
-      __device__ cudf::column_device_view get_column() { return column; }
+      __device__ cudf::column_device_view get_column() { return _column; }
     };
 
     typedef col_stack_element* col_stack_element_ptr;

From 8afa2a1e109f1026eeaadbe7c5c9c4b25b9bd0b4 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 8 Nov 2024 19:35:11 +0800
Subject: [PATCH 04/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 8 --------
 1 file changed, 8 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 810f84d15a..95e5e38c76 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -214,14 +214,6 @@ class hive_device_row_hasher {
      * This structure is used to keep track of the current column being processed, the index of the
      * next child column to process, and current hash value for the column.
      *
-     * @param _column The current column being processed.
-     * @param _child_idx The index of the next child column to process.
-     * @param _cur_hash The current hash value for the column.
-     *
-     * @note The default constructor is deleted to prevent uninitialized usage.
-     *
-     * @constructor
-     * @param col The column device view to be processed.
      */
     class col_stack_element {
      private:

From 5404d5f66cec7aaa58aaa8706fc5abb7d1d37b3e Mon Sep 17 00:00:00 2001
From: Yan Feng <fengyan_@mail.ustc.edu.cn>
Date: Wed, 13 Nov 2024 15:46:55 +0800
Subject: [PATCH 05/25] Fix stack_wrapper alignment

Signed-off-by: Yan Feng <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 95e5e38c76..19d0040538 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -213,9 +213,8 @@ class hive_device_row_hasher {
      *
      * This structure is used to keep track of the current column being processed, the index of the
      * next child column to process, and current hash value for the column.
-     *
      */
-    class col_stack_element {
+    struct col_stack_element {
      private:
       cudf::column_device_view _column;  // current column
       hive_hash_value_t _cur_hash;       // current hash value of the column
@@ -382,7 +381,7 @@ class hive_device_row_hasher {
       // The default constructor of `col_stack_element` is deleted, so it can not allocate an array
       // of `col_stack_element` directly.
       // Instead leverage the byte array to create the col_stack_element array.
-      uint8_t stack_wrapper[MAX_NESTED_DEPTH * sizeof(col_stack_element)];
+      alignas(col_stack_element) char stack_wrapper[sizeof(col_stack_element) * MAX_NESTED_DEPTH];
       col_stack_element_ptr col_stack = reinterpret_cast<col_stack_element_ptr>(stack_wrapper);
       int stack_size                  = 0;
 

From 058fd47d7dff2241bee7e2adc2f958677feec07c Mon Sep 17 00:00:00 2001
From: Yan Feng <fengyan_@mail.ustc.edu.cn>
Date: Mon, 18 Nov 2024 11:06:13 +0800
Subject: [PATCH 06/25] Refactor variable names for better understanding

Signed-off-by: Yan Feng <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu | 105 +++++++++++++++-------------------
 1 file changed, 45 insertions(+), 60 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 19d0040538..a9b94c4bf2 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -209,47 +209,44 @@ class hive_device_row_hasher {
     }
 
     /**
-     * @brief A structure representing an element in the stack used for processing columns.
-     *
-     * This structure is used to keep track of the current column being processed, the index of the
-     * next child column to process, and current hash value for the column.
+     * @brief A structure to keep track of the computation for nested types.
      */
-    struct col_stack_element {
+    struct col_stack_frame {
      private:
-      cudf::column_device_view _column;  // current column
-      hive_hash_value_t _cur_hash;       // current hash value of the column
-      int _child_idx;  // index of the child column to process next, initialized as 0
+      cudf::column_device_view _element;  // the column has only one row
+      hive_hash_value_t _cur_hash;        // current hash value of the column
+      int _idx_to_process;                // the index of child or element to process next
 
      public:
-      __device__ col_stack_element() =
+      __device__ col_stack_frame() =
         delete;  // Because the default constructor of `cudf::column_device_view` is deleted
 
-      __device__ col_stack_element(cudf::column_device_view col)
-        : _column(col), _child_idx(0), _cur_hash(HIVE_INIT_HASH)
+      __device__ col_stack_frame(cudf::column_device_view element)
+        : _element(element), _idx_to_process(0), _cur_hash(HIVE_INIT_HASH)
       {
       }
 
-      __device__ void update_cur_hash(hive_hash_value_t child_hash)
+      __device__ void update_cur_hash(hive_hash_value_t hash)
       {
-        _cur_hash = _cur_hash * HIVE_HASH_FACTOR + child_hash;
+        _cur_hash = _cur_hash * HIVE_HASH_FACTOR + hash;
       }
 
       __device__ hive_hash_value_t get_hash() { return _cur_hash; }
 
-      __device__ int get_and_inc_child_idx() { return _child_idx++; }
+      __device__ int get_and_inc_idx_to_process() { return _idx_to_process++; }
 
-      __device__ int cur_child_idx() { return _child_idx; }
+      __device__ int get_idx_to_process() { return _idx_to_process; }
 
-      __device__ cudf::column_device_view get_column() { return _column; }
+      __device__ cudf::column_device_view get_element() { return _element; }
     };
 
-    typedef col_stack_element* col_stack_element_ptr;
+    typedef col_stack_frame* col_stack_frame_ptr;
 
     /**
      * @brief Functor to compute the hive hash value for a nested column.
      *
-     * This functor produces the same result as "HiveHash" in Spark for nested types.
-     * The pseudocode of Spark's HiveHash function is as follows:
+     * This functor produces the same result as "HiveHash" in Spark for structs and lists.
+     * The pseudocode of Spark's HiveHash function for structs and lists is as follows:
      *
      * hive_hash_value_t hive_hash(NestedType element) {
      *    hive_hash_value_t hash = HIVE_INIT_HASH;
@@ -332,12 +329,13 @@ class hive_device_row_hasher {
      * |3    |null|
      * length: 4
      *
-     * Since the underlying data loses the null information of the top-level list column,
-     * It will produce different results than Spark to compute the hash value using the
-     * underlying data merely.
-     *
+     * Since the underlying data loses the null information of the top-level list column, computing
+     * hash values using the underlying data merely can yield different results compared to Spark.
      * For example, [[1, 0], [2, null]] has the same underlying data as [[1, 0], null, [2, null]].
-     * However, they have different hive hash values.
+     * However, they have different hive hash values, as null values affect the hash value.
+     *
+     * And sublist structure factors in the hash value. For example, [[1, 0], [2]] and [[1], [0, 2]]
+     * have different hive hash values.
      *
      * The computation process for lists columns in this solution is as follows:
      *            L1              List<list<int>>
@@ -352,20 +350,6 @@ class hive_device_row_hasher {
      *
      * Note: L2、i1、i2 are all temporary columns, which would not be pushed into the stack.
      *
-     *  Int level i1
-     * |Index|Int |
-     * |-----|----|
-     * |0    |1   |
-     * |1    |0   |
-     * length: 2
-     *
-     * Int level i2
-     * |Index|Int |
-     * |-----|----|
-     * |0    |2   |
-     * |1    |null|
-     * length: 2
-     *
      * @tparam T The type of the column.
      * @param col The column device view.
      * @param row_index The index of the row to compute the hash for.
@@ -378,46 +362,47 @@ class hive_device_row_hasher {
                                             cudf::size_type row_index) const noexcept
     {
       cudf::column_device_view curr_col = col.slice(row_index, 1);
-      // The default constructor of `col_stack_element` is deleted, so it can not allocate an array
-      // of `col_stack_element` directly.
-      // Instead leverage the byte array to create the col_stack_element array.
-      alignas(col_stack_element) char stack_wrapper[sizeof(col_stack_element) * MAX_NESTED_DEPTH];
-      col_stack_element_ptr col_stack = reinterpret_cast<col_stack_element_ptr>(stack_wrapper);
-      int stack_size                  = 0;
+      // The default constructor of `col_stack_frame` is deleted, so it can not allocate an array
+      // of `col_stack_frame` directly.
+      // Instead leverage the byte array to create the col_stack_frame array.
+      alignas(col_stack_frame) char stack_wrapper[sizeof(col_stack_frame) * MAX_NESTED_DEPTH];
+      col_stack_frame_ptr col_stack = reinterpret_cast<col_stack_frame_ptr>(stack_wrapper);
+      int stack_size                = 0;
 
-      col_stack[stack_size++] = col_stack_element(curr_col);
+      col_stack[stack_size++] = col_stack_frame(curr_col);
 
       while (stack_size > 0) {
-        col_stack_element& element = col_stack[stack_size - 1];
-        curr_col                   = element.get_column();
+        col_stack_frame& top = col_stack[stack_size - 1];
+        curr_col             = top.get_element();
         // Do not pop it until it is processed. The definition of `processed` is:
-        // - For nested types, it is when the children are processed.
+        // - For structs, it is when all child columns are processed.
+        // - For lists, it is when all elements in the list are processed.
         // - For primitive types, it is when the hash value is computed.
         if (curr_col.type().id() == cudf::type_id::STRUCT) {
-          // All child columns processed, pop the element and update `cur_hash` of its parent column
-          if (element.cur_child_idx() == curr_col.num_child_columns()) {
-            if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(element.get_hash()); }
+          if (top.get_idx_to_process() == curr_col.num_child_columns()) {
+            if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(top.get_hash()); }
           } else {
-            // Push the next child column into the stack
+            // Push the next child into the stack
             col_stack[stack_size++] =
-              col_stack_element(cudf::detail::structs_column_device_view(curr_col).get_sliced_child(
-                element.get_and_inc_child_idx()));
+              col_stack_frame(cudf::detail::structs_column_device_view(curr_col).get_sliced_child(
+                top.get_and_inc_idx_to_process()));
           }
         } else if (curr_col.type().id() == cudf::type_id::LIST) {
-          // lists_column_device_view has a different interface from structs_column_device_view
+          // Get the child column of the list column
           curr_col = cudf::detail::lists_column_device_view(curr_col).get_sliced_child();
-          if (element.cur_child_idx() == curr_col.size()) {
-            if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(element.get_hash()); }
+          if (top.get_idx_to_process() == curr_col.size()) {
+            if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(top.get_hash()); }
           } else {
+            // Push the next element into the stack
             col_stack[stack_size++] =
-              col_stack_element(curr_col.slice(element.get_and_inc_child_idx(), 1));
+              col_stack_frame(curr_col.slice(top.get_and_inc_idx_to_process(), 1));
           }
         } else {
           // There is only one element in the column for primitive types
           auto hash = cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
             curr_col.type(), this->hash_functor, curr_col, 0);
-          element.update_cur_hash(hash);
-          if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(element.get_hash()); }
+          top.update_cur_hash(hash);
+          if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(top.get_hash()); }
         }
       }
       return col_stack[0].get_hash();

From e7749c8a0a59a1de7805166c671b5b64efa67573 Mon Sep 17 00:00:00 2001
From: Yan Feng <fengyan_@mail.ustc.edu.cn>
Date: Tue, 19 Nov 2024 13:04:00 +0800
Subject: [PATCH 07/25] Optimize computation for single-level lists

Signed-off-by: Yan Feng <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu                 | 39 +++++++++++++++----
 .../com/nvidia/spark/rapids/jni/HashTest.java |  3 +-
 2 files changed, 33 insertions(+), 9 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index a9b94c4bf2..d99f84d971 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -345,10 +345,11 @@ class hive_device_row_hasher {
      *    L2[0] L2[1] L2[2]       List<int>
      *     |            |
      *    i1           i2         Int
-     *  /    \        /   \
-     * i1[0] i1[1]  i2[0] i2[1]   Int
      *
      * Note: L2、i1、i2 are all temporary columns, which would not be pushed into the stack.
+     * There is an optimization for the list column. If the child column is of primitive type, the
+     * hash value of the list column can be directly computed. Thus we can decrease the depth of
+     * the stack by one.
      *
      * @tparam T The type of the column.
      * @param col The column device view.
@@ -389,13 +390,30 @@ class hive_device_row_hasher {
           }
         } else if (curr_col.type().id() == cudf::type_id::LIST) {
           // Get the child column of the list column
-          curr_col = cudf::detail::lists_column_device_view(curr_col).get_sliced_child();
-          if (top.get_idx_to_process() == curr_col.size()) {
+          cudf::column_device_view child_col =
+            cudf::detail::lists_column_device_view(curr_col).get_sliced_child();
+          // If the child column is of primitive type, directly compute the hash value of the list
+          if (child_col.type().id() != cudf::type_id::LIST &&
+              child_col.type().id() != cudf::type_id::STRUCT) {
+            auto single_level_list_hash = cudf::detail::accumulate(
+              thrust::counting_iterator(0),
+              thrust::counting_iterator(child_col.size()),
+              HIVE_INIT_HASH,
+              [child_col, hasher = this->hash_functor] __device__(auto hash, auto element_index) {
+                auto cur_hash = cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
+                  child_col.type(), hasher, child_col, element_index);
+                return HIVE_HASH_FACTOR * hash + cur_hash;
+              });
+            top.update_cur_hash(single_level_list_hash);
             if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(top.get_hash()); }
           } else {
-            // Push the next element into the stack
-            col_stack[stack_size++] =
-              col_stack_frame(curr_col.slice(top.get_and_inc_idx_to_process(), 1));
+            if (top.get_idx_to_process() == child_col.size()) {
+              if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(top.get_hash()); }
+            } else {
+              // Push the next element into the stack
+              col_stack[stack_size++] =
+                col_stack_frame(child_col.slice(top.get_and_inc_idx_to_process(), 1));
+            }
           }
         } else {
           // There is only one element in the column for primitive types
@@ -422,7 +440,12 @@ void check_nested_depth(cudf::table_view const& input, int max_nested_depth)
 
   column_checker_fn_t get_nested_depth = [&](cudf::column_view const& col) {
     if (col.type().id() == cudf::type_id::LIST) {
-      return 1 + get_nested_depth(cudf::lists_column_view(col).child());
+      auto const& child_col = cudf::lists_column_view(col).child();
+      if (child_col.type().id() != cudf::type_id::STRUCT &&
+          child_col.type().id() != cudf::type_id::LIST) {
+        return 1;
+      }
+      return 1 + get_nested_depth(child_col);
     } else if (col.type().id() == cudf::type_id::STRUCT) {
       int max_child_depth = 0;
       for (auto child = col.child_begin(); child != col.child_end(); ++child) {
diff --git a/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java b/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
index 20fba6f383..d3ef6b5c54 100644
--- a/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
+++ b/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
@@ -693,7 +693,8 @@ void testHiveHashNestedDepthExceedsLimit() {
          ColumnView structs3 = ColumnView.makeStructView(structs2, bools);
          ColumnView structs4 = ColumnView.makeStructView(structs3);
          ColumnView structs5 = ColumnView.makeStructView(structs4, floats);
-         ColumnView nestedResult = ColumnView.makeStructView(structs5);) {
+         ColumnView structs6 = ColumnView.makeStructView(structs5);
+         ColumnView nestedResult = ColumnView.makeStructView(structs6);) {
       assertThrows(CudfException.class, () -> Hash.hiveHash(new ColumnView[]{nestedResult}));
     }
   }

From e50be83e27dfdcd5e060db5b9f9af9345b140e44 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 09:50:54 +0800
Subject: [PATCH 08/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index d99f84d971..8a40b01d1b 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -246,7 +246,7 @@ class hive_device_row_hasher {
      * @brief Functor to compute the hive hash value for a nested column.
      *
      * This functor produces the same result as "HiveHash" in Spark for structs and lists.
-     * The pseudocode of Spark's HiveHash function for structs and lists is as follows:
+     * The pseudocode of Spark's HiveHash function for structs is as follows:
      *
      * hive_hash_value_t hive_hash(NestedType element) {
      *    hive_hash_value_t hash = HIVE_INIT_HASH;
@@ -256,6 +256,9 @@ class hive_device_row_hasher {
      *    return hash;
      * }
      *
+     * In the cases of lists, the hash value is computed by a similar way but we iterate through the list 
+     * elements instead of through the child columns' elements.
+     *
      * This functor uses a stack to simulate the recursive process of the above pseudocode.
      * When an element is popped from the stack, it means that the hash value of it has been
      * computed. Therefore, we should update the parent's `cur_hash` upon popping the element.

From 00d27a7191c0e326c67b197fcf8ab03dc415736e Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 09:51:22 +0800
Subject: [PATCH 09/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 8a40b01d1b..62780601a4 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -297,8 +297,7 @@ class hive_device_row_hasher {
      * separately.
      *
      * For example, consider the following nested column: `List<List<int>>`
-     * lists_column = [ [[1, 0], null, [2, null]] ], which has only one item
-     * [[1, 0], null, [2, null]
+     * [[1, 0], null, [2, null]]
      *
      *     L1
      *     |

From d2f536c97cb931dc78961679076f373079721407 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 10:25:08 +0800
Subject: [PATCH 10/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 62780601a4..dd569e61ea 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -268,11 +268,12 @@ class hive_device_row_hasher {
      * 1. Initialize the stack and push the root column into the stack.
      * 2. While the stack is not empty:
      *    a. Get the top element of the stack. Don't pop it until it is processed.
-     *    b. If the column is a nested column:
+     *    b. If the column is a structs column:
      *        i.  If all child columns are processed, pop the element and update `cur_hash` of its
      *            parent column.
      *        ii. Otherwise, push the next child column into the stack.
-     *    c. If the column is a primitive column, compute the hash value, pop the element,
+     *    c. If the column is a lists column, process it by a similar way as structs column but iterating through the list elements instead of child columns' elements.
+     *    d. If the column is a primitive column, compute the hash value, pop the element,
      *       and update `cur_hash` of its parent column.
      * 3. Return the hash value of the root column.
      *

From 3b210a2e3ac6657d70f7b71f0139bf8be3137717 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 10:25:24 +0800
Subject: [PATCH 11/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index dd569e61ea..a76f441aec 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -345,9 +345,9 @@ class hive_device_row_hasher {
      *            |
      *            L2              List<int>
      *      /     |     \
-     *    L2[0] L2[1] L2[2]       List<int>
+     *    L2[0] L2[1] L2[2]       int
      *     |            |
-     *    i1           i2         Int
+     *    i1           i2         int
      *
      * Note: L2、i1、i2 are all temporary columns, which would not be pushed into the stack.
      * There is an optimization for the list column. If the child column is of primitive type, the

From 6653066e2f055e6b6f66286cb3cf094907c79b58 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 10:28:08 +0800
Subject: [PATCH 12/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index a76f441aec..1c3da352d0 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -370,7 +370,7 @@ class hive_device_row_hasher {
       // of `col_stack_frame` directly.
       // Instead leverage the byte array to create the col_stack_frame array.
       alignas(col_stack_frame) char stack_wrapper[sizeof(col_stack_frame) * MAX_NESTED_DEPTH];
-      col_stack_frame_ptr col_stack = reinterpret_cast<col_stack_frame_ptr>(stack_wrapper);
+      auto col_stack = reinterpret_cast<col_stack_frame*>(stack_wrapper);
       int stack_size                = 0;
 
       col_stack[stack_size++] = col_stack_frame(curr_col);

From ae903f0e003ea19a62344a42add499bfbf7c5b90 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 10:28:57 +0800
Subject: [PATCH 13/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 1 -
 1 file changed, 1 deletion(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 1c3da352d0..5f516bcefb 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -240,7 +240,6 @@ class hive_device_row_hasher {
       __device__ cudf::column_device_view get_element() { return _element; }
     };
 
-    typedef col_stack_frame* col_stack_frame_ptr;
 
     /**
      * @brief Functor to compute the hive hash value for a nested column.

From eb63abcf061279531d7d44fcc98fc6e209702594 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 14:20:46 +0800
Subject: [PATCH 14/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 5f516bcefb..70d801544d 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -484,8 +484,7 @@ std::unique_ptr<cudf::column> hive_hash(cudf::table_view const& input,
   // Return early if there's nothing to hash
   if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
 
-  // Nested depth cannot exceed 8
-  constexpr int max_nested_depth = 8;
+    constexpr int  MAX_NESTED_DEPTH = 8;
   check_nested_depth(input, max_nested_depth);
 
   bool const nullable   = has_nested_nulls(input);

From e20792e6693176670878f10e46d28c75b71849e2 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 14:20:56 +0800
Subject: [PATCH 15/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 1 +
 1 file changed, 1 insertion(+)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 70d801544d..46ac79af4f 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -330,6 +330,7 @@ class hive_device_row_hasher {
      * |2    |2   |
      * |3    |null|
      * length: 4
+     * null_mask: 0111
      *
      * Since the underlying data loses the null information of the top-level list column, computing
      * hash values using the underlying data merely can yield different results compared to Spark.

From 994ec3404805901fd8d0cb7f5a4cef8dd6ecf426 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 14:21:45 +0800
Subject: [PATCH 16/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 46ac79af4f..1d929700f5 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -213,7 +213,7 @@ class hive_device_row_hasher {
      */
     struct col_stack_frame {
      private:
-      cudf::column_device_view _element;  // the column has only one row
+      cudf::column_device_view _column;
       hive_hash_value_t _cur_hash;        // current hash value of the column
       int _idx_to_process;                // the index of child or element to process next
 

From 90a4aaead5ba1718d21116181e904366e257e166 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Fri, 22 Nov 2024 14:21:52 +0800
Subject: [PATCH 17/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 1d929700f5..a80281560c 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -237,7 +237,7 @@ class hive_device_row_hasher {
 
       __device__ int get_idx_to_process() { return _idx_to_process; }
 
-      __device__ cudf::column_device_view get_element() { return _element; }
+      __device__ cudf::column_device_view get_column() { return _column; }
     };
 
 

From ccdd64b32e78f2d297897e94290ee6c4aa94be77 Mon Sep 17 00:00:00 2001
From: Yan Feng <fengyan_@mail.ustc.edu.cn>
Date: Fri, 22 Nov 2024 17:46:59 +0800
Subject: [PATCH 18/25] Fix minor details

Signed-off-by: Yan Feng <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu | 41 +++++++++++++++++++----------------
 1 file changed, 22 insertions(+), 19 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index a80281560c..4e6cfad107 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -37,6 +37,8 @@ using hive_hash_value_t = int32_t;
 constexpr hive_hash_value_t HIVE_HASH_FACTOR = 31;
 constexpr hive_hash_value_t HIVE_INIT_HASH   = 0;
 
+constexpr int  MAX_NESTED_DEPTH = 8;
+
 hive_hash_value_t __device__ inline compute_int(int32_t key) { return key; }
 
 hive_hash_value_t __device__ inline compute_long(int64_t key)
@@ -157,7 +159,7 @@ hive_hash_value_t __device__ inline hive_hash_function<cudf::timestamp_us>::oper
  * @tparam hash_function Hash functor to use for hashing elements. Must be hive_hash_function.
  * @tparam Nullate A cudf::nullate type describing whether to check for nulls.
  */
-template <template <typename> class hash_function, typename Nullate, int MAX_NESTED_DEPTH>
+template <template <typename> class hash_function, typename Nullate>
 class hive_device_row_hasher {
  public:
   CUDF_HOST_DEVICE hive_device_row_hasher(Nullate check_nulls, cudf::table_device_view t) noexcept
@@ -221,8 +223,8 @@ class hive_device_row_hasher {
       __device__ col_stack_frame() =
         delete;  // Because the default constructor of `cudf::column_device_view` is deleted
 
-      __device__ col_stack_frame(cudf::column_device_view element)
-        : _element(element), _idx_to_process(0), _cur_hash(HIVE_INIT_HASH)
+      __device__ col_stack_frame(cudf::column_device_view col)
+        : _column(col), _idx_to_process(0), _cur_hash(HIVE_INIT_HASH)
       {
       }
 
@@ -240,7 +242,6 @@ class hive_device_row_hasher {
       __device__ cudf::column_device_view get_column() { return _column; }
     };
 
-
     /**
      * @brief Functor to compute the hive hash value for a nested column.
      *
@@ -255,8 +256,8 @@ class hive_device_row_hasher {
      *    return hash;
      * }
      *
-     * In the cases of lists, the hash value is computed by a similar way but we iterate through the list 
-     * elements instead of through the child columns' elements.
+     * In the cases of lists, the hash value is computed by a similar way but we iterate through the
+     * list elements instead of through the child columns' elements.
      *
      * This functor uses a stack to simulate the recursive process of the above pseudocode.
      * When an element is popped from the stack, it means that the hash value of it has been
@@ -271,7 +272,8 @@ class hive_device_row_hasher {
      *        i.  If all child columns are processed, pop the element and update `cur_hash` of its
      *            parent column.
      *        ii. Otherwise, push the next child column into the stack.
-     *    c. If the column is a lists column, process it by a similar way as structs column but iterating through the list elements instead of child columns' elements.
+     *    c. If the column is a lists column, process it by a similar way as structs column but
+     *       iterating through the list elements instead of child columns' elements.
      *    d. If the column is a primitive column, compute the hash value, pop the element,
      *       and update `cur_hash` of its parent column.
      * 3. Return the hash value of the root column.
@@ -296,8 +298,10 @@ class hive_device_row_hasher {
      * As lists columns have a different interface from structs columns, we need to handle them
      * separately.
      *
-     * For example, consider the following nested column: `List<List<int>>`
-     * [[1, 0], null, [2, null]]
+     * For example, consider that the input column is of type `List<List<int>>`.
+     * Assume that the element at `row_index` is: [[1, 0], null, [2, null]]
+     * Since the stack_frame should contain a column that consists of only one row, the input column
+     * should be sliced. The sliced column is noted as L1 here.
      *
      *     L1
      *     |
@@ -323,14 +327,14 @@ class hive_device_row_hasher {
      * null_mask: 101
      *
      * Int level i:
-     * |Index|Int |
+     * |Index|int |
      * |-----|----|
      * |0    |1   |
      * |1    |0   |
      * |2    |2   |
      * |3    |null|
      * length: 4
-     * null_mask: 0111
+     * null_mask: 1110
      *
      * Since the underlying data loses the null information of the top-level list column, computing
      * hash values using the underlying data merely can yield different results compared to Spark.
@@ -371,13 +375,13 @@ class hive_device_row_hasher {
       // Instead leverage the byte array to create the col_stack_frame array.
       alignas(col_stack_frame) char stack_wrapper[sizeof(col_stack_frame) * MAX_NESTED_DEPTH];
       auto col_stack = reinterpret_cast<col_stack_frame*>(stack_wrapper);
-      int stack_size                = 0;
+      int stack_size = 0;
 
       col_stack[stack_size++] = col_stack_frame(curr_col);
 
       while (stack_size > 0) {
         col_stack_frame& top = col_stack[stack_size - 1];
-        curr_col             = top.get_element();
+        curr_col             = top.get_column();
         // Do not pop it until it is processed. The definition of `processed` is:
         // - For structs, it is when all child columns are processed.
         // - For lists, it is when all elements in the list are processed.
@@ -437,7 +441,7 @@ class hive_device_row_hasher {
   cudf::table_device_view const _table;
 };
 
-void check_nested_depth(cudf::table_view const& input, int max_nested_depth)
+void check_nested_depth(cudf::table_view const& input)
 {
   using column_checker_fn_t = std::function<int(cudf::column_view const&)>;
 
@@ -462,11 +466,11 @@ void check_nested_depth(cudf::table_view const& input, int max_nested_depth)
 
   for (auto i = 0; i < input.num_columns(); i++) {
     cudf::column_view const& col = input.column(i);
-    CUDF_EXPECTS(get_nested_depth(col) <= max_nested_depth,
+    CUDF_EXPECTS(get_nested_depth(col) <= MAX_NESTED_DEPTH,
                  "The " + std::to_string(i) +
                    "-th column exceeds the maximum allowed nested depth. " +
                    "Current depth: " + std::to_string(get_nested_depth(col)) + ", " +
-                   "Maximum allowed depth: " + std::to_string(max_nested_depth));
+                   "Maximum allowed depth: " + std::to_string(MAX_NESTED_DEPTH));
   }
 }
 
@@ -485,8 +489,7 @@ std::unique_ptr<cudf::column> hive_hash(cudf::table_view const& input,
   // Return early if there's nothing to hash
   if (input.num_columns() == 0 || input.num_rows() == 0) { return output; }
 
-    constexpr int  MAX_NESTED_DEPTH = 8;
-  check_nested_depth(input, max_nested_depth);
+  check_nested_depth(input);
 
   bool const nullable   = has_nested_nulls(input);
   auto const input_view = cudf::table_device_view::create(input, stream);
@@ -497,7 +500,7 @@ std::unique_ptr<cudf::column> hive_hash(cudf::table_view const& input,
     rmm::exec_policy(stream),
     output_view.begin<hive_hash_value_t>(),
     output_view.end<hive_hash_value_t>(),
-    hive_device_row_hasher<hive_hash_function, bool, max_nested_depth>(nullable, *input_view));
+    hive_device_row_hasher<hive_hash_function, bool>(nullable, *input_view));
 
   return output;
 }

From 1d6c16b9eadde516f09d1edbd9da2e32dbcd1822 Mon Sep 17 00:00:00 2001
From: Yan Feng <fengyan_@mail.ustc.edu.cn>
Date: Fri, 22 Nov 2024 17:55:03 +0800
Subject: [PATCH 19/25] Format C++ code

Signed-off-by: Yan Feng <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu | 17 ++++++++---------
 1 file changed, 8 insertions(+), 9 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 4e6cfad107..ccaed15483 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -37,7 +37,7 @@ using hive_hash_value_t = int32_t;
 constexpr hive_hash_value_t HIVE_HASH_FACTOR = 31;
 constexpr hive_hash_value_t HIVE_INIT_HASH   = 0;
 
-constexpr int  MAX_NESTED_DEPTH = 8;
+constexpr int MAX_NESTED_DEPTH = 8;
 
 hive_hash_value_t __device__ inline compute_int(int32_t key) { return key; }
 
@@ -215,9 +215,9 @@ class hive_device_row_hasher {
      */
     struct col_stack_frame {
      private:
-      cudf::column_device_view _column;
-      hive_hash_value_t _cur_hash;        // current hash value of the column
-      int _idx_to_process;                // the index of child or element to process next
+      cudf::column_device_view _column;  // the column has only one row
+      hive_hash_value_t _cur_hash;       // current hash value of the column
+      int _idx_to_process;               // the index of child or element to process next
 
      public:
       __device__ col_stack_frame() =
@@ -496,11 +496,10 @@ std::unique_ptr<cudf::column> hive_hash(cudf::table_view const& input,
   auto output_view      = output->mutable_view();
 
   // Compute the hash value for each row
-  thrust::tabulate(
-    rmm::exec_policy(stream),
-    output_view.begin<hive_hash_value_t>(),
-    output_view.end<hive_hash_value_t>(),
-    hive_device_row_hasher<hive_hash_function, bool>(nullable, *input_view));
+  thrust::tabulate(rmm::exec_policy(stream),
+                   output_view.begin<hive_hash_value_t>(),
+                   output_view.end<hive_hash_value_t>(),
+                   hive_device_row_hasher<hive_hash_function, bool>(nullable, *input_view));
 
   return output;
 }

From 3e9924fd58dc07716797ad404ea6d7ab9116c43f Mon Sep 17 00:00:00 2001
From: Yan Feng <fengyan_@mail.ustc.edu.cn>
Date: Fri, 22 Nov 2024 18:07:17 +0800
Subject: [PATCH 20/25] Update copyright

Signed-off-by: Yan Feng <fengyan_@mail.ustc.edu.cn>
---
 src/test/java/com/nvidia/spark/rapids/jni/HashTest.java | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java b/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
index d3ef6b5c54..19172a8d33 100644
--- a/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
+++ b/src/test/java/com/nvidia/spark/rapids/jni/HashTest.java
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2023, NVIDIA CORPORATION.
+ * Copyright (c) 2023-2024, 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 ced2632bdeca7ee58725fb65d8687f7a2274c742 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Sat, 23 Nov 2024 14:26:45 +0800
Subject: [PATCH 21/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index ccaed15483..349803af0a 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -224,7 +224,7 @@ class hive_device_row_hasher {
         delete;  // Because the default constructor of `cudf::column_device_view` is deleted
 
       __device__ col_stack_frame(cudf::column_device_view col)
-        : _column(col), _idx_to_process(0), _cur_hash(HIVE_INIT_HASH)
+        : _column(std::move(col)), _idx_to_process(0), _cur_hash(HIVE_INIT_HASH)
       {
       }
 

From cccdb8c4fa931c5d28f7dd01f458f3299da4af5e Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Sat, 23 Nov 2024 14:26:52 +0800
Subject: [PATCH 22/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 349803af0a..c0d61cdde1 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -447,7 +447,7 @@ void check_nested_depth(cudf::table_view const& input)
 
   column_checker_fn_t get_nested_depth = [&](cudf::column_view const& col) {
     if (col.type().id() == cudf::type_id::LIST) {
-      auto const& child_col = cudf::lists_column_view(col).child();
+      auto const child_col = cudf::lists_column_view(col).child();
       if (child_col.type().id() != cudf::type_id::STRUCT &&
           child_col.type().id() != cudf::type_id::LIST) {
         return 1;

From 7b4d647698144ad2042038b1c771f49e7b3bc2e0 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Sat, 23 Nov 2024 14:27:03 +0800
Subject: [PATCH 23/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index c0d61cdde1..3ef72ef877 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -334,7 +334,7 @@ class hive_device_row_hasher {
      * |2    |2   |
      * |3    |null|
      * length: 4
-     * null_mask: 1110
+* null_mask: 0111
      *
      * Since the underlying data loses the null information of the top-level list column, computing
      * hash values using the underlying data merely can yield different results compared to Spark.

From ebdc2dd278b4e0a6bd4d46046e2f3651634b42b7 Mon Sep 17 00:00:00 2001
From: ustcfy <96854327+ustcfy@users.noreply.github.com>
Date: Sat, 23 Nov 2024 14:54:37 +0800
Subject: [PATCH 24/25] Update src/main/cpp/src/hive_hash.cu

Co-authored-by: Nghia Truong <7416935+ttnghia@users.noreply.github.com>
---
 src/main/cpp/src/hive_hash.cu | 15 +++++++++++----
 1 file changed, 11 insertions(+), 4 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index 3ef72ef877..e06e5b2395 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -390,10 +390,17 @@ class hive_device_row_hasher {
           if (top.get_idx_to_process() == curr_col.num_child_columns()) {
             if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(top.get_hash()); }
           } else {
-            // Push the next child into the stack
-            col_stack[stack_size++] =
-              col_stack_frame(cudf::detail::structs_column_device_view(curr_col).get_sliced_child(
-                top.get_and_inc_idx_to_process()));
+    auto const scv = cudf::detail::structs_column_device_view(curr_col);
+    for(auto idx = top.get_idx_to_process(); idx < scv.num_child_columns(); ++idx) {
+      auto const child = scv.get_sliced_child(idx);
+      top.get_and_inc_idx_to_process();
+      if(child is not nested) {
+        // update hash of current row
+      } else {
+        col_stack[stack_size++] = col_stack_frame(child);
+        break;
+      }
+    }
           }
         } else if (curr_col.type().id() == cudf::type_id::LIST) {
           // Get the child column of the list column

From 80688ec5e0e71e8ef3853c3c372c6fb6ea7fb2fc Mon Sep 17 00:00:00 2001
From: Yan Feng <fengyan_@mail.ustc.edu.cn>
Date: Sun, 24 Nov 2024 01:00:24 +0800
Subject: [PATCH 25/25] Primitive types do not enter the stack

Signed-off-by: Yan Feng <fengyan_@mail.ustc.edu.cn>
---
 src/main/cpp/src/hive_hash.cu | 65 +++++++++++++++--------------------
 1 file changed, 27 insertions(+), 38 deletions(-)

diff --git a/src/main/cpp/src/hive_hash.cu b/src/main/cpp/src/hive_hash.cu
index e06e5b2395..c10c424b8b 100644
--- a/src/main/cpp/src/hive_hash.cu
+++ b/src/main/cpp/src/hive_hash.cu
@@ -271,11 +271,9 @@ class hive_device_row_hasher {
      *    b. If the column is a structs column:
      *        i.  If all child columns are processed, pop the element and update `cur_hash` of its
      *            parent column.
-     *        ii. Otherwise, push the next child column into the stack.
+     *        ii. Otherwise, process the next child column.
      *    c. If the column is a lists column, process it by a similar way as structs column but
      *       iterating through the list elements instead of child columns' elements.
-     *    d. If the column is a primitive column, compute the hash value, pop the element,
-     *       and update `cur_hash` of its parent column.
      * 3. Return the hash value of the root column.
      *
      * For example, consider the following nested column: `Struct<Struct<int, float>, decimal>`
@@ -286,14 +284,12 @@ class hive_device_row_hasher {
      *   / \
      *  i   f
      *
-     * The `pop` order of the stack is: i, f, S2, d, S1.
-     * - When i   is popped, S2's cur_hash is updated to `hash(i)`.
-     * - When f   is popped, S2's cur_hash is updated to `hash(i) * HIVE_HASH_FACTOR + hash(f)`,
-     *   which is the hash value of S2.
-     * - When S2  is popped, S1's cur_hash is updated to `hash(S2)`.
-     * - When d   is popped, S1's cur_hash is updated to `hash(S2) * HIVE_HASH_FACTOR + hash(d)`,
-     *   which is the hash value of S1.
-     * - When S1  is popped, the hash value of the root column is returned.
+     * - First, S1 is pushed into the stack. Then, S2 is pushed into the stack.
+     * - S2's hash value can be computed directly because its children are of primitive types.
+     *   When S2 is popped, S1's `cur_hash` is updated to S2's hash value.
+     * - Now the top of the stack is S1. The next child to process is d. S1's `cur_hash` is updated
+     *   to `hash(S2) * HIVE_HASH_FACTOR + hash(d)`, which is the hash value of S1.
+     * - When S1 is popped, the hash value of the root column is returned.
      *
      * As lists columns have a different interface from structs columns, we need to handle them
      * separately.
@@ -334,7 +330,7 @@ class hive_device_row_hasher {
      * |2    |2   |
      * |3    |null|
      * length: 4
-* null_mask: 0111
+     * null_mask: 0111
      *
      * Since the underlying data loses the null information of the top-level list column, computing
      * hash values using the underlying data merely can yield different results compared to Spark.
@@ -354,9 +350,8 @@ class hive_device_row_hasher {
      *    i1           i2         int
      *
      * Note: L2、i1、i2 are all temporary columns, which would not be pushed into the stack.
-     * There is an optimization for the list column. If the child column is of primitive type, the
-     * hash value of the list column can be directly computed. Thus we can decrease the depth of
-     * the stack by one.
+     * If the child column is of primitive type, the hash value of the list column can be directly
+     * computed.
      *
      * @tparam T The type of the column.
      * @param col The column device view.
@@ -385,22 +380,26 @@ class hive_device_row_hasher {
         // Do not pop it until it is processed. The definition of `processed` is:
         // - For structs, it is when all child columns are processed.
         // - For lists, it is when all elements in the list are processed.
-        // - For primitive types, it is when the hash value is computed.
         if (curr_col.type().id() == cudf::type_id::STRUCT) {
           if (top.get_idx_to_process() == curr_col.num_child_columns()) {
             if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(top.get_hash()); }
           } else {
-    auto const scv = cudf::detail::structs_column_device_view(curr_col);
-    for(auto idx = top.get_idx_to_process(); idx < scv.num_child_columns(); ++idx) {
-      auto const child = scv.get_sliced_child(idx);
-      top.get_and_inc_idx_to_process();
-      if(child is not nested) {
-        // update hash of current row
-      } else {
-        col_stack[stack_size++] = col_stack_frame(child);
-        break;
-      }
-    }
+            auto const structcv = cudf::detail::structs_column_device_view(curr_col);
+            while (top.get_idx_to_process() < curr_col.num_child_columns()) {
+              auto idx             = top.get_and_inc_idx_to_process();
+              auto const child_col = structcv.get_sliced_child(idx);
+              // If the child is of primitive type, accumulate child hash into struct hash
+              if (child_col.type().id() != cudf::type_id::LIST &&
+                  child_col.type().id() != cudf::type_id::STRUCT) {
+                auto child_hash =
+                  cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
+                    child_col.type(), this->hash_functor, child_col, 0);
+                top.update_cur_hash(child_hash);
+              } else {
+                col_stack[stack_size++] = col_stack_frame(child_col);
+                break;
+              }
+            }
           }
         } else if (curr_col.type().id() == cudf::type_id::LIST) {
           // Get the child column of the list column
@@ -429,12 +428,6 @@ class hive_device_row_hasher {
                 col_stack_frame(child_col.slice(top.get_and_inc_idx_to_process(), 1));
             }
           }
-        } else {
-          // There is only one element in the column for primitive types
-          auto hash = cudf::type_dispatcher<cudf::experimental::dispatch_void_if_nested>(
-            curr_col.type(), this->hash_functor, curr_col, 0);
-          top.update_cur_hash(hash);
-          if (--stack_size > 0) { col_stack[stack_size - 1].update_cur_hash(top.get_hash()); }
         }
       }
       return col_stack[0].get_hash();
@@ -455,10 +448,6 @@ void check_nested_depth(cudf::table_view const& input)
   column_checker_fn_t get_nested_depth = [&](cudf::column_view const& col) {
     if (col.type().id() == cudf::type_id::LIST) {
       auto const child_col = cudf::lists_column_view(col).child();
-      if (child_col.type().id() != cudf::type_id::STRUCT &&
-          child_col.type().id() != cudf::type_id::LIST) {
-        return 1;
-      }
       return 1 + get_nested_depth(child_col);
     } else if (col.type().id() == cudf::type_id::STRUCT) {
       int max_child_depth = 0;
@@ -467,7 +456,7 @@ void check_nested_depth(cudf::table_view const& input)
       }
       return 1 + max_child_depth;
     } else {  // Primitive type
-      return 1;
+      return 0;
     }
   };