From cca009192dbd67616ffb53ab1d708a1d12868625 Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Fri, 20 Sep 2024 02:47:53 +0000
Subject: [PATCH 01/11] Reapply Refactor mixed_semi_join using cuco::static_set

---
 cpp/src/join/join_common_utils.hpp       |  6 --
 cpp/src/join/mixed_join_common_utils.cuh | 33 +++++++++
 cpp/src/join/mixed_join_kernels_semi.cu  | 35 ++++-----
 cpp/src/join/mixed_join_kernels_semi.cuh |  6 +-
 cpp/src/join/mixed_join_semi.cu          | 90 +++++++-----------------
 cpp/tests/join/mixed_join_tests.cu       | 30 ++++++++
 6 files changed, 109 insertions(+), 91 deletions(-)

diff --git a/cpp/src/join/join_common_utils.hpp b/cpp/src/join/join_common_utils.hpp
index 86402a0e7de..573101cefd9 100644
--- a/cpp/src/join/join_common_utils.hpp
+++ b/cpp/src/join/join_common_utils.hpp
@@ -22,7 +22,6 @@
 #include <cudf/table/row_operators.cuh>
 #include <cudf/table/table_view.hpp>
 
-#include <cuco/static_map.cuh>
 #include <cuco/static_multimap.cuh>
 #include <cuda/atomic>
 
@@ -51,11 +50,6 @@ using mixed_multimap_type =
                         cudf::detail::cuco_allocator<char>,
                         cuco::legacy::double_hashing<1, hash_type, hash_type>>;
 
-using semi_map_type = cuco::legacy::static_map<hash_value_type,
-                                               size_type,
-                                               cuda::thread_scope_device,
-                                               cudf::detail::cuco_allocator<char>>;
-
 using row_hash_legacy =
   cudf::row_hasher<cudf::hashing::detail::default_hash, cudf::nullate::DYNAMIC>;
 
diff --git a/cpp/src/join/mixed_join_common_utils.cuh b/cpp/src/join/mixed_join_common_utils.cuh
index 19701816867..89c13285cfe 100644
--- a/cpp/src/join/mixed_join_common_utils.cuh
+++ b/cpp/src/join/mixed_join_common_utils.cuh
@@ -25,6 +25,7 @@
 #include <rmm/device_uvector.hpp>
 
 #include <cub/cub.cuh>
+#include <cuco/static_set.cuh>
 
 namespace cudf {
 namespace detail {
@@ -160,6 +161,38 @@ struct pair_expression_equality : public expression_equality<has_nulls> {
   }
 };
 
+/**
+ * @brief Equality comparator that composes two row_equality comparators.
+ */
+struct double_row_equality_comparator {
+  row_equality const equality_comparator;
+  row_equality const conditional_comparator;
+
+  __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept
+  {
+    using experimental::row::lhs_index_type;
+    using experimental::row::rhs_index_type;
+
+    return equality_comparator(lhs_index_type{lhs_row_index}, rhs_index_type{rhs_row_index}) &&
+           conditional_comparator(lhs_index_type{lhs_row_index}, rhs_index_type{rhs_row_index});
+  }
+};
+
+// A CUDA Cooperative Group of 4 threads for the hash set.
+auto constexpr DEFAULT_MIXED_JOIN_CG_SIZE = 4;
+
+// The hash set type used by mixed_semi_join with the build_table.
+using hash_set_type = cuco::static_set<size_type,
+                                       cuco::extent<size_t>,
+                                       cuda::thread_scope_device,
+                                       double_row_equality_comparator,
+                                       cuco::linear_probing<DEFAULT_MIXED_JOIN_CG_SIZE, row_hash>,
+                                       cudf::detail::cuco_allocator<char>,
+                                       cuco::storage<1>>;
+
+// The hash_set_ref_type used by mixed_semi_join kerenels for probing.
+using hash_set_ref_type = hash_set_type::ref_type<cuco::contains_tag>;
+
 }  // namespace detail
 
 }  // namespace cudf
diff --git a/cpp/src/join/mixed_join_kernels_semi.cu b/cpp/src/join/mixed_join_kernels_semi.cu
index 7459ac3e99c..f2c5ff13638 100644
--- a/cpp/src/join/mixed_join_kernels_semi.cu
+++ b/cpp/src/join/mixed_join_kernels_semi.cu
@@ -38,12 +38,16 @@ CUDF_KERNEL void __launch_bounds__(block_size)
                   table_device_view right_table,
                   table_device_view probe,
                   table_device_view build,
-                  row_hash const hash_probe,
                   row_equality const equality_probe,
-                  cudf::detail::semi_map_type::device_view hash_table_view,
+                  hash_set_ref_type set_ref,
                   cudf::device_span<bool> left_table_keep_mask,
                   cudf::ast::detail::expression_device_view device_expression_data)
 {
+  auto constexpr cg_size = hash_set_ref_type::cg_size;
+
+  auto const tile =
+    cooperative_groups::tiled_partition<cg_size>(cooperative_groups::this_thread_block());
+
   // Normally the casting of a shared memory array is used to create multiple
   // arrays of different types from the shared memory buffer, but here it is
   // used to circumvent conflicts between arrays of different types between
@@ -52,24 +56,24 @@ CUDF_KERNEL void __launch_bounds__(block_size)
   cudf::ast::detail::IntermediateDataType<has_nulls>* intermediate_storage =
     reinterpret_cast<cudf::ast::detail::IntermediateDataType<has_nulls>*>(raw_intermediate_storage);
   auto thread_intermediate_storage =
-    &intermediate_storage[threadIdx.x * device_expression_data.num_intermediates];
-
-  cudf::size_type const left_num_rows  = left_table.num_rows();
-  cudf::size_type const right_num_rows = right_table.num_rows();
-  auto const outer_num_rows            = left_num_rows;
+    &intermediate_storage[tile.meta_group_rank() * device_expression_data.num_intermediates];
 
-  cudf::size_type outer_row_index = threadIdx.x + blockIdx.x * block_size;
+  cudf::size_type const outer_num_rows = left_table.num_rows();
+  auto const outer_row_index = cudf::detail::grid_1d::global_thread_id<block_size>() / cg_size;
 
   auto evaluator = cudf::ast::detail::expression_evaluator<has_nulls>(
     left_table, right_table, device_expression_data);
 
   if (outer_row_index < outer_num_rows) {
+    // Make sure to swap_tables here as hash_set will use probe table as the left one.
+    auto constexpr swap_tables = true;
     // Figure out the number of elements for this key.
     auto equality = single_expression_equality<has_nulls>{
-      evaluator, thread_intermediate_storage, false, equality_probe};
+      evaluator, thread_intermediate_storage, swap_tables, equality_probe};
 
-    left_table_keep_mask[outer_row_index] =
-      hash_table_view.contains(outer_row_index, hash_probe, equality);
+    auto const set_ref_equality = set_ref.with_key_eq(equality);
+    auto const result           = set_ref_equality.contains(tile, outer_row_index);
+    if (tile.thread_rank() == 0) left_table_keep_mask[outer_row_index] = result;
   }
 }
 
@@ -78,9 +82,8 @@ void launch_mixed_join_semi(bool has_nulls,
                             table_device_view right_table,
                             table_device_view probe,
                             table_device_view build,
-                            row_hash const hash_probe,
                             row_equality const equality_probe,
-                            cudf::detail::semi_map_type::device_view hash_table_view,
+                            hash_set_ref_type set_ref,
                             cudf::device_span<bool> left_table_keep_mask,
                             cudf::ast::detail::expression_device_view device_expression_data,
                             detail::grid_1d const config,
@@ -94,9 +97,8 @@ void launch_mixed_join_semi(bool has_nulls,
         right_table,
         probe,
         build,
-        hash_probe,
         equality_probe,
-        hash_table_view,
+        set_ref,
         left_table_keep_mask,
         device_expression_data);
   } else {
@@ -106,9 +108,8 @@ void launch_mixed_join_semi(bool has_nulls,
         right_table,
         probe,
         build,
-        hash_probe,
         equality_probe,
-        hash_table_view,
+        set_ref,
         left_table_keep_mask,
         device_expression_data);
   }
diff --git a/cpp/src/join/mixed_join_kernels_semi.cuh b/cpp/src/join/mixed_join_kernels_semi.cuh
index 43714ffb36a..b08298e64e4 100644
--- a/cpp/src/join/mixed_join_kernels_semi.cuh
+++ b/cpp/src/join/mixed_join_kernels_semi.cuh
@@ -45,9 +45,8 @@ 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] hash_table_view The hash table built from `build`.
+ * @param[in] set_ref The hash table device view built from `build`.
  * @param[out] left_table_keep_mask The result of the join operation with "true" element indicating
  * the corresponding index from left table is present in output
  * @param[in] device_expression_data Container of device data required to evaluate the desired
@@ -58,9 +57,8 @@ void launch_mixed_join_semi(bool has_nulls,
                             table_device_view right_table,
                             table_device_view probe,
                             table_device_view build,
-                            row_hash const hash_probe,
                             row_equality const equality_probe,
-                            cudf::detail::semi_map_type::device_view hash_table_view,
+                            hash_set_ref_type set_ref,
                             cudf::device_span<bool> left_table_keep_mask,
                             cudf::ast::detail::expression_device_view device_expression_data,
                             detail::grid_1d const config,
diff --git a/cpp/src/join/mixed_join_semi.cu b/cpp/src/join/mixed_join_semi.cu
index cfb785e242c..719b1d47105 100644
--- a/cpp/src/join/mixed_join_semi.cu
+++ b/cpp/src/join/mixed_join_semi.cu
@@ -46,45 +46,6 @@
 namespace cudf {
 namespace detail {
 
-namespace {
-/**
- * @brief Device functor to create a pair of hash value and index for a given row.
- */
-struct make_pair_function_semi {
-  __device__ __forceinline__ cudf::detail::pair_type operator()(size_type i) const noexcept
-  {
-    // The value is irrelevant since we only ever use the hash map to check for
-    // membership of a particular row index.
-    return cuco::make_pair(static_cast<hash_value_type>(i), 0);
-  }
-};
-
-/**
- * @brief Equality comparator that composes two row_equality comparators.
- */
-class double_row_equality {
- public:
-  double_row_equality(row_equality equality_comparator, row_equality conditional_comparator)
-    : _equality_comparator{equality_comparator}, _conditional_comparator{conditional_comparator}
-  {
-  }
-
-  __device__ bool operator()(size_type lhs_row_index, size_type rhs_row_index) const noexcept
-  {
-    using experimental::row::lhs_index_type;
-    using experimental::row::rhs_index_type;
-
-    return _equality_comparator(lhs_index_type{lhs_row_index}, rhs_index_type{rhs_row_index}) &&
-           _conditional_comparator(lhs_index_type{lhs_row_index}, rhs_index_type{rhs_row_index});
-  }
-
- private:
-  row_equality _equality_comparator;
-  row_equality _conditional_comparator;
-};
-
-}  // namespace
-
 std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
   table_view const& left_equality,
   table_view const& right_equality,
@@ -96,7 +57,7 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
   rmm::cuda_stream_view stream,
   rmm::device_async_resource_ref mr)
 {
-  CUDF_EXPECTS((join_type != join_kind::INNER_JOIN) && (join_type != join_kind::LEFT_JOIN) &&
+  CUDF_EXPECTS((join_type != join_kind::INNER_JOIN) and (join_type != join_kind::LEFT_JOIN) and
                  (join_type != join_kind::FULL_JOIN),
                "Inner, left, and full joins should use mixed_join.");
 
@@ -137,7 +98,7 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
   // output column and follow the null-supporting expression evaluation code
   // path.
   auto const has_nulls = cudf::nullate::DYNAMIC{
-    cudf::has_nulls(left_equality) || cudf::has_nulls(right_equality) ||
+    cudf::has_nulls(left_equality) or cudf::has_nulls(right_equality) or
     binary_predicate.may_evaluate_null(left_conditional, right_conditional, stream)};
 
   auto const parser = ast::detail::expression_parser{
@@ -156,27 +117,20 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
   auto right_conditional_view = table_device_view::create(right_conditional, stream);
 
   auto const preprocessed_build =
-    experimental::row::equality::preprocessed_table::create(build, stream);
+    cudf::experimental::row::equality::preprocessed_table::create(build, stream);
   auto const preprocessed_probe =
-    experimental::row::equality::preprocessed_table::create(probe, stream);
+    cudf::experimental::row::equality::preprocessed_table::create(probe, stream);
   auto const row_comparator =
-    cudf::experimental::row::equality::two_table_comparator{preprocessed_probe, preprocessed_build};
+    cudf::experimental::row::equality::two_table_comparator{preprocessed_build, preprocessed_probe};
   auto const equality_probe = row_comparator.equal_to<false>(has_nulls, compare_nulls);
 
-  semi_map_type hash_table{
-    compute_hash_table_size(build.num_rows()),
-    cuco::empty_key{std::numeric_limits<hash_value_type>::max()},
-    cuco::empty_value{cudf::detail::JoinNoneValue},
-    cudf::detail::cuco_allocator<char>{rmm::mr::polymorphic_allocator<char>{}, stream},
-    stream.value()};
-
   // Create hash table containing all keys found in right table
   // TODO: To add support for nested columns we will need to flatten in many
   // places. However, this probably isn't worth adding any time soon since we
   // won't be able to support AST conditions for those types anyway.
   auto const build_nulls    = cudf::nullate::DYNAMIC{cudf::has_nulls(build)};
   auto const row_hash_build = cudf::experimental::row::hash::row_hasher{preprocessed_build};
-  auto const hash_build     = row_hash_build.device_hasher(build_nulls);
+
   // Since we may see multiple rows that are identical in the equality tables
   // but differ in the conditional tables, the equality comparator used for
   // insertion must account for both sets of tables. An alternative solution
@@ -191,20 +145,28 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
   auto const equality_build_equality =
     row_comparator_build.equal_to<false>(build_nulls, compare_nulls);
   auto const preprocessed_build_condtional =
-    experimental::row::equality::preprocessed_table::create(right_conditional, stream);
+    cudf::experimental::row::equality::preprocessed_table::create(right_conditional, stream);
   auto const row_comparator_conditional_build =
     cudf::experimental::row::equality::two_table_comparator{preprocessed_build_condtional,
                                                             preprocessed_build_condtional};
   auto const equality_build_conditional =
     row_comparator_conditional_build.equal_to<false>(build_nulls, compare_nulls);
-  double_row_equality equality_build{equality_build_equality, equality_build_conditional};
-  make_pair_function_semi pair_func_build{};
 
-  auto iter = cudf::detail::make_counting_transform_iterator(0, pair_func_build);
+  hash_set_type row_set{
+    {compute_hash_table_size(build.num_rows())},
+    cuco::empty_key{JoinNoneValue},
+    {equality_build_equality, equality_build_conditional},
+    {row_hash_build.device_hasher(build_nulls)},
+    {},
+    {},
+    cudf::detail::cuco_allocator<char>{rmm::mr::polymorphic_allocator<char>{}, stream},
+    {stream.value()}};
+
+  auto iter = thrust::make_counting_iterator(0);
 
   // skip rows that are null here.
   if ((compare_nulls == null_equality::EQUAL) or (not nullable(build))) {
-    hash_table.insert(iter, iter + right_num_rows, hash_build, equality_build, stream.value());
+    row_set.insert(iter, iter + right_num_rows, stream.value());
   } else {
     thrust::counting_iterator<cudf::size_type> stencil(0);
     auto const [row_bitmask, _] =
@@ -212,18 +174,19 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
     row_is_valid pred{static_cast<bitmask_type const*>(row_bitmask.data())};
 
     // insert valid rows
-    hash_table.insert_if(
-      iter, iter + right_num_rows, stencil, pred, hash_build, equality_build, stream.value());
+    row_set.insert_if(iter, iter + right_num_rows, stencil, pred, stream.value());
   }
 
-  auto hash_table_view = hash_table.get_device_view();
-
   detail::grid_1d const config(outer_num_rows, DEFAULT_JOIN_BLOCK_SIZE);
-  auto const shmem_size_per_block = parser.shmem_per_thread * config.num_threads_per_block;
+  auto const shmem_size_per_block =
+    parser.shmem_per_thread *
+    cuco::detail::int_div_ceil(config.num_threads_per_block, hash_set_type::cg_size);
 
   auto const row_hash   = cudf::experimental::row::hash::row_hasher{preprocessed_probe};
   auto const hash_probe = row_hash.device_hasher(has_nulls);
 
+  hash_set_ref_type const row_set_ref = row_set.ref(cuco::contains).with_hash_function(hash_probe);
+
   // Vector used to indicate indices from left/probe table which are present in output
   auto left_table_keep_mask = rmm::device_uvector<bool>(probe.num_rows(), stream);
 
@@ -232,9 +195,8 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
                          *right_conditional_view,
                          *probe_view,
                          *build_view,
-                         hash_probe,
                          equality_probe,
-                         hash_table_view,
+                         row_set_ref,
                          cudf::device_span<bool>(left_table_keep_mask),
                          parser.device_expression_data,
                          config,
diff --git a/cpp/tests/join/mixed_join_tests.cu b/cpp/tests/join/mixed_join_tests.cu
index 6c147c8a128..08a0136700d 100644
--- a/cpp/tests/join/mixed_join_tests.cu
+++ b/cpp/tests/join/mixed_join_tests.cu
@@ -778,6 +778,21 @@ TYPED_TEST(MixedLeftSemiJoinTest, BasicEquality)
              {1});
 }
 
+TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMap)
+{
+  auto const col_ref_left_1  = cudf::ast::column_reference(0, cudf::ast::table_reference::LEFT);
+  auto const col_ref_right_1 = cudf::ast::column_reference(0, cudf::ast::table_reference::RIGHT);
+  auto left_one_greater_right_one =
+    cudf::ast::operation(cudf::ast::ast_operator::GREATER, col_ref_left_1, col_ref_right_1);
+
+  this->test({{2, 3, 9, 0, 1, 7, 4, 6, 5, 8}, {1, 2, 3, 4, 5, 6, 7, 8, 9, 0}},
+             {{6, 5, 9, 8, 10, 32}, {0, 1, 2, 3, 4, 5}, {7, 8, 9, 0, 1, 2}},
+             {0},
+             {1},
+             left_one_greater_right_one,
+             {2, 7, 8});
+}
+
 TYPED_TEST(MixedLeftSemiJoinTest, BasicEqualityDuplicates)
 {
   this->test({{0, 1, 2, 1}, {3, 4, 5, 6}, {10, 20, 30, 40}},
@@ -900,3 +915,18 @@ TYPED_TEST(MixedLeftAntiJoinTest, AsymmetricLeftLargerEquality)
              left_zero_eq_right_zero,
              {0, 1, 3});
 }
+
+TYPED_TEST(MixedLeftAntiJoinTest, MixedLeftAntiJoinGatherMap)
+{
+  auto const col_ref_left_1  = cudf::ast::column_reference(0, cudf::ast::table_reference::LEFT);
+  auto const col_ref_right_1 = cudf::ast::column_reference(0, cudf::ast::table_reference::RIGHT);
+  auto left_one_greater_right_one =
+    cudf::ast::operation(cudf::ast::ast_operator::GREATER, col_ref_left_1, col_ref_right_1);
+
+  this->test({{2, 3, 9, 0, 1, 7, 4, 6, 5, 8}, {1, 2, 3, 4, 5, 6, 7, 8, 9, 0}},
+             {{6, 5, 9, 8, 10, 32}, {0, 1, 2, 3, 4, 5}, {7, 8, 9, 0, 1, 2}},
+             {0},
+             {1},
+             left_one_greater_right_one,
+             {0, 1, 3, 4, 5, 6, 9});
+}

From 36e1cd768a805fbb53b4528ab61f132d3105566e Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Fri, 20 Sep 2024 02:48:13 +0000
Subject: [PATCH 02/11] Fix for mixed_semi_join

---
 cpp/src/join/mixed_join_kernels_semi.cu |  21 +++--
 cpp/src/join/mixed_join_semi.cu         |   2 +-
 cpp/tests/join/mixed_join_tests.cu      | 110 ++++++++++++++++++++++++
 3 files changed, 123 insertions(+), 10 deletions(-)

diff --git a/cpp/src/join/mixed_join_kernels_semi.cu b/cpp/src/join/mixed_join_kernels_semi.cu
index f2c5ff13638..9f48717c322 100644
--- a/cpp/src/join/mixed_join_kernels_semi.cu
+++ b/cpp/src/join/mixed_join_kernels_semi.cu
@@ -59,20 +59,23 @@ CUDF_KERNEL void __launch_bounds__(block_size)
     &intermediate_storage[tile.meta_group_rank() * device_expression_data.num_intermediates];
 
   cudf::size_type const outer_num_rows = left_table.num_rows();
-  auto const outer_row_index = cudf::detail::grid_1d::global_thread_id<block_size>() / cg_size;
 
   auto evaluator = cudf::ast::detail::expression_evaluator<has_nulls>(
     left_table, right_table, device_expression_data);
 
-  if (outer_row_index < outer_num_rows) {
-    // Make sure to swap_tables here as hash_set will use probe table as the left one.
-    auto constexpr swap_tables = true;
-    // Figure out the number of elements for this key.
-    auto equality = single_expression_equality<has_nulls>{
-      evaluator, thread_intermediate_storage, swap_tables, equality_probe};
+  // Make sure to swap_tables here as hash_set will use probe table as the left one.
+  auto constexpr swap_tables = true;
+  auto equality              = single_expression_equality<has_nulls>{
+    evaluator, thread_intermediate_storage, swap_tables, equality_probe};
 
-    auto const set_ref_equality = set_ref.with_key_eq(equality);
-    auto const result           = set_ref_equality.contains(tile, outer_row_index);
+  // Create set ref with the new equality comparator.
+  auto const set_ref_equality = set_ref.with_key_eq(equality);
+
+  // Find all the rows in the left table that are in the hash table.
+  for (auto outer_row_index = cudf::detail::grid_1d::global_thread_id<block_size>() / cg_size;
+       outer_row_index < outer_num_rows;
+       outer_row_index += cudf::detail::grid_1d::grid_stride<block_size>() / cg_size) {
+    auto const result = set_ref_equality.contains(tile, outer_row_index);
     if (tile.thread_rank() == 0) left_table_keep_mask[outer_row_index] = result;
   }
 }
diff --git a/cpp/src/join/mixed_join_semi.cu b/cpp/src/join/mixed_join_semi.cu
index 719b1d47105..230e9978a6d 100644
--- a/cpp/src/join/mixed_join_semi.cu
+++ b/cpp/src/join/mixed_join_semi.cu
@@ -177,7 +177,7 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
     row_set.insert_if(iter, iter + right_num_rows, stencil, pred, stream.value());
   }
 
-  detail::grid_1d const config(outer_num_rows, DEFAULT_JOIN_BLOCK_SIZE);
+  detail::grid_1d const config(outer_num_rows * hash_set_type::cg_size, DEFAULT_JOIN_BLOCK_SIZE);
   auto const shmem_size_per_block =
     parser.shmem_per_thread *
     cuco::detail::int_div_ceil(config.num_threads_per_block, hash_set_type::cg_size);
diff --git a/cpp/tests/join/mixed_join_tests.cu b/cpp/tests/join/mixed_join_tests.cu
index 08a0136700d..8ab81d311aa 100644
--- a/cpp/tests/join/mixed_join_tests.cu
+++ b/cpp/tests/join/mixed_join_tests.cu
@@ -793,6 +793,116 @@ TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMap)
              {2, 7, 8});
 }
 
+TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMapLarge)
+{
+  using T = double;
+
+  auto const random_data = [](size_t size) {
+    std::vector<T> values(size);
+    using uniform_distribution =
+      typename std::conditional_t<std::is_same_v<T, bool>,
+                                  std::bernoulli_distribution,
+                                  std::conditional_t<std::is_floating_point_v<T>,
+                                                     std::uniform_real_distribution<T>,
+                                                     std::uniform_int_distribution<T>>>;
+
+    static constexpr auto seed = 0xf00d;
+    static std::mt19937 engine{seed};
+    static uniform_distribution dist{};
+    std::generate_n(values.begin(), size, [&]() { return T{dist(engine)}; });
+
+    return values;
+  };
+
+  auto const random_validity = [&](size_t size) {
+    std::vector<bool> validity(size);
+    std::generate_n(validity.begin(), size, [&]() {
+      constexpr auto seed = 0xcafe;
+      std::mt19937 engine{seed};
+      std::bernoulli_distribution dist{};
+      return dist(engine);
+    });
+    return validity;
+  };
+
+  std::vector<std::pair<std::vector<T>, std::vector<bool>>> lefts = {
+    {random_data(500), random_validity(500)}, {random_data(500), random_validity(500)}};
+  std::vector<cudf::test::fixed_width_column_wrapper<T>> left_wrappers;
+  std::vector<cudf::column_view> left_columns;
+  for (auto [data, valids] : lefts) {
+    left_wrappers.emplace_back(
+      cudf::test::fixed_width_column_wrapper<T>(data.begin(), data.end(), valids.begin()));
+    left_columns.emplace_back(left_wrappers.back());
+  };
+
+  std::vector<std::pair<std::vector<T>, std::vector<bool>>> rights = {
+    {random_data(250), random_validity(250)}, {random_data(250), random_validity(250)}};
+  std::vector<cudf::test::fixed_width_column_wrapper<T>> right_wrappers;
+  std::vector<cudf::column_view> right_columns;
+  for (auto [data, valids] : rights) {
+    right_wrappers.emplace_back(
+      cudf::test::fixed_width_column_wrapper<T>(data.begin(), data.end(), valids.begin()));
+    right_columns.emplace_back(left_wrappers.back());
+  };
+
+  // Left and right table views.
+  auto const left_table  = cudf::table_view{left_columns};
+  auto const right_table = cudf::table_view{right_columns};
+
+  // Use the zeroth column for equality.
+  auto const left_equality  = left_table.select({0});
+  auto const right_equality = right_table.select({0});
+
+  // Column references for equality column.
+  auto const col_ref_left_0  = cudf::ast::column_reference(0, cudf::ast::table_reference::LEFT);
+  auto const col_ref_right_0 = cudf::ast::column_reference(0, cudf::ast::table_reference::RIGHT);
+  auto left_zero_eq_right_zero =
+    cudf::ast::operation(cudf::ast::ast_operator::EQUAL, col_ref_left_0, col_ref_right_0);
+
+  // Expected size of left_semi_join with only zeroth column equality.
+  auto const expected_num_idx_left_zero_eq_right_zero =
+    cudf::conditional_left_semi_join_size(left_table, right_table, left_zero_eq_right_zero);
+
+  // Actual size of mixed_left_semi_join with only zeroth column equality.
+  auto const num_idx_left_zero_eq_right_zero =
+    cudf::mixed_left_semi_join(left_equality,
+                               right_equality,
+                               left_table,
+                               right_table,
+                               left_zero_eq_right_zero,
+                               cudf::null_equality::UNEQUAL)
+      ->size();
+
+  // Expected and actual sizes must match.
+  EXPECT_EQ(expected_num_idx_left_zero_eq_right_zero, num_idx_left_zero_eq_right_zero);
+
+  // Common column references for conditional column.
+  auto const col_ref_left_1  = cudf::ast::column_reference(1, cudf::ast::table_reference::LEFT);
+  auto const col_ref_right_1 = cudf::ast::column_reference(1, cudf::ast::table_reference::RIGHT);
+  auto left_one_gt_right_one =
+    cudf::ast::operation(cudf::ast::ast_operator::GREATER, col_ref_left_1, col_ref_right_1);
+
+  auto combined_condition = cudf::ast::operation(
+    cudf::ast::ast_operator::LOGICAL_AND, left_zero_eq_right_zero, left_one_gt_right_one);
+
+  // Expected size of left_semi_join with zeroth col equality and first col conditional.
+  auto const expected_num_idx_left_one_greater_right_one =
+    cudf::conditional_left_semi_join_size(left_table, right_table, combined_condition);
+
+  // Actual size of left_semi_join with zeroth col equality and first col conditional.
+  auto const num_idx_left_one_greater_right_one =
+    cudf::mixed_left_semi_join(left_equality,
+                               right_equality,
+                               left_table,
+                               right_table,
+                               left_one_gt_right_one,
+                               cudf::null_equality::UNEQUAL)
+      ->size();
+
+  // Expected and actual sizes must match.
+  EXPECT_EQ(expected_num_idx_left_one_greater_right_one, num_idx_left_one_greater_right_one);
+}
+
 TYPED_TEST(MixedLeftSemiJoinTest, BasicEqualityDuplicates)
 {
   this->test({{0, 1, 2, 1}, {3, 4, 5, 6}, {10, 20, 30, 40}},

From 66b3f889b016e0af483d1a68ac8da9c1b6c16283 Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Fri, 20 Sep 2024 04:15:42 +0000
Subject: [PATCH 03/11] Perf tuning.

---
 cpp/src/join/mixed_join_common_utils.cuh | 17 ++++++-----
 cpp/tests/join/mixed_join_tests.cu       | 37 ++++++++++++------------
 2 files changed, 27 insertions(+), 27 deletions(-)

diff --git a/cpp/src/join/mixed_join_common_utils.cuh b/cpp/src/join/mixed_join_common_utils.cuh
index 89c13285cfe..aed93e20823 100644
--- a/cpp/src/join/mixed_join_common_utils.cuh
+++ b/cpp/src/join/mixed_join_common_utils.cuh
@@ -179,16 +179,17 @@ struct double_row_equality_comparator {
 };
 
 // A CUDA Cooperative Group of 4 threads for the hash set.
-auto constexpr DEFAULT_MIXED_JOIN_CG_SIZE = 4;
+auto constexpr DEFAULT_MIXED_SEMI_JOIN_CG_SIZE = 1;
 
 // The hash set type used by mixed_semi_join with the build_table.
-using hash_set_type = cuco::static_set<size_type,
-                                       cuco::extent<size_t>,
-                                       cuda::thread_scope_device,
-                                       double_row_equality_comparator,
-                                       cuco::linear_probing<DEFAULT_MIXED_JOIN_CG_SIZE, row_hash>,
-                                       cudf::detail::cuco_allocator<char>,
-                                       cuco::storage<1>>;
+using hash_set_type =
+  cuco::static_set<size_type,
+                   cuco::extent<size_t>,
+                   cuda::thread_scope_device,
+                   double_row_equality_comparator,
+                   cuco::linear_probing<DEFAULT_MIXED_SEMI_JOIN_CG_SIZE, row_hash>,
+                   cudf::detail::cuco_allocator<char>,
+                   cuco::storage<1>>;
 
 // The hash_set_ref_type used by mixed_semi_join kerenels for probing.
 using hash_set_ref_type = hash_set_type::ref_type<cuco::contains_tag>;
diff --git a/cpp/tests/join/mixed_join_tests.cu b/cpp/tests/join/mixed_join_tests.cu
index 8ab81d311aa..d43bf4dfca9 100644
--- a/cpp/tests/join/mixed_join_tests.cu
+++ b/cpp/tests/join/mixed_join_tests.cu
@@ -795,53 +795,52 @@ TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMap)
 
 TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMapLarge)
 {
-  using T = double;
+  using T1 = double;
 
   auto const random_data = [](size_t size) {
-    std::vector<T> values(size);
+    std::vector<T1> values(size);
     using uniform_distribution =
-      typename std::conditional_t<std::is_same_v<T, bool>,
+      typename std::conditional_t<std::is_same_v<T1, bool>,
                                   std::bernoulli_distribution,
-                                  std::conditional_t<std::is_floating_point_v<T>,
-                                                     std::uniform_real_distribution<T>,
-                                                     std::uniform_int_distribution<T>>>;
+                                  std::conditional_t<std::is_floating_point_v<T1>,
+                                                     std::uniform_real_distribution<T1>,
+                                                     std::uniform_int_distribution<T1>>>;
 
     static constexpr auto seed = 0xf00d;
     static std::mt19937 engine{seed};
     static uniform_distribution dist{};
-    std::generate_n(values.begin(), size, [&]() { return T{dist(engine)}; });
+    std::generate(values.begin(), values.end(), [&]() { return T1{dist(engine)}; });
 
     return values;
   };
 
   auto const random_validity = [&](size_t size) {
     std::vector<bool> validity(size);
-    std::generate_n(validity.begin(), size, [&]() {
-      constexpr auto seed = 0xcafe;
-      std::mt19937 engine{seed};
-      std::bernoulli_distribution dist{};
-      return dist(engine);
-    });
+    static constexpr auto seed = 0xcafe;
+    static std::mt19937 engine{seed};
+    static std::bernoulli_distribution dist{};
+    std::generate(validity.begin(), validity.end(), [&]() { return dist(engine); });
+
     return validity;
   };
 
-  std::vector<std::pair<std::vector<T>, std::vector<bool>>> lefts = {
+  std::vector<std::pair<std::vector<T1>, std::vector<bool>>> lefts = {
     {random_data(500), random_validity(500)}, {random_data(500), random_validity(500)}};
-  std::vector<cudf::test::fixed_width_column_wrapper<T>> left_wrappers;
+  std::vector<cudf::test::fixed_width_column_wrapper<T1>> left_wrappers;
   std::vector<cudf::column_view> left_columns;
   for (auto [data, valids] : lefts) {
     left_wrappers.emplace_back(
-      cudf::test::fixed_width_column_wrapper<T>(data.begin(), data.end(), valids.begin()));
+      cudf::test::fixed_width_column_wrapper<T1>(data.begin(), data.end(), valids.begin()));
     left_columns.emplace_back(left_wrappers.back());
   };
 
-  std::vector<std::pair<std::vector<T>, std::vector<bool>>> rights = {
+  std::vector<std::pair<std::vector<T1>, std::vector<bool>>> rights = {
     {random_data(250), random_validity(250)}, {random_data(250), random_validity(250)}};
-  std::vector<cudf::test::fixed_width_column_wrapper<T>> right_wrappers;
+  std::vector<cudf::test::fixed_width_column_wrapper<T1>> right_wrappers;
   std::vector<cudf::column_view> right_columns;
   for (auto [data, valids] : rights) {
     right_wrappers.emplace_back(
-      cudf::test::fixed_width_column_wrapper<T>(data.begin(), data.end(), valids.begin()));
+      cudf::test::fixed_width_column_wrapper<T1>(data.begin(), data.end(), valids.begin()));
     right_columns.emplace_back(left_wrappers.back());
   };
 

From 32b1f2902ac912d76c3c2965af35c50bdeef77bf Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Fri, 20 Sep 2024 04:35:01 +0000
Subject: [PATCH 04/11] Minor improvement

---
 cpp/src/join/mixed_join_kernels_semi.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/src/join/mixed_join_kernels_semi.cu b/cpp/src/join/mixed_join_kernels_semi.cu
index 9f48717c322..2baf5279cdb 100644
--- a/cpp/src/join/mixed_join_kernels_semi.cu
+++ b/cpp/src/join/mixed_join_kernels_semi.cu
@@ -53,7 +53,7 @@ CUDF_KERNEL void __launch_bounds__(block_size)
   // used to circumvent conflicts between arrays of different types between
   // different template instantiations due to the extern specifier.
   extern __shared__ char raw_intermediate_storage[];
-  cudf::ast::detail::IntermediateDataType<has_nulls>* intermediate_storage =
+  auto intermediate_storage =
     reinterpret_cast<cudf::ast::detail::IntermediateDataType<has_nulls>*>(raw_intermediate_storage);
   auto thread_intermediate_storage =
     &intermediate_storage[tile.meta_group_rank() * device_expression_data.num_intermediates];

From 37c1e045f94fc2145d73b491a0e09f6bf2d714f0 Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Fri, 20 Sep 2024 04:46:50 +0000
Subject: [PATCH 05/11] Add missing consts

---
 cpp/src/join/mixed_join_kernels_semi.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/cpp/src/join/mixed_join_kernels_semi.cu b/cpp/src/join/mixed_join_kernels_semi.cu
index 2baf5279cdb..0e008eec0a6 100644
--- a/cpp/src/join/mixed_join_kernels_semi.cu
+++ b/cpp/src/join/mixed_join_kernels_semi.cu
@@ -60,12 +60,12 @@ CUDF_KERNEL void __launch_bounds__(block_size)
 
   cudf::size_type const outer_num_rows = left_table.num_rows();
 
-  auto evaluator = cudf::ast::detail::expression_evaluator<has_nulls>(
+  auto const evaluator = cudf::ast::detail::expression_evaluator<has_nulls>(
     left_table, right_table, device_expression_data);
 
   // Make sure to swap_tables here as hash_set will use probe table as the left one.
   auto constexpr swap_tables = true;
-  auto equality              = single_expression_equality<has_nulls>{
+  auto const equality        = single_expression_equality<has_nulls>{
     evaluator, thread_intermediate_storage, swap_tables, equality_probe};
 
   // Create set ref with the new equality comparator.

From 6063f756d3d0ec11bf84ef7f1aba5e40e0dce6fe Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Fri, 20 Sep 2024 04:57:42 +0000
Subject: [PATCH 06/11] Perf improvements

---
 cpp/src/join/mixed_join_semi.cu | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/cpp/src/join/mixed_join_semi.cu b/cpp/src/join/mixed_join_semi.cu
index 230e9978a6d..fd10b9b725d 100644
--- a/cpp/src/join/mixed_join_semi.cu
+++ b/cpp/src/join/mixed_join_semi.cu
@@ -166,7 +166,7 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
 
   // skip rows that are null here.
   if ((compare_nulls == null_equality::EQUAL) or (not nullable(build))) {
-    row_set.insert(iter, iter + right_num_rows, stream.value());
+    row_set.insert_async(iter, iter + right_num_rows, stream.value());
   } else {
     thrust::counting_iterator<cudf::size_type> stencil(0);
     auto const [row_bitmask, _] =
@@ -174,7 +174,7 @@ std::unique_ptr<rmm::device_uvector<size_type>> mixed_join_semi(
     row_is_valid pred{static_cast<bitmask_type const*>(row_bitmask.data())};
 
     // insert valid rows
-    row_set.insert_if(iter, iter + right_num_rows, stencil, pred, stream.value());
+    row_set.insert_if_async(iter, iter + right_num_rows, stencil, pred, stream.value());
   }
 
   detail::grid_1d const config(outer_num_rows * hash_set_type::cg_size, DEFAULT_JOIN_BLOCK_SIZE);

From 18c3e20656796fb98d43ef4d87d8f95be50e1756 Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Fri, 20 Sep 2024 18:37:18 +0000
Subject: [PATCH 07/11] MInor comments

---
 cpp/src/join/mixed_join_kernels_semi.cu | 21 ++++++++++++---------
 1 file changed, 12 insertions(+), 9 deletions(-)

diff --git a/cpp/src/join/mixed_join_kernels_semi.cu b/cpp/src/join/mixed_join_kernels_semi.cu
index 0e008eec0a6..e3c51d2a76a 100644
--- a/cpp/src/join/mixed_join_kernels_semi.cu
+++ b/cpp/src/join/mixed_join_kernels_semi.cu
@@ -45,8 +45,7 @@ CUDF_KERNEL void __launch_bounds__(block_size)
 {
   auto constexpr cg_size = hash_set_ref_type::cg_size;
 
-  auto const tile =
-    cooperative_groups::tiled_partition<cg_size>(cooperative_groups::this_thread_block());
+  auto const tile = cg::tiled_partition<cg_size>(cg::this_thread_block());
 
   // Normally the casting of a shared memory array is used to create multiple
   // arrays of different types from the shared memory buffer, but here it is
@@ -56,25 +55,29 @@ CUDF_KERNEL void __launch_bounds__(block_size)
   auto intermediate_storage =
     reinterpret_cast<cudf::ast::detail::IntermediateDataType<has_nulls>*>(raw_intermediate_storage);
   auto thread_intermediate_storage =
-    &intermediate_storage[tile.meta_group_rank() * device_expression_data.num_intermediates];
-
-  cudf::size_type const outer_num_rows = left_table.num_rows();
+    intermediate_storage + (tile.meta_group_rank() * device_expression_data.num_intermediates);
 
+  // Equality evaluator to use
   auto const evaluator = cudf::ast::detail::expression_evaluator<has_nulls>(
     left_table, right_table, device_expression_data);
 
-  // Make sure to swap_tables here as hash_set will use probe table as the left one.
+  // Make sure to swap_tables here as hash_set will use probe table as the left one
   auto constexpr swap_tables = true;
   auto const equality        = single_expression_equality<has_nulls>{
     evaluator, thread_intermediate_storage, swap_tables, equality_probe};
 
-  // Create set ref with the new equality comparator.
+  // Create set ref with the new equality comparator
   auto const set_ref_equality = set_ref.with_key_eq(equality);
 
-  // Find all the rows in the left table that are in the hash table.
+  // Total number of rows to query the set
+  auto const outer_num_rows = left_table.num_rows();
+  // Grid stride for the tile
+  auto const cg_grid_stride = cudf::detail::grid_1d::grid_stride<block_size>() / cg_size;
+
+  // Find all the rows in the left table that are in the hash table
   for (auto outer_row_index = cudf::detail::grid_1d::global_thread_id<block_size>() / cg_size;
        outer_row_index < outer_num_rows;
-       outer_row_index += cudf::detail::grid_1d::grid_stride<block_size>() / cg_size) {
+       outer_row_index += cg_grid_stride) {
     auto const result = set_ref_equality.contains(tile, outer_row_index);
     if (tile.thread_rank() == 0) left_table_keep_mask[outer_row_index] = result;
   }

From a2ddf7f0d6e8ec636eb5708f2658f518a35b4648 Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Tue, 24 Sep 2024 15:27:42 -0700
Subject: [PATCH 08/11] Update cpp/src/join/mixed_join_kernels_semi.cu

Co-authored-by: Yunsong Wang <yunsongw@nvidia.com>
---
 cpp/src/join/mixed_join_kernels_semi.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/src/join/mixed_join_kernels_semi.cu b/cpp/src/join/mixed_join_kernels_semi.cu
index e3c51d2a76a..bd8c80652a0 100644
--- a/cpp/src/join/mixed_join_kernels_semi.cu
+++ b/cpp/src/join/mixed_join_kernels_semi.cu
@@ -79,7 +79,7 @@ CUDF_KERNEL void __launch_bounds__(block_size)
        outer_row_index < outer_num_rows;
        outer_row_index += cg_grid_stride) {
     auto const result = set_ref_equality.contains(tile, outer_row_index);
-    if (tile.thread_rank() == 0) left_table_keep_mask[outer_row_index] = result;
+    if (tile.thread_rank() == 0) { left_table_keep_mask[outer_row_index] = result; }
   }
 }
 

From 4c4821a6cf6a68fd6208dc09c7609374e9ad924b Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Thu, 26 Sep 2024 02:46:16 +0000
Subject: [PATCH 09/11] Apply suggestions from reviewer comments

---
 cpp/src/join/mixed_join_common_utils.cuh |   2 +-
 cpp/tests/join/mixed_join_tests.cu       | 146 ++++++++++++-----------
 2 files changed, 78 insertions(+), 70 deletions(-)

diff --git a/cpp/src/join/mixed_join_common_utils.cuh b/cpp/src/join/mixed_join_common_utils.cuh
index aed93e20823..4a52cfe098a 100644
--- a/cpp/src/join/mixed_join_common_utils.cuh
+++ b/cpp/src/join/mixed_join_common_utils.cuh
@@ -178,7 +178,7 @@ struct double_row_equality_comparator {
   }
 };
 
-// A CUDA Cooperative Group of 4 threads for the hash set.
+// A CUDA Cooperative Group of 1 thread for the hash set for mixed semi.
 auto constexpr DEFAULT_MIXED_SEMI_JOIN_CG_SIZE = 1;
 
 // The hash set type used by mixed_semi_join with the build_table.
diff --git a/cpp/tests/join/mixed_join_tests.cu b/cpp/tests/join/mixed_join_tests.cu
index d43bf4dfca9..6b720484771 100644
--- a/cpp/tests/join/mixed_join_tests.cu
+++ b/cpp/tests/join/mixed_join_tests.cu
@@ -797,35 +797,16 @@ TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMapLarge)
 {
   using T1 = double;
 
-  auto const random_data = [](size_t size) {
-    std::vector<T1> values(size);
-    using uniform_distribution =
-      typename std::conditional_t<std::is_same_v<T1, bool>,
-                                  std::bernoulli_distribution,
-                                  std::conditional_t<std::is_floating_point_v<T1>,
-                                                     std::uniform_real_distribution<T1>,
-                                                     std::uniform_int_distribution<T1>>>;
-
-    static constexpr auto seed = 0xf00d;
-    static std::mt19937 engine{seed};
-    static uniform_distribution dist{};
-    std::generate(values.begin(), values.end(), [&]() { return T1{dist(engine)}; });
-
-    return values;
-  };
-
-  auto const random_validity = [&](size_t size) {
-    std::vector<bool> validity(size);
-    static constexpr auto seed = 0xcafe;
-    static std::mt19937 engine{seed};
-    static std::bernoulli_distribution dist{};
-    std::generate(validity.begin(), validity.end(), [&]() { return dist(engine); });
+  // Column size
+  auto constexpr N = 1000;
 
-    return validity;
-  };
+  // Generate column data for left and right tables
+  auto const [left_col0, right_col0] = gen_random_nullable_repeated_columns<T1>(N, 200);
+  auto const [left_col1, right_col1] = gen_random_nullable_repeated_columns<T1>(N, 100);
 
+  // Setup data and nulls for the left table
   std::vector<std::pair<std::vector<T1>, std::vector<bool>>> lefts = {
-    {random_data(500), random_validity(500)}, {random_data(500), random_validity(500)}};
+    {left_col0.first, left_col0.second}, {left_col1.first, left_col1.second}};
   std::vector<cudf::test::fixed_width_column_wrapper<T1>> left_wrappers;
   std::vector<cudf::column_view> left_columns;
   for (auto [data, valids] : lefts) {
@@ -834,8 +815,9 @@ TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMapLarge)
     left_columns.emplace_back(left_wrappers.back());
   };
 
+  // Setup data and nulls for the right table
   std::vector<std::pair<std::vector<T1>, std::vector<bool>>> rights = {
-    {random_data(250), random_validity(250)}, {random_data(250), random_validity(250)}};
+    {right_col0.first, right_col0.second}, {right_col1.first, right_col1.second}};
   std::vector<cudf::test::fixed_width_column_wrapper<T1>> right_wrappers;
   std::vector<cudf::column_view> right_columns;
   for (auto [data, valids] : rights) {
@@ -848,7 +830,7 @@ TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMapLarge)
   auto const left_table  = cudf::table_view{left_columns};
   auto const right_table = cudf::table_view{right_columns};
 
-  // Use the zeroth column for equality.
+  // Using the zeroth column for equality.
   auto const left_equality  = left_table.select({0});
   auto const right_equality = right_table.select({0});
 
@@ -858,48 +840,74 @@ TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMapLarge)
   auto left_zero_eq_right_zero =
     cudf::ast::operation(cudf::ast::ast_operator::EQUAL, col_ref_left_0, col_ref_right_0);
 
-  // Expected size of left_semi_join with only zeroth column equality.
-  auto const expected_num_idx_left_zero_eq_right_zero =
-    cudf::conditional_left_semi_join_size(left_table, right_table, left_zero_eq_right_zero);
-
-  // Actual size of mixed_left_semi_join with only zeroth column equality.
-  auto const num_idx_left_zero_eq_right_zero =
-    cudf::mixed_left_semi_join(left_equality,
-                               right_equality,
-                               left_table,
-                               right_table,
-                               left_zero_eq_right_zero,
-                               cudf::null_equality::UNEQUAL)
-      ->size();
-
-  // Expected and actual sizes must match.
-  EXPECT_EQ(expected_num_idx_left_zero_eq_right_zero, num_idx_left_zero_eq_right_zero);
-
-  // Common column references for conditional column.
-  auto const col_ref_left_1  = cudf::ast::column_reference(1, cudf::ast::table_reference::LEFT);
-  auto const col_ref_right_1 = cudf::ast::column_reference(1, cudf::ast::table_reference::RIGHT);
-  auto left_one_gt_right_one =
-    cudf::ast::operation(cudf::ast::ast_operator::GREATER, col_ref_left_1, col_ref_right_1);
+  // Mixed semi join with zeroth column equality
+  {
+    // Expected left_semi_join result
+    auto const expected_mixed_semi_join =
+      cudf::conditional_left_semi_join(left_table, right_table, left_zero_eq_right_zero);
+
+    // Actual mixed_left_semi_join result
+    auto const mixed_semi_join = cudf::mixed_left_semi_join(left_equality,
+                                                            right_equality,
+                                                            left_table,
+                                                            right_table,
+                                                            left_zero_eq_right_zero,
+                                                            cudf::null_equality::UNEQUAL);
+
+    // Copy data back to host for comparisons
+    auto expected_indices = cudf::detail::make_std_vector_async<int32_t>(
+      cudf::device_span<int32_t>(*expected_mixed_semi_join), cudf::get_default_stream());
+    auto result_indices = cudf::detail::make_std_vector_sync<int32_t>(
+      cudf::device_span<int32_t>(*mixed_semi_join), cudf::get_default_stream());
+
+    // Sort the indices for 1-1 comparison
+    std::sort(expected_indices.begin(), expected_indices.end());
+    std::sort(result_indices.begin(), result_indices.end());
+
+    // Expected and actual vectors must match.
+    EXPECT_EQ(expected_mixed_semi_join->size(), mixed_semi_join->size());
+    EXPECT_TRUE(
+      std::equal(expected_indices.begin(), expected_indices.end(), result_indices.begin()));
+  }
 
-  auto combined_condition = cudf::ast::operation(
-    cudf::ast::ast_operator::LOGICAL_AND, left_zero_eq_right_zero, left_one_gt_right_one);
-
-  // Expected size of left_semi_join with zeroth col equality and first col conditional.
-  auto const expected_num_idx_left_one_greater_right_one =
-    cudf::conditional_left_semi_join_size(left_table, right_table, combined_condition);
-
-  // Actual size of left_semi_join with zeroth col equality and first col conditional.
-  auto const num_idx_left_one_greater_right_one =
-    cudf::mixed_left_semi_join(left_equality,
-                               right_equality,
-                               left_table,
-                               right_table,
-                               left_one_gt_right_one,
-                               cudf::null_equality::UNEQUAL)
-      ->size();
-
-  // Expected and actual sizes must match.
-  EXPECT_EQ(expected_num_idx_left_one_greater_right_one, num_idx_left_one_greater_right_one);
+  // Mixed semi join with zeroth column equality and first column GREATER conditional
+  {
+    // Column references for conditional column.
+    auto const col_ref_left_1  = cudf::ast::column_reference(1, cudf::ast::table_reference::LEFT);
+    auto const col_ref_right_1 = cudf::ast::column_reference(1, cudf::ast::table_reference::RIGHT);
+    auto left_one_gt_right_one =
+      cudf::ast::operation(cudf::ast::ast_operator::GREATER, col_ref_left_1, col_ref_right_1);
+
+    // Expected left_semi_join result
+    auto const expected_mixed_semi_join = cudf::conditional_left_semi_join(
+      left_table,
+      right_table,
+      cudf::ast::operation(
+        cudf::ast::ast_operator::LOGICAL_AND, left_zero_eq_right_zero, left_one_gt_right_one));
+
+    // Actual left_semi_join result
+    auto const mixed_semi_join = cudf::mixed_left_semi_join(left_equality,
+                                                            right_equality,
+                                                            left_table,
+                                                            right_table,
+                                                            left_one_gt_right_one,
+                                                            cudf::null_equality::UNEQUAL);
+
+    // Copy data back to host for comparisons
+    auto expected_indices = cudf::detail::make_std_vector_async<int32_t>(
+      cudf::device_span<int32_t>(*expected_mixed_semi_join), cudf::get_default_stream());
+    auto result_indices = cudf::detail::make_std_vector_sync<int32_t>(
+      cudf::device_span<int32_t>(*mixed_semi_join), cudf::get_default_stream());
+
+    // Sort the indices for 1-1 comparison
+    std::sort(expected_indices.begin(), expected_indices.end());
+    std::sort(result_indices.begin(), result_indices.end());
+
+    // Expected and actual vectors must match.
+    EXPECT_EQ(expected_mixed_semi_join->size(), mixed_semi_join->size());
+    EXPECT_TRUE(
+      std::equal(expected_indices.begin(), expected_indices.end(), result_indices.begin()));
+  }
 }
 
 TYPED_TEST(MixedLeftSemiJoinTest, BasicEqualityDuplicates)

From d618b55b0f38f5dcd9379803a883417260021c76 Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Wed, 25 Sep 2024 19:47:53 -0700
Subject: [PATCH 10/11] Update cpp/tests/join/mixed_join_tests.cu

---
 cpp/tests/join/mixed_join_tests.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/tests/join/mixed_join_tests.cu b/cpp/tests/join/mixed_join_tests.cu
index 6b720484771..fb09e1572d5 100644
--- a/cpp/tests/join/mixed_join_tests.cu
+++ b/cpp/tests/join/mixed_join_tests.cu
@@ -797,7 +797,7 @@ TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMapLarge)
 {
   using T1 = double;
 
-  // Column size
+  // Number of rows in each column
   auto constexpr N = 1000;
 
   // Generate column data for left and right tables

From 67ce08004e044de3c81ede78d3dfeed49c7186bf Mon Sep 17 00:00:00 2001
From: Muhammad Haseeb <14217455+mhaseeb123@users.noreply.github.com>
Date: Thu, 26 Sep 2024 02:49:27 +0000
Subject: [PATCH 11/11] Increase number of rows in tables for better testing.

---
 cpp/tests/join/mixed_join_tests.cu | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/tests/join/mixed_join_tests.cu b/cpp/tests/join/mixed_join_tests.cu
index fb09e1572d5..9041969bec7 100644
--- a/cpp/tests/join/mixed_join_tests.cu
+++ b/cpp/tests/join/mixed_join_tests.cu
@@ -798,7 +798,7 @@ TYPED_TEST(MixedLeftSemiJoinTest, MixedLeftSemiJoinGatherMapLarge)
   using T1 = double;
 
   // Number of rows in each column
-  auto constexpr N = 1000;
+  auto constexpr N = 10000;
 
   // Generate column data for left and right tables
   auto const [left_col0, right_col0] = gen_random_nullable_repeated_columns<T1>(N, 200);