Occupancy improvement for Hash table build

cpp/src/join/hash_join.cu

@@ -42,6 +42,7 @@
 #include <cstddef>
 #include <iostream>
 #include <numeric>
+#include <variant>
 
 namespace cudf {
 namespace detail {
@@ -97,33 +98,44 @@ std::size_t compute_join_output_size(
   }
 
   auto const probe_nulls = cudf::nullate::DYNAMIC{has_nulls};
+  std::unordered_set<cudf::type_id> probe_column_types;
+  for (auto col : probe_table) {
+    probe_column_types.insert(col.type().id());
+  }
 
   auto const row_hash           = cudf::experimental::row::hash::row_hasher{preprocessed_probe};
-  auto const hash_probe         = row_hash.device_hasher(probe_nulls);
+  auto const hash_probe         = row_hash.device_hasher(probe_column_types, probe_nulls);
   auto const empty_key_sentinel = hash_table.get_empty_key_sentinel();
-  auto const iter               = cudf::detail::make_counting_transform_iterator(
-    0, make_pair_function{hash_probe, empty_key_sentinel});
 
   auto const row_comparator =
     cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build};
   auto const comparator_helper = [&](auto device_comparator) {
     pair_equality equality{device_comparator};
 
     if (join == join_kind::LEFT_JOIN) {
-      return hash_table.pair_count_outer(
-        iter, iter + probe_table_num_rows, equality, stream.value());
+      return std::visit(
+        [&](auto&& hasher) {
+          auto const iter = cudf::detail::make_counting_transform_iterator(
+            0, make_pair_function{hasher, empty_key_sentinel});
+          return hash_table.pair_count_outer(
+            iter, iter + probe_table_num_rows, equality, stream.value());
+        },
+        hash_probe);
     } else {
-      return hash_table.pair_count(iter, iter + probe_table_num_rows, equality, stream.value());
+      return std::visit(
+        [&](auto&& hasher) {
+          auto const iter = cudf::detail::make_counting_transform_iterator(
+            0, make_pair_function{hasher, empty_key_sentinel});
+          return hash_table.pair_count(iter, iter + probe_table_num_rows, equality, stream.value());
+        },
+        hash_probe);
     }
   };
 
-  if (cudf::detail::has_nested_columns(probe_table)) {
-    auto const device_comparator = row_comparator.equal_to<true>(has_nulls, nulls_equal);
-    return comparator_helper(device_comparator);
-  } else {
-    auto const device_comparator = row_comparator.equal_to<false>(has_nulls, nulls_equal);
-    return comparator_helper(device_comparator);
-  }
+  auto const device_comparator =
+    row_comparator.equal_to(probe_column_types, has_nulls, nulls_equal);
+  return std::visit([&](auto&& comparator) { return comparator_helper(comparator); },
+                    device_comparator);
 }
 
 /**
@@ -187,12 +199,14 @@ probe_join_hash_table(
   auto right_indices = std::make_unique<rmm::device_uvector<size_type>>(join_size, stream, mr);
 
   auto const probe_nulls = cudf::nullate::DYNAMIC{has_nulls};
+  std::unordered_set<cudf::type_id> probe_column_types;
+  for (auto col : probe_table) {
+    probe_column_types.insert(col.type().id());
+  }
 
   auto const row_hash           = cudf::experimental::row::hash::row_hasher{preprocessed_probe};
-  auto const hash_probe         = row_hash.device_hasher(probe_nulls);
+  auto const hash_probe         = row_hash.device_hasher(probe_column_types, probe_nulls);
   auto const empty_key_sentinel = hash_table.get_empty_key_sentinel();
-  auto const iter               = cudf::detail::make_counting_transform_iterator(
-    0, make_pair_function{hash_probe, empty_key_sentinel});
 
   cudf::size_type const probe_table_num_rows = probe_table.num_rows();
 
@@ -207,36 +221,42 @@ probe_join_hash_table(
     pair_equality equality{device_comparator};
 
     if (join == cudf::detail::join_kind::FULL_JOIN or join == cudf::detail::join_kind::LEFT_JOIN) {
-      [[maybe_unused]] auto [out1_zip_end, out2_zip_end] =
-        hash_table.pair_retrieve_outer(iter,
-                                       iter + probe_table_num_rows,
-                                       out1_zip_begin,
-                                       out2_zip_begin,
-                                       equality,
-                                       stream.value());
+      [[maybe_unused]] auto [out1_zip_end, out2_zip_end] = std::visit(
+        [&](auto&& hasher) {
+          auto const iter = cudf::detail::make_counting_transform_iterator(
+            0, make_pair_function{hasher, empty_key_sentinel});
+          return hash_table.pair_retrieve_outer(iter,
+                                                iter + probe_table_num_rows,
+                                                out1_zip_begin,
+                                                out2_zip_begin,
+                                                equality,
+                                                stream.value());
+        },
+        hash_probe);
 
       if (join == cudf::detail::join_kind::FULL_JOIN) {
         auto const actual_size = thrust::distance(out1_zip_begin, out1_zip_end);
         left_indices->resize(actual_size, stream);
         right_indices->resize(actual_size, stream);
       }
     } else {
-      hash_table.pair_retrieve(iter,
-                               iter + probe_table_num_rows,
-                               out1_zip_begin,
-                               out2_zip_begin,
-                               equality,
-                               stream.value());
+      std::visit(
+        [&](auto&& hasher) {
+          auto const iter = cudf::detail::make_counting_transform_iterator(
+            0, make_pair_function{hasher, empty_key_sentinel});
+          hash_table.pair_retrieve(iter,
+                                   iter + probe_table_num_rows,
+                                   out1_zip_begin,
+                                   out2_zip_begin,
+                                   equality,
+                                   stream.value());
+        },
+        hash_probe);
     }
   };
-
-  if (cudf::detail::has_nested_columns(probe_table)) {
-    auto const device_comparator = row_comparator.equal_to<true>(probe_nulls, compare_nulls);
-    comparator_helper(device_comparator);
-  } else {
-    auto const device_comparator = row_comparator.equal_to<false>(probe_nulls, compare_nulls);
-    comparator_helper(device_comparator);
-  }
+  auto const device_comparator =
+    row_comparator.equal_to(probe_column_types, probe_nulls, compare_nulls);
+  std::visit([&](auto&& comparator) { comparator_helper(comparator); }, device_comparator);
 
   return std::pair(std::move(left_indices), std::move(right_indices));
 }
@@ -289,12 +309,14 @@ std::size_t get_full_join_size(
   auto right_indices = std::make_unique<rmm::device_uvector<size_type>>(join_size, stream, mr);
 
   auto const probe_nulls = cudf::nullate::DYNAMIC{has_nulls};
+  std::unordered_set<cudf::type_id> probe_column_types;
+  for (auto col : probe_table) {
+    probe_column_types.insert(col.type().id());
+  }
 
   auto const row_hash           = cudf::experimental::row::hash::row_hasher{preprocessed_probe};
-  auto const hash_probe         = row_hash.device_hasher(probe_nulls);
+  auto const hash_probe         = row_hash.device_hasher(probe_column_types, probe_nulls);
   auto const empty_key_sentinel = hash_table.get_empty_key_sentinel();
-  auto const iter               = cudf::detail::make_counting_transform_iterator(
-    0, make_pair_function{hash_probe, empty_key_sentinel});
 
   cudf::size_type const probe_table_num_rows = probe_table.num_rows();
 
@@ -307,16 +329,22 @@ std::size_t get_full_join_size(
     cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build};
   auto const comparator_helper = [&](auto device_comparator) {
     pair_equality equality{device_comparator};
-    hash_table.pair_retrieve_outer(
-      iter, iter + probe_table_num_rows, out1_zip_begin, out2_zip_begin, equality, stream.value());
+    std::visit(
+      [&](auto&& hasher) {
+        auto const iter = cudf::detail::make_counting_transform_iterator(
+          0, make_pair_function{hasher, empty_key_sentinel});
+        hash_table.pair_retrieve_outer(iter,
+                                       iter + probe_table_num_rows,
+                                       out1_zip_begin,
+                                       out2_zip_begin,
+                                       equality,
+                                       stream.value());
+      },
+      hash_probe);
   };
-  if (cudf::detail::has_nested_columns(probe_table)) {
-    auto const device_comparator = row_comparator.equal_to<true>(probe_nulls, compare_nulls);
-    comparator_helper(device_comparator);
-  } else {
-    auto const device_comparator = row_comparator.equal_to<false>(probe_nulls, compare_nulls);
-    comparator_helper(device_comparator);
-  }
+  auto const device_comparator =
+    row_comparator.equal_to(probe_column_types, probe_nulls, compare_nulls);
+  std::visit([&](auto&& comparator) { comparator_helper(comparator); }, device_comparator);
 
   // Release intermediate memory allocation
   left_indices->resize(0, stream);

cpp/src/join/join_common_utils.cuh

@@ -177,23 +177,37 @@ void build_join_hash_table(
   CUDF_EXPECTS(0 != build.num_columns(), "Selected build dataset is empty");
   CUDF_EXPECTS(0 != build.num_rows(), "Build side table has no rows");
 
+  std::unordered_set<cudf::type_id> build_column_types;
+  for (auto col : build) {
+    build_column_types.insert(col.type().id());
+  }
+
   auto const row_hash   = experimental::row::hash::row_hasher{preprocessed_build};
-  auto const hash_build = row_hash.device_hasher(nullate::DYNAMIC{has_nulls});
+  auto const hash_build = row_hash.device_hasher(build_column_types, nullate::DYNAMIC{has_nulls});
 
   auto const empty_key_sentinel = hash_table.get_empty_key_sentinel();
-  make_pair_function pair_func{hash_build, empty_key_sentinel};
-
-  auto const iter = cudf::detail::make_counting_transform_iterator(0, pair_func);
 
   size_type const build_table_num_rows{build.num_rows()};
   if (nulls_equal == cudf::null_equality::EQUAL or (not nullable(build))) {
-    hash_table.insert(iter, iter + build_table_num_rows, stream.value());
+    std::visit(
+      [&](auto&& hasher) {
+        make_pair_function pair_func{hasher, empty_key_sentinel};
+        auto const iter = cudf::detail::make_counting_transform_iterator(0, pair_func);
+        hash_table.insert(iter, iter + build_table_num_rows, stream.value());
+      },
+      hash_build);
   } else {
     thrust::counting_iterator<size_type> stencil(0);
     row_is_valid pred{bitmask};
 
-    // insert valid rows
-    hash_table.insert_if(iter, iter + build_table_num_rows, stencil, pred, stream.value());
+    std::visit(
+      [&](auto&& hasher) {
+        make_pair_function pair_func{hasher, empty_key_sentinel};
+        auto const iter = cudf::detail::make_counting_transform_iterator(0, pair_func);
+        // insert valid rows
+        hash_table.insert_if(iter, iter + build_table_num_rows, stencil, pred, stream.value());
+      },
+      hash_build);
   }
 }
 

cpp/src/join/mixed_join_common_utils.cuh

@@ -31,7 +31,18 @@ namespace detail {
 
 using row_hash =
   cudf::experimental::row::hash::device_row_hasher<cudf::hashing::detail::default_hash,
-                                                   cudf::nullate::DYNAMIC>;
+                                                   cudf::nullate::DYNAMIC,
+                                                   cudf::experimental::type_identity_t>;
+
+using row_hash_no_nested =
+  cudf::experimental::row::hash::device_row_hasher<cudf::hashing::detail::default_hash,
+                                                   cudf::nullate::DYNAMIC,
+                                                   cudf::experimental::dispatch_void_if_nested_t>;
+
+using row_hash_no_compound =
+  cudf::experimental::row::hash::device_row_hasher<cudf::hashing::detail::default_hash,
+                                                   cudf::nullate::DYNAMIC,
+                                                   cudf::experimental::dispatch_void_if_compound_t>;
 
 // // This alias is used by mixed_joins, which support only non-nested types
 using row_equality = cudf::experimental::row::equality::strong_index_comparator_adapter<

cpp/src/join/mixed_join_kernels_semi.cu

@@ -33,13 +33,13 @@ namespace cg = cooperative_groups;
 
 #pragma GCC diagnostic ignored "-Wattributes"
 
-template <cudf::size_type block_size, bool has_nulls>
+template <cudf::size_type block_size, bool has_nulls, typename HashProbe>
 __attribute__((visibility("hidden"))) __launch_bounds__(block_size) __global__
   void mixed_join_semi(table_device_view left_table,
                        table_device_view right_table,
                        table_device_view probe,
                        table_device_view build,
-                       row_hash const hash_probe,
+                       HashProbe const hash_probe,
                        row_equality const equality_probe,
                        cudf::detail::semi_map_type::device_view hash_table_view,
                        cudf::device_span<bool> left_table_keep_mask,
@@ -74,7 +74,7 @@ __attribute__((visibility("hidden"))) __launch_bounds__(block_size) __global__
   }
 }
 
-template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, true>(
+template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, true, row_hash>(
   table_device_view left_table,
   table_device_view right_table,
   table_device_view probe,
@@ -85,7 +85,7 @@ template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, true>(
   cudf::device_span<bool> left_table_keep_mask,
   cudf::ast::detail::expression_device_view device_expression_data);
 
-template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, false>(
+template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, false, row_hash>(
   table_device_view left_table,
   table_device_view right_table,
   table_device_view probe,
@@ -96,6 +96,50 @@ template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, false>(
   cudf::device_span<bool> left_table_keep_mask,
   cudf::ast::detail::expression_device_view device_expression_data);
 
+template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, true, row_hash_no_nested>(
+  table_device_view left_table,
+  table_device_view right_table,
+  table_device_view probe,
+  table_device_view build,
+  row_hash_no_nested const hash_probe,
+  row_equality const equality_probe,
+  cudf::detail::semi_map_type::device_view hash_table_view,
+  cudf::device_span<bool> left_table_keep_mask,
+  cudf::ast::detail::expression_device_view device_expression_data);
+
+template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, false, row_hash_no_nested>(
+  table_device_view left_table,
+  table_device_view right_table,
+  table_device_view probe,
+  table_device_view build,
+  row_hash_no_nested const hash_probe,
+  row_equality const equality_probe,
+  cudf::detail::semi_map_type::device_view hash_table_view,
+  cudf::device_span<bool> left_table_keep_mask,
+  cudf::ast::detail::expression_device_view device_expression_data);
+
+template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, true, row_hash_no_compound>(
+  table_device_view left_table,
+  table_device_view right_table,
+  table_device_view probe,
+  table_device_view build,
+  row_hash_no_compound const hash_probe,
+  row_equality const equality_probe,
+  cudf::detail::semi_map_type::device_view hash_table_view,
+  cudf::device_span<bool> left_table_keep_mask,
+  cudf::ast::detail::expression_device_view device_expression_data);
+
+template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, false, row_hash_no_compound>(
+  table_device_view left_table,
+  table_device_view right_table,
+  table_device_view probe,
+  table_device_view build,
+  row_hash_no_compound const hash_probe,
+  row_equality const equality_probe,
+  cudf::detail::semi_map_type::device_view hash_table_view,
+  cudf::device_span<bool> left_table_keep_mask,
+  cudf::ast::detail::expression_device_view device_expression_data);
+
 }  // namespace detail
 
 }  // namespace cudf

cpp/src/join/mixed_join_kernels_semi.cuh

@@ -51,12 +51,12 @@ namespace detail {
  * @param[in] device_expression_data Container of device data required to evaluate the desired
  * expression.
  */
-template <cudf::size_type block_size, bool has_nulls>
+template <cudf::size_type block_size, bool has_nulls, typename HashProbe>
 __global__ void mixed_join_semi(table_device_view left_table,
                                 table_device_view right_table,
                                 table_device_view probe,
                                 table_device_view build,
-                                row_hash const hash_probe,
+                                HashProbe const hash_probe,
                                 row_equality const equality_probe,
                                 cudf::detail::semi_map_type::device_view hash_table_view,
                                 cudf::device_span<bool> left_table_keep_mask,