Fix FastIntDiv tparam - #8299
Conversation
📝 WalkthroughSummary by CodeRabbit
Walkthrough
ChangesFastIntDiv explicit template parameter typing
Estimated code review effort🎯 1 (Trivial) | ⏱️ ~3 minutes 🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Comment |
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/hdbscan/detail/soft_clustering.cuh`:
- Around line 175-183: Guard the narrowing of n_selected_clusters before
launching merge_height_kernel and the other affected entry point in
soft_clustering.cuh: both currently pass a size_t through static_cast<value_idx>
for the kernel count and FastIntDiv divisor, which can truncate large values and
break indexing. Replace the unchecked cast with
ML::narrow_cast<value_idx>(n_selected_clusters) at the call sites (or widen the
API if value_idx cannot safely hold the full range), and use the narrowed value
consistently for the launch arguments in merge_height_kernel and the
corresponding second launch site.
🪄 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: 0c1bb20f-834c-462b-8037-faa97ee25d71
📒 Files selected for processing (2)
cpp/src/hdbscan/detail/kernels/soft_clustering.cuhcpp/src/hdbscan/detail/soft_clustering.cuh
| merge_height_kernel<<<n_blocks, tpb, 0, stream>>>( | ||
| merge_heights, | ||
| lambdas, | ||
| index_into_children, | ||
| parents, | ||
| m, | ||
| static_cast<value_idx>(n_selected_clusters), | ||
| raft::util::FastIntDiv<value_idx>(static_cast<value_idx>(n_selected_clusters)), | ||
| selected_clusters); |
There was a problem hiding this comment.
🩺 Stability & Availability | 🟠 Major | ⚡ Quick win
HIGH: Guard the size_t→value_idx cast before these launches.
Both entry points take n_selected_clusters as size_t, but the new launch arguments narrow it with static_cast<value_idx> for both the kernel count and the FastIntDiv divisor. If that count ever exceeds value_idx, the kernel sees a truncated divisor/count and its row/column indexing becomes invalid.
Suggested fix
+ auto n_selected_clusters_idx = ML::narrow_cast<value_idx>(n_selected_clusters);
merge_height_kernel<<<n_blocks, tpb, 0, stream>>>(
merge_heights,
lambdas,
...
- static_cast<value_idx>(n_selected_clusters),
- raft::util::FastIntDiv<value_idx>(static_cast<value_idx>(n_selected_clusters)),
+ n_selected_clusters_idx,
+ raft::util::FastIntDiv<value_idx>(n_selected_clusters_idx),
selected_clusters);As per path instructions, host-side count/dimension narrowing must use ML::narrow_cast<...>(...) or widen the API.
Also applies to: 281-291
🤖 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/hdbscan/detail/soft_clustering.cuh` around lines 175 - 183, Guard the
narrowing of n_selected_clusters before launching merge_height_kernel and the
other affected entry point in soft_clustering.cuh: both currently pass a size_t
through static_cast<value_idx> for the kernel count and FastIntDiv divisor,
which can truncate large values and break indexing. Replace the unchecked cast
with ML::narrow_cast<value_idx>(n_selected_clusters) at the call sites (or widen
the API if value_idx cannot safely hold the full range), and use the narrowed
value consistently for the launch arguments in merge_height_kernel and the
corresponding second launch site.
Source: Path instructions
|
/merge |
Add explicit template parameters for FastIntDiv, follow up to NVIDIA/raft#3059.