Simplify RF lower-bound and index types - #8383
Conversation
…ower-bound-cleanup
…ower-bound-cleanup # Conflicts: # cpp/src/randomforest/randomforest.cuh
…ower-bound-cleanup # Conflicts: # cpp/src/decisiontree/batched-levelalgo/builder.cuh # cpp/src/decisiontree/batched-levelalgo/dataset.h # cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh # cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh # cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu # cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu # cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu # cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu # cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu # cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu # cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu # cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu # cpp/src/decisiontree/batched-levelalgo/objectives.cuh # cpp/src/decisiontree/batched-levelalgo/split.cuh # cpp/src/decisiontree/decisiontree.cuh # cpp/tests/sg/rf_test.cu
|
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
There was a problem hiding this comment.
Pull request overview
This PR refactors the C++ batched-level random forest / decision tree training path to simplify internal index typing (moving many dataset- and node-facing indices to std::int64_t) and to replace a custom quantile-bin lower-bound routine with cuda::std::lower_bound (plus an inline clamp to the last bin).
Changes:
- Replaces the custom quantile-bin lower-bound helper with a
cuda::std::lower_boundover a 32-bitcuda::counting_iteratorbin domain and clamps “above last quantile” to the last bin. - Simplifies/removes
IdxTtemplate plumbing across batched-levelalgo components (Split,Dataset,Quantiles, objectives, kernels) and standardizes many indices tostd::int64_t. - Updates RF/DT tests and explicit kernel instantiations to match the new types, and removes the old lower-bound test.
Reviewed changes
Copilot reviewed 22 out of 22 changed files in this pull request and generated 4 comments.
Show a summary per file
| File | Description |
|---|---|
| cpp/tests/sg/rf_test.cu | Updates RF unit tests for new Split/objective signatures and removes the lower-bound test. |
| cpp/src/randomforest/randomforest.cuh | Moves sampled row-id buffers to std::int64_t and updates sampling paths accordingly. |
| cpp/src/decisiontree/decisiontree.cuh | Adjusts internal DT training interface to use std::int64_t row/col and updates objective dispatch types. |
| cpp/src/decisiontree/batched-levelalgo/split.cuh | Removes IdxT from Split, hard-codes split/node indices to std::int64_t, and updates helpers. |
| cpp/src/decisiontree/batched-levelalgo/quantiles.h | Removes IdxT from Quantiles and makes n_bins_array an int*. |
| cpp/src/decisiontree/batched-levelalgo/quantiles.cuh | Updates quantile result view to the new Quantiles<T> type and minor iterator cleanup. |
| cpp/src/decisiontree/batched-levelalgo/objectives.cuh | Removes IdxT template parameter from objectives and shifts key indices to std::int64_t. |
| cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu | Updates explicit instantiations to new template signatures/types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu | Updates explicit instantiations to new template signatures/types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu | Updates explicit instantiations to new template signatures/types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu | Updates explicit instantiations to new template signatures/types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu | Updates explicit instantiations to new template signatures/types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu | Updates explicit instantiations to new template signatures/types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/node-split.cu | Updates explicit instantiations for node-splitting to new dataset/split/workload types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu | Updates explicit instantiations to new template signatures/types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu | Updates explicit instantiations to new template signatures/types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh | Removes the custom lower_bound helper and updates kernel declarations to new types. |
| cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh | Implements bin lookup via cuda::std::lower_bound + clamp and updates kernels to new types. |
| cpp/src/decisiontree/batched-levelalgo/dataset.h | Removes IdxT from Dataset and hard-codes row/col indices to std::int64_t. |
| cpp/src/decisiontree/batched-levelalgo/builder.cuh | Propagates new index types through builder state/workspace and compute-split launch plumbing. |
| cpp/include/cuml/tree/flatnode.h | Changes SparseTreeNode index fields/accessors to std::int64_t and removes the IdxT template parameter. |
| cpp/include/cuml/tree/decisiontree.hpp | Adds <cstdint> include (supports new std::int64_t usage in public headers). |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
|
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:
📝 WalkthroughWalkthroughThe batched-level decision-tree algorithm removes configurable index-type template parameters and standardizes row, column, node, split, and workload metadata on ChangesIndex type consolidation
Estimated code review effort: 5 (Critical) | ~90 minutes Possibly related PRs
Suggested reviewers: 🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 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: 4
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cpp/src/decisiontree/batched-levelalgo/builder.cuh (1)
209-242: 🩺 Stability & Availability | 🟠 Major | 🏗️ Heavy liftFinish the 64-bit migration with checked host arithmetic.
The widened dimensions still pass through unsafe calculations:
- Line 235 narrows
row_ids->size()fromsize_tusingstatic_cast.- Line 242 and Lines 381-387 narrow and accumulate block counts in
int, potentially underallocatingWorkloadInfo.- Workspace byte counts and offsets use unchecked products/additions.
- Lines 402-422 use overflow-prone ceil-division and
size_t-to-int64_tsampling-offset arithmetic.Use
ML::checked_mul/add/sub/divandML::narrow_cast, with explicit dimension bounds before allocating or launching. As per path instructions, host-side size, offset, and launch arithmetic must use checked helpers, and narrowed counts requireML::narrow_cast.Also applies to: 284-350, 376-422
🤖 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/decisiontree/batched-levelalgo/builder.cuh` around lines 209 - 242, Complete the 64-bit migration in Builder by replacing unchecked size, offset, block-count, workspace, and ceil-division arithmetic across the constructor and code around dataset setup and sampling with ML::checked_mul/add/sub/div. Validate dimensions and launch counts before allocation or kernel launch, convert row_ids->size() and other narrowed values only through ML::narrow_cast, and keep block accumulation in a checked wide type before narrowing for WorkloadInfo.Source: Path instructions
🤖 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/kernels/builder_kernels.cuh`:
- Around line 101-104: Update the node ID handling near fnv1a32_hash so the full
size_t value from work_items[node_idx].idx is preserved when deriving rng_seed;
widen the hash input/API as needed, or use ML::narrow_cast<uint32_t> only with
an explicit validated 32-bit limit. Do not silently assign NodeWorkItem::idx to
uint32_t.
- Around line 83-93: Update sample_features to validate that k is non-negative
before converting it to an unsigned size, then use checked conversion for k and
ML::checked_mul for work_items_size multiplied by the converted k. Preserve
n_column_samples as a safely validated count before passing it to any Thrust or
CUDA operations.
In `@cpp/tests/sg/rf_test.cu`:
- Line 2586: Update the sampled-column allocation and host indexing around the
device buffer declaration and the corresponding lines 2641-2651 to compute a
single checked std::size_t element count from params.n_nodes and params.k, then
reuse it for device and host buffer sizing. Replace the unchecked host offset
arithmetic with the project’s checked arithmetic helper, preserving the existing
indexing behavior.
- Around line 2622-2623: Update the CUDA launch in
excess_sample_with_replacement_kernel to pass params.n_nodes through
ML::narrow_cast<ML::cuda_launch_t> before using it as the grid dimension,
preserving the existing BLOCK_THREADS and stream configuration.
---
Outside diff comments:
In `@cpp/src/decisiontree/batched-levelalgo/builder.cuh`:
- Around line 209-242: Complete the 64-bit migration in Builder by replacing
unchecked size, offset, block-count, workspace, and ceil-division arithmetic
across the constructor and code around dataset setup and sampling with
ML::checked_mul/add/sub/div. Validate dimensions and launch counts before
allocation or kernel launch, convert row_ids->size() and other narrowed values
only through ML::narrow_cast, and keep block accumulation in a checked wide type
before narrowing for WorkloadInfo.
🪄 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: 31d001b6-b593-47c2-a6b5-71d799d00a6d
📒 Files selected for processing (22)
cpp/include/cuml/tree/decisiontree.hppcpp/include/cuml/tree/flatnode.hcpp/src/decisiontree/batched-levelalgo/builder.cuhcpp/src/decisiontree/batched-levelalgo/dataset.hcpp/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/node-split.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.cucpp/src/decisiontree/batched-levelalgo/objectives.cuhcpp/src/decisiontree/batched-levelalgo/quantiles.cuhcpp/src/decisiontree/batched-levelalgo/quantiles.hcpp/src/decisiontree/batched-levelalgo/split.cuhcpp/src/decisiontree/decisiontree.cuhcpp/src/randomforest/randomforest.cuhcpp/tests/sg/rf_test.cu
| excess_sample_with_replacement_kernel<std::int64_t, MAX_SAMPLES_PER_THREAD, BLOCK_THREADS> | ||
| <<<params.n_nodes, BLOCK_THREADS, 0, stream>>>(d_colids.data(), |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
Narrow CUDA launch dimensions explicitly.
params.n_nodes is implicitly converted to CUDA’s unsigned launch type. Use ML::narrow_cast<ML::cuda_launch_t> so invalid or oversized test parameters fail predictably.
Proposed fix
+ const auto grid =
+ ML::narrow_cast<ML::cuda_launch_t>(params.n_nodes);
+ const auto block =
+ ML::narrow_cast<ML::cuda_launch_t>(BLOCK_THREADS);
excess_sample_with_replacement_kernel<std::int64_t, MAX_SAMPLES_PER_THREAD, BLOCK_THREADS>
- <<<params.n_nodes, BLOCK_THREADS, 0, stream>>>(d_colids.data(),
+ <<<grid, block, 0, stream>>>(d_colids.data(),As per coding guidelines, values used in CUDA launch configurations must use ML::narrow_cast<ML::cuda_launch_t>.
📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| excess_sample_with_replacement_kernel<std::int64_t, MAX_SAMPLES_PER_THREAD, BLOCK_THREADS> | |
| <<<params.n_nodes, BLOCK_THREADS, 0, stream>>>(d_colids.data(), | |
| const auto grid = | |
| ML::narrow_cast<ML::cuda_launch_t>(params.n_nodes); | |
| const auto block = | |
| ML::narrow_cast<ML::cuda_launch_t>(BLOCK_THREADS); | |
| excess_sample_with_replacement_kernel<std::int64_t, MAX_SAMPLES_PER_THREAD, BLOCK_THREADS> | |
| <<<grid, block, 0, stream>>>(d_colids.data(), |
🤖 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/sg/rf_test.cu` around lines 2622 - 2623, Update the CUDA launch in
excess_sample_with_replacement_kernel to pass params.n_nodes through
ML::narrow_cast<ML::cuda_launch_t> before using it as the grid dimension,
preserving the existing BLOCK_THREADS and stream configuration.
Source: Coding guidelines
…ower-bound-cleanup
|
Not sure whats up with CI here. |
It is possible that the changes in this PR result in the |
…ower-bound-cleanup # Conflicts: # cpp/src/decisiontree/batched-levelalgo/builder.cuh # cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh # cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh # cpp/src/decisiontree/batched-levelalgo/kernels/classification-double.cu # cpp/src/decisiontree/batched-levelalgo/kernels/classification-float.cu # cpp/src/decisiontree/batched-levelalgo/kernels/node-split.cu # cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu # cpp/src/decisiontree/batched-levelalgo/kernels/regression-float.cu # cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-double.cu # cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu # cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-double.cu # cpp/src/decisiontree/batched-levelalgo/kernels/weighted-regression-float.cu # cpp/src/decisiontree/batched-levelalgo/objectives.cuh # cpp/src/decisiontree/batched-levelalgo/split.cuh # cpp/tests/sg/rf_test.cu
…ower-bound-cleanup
There was a problem hiding this comment.
Actionable comments posted: 1
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh (1)
150-163: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick winUse checked arithmetic for the packed operation count.
Line 163 can wrap
len * reduction_buffer_size_v<BinT>, causing incomplete or out-of-bounds packing.Proposed fix
+ const auto packed_len = + ML::checked_mul<std::size_t>(len, reduction_buffer_size_v<BinT>); raft::linalg::writeOnlyUnaryOp<double, decltype(op), std::size_t, 256>( - out, len * reduction_buffer_size_v<BinT>, op, stream); + out, packed_len, op, stream);As per path instructions, host-side size arithmetic passed as a
size_tcount requires checked arithmetic.🤖 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/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh` around lines 150 - 163, Update packHistograms to compute len * reduction_buffer_size_v<BinT> with checked size_t arithmetic before passing it to raft::linalg::writeOnlyUnaryOp; preserve the existing operation and ensure overflow is detected rather than wrapping into an incomplete or unsafe count.Sources: Coding guidelines, Path instructions
🤖 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/kernels/builder_kernels_impl.cuh`:
- Around line 210-212: Update the nodeSplitCopyBackKernel launch in the
copy-back section to narrow n_blocks_dimx explicitly with
ML::narrow_cast<ML::cuda_launch_t>, matching the checked CUDA launches
immediately above. Leave the kernel arguments and other launch dimensions
unchanged.
---
Outside diff comments:
In `@cpp/src/decisiontree/batched-levelalgo/kernels/builder_kernels.cuh`:
- Around line 150-163: Update packHistograms to compute len *
reduction_buffer_size_v<BinT> with checked size_t arithmetic before passing it
to raft::linalg::writeOnlyUnaryOp; preserve the existing operation and ensure
overflow is detected rather than wrapping into an incomplete or unsafe count.
🪄 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: 7d8670d2-9980-4a7c-9c71-344c8f324eff
📒 Files selected for processing (15)
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/node-split.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.cucpp/src/decisiontree/batched-levelalgo/objectives.cuhcpp/src/decisiontree/batched-levelalgo/split.cuhcpp/tests/sg/rf_test.cu
🚧 Files skipped from review as they are similar to previous changes (7)
- cpp/src/decisiontree/batched-levelalgo/kernels/regression-double.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/weighted-classification-float.cu
- cpp/src/decisiontree/batched-levelalgo/kernels/node-split.cu
- cpp/src/decisiontree/batched-levelalgo/split.cuh
- cpp/src/decisiontree/batched-levelalgo/objectives.cuh
- cpp/src/decisiontree/batched-levelalgo/builder.cuh
- cpp/tests/sg/rf_test.cu
| // The original row_ids buffer remains the source during the scan, so copy back after it finishes. | ||
| nodeSplitCopyBackKernel<DataT, LabelT, IdxT, TPB><<<n_blocks_dimx, TPB, 0, builder_stream>>>( | ||
| nodeSplitCopyBackKernel<DataT, LabelT, TPB><<<n_blocks_dimx, TPB, 0, builder_stream>>>( | ||
| dataset, work_items, splits, workload_info, partition_row_ids); |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
Narrow the copy-back grid dimension explicitly.
Line 211 implicitly converts size_t n_blocks_dimx to CUDA’s launch type, unlike the preceding checked launches.
Proposed fix
- nodeSplitCopyBackKernel<DataT, LabelT, TPB><<<n_blocks_dimx, TPB, 0, builder_stream>>>(
+ nodeSplitCopyBackKernel<DataT, LabelT, TPB>
+ <<<ML::narrow_cast<ML::cuda_launch_t>(n_blocks_dimx), TPB, 0, builder_stream>>>(
dataset, work_items, splits, workload_info, partition_row_ids);As per path instructions, values passed directly as CUDA grid dimensions require ML::narrow_cast<ML::cuda_launch_t>.
📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| // The original row_ids buffer remains the source during the scan, so copy back after it finishes. | |
| nodeSplitCopyBackKernel<DataT, LabelT, IdxT, TPB><<<n_blocks_dimx, TPB, 0, builder_stream>>>( | |
| nodeSplitCopyBackKernel<DataT, LabelT, TPB><<<n_blocks_dimx, TPB, 0, builder_stream>>>( | |
| dataset, work_items, splits, workload_info, partition_row_ids); | |
| // The original row_ids buffer remains the source during the scan, so copy back after it finishes. | |
| nodeSplitCopyBackKernel<DataT, LabelT, TPB> | |
| <<<ML::narrow_cast<ML::cuda_launch_t>(n_blocks_dimx), TPB, 0, builder_stream>>>( | |
| dataset, work_items, splits, workload_info, partition_row_ids); |
🤖 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/decisiontree/batched-levelalgo/kernels/builder_kernels_impl.cuh`
around lines 210 - 212, Update the nodeSplitCopyBackKernel launch in the
copy-back section to narrow n_blocks_dimx explicitly with
ML::narrow_cast<ML::cuda_launch_t>, matching the checked CUDA launches
immediately above. Leave the kernel arguments and other launch dimensions
unchanged.
Sources: Coding guidelines, Path instructions
viclafargue
left a comment
viclafargue
left a comment
There was a problem hiding this comment.
LGTM. Just 3 minor concerns.
| } else if (bootstrap_) { | ||
| // Draw bootstrap rows uniformly when there are no sample weights. | ||
| raft::random::uniformInt<int>( | ||
| raft::random::uniformInt<std::int64_t>( |
There was a problem hiding this comment.
Going from uniformInt<int> to uniformInt<std::int64_t> will probably cause the same seed to produce different results making RF trees trained with earlier RAPIDS version non-reproducible.
There was a problem hiding this comment.
Thats fine but good to note.
| const uint32_t nodeid = work_items[node_idx].idx; | ||
| uint32_t rng_seed = fnv1a32_hash(seed, treeid, nodeid); | ||
| auto nodeid = work_items[node_idx].idx; | ||
| uint32_t rng_seed = fnv1a32_hash(seed, treeid, nodeid); |
There was a problem hiding this comment.
Similar issue over here with treeid and nodeid both now being on 64 bits.
| double sample_weight_sum_; | ||
| rmm::device_uvector<double> sample_weight_cdf_; | ||
| std::deque<rmm::device_uvector<int>> selected_rows_; | ||
| std::deque<rmm::device_uvector<std::int64_t>> selected_rows_; |
There was a problem hiding this comment.
Widening the stored IDs to int64_t does not enable larger inputs as long as n_rows_ and n_sampled_rows_ are stored as int. We should maybe store the IDs as int and while retaining 64 bits format to perform the necessary operations (counts, strides, products, and pointer offsets)?
There was a problem hiding this comment.
I will make n_rows and n_sampled_rows_ also int64_t. My plan is to make most everything 64 bit ints and then evaluate from there if anything needs optimising. 64 bit ints should be the default.
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/randomforest/randomforest.cuh (1)
149-158: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick winUse a checked conversion before resizing
selected_rows.The new
std::int64_tcounting iterator makesn_selecteda signed 64-bit count, but it is passed directly toselected_rows.resize(...)at Line 158. ApplyML::narrow_cast<std::size_t>(n_selected)after the positivity check.As per path instructions, host-side counts passed to
device_uvector::resizerequire checked narrowing.Proposed fix
auto n_selected = selected_rows_end - selected_rows.begin(); ASSERT(n_selected > 0, "sample_weight values must contain at least one positive value"); -selected_rows.resize(n_selected, stream); +selected_rows.resize(ML::narrow_cast<std::size_t>(n_selected), stream);🤖 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/randomforest/randomforest.cuh` around lines 149 - 158, In the selected-row allocation flow, update the resize call after the n_selected positivity assertion to pass ML::narrow_cast<std::size_t>(n_selected) instead of the signed 64-bit count directly. Keep the existing assertion and selected_rows behavior unchanged.Source: Path instructions
🤖 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/randomforest/randomforest.cuh`:
- Around line 149-158: In the selected-row allocation flow, update the resize
call after the n_selected positivity assertion to pass
ML::narrow_cast<std::size_t>(n_selected) instead of the signed 64-bit count
directly. Keep the existing assertion and selected_rows behavior unchanged.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: e4edbf01-4e78-4486-ba6e-0c660835e200
📒 Files selected for processing (1)
cpp/src/randomforest/randomforest.cuh
|
/merge |
|
Merging this now to unblock the random forest development. I created #8407 to follow up with the change to |
Description
This PR cleans up the random forest batched-level decision tree implementation around split-bin lookup and internal index types.
Changes:
lower_boundhelper withcuda::std::lower_bound.cuda::counting_iterator<int, int>for bin lookup so the lower-bound search keeps 32-bit iterator arithmetic.IdxTtemplate plumbing from RF internals.std::int64_twhere they represent rows, columns, nodes, or partition offsets.num_outputs,n_classes, and bin ids asint.The Python API is unchanged. The C++ batched-level RF internals are not treated as a stable public API, so the template simplification should be safe.
Performance
The custom lower-bound implementation did not show a meaningful whole-algorithm benefit over the CUDA standard lower-bound path. The final version keeps the lower-bound iterator arithmetic at 32-bit width for the bin search while using 64-bit indexes for larger dataset-facing RF internals.
Testing
pre-commithooksgit diff --checkcmake --build /home/rorym/cuml-builds/codex-enh-rf-idxt-lower-bound-cleanup/cpp-rf-bench --target cuml -j8