Add a direct_inner_join API for pre-hashed distinct UINT32 keys

[PointKernel]

Description

Closes #23146

This PR adds a direct_inner_join free function to libcudf, the first step of the perfect hash join effort in #23126. The keys are a single UINT32 column per side, produced by a prior perfect hashing pass such as cudf::key_remapping, dictionary encoding, or dense integer primary keys. The right keys act as a perfect hash of the right rows: a lookup table of caller-specified capacity entries maps each key value to its row index and each left key probes that table directly, so the join performs no hashing or key comparison at all.

• The caller controls the memory footprint via the explicit capacity argument. All key values must be in [0, capacity) and the right keys must be distinct; behavior is undefined otherwise.
• Inner join only, as a free function: no table reuse across probes is needed, so there is no join object.
• The build scatters right row indices into the lookup table with cub::DeviceTransform::Fill + cub::DeviceFor::Bulk; the probe is a single cub-based copy_if pass emitting the matched index pairs.
• A new JOIN_NVBENCH benchmark compares inner_join, distinct_hash_join, and direct_inner_join on identical conforming input; results in the comment below.

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]

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.

[PointKernel]

Benchmark sweep of the new direct_inner_join bench, comparing the three inner join implementations on identical input satisfying direct_inner_join's preconditions: single UINT32 key column per side, distinct right keys in [0, capacity) with no nulls:

## direct_inner_join

### [0] NVIDIA GH200 480GB

|   algorithm   | left_size | right_size | skip_large_sizes | input_bytes | Samples |  CPU Time  | Noise |  GPU Time  | Noise |  Elem/s  | GlobalMem BW | BWUtil |
|---------------|-----------|------------|------------------|-------------|---------|------------|-------|------------|-------|----------|--------------|--------|
|          hash |      1000 |       1000 |                1 |        8000 |   4000x | 129.934 us | 3.35% | 125.068 us | 3.47% |  63.965M |  63.965 MB/s |  0.00% |
| distinct_hash |      1000 |       1000 |                1 |        8000 |   6544x |  81.248 us | 2.37% |  76.442 us | 2.46% | 104.655M | 104.655 MB/s |  0.00% |
|        direct |      1000 |       1000 |                1 |        8000 |  13648x |  41.675 us | 2.92% |  36.669 us | 3.38% | 218.166M | 218.166 MB/s |  0.01% |
|          hash |    100000 |       1000 |                1 |      404000 |   3920x | 132.442 us | 5.82% | 127.595 us | 6.02% |   3.166G |   3.166 GB/s |  0.08% |
| distinct_hash |    100000 |       1000 |                1 |      404000 |   6352x |  83.566 us | 2.35% |  78.756 us | 2.44% |   5.130G |   5.130 GB/s |  0.13% |
|        direct |    100000 |       1000 |                1 |      404000 |  13616x |  41.758 us | 5.98% |  36.754 us | 6.81% |  10.992G |  10.992 GB/s |  0.27% |
|          hash |  10000000 |       1000 |                1 |    40004000 |    544x |   1.601 ms | 0.75% |   1.596 ms | 0.75% |  25.071G |  25.071 GB/s |  0.62% |
| distinct_hash |  10000000 |       1000 |                1 |    40004000 |   1072x | 475.370 us | 0.59% | 470.516 us | 0.59% |  85.022G |  85.022 GB/s |  2.11% |
|        direct |  10000000 |       1000 |                1 |    40004000 |   5440x |  97.045 us | 3.03% |  92.012 us | 3.20% | 434.770G | 434.770 GB/s | 10.81% |
|          hash |    100000 |     100000 |                1 |      800000 |   3664x | 141.707 us | 3.70% | 136.857 us | 3.81% |   5.846G |   5.846 GB/s |  0.15% |
| distinct_hash |    100000 |     100000 |                1 |      800000 |   5552x |  94.999 us | 2.28% |  90.180 us | 2.23% |   8.871G |   8.871 GB/s |  0.22% |
|        direct |    100000 |     100000 |                1 |      800000 |  10784x |  51.421 us | 2.63% |  46.416 us | 2.95% |  17.236G |  17.236 GB/s |  0.43% |
|          hash |  10000000 |     100000 |                1 |    40400000 |    302x |   1.666 ms | 0.33% |   1.661 ms | 0.33% |  24.325G |  24.325 GB/s |  0.60% |
| distinct_hash |  10000000 |     100000 |                1 |    40400000 |    992x | 512.977 us | 0.59% | 508.142 us | 0.59% |  79.505G |  79.505 GB/s |  1.98% |
|        direct |  10000000 |     100000 |                1 |    40400000 |   2704x | 190.432 us | 1.09% | 185.457 us | 1.13% | 217.840G | 217.840 GB/s |  5.42% |
|          hash |  10000000 |   10000000 |                1 |    80000000 |    191x |   2.626 ms | 0.21% |   2.621 ms | 0.21% |  30.527G |  30.527 GB/s |  0.76% |
| distinct_hash |  10000000 |   10000000 |                1 |    80000000 |    356x |   1.412 ms | 0.28% |   1.407 ms | 0.28% |  56.877G |  56.877 GB/s |  1.41% |
|        direct |  10000000 |   10000000 |                1 |    80000000 |   1120x | 452.128 us | 0.64% | 447.107 us | 0.65% | 178.928G | 178.928 GB/s |  4.45% |

[PointKernel]

/ok to test 29e7764

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: cc499613-432a-4120-a1ab-69cdd7ca6b81

📥 Commits

Reviewing files that changed from the base of the PR and between 29e7764 and bae090e.

📒 Files selected for processing (5)

• cpp/CMakeLists.txt
• cpp/benchmarks/CMakeLists.txt
• cpp/src/join/direct_join.cu
• cpp/tests/CMakeLists.txt
• cpp/tests/streams/join_test.cpp

🚧 Files skipped from review as they are similar to previous changes (4)

• cpp/tests/CMakeLists.txt
• cpp/CMakeLists.txt
• cpp/tests/streams/join_test.cpp
• cpp/src/join/direct_join.cu

---

📝 Walkthrough

Summary by CodeRabbit

• New Features
  • Added a new direct inner join API for UINT32 join keys, returning matching left/right row indices.
  • Added a new benchmark to compare join variants under direct inner-join-compatible workloads.
• Tests
  • Added unit tests validating correct matches, empty input behavior, invalid key type handling, null handling, and insufficient capacity errors.
  • Added a stream-level test exercising the new direct inner join path.

Walkthrough

Adds a new cudf::direct_inner_join free function for inner-joining UINT32 key columns through a capacity-sized lookup table. Includes its public declaration, CUDA implementation, benchmark comparisons, correctness tests, and build wiring.

Changes

Direct inner join feature

Layer / File(s)	Summary
Public API declaration cpp/include/cudf/join/direct_join.hpp	Declares cudf::direct_inner_join with key, capacity, exception, stream, and memory-resource contracts.
CUDA implementation cpp/src/join/direct_join.cu, cpp/CMakeLists.txt	Implements lookup-table scatter/probing, input validation, matched-index output, and library build integration.
Correctness tests cpp/tests/join/direct_join_tests.cpp, cpp/tests/CMakeLists.txt, cpp/tests/streams/join_test.cpp	Adds match, empty-input, invalid-type, null-key, capacity, and stream invocation tests.
Benchmark comparison cpp/benchmarks/join/direct_join.cu, cpp/benchmarks/CMakeLists.txt	Adds configurable NVBench comparisons for hash, distinct_hash, and direct inner joins.

Estimated code review effort: 3 (Moderate) | ~30 minutes

Possibly related issues

• cudf issue 23126: Covers the broader perfect-hash join API that this direct UINT32 inner-join implementation contributes toward.

Suggested labels: improvement

Suggested reviewers: igorpeshansky, kingcrimsontianyu

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly summarizes the main change: adding a direct_inner_join API for pre-hashed UINT32 keys.
Description check	✅ Passed	The description is directly related to the change and describes the new API, constraints, tests, and benchmark.
Linked Issues check	✅ Passed	The PR implements the requested free function for single UINT32 keys with distinct right keys, inner join only, and no join object.
Out of Scope Changes check	✅ Passed	The extra benchmark, tests, and CMake updates are all aligned with the stated objectives and do not appear out of scope.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

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

[coderabbitai]

Actionable comments posted: 2

🧹 Nitpick comments (1)

cpp/tests/join/direct_join_tests.cpp (1)

60-79: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Consider adding a sliced-column test case.

Current tests cover dense/sparse/empty/invalid-type/null/capacity cases well, but none exercise a sliced column_view (non-zero offset) for left_keys/right_keys, which the coding guidelines call out as an edge case worth covering for join-style APIs.

🤖 Prompt for 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.

In `@cpp/tests/join/direct_join_tests.cpp` around lines 60 - 79, Add a join test
that uses a sliced column_view with a non-zero offset for either left_keys or
right_keys in DirectJoinTest. Extend the DirectJoinTest coverage alongside
DenseKeys and SparseKeys so compare_to_reference is exercised with a sliced
input, ensuring the join path handles non-zero column offsets correctly.

Source: Coding guidelines

🤖 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/src/join/direct_join.cu`:
- Around line 93-96: The temporary lookup buffer in direct_join.cu is being
allocated with the default RMM resource instead of cuDF’s tracked resource.
Update the lookup allocation in the direct join path to pass
cudf::get_current_device_resource_ref() when constructing the
rmm::device_uvector<size_type> in the block that builds the lookup table,
keeping the rest of the DeviceTransform::Fill logic unchanged.

In `@cpp/tests/streams/join_test.cpp`:
- Around line 90-95: The DirectInnerJoin test is using an out-of-contract key
because direct_inner_join is called with capacity 4 while left_keys includes 5,
which can index past the lookup table. Update JoinTest.DirectInnerJoin so the
keys stay within [0, capacity) or increase the capacity to cover the largest
key, keeping the test data consistent with the direct_inner_join contract.

---

Nitpick comments:
In `@cpp/tests/join/direct_join_tests.cpp`:
- Around line 60-79: Add a join test that uses a sliced column_view with a
non-zero offset for either left_keys or right_keys in DirectJoinTest. Extend the
DirectJoinTest coverage alongside DenseKeys and SparseKeys so
compare_to_reference is exercised with a sliced input, ensuring the join path
handles non-zero column offsets correctly.

🪄 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: 3fad77ec-2416-44d9-9a06-600dfe06b158

📥 Commits

Reviewing files that changed from the base of the PR and between 175fb17 and 29e7764.

📒 Files selected for processing (8)

• cpp/CMakeLists.txt
• cpp/benchmarks/CMakeLists.txt
• cpp/benchmarks/join/direct_join.cu
• cpp/include/cudf/join/direct_join.hpp
• cpp/src/join/direct_join.cu
• cpp/tests/CMakeLists.txt
• cpp/tests/join/direct_join_tests.cpp
• cpp/tests/streams/join_test.cpp

[GregoryKimball]

@shrshi and @kingcrimsontianyu would you please share your review?

[kingcrimsontianyu]

Looks great. Left a minor question.

[shrshi]

This is awesome! Very clean implementation :)

[PointKernel]

/merge

[mhaseeb123]

Approving CMake