From 963c8d061df5ac11f7202f278d208b38c0a17ec0 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 4 Aug 2026 00:01:03 +0000 Subject: [PATCH 1/4] Optimize hash full join finalization --- cpp/src/join/hash_join/dispatch.cuh | 21 ++++++- cpp/src/join/hash_join/full_join_size_impl.cu | 55 +++++-------------- cpp/src/join/hash_join/retrieve_impl.cuh | 36 ++++++++++-- cpp/src/join/join_common_utils.hpp | 5 +- cpp/src/join/join_utils.cu | 40 ++++++++------ cpp/tests/join/join_tests.cpp | 8 +++ 6 files changed, 100 insertions(+), 65 deletions(-) diff --git a/cpp/src/join/hash_join/dispatch.cuh b/cpp/src/join/hash_join/dispatch.cuh index 2092817e2931..5137c470b2da 100644 --- a/cpp/src/join/hash_join/dispatch.cuh +++ b/cpp/src/join/hash_join/dispatch.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -52,6 +52,25 @@ struct output_fn { } }; +/** + * @brief Extracts a right-side row index and marks it as matched. + * + * This is used while retrieving a full join to build the right-side match set without a + * subsequent pass over the join output. + */ +struct mark_matched_output_fn { + size_type* right_matches; + size_type right_table_num_rows; + + __device__ cudf::size_type operator()( + cuco::pair const& slot) const + { + auto const index = slot.second; + if (index >= 0 && index < right_table_num_rows) { right_matches[index] = 1; } + return index; + } +}; + /** * @brief Equality comparator for cuco hash table probing with primitive row equality. */ diff --git a/cpp/src/join/hash_join/full_join_size_impl.cu b/cpp/src/join/hash_join/full_join_size_impl.cu index 0e9afcde668d..d5adb09ced2c 100644 --- a/cpp/src/join/hash_join/full_join_size_impl.cu +++ b/cpp/src/join/hash_join/full_join_size_impl.cu @@ -6,48 +6,22 @@ #include "retrieve_impl.cuh" #include +#include #include -#include -#include - #include -#include -#include -#include - -#include namespace cudf::detail { namespace { -std::size_t compute_left_join_complement_size(cudf::device_span right_indices, - size_type left_table_row_count, - size_type right_table_row_count, +std::size_t compute_left_join_complement_size(cudf::device_span right_matches, rmm::cuda_stream_view stream) { - if (left_table_row_count == 0) { return right_table_row_count; } - - auto invalid_index_map = - std::make_unique>(right_table_row_count, stream); - thrust::uninitialized_fill( - rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - invalid_index_map->begin(), - invalid_index_map->end(), - int32_t{1}); - - valid_range valid(0, right_table_row_count); - - thrust::scatter_if(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::make_constant_iterator(0), - cuda::make_constant_iterator(0) + right_indices.size(), - right_indices.begin(), - right_indices.begin(), - invalid_index_map->begin(), - valid); - return cudf::detail::count_if( - invalid_index_map->begin(), invalid_index_map->end(), cuda::std::identity{}, stream); + right_matches.begin(), + right_matches.end(), + [] __device__(size_type is_matched) { return is_matched == 0; }, + stream); } } // namespace @@ -71,12 +45,13 @@ std::size_t get_full_join_size( compare_nulls, stream); - if (join_size == 0) { return join_size; } + if (join_size == 0) { return right_table.num_rows(); } - auto right_indices = std::make_unique>(join_size, stream, mr); - - auto const out_build_begin = - cuda::make_transform_output_iterator(right_indices->begin(), output_fn{}); + auto right_matches = + cudf::detail::make_zeroed_device_uvector_async(right_table.num_rows(), stream, mr); + auto const out_build_begin = cuda::make_transform_output_iterator( + cuda::make_discard_iterator(), + mark_matched_output_fn{right_matches.data(), right_table.num_rows()}); retrieve_left_join_build_indices(right_table, left_table, @@ -88,11 +63,7 @@ std::size_t get_full_join_size( out_build_begin, stream); - auto const left_table_row_count = left_table.num_rows(); - auto const right_table_row_count = right_table.num_rows(); - - return join_size + compute_left_join_complement_size( - *right_indices, left_table_row_count, right_table_row_count, stream); + return join_size + compute_left_join_complement_size(right_matches, stream); } } // namespace cudf::detail diff --git a/cpp/src/join/hash_join/retrieve_impl.cuh b/cpp/src/join/hash_join/retrieve_impl.cuh index 5efe69afe850..c208fdc86c47 100644 --- a/cpp/src/join/hash_join/retrieve_impl.cuh +++ b/cpp/src/join/hash_join/retrieve_impl.cuh @@ -12,6 +12,7 @@ #include #include +#include #include #include #include @@ -37,6 +38,7 @@ probe_join_hash_table( bool has_nulls, null_equality compare_nulls, std::optional output_size, + cudf::device_span right_matches, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr) { @@ -61,16 +63,34 @@ probe_join_hash_table( std::make_unique>(0, stream, mr)); } - auto left_indices = std::make_unique>(join_size, stream, mr); - auto right_indices = std::make_unique>(join_size, stream, mr); + // Without a supplied full-join size, reserve room for the largest possible right complement. + // This prevents finalization from reallocating and copying both left-join output vectors. + auto const allocation_size = Join == join_kind::FULL_JOIN && !output_size + ? join_size + static_cast(right_table.num_rows()) + : join_size; + auto left_indices = std::make_unique>(allocation_size, stream, mr); + auto right_indices = + std::make_unique>(allocation_size, stream, mr); + left_indices->resize(join_size, stream); + right_indices->resize(join_size, stream); cudf::prefetch::detail::prefetch(*left_indices, stream); cudf::prefetch::detail::prefetch(*right_indices, stream); auto const left_table_num_rows = left_table.num_rows(); auto const out_probe_begin = cuda::make_transform_output_iterator(left_indices->begin(), output_fn{}); - auto const out_build_begin = - cuda::make_transform_output_iterator(right_indices->begin(), output_fn{}); + auto const out_build_begin = [&] { + if constexpr (Join == join_kind::FULL_JOIN) { + CUDF_EXPECTS(right_matches.size() == right_table.num_rows(), + "full join requires one match flag per right row", + std::invalid_argument); + return cuda::make_transform_output_iterator( + right_indices->begin(), + mark_matched_output_fn{right_matches.data(), right_table.num_rows()}); + } else { + return cuda::make_transform_output_iterator(right_indices->begin(), output_fn{}); + } + }(); auto retrieve_results = [&](auto equality, auto d_hasher) { auto const iter = cudf::detail::make_counting_transform_iterator(0, pair_fn{d_hasher}); @@ -176,6 +196,11 @@ hash_join::join_retrieve(cudf::table_view const& left, auto const preprocessed_left = cudf::detail::row::equality::preprocessed_table::create(left, stream); + auto right_matches = cudf::detail::make_zeroed_device_uvector_async( + Join == join_kind::FULL_JOIN ? _right.num_rows() : 0, + stream, + cudf::get_current_device_resource_ref()); + auto join_indices = cudf::detail::probe_join_hash_table(_right, left, _preprocessed_right, @@ -184,12 +209,13 @@ hash_join::join_retrieve(cudf::table_view const& left, _has_nulls, _nulls_equal, output_size, + right_matches, stream, mr); if constexpr (Join == join_kind::FULL_JOIN) { return detail::finalize_full_join( - std::move(join_indices), left.num_rows(), _right.num_rows(), stream, mr); + std::move(join_indices), left.num_rows(), _right.num_rows(), stream, mr, right_matches); } else { return join_indices; } diff --git a/cpp/src/join/join_common_utils.hpp b/cpp/src/join/join_common_utils.hpp index a5bf0ce6785c..63dfd802616c 100644 --- a/cpp/src/join/join_common_utils.hpp +++ b/cpp/src/join/join_common_utils.hpp @@ -67,6 +67,8 @@ VectorPair get_trivial_left_join_indices(table_view const& left, * @param right_table_num_rows Number of rows in the right table. * @param stream CUDA stream used for device memory operations and kernel launches. * @param mr Device memory resource used to allocate working storage. + * @param right_matches Optional precomputed flags indicating which right rows matched. When empty, + * the flags are derived from `indices.second`. * * @return `[left_indices, right_indices]` of the complete full-join output. */ @@ -74,7 +76,8 @@ VectorPair finalize_full_join(VectorPair&& indices, size_type left_table_num_rows, size_type right_table_num_rows, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); + rmm::device_async_resource_ref mr, + cudf::device_span right_matches = {}); /** * @brief Finalize a full-join result from per-partition index spans. diff --git a/cpp/src/join/join_utils.cu b/cpp/src/join/join_utils.cu index a238ec5c44e6..401f065d3345 100644 --- a/cpp/src/join/join_utils.cu +++ b/cpp/src/join/join_utils.cu @@ -88,7 +88,8 @@ VectorPair finalize_full_join(VectorPair&& indices, size_type left_table_num_rows, size_type right_table_num_rows, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + rmm::device_async_resource_ref mr, + cudf::device_span right_matches) { auto [left_out, right_out] = std::move(indices); CUDF_EXPECTS(left_out->size() == right_out->size(), @@ -122,20 +123,27 @@ VectorPair finalize_full_join(VectorPair&& indices, left_out->resize(upper, stream); right_out->resize(upper, stream); - // Mark matched right rows in an int32 flag array (one word per right row). Redundant stores - // of the same value are idempotent, so no atomics are needed. Word-sized stores coalesce into - // full 128-byte transactions per warp; byte-sized flags cost ~2–3× here because partial-word - // stores from dense scatters serialize within each 32-bit sector. - auto flags = cudf::detail::make_zeroed_device_uvector_async( - right_table_num_rows, stream, cudf::get_current_device_resource_ref()); - - thrust::scatter_if(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), - cuda::make_constant_iterator(size_type{1}), - cuda::make_constant_iterator(size_type{1}) + match_total, - right_out->begin(), - right_out->begin(), - flags.begin(), - valid_range{0, right_table_num_rows}); + CUDF_EXPECTS(right_matches.empty() || right_matches.size() == right_table_num_rows, + "right match flags must be empty or have one entry per right row", + std::invalid_argument); + + // Hash joins mark right rows as part of retrieval and pass those flags here, eliminating an + // output-sized scatter. Other join implementations use this fallback to derive the same flags + // from their materialized right indices. + auto computed_matches = cudf::detail::make_zeroed_device_uvector_async( + right_matches.empty() ? right_table_num_rows : 0, + stream, + cudf::get_current_device_resource_ref()); + if (right_matches.empty()) { + thrust::scatter_if(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::make_constant_iterator(size_type{1}), + cuda::make_constant_iterator(size_type{1}) + match_total, + right_out->begin(), + right_out->begin(), + computed_matches.begin(), + valid_range{0, right_table_num_rows}); + } + auto const match_flags = right_matches.empty() ? computed_matches.data() : right_matches.data(); // Fused compaction: for each unmatched right row, emit (JoinNoMatch, right_idx) into // (left_out_tail, right_out_tail) in a single CUB DeviceSelect pass. @@ -147,7 +155,7 @@ VectorPair finalize_full_join(VectorPair&& indices, cudf::detail::copy_if(cuda::counting_iterator{0}, cuda::counting_iterator{right_table_num_rows}, out_iter, - unmatched_flag{flags.data()}, + unmatched_flag{match_flags}, stream); auto const comp_size = cuda::std::distance(out_iter, new_end); diff --git a/cpp/tests/join/join_tests.cpp b/cpp/tests/join/join_tests.cpp index d328777418b8..a29d154b8efa 100644 --- a/cpp/tests/join/join_tests.cpp +++ b/cpp/tests/join/join_tests.cpp @@ -1713,6 +1713,14 @@ TEST_F(JoinTest, EmptyLeftTableFullJoin) auto sorted_gold = cudf::gather(gold.view(), *gold_sort_order); CUDF_TEST_EXPECT_TABLES_EQUIVALENT(*sorted_gold, *sorted_result); + + auto hash_joiner = cudf::hash_join(rhs, cudf::null_equality::EQUAL); + auto const output_size = hash_joiner.full_join_size(lhs); + EXPECT_EQ(output_size, rhs.num_rows()); + + auto const [left_indices, right_indices] = hash_joiner.full_join(lhs, output_size); + EXPECT_EQ(left_indices->size(), output_size); + EXPECT_EQ(right_indices->size(), output_size); } // Empty Right Table From 6a474cdab4328938c84928cf149f8ec15d6f804e Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 4 Aug 2026 00:59:04 +0000 Subject: [PATCH 2/4] Fix full join match flag size checks --- cpp/src/join/hash_join/retrieve_impl.cuh | 2 +- cpp/src/join/join_utils.cu | 7 ++++--- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/cpp/src/join/hash_join/retrieve_impl.cuh b/cpp/src/join/hash_join/retrieve_impl.cuh index c208fdc86c47..188068018c4e 100644 --- a/cpp/src/join/hash_join/retrieve_impl.cuh +++ b/cpp/src/join/hash_join/retrieve_impl.cuh @@ -81,7 +81,7 @@ probe_join_hash_table( cuda::make_transform_output_iterator(left_indices->begin(), output_fn{}); auto const out_build_begin = [&] { if constexpr (Join == join_kind::FULL_JOIN) { - CUDF_EXPECTS(right_matches.size() == right_table.num_rows(), + CUDF_EXPECTS(right_matches.size() == static_cast(right_table.num_rows()), "full join requires one match flag per right row", std::invalid_argument); return cuda::make_transform_output_iterator( diff --git a/cpp/src/join/join_utils.cu b/cpp/src/join/join_utils.cu index 401f065d3345..b76696c9427c 100644 --- a/cpp/src/join/join_utils.cu +++ b/cpp/src/join/join_utils.cu @@ -123,9 +123,10 @@ VectorPair finalize_full_join(VectorPair&& indices, left_out->resize(upper, stream); right_out->resize(upper, stream); - CUDF_EXPECTS(right_matches.empty() || right_matches.size() == right_table_num_rows, - "right match flags must be empty or have one entry per right row", - std::invalid_argument); + CUDF_EXPECTS( + right_matches.empty() || right_matches.size() == static_cast(right_table_num_rows), + "right match flags must be empty or have one entry per right row", + std::invalid_argument); // Hash joins mark right rows as part of retrieval and pass those flags here, eliminating an // output-sized scatter. Other join implementations use this fallback to derive the same flags From 53d80f704e872c35375dad16567e1c6b05bc4dd8 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 4 Aug 2026 22:05:39 +0000 Subject: [PATCH 3/4] Address hash full join review feedback --- cpp/src/join/conditional_join.cu | 4 ++-- cpp/src/join/hash_join/full_join_size_impl.cu | 9 ++++----- cpp/src/join/hash_join/retrieve_impl.cuh | 2 +- cpp/src/join/hash_join/size_impl.cuh | 10 ++++------ cpp/src/join/join_common_utils.hpp | 9 +++++---- cpp/src/join/join_utils.cu | 17 ++++++++--------- cpp/src/join/mixed_join.cu | 8 ++++++-- 7 files changed, 30 insertions(+), 29 deletions(-) diff --git a/cpp/src/join/conditional_join.cu b/cpp/src/join/conditional_join.cu index 8e2688b99f3b..44b7734b1951 100644 --- a/cpp/src/join/conditional_join.cu +++ b/cpp/src/join/conditional_join.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -277,7 +277,7 @@ conditional_join(table_view const& left, // by any row in the left table. if (join_type == join_kind::FULL_JOIN) { join_indices = detail::finalize_full_join( - std::move(join_indices), left.num_rows(), right.num_rows(), stream, mr); + std::move(join_indices), left.num_rows(), right.num_rows(), std::nullopt, stream, mr); } return join_indices; } diff --git a/cpp/src/join/hash_join/full_join_size_impl.cu b/cpp/src/join/hash_join/full_join_size_impl.cu index d5adb09ced2c..db4c2a6ef408 100644 --- a/cpp/src/join/hash_join/full_join_size_impl.cu +++ b/cpp/src/join/hash_join/full_join_size_impl.cu @@ -20,7 +20,7 @@ std::size_t compute_left_join_complement_size(cudf::device_span return cudf::detail::count_if( right_matches.begin(), right_matches.end(), - [] __device__(size_type is_matched) { return is_matched == 0; }, + [] __device__(size_type is_matched) -> bool { return is_matched == 0; }, stream); } } // namespace @@ -33,8 +33,7 @@ std::size_t get_full_join_size( cudf::detail::hash_table_t const& hash_table, bool has_nulls, null_equality compare_nulls, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + rmm::cuda_stream_view stream) { std::size_t join_size = compute_join_output_size(right_table, left_table, @@ -47,8 +46,8 @@ std::size_t get_full_join_size( if (join_size == 0) { return right_table.num_rows(); } - auto right_matches = - cudf::detail::make_zeroed_device_uvector_async(right_table.num_rows(), stream, mr); + auto right_matches = cudf::detail::make_zeroed_device_uvector_async( + right_table.num_rows(), stream, cudf::get_current_device_resource_ref()); auto const out_build_begin = cuda::make_transform_output_iterator( cuda::make_discard_iterator(), mark_matched_output_fn{right_matches.data(), right_table.num_rows()}); diff --git a/cpp/src/join/hash_join/retrieve_impl.cuh b/cpp/src/join/hash_join/retrieve_impl.cuh index 188068018c4e..4ecac53c4def 100644 --- a/cpp/src/join/hash_join/retrieve_impl.cuh +++ b/cpp/src/join/hash_join/retrieve_impl.cuh @@ -215,7 +215,7 @@ hash_join::join_retrieve(cudf::table_view const& left, if constexpr (Join == join_kind::FULL_JOIN) { return detail::finalize_full_join( - std::move(join_indices), left.num_rows(), _right.num_rows(), stream, mr, right_matches); + std::move(join_indices), left.num_rows(), _right.num_rows(), right_matches, stream, mr); } else { return join_indices; } diff --git a/cpp/src/join/hash_join/size_impl.cuh b/cpp/src/join/hash_join/size_impl.cuh index 3e20ebc7367e..44f2e2950665 100644 --- a/cpp/src/join/hash_join/size_impl.cuh +++ b/cpp/src/join/hash_join/size_impl.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #pragma once @@ -20,8 +20,7 @@ std::size_t get_full_join_size( cudf::detail::hash_table_t const& hash_table, bool has_nulls, null_equality compare_nulls, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr); + rmm::cuda_stream_view stream); template std::size_t compute_join_output_size( @@ -97,7 +96,7 @@ template template std::size_t hash_join::join_size(cudf::table_view const& left, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) const + rmm::device_async_resource_ref) const { static_assert(Join == join_kind::FULL_JOIN); @@ -119,8 +118,7 @@ std::size_t hash_join::join_size(cudf::table_view const& left, _impl->_hash_table, _has_nulls, _nulls_equal, - stream, - mr); + stream); } } // namespace cudf::detail diff --git a/cpp/src/join/join_common_utils.hpp b/cpp/src/join/join_common_utils.hpp index 63dfd802616c..f6fbfbd80aba 100644 --- a/cpp/src/join/join_common_utils.hpp +++ b/cpp/src/join/join_common_utils.hpp @@ -14,6 +14,7 @@ #include #include +#include #include namespace cudf::detail { @@ -65,19 +66,19 @@ VectorPair get_trivial_left_join_indices(table_view const& left, * @param left_table_num_rows Number of rows in the left table (0 → every right row is * unmatched, fast path). * @param right_table_num_rows Number of rows in the right table. + * @param right_matches Optional precomputed flags indicating which right rows matched. When absent, + * the flags are derived from `indices.second`. * @param stream CUDA stream used for device memory operations and kernel launches. * @param mr Device memory resource used to allocate working storage. - * @param right_matches Optional precomputed flags indicating which right rows matched. When empty, - * the flags are derived from `indices.second`. * * @return `[left_indices, right_indices]` of the complete full-join output. */ VectorPair finalize_full_join(VectorPair&& indices, size_type left_table_num_rows, size_type right_table_num_rows, + std::optional> right_matches, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr, - cudf::device_span right_matches = {}); + rmm::device_async_resource_ref mr); /** * @brief Finalize a full-join result from per-partition index spans. diff --git a/cpp/src/join/join_utils.cu b/cpp/src/join/join_utils.cu index b76696c9427c..5e22c7c95e36 100644 --- a/cpp/src/join/join_utils.cu +++ b/cpp/src/join/join_utils.cu @@ -87,9 +87,9 @@ struct to_no_match_pair { VectorPair finalize_full_join(VectorPair&& indices, size_type left_table_num_rows, size_type right_table_num_rows, + std::optional> right_matches, rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr, - cudf::device_span right_matches) + rmm::device_async_resource_ref mr) { auto [left_out, right_out] = std::move(indices); CUDF_EXPECTS(left_out->size() == right_out->size(), @@ -124,18 +124,16 @@ VectorPair finalize_full_join(VectorPair&& indices, right_out->resize(upper, stream); CUDF_EXPECTS( - right_matches.empty() || right_matches.size() == static_cast(right_table_num_rows), - "right match flags must be empty or have one entry per right row", + !right_matches || right_matches->size() == static_cast(right_table_num_rows), + "right match flags must be absent or have one entry per right row", std::invalid_argument); // Hash joins mark right rows as part of retrieval and pass those flags here, eliminating an // output-sized scatter. Other join implementations use this fallback to derive the same flags // from their materialized right indices. auto computed_matches = cudf::detail::make_zeroed_device_uvector_async( - right_matches.empty() ? right_table_num_rows : 0, - stream, - cudf::get_current_device_resource_ref()); - if (right_matches.empty()) { + right_matches ? 0 : right_table_num_rows, stream, cudf::get_current_device_resource_ref()); + if (!right_matches) { thrust::scatter_if(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), cuda::make_constant_iterator(size_type{1}), cuda::make_constant_iterator(size_type{1}) + match_total, @@ -144,7 +142,7 @@ VectorPair finalize_full_join(VectorPair&& indices, computed_matches.begin(), valid_range{0, right_table_num_rows}); } - auto const match_flags = right_matches.empty() ? computed_matches.data() : right_matches.data(); + auto const match_flags = right_matches ? right_matches->data() : computed_matches.data(); // Fused compaction: for each unmatched right row, emit (JoinNoMatch, right_idx) into // (left_out_tail, right_out_tail) in a single CUB DeviceSelect pass. @@ -224,6 +222,7 @@ VectorPair finalize_full_join( return finalize_full_join(std::pair(std::move(left_out), std::move(right_out)), left_table_num_rows, right_table_num_rows, + std::nullopt, stream, mr); } diff --git a/cpp/src/join/mixed_join.cu b/cpp/src/join/mixed_join.cu index baeb35670a36..64c9496f08ae 100644 --- a/cpp/src/join/mixed_join.cu +++ b/cpp/src/join/mixed_join.cu @@ -102,8 +102,12 @@ mixed_join(table_view const& left_equality, std::nullopt, stream, mr); - return finalize_full_join( - std::move(left_outer), left_conditional.num_rows(), right_conditional.num_rows(), stream, mr); + return finalize_full_join(std::move(left_outer), + left_conditional.num_rows(), + right_conditional.num_rows(), + std::nullopt, + stream, + mr); } auto const hash_joiner = cudf::hash_join{right_equality, compare_nulls, stream}; From 847d2fab8c84735a32b052a25e41249cea9a8429 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Fri, 14 Aug 2026 21:01:28 +0000 Subject: [PATCH 4/4] Strengthen empty full join index test --- cpp/tests/join/join_tests.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/cpp/tests/join/join_tests.cpp b/cpp/tests/join/join_tests.cpp index 6f0e0ebed41a..cfdb5f9975f8 100644 --- a/cpp/tests/join/join_tests.cpp +++ b/cpp/tests/join/join_tests.cpp @@ -1783,6 +1783,16 @@ TEST_F(JoinTest, EmptyLeftTableFullJoin) auto const [left_indices, right_indices] = hash_joiner.full_join(lhs, output_size); EXPECT_EQ(left_indices->size(), output_size); EXPECT_EQ(right_indices->size(), output_size); + + column_wrapper expected_left_indices{ + {NoneValue, NoneValue, NoneValue, NoneValue, NoneValue}}; + column_wrapper expected_right_indices{{0, 1, 2, 3, 4}}; + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected_left_indices, + cudf::column_view{cudf::device_span{*left_indices}}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL( + expected_right_indices, + cudf::column_view{cudf::device_span{*right_indices}}); } // Empty Right Table