From 0d34c505e681cd3d5fa0ea614162b499d85e1995 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Thu, 28 May 2026 16:55:01 +0000 Subject: [PATCH 1/7] Add filter_join_indices_size --- cpp/CMakeLists.txt | 18 ++- cpp/include/cudf/join/join.hpp | 39 +++++ .../filter_join_indices.cu | 145 +++++++++++++++++- .../filter_join_indices_jit.cu | 4 +- .../filter_join_indices_jit_kernel.cuh | 0 .../filter_join_indices_kernel.cuh | 4 +- .../filter_join_indices_kernel.hpp | 0 .../filter_join_indices_kernel_complex.cu | 6 +- ...filter_join_indices_kernel_null_complex.cu | 6 +- ...lter_join_indices_kernel_null_primitive.cu | 6 +- .../filter_join_indices_kernel_primitive.cu | 6 +- .../filter_join_indices_size_kernel.cuh | 125 +++++++++++++++ .../filter_join_indices_size_kernel.hpp | 67 ++++++++ ...filter_join_indices_size_kernel_complex.cu | 22 +++ ...r_join_indices_size_kernel_null_complex.cu | 22 +++ ...join_indices_size_kernel_null_primitive.cu | 22 +++ ...lter_join_indices_size_kernel_primitive.cu | 22 +++ cpp/tests/join/mixed_join_tests.cu | 30 ++++ 18 files changed, 520 insertions(+), 24 deletions(-) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices.cu (70%) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices_jit.cu (99%) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices_jit_kernel.cuh (100%) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices_kernel.cuh (96%) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices_kernel.hpp (100%) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices_kernel_complex.cu (73%) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices_kernel_null_complex.cu (73%) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices_kernel_null_primitive.cu (73%) rename cpp/src/join/{ => filter_join_indices}/filter_join_indices_kernel_primitive.cu (73%) create mode 100644 cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh create mode 100644 cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.hpp create mode 100644 cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_complex.cu create mode 100644 cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_complex.cu create mode 100644 cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_primitive.cu create mode 100644 cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_primitive.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 63c292646773..6067396a5635 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -675,11 +675,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_size_kernel_complex.cu + src/join/filter_join_indices/filter_join_indices_size_kernel_null_complex.cu + src/join/filter_join_indices/filter_join_indices_size_kernel_null_primitive.cu + src/join/filter_join_indices/filter_join_indices_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 @@ -702,11 +707,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 16ed0ad0e2ec..a29f695594bf 100644 --- a/cpp/include/cudf/join/join.hpp +++ b/cpp/include/cudf/join/join.hpp @@ -346,6 +346,45 @@ 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. + * + * @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. + * @param mr Device memory resource used to allocate intermediate device memory. + * + * @return The exact number of pairs that `filter_join_indices` would produce. + */ +std::size_t filter_join_indices_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(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + /** * @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 70% rename from cpp/src/join/filter_join_indices.cu rename to cpp/src/join/filter_join_indices/filter_join_indices.cu index cb0e1670d57c..dce237aab2c6 100644 --- a/cpp/src/join/filter_join_indices.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices.cu @@ -3,7 +3,8 @@ * 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_size_kernel.hpp" #include #include @@ -355,6 +356,134 @@ filter_join_indices(cudf::table_view const& left, } } +std::size_t filter_join_indices_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, + rmm::device_async_resource_ref mr) +{ + // 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_size only supports INNER_JOIN, LEFT_JOIN, and FULL_JOIN.", + std::invalid_argument); + + if (left_indices.empty()) { 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"); + + 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); + + int device_id; + CUDF_CUDA_TRY(cudaGetDevice(&device_id)); + int shmem_limit_per_block; + CUDF_CUDA_TRY( + cudaDeviceGetAttribute(&shmem_limit_per_block, cudaDevAttrMaxSharedMemoryPerBlock, device_id)); + + auto const block_size = + parser.shmem_per_thread != 0 + ? std::min(MAX_BLOCK_SIZE, shmem_limit_per_block / parser.shmem_per_thread) + : MAX_BLOCK_SIZE; + + detail::grid_1d const config(left_indices.size(), block_size); + auto const shmem_per_block = parser.shmem_per_thread * config.num_threads_per_block; + + // The count kernel uses a single atomic counter. Allocate device_scalar zero-initialized. + cudf::detail::device_scalar d_count(std::size_t{0}, stream, mr); + + // For LEFT_JOIN, allocate per-left-row mark buffer; for others, pass nullptr. + rmm::device_uvector left_passing_marks( + join_kind == join_kind::LEFT_JOIN ? static_cast(left.num_rows()) : 0, stream); + if (join_kind == join_kind::LEFT_JOIN) { + CUDF_CUDA_TRY(cudaMemsetAsync( + left_passing_marks.data(), 0, left_passing_marks.size() * sizeof(bool), stream.value())); + } + auto* const marks_ptr = join_kind == join_kind::LEFT_JOIN ? left_passing_marks.data() : nullptr; + + if (has_nulls && has_complex_type) { + launch_filter_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); + } else if (has_nulls && !has_complex_type) { + launch_filter_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); + } else if (!has_nulls && has_complex_type) { + launch_filter_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); + } else { + launch_filter_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(), + [] __device__(bool m) -> bool { return m; }, + stream); + auto const num_invalid = + static_cast(left.num_rows()) - static_cast(num_filter_passing); + return num_predicate_passing + num_invalid; + } + default: CUDF_FAIL("Unsupported join kind for filter_join_indices_size"); + } +} + } // namespace detail // Public API implementation @@ -374,4 +503,18 @@ 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_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, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::filter_join_indices_size( + left, right, left_indices, right_indices, predicate, join_kind, stream, mr); +} + } // 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 15b532cac030..b2ead7bc3f29 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 96% 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..f09c5ce5f3ee 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 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_size_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh new file mode 100644 index 000000000000..30a2cb8863bd --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh @@ -0,0 +1,125 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ +#pragma once + +#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include + +namespace cudf::detail { + +/** + * @brief Kernel that counts the per-join-kind output size without materializing + * a per-pair boolean buffer. + */ +template +CUDF_KERNEL __launch_bounds__(max_block_size) void filter_join_indices_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, + unsigned long long* 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]; + + 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}; + + for (cudf::size_type 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 && both_valid) { atomicAdd(count_out, 1ULL); } + break; + case cudf::join_kind::LEFT_JOIN: + if (predicate_pass) { + atomicAdd(count_out, 1ULL); + // 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) { atomicAdd(count_out, 1ULL); } + break; + default: break; + } + } +} + +template +void launch_filter_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) +{ + static_assert(sizeof(std::size_t) == sizeof(unsigned long long), + "filter_size_kernel assumes size_t and unsigned long long are layout-compatible " + "for atomic counters."); + filter_join_indices_size_kernel + <<>>( + left_table, + right_table, + left_indices, + right_indices, + device_expression_data, + join_kind, + reinterpret_cast(count_out), + left_passing_marks); +} + +} // namespace cudf::detail diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.hpp b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.hpp new file mode 100644 index 000000000000..afa6b9cc6842 --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.hpp @@ -0,0 +1,67 @@ +/* + * 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_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_size_kernel_complex.cu b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_complex.cu new file mode 100644 index 000000000000..6548ea97299a --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_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_size_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" + +namespace cudf::detail { +template void launch_filter_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_size_kernel_null_complex.cu b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_complex.cu new file mode 100644 index 000000000000..308811ddb39f --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_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_size_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" + +namespace cudf::detail { +template void launch_filter_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_size_kernel_null_primitive.cu b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_primitive.cu new file mode 100644 index 000000000000..c6d2fe650e93 --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_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_size_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" + +namespace cudf::detail { +template void launch_filter_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_size_kernel_primitive.cu b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_primitive.cu new file mode 100644 index 000000000000..70e437ac347b --- /dev/null +++ b/cpp/src/join/filter_join_indices/filter_join_indices_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_size_kernel.cuh" +#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" + +namespace cudf::detail { +template void launch_filter_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 683035167396..0dc372e89214 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_size matches the materialized output size. + auto const fji_size = cudf::filter_join_indices_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( @@ -1089,6 +1099,16 @@ struct MixedLeftJoinTest : public MixedJoinPairReturnTest { cudf::join_kind::LEFT_JOIN); this->compare_join_results(mixed_result, ast_filter_result); + // Verify filter_join_indices_size matches the materialized output size. + auto const fji_size = cudf::filter_join_indices_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( @@ -1367,6 +1387,16 @@ struct MixedFullJoinTest : public MixedJoinPairReturnTest { cudf::join_kind::FULL_JOIN); this->compare_join_results(mixed_result, ast_filter_result); + // Verify filter_join_indices_size matches the materialized output size. + auto const fji_size = cudf::filter_join_indices_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( From 473f2c6f840e9488ed7118888b8f5472070c744a Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Thu, 28 May 2026 17:13:49 +0000 Subject: [PATCH 2/7] Kernel cleanups --- .../filter_join_indices_size_kernel.cuh | 69 ++++++++++++------- 1 file changed, 44 insertions(+), 25 deletions(-) diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh index 30a2cb8863bd..edd56f352703 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh +++ b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh @@ -17,23 +17,32 @@ #include +#include +#include + #include namespace cudf::detail { /** - * @brief Kernel that counts the per-join-kind output size without materializing + * @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__(max_block_size) void filter_join_indices_size_kernel( +template +CUDF_KERNEL __launch_bounds__(MAX_BLOCK_SIZE) void filter_join_indices_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, - unsigned long long* count_out, + std::size_t* count_out, bool* left_passing_marks) { extern __shared__ char raw_intermediate_storage[]; @@ -42,12 +51,17 @@ CUDF_KERNEL __launch_bounds__(max_block_size) void filter_join_indices_size_kern 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}; + std::size_t thread_local_count = 0; + for (cudf::size_type i = tid; i < static_cast(left_indices.size()); i += stride) { auto const left_row_index = left_indices[i]; @@ -71,13 +85,13 @@ CUDF_KERNEL __launch_bounds__(max_block_size) void filter_join_indices_size_kern switch (join_kind) { case cudf::join_kind::INNER_JOIN: - if (predicate_pass && both_valid) { atomicAdd(count_out, 1ULL); } + if (predicate_pass && both_valid) { ++thread_local_count; } break; case cudf::join_kind::LEFT_JOIN: if (predicate_pass) { - atomicAdd(count_out, 1ULL); - // 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 + ++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()) { @@ -87,30 +101,35 @@ CUDF_KERNEL __launch_bounds__(max_block_size) void filter_join_indices_size_kern break; case cudf::join_kind::FULL_JOIN: // Count failed matches: predicate false AND both indices valid. - if (both_valid && !predicate_pass) { atomicAdd(count_out, 1ULL); } + if (both_valid && !predicate_pass) { ++thread_local_count; } break; default: break; } } + + 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_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) +void launch_filter_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) { - static_assert(sizeof(std::size_t) == sizeof(unsigned long long), - "filter_size_kernel assumes size_t and unsigned long long are layout-compatible " - "for atomic counters."); - filter_join_indices_size_kernel + filter_join_indices_size_kernel <<>>( left_table, right_table, @@ -118,7 +137,7 @@ void launch_filter_size_kernel(cudf::table_device_view const& left_table, right_indices, device_expression_data, join_kind, - reinterpret_cast(count_out), + count_out, left_passing_marks); } From c2eed6e43a36b6ff36d5424858d013af83b7be9f Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Thu, 28 May 2026 18:05:02 +0000 Subject: [PATCH 3/7] Fix runtime block size determination --- .../filter_join_indices.cu | 17 ++------- .../filter_join_indices_size_kernel.cuh | 38 +++++++++---------- 2 files changed, 23 insertions(+), 32 deletions(-) diff --git a/cpp/src/join/filter_join_indices/filter_join_indices.cu b/cpp/src/join/filter_join_indices/filter_join_indices.cu index dce237aab2c6..27b1364bf610 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices.cu @@ -5,6 +5,7 @@ #include "join/filter_join_indices/filter_join_indices_kernel.cuh" #include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" +#include "join/join_common_utils.hpp" #include #include @@ -375,6 +376,7 @@ std::size_t filter_join_indices_size(cudf::table_view const& left, 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); @@ -389,19 +391,8 @@ std::size_t filter_join_indices_size(cudf::table_view const& left, auto left_table = table_device_view::create(left, stream); auto right_table = table_device_view::create(right, stream); - int device_id; - CUDF_CUDA_TRY(cudaGetDevice(&device_id)); - int shmem_limit_per_block; - CUDF_CUDA_TRY( - cudaDeviceGetAttribute(&shmem_limit_per_block, cudaDevAttrMaxSharedMemoryPerBlock, device_id)); - - auto const block_size = - parser.shmem_per_thread != 0 - ? std::min(MAX_BLOCK_SIZE, shmem_limit_per_block / parser.shmem_per_thread) - : MAX_BLOCK_SIZE; - - detail::grid_1d const config(left_indices.size(), block_size); - auto const shmem_per_block = parser.shmem_per_thread * config.num_threads_per_block; + 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, mr); diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh index edd56f352703..c245a4ee25f5 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh +++ b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh @@ -5,6 +5,7 @@ #pragma once #include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" +#include "join/join_common_utils.hpp" #include #include @@ -19,6 +20,7 @@ #include #include +#include #include @@ -35,7 +37,7 @@ namespace cudf::detail { * entries. */ template -CUDF_KERNEL __launch_bounds__(MAX_BLOCK_SIZE) void filter_join_indices_size_kernel( +CUDF_KERNEL __launch_bounds__(DEFAULT_JOIN_BLOCK_SIZE) void filter_join_indices_size_kernel( cudf::table_device_view left_table, cudf::table_device_view right_table, cudf::device_span left_indices, @@ -51,7 +53,7 @@ CUDF_KERNEL __launch_bounds__(MAX_BLOCK_SIZE) void filter_join_indices_size_kern auto thread_intermediate_storage = &intermediate_storage[threadIdx.x * device_expression_data.num_intermediates]; - using BlockReduce = cub::BlockReduce; + using BlockReduce = cub::BlockReduce; __shared__ typename BlockReduce::TempStorage temp_storage; auto const tid = cudf::detail::grid_1d::global_thread_id(); @@ -60,10 +62,9 @@ CUDF_KERNEL __launch_bounds__(MAX_BLOCK_SIZE) void filter_join_indices_size_kern auto evaluator = cudf::ast::detail::expression_evaluator{ left_table, right_table, device_expression_data}; - std::size_t thread_local_count = 0; + cuda::std::size_t thread_local_count = 0; - for (cudf::size_type i = tid; i < static_cast(left_indices.size()); - i += stride) { + 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]; @@ -107,27 +108,26 @@ CUDF_KERNEL __launch_bounds__(MAX_BLOCK_SIZE) void filter_join_indices_size_kern } } - std::size_t const block_sum = BlockReduce(temp_storage).Sum(thread_local_count); + 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}; + cuda::atomic_ref count_ref{*count_out}; count_ref.fetch_add(block_sum, cuda::memory_order_relaxed); } } template -void launch_filter_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) +void launch_filter_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_size_kernel <<>>( From 91ea31a3b18b4dda1b5e8a3a466b7dff620ffd42 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 1 Jun 2026 22:29:16 +0000 Subject: [PATCH 4/7] Updates --- cpp/include/cudf/join/join.hpp | 6 ++---- .../join/filter_join_indices/filter_join_indices.cu | 11 +++++------ .../filter_join_indices_size_kernel.cuh | 2 +- 3 files changed, 8 insertions(+), 11 deletions(-) diff --git a/cpp/include/cudf/join/join.hpp b/cpp/include/cudf/join/join.hpp index a29f695594bf..2818402643fd 100644 --- a/cpp/include/cudf/join/join.hpp +++ b/cpp/include/cudf/join/join.hpp @@ -371,19 +371,17 @@ filter_join_indices(cudf::table_view const& left, * @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. - * @param mr Device memory resource used to allocate intermediate device memory. * * @return The exact number of pairs that `filter_join_indices` would produce. */ -std::size_t filter_join_indices_size( +[[nodiscard]] std::size_t filter_join_indices_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(), - rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + 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/filter_join_indices.cu b/cpp/src/join/filter_join_indices/filter_join_indices.cu index 27b1364bf610..c048070b2ca9 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices.cu @@ -363,8 +363,7 @@ std::size_t filter_join_indices_size(cudf::table_view const& left, cudf::device_span right_indices, ast::expression const& predicate, join_kind join_kind, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + rmm::cuda_stream_view stream) { // Validate inputs (same constraints as filter_join_indices) CUDF_EXPECTS(left_indices.size() == right_indices.size(), @@ -395,7 +394,8 @@ std::size_t filter_join_indices_size(cudf::table_view const& left, 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, mr); + cudf::detail::device_scalar d_count( + std::size_t{0}, stream, cudf::get_current_device_resource_ref()); // For LEFT_JOIN, allocate per-left-row mark buffer; for others, pass nullptr. rmm::device_uvector left_passing_marks( @@ -500,12 +500,11 @@ std::size_t filter_join_indices_size(cudf::table_view const& left, cudf::device_span right_indices, ast::expression const& predicate, cudf::join_kind join_kind, - rmm::cuda_stream_view stream, - rmm::device_async_resource_ref mr) + rmm::cuda_stream_view stream) { CUDF_FUNC_RANGE(); return detail::filter_join_indices_size( - left, right, left_indices, right_indices, predicate, join_kind, stream, mr); + left, right, left_indices, right_indices, predicate, join_kind, stream); } } // namespace cudf diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh index c245a4ee25f5..3f1990310482 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh +++ b/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh @@ -86,7 +86,7 @@ CUDF_KERNEL __launch_bounds__(DEFAULT_JOIN_BLOCK_SIZE) void filter_join_indices_ switch (join_kind) { case cudf::join_kind::INNER_JOIN: - if (predicate_pass && both_valid) { ++thread_local_count; } + if (predicate_pass) { ++thread_local_count; } break; case cudf::join_kind::LEFT_JOIN: if (predicate_pass) { From a3058e6786f947c34e413f039a3492563689dd03 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 2 Jun 2026 21:38:59 +0000 Subject: [PATCH 5/7] Renaming --- cpp/CMakeLists.txt | 8 +- cpp/include/cudf/join/join.hpp | 2 +- .../filter_join_indices.cu | 131 +++++++++--------- .../filter_join_indices_kernel.cuh | 1 + ...ilter_join_indices_output_size_kernel.cuh} | 30 ++-- ...ilter_join_indices_output_size_kernel.hpp} | 23 +-- ...oin_indices_output_size_kernel_complex.cu} | 6 +- ...ndices_output_size_kernel_null_complex.cu} | 6 +- ...ices_output_size_kernel_null_primitive.cu} | 6 +- ...n_indices_output_size_kernel_primitive.cu} | 6 +- cpp/tests/join/mixed_join_tests.cu | 12 +- 11 files changed, 118 insertions(+), 113 deletions(-) rename cpp/src/join/filter_join_indices/{filter_join_indices_size_kernel.cuh => filter_join_indices_output_size_kernel.cuh} (84%) rename cpp/src/join/filter_join_indices/{filter_join_indices_size_kernel.hpp => filter_join_indices_output_size_kernel.hpp} (76%) rename cpp/src/join/filter_join_indices/{filter_join_indices_size_kernel_primitive.cu => filter_join_indices_output_size_kernel_complex.cu} (74%) rename cpp/src/join/filter_join_indices/{filter_join_indices_size_kernel_null_primitive.cu => filter_join_indices_output_size_kernel_null_complex.cu} (74%) rename cpp/src/join/filter_join_indices/{filter_join_indices_size_kernel_complex.cu => filter_join_indices_output_size_kernel_null_primitive.cu} (74%) rename cpp/src/join/filter_join_indices/{filter_join_indices_size_kernel_null_complex.cu => filter_join_indices_output_size_kernel_primitive.cu} (74%) diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 3a5268717b33..3ae9865881d0 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -681,10 +681,10 @@ add_library( 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_size_kernel_complex.cu - src/join/filter_join_indices/filter_join_indices_size_kernel_null_complex.cu - src/join/filter_join_indices/filter_join_indices_size_kernel_null_primitive.cu - src/join/filter_join_indices/filter_join_indices_size_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 diff --git a/cpp/include/cudf/join/join.hpp b/cpp/include/cudf/join/join.hpp index 2818402643fd..82ed8657fc8f 100644 --- a/cpp/include/cudf/join/join.hpp +++ b/cpp/include/cudf/join/join.hpp @@ -374,7 +374,7 @@ filter_join_indices(cudf::table_view const& left, * * @return The exact number of pairs that `filter_join_indices` would produce. */ -[[nodiscard]] std::size_t filter_join_indices_size( +[[nodiscard]] std::size_t filter_join_indices_output_size( cudf::table_view const& left, cudf::table_view const& right, cudf::device_span left_indices, diff --git a/cpp/src/join/filter_join_indices/filter_join_indices.cu b/cpp/src/join/filter_join_indices/filter_join_indices.cu index c048070b2ca9..1ce37ea9a095 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices.cu @@ -4,7 +4,7 @@ */ #include "join/filter_join_indices/filter_join_indices_kernel.cuh" -#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.hpp" #include "join/join_common_utils.hpp" #include @@ -357,22 +357,23 @@ filter_join_indices(cudf::table_view const& left, } } -std::size_t filter_join_indices_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) +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_size only supports INNER_JOIN, LEFT_JOIN, and FULL_JOIN.", - 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; } @@ -407,53 +408,53 @@ std::size_t filter_join_indices_size(cudf::table_view const& left, auto* const marks_ptr = join_kind == join_kind::LEFT_JOIN ? left_passing_marks.data() : nullptr; if (has_nulls && has_complex_type) { - launch_filter_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); + 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); } else if (has_nulls && !has_complex_type) { - launch_filter_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); + 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); } else if (!has_nulls && has_complex_type) { - launch_filter_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); + 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); } else { - launch_filter_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); + 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); @@ -471,7 +472,7 @@ std::size_t filter_join_indices_size(cudf::table_view const& left, static_cast(left.num_rows()) - static_cast(num_filter_passing); return num_predicate_passing + num_invalid; } - default: CUDF_FAIL("Unsupported join kind for filter_join_indices_size"); + default: CUDF_FAIL("Unsupported join kind for filter_join_indices_output_size"); } } @@ -494,16 +495,16 @@ 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_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) +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_size( + return detail::filter_join_indices_output_size( left, right, left_indices, right_indices, predicate, join_kind, stream); } diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_kernel.cuh index f09c5ce5f3ee..9ea25a4336f5 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_kernel.cuh +++ b/cpp/src/join/filter_join_indices/filter_join_indices_kernel.cuh @@ -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/filter_join_indices_size_kernel.cuh b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh similarity index 84% rename from cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh rename to cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh index 3f1990310482..e158e3883662 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh @@ -4,7 +4,7 @@ */ #pragma once -#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" +#include "join/filter_join_indices/filter_join_indices_output_size_kernel.hpp" #include "join/join_common_utils.hpp" #include @@ -37,7 +37,7 @@ namespace cudf::detail { * entries. */ template -CUDF_KERNEL __launch_bounds__(DEFAULT_JOIN_BLOCK_SIZE) void filter_join_indices_size_kernel( +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, @@ -117,19 +117,20 @@ CUDF_KERNEL __launch_bounds__(DEFAULT_JOIN_BLOCK_SIZE) void filter_join_indices_ } template -void launch_filter_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) +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_size_kernel + filter_join_indices_output_size_kernel <<>>( left_table, right_table, @@ -139,6 +140,7 @@ void launch_filter_size_kernel(cudf::table_device_view const& left_table, 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_size_kernel.hpp b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.hpp similarity index 76% rename from cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.hpp rename to cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.hpp index afa6b9cc6842..b5b25e5e7a0e 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.hpp +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.hpp @@ -52,16 +52,17 @@ namespace cudf::detail { * @param[in] stream CUDA stream on which to launch the kernel */ template -void launch_filter_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); +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_size_kernel_primitive.cu b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cu similarity index 74% rename from cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_primitive.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cu index 70e437ac347b..3bb93635d553 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_primitive.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cu @@ -3,11 +3,11 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "join/filter_join_indices/filter_join_indices_size_kernel.cuh" -#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" +#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_size_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, diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_primitive.cu b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cu similarity index 74% rename from cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_primitive.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cu index c6d2fe650e93..195feaeeb65a 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_primitive.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cu @@ -3,11 +3,11 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "join/filter_join_indices/filter_join_indices_size_kernel.cuh" -#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" +#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_size_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, diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_complex.cu b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cu similarity index 74% rename from cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_complex.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cu index 6548ea97299a..babeb76a3f82 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_complex.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cu @@ -3,11 +3,11 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "join/filter_join_indices/filter_join_indices_size_kernel.cuh" -#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" +#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_size_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, diff --git a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_complex.cu b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu similarity index 74% rename from cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_complex.cu rename to cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu index 308811ddb39f..f695652818fb 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_complex.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu @@ -3,11 +3,11 @@ * SPDX-License-Identifier: Apache-2.0 */ -#include "join/filter_join_indices/filter_join_indices_size_kernel.cuh" -#include "join/filter_join_indices/filter_join_indices_size_kernel.hpp" +#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_size_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, diff --git a/cpp/tests/join/mixed_join_tests.cu b/cpp/tests/join/mixed_join_tests.cu index 92c67336717f..4000af524064 100644 --- a/cpp/tests/join/mixed_join_tests.cu +++ b/cpp/tests/join/mixed_join_tests.cu @@ -433,8 +433,8 @@ struct MixedInnerJoinTest : public MixedJoinPairReturnTest { cudf::join_kind::INNER_JOIN); this->compare_join_results(mixed_result, ast_filter_result); - // Verify filter_join_indices_size matches the materialized output size. - auto const fji_size = cudf::filter_join_indices_size( + // 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), @@ -1101,8 +1101,8 @@ struct MixedLeftJoinTest : public MixedJoinPairReturnTest { cudf::join_kind::LEFT_JOIN); this->compare_join_results(mixed_result, ast_filter_result); - // Verify filter_join_indices_size matches the materialized output size. - auto const fji_size = cudf::filter_join_indices_size( + // 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), @@ -1389,8 +1389,8 @@ struct MixedFullJoinTest : public MixedJoinPairReturnTest { cudf::join_kind::FULL_JOIN); this->compare_join_results(mixed_result, ast_filter_result); - // Verify filter_join_indices_size matches the materialized output size. - auto const fji_size = cudf::filter_join_indices_size( + // 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), From 10959229accf8d19b51df137d90b7687beffeb76 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Mon, 8 Jun 2026 21:09:03 +0000 Subject: [PATCH 6/7] Document and enforce Boolean predicate validation in filter_join_indices --- cpp/include/cudf/join/join.hpp | 2 ++ cpp/src/join/filter_join_indices/filter_join_indices.cu | 6 ++++-- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/cpp/include/cudf/join/join.hpp b/cpp/include/cudf/join/join.hpp index 82ed8657fc8f..e83d2d3758be 100644 --- a/cpp/include/cudf/join/join.hpp +++ b/cpp/include/cudf/join/join.hpp @@ -322,6 +322,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). @@ -363,6 +364,7 @@ filter_join_indices(cudf::table_view const& left, * * @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). diff --git a/cpp/src/join/filter_join_indices/filter_join_indices.cu b/cpp/src/join/filter_join_indices/filter_join_indices.cu index 1ce37ea9a095..4a96d046c8a5 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices.cu @@ -78,7 +78,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(); @@ -384,7 +385,8 @@ std::size_t filter_join_indices_output_size(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); auto const has_complex_type = parser.has_complex_type(); From fdfd4ed087e063648c7c399ec8d5c6145d4413d8 Mon Sep 17 00:00:00 2001 From: Yunsong Wang Date: Tue, 9 Jun 2026 18:25:09 +0000 Subject: [PATCH 7/7] Cleanups with existing utiks --- .../filter_join_indices.cu | 89 ++++++------------- 1 file changed, 27 insertions(+), 62 deletions(-) diff --git a/cpp/src/join/filter_join_indices/filter_join_indices.cu b/cpp/src/join/filter_join_indices/filter_join_indices.cu index 4a96d046c8a5..79ed59797fd4 100644 --- a/cpp/src/join/filter_join_indices/filter_join_indices.cu +++ b/cpp/src/join/filter_join_indices/filter_join_indices.cu @@ -15,7 +15,9 @@ #include #include #include +#include #include +#include #include #include #include @@ -33,6 +35,7 @@ #include #include #include +#include #include #include @@ -400,64 +403,30 @@ std::size_t filter_join_indices_output_size(cudf::table_view const& left, cudf::detail::device_scalar d_count( std::size_t{0}, stream, cudf::get_current_device_resource_ref()); - // For LEFT_JOIN, allocate per-left-row mark buffer; for others, pass nullptr. - rmm::device_uvector left_passing_marks( - join_kind == join_kind::LEFT_JOIN ? static_cast(left.num_rows()) : 0, stream); - if (join_kind == join_kind::LEFT_JOIN) { - CUDF_CUDA_TRY(cudaMemsetAsync( - left_passing_marks.data(), 0, left_passing_marks.size() * sizeof(bool), stream.value())); - } + // 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; - if (has_nulls && has_complex_type) { - 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); - } else if (has_nulls && !has_complex_type) { - 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); - } else if (!has_nulls && has_complex_type) { - 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); - } else { - 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); - } + 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); @@ -466,12 +435,8 @@ std::size_t filter_join_indices_output_size(cudf::table_view const& left, 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(), - [] __device__(bool m) -> bool { return m; }, - stream); - auto const num_invalid = - static_cast(left.num_rows()) - static_cast(num_filter_passing); + 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");