Reduce distinct_helpers build time - #23322
Conversation
github-actions
Bot
added
libcudf
PointKernel
changed the title
PointKernel
added
non-breaking
|
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 (1)
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesDistinct reduction pipeline
Filtered join source split
Estimated code review effort: 4 (Complex) | ~45 minutes Possibly related PRs
Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches 💡 1🛠️ Fix failing CI checks 💡
🧪 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: 1
🧹 Nitpick comments (3)
cpp/src/stream_compaction/distinct_helpers.cu (1)
17-33: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winMissing direct include for
thrust::uninitialized_fill.Line 29 calls
thrust::uninitialized_fill, but only<thrust/fill.h>is included (line 20);thrust::uninitialized_fillis declared in<thrust/uninitialized_fill.h>. This likely compiles only via a transitive include today, which is fragile against unrelated header changes elsewhere.As per coding guidelines, "include headers directly for every used symbol without unused or incorrectly styled includes."
📝 Proposed fix
`#include` <cuda/functional> `#include` <cuda/iterator> `#include` <cuda/std/iterator> -#include <thrust/fill.h> +#include <thrust/uninitialized_fill.h>🤖 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/stream_compaction/distinct_helpers.cu` around lines 17 - 33, Add the direct <thrust/uninitialized_fill.h> include in the file containing initialize_reduction_results, and remove or retain <thrust/fill.h> based on whether another symbol uses it. Ensure thrust::uninitialized_fill is declared by its own header rather than relying on transitive includes.Source: Coding guidelines
cpp/src/stream_compaction/distinct_helpers.hpp (1)
114-139: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winDoxygen doc block now only covers
reduce_by_row_keep_any.The comment at lines 86-113 was written for a single reduction function and is now only adjacent to
reduce_by_row_keep_any.reduce_by_row_keep_first_last_none(120-126) andreduce_by_row(128-139) are left without their own doc comments, so doxygen output/linting will treat them as undocumented.As per coding guidelines, "Use doxygen as a documentation generator and linter for C++ and CUDA code."
📝 Proposed fix
+/** + * `@copydoc` cudf::detail::reduce_by_row_keep_any + */ template <typename Set> rmm::device_uvector<size_type> reduce_by_row_keep_first_last_none( Set& set, size_type num_rows, duplicate_keep_option keep, rmm::cuda_stream_view stream, rmm::device_async_resource_ref mr); +/** + * `@copydoc` cudf::detail::reduce_by_row_keep_any + */ template <typename Set> rmm::device_uvector<size_type> reduce_by_row(Set& set, ...);🤖 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/stream_compaction/distinct_helpers.hpp` around lines 114 - 139, Add Doxygen documentation for reduce_by_row_keep_first_last_none and reduce_by_row in distinct_helpers.hpp, ensuring each function’s parameters, behavior, and return value are documented. Keep the existing documentation scoped to reduce_by_row_keep_any and align the new comments with the project’s Doxygen conventions.Source: Coding guidelines
cpp/src/stream_compaction/distinct_helpers.cuh (1)
55-74: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winDeclare the device lambda’s return type explicitly.
Proposed fix
- [set_ref, keep, reduction_results = reduction_results.begin()] __device__( - size_type const idx) mutable { + [set_ref, keep, reduction_results = reduction_results.begin()] __device__( + size_type const idx) mutable -> void {As per coding guidelines, “Declare explicit return types for extended device lambdas passed to device algorithms.”
🤖 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/stream_compaction/distinct_helpers.cuh` around lines 55 - 74, Update the device lambda passed to thrust::for_each in the distinct compaction logic to declare an explicit void return type, while preserving its existing capture list, mutability, and body 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.
Inline comments:
In `@cpp/src/stream_compaction/distinct_helpers.cuh`:
- Around line 49-51: Update the reduction_results allocation in the distinct
compaction flow to use cudf::get_current_device_resource_ref() instead of the
caller-provided mr, while keeping output_indices allocated with mr. Leave
initialize_reduction_results and the surrounding stream behavior unchanged.
---
Nitpick comments:
In `@cpp/src/stream_compaction/distinct_helpers.cu`:
- Around line 17-33: Add the direct <thrust/uninitialized_fill.h> include in the
file containing initialize_reduction_results, and remove or retain
<thrust/fill.h> based on whether another symbol uses it. Ensure
thrust::uninitialized_fill is declared by its own header rather than relying on
transitive includes.
In `@cpp/src/stream_compaction/distinct_helpers.cuh`:
- Around line 55-74: Update the device lambda passed to thrust::for_each in the
distinct compaction logic to declare an explicit void return type, while
preserving its existing capture list, mutability, and body behavior.
In `@cpp/src/stream_compaction/distinct_helpers.hpp`:
- Around line 114-139: Add Doxygen documentation for
reduce_by_row_keep_first_last_none and reduce_by_row in distinct_helpers.hpp,
ensuring each function’s parameters, behavior, and return value are documented.
Keep the existing documentation scoped to reduce_by_row_keep_any and align the
new comments with the project’s Doxygen conventions.
🪄 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: fb962803-9e5f-4ba7-8243-3dda1deacc3d
📒 Files selected for processing (15)
cpp/CMakeLists.txtcpp/src/hash/murmurhash3_x86_32.cucpp/src/hash/murmurhash3_x86_32.cuhcpp/src/stream_compaction/distinct.cucpp/src/stream_compaction/distinct_helpers.cucpp/src/stream_compaction/distinct_helpers.cuhcpp/src/stream_compaction/distinct_helpers.hppcpp/src/stream_compaction/distinct_helpers_flat_nan_equal_any.cucpp/src/stream_compaction/distinct_helpers_flat_nan_equal_ordered.cucpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_any.cucpp/src/stream_compaction/distinct_helpers_flat_nan_unequal_ordered.cucpp/src/stream_compaction/distinct_helpers_nested_nan_equal_any.cucpp/src/stream_compaction/distinct_helpers_nested_nan_equal_ordered.cucpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_any.cucpp/src/stream_compaction/distinct_helpers_nested_nan_unequal_ordered.cu
bdice
left a comment
bdice
left a comment
There was a problem hiding this comment.
Nice. All this makes sense.
| thrust::for_each(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), | ||
| cuda::counting_iterator<cudf::size_type>{0}, | ||
| cuda::counting_iterator{num_rows}, | ||
| [set_ref, keep, reduction_results = reduction_results.begin()] __device__( |
There was a problem hiding this comment.
Device lambdas should declare trailing return types, but I am unsure if that also applies when void. We typically discover these requirements when we update compilers or CCCL, not from cuDF source changes...
|
I would be curious to know how the aggregate build time compares to the original single .cu build time. |
|
@davidwendt I reran both sides from empty build directories with identical Release/SM100/CUDA 13.0 settings and all compiler caches disabled (no
I couldn't find the original Slack discussion, but with sccache we effectively have unlimited parallelism, so what really matters is the build time of the longest single TU. That's why similar optimization efforts are all evaluated based on the longest TU build time before and after the change. |
|
One concern is that this doubles the binary size. Reviewers, do you think we should put this on hold or merge it as is? |
I assume this is the .o sizes. How does it affect the libcudf.so size? |
Yes, it's just the TU .o size: TU object size: 2.095 MiB for the original distinct_helpers.cu.o vs. 4.659 MiB summed across the nine split objects (2.224x / +122.4%). I had the agent running all day yesterday. The best result it found without increasing binary size was using the custom dispatcher and custom kernel, which reduced the build time from 662s to around 400s. That said, for build-time-critical TUs, increasing the TU size by 2 MB seems like a reasonable tradeoff given the amount of build time it saves. For non-build-time-critical TUs, though, we should avoid increasing the |
4 participants
Description
Related to #21973
This PR reduces build time by splitting flat and nested row operators, NaN equality comparators, and
KEEP_ANYand ordered operations into separate TUs.For nested
KEEP_ANY, this PR materializes MurmurHash values in a separate pass and uses lightweight hash lookups during cuco insertion. Ordered nested operations continue hashing inline to avoid a runtime regression.The aggregate CUDA compile time is effectively unchanged (662s for the original TU vs 665s summed across the nine split TUs), while the longest individual compile drops from 662s to 103s. Nested LIST
KEEP_ANYbenchmarks improve by 18–20%, with no measurable regressions across 24 cases.Checklist