Add specialized dispatch to improve occupancy for hash table operations in distinct join

cpp/CMakeLists.txt

@@ -462,8 +462,10 @@ add_library(
   src/join/mixed_join.cu
   src/join/mixed_join_kernel.cu
   src/join/mixed_join_kernel_nulls.cu
-  src/join/mixed_join_kernels_semi.cu
   src/join/mixed_join_semi.cu
+  src/join/mixed_join_kernels_semi.cu
+  src/join/mixed_join_kernels_semi_nested.cu
+  src/join/mixed_join_kernels_semi_compound.cu
   src/join/mixed_join_size_kernel.cu
   src/join/mixed_join_size_kernel_nulls.cu
   src/join/semi_join.cu

cpp/include/cudf/detail/distinct_hash_join.cuh

@@ -29,6 +29,7 @@
 #include <memory>
 #include <type_traits>
 #include <utility>
+#include <variant>
 
 namespace cudf::detail {
 
@@ -85,22 +86,42 @@ struct hasher_adapter {
 template <cudf::has_nested HasNested>
 struct distinct_hash_join {
  private:
-  /// Device row equal type
-  using d_equal_type = cudf::experimental::row::equality::strong_index_comparator_adapter<
-    cudf::experimental::row::equality::device_row_comparator<HasNested == cudf::has_nested::YES,
-                                                             cudf::nullate::DYNAMIC>>;
+  using row_comparator = cudf::experimental::row::equality::device_row_comparator<
+    true,
+    cudf::nullate::DYNAMIC,
+    cudf::experimental::row::equality::nan_equal_physical_equality_comparator,
+    cudf::experimental::type_identity_t>;
+
+  using row_comparator_no_nested = cudf::experimental::row::equality::device_row_comparator<
+    false,
+    cudf::nullate::DYNAMIC,
+    cudf::experimental::row::equality::nan_equal_physical_equality_comparator,
+    cudf::experimental::dispatch_void_if_nested_t>;
+
+  using row_comparator_no_compound = cudf::experimental::row::equality::device_row_comparator<
+    false,
+    cudf::nullate::DYNAMIC,
+    cudf::experimental::row::equality::nan_equal_physical_equality_comparator,
+    cudf::experimental::dispatch_void_if_compound_t>;
+
   using hasher              = hasher_adapter<thrust::identity<hash_value_type>>;
   using probing_scheme_type = cuco::linear_probing<1, hasher>;
   using cuco_storage_type   = cuco::storage<1>;
 
   /// Hash table type
-  using hash_table_type = cuco::static_set<cuco::pair<hash_value_type, rhs_index_type>,
-                                           cuco::extent<size_type>,
-                                           cuda::thread_scope_device,
-                                           comparator_adapter<d_equal_type>,
-                                           probing_scheme_type,
-                                           cudf::detail::cuco_allocator,
-                                           cuco_storage_type>;
+  template <typename Comparator>
+  using static_set_with_comparator = cuco::static_set<
+    cuco::pair<hash_value_type, rhs_index_type>,
+    cuco::extent<size_type>,
+    cuda::thread_scope_device,
+    comparator_adapter<
+      cudf::experimental::row::equality::strong_index_comparator_adapter<Comparator>>,
+    probing_scheme_type,
+    cudf::detail::cuco_allocator,
+    cuco_storage_type>;
+  using hash_table_type = std::variant<static_set_with_comparator<row_comparator>,
+                                       static_set_with_comparator<row_comparator_no_nested>,
+                                       static_set_with_comparator<row_comparator_no_compound>>;
 
   bool _has_nulls;  ///< true if nulls are present in either build table or probe table
   cudf::null_equality _nulls_equal;  ///< whether to consider nulls as equal
@@ -109,8 +130,8 @@ struct distinct_hash_join {
   std::shared_ptr<cudf::experimental::row::equality::preprocessed_table>
     _preprocessed_build;  ///< input table preprocssed for row operators
   std::shared_ptr<cudf::experimental::row::equality::preprocessed_table>
-    _preprocessed_probe;        ///< input table preprocssed for row operators
-  hash_table_type _hash_table;  ///< hash table built on `_build`
+    _preprocessed_probe;                         ///< input table preprocssed for row operators
+  std::unique_ptr<hash_table_type> _hash_table;  ///< hash table built on `_build`
 
  public:
   distinct_hash_join()                                     = delete;