Split random forest histogram building and split scoring kernels - #8370
Merged
rapids-bot[bot] merged 11 commits intoJul 14, 2026
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. |
RAMitchell
added
improvement
RAMitchell
marked this pull request as ready for review
July 13, 2026 11:10
ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (17)
🚧 Files skipped from review as they are similar to previous changes (11)
📝 WalkthroughSummary by CodeRabbit
WalkthroughThe batched decision-tree builder removes ChangesBatched split pipeline
Estimated code review effort: 4 (Complex) | ~45 minutes Suggested reviewers: 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 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
🤖 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/decisiontree/batched-levelalgo/builder.cuh`:
- Around line 550-551: Guard the dim3 launch dimensions in the histogram_grid
and split_grid initialization by converting each size_t/int dimension through
ML::narrow_cast<ML::cuda_launch_t> before constructing dim3. Preserve the
existing grid dimensions while ensuring oversized values fail on the host
instead of truncating to unsigned int.
🪄 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: 61575d02-2dbc-4264-a9fd-a31c1fdb2922
📒 Files selected for processing (11)
cpp/src/decisiontree/batched-levelalgo/builder.cuhcpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuhcpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuhcpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cucpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cucpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cucpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu
There was a problem hiding this comment.
Pull request overview
This PR refactors the batched-level random forest split computation into two separate CUDA kernel launches: one for histogram construction and another for CDF conversion + split scoring. This separation creates a clean synchronization point needed for upcoming distributed histogram all-reduce integration while keeping the existing shared-memory histogram fast path.
Changes:
- Split the previous monolithic split kernel into
buildHistogramsKernelandfindBestSplitsKernel, and update the host launcher API accordingly. - Remove the inter-block completion handshake (
done_count) and rework workload bookkeeping to support the two-kernel flow. - Update all explicit template instantiations and shared-memory configuration plumbing to match the new launcher signature.
Reviewed changes
Copilot reviewed 11 out of 11 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh | Implements the new two-kernel split flow and updates histogram/CDF handling. |
| cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh | Updates public kernel launcher declarations and introduces SharedMemoryConfig. |
| cpp/src/decisiontree/batched-levelalgo/builder.cuh | Removes done_count, adjusts workspace sizing, workload mapping, and launches the split kernels with separate grids. |
| cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu | Updates explicit instantiation to launchComputeSplitKernels and new parameters. |
| cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu | Updates explicit instantiation to launchComputeSplitKernels and new parameters. |
| cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu | Updates explicit instantiation to launchComputeSplitKernels and new parameters. |
| cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu | Updates explicit instantiation to launchComputeSplitKernels and new parameters. |
| cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu | Updates explicit instantiation to launchComputeSplitKernels and new parameters. |
| cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu | Updates explicit instantiation to launchComputeSplitKernels and new parameters. |
| cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu | Updates explicit instantiation to launchComputeSplitKernels and new parameters. |
| cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu | Updates explicit instantiation to launchComputeSplitKernels and new parameters. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Contributor
Author
|
chyunsu3
approved these changes
Jul 14, 2026
chyunsu3
left a comment
chyunsu3
left a comment
Contributor
There was a problem hiding this comment.
LGTM. Nice job cleaning up the code.
Comment on lines
-343
to
-348
| __threadfence(); // for commit guarantee before the last block scores the split | ||
| if (!use_global_memory_histogram) { | ||
| __syncthreads(); | ||
|
|
||
| bool last = MLCommon::signalDone( | ||
| done_count + nid * gridDim.y + blockIdx.y, num_blocks, offset_blockid == 0, shared_done); | ||
| if (!last) return; |
Contributor
There was a problem hiding this comment.
Signaling a completion is no longer needed because we now have two kernels, with a natural point for synchronization in between. Is my understanding correct?
Contributor
Author
There was a problem hiding this comment.
Yes the kernel boundary is now the sync.
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 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 `@python/cuml/cuml/decomposition/pca.pyx`:
- Around line 482-483: Remove the ensure_min_features=2 guard from the input
validation in python/cuml/cuml/decomposition/pca.pyx lines 482-483,
python/cuml/cuml/decomposition/tsvd.pyx lines 316-317, and
python/cuml/cuml/manifold/t_sne.pyx lines 598-599 so valid one-feature inputs
are accepted; if retaining the restriction instead, document it and add coverage
for the enforced behavior at all three sites.
🪄 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: d17cb377-34b1-4f23-809a-cf7a29f0bbab
📒 Files selected for processing (17)
cpp/src/decisiontree/batched-levelalgo/builder.cuhcpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuhcpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuhcpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cucpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cucpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cucpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cupython/cuml/cuml/_thirdparty/sklearn/preprocessing/_data.pypython/cuml/cuml/decomposition/pca.pyxpython/cuml/cuml/decomposition/tsvd.pyxpython/cuml/cuml/manifold/t_sne.pyxpython/cuml/cuml/neighbors/kneighbors_regressor.pyxpython/cuml/tests/test_sklearn_compatibility.py
🚧 Files skipped from review as they are similar to previous changes (11)
- cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh
- cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh
- cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu
- cpp/src/decisiontree/batched-levelalgo/builder.cuh
There was a problem hiding this comment.
Caution
Inline review comments failed to post. This is likely due to GitHub's internal server error or limits when posting large numbers of comments. If you are seeing this consistently it is likely a permissions issue. Please check "Moderation" -> "Code review limits" under your organization settings.
Actionable comments posted: 1
🤖 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 `@python/cuml/cuml/decomposition/pca.pyx`:
- Around line 482-483: Remove the ensure_min_features=2 guard from the input
validation in python/cuml/cuml/decomposition/pca.pyx lines 482-483,
python/cuml/cuml/decomposition/tsvd.pyx lines 316-317, and
python/cuml/cuml/manifold/t_sne.pyx lines 598-599 so valid one-feature inputs
are accepted; if retaining the restriction instead, document it and add coverage
for the enforced behavior at all three sites.
🪄 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: d17cb377-34b1-4f23-809a-cf7a29f0bbab
📒 Files selected for processing (17)
cpp/src/decisiontree/batched-levelalgo/builder.cuhcpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuhcpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuhcpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cucpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cucpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cucpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cucpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cupython/cuml/cuml/_thirdparty/sklearn/preprocessing/_data.pypython/cuml/cuml/decomposition/pca.pyxpython/cuml/cuml/decomposition/tsvd.pyxpython/cuml/cuml/manifold/t_sne.pyxpython/cuml/cuml/neighbors/kneighbors_regressor.pyxpython/cuml/tests/test_sklearn_compatibility.py
🚧 Files skipped from review as they are similar to previous changes (11)
- cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh
- cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh
- cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu
- cpp/src/decisiontree/batched-levelalgo/builder.cuh
🛑 Comments failed to post (1)
python/cuml/cuml/decomposition/pca.pyx (1)
482-483: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🧩 Analysis chain
🏁 Script executed:
#!/bin/bash set -euo pipefail # Inspect the relevant source ranges and search for related validation/tests/docs. for f in \ python/cuml/cuml/decomposition/pca.pyx \ python/cuml/cuml/decomposition/tsvd.pyx \ python/cuml/cuml/manifold/t_sne.pyx do echo "===== $f =====" wc -l "$f" done echo "===== PCA excerpt =====" sed -n '450,510p' python/cuml/cuml/decomposition/pca.pyx echo "===== TSVD excerpt =====" sed -n '280,340p' python/cuml/cuml/decomposition/tsvd.pyx echo "===== t-SNE excerpt =====" sed -n '560,630p' python/cuml/cuml/manifold/t_sne.pyx echo "===== Search for ensure_min_features / one-feature handling =====" rg -n "ensure_min_features\s*=\s*2|ensure_min_features|min_features|n_features == 1|one feature|1-feature|single feature" python/cuml/cuml -g '!**/build/**' -g '!**/dist/**' echo "===== Search for docs/tests mentioning feature minimums =====" rg -n "PCA|TruncatedSVD|TSVD|t-SNE|t_sne|ensure_min_features|one feature|single feature|n_features" python -g '!**/build/**' -g '!**/dist/**'Repository: rapidsai/cuml
Length of output: 50370
🏁 Script executed:
#!/bin/bash set -euo pipefail # Show only the exact validation blocks and nearby comments. nl -ba python/cuml/cuml/decomposition/pca.pyx | sed -n '470,490p' echo nl -ba python/cuml/cuml/decomposition/tsvd.pyx | sed -n '305,325p' echo nl -ba python/cuml/cuml/manifold/t_sne.pyx | sed -n '588,608p' echo # Look for any estimator docs that state the minimum number of features. rg -n "minimum.*feature|at least 2 features|ensure_min_features|n_features" python/cuml/cuml/decomposition python/cuml/cuml/manifold -g '!**/build/**' -g '!**/dist/**'Repository: rapidsai/cuml
Length of output: 191
🏁 Script executed:
#!/bin/bash set -euo pipefail echo "===== Relevant validation docs =====" sed -n '100,165p' python/cuml/cuml/internals/validation.py echo "===== PCA/TSVD/t-SNE docstrings around fit =====" sed -n '450,500p' python/cuml/cuml/decomposition/pca.pyx echo sed -n '300,330p' python/cuml/cuml/decomposition/tsvd.pyx echo sed -n '585,610p' python/cuml/cuml/manifold/t_sne.pyx echo "===== Tests mentioning single-feature inputs or feature minima =====" rg -n "single feature|one feature|n_features=1|n_features == 1|minimum of 2|ensure_min_features=2" python/cuml/cuml_accel_tests python/cuml/tests python/cuml/cuml -g '!**/build/**' -g '!**/dist/**' | head -n 200Repository: rapidsai/cuml
Length of output: 17565
Remove or document the 2-feature minimum
ensure_min_features=2blocks valid 1-feature inputs inpython/cuml/cuml/decomposition/pca.pyx#L482-L483,python/cuml/cuml/decomposition/tsvd.pyx#L316-L317, andpython/cuml/cuml/manifold/t_sne.pyx#L598-L599. If that restriction is intentional, add docs and tests; otherwise drop the guard.📍 Affects 3 files
python/cuml/cuml/decomposition/pca.pyx#L482-L483(this comment)python/cuml/cuml/decomposition/tsvd.pyx#L316-L317python/cuml/cuml/manifold/t_sne.pyx#L598-L599🤖 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 `@python/cuml/cuml/decomposition/pca.pyx` around lines 482 - 483, Remove the ensure_min_features=2 guard from the input validation in python/cuml/cuml/decomposition/pca.pyx lines 482-483, python/cuml/cuml/decomposition/tsvd.pyx lines 316-317, and python/cuml/cuml/manifold/t_sne.pyx lines 598-599 so valid one-feature inputs are accepted; if retaining the restriction instead, document it and add coverage for the enforced behavior at all three sites.
…rve-shmem' into codex/enh-rf-split-kernels-preserve-shmem
Contributor
Author
|
/merge |
This was referenced Jul 14, 2026
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
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
Splits random forest split computation into two kernel launches:
buildHistogramsKernel: builds per-node/per-feature histogramsfindBestSplitsKernel: converts histograms to CDFs and scores split candidatesThis is preparation for distributed histogram construction. In the distributed path, each rank will build partial histograms locally, then an all-reduce must take place before split scoring can run. Splitting histogram construction and split scoring gives that all-reduce a natural synchronization point between the two kernels.
The histogram kernel preserves the existing shared-memory fast path for histogram accumulation and quantile lookup. The split-scoring kernel no longer uses dynamic shared memory for quantiles, since it only indexes quantile thresholds by bin rather than searching them per row.
Notes
Testing
SG_RF_TESTRFEquivalentSplitRangeTest.*:RfTest.EquivalentSplitRangePersistsThroughBuilder