Skip to content

Rewrite mixed inner/left/full join with post-filtering - #23012

Merged
rapids-bot[bot] merged 23 commits into
NVIDIA:release/26.08from
PointKernel:mixed-join-via-filter-indices
Jul 22, 2026
Merged

Rewrite mixed inner/left/full join with post-filtering#23012
rapids-bot[bot] merged 23 commits into
NVIDIA:release/26.08from
PointKernel:mixed-join-via-filter-indices

Conversation

@PointKernel

@PointKernel PointKernel commented Jun 26, 2026

Copy link
Copy Markdown
Member

Description

Part of #22124

This PR rewrites the mixed inner/left/full joins as equality-then-filter: run the keys through cudf::hash_join, then apply the conditional predicate to the index pairs via a new filter_join_indices primitive (AST + JIT, plus filter_join_indices_output_size). mixed_full_join becomes a left join plus finalize_full_join.

This replaces the fused mixed-join kernels, reusing hash_join instead of duplicating it.

Checklist

  • I am familiar with the Contributing Guidelines.
  • New or existing tests cover these changes.
  • The documentation is up to date with these changes.

@copy-pr-bot

copy-pr-bot Bot commented Jun 26, 2026

Copy link
Copy Markdown

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

@github-actions github-actions Bot added libcudf Affects libcudf (C++/CUDA) code. CMake CMake build issue labels Jun 26, 2026
@PointKernel PointKernel added improvement Improvement / enhancement to an existing function non-breaking Non-breaking change breaking Breaking change and removed non-breaking Non-breaking change labels Jun 26, 2026
@PointKernel

Copy link
Copy Markdown
Member Author

/ok to test edd3b59

@PointKernel

Copy link
Copy Markdown
Member Author

/ok to test 130ffd2

@PointKernel

Copy link
Copy Markdown
Member Author

/ok to test e9cc88d

@PointKernel

Copy link
Copy Markdown
Member Author

/ok to test f2705ac

@PointKernel
PointKernel marked this pull request as ready for review June 29, 2026 20:00
@PointKernel
PointKernel requested review from a team as code owners June 29, 2026 20:00
@coderabbitai

coderabbitai Bot commented Jun 29, 2026

Copy link
Copy Markdown

Review Change Stack

Note

Reviews paused

It looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the reviews.auto_review.auto_pause_after_reviewed_commits setting.

Use the following commands to manage reviews:

  • @coderabbitai resume to resume automatic reviews.
  • @coderabbitai review to trigger a single review.

Use the checkboxes below for quick actions:

  • ▶️ Resume reviews
  • 🔍 Trigger review
📝 Walkthrough

Walkthrough

The PR replaces bespoke mixed-join execution with hash_join and filter_join_indices, adds optional precomputed output sizing, changes output-size APIs to return contribution counts, removes obsolete build inputs, and updates mixed-join tests.

Changes

Mixed Join Refactor via filter_join_indices

Layer / File(s) Summary
Public filter-join API changes
cpp/include/cudf/join/join.hpp, cpp/include/cudf/detail/join/join.hpp
Adds output-size-aware filtering and changes output-size computation to return total size plus contribution counts.
Per-output sizing kernel
cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.*, cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_*.cu
Replaces scalar counting and LEFT_JOIN marking with per-output contribution counts and updates launcher instantiations.
Filter-join sizing and execution
cpp/src/join/filter_join_indices/filter_join_indices.cu
Threads optional output sizes through join execution, allocates and reduces contribution counts, and updates public wrappers.
Hash-join-backed mixed joins
cpp/src/join/mixed_join.cu, cpp/src/join/mixed_join_common_utils.cuh, cpp/src/join/mixed_filter_join_common_utils.cuh, cpp/src/join/mixed_join_semi.cu
Delegates mixed-join probing and conditional filtering to hash_join and filter_join_indices, while removing obsolete helper types and kernels.
Build cleanup and validation
cpp/CMakeLists.txt, cpp/tests/join/mixed_join_tests.cu, cpp/tests/streams/join_test.cpp
Removes obsolete mixed-join build inputs and updates sizing, full-join, stream-call, and unmatched-deduplication coverage.

Estimated code review effort: 5 (Critical) | ~120 minutes

Possibly related issues

Possibly related PRs

  • rapidsai/cudf#22694: Directly relates to the output-size API and kernel plumbing refactored here.

Suggested labels: 4 - Needs Review

Suggested reviewers: vyasr, vuule, lamarrr

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name Status Explanation Resolution
Docstring Coverage ⚠️ Warning Docstring coverage is 25.00% which is insufficient. The required threshold is 80.00%. Write docstrings for the functions missing them to satisfy the coverage threshold.
✅ Passed checks (4 passed)
Check name Status Explanation
Linked Issues check ✅ Passed Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check ✅ Passed Check skipped because no linked issues were found for this pull request.
Title check ✅ Passed The title clearly summarizes the main change: rewiring mixed inner, left, and full joins to post-filtering.
Description check ✅ Passed The description directly matches the implemented rewrite, including hash_join, filter_join_indices, and mixed_full_join changes.
✨ Finishing Touches
🧪 Generate unit tests (beta)
  • Create PR with unit tests

Comment @coderabbitai help to get the list of available commands.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 3

🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.

Inline comments:
In `@cpp/include/cudf/join/join.hpp`:
- Around line 425-434: The public API change in filter_join_indices_output_size
is source-breaking because it alters both the return type and the signature in a
public header. Preserve the existing contract by keeping the current
filter_join_indices_output_size declaration as a deprecated wrapper, or
introduce the new behavior under a different name and have the old function
forward to it. Make sure the symbols filter_join_indices_output_size,
cudf::table_view, and the join_kind-based overload remain available for
downstream callers during the transition.
- Around line 372-390: The public filter_join_indices overload currently trusts
output_size, which can lead to out-of-bounds writes or unsafe size arithmetic in
the implementation. Update filter_join_indices to validate the hint before use,
especially for the INNER_JOIN path and the LEFT/FULL size calculations, and fail
cleanly on mismatch rather than proceeding with unchecked allocation/copy
bounds. If this contract cannot be enforced safely, keep the size-hinting fast
path internal instead of exposing it in the public API.

In `@cpp/tests/join/mixed_join_tests.cu`:
- Around line 274-281: The mixed join test is only checking the sum of
`actual_counts`, which can miss incorrect per-row count placement; keep the
existing total check in `mixed_join_tests.cu`, but also validate `actual_counts`
element-wise against `expected_counts` in the test that uses `thrust::reduce`
and `result_size`. Update the assertions around the `actual_counts` vector so
the contract for output counts is verified entry-by-entry, not just by total.
🪄 Autofix (Beta)

Fix all unresolved CodeRabbit comments on this PR:

  • Push a commit to this branch (recommended)
  • Create a new PR with the fixes

ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 9c54aefe-bb98-416a-b078-7748b1a1a542

📥 Commits

Reviewing files that changed from the base of the PR and between b314d96 and f2705ac.

📒 Files selected for processing (23)
  • cpp/CMakeLists.txt
  • cpp/include/cudf/detail/join/join.hpp
  • cpp/include/cudf/join/join.hpp
  • cpp/src/join/filter_join_indices/filter_join_indices.cu
  • cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh
  • cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.hpp
  • cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cu
  • cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cu
  • cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cu
  • cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu
  • cpp/src/join/mixed_filter_join_common_utils.cuh
  • cpp/src/join/mixed_join.cu
  • cpp/src/join/mixed_join_common_utils.cuh
  • cpp/src/join/mixed_join_kernel.cu
  • cpp/src/join/mixed_join_kernel.cuh
  • cpp/src/join/mixed_join_kernel.hpp
  • cpp/src/join/mixed_join_kernel_nulls.cu
  • cpp/src/join/mixed_join_semi.cu
  • cpp/src/join/mixed_join_size_kernel.cu
  • cpp/src/join/mixed_join_size_kernel.cuh
  • cpp/src/join/mixed_join_size_kernel.hpp
  • cpp/src/join/mixed_join_size_kernel_nulls.cu
  • cpp/tests/join/mixed_join_tests.cu
💤 Files with no reviewable changes (9)
  • cpp/src/join/mixed_join_kernel.hpp
  • cpp/src/join/mixed_join_size_kernel.hpp
  • cpp/src/join/mixed_join_size_kernel.cu
  • cpp/src/join/mixed_join_size_kernel.cuh
  • cpp/src/join/mixed_join_kernel.cuh
  • cpp/src/join/mixed_join_kernel.cu
  • cpp/src/join/mixed_join_size_kernel_nulls.cu
  • cpp/CMakeLists.txt
  • cpp/src/join/mixed_join_kernel_nulls.cu

Comment thread cpp/include/cudf/join/join.hpp Outdated
Comment thread cpp/include/cudf/join/join.hpp
Comment thread cpp/tests/join/mixed_join_tests.cu
@PointKernel

PointKernel commented Jun 29, 2026

Copy link
Copy Markdown
Member Author

Performance impact of the post-filter mixed join. The slowdown is significant for small input sizes. In contrast, for large datasets, the optimization improves performance by several milliseconds.

# mixed_inner_join

## [0] NVIDIA GH200 480GB

|  Nullable  |  NullEquality  |  DataType  |  left_size  |  right_size  |  skip_large_sizes  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |         Diff |   %Diff |  Status  |
|------------|----------------|------------|-------------|--------------|--------------------|------------|-------------|------------|-------------|--------------|---------|----------|
|     0      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |         1          | 172.559 us |       2.27% | 189.602 us |       6.01% |    17.043 us |   9.88% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |         1          | 187.236 us |       2.82% | 212.595 us |       2.58% |    25.359 us |  13.54% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |         1          |   3.108 ms |       0.69% |   1.804 ms |       0.94% | -1304.322 us | -41.96% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |         1          | 198.391 us |       2.24% | 214.141 us |       2.28% |    15.750 us |   7.94% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |         1          |   3.624 ms |       0.14% |   1.954 ms |       0.40% | -1670.492 us | -46.09% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |         1          |   5.379 ms |       0.08% |   2.886 ms |       0.21% | -2492.647 us | -46.34% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |         1          | 188.708 us |       2.65% | 207.531 us |       3.29% |    18.822 us |   9.97% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |         1          |   2.841 ms |       0.34% |   1.577 ms |       0.47% | -1263.806 us | -44.48% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |         1          | 200.231 us |       2.20% | 211.185 us |       2.97% |    10.953 us |   5.47% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |         1          |   3.317 ms |       0.50% |   1.672 ms |       0.41% | -1645.244 us | -49.60% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |         1          |   4.594 ms |       0.10% |   2.424 ms |       0.36% | -2169.598 us | -47.23% |   FAST   |

# mixed_left_join

## [0] NVIDIA GH200 480GB

|  Nullable  |  NullEquality  |  DataType  |  left_size  |  right_size  |  skip_large_sizes  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |         Diff |   %Diff |  Status  |
|------------|----------------|------------|-------------|--------------|--------------------|------------|-------------|------------|-------------|--------------|---------|----------|
|     0      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |         1          | 171.806 us |       2.23% | 210.581 us |       3.22% |    38.775 us |  22.57% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |         1          | 187.319 us |       2.73% | 235.757 us |       2.76% |    48.439 us |  25.86% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |         1          |   3.146 ms |       0.76% |   2.565 ms |       0.46% |  -580.923 us | -18.47% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |         1          | 198.430 us |       2.21% | 237.643 us |       2.45% |    39.213 us |  19.76% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |         1          |   3.674 ms |       0.17% |   2.726 ms |       0.31% |  -947.858 us | -25.80% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |         1          |   5.445 ms |       0.11% |   3.670 ms |       0.23% | -1775.244 us | -32.60% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |         1          | 173.580 us |       2.15% | 227.574 us |       3.22% |    53.994 us |  31.11% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |         1          | 189.698 us |       2.65% | 264.602 us |       3.80% |    74.904 us |  39.49% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |         1          |   2.904 ms |       0.32% |   2.648 ms |       0.32% |  -255.467 us |  -8.80% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |         1          | 201.474 us |       2.12% | 266.238 us |       1.88% |    64.763 us |  32.14% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |         1          |   3.386 ms |       0.18% |   2.741 ms |       0.27% |  -644.405 us | -19.03% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |         1          |   4.692 ms |       0.09% |   3.514 ms |       0.20% | -1177.481 us | -25.10% |   FAST   |

# mixed_full_join

## [0] NVIDIA GH200 480GB

|  Nullable  |  NullEquality  |  DataType  |  left_size  |  right_size  |  skip_large_sizes  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |         Diff |   %Diff |  Status  |
|------------|----------------|------------|-------------|--------------|--------------------|------------|-------------|------------|-------------|--------------|---------|----------|
|     0      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |         1          | 215.693 us |      22.83% | 248.333 us |       2.93% |    32.640 us |  15.13% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |         1          | 233.240 us |       1.85% | 279.900 us |       2.19% |    46.661 us |  20.01% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |         1          |   3.480 ms |       0.72% |   2.900 ms |       0.39% |  -580.045 us | -16.67% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |         1          | 240.538 us |       1.44% | 280.086 us |       1.98% |    39.548 us |  16.44% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |         1          |   3.819 ms |       0.21% |   2.872 ms |       0.27% |  -947.415 us | -24.81% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |         1          |   5.653 ms |       0.11% |   3.882 ms |       0.44% | -1770.964 us | -31.33% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |         1          | 212.183 us |       2.32% | 269.147 us |       2.41% |    56.964 us |  26.85% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |         1          | 233.751 us |       2.15% | 298.833 us |       2.79% |    65.082 us |  27.84% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |         1          |   3.189 ms |       0.34% |   2.912 ms |       0.33% |  -276.821 us |  -8.68% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |         1          | 244.179 us |       1.42% | 300.409 us |       2.46% |    56.229 us |  23.03% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |         1          |   3.509 ms |       0.15% |   2.860 ms |       0.25% |  -649.248 us | -18.50% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |         1          |   4.881 ms |       0.11% |   3.700 ms |       0.20% | -1180.939 us | -24.19% |   FAST   |

# mixed_inner_join_complex_ast

## [0] NVIDIA GH200 480GB

|  Nullable  |  NullEquality  |  DataType  |  left_size  |  right_size  |  ast_levels  |  skip_large_sizes  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |          Diff |   %Diff |  Status  |
|------------|----------------|------------|-------------|--------------|--------------|--------------------|------------|-------------|------------|-------------|---------------|---------|----------|
|     0      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |      1       |         1          | 168.497 us |       2.99% | 188.148 us |       3.77% |     19.651 us |  11.66% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |      1       |         1          | 186.648 us |       2.78% | 212.969 us |       2.42% |     26.322 us |  14.10% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |      1       |         1          |   3.111 ms |       0.74% |   1.799 ms |       0.67% |  -1312.330 us | -42.19% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |      1       |         1          | 197.925 us |       2.24% | 215.645 us |       1.97% |     17.720 us |   8.95% |   SLOW   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |      1       |         1          |   3.626 ms |       0.14% |   1.947 ms |       0.27% |  -1678.559 us | -46.30% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |      1       |         1          |   5.379 ms |       0.08% |   2.887 ms |       0.23% |  -2491.350 us | -46.32% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |      5       |         1          | 220.161 us |       2.78% | 196.390 us |       3.76% |    -23.771 us | -10.80% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |      5       |         1          | 285.577 us |       2.45% | 221.504 us |       2.48% |    -64.073 us | -22.44% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |      5       |         1          |   9.577 ms |       0.31% |   1.909 ms |       0.61% |  -7668.405 us | -80.07% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |      5       |         1          | 291.851 us |       2.23% | 224.908 us |       2.04% |    -66.943 us | -22.94% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |      5       |         1          |   9.871 ms |       0.13% |   2.073 ms |       0.26% |  -7797.853 us | -79.00% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |      5       |         1          |  11.600 ms |       0.11% |   2.998 ms |       0.23% |  -8601.422 us | -74.15% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |      10      |         1          | 277.096 us |       2.99% | 203.596 us |       3.59% |    -73.500 us | -26.53% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |      10      |         1          | 401.719 us |       3.01% | 229.907 us |       3.58% |   -171.811 us | -42.77% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |      10      |         1          |  17.600 ms |       0.36% |   2.049 ms |       0.57% | -15551.106 us | -88.36% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |      10      |         1          | 401.447 us |       2.10% | 231.623 us |       2.26% |   -169.824 us | -42.30% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |      10      |         1          |  17.530 ms |       0.14% |   2.222 ms |       0.27% | -15308.469 us | -87.33% |   FAST   |
|     0      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |      10      |         1          |  19.331 ms |       0.08% |   3.150 ms |       0.21% | -16181.247 us | -83.71% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |      1       |         1          | 169.713 us |       3.11% | 184.159 us |       3.73% |     14.446 us |   8.51% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |      1       |         1          | 188.472 us |       4.29% | 208.310 us |       3.13% |     19.838 us |  10.53% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |      1       |         1          |   2.842 ms |       0.33% |   1.578 ms |       0.50% |  -1264.026 us | -44.48% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |      1       |         1          | 199.909 us |       2.21% | 211.955 us |       2.98% |     12.046 us |   6.03% |   SLOW   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |      1       |         1          |   3.317 ms |       0.16% |   1.672 ms |       0.43% |  -1644.518 us | -49.58% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |      1       |         1          |   4.593 ms |       0.09% |   2.425 ms |       0.27% |  -2168.072 us | -47.21% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |      5       |         1          | 224.559 us |       3.68% | 193.649 us |       3.72% |    -30.910 us | -13.76% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |      5       |         1          | 283.811 us |       2.48% | 218.918 us |       3.06% |    -64.893 us | -22.86% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |      5       |         1          |   8.420 ms |       0.22% |   1.702 ms |       0.42% |  -6718.072 us | -79.79% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |      5       |         1          | 293.366 us |       2.05% | 222.721 us |       2.89% |    -70.645 us | -24.08% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |      5       |         1          |   8.860 ms |       0.15% |   1.815 ms |       0.40% |  -7045.918 us | -79.52% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |      5       |         1          |  10.209 ms |       0.10% |   2.585 ms |       0.26% |  -7624.302 us | -74.68% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |    1000     |     1000     |      10      |         1          | 276.633 us |       5.38% | 200.963 us |       3.79% |    -75.671 us | -27.35% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |     1000     |      10      |         1          | 396.510 us |       2.92% | 228.645 us |       2.91% |   -167.865 us | -42.34% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |     1000     |      10      |         1          |  15.060 ms |       0.18% |   1.861 ms |       0.41% | -13199.122 us | -87.65% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |   100000    |    100000    |      10      |         1          | 403.073 us |       2.19% | 232.424 us |       2.77% |   -170.649 us | -42.34% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |    100000    |      10      |         1          |  15.375 ms |       0.12% |   1.994 ms |       0.39% | -13381.439 us | -87.03% |   FAST   |
|     1      | NULLS_UNEQUAL  |   INT32    |  10000000   |   10000000   |      10      |         1          |  16.784 ms |       0.11% |   2.793 ms |       0.25% | -13990.608 us | -83.36% |   FAST   |

@PointKernel
PointKernel requested a review from lamarrr July 15, 2026 23:25

@shrshi shrshi left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks great overall! Some clarifying questions:

* @param mr Device memory resource used to allocate the returned contribution counts.
*
* @return The exact number of pairs that `filter_join_indices` would produce.
* @return A pair containing the exact number of pairs that `filter_join_indices` would produce

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Two questions here: (i) Can you help me understand why we need the per-output contribution vector, and (ii) Is it binary vector to indicate which rows are present in the filtered output? Should we add the comment in filter_join_indices_output_size_kernel here as well?

@PointKernel PointKernel Jul 16, 2026

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

(i) why we need the per-output contribution vector

Good question! It's only there to keep the existing mixed-join two-pass interface source-compatible:

https://github.com/rapidsai/cudf/blob/9131304f6ba37ceaceb0bf187f3265df6b82bfc7/cpp/include/cudf/join/mixed_join.hpp#L41

The legacy kernel used the per-row counts to compute write offsets in the retrieve pass; this filter-based rewrite doesn't, detail::mixed_join reads only the scalar output_size_data->first:

https://github.com/rapidsai/cudf/blob/f1408960e3fcf95bf0b16270a6a4d79af992e7b5/cpp/src/join/mixed_join.cu#L113-L114

I'll check whether Spark actually reads the vector as their early chunked-probing attempt. If nobody does, I'll file a follow-up to deprecate it and return the total only; if chunked probing is genuinely wanted, the proper fix is the match-context design you introduced, shared across joins, rather than this passthrough:

https://github.com/rapidsai/cudf/blob/9131304f6ba37ceaceb0bf187f3265df6b82bfc7/cpp/include/cudf/join/join.hpp#L77

https://github.com/rapidsai/cudf/blob/9131304f6ba37ceaceb0bf187f3265df6b82bfc7/cpp/include/cudf/join/hash_join.hpp#L275

(ii) Should we add the comment in filter_join_indices_output_size_kernel here as well?

Done, added the per-join-kind layout to this doc.

@mhaseeb123 mhaseeb123 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Approving CMake only.

@PointKernel
PointKernel changed the base branch from main to release/26.08 July 17, 2026 18:18
Comment thread cpp/include/cudf/join/join.hpp
@PointKernel
PointKernel dismissed lamarrr’s stale review July 22, 2026 05:57

dismissing stale request

@PointKernel

Copy link
Copy Markdown
Member Author

/merge

@rapids-bot
rapids-bot Bot merged commit 7762df3 into NVIDIA:release/26.08 Jul 22, 2026
259 of 261 checks passed
@PointKernel
PointKernel deleted the mixed-join-via-filter-indices branch July 22, 2026 05:57
@GregoryKimball GregoryKimball moved this from Burndown to Landed in libcudf Jul 22, 2026
@GregoryKimball GregoryKimball removed this from libcudf Aug 21, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

breaking Breaking change CMake CMake build issue improvement Improvement / enhancement to an existing function libcudf Affects libcudf (C++/CUDA) code.

Projects

Status: Landed

Development

Successfully merging this pull request may close these issues.

8 participants