Add filter_join_indices_output_size - #22694
Merged
Merged
Conversation
|
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
Bot
added
libcudf
Member
Author
|
/ok to test 046f944 |
PointKernel
added
feature request
PointKernel
changed the title
Member
Author
|
Could use some help and advice on the naming here. The existing For this new API, the goal is to return the number of output elements produced by However, this creates an odd situation where the intent of the API becomes fairly obscure from the name alone. Any suggestions? Updates: switched to |
|
Note Reviews pausedIt 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 Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (3)
🚧 Files skipped from review as they are similar to previous changes (2)
📝 WalkthroughSummary by CodeRabbit
WalkthroughAdds a public API ChangesFilter Join Indices Size API Implementation
Estimated code review effort🎯 3 (Moderate) | ⏱️ ~25 minutes Possibly related PRs
Suggested labels
Suggested reviewers
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches📝 Generate docstrings
🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 5
🤖 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 378-386: The declaration of the non-void, side-effect-free API
filter_join_indices_size should be marked [[nodiscard]]; update the function
prototype for filter_join_indices_size (the overload taking cudf::table_view
left/right, device_span left_indices/right_indices, cudf::ast::expression
predicate, cudf::join_kind join_kind, and optional stream/mr) to include the
[[nodiscard]] attribute so callers are warned if the returned size is ignored.
In `@cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh`:
- Around line 87-90: The INNER_JOIN size path currently requires both_valid
before incrementing thread_local_count in the switch (case
cudf::join_kind::INNER_JOIN), which causes undercounting of preserved
JoinNoMatch pass-through pairs; update the check so that thread_local_count is
incremented when predicate_pass is true (remove the both_valid requirement) so
filter_join_indices_size matches the behavior of filter_join_indices and the
materialized path for INNER_JOIN.
- Around line 132-142: The kernel launch of filter_join_indices_size_kernel
(invoked from launch_filter_size_kernel) lacks an immediate CUDA error check;
add a post-launch check that captures launch failures by calling the appropriate
CUDA error-check helper used in this codebase (e.g., CUDA_TRY / CUDA_CHECK or
checking cudaGetLastError()) right after the <<<...>>> launch and log/propagate
the error (or abort) immediately; if a stream is used, also consider checking
the stream via cudaStreamSynchronize(stream.value()) or cudaGetLastError() on
the stream to ensure shared-memory/launch configuration errors are caught early.
In `@cpp/src/join/filter_join_indices/filter_join_indices.cu`:
- Around line 394-396: The current fixed DEFAULT_JOIN_BLOCK_SIZE can request
more threads or dynamic shared memory than the device supports (config and
shmem_per_block), so before computing detail::grid_1d config(...) and
shmem_per_block use device properties to bound both the block size and dynamic
shared memory: query cudaDeviceProp.maxThreadsPerBlock and .sharedMemPerBlock,
choose block_size = max(1, min(DEFAULT_JOIN_BLOCK_SIZE, maxThreadsPerBlock,
left_indices.size())), compute shmem_per_block = min(parser.shmem_per_thread *
block_size, maxSharedMemPerBlock) and ensure grid_1d is constructed with that
block_size and that the launch uses the capped shmem_per_block; also validate
non-zero threads/blocks and fail gracefully if requirements exceed device
limits.
- Line 398: The temporary device counter is being allocated with the passed-in
memory_resource (mr) instead of the current device resource; change the
allocation of the device_scalar (d_count) to use
cudf::get_current_device_resource_ref() as the memory resource argument so
internal temp workspace uses the current device resource (i.e., update the
device_scalar<std::size_t> d_count(...) call to pass
cudf::get_current_device_resource_ref() and keep the same stream and size).
🪄 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: a598a257-f1bd-4943-bdee-fd2255779779
📒 Files selected for processing (18)
cpp/CMakeLists.txtcpp/include/cudf/join/join.hppcpp/src/join/filter_join_indices/filter_join_indices.cucpp/src/join/filter_join_indices/filter_join_indices_jit.cucpp/src/join/filter_join_indices/filter_join_indices_jit_kernel.cuhcpp/src/join/filter_join_indices/filter_join_indices_kernel.cuhcpp/src/join/filter_join_indices/filter_join_indices_kernel.hppcpp/src/join/filter_join_indices/filter_join_indices_kernel_complex.cucpp/src/join/filter_join_indices/filter_join_indices_kernel_null_complex.cucpp/src/join/filter_join_indices/filter_join_indices_kernel_null_primitive.cucpp/src/join/filter_join_indices/filter_join_indices_kernel_primitive.cucpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuhcpp/src/join/filter_join_indices/filter_join_indices_size_kernel.hppcpp/src/join/filter_join_indices/filter_join_indices_size_kernel_complex.cucpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_complex.cucpp/src/join/filter_join_indices/filter_join_indices_size_kernel_null_primitive.cucpp/src/join/filter_join_indices/filter_join_indices_size_kernel_primitive.cucpp/tests/join/mixed_join_tests.cu
There was a problem hiding this comment.
♻️ Duplicate comments (1)
cpp/src/join/filter_join_indices/filter_join_indices.cu (1)
393-394:⚠️ Potential issue | 🟠 Major | ⚡ Quick winBlock size / dynamic shared memory is not bounded by device limits in the size path.
Unlike
filter_join_indices(Lines 101-107), this size path fixes the block size atDEFAULT_JOIN_BLOCK_SIZEand computesshmem_per_blockwithout clamping tocudaDevAttrMaxSharedMemoryPerBlock. For heavy predicates with largeshmem_per_thread, the launch can request more dynamic shared memory than the device allows, causing invalid-launch failures.Suggested fix mirroring `filter_join_indices`
- 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; + 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(DEFAULT_JOIN_BLOCK_SIZE, shmem_limit_per_block / parser.shmem_per_thread) + : DEFAULT_JOIN_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;As per coding guidelines, "Verify kernel launches have valid grid/block dimensions (non-zero blocks/threads)."
🤖 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/join/filter_join_indices/filter_join_indices.cu` around lines 393 - 394, The size-path launch uses a fixed DEFAULT_JOIN_BLOCK_SIZE and computes shmem_per_block from parser.shmem_per_thread without clamping, so update the logic to mirror the other path: query cudaDevAttrMaxSharedMemoryPerBlock, compute max_threads_allowed = max_shared_mem / max(1, parser.shmem_per_thread), clamp block size to min(DEFAULT_JOIN_BLOCK_SIZE, max_threads_allowed) and recompute shmem_per_block = parser.shmem_per_thread * block_size (ensuring shmem_per_block <= max_shared_mem), and finally validate the block size and shmem_per_block are non-zero/within device limits before launching; reference the variables/configs used here: DEFAULT_JOIN_BLOCK_SIZE, parser.shmem_per_thread, config, and shmem_per_block.
🤖 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.
Duplicate comments:
In `@cpp/src/join/filter_join_indices/filter_join_indices.cu`:
- Around line 393-394: The size-path launch uses a fixed DEFAULT_JOIN_BLOCK_SIZE
and computes shmem_per_block from parser.shmem_per_thread without clamping, so
update the logic to mirror the other path: query
cudaDevAttrMaxSharedMemoryPerBlock, compute max_threads_allowed = max_shared_mem
/ max(1, parser.shmem_per_thread), clamp block size to
min(DEFAULT_JOIN_BLOCK_SIZE, max_threads_allowed) and recompute shmem_per_block
= parser.shmem_per_thread * block_size (ensuring shmem_per_block <=
max_shared_mem), and finally validate the block size and shmem_per_block are
non-zero/within device limits before launching; reference the variables/configs
used here: DEFAULT_JOIN_BLOCK_SIZE, parser.shmem_per_thread, config, and
shmem_per_block.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 3e8a0552-7ee9-4fba-bb11-758c064e8f80
📒 Files selected for processing (5)
cpp/CMakeLists.txtcpp/include/cudf/join/join.hppcpp/src/join/filter_join_indices/filter_join_indices.cucpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuhcpp/tests/join/mixed_join_tests.cu
💤 Files with no reviewable changes (1)
- cpp/CMakeLists.txt
🚧 Files skipped from review as they are similar to previous changes (2)
- cpp/tests/join/mixed_join_tests.cu
- cpp/src/join/filter_join_indices/filter_join_indices_size_kernel.cuh
PointKernel
changed the title
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh (1)
40-48:⚠️ Potential issue | 🟠 Major | 🏗️ Heavy liftMake LEFT_JOIN
left_passing_marksmarking race-free.In
cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh, theLEFT_JOINpath doesleft_passing_marks[left_row_index] = true;for every predicate-passing pair; multiple threads can target the sameleft_row_index, so this plain non-atomic store to a sharedboolelement is a data race. The host then derives the output size bycount_ifoverleft_passing_marksincpp/src/join/filter_join_indices/filter_join_indices.cu(LEFT_JOIN case), so this must be race-free. Use an atomic/set-once representation for the marks (e.g., switch touint8_t/unsigned intand set viaatomicOr(..., 1)/cuda::atomic_ref), rather than writingbooldirectly.🤖 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/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh` around lines 40 - 48, The LEFT_JOIN path in filter_join_indices_output_size_kernel currently does non-atomic writes to left_passing_marks (left_passing_marks[left_row_index] = true), causing data races; change the mark representation to an atomic-safe integer type (e.g., uint8_t or unsigned int) and set it atomically in the kernel (use atomicOr or cuda::atomic_ref to set bit/value to 1) instead of writing a bool; update all callers/allocations that provide left_passing_marks to allocate and zero-initialize the new integer type and update the host-side count logic (the count_if in filter_join_indices.cu) to treat non-zero as passed; keep the symbol names filter_join_indices_output_size_kernel and left_passing_marks so the changes are localized.
🤖 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.
Outside diff comments:
In `@cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuh`:
- Around line 40-48: The LEFT_JOIN path in
filter_join_indices_output_size_kernel currently does non-atomic writes to
left_passing_marks (left_passing_marks[left_row_index] = true), causing data
races; change the mark representation to an atomic-safe integer type (e.g.,
uint8_t or unsigned int) and set it atomically in the kernel (use atomicOr or
cuda::atomic_ref to set bit/value to 1) instead of writing a bool; update all
callers/allocations that provide left_passing_marks to allocate and
zero-initialize the new integer type and update the host-side count logic (the
count_if in filter_join_indices.cu) to treat non-zero as passed; keep the symbol
names filter_join_indices_output_size_kernel and left_passing_marks so the
changes are localized.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 20f32451-0f39-4b30-808a-d09c42e23d7f
📒 Files selected for processing (11)
cpp/CMakeLists.txtcpp/include/cudf/join/join.hppcpp/src/join/filter_join_indices/filter_join_indices.cucpp/src/join/filter_join_indices/filter_join_indices_kernel.cuhcpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.cuhcpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel.hppcpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_complex.cucpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_complex.cucpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_null_primitive.cucpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cucpp/tests/join/mixed_join_tests.cu
✅ Files skipped from review due to trivial changes (1)
- cpp/src/join/filter_join_indices/filter_join_indices_output_size_kernel_primitive.cu
🚧 Files skipped from review as they are similar to previous changes (3)
- cpp/src/join/filter_join_indices/filter_join_indices_kernel.cuh
- cpp/tests/join/mixed_join_tests.cu
- cpp/CMakeLists.txt
bdice
approved these changes
Jun 8, 2026
bdice
left a comment
bdice
left a comment
Contributor
There was a problem hiding this comment.
All seems fine to me.
I had one small thought for future work: we could process fewer bytes and potentially use smaller/faster kernels if we can pre-apply filters that only use one table (left/right).
lamarrr
approved these changes
Jun 8, 2026
lamarrr
reviewed
Jun 8, 2026
vuule
reviewed
Jun 8, 2026
mhaseeb123
approved these changes
Jun 9, 2026
mhaseeb123
approved these changes
Jun 9, 2026
mhaseeb123
left a comment
mhaseeb123
left a comment
Contributor
There was a problem hiding this comment.
Nothing further than @vuule's comments
vuule
approved these changes
Jun 9, 2026
Member
Author
|
/merge |
3 tasks
3 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
Required by #22124
This PR introduces a new API,
filter_join_indices_output_size, which returns the output size offilter_join_indiceswithout materializing the join indices output.Checklist