Reduce filtered join build time

[PointKernel]

Description

Related to #21973

join/filtered_join.cu was the second most expensive libcudf TU in a cold build. This PR reduces its build time by splitting primitive, flat, and nested row operators into separate TUs.

For nested rows, this PR materializes MurmurHash values in a separate pass and uses lightweight hash lookups during CUCO insertion and probing. This trades an extra hash pass and a temporary 32-bit hash per row for keeping recursive nested hashing out of the template-heavy cuco kernels, reducing compile complexity and register pressure. The existing preprocessed row metadata is reused, while primitive and flat paths continue hashing inline.

The longest CUDA compile drops from 432s to 91s. The combined affected object size drops by 15%. Nested LIST and STRUCT benchmarks improve by up to 44%, with no regressions across 32 cases.

Checklist

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

[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: 5ca997ea-52b8-400d-a026-af5b14ad689a

📥 Commits

Reviewing files that changed from the base of the PR and between ea83b98 and f331ad5.

📒 Files selected for processing (2)

• cpp/include/cudf/detail/join/distinct_filtered_join.cuh
• cpp/include/cudf/detail/join/filtered_join.cuh

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

• cpp/include/cudf/detail/join/distinct_filtered_join.cuh
• cpp/include/cudf/detail/join/filtered_join.cuh

---

📝 Walkthrough

Summary by CodeRabbit

• New Features
  • Improved filtered semi-join and anti-join support for primitive, flat, and nested data, including more accurate handling of nested nulls and null-equality behavior.
  • Added support for hashing from preprocessed join inputs to streamline distinct filtered join evaluation.
• Performance
  • Enhanced execution by automatically selecting the most suitable join strategy (primitive, flat, or nested) based on input structure.
  • Improved filtered join hash-table build and probe paths, including better handling of empty inputs.

Walkthrough

Filtered joins are reorganized into split primitive, flat, and nested CUDA paths. The implementation adds preprocessed-table MurmurHash3 support, mode-specific CUCO probing, nested-null validity masks, semi/anti join result handling, and build integration.

Changes

Filtered Join Refactor

Layer / File(s)	Summary
Preprocessed hashing interface cpp/src/hash/murmurhash3_x86_32.*	MurmurHash3 shares an implementation for preprocessed tables and exposes a matching overload.
Join mode contracts and shared kernels cpp/include/cudf/detail/join/*.cuh, cpp/src/join/filtered_join/filtered_join_common.cuh	Join interfaces define row-operator modes, mode-specific helpers, bucket sizing, validity predicates, and shared insertion/query kernels.
Join initialization and dispatch cpp/src/join/filtered_join/filtered_join.cu	Construction selects the row-operator mode, allocates storage, preprocesses the right table, and dispatches insertion.
Mode-specific probing paths cpp/src/join/filtered_join/filtered_join_{primitive,flat,nested,nested_query}.cu	Separate CUDA units implement CUCO insertion and membership queries for primitive, flat, and nested representations.
Semi/anti results and build wiring cpp/src/join/filtered_join/filtered_join.cu, cpp/CMakeLists.txt	Semi and anti joins handle empty inputs, filter membership results, and compile the split sources into the cudf target.

Estimated code review effort: 4 (Complex) | ~45 minutes

Possibly related PRs

• rapidsai/cudf#23322: Refactors the same MurmurHash3 implementation around preprocessed-table hashing.
• rapidsai/cudf#23379: Updates the filtered-join load-factor to bucket-storage path.

Suggested labels: libcudf, CMake, improvement, non-breaking, Performance

Suggested reviewers: ttnghia, kingcrimsontianyu, mhaseeb123

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Docstring Coverage	⚠️ Warning	Docstring coverage is 50.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
Title check	✅ Passed	The title clearly summarizes the main change: reducing filtered join build time.
Description check	✅ Passed	The description directly explains the filtered join TU split and the nested hash refactor.
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.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

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

[coderabbitai]

🧹 Nitpick comments (1)

cpp/src/hash/murmurhash3_x86_32.cu (1)

45-48: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

Declare the device lambda’s return type.

Line 47 passes an extended device lambda to a device algorithm without an explicit return type. As per coding guidelines, “Declare explicit return types for extended device lambdas passed to device algorithms.”

Proposed fix

-    [output_begin, hasher] __device__(size_type i) mutable { output_begin[i] = hasher(i); },
+    [output_begin, hasher] __device__(size_type i) mutable -> void { output_begin[i] = hasher(i); },

🤖 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/src/hash/murmurhash3_x86_32.cu` around lines 45 - 48, Update the extended
device lambda passed to cub::DeviceFor::Bulk in the MurmurHash3 processing call
to declare its return type explicitly, while preserving its existing mutable
capture and output assignment behavior.

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.

Nitpick comments:
In `@cpp/src/hash/murmurhash3_x86_32.cu`:
- Around line 45-48: Update the extended device lambda passed to
cub::DeviceFor::Bulk in the MurmurHash3 processing call to declare its return
type explicitly, while preserving its existing mutable capture and output
assignment behavior.

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 33fb40a4-83b6-46d4-bbeb-50718e2e0f95

📥 Commits

Reviewing files that changed from the base of the PR and between 16f187b and ea83b98.

📒 Files selected for processing (12)

• cpp/CMakeLists.txt
• cpp/include/cudf/detail/join/distinct_filtered_join.cuh
• cpp/include/cudf/detail/join/filtered_join.cuh
• cpp/src/hash/murmurhash3_x86_32.cu
• cpp/src/hash/murmurhash3_x86_32.cuh
• cpp/src/join/filtered_join.cu
• cpp/src/join/filtered_join/filtered_join.cu
• cpp/src/join/filtered_join/filtered_join_common.cuh
• cpp/src/join/filtered_join/filtered_join_flat.cu
• cpp/src/join/filtered_join/filtered_join_nested.cu
• cpp/src/join/filtered_join/filtered_join_nested_query.cu
• cpp/src/join/filtered_join/filtered_join_primitive.cu

💤 Files with no reviewable changes (1)

• cpp/src/join/filtered_join.cu

[mhaseeb123]

🚀

[kingcrimsontianyu]

Looks good!

[PointKernel]

/merge