diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 2061bb3e993f..d1fa5aa5652e 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -740,11 +740,16 @@ add_library( src/join/conditional_join.cu src/join/cross_join.cu src/join/distinct_hash_join.cu - src/join/filter_join_indices.cu - src/join/filter_join_indices_kernel_complex.cu - src/join/filter_join_indices_kernel_null_complex.cu - src/join/filter_join_indices_kernel_null_primitive.cu - src/join/filter_join_indices_kernel_primitive.cu + src/join/filter_join_indices/filter_join_indices.cu + src/join/filter_join_indices/filter_join_indices_jit.cu + src/join/filter_join_indices/filter_join_indices_kernel_complex.cu + src/join/filter_join_indices/filter_join_indices_kernel_null_complex.cu + src/join/filter_join_indices/filter_join_indices_kernel_null_primitive.cu + src/join/filter_join_indices/filter_join_indices_kernel_primitive.cu + src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cu + src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cu + src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cu + src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu src/join/filtered_join.cu src/join/hash_join/finalize_partitioned_full_join.cpp src/join/hash_join/full_join_match_context.cpp @@ -767,11 +772,10 @@ add_library( src/join/hash_join/partitioned_left_join.cu src/join/hash_join/partitioned_retrieve.cu src/join/hash_join/partitioned_retrieve_outer.cu - src/join/mark_join.cu - src/join/filter_join_indices_jit.cu src/join/join.cu src/join/join_utils.cu src/join/key_remapping.cu + src/join/mark_join.cu src/join/mixed_join.cu src/join/mixed_join_kernel.cu src/join/mixed_join_kernel_nulls.cu diff --git a/cpp/include/cudf/join/join.hpp b/cpp/include/cudf/join/join.hpp index 348086e7381b..15530227cf5b 100644 --- a/cpp/include/cudf/join/join.hpp +++ b/cpp/include/cudf/join/join.hpp @@ -323,6 +323,7 @@ std::unique_ptr cross_join( * * @throw std::invalid_argument if join_kind is not INNER_JOIN, LEFT_JOIN, or FULL_JOIN. * @throw std::invalid_argument if left_indices and right_indices have different sizes. + * @throw std::invalid_argument if predicate does not produce a Boolean output. * * @param left The left table for predicate evaluation (conditional columns only). * @param right The right table for predicate evaluation (conditional columns only). @@ -347,6 +348,44 @@ filter_join_indices(cudf::table_view const& left, rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); +/** + * @brief Returns the exact output size of `filter_join_indices` without materializing + * the filtered index vectors. + * + * Runs the same predicate evaluation as `filter_join_indices` but skips the index + * materialization step, returning only the total number of pairs that would be + * emitted. The semantics per `join_kind` match `filter_join_indices`: + * - INNER_JOIN: number of pairs where the predicate evaluates to true. + * - LEFT_JOIN: predicate-passing pairs plus one entry per left row with no passing match. + * - FULL_JOIN: input pairs plus one extra entry per pair whose predicate failed + * (because failed matches split into `(left, JoinNoMatch)` and `(JoinNoMatch, right)`). + * + * The returned size may be passed as a precomputed hint to APIs that compose + * `filter_join_indices` (for example, the mixed join APIs). + * + * @throw std::invalid_argument if `join_kind` is not INNER_JOIN, LEFT_JOIN, or FULL_JOIN. + * @throw std::invalid_argument if `left_indices` and `right_indices` have different sizes. + * @throw std::invalid_argument if `predicate` does not produce a Boolean output. + * + * @param left The left table for predicate evaluation (conditional columns only). + * @param right The right table for predicate evaluation (conditional columns only). + * @param left_indices Device span of row indices in the left table. + * @param right_indices Device span of row indices in the right table. + * @param predicate An AST expression that returns a boolean for each pair of rows. + * @param join_kind The type of join operation. Must be INNER_JOIN, LEFT_JOIN, or FULL_JOIN. + * @param stream CUDA stream used for kernel launches and memory operations. + * + * @return The exact number of pairs that `filter_join_indices` would produce. + */ +[[nodiscard]] std::size_t filter_join_indices_output_size( + cudf::table_view const& left, + cudf::table_view const& right, + cudf::device_span left_indices, + cudf::device_span right_indices, + cudf::ast::expression const& predicate, + cudf::join_kind join_kind, + rmm::cuda_stream_view stream = cudf::get_default_stream()); + /** * @brief JIT-based filtering of join result indices using string predicate. * diff --git a/cpp/src/join/filter_join_indices.cu b/cpp/src/join/filter_join_indices/filter_join_indices.cu similarity index 78% rename from cpp/src/join/filter_join_indices.cu rename to cpp/src/join/filter_join_indices/filter_join_indices.cu index cb0e1670d57c..79ed59797fd4 100644 --- a/cpp/src/join/filter_join_indices.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices.cu @@ -3,7 +3,9 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "filter_join_indices_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.hpp" +#include "join/join_common_utils.hpp" #include #include @@ -13,7 +15,9 @@ #include #include #include +#include #include +#include #include #include #include @@ -31,6 +35,7 @@ #include #include #include +#include #include #include @@ -76,7 +81,8 @@ filter_join_indices(cudf::table_view const& left, predicate, left, right, has_nulls, stream, cudf::get_current_device_resource_ref()}; CUDF_EXPECTS(parser.output_type().id() == type_id::BOOL8, - "The predicate expression must produce a Boolean output"); + "The predicate expression must produce a Boolean output", + std::invalid_argument); // Check if expression contains complex types auto const has_complex_type = parser.has_complex_type(); @@ -355,6 +361,88 @@ filter_join_indices(cudf::table_view const& left, } } +std::size_t filter_join_indices_output_size(cudf::table_view const& left, + cudf::table_view const& right, + cudf::device_span left_indices, + cudf::device_span right_indices, + ast::expression const& predicate, + join_kind join_kind, + rmm::cuda_stream_view stream) +{ + // Validate inputs (same constraints as filter_join_indices) + CUDF_EXPECTS(left_indices.size() == right_indices.size(), + "Left and right index arrays must have the same size", + std::invalid_argument); + CUDF_EXPECTS( + join_kind == join_kind::INNER_JOIN || join_kind == join_kind::LEFT_JOIN || + join_kind == join_kind::FULL_JOIN, + "filter_join_indices_output_size only supports INNER_JOIN, LEFT_JOIN, and FULL_JOIN.", + std::invalid_argument); + + if (left_indices.empty()) { return 0; } + if (join_kind == join_kind::LEFT_JOIN && left.num_rows() == 0) { return 0; } + + auto const has_nulls = predicate.may_evaluate_null(left, right, stream); + + auto const parser = ast::detail::expression_parser{ + predicate, left, right, has_nulls, stream, cudf::get_current_device_resource_ref()}; + + CUDF_EXPECTS(parser.output_type().id() == type_id::BOOL8, + "The predicate expression must produce a Boolean output", + std::invalid_argument); + + auto const has_complex_type = parser.has_complex_type(); + + auto left_table = table_device_view::create(left, stream); + auto right_table = table_device_view::create(right, stream); + + detail::grid_1d const config(left_indices.size(), DEFAULT_JOIN_BLOCK_SIZE); + auto const shmem_per_block = parser.shmem_per_thread * DEFAULT_JOIN_BLOCK_SIZE; + + // The count kernel uses a single atomic counter. Allocate device_scalar zero-initialized. + cudf::detail::device_scalar d_count( + std::size_t{0}, stream, cudf::get_current_device_resource_ref()); + + // For LEFT_JOIN, allocate a zeroed per-left-row mark buffer; for others, pass nullptr. + auto left_passing_marks = cudf::detail::make_zeroed_device_uvector_async( + join_kind == join_kind::LEFT_JOIN ? static_cast(left.num_rows()) : 0, + stream, + cudf::get_current_device_resource_ref()); + auto* const marks_ptr = join_kind == join_kind::LEFT_JOIN ? left_passing_marks.data() : nullptr; + + cudf::detail::dispatch_bool(has_nulls, [&](auto has_nulls_c) { + cudf::detail::dispatch_bool(has_complex_type, [&](auto has_complex_c) { + launch_filter_output_size_kernel( + *left_table, + *right_table, + left_indices, + right_indices, + parser.device_expression_data, + config, + shmem_per_block, + join_kind, + d_count.data(), + marks_ptr, + stream); + }); + }); + + auto const num_predicate_passing = d_count.value(stream); + + switch (join_kind) { + case join_kind::INNER_JOIN: return num_predicate_passing; + case join_kind::FULL_JOIN: return left_indices.size() + num_predicate_passing; + case join_kind::LEFT_JOIN: { + auto const num_filter_passing = cudf::detail::count_if( + left_passing_marks.begin(), left_passing_marks.end(), cuda::std::identity{}, stream); + auto const num_invalid = static_cast(left.num_rows()) - num_filter_passing; + return num_predicate_passing + num_invalid; + } + default: CUDF_FAIL("Unsupported join kind for filter_join_indices_output_size"); + } +} + } // namespace detail // Public API implementation @@ -374,4 +462,17 @@ filter_join_indices(cudf::table_view const& left, left, right, left_indices, right_indices, predicate, join_kind, stream, mr); } +std::size_t filter_join_indices_output_size(cudf::table_view const& left, + cudf::table_view const& right, + cudf::device_span left_indices, + cudf::device_span right_indices, + ast::expression const& predicate, + cudf::join_kind join_kind, + rmm::cuda_stream_view stream) +{ + CUDF_FUNC_RANGE(); + return detail::filter_join_indices_output_size( + left, right, left_indices, right_indices, predicate, join_kind, stream); +} + } // namespace cudf diff --git a/cpp/src/join/filter_join_indices_jit.cu b/cpp/src/join/filter_join_indices/filter_join_indices_jit.cu similarity index 99% rename from cpp/src/join/filter_join_indices_jit.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_jit.cu index ae57a4263058..679d429d9895 100644 --- a/cpp/src/join/filter_join_indices_jit.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_jit.cu @@ -3,8 +3,8 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "filter_join_indices_jit_kernel.cuh" -#include "jit/filter_join_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_jit_kernel.cuh" +#include "join/jit/filter_join_kernel.cuh" #include #include diff --git a/cpp/src/join/filter_join_indices_jit_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_jit_kernel.cuh similarity index 100% rename from cpp/src/join/filter_join_indices_jit_kernel.cuh rename to cpp/src/join/filter_join_indices/filter_join_indices_jit_kernel.cuh diff --git a/cpp/src/join/filter_join_indices_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_kernel.cuh similarity index 95% rename from cpp/src/join/filter_join_indices_kernel.cuh rename to cpp/src/join/filter_join_indices/filter_join_indices_kernel.cuh index 20d40e166eea..9ea25a4336f5 100644 --- a/cpp/src/join/filter_join_indices_kernel.cuh +++ b/cpp/src/join/filter_join_indices/filter_join_indices_kernel.cuh @@ -1,10 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ #pragma once -#include "filter_join_indices_kernel.hpp" +#include "join/filter_join_indices/filter_join_indices_kernel.hpp" #include #include @@ -102,6 +102,7 @@ void launch_filter_gather_map_kernel( right_indices, device_expression_data, predicate_results); + CUDF_CUDA_TRY(cudaGetLastError()); } } // namespace cudf::detail diff --git a/cpp/src/join/filter_join_indices_kernel.hpp b/cpp/src/join/filter_join_indices/filter_join_indices_kernel.hpp similarity index 100% rename from cpp/src/join/filter_join_indices_kernel.hpp rename to cpp/src/join/filter_join_indices/filter_join_indices_kernel.hpp diff --git a/cpp/src/join/filter_join_indices_kernel_complex.cu b/cpp/src/join/filter_join_indices/filter_join_indices_kernel_complex.cu similarity index 73% rename from cpp/src/join/filter_join_indices_kernel_complex.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_kernel_complex.cu index 6152c94439fb..e12ea20a0651 100644 --- a/cpp/src/join/filter_join_indices_kernel_complex.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_kernel_complex.cu @@ -1,10 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ -#include "filter_join_indices_kernel.cuh" -#include "filter_join_indices_kernel.hpp" +#include "join/filter_join_indices/filter_join_indices_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_kernel.hpp" namespace cudf::detail { template void launch_filter_gather_map_kernel( diff --git a/cpp/src/join/filter_join_indices_kernel_null_complex.cu b/cpp/src/join/filter_join_indices/filter_join_indices_kernel_null_complex.cu similarity index 73% rename from cpp/src/join/filter_join_indices_kernel_null_complex.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_kernel_null_complex.cu index 71547e69c58f..531c2ac282ee 100644 --- a/cpp/src/join/filter_join_indices_kernel_null_complex.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_kernel_null_complex.cu @@ -1,10 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ -#include "filter_join_indices_kernel.cuh" -#include "filter_join_indices_kernel.hpp" +#include "join/filter_join_indices/filter_join_indices_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_kernel.hpp" namespace cudf::detail { template void launch_filter_gather_map_kernel( diff --git a/cpp/src/join/filter_join_indices_kernel_null_primitive.cu b/cpp/src/join/filter_join_indices/filter_join_indices_kernel_null_primitive.cu similarity index 73% rename from cpp/src/join/filter_join_indices_kernel_null_primitive.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_kernel_null_primitive.cu index de8f4b7c4f6f..0b36ad324b56 100644 --- a/cpp/src/join/filter_join_indices_kernel_null_primitive.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_kernel_null_primitive.cu @@ -1,10 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ -#include "filter_join_indices_kernel.cuh" -#include "filter_join_indices_kernel.hpp" +#include "join/filter_join_indices/filter_join_indices_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_kernel.hpp" namespace cudf::detail { template void launch_filter_gather_map_kernel( diff --git a/cpp/src/join/filter_join_indices_kernel_primitive.cu b/cpp/src/join/filter_join_indices/filter_join_indices_kernel_primitive.cu similarity index 73% rename from cpp/src/join/filter_join_indices_kernel_primitive.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_kernel_primitive.cu index 4989b2f6aa15..ab9252f243ac 100644 --- a/cpp/src/join/filter_join_indices_kernel_primitive.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_kernel_primitive.cu @@ -1,10 +1,10 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ -#include "filter_join_indices_kernel.cuh" -#include "filter_join_indices_kernel.hpp" +#include "join/filter_join_indices/filter_join_indices_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_kernel.hpp" namespace cudf::detail { template void launch_filter_gather_map_kernel( diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh new file mode 100644 index 000000000000..e158e3883662 --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh @@ -0,0 +1,146 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.hpp" +#include "join/join_common_utils.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include +#include +#include + +#include + +namespace cudf::detail { + +/** + * @brief Counts the per-join-kind output size of `filter_join_indices` without materializing + * a per-pair boolean buffer. + * + * Each thread accumulates a private partial count, the block aggregates with CUB, and each + * block adds its block-sum to `*count_out` exactly once via `cuda::atomic_ref`. For LEFT_JOIN, + * `left_passing_marks[left_row_index]` is additionally set to `true` for every left row that + * contributes to the count, which lets the host derive the number of synthetic JoinNoMatch + * entries. + */ +template +CUDF_KERNEL __launch_bounds__(DEFAULT_JOIN_BLOCK_SIZE) void filter_join_indices_output_size_kernel( + cudf::table_device_view left_table, + cudf::table_device_view right_table, + cudf::device_span left_indices, + cudf::device_span right_indices, + cudf::ast::detail::expression_device_view device_expression_data, + cudf::join_kind join_kind, + std::size_t* count_out, + bool* left_passing_marks) +{ + extern __shared__ char raw_intermediate_storage[]; + auto* intermediate_storage = + reinterpret_cast*>(raw_intermediate_storage); + auto thread_intermediate_storage = + &intermediate_storage[threadIdx.x * device_expression_data.num_intermediates]; + + using BlockReduce = cub::BlockReduce; + __shared__ typename BlockReduce::TempStorage temp_storage; + + auto const tid = cudf::detail::grid_1d::global_thread_id(); + auto const stride = cudf::detail::grid_1d::grid_stride(); + + auto evaluator = cudf::ast::detail::expression_evaluator{ + left_table, right_table, device_expression_data}; + + cuda::std::size_t thread_local_count = 0; + + for (auto i = tid; i < static_cast(left_indices.size()); i += stride) { + auto const left_row_index = left_indices[i]; + auto const right_row_index = right_indices[i]; + + bool const has_non_match = + (left_row_index == cudf::JoinNoMatch || right_row_index == cudf::JoinNoMatch); + + bool predicate_pass = false; + bool both_valid = false; + if (has_non_match) { + // Outer-join unmatched pair: treat as passing so it is preserved in the output count. + predicate_pass = true; + } else if (left_row_index >= 0 && left_row_index < left_table.num_rows() && + right_row_index >= 0 && right_row_index < right_table.num_rows()) { + auto result = cudf::ast::detail::value_expression_result{}; + evaluator.evaluate(result, left_row_index, right_row_index, 0, thread_intermediate_storage); + predicate_pass = result.is_valid() && result.value(); + both_valid = true; + } + + switch (join_kind) { + case cudf::join_kind::INNER_JOIN: + if (predicate_pass) { ++thread_local_count; } + break; + case cudf::join_kind::LEFT_JOIN: + if (predicate_pass) { + ++thread_local_count; + // Mark the left row as "passing" so the host can derive how many left rows need a + // synthetic JoinNoMatch entry. For matched-passing pairs and for pre-existing + // (left, JoinNoMatch) entries from upstream hash_join.left_join the left index is a + // valid row index in [0, left_table.num_rows()). + if (left_row_index >= 0 && left_row_index < left_table.num_rows()) { + left_passing_marks[left_row_index] = true; + } + } + break; + case cudf::join_kind::FULL_JOIN: + // Count failed matches: predicate false AND both indices valid. + if (both_valid && !predicate_pass) { ++thread_local_count; } + break; + default: break; + } + } + + cuda::std::size_t const block_sum = BlockReduce(temp_storage).Sum(thread_local_count); + + if (threadIdx.x == 0) { + cuda::atomic_ref count_ref{*count_out}; + count_ref.fetch_add(block_sum, cuda::memory_order_relaxed); + } +} + +template +void launch_filter_output_size_kernel( + cudf::table_device_view const& left_table, + cudf::table_device_view const& right_table, + cudf::device_span left_indices, + cudf::device_span right_indices, + cudf::ast::detail::expression_device_view device_expression_data, + cudf::detail::grid_1d const& config, + std::size_t shmem_per_block, + cudf::join_kind join_kind, + std::size_t* count_out, + bool* left_passing_marks, + rmm::cuda_stream_view stream) +{ + filter_join_indices_output_size_kernel + <<>>( + left_table, + right_table, + left_indices, + right_indices, + device_expression_data, + join_kind, + count_out, + left_passing_marks); + CUDF_CUDA_TRY(cudaGetLastError()); +} + +} // namespace cudf::detail diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.hpp b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.hpp new file mode 100644 index 000000000000..b5b25e5e7a0e --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.hpp @@ -0,0 +1,68 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include "join/filter_join_indices/filter_join_indices_kernel.hpp" + +#include +#include +#include +#include +#include +#include + +#include + +#include + +namespace cudf::detail { + +/** + * @brief Launches a kernel that counts the per-join-kind output size for `filter_join_indices`. + * + * For INNER_JOIN this is the number of pairs whose predicate evaluates to true. + * For LEFT_JOIN this is the number of input pairs whose predicate evaluates to true + * (including pre-existing unmatched pairs that are preserved); additionally, + * `left_passing_marks[left_row_index]` is set to `true` for every left row that + * contributes to that count (used by the host code to derive the number of left + * rows that need a synthetic JoinNoMatch entry). + * For FULL_JOIN this is the number of failed matched pairs (predicate false and + * both indices valid), which is added on top of `left_indices.size()` host-side. + * + * The kernel avoids materializing a per-pair boolean buffer; it folds the count + * directly into `count_out` via atomic increments. + * + * @tparam has_nulls Indicates whether the expression may evaluate to null + * @tparam has_complex_type Indicates whether the expression may contain complex types + * + * @param[in] left_table Device view of the left table + * @param[in] right_table Device view of the right table + * @param[in] left_indices Device span of left table indices + * @param[in] right_indices Device span of right table indices + * @param[in] device_expression_data Device data required to evaluate the expression + * @param[in] config Grid configuration for kernel launch + * @param[in] shmem_per_block Amount of shared memory to allocate per block + * @param[in] join_kind The join kind. Must be INNER_JOIN, LEFT_JOIN, or FULL_JOIN. + * @param[out] count_out Atomic counter for the per-kind count described above + * @param[out] left_passing_marks Byte buffer of size `left_table.num_rows()` used by LEFT_JOIN + * to mark left rows whose entries contribute to `count_out`. Must be zero-initialized + * before the kernel launch and may be `nullptr` for INNER_JOIN and FULL_JOIN. + * @param[in] stream CUDA stream on which to launch the kernel + */ +template +void launch_filter_output_size_kernel( + cudf::table_device_view const& left_table, + cudf::table_device_view const& right_table, + cudf::device_span left_indices, + cudf::device_span right_indices, + cudf::ast::detail::expression_device_view device_expression_data, + cudf::detail::grid_1d const& config, + std::size_t shmem_per_block, + cudf::join_kind join_kind, + std::size_t* count_out, + bool* left_passing_marks, + rmm::cuda_stream_view stream); + +} // namespace cudf::detail diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cu b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cu new file mode 100644 index 000000000000..3bb93635d553 --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.hpp" + +namespace cudf::detail { +template void launch_filter_output_size_kernel( + cudf::table_device_view const& left_table, + cudf::table_device_view const& right_table, + cudf::device_span left_indices, + cudf::device_span right_indices, + cudf::ast::detail::expression_device_view device_expression_data, + cudf::detail::grid_1d const& config, + std::size_t shmem_per_block, + cudf::join_kind join_kind, + std::size_t* count_out, + bool* left_passing_marks, + rmm::cuda_stream_view stream); +} // namespace cudf::detail diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cu b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cu new file mode 100644 index 000000000000..195feaeeb65a --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.hpp" + +namespace cudf::detail { +template void launch_filter_output_size_kernel( + cudf::table_device_view const& left_table, + cudf::table_device_view const& right_table, + cudf::device_span left_indices, + cudf::device_span right_indices, + cudf::ast::detail::expression_device_view device_expression_data, + cudf::detail::grid_1d const& config, + std::size_t shmem_per_block, + cudf::join_kind join_kind, + std::size_t* count_out, + bool* left_passing_marks, + rmm::cuda_stream_view stream); +} // namespace cudf::detail diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cu b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cu new file mode 100644 index 000000000000..babeb76a3f82 --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.hpp" + +namespace cudf::detail { +template void launch_filter_output_size_kernel( + cudf::table_device_view const& left_table, + cudf::table_device_view const& right_table, + cudf::device_span left_indices, + cudf::device_span right_indices, + cudf::ast::detail::expression_device_view device_expression_data, + cudf::detail::grid_1d const& config, + std::size_t shmem_per_block, + cudf::join_kind join_kind, + std::size_t* count_out, + bool* left_passing_marks, + rmm::cuda_stream_view stream); +} // namespace cudf::detail diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu new file mode 100644 index 000000000000..f695652818fb --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu @@ -0,0 +1,22 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.hpp" + +namespace cudf::detail { +template void launch_filter_output_size_kernel( + cudf::table_device_view const& left_table, + cudf::table_device_view const& right_table, + cudf::device_span left_indices, + cudf::device_span right_indices, + cudf::ast::detail::expression_device_view device_expression_data, + cudf::detail::grid_1d const& config, + std::size_t shmem_per_block, + cudf::join_kind join_kind, + std::size_t* count_out, + bool* left_passing_marks, + rmm::cuda_stream_view stream); +} // namespace cudf::detail diff --git a/cpp/tests/join/mixed_join_tests.cu b/cpp/tests/join/mixed_join_tests.cu index a7c3a2603a27..4000af524064 100644 --- a/cpp/tests/join/mixed_join_tests.cu +++ b/cpp/tests/join/mixed_join_tests.cu @@ -433,6 +433,16 @@ struct MixedInnerJoinTest : public MixedJoinPairReturnTest { cudf::join_kind::INNER_JOIN); this->compare_join_results(mixed_result, ast_filter_result); + // Verify filter_join_indices_output_size matches the materialized output size. + auto const fji_size = cudf::filter_join_indices_output_size( + left_conditional, + right_conditional, + cudf::device_span(*hash_join_result.first), + cudf::device_span(*hash_join_result.second), + predicate, + cudf::join_kind::INNER_JOIN); + EXPECT_EQ(fji_size, ast_filter_result.first->size()); + // Verify JIT filter_join_indices if provided if (!jit_predicate.empty()) { auto jit_filter_result = cudf::filter_join_indices_jit( @@ -1091,6 +1101,16 @@ struct MixedLeftJoinTest : public MixedJoinPairReturnTest { cudf::join_kind::LEFT_JOIN); this->compare_join_results(mixed_result, ast_filter_result); + // Verify filter_join_indices_output_size matches the materialized output size. + auto const fji_size = cudf::filter_join_indices_output_size( + left_conditional, + right_conditional, + cudf::device_span(*hash_join_result.first), + cudf::device_span(*hash_join_result.second), + predicate, + cudf::join_kind::LEFT_JOIN); + EXPECT_EQ(fji_size, ast_filter_result.first->size()); + // Verify JIT filter_join_indices if provided if (!jit_predicate.empty()) { auto jit_filter_result = cudf::filter_join_indices_jit( @@ -1369,6 +1389,16 @@ struct MixedFullJoinTest : public MixedJoinPairReturnTest { cudf::join_kind::FULL_JOIN); this->compare_join_results(mixed_result, ast_filter_result); + // Verify filter_join_indices_output_size matches the materialized output size. + auto const fji_size = cudf::filter_join_indices_output_size( + left_conditional, + right_conditional, + cudf::device_span(*hash_join_result.first), + cudf::device_span(*hash_join_result.second), + predicate, + cudf::join_kind::FULL_JOIN); + EXPECT_EQ(fji_size, ast_filter_result.first->size()); + // Verify JIT filter_join_indices if provided if (!jit_predicate.empty()) { auto jit_filter_result = cudf::filter_join_indices_jit(