Implement mixed equality/conditional joins

cpp/include/cudf/ast/detail/expression_evaluator.cuh

@@ -426,9 +426,10 @@ struct expression_evaluator {
    * @param row_index Row index of all input and output data column(s).
    */
   template <typename ResultSubclass, typename T, bool result_has_nulls>
-  CUDF_DFI void evaluate(expression_result<ResultSubclass, T, result_has_nulls>& output_object,
-                         cudf::size_type const row_index,
-                         IntermediateDataType<has_nulls>* thread_intermediate_storage) const
+  __device__ __forceinline__ void evaluate(
+    expression_result<ResultSubclass, T, result_has_nulls>& output_object,
+    cudf::size_type const row_index,
+    IntermediateDataType<has_nulls>* thread_intermediate_storage) const
   {
     evaluate(output_object, row_index, row_index, row_index, thread_intermediate_storage);
   }
@@ -446,11 +447,12 @@ struct expression_evaluator {
    * @param output_row_index The row in the output to insert the result.
    */
   template <typename ResultSubclass, typename T, bool result_has_nulls>
-  CUDF_DFI void evaluate(expression_result<ResultSubclass, T, result_has_nulls>& output_object,
-                         cudf::size_type const left_row_index,
-                         cudf::size_type const right_row_index,
-                         cudf::size_type const output_row_index,
-                         IntermediateDataType<has_nulls>* thread_intermediate_storage) const
+  __device__ __forceinline__ void evaluate(
+    expression_result<ResultSubclass, T, result_has_nulls>& output_object,
+    cudf::size_type const left_row_index,
+    cudf::size_type const right_row_index,
+    cudf::size_type const output_row_index,
+    IntermediateDataType<has_nulls>* thread_intermediate_storage) const
   {
     cudf::size_type operator_source_index{0};
     for (cudf::size_type operator_index = 0; operator_index < plan.operators.size();

cpp/include/cudf/join.hpp

@@ -897,24 +897,25 @@ std::unique_ptr<rmm::device_uvector<size_type>> conditional_left_anti_join(
  * @param right_conditional The right table used for the conditional join.
  * @param binary_predicate The condition on which to join.
  * @param compare_nulls Whether or not null values join to each other or not.
- * @param output_size An optional pair of values indicating the exact output size and the number of
- * matches for each row in the larger of the two input tables, left or right (may be precomputed
- * using the corresponding mixed_inner_join_size API).
+ * @param output_size_data An optional pair of values indicating the exact output size and the
+ * number of matches for each row in the larger of the two input tables, left or right (may be
+ * precomputed using the corresponding mixed_inner_join_size API).
  * @param mr Device memory resource used to allocate the returned table and columns' device memory
  *
  * @return A pair of vectors [`left_indices`, `right_indices`] that can be used to construct
  * the result of performing a mixed inner join between the four input tables.
  */
 std::pair<std::unique_ptr<rmm::device_uvector<size_type>>,
           std::unique_ptr<rmm::device_uvector<size_type>>>
-mixed_inner_join(table_view const& left_equality,
-                 table_view const& right_equality,
-                 table_view const& left_conditional,
-                 table_view const& right_conditional,
-                 ast::expression const& binary_predicate,
-                 null_equality compare_nulls = null_equality::EQUAL,
-                 std::optional<std::pair<std::size_t, column_view>> output_size_data = {},
-                 rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+mixed_inner_join(
+  table_view const& left_equality,
+  table_view const& right_equality,
+  table_view const& left_conditional,
+  table_view const& right_conditional,
+  ast::expression const& binary_predicate,
+  null_equality compare_nulls = null_equality::EQUAL,
+  std::optional<std::pair<std::size_t, device_span<size_type>>> output_size_data = {},
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Returns a pair of row index vectors corresponding to all pairs of
@@ -956,9 +957,9 @@ mixed_inner_join(table_view const& left_equality,
  * @param right_conditional The right table used for the conditional join.
  * @param binary_predicate The condition on which to join.
  * @param compare_nulls Whether or not null values join to each other or not.
- * @param output_size An optional pair of values indicating the exact output size and the number of
- * matches for each row in the larger of the two input tables, left or right (may be precomputed
- * using the corresponding mixed_left_join_size API).
+ * @param output_size_data An optional pair of values indicating the exact output size and the
+ * number of matches for each row in the larger of the two input tables, left or right (may be
+ * precomputed using the corresponding mixed_left_join_size API).
  * @param mr Device memory resource used to allocate the returned table and columns' device memory
  *
  * @return A pair of vectors [`left_indices`, `right_indices`] that can be used to construct
@@ -972,7 +973,7 @@ mixed_left_join(table_view const& left_equality,
                 table_view const& right_conditional,
                 ast::expression const& binary_predicate,
                 null_equality compare_nulls = null_equality::EQUAL,
-                std::optional<std::pair<std::size_t, column_view>> output_size_data = {},
+                std::optional<std::pair<std::size_t, device_span<size_type>>> output_size_data = {},
                 rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
@@ -1015,9 +1016,9 @@ mixed_left_join(table_view const& left_equality,
  * @param right_conditional The right table used for the conditional join.
  * @param binary_predicate The condition on which to join.
  * @param compare_nulls Whether or not null values join to each other or not.
- * @param output_size An optional pair of values indicating the exact output size and the number of
- * matches for each row in the larger of the two input tables, left or right (may be precomputed
- * using the corresponding mixed_full_join_size API).
+ * @param output_size_data An optional pair of values indicating the exact output size and the
+ * number of matches for each row in the larger of the two input tables, left or right (may be
+ * precomputed using the corresponding mixed_full_join_size API).
  * @param mr Device memory resource used to allocate the returned table and columns' device memory
  *
  * @return A pair of vectors [`left_indices`, `right_indices`] that can be used to construct
@@ -1031,7 +1032,7 @@ mixed_full_join(table_view const& left_equality,
                 table_view const& right_conditional,
                 ast::expression const& binary_predicate,
                 null_equality compare_nulls = null_equality::EQUAL,
-                std::optional<std::pair<std::size_t, column_view>> output_size_data = {},
+                std::optional<std::pair<std::size_t, device_span<size_type>>> output_size_data = {},
                 rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
@@ -1068,7 +1069,7 @@ mixed_full_join(table_view const& left_equality,
  * be relied upon, simply passed to the corresponding `mixed_inner_join` API as
  * is.
  */
-std::pair<std::size_t, std::unique_ptr<column>> mixed_inner_join_size(
+std::pair<std::size_t, std::unique_ptr<rmm::device_uvector<size_type>>> mixed_inner_join_size(
   table_view const& left_equality,
   table_view const& right_equality,
   table_view const& left_conditional,
@@ -1111,7 +1112,7 @@ std::pair<std::size_t, std::unique_ptr<column>> mixed_inner_join_size(
  * be relied upon, simply passed to the corresponding `mixed_left_join` API as
  * is.
  */
-std::pair<std::size_t, std::unique_ptr<column>> mixed_left_join_size(
+std::pair<std::size_t, std::unique_ptr<rmm::device_uvector<size_type>>> mixed_left_join_size(
   table_view const& left_equality,
   table_view const& right_equality,
   table_view const& left_conditional,

cpp/include/cudf/table/row_operators.cuh

@@ -492,11 +492,11 @@ template <template <typename> class hash_function, typename Nullate>
 class row_hasher {
  public:
   row_hasher() = delete;
-  CUDA_HOST_DEVICE_CALLABLE row_hasher(Nullate has_nulls, table_device_view t)
+  CUDF_HOST_DEVICE row_hasher(Nullate has_nulls, table_device_view t)
     : _table{t}, _has_nulls{has_nulls}
   {
   }
-  CUDA_HOST_DEVICE_CALLABLE row_hasher(Nullate has_nulls, table_device_view t, uint32_t seed)
+  CUDF_HOST_DEVICE row_hasher(Nullate has_nulls, table_device_view t, uint32_t seed)
     : _table{t}, _seed(seed), _has_nulls{has_nulls}
   {
   }

cpp/src/join/hash_join.cuh

@@ -58,8 +58,8 @@ constexpr auto remap_sentinel_hash(H hash, S sentinel)
  */
 class make_pair_function {
  public:
-  CUDA_HOST_DEVICE_CALLABLE make_pair_function(row_hash const& hash,
-                                               hash_value_type const empty_key_sentinel)
+  CUDF_HOST_DEVICE make_pair_function(row_hash const& hash,
+                                      hash_value_type const empty_key_sentinel)
     : _hash{hash}, _empty_key_sentinel{empty_key_sentinel}
   {
   }