Update mixed_join to use experimental row hasher and comparator

cpp/src/join/mixed_join.cu

@@ -109,9 +109,9 @@ mixed_join(
   // If evaluating the expression may produce null outputs we create a nullable
   // output column and follow the null-supporting expression evaluation code
   // path.
-  auto const has_nulls =
+  auto const has_nulls = cudf::nullate::DYNAMIC{
     cudf::has_nulls(left_equality) || cudf::has_nulls(right_equality) ||
-    binary_predicate.may_evaluate_null(left_conditional, right_conditional, stream);
+    binary_predicate.may_evaluate_null(left_conditional, right_conditional, stream)};
 
   auto const parser = ast::detail::expression_parser{
     binary_predicate, left_conditional, right_conditional, has_nulls, stream, mr};
@@ -125,8 +125,6 @@ mixed_join(
   auto& build     = swap_tables ? left_equality : right_equality;
   auto probe_view = table_device_view::create(probe, stream);
   auto build_view = table_device_view::create(build, stream);
-  row_equality equality_probe{
-    cudf::nullate::DYNAMIC{has_nulls}, *probe_view, *build_view, compare_nulls};
 
   // Don't use multimap_type because we want a CG size of 1.
   mixed_multimap_type hash_table{
@@ -168,6 +166,14 @@ mixed_join(
   std::optional<rmm::device_uvector<size_type>> matches_per_row{};
   device_span<size_type const> matches_per_row_span{};
 
+  auto const preprocessed_probe =
+    experimental::row::equality::preprocessed_table::create(probe, stream);
+  auto const row_hash   = cudf::experimental::row::hash::row_hasher{preprocessed_probe};
+  auto const hash_probe = row_hash.device_hasher(has_nulls);
+  auto const row_comparator =
+    cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build};
+  auto const equality_probe = row_comparator.equal_to<false>(has_nulls, compare_nulls);
+
   if (output_size_data.has_value()) {
     join_size            = output_size_data->first;
     matches_per_row_span = output_size_data->second;
@@ -190,6 +196,7 @@ mixed_join(
           *right_conditional_view,
           *probe_view,
           *build_view,
+          hash_probe,
           equality_probe,
           kernel_join_type,
           hash_table_view,
@@ -204,6 +211,7 @@ mixed_join(
           *right_conditional_view,
           *probe_view,
           *build_view,
+          hash_probe,
           equality_probe,
           kernel_join_type,
           hash_table_view,
@@ -247,6 +255,7 @@ mixed_join(
         *right_conditional_view,
         *probe_view,
         *build_view,
+        hash_probe,
         equality_probe,
         kernel_join_type,
         hash_table_view,
@@ -262,6 +271,7 @@ mixed_join(
         *right_conditional_view,
         *probe_view,
         *build_view,
+        hash_probe,
         equality_probe,
         kernel_join_type,
         hash_table_view,
@@ -364,9 +374,9 @@ compute_mixed_join_output_size(table_view const& left_equality,
   // If evaluating the expression may produce null outputs we create a nullable
   // output column and follow the null-supporting expression evaluation code
   // path.
-  auto const has_nulls =
+  auto const has_nulls = cudf::nullate::DYNAMIC{
     cudf::has_nulls(left_equality) || cudf::has_nulls(right_equality) ||
-    binary_predicate.may_evaluate_null(left_conditional, right_conditional, stream);
+    binary_predicate.may_evaluate_null(left_conditional, right_conditional, stream)};
 
   auto const parser = ast::detail::expression_parser{
     binary_predicate, left_conditional, right_conditional, has_nulls, stream, mr};
@@ -380,8 +390,6 @@ compute_mixed_join_output_size(table_view const& left_equality,
   auto& build     = swap_tables ? left_equality : right_equality;
   auto probe_view = table_device_view::create(probe, stream);
   auto build_view = table_device_view::create(build, stream);
-  row_equality equality_probe{
-    cudf::nullate::DYNAMIC{has_nulls}, *probe_view, *build_view, compare_nulls};
 
   // Don't use multimap_type because we want a CG size of 1.
   mixed_multimap_type hash_table{
@@ -417,6 +425,14 @@ compute_mixed_join_output_size(table_view const& left_equality,
   // Allocate storage for the counter used to get the size of the join output
   rmm::device_scalar<std::size_t> size(0, stream, mr);
 
+  auto const preprocessed_probe =
+    experimental::row::equality::preprocessed_table::create(probe, stream);
+  auto const row_hash   = cudf::experimental::row::hash::row_hasher{preprocessed_probe};
+  auto const hash_probe = row_hash.device_hasher(has_nulls);
+  auto const row_comparator =
+    cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build};
+  auto const equality_probe = row_comparator.equal_to<false>(has_nulls, compare_nulls);
+
   // Determine number of output rows without actually building the output to simply
   // find what the size of the output will be.
   if (has_nulls) {
@@ -426,6 +442,7 @@ compute_mixed_join_output_size(table_view const& left_equality,
         *right_conditional_view,
         *probe_view,
         *build_view,
+        hash_probe,
         equality_probe,
         join_type,
         hash_table_view,
@@ -440,6 +457,7 @@ compute_mixed_join_output_size(table_view const& left_equality,
         *right_conditional_view,
         *probe_view,
         *build_view,
+        hash_probe,
         equality_probe,
         join_type,
         hash_table_view,

cpp/src/join/mixed_join_common_utils.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
 
 #include <cudf/ast/detail/expression_evaluator.cuh>
 #include <cudf/detail/utilities/cuda.cuh>
+#include <cudf/table/experimental/row_operators.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
@@ -28,6 +29,14 @@
 namespace cudf {
 namespace detail {
 
+using row_hash_experimental =
+  cudf::experimental::row::hash::device_row_hasher<default_hash, cudf::nullate::DYNAMIC>;
+
+// // This alias is used by mixed_joins, which support only non-nested types
+using row_equality_experimental =
+  cudf::experimental::row::equality::strong_index_comparator_adapter<
+    cudf::experimental::row::equality::device_row_comparator<false, cudf::nullate::DYNAMIC>>;
+
 /**
  * @brief Equality comparator for use with cuco map methods that require expression evaluation.
  *
@@ -41,7 +50,7 @@ struct expression_equality {
     cudf::ast::detail::expression_evaluator<has_nulls> const& evaluator,
     cudf::ast::detail::IntermediateDataType<has_nulls>* thread_intermediate_storage,
     bool const swap_tables,
-    row_equality const& equality_probe)
+    row_equality_experimental const& equality_probe)
     : evaluator{evaluator},
       thread_intermediate_storage{thread_intermediate_storage},
       swap_tables{swap_tables},
@@ -52,7 +61,7 @@ struct expression_equality {
   cudf::ast::detail::IntermediateDataType<has_nulls>* thread_intermediate_storage;
   cudf::ast::detail::expression_evaluator<has_nulls> const& evaluator;
   bool const swap_tables;
-  row_equality const& equality_probe;
+  row_equality_experimental const& equality_probe;
 };
 
 /**
@@ -79,12 +88,15 @@ struct single_expression_equality : expression_equality<has_nulls> {
   __device__ __forceinline__ bool operator()(hash_value_type const build_row_index,
                                              hash_value_type const probe_row_index) const noexcept
   {
+    using cudf::experimental::row::lhs_index_type;
+    using cudf::experimental::row::rhs_index_type;
+
     auto output_dest = cudf::ast::detail::value_expression_result<bool, has_nulls>();
     // Two levels of checks:
     // 1. The contents of the columns involved in the equality condition are equal.
     // 2. The predicate evaluated on the relevant columns (already encoded in the evaluator)
     // evaluates to true.
-    if (this->equality_probe(probe_row_index, build_row_index)) {
+    if (this->equality_probe(lhs_index_type{probe_row_index}, rhs_index_type{build_row_index})) {
       auto const lrow_idx = this->swap_tables ? build_row_index : probe_row_index;
       auto const rrow_idx = this->swap_tables ? probe_row_index : build_row_index;
       this->evaluator.evaluate(output_dest,
@@ -127,14 +139,17 @@ struct pair_expression_equality : public expression_equality<has_nulls> {
   __device__ __forceinline__ bool operator()(pair_type const& build_row,
                                              pair_type const& probe_row) const noexcept
   {
+    using cudf::experimental::row::lhs_index_type;
+    using cudf::experimental::row::rhs_index_type;
+
     auto output_dest = cudf::ast::detail::value_expression_result<bool, has_nulls>();
     // Three levels of checks:
     // 1. Row hashes of the columns involved in the equality condition are equal.
     // 2. The contents of the columns involved in the equality condition are equal.
     // 3. The predicate evaluated on the relevant columns (already encoded in the evaluator)
     // evaluates to true.
     if ((probe_row.first == build_row.first) &&
-        this->equality_probe(probe_row.second, build_row.second)) {
+        this->equality_probe(lhs_index_type{probe_row.second}, rhs_index_type{build_row.second})) {
       auto const lrow_idx = this->swap_tables ? build_row.second : probe_row.second;
       auto const rrow_idx = this->swap_tables ? probe_row.second : build_row.second;
       this->evaluator.evaluate(

cpp/src/join/mixed_join_kernel.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -24,7 +24,8 @@ template __global__ void mixed_join<DEFAULT_JOIN_BLOCK_SIZE, false>(
   table_device_view right_table,
   table_device_view probe,
   table_device_view build,
-  row_equality const equality_probe,
+  row_hash_experimental const hash_probe,
+  row_equality_experimental const equality_probe,
   join_kind const join_type,
   cudf::detail::mixed_multimap_type::device_view hash_table_view,
   size_type* join_output_l,

cpp/src/join/mixed_join_kernel.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -42,7 +42,8 @@ __launch_bounds__(block_size) __global__
                   table_device_view right_table,
                   table_device_view probe,
                   table_device_view build,
-                  row_equality const equality_probe,
+                  row_hash_experimental const hash_probe,
+                  row_equality_experimental const equality_probe,
                   join_kind const join_type,
                   cudf::detail::mixed_multimap_type::device_view hash_table_view,
                   size_type* join_output_l,
@@ -70,7 +71,6 @@ __launch_bounds__(block_size) __global__
   auto evaluator = cudf::ast::detail::expression_evaluator<has_nulls>(
     left_table, right_table, device_expression_data);
 
-  row_hash hash_probe{nullate::DYNAMIC{has_nulls}, probe};
   auto const empty_key_sentinel = hash_table_view.get_empty_key_sentinel();
   make_pair_function pair_func{hash_probe, empty_key_sentinel};
 

cpp/src/join/mixed_join_kernel_nulls.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -24,7 +24,8 @@ template __global__ void mixed_join<DEFAULT_JOIN_BLOCK_SIZE, true>(
   table_device_view right_table,
   table_device_view probe,
   table_device_view build,
-  row_equality const equality_probe,
+  row_hash_experimental const hash_probe,
+  row_equality_experimental const equality_probe,
   join_kind const join_type,
   cudf::detail::mixed_multimap_type::device_view hash_table_view,
   size_type* join_output_l,

cpp/src/join/mixed_join_kernels.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -41,6 +41,7 @@ namespace detail {
  * @param[in] right_table The right table
  * @param[in] probe The table with which to probe the hash table for matches.
  * @param[in] build The table with which the hash table was built.
+ * @param[in] hash_probe The hasher used for the probe table.
  * @param[in] equality_probe The equality comparator used when probing the hash table.
  * @param[in] join_type The type of join to be performed
  * @param[in] hash_table_view The hash table built from `build`.
@@ -62,7 +63,8 @@ __global__ void compute_mixed_join_output_size(
   table_device_view right_table,
   table_device_view probe,
   table_device_view build,
-  row_equality const equality_probe,
+  row_hash_experimental const hash_probe,
+  row_equality_experimental const equality_probe,
   join_kind const join_type,
   cudf::detail::mixed_multimap_type::device_view hash_table_view,
   ast::detail::expression_device_view device_expression_data,
@@ -87,6 +89,7 @@ __global__ void compute_mixed_join_output_size(
  * @param[in] right_table The right table
  * @param[in] probe The table with which to probe the hash table for matches.
  * @param[in] build The table with which the hash table was built.
+ * @param[in] hash_probe The hasher used for the probe table.
  * @param[in] equality_probe The equality comparator used when probing the hash table.
  * @param[in] join_type The type of join to be performed
  * @param[in] hash_table_view The hash table built from `build`.
@@ -105,7 +108,8 @@ __global__ void mixed_join(table_device_view left_table,
                            table_device_view right_table,
                            table_device_view probe,
                            table_device_view build,
-                           row_equality const equality_probe,
+                           row_hash_experimental const hash_probe,
+                           row_equality_experimental const equality_probe,
                            join_kind const join_type,
                            cudf::detail::mixed_multimap_type::device_view hash_table_view,
                            size_type* join_output_l,

cpp/src/join/mixed_join_kernels_semi.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -37,7 +37,8 @@ __launch_bounds__(block_size) __global__
                        table_device_view right_table,
                        table_device_view probe,
                        table_device_view build,
-                       row_equality const equality_probe,
+                       row_hash_experimental const hash_probe,
+                       row_equality_experimental const equality_probe,
                        join_kind const join_type,
                        cudf::detail::semi_map_type::device_view hash_table_view,
                        size_type* join_output_l,
@@ -64,8 +65,6 @@ __launch_bounds__(block_size) __global__
   auto evaluator = cudf::ast::detail::expression_evaluator<has_nulls>(
     left_table, right_table, device_expression_data);
 
-  row_hash hash_probe{nullate::DYNAMIC{has_nulls}, probe};
-
   if (outer_row_index < outer_num_rows) {
     // Figure out the number of elements for this key.
     auto equality = single_expression_equality<has_nulls>{
@@ -83,7 +82,8 @@ template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, true>(
   table_device_view right_table,
   table_device_view probe,
   table_device_view build,
-  row_equality const equality_probe,
+  row_hash_experimental const hash_probe,
+  row_equality_experimental const equality_probe,
   join_kind const join_type,
   cudf::detail::semi_map_type::device_view hash_table_view,
   size_type* join_output_l,
@@ -96,7 +96,8 @@ template __global__ void mixed_join_semi<DEFAULT_JOIN_BLOCK_SIZE, false>(
   table_device_view right_table,
   table_device_view probe,
   table_device_view build,
-  row_equality const equality_probe,
+  row_hash_experimental const hash_probe,
+  row_equality_experimental const equality_probe,
   join_kind const join_type,
   cudf::detail::semi_map_type::device_view hash_table_view,
   size_type* join_output_l,

cpp/src/join/mixed_join_kernels_semi.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2022, NVIDIA CORPORATION.
+ * Copyright (c) 2022-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -41,6 +41,7 @@ namespace detail {
  * @param[in] right_table The right table
  * @param[in] probe The table with which to probe the hash table for matches.
  * @param[in] build The table with which the hash table was built.
+ * @param[in] hash_probe The hasher used for the probe table.
  * @param[in] equality_probe The equality comparator used when probing the hash table.
  * @param[in] join_type The type of join to be performed
  * @param[in] hash_table_view The hash table built from `build`.
@@ -62,7 +63,8 @@ __global__ void compute_mixed_join_output_size_semi(
   table_device_view right_table,
   table_device_view probe,
   table_device_view build,
-  row_equality const equality_probe,
+  row_hash_experimental const hash_probe,
+  row_equality_experimental const equality_probe,
   join_kind const join_type,
   cudf::detail::semi_map_type::device_view hash_table_view,
   ast::detail::expression_device_view device_expression_data,
@@ -87,6 +89,7 @@ __global__ void compute_mixed_join_output_size_semi(
  * @param[in] right_table The right table
  * @param[in] probe The table with which to probe the hash table for matches.
  * @param[in] build The table with which the hash table was built.
+ * @param[in] hash_probe The hasher used for the probe table.
  * @param[in] equality_probe The equality comparator used when probing the hash table.
  * @param[in] join_type The type of join to be performed
  * @param[in] hash_table_view The hash table built from `build`.
@@ -104,7 +107,8 @@ __global__ void mixed_join_semi(table_device_view left_table,
                                 table_device_view right_table,
                                 table_device_view probe,
                                 table_device_view build,
-                                row_equality const equality_probe,
+                                row_hash_experimental const hash_probe,
+                                row_equality_experimental const equality_probe,
                                 join_kind const join_type,
                                 cudf::detail::semi_map_type::device_view hash_table_view,
                                 size_type* join_output_l,