Skip to content

Fix FastIntDiv tparam #8299

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
rapids-bot merged 1 commit into from
Jun 29, 2026
Merged

Fix FastIntDiv tparam #8299

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions cpp/src/hdbscan/detail/kernels/soft_clustering.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@
/*
* SPDX-FileCopyrightText: Copyright (c) 2022-2024, NVIDIA CORPORATION.
* SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
* SPDX-License-Identifier: Apache-2.0
*/
#pragma once
Expand All @@ -18,7 +18,7 @@ CUML_KERNEL void merge_height_kernel(value_t* heights,
value_idx* parents,
size_t m,
value_idx n_selected_clusters,
raft::util::FastIntDiv n,
raft::util::FastIntDiv<value_idx> n,
value_idx* selected_clusters)
{
value_idx idx = blockDim.x * blockIdx.x + threadIdx.x;
Expand Down Expand Up @@ -62,7 +62,7 @@ CUML_KERNEL void merge_height_kernel(value_t* heights,
value_idx* parents,
size_t n_prediction_points,
value_idx n_selected_clusters,
raft::util::FastIntDiv n,
raft::util::FastIntDiv<value_idx> n,
value_idx* selected_clusters)
{
value_idx idx = blockDim.x * blockIdx.x + threadIdx.x;
Expand Down
40 changes: 21 additions & 19 deletions cpp/src/hdbscan/detail/soft_clustering.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -99,7 +99,7 @@ void dist_membership_vector(const raft::handle_t& handle,

auto reduction_op = [dist = dist.data(),
batch_offset,
divisor = raft::util::FastIntDiv(n_selected_clusters),
divisor = raft::util::FastIntDiv<value_idx>(n_selected_clusters),
n_selected_clusters,
n_exemplars,
exemplar_label_offsets] __device__(auto idx) {
Expand Down Expand Up @@ -172,14 +172,15 @@ void all_points_outlier_membership_vector(
auto n_leaves = condensed_tree.get_n_leaves();

int n_blocks = raft::ceildiv(int(m * n_selected_clusters), tpb);
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(n_selected_clusters),
selected_clusters);
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);
Comment on lines +175 to +183

@coderabbitai coderabbitai Bot Jun 29, 2026

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🩺 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


auto leaf_max_lambdas = raft::make_device_vector<value_t, value_idx>(handle, n_leaves);

Expand Down Expand Up @@ -277,16 +278,17 @@ void outlier_membership_vector(const raft::handle_t& handle,

// Using the nearest neighbor indices, compute outlier membership
int n_blocks = raft::ceildiv(int(n_prediction_points * n_selected_clusters), tpb);
merge_height_kernel<<<n_blocks, tpb, 0, stream>>>(merge_heights,
lambdas,
prediction_lambdas,
min_mr_inds,
index_into_children,
parents,
n_prediction_points,
static_cast<value_idx>(n_selected_clusters),
raft::util::FastIntDiv(n_selected_clusters),
selected_clusters);
merge_height_kernel<<<n_blocks, tpb, 0, stream>>>(
merge_heights,
lambdas,
prediction_lambdas,
min_mr_inds,
index_into_children,
parents,
n_prediction_points,
static_cast<value_idx>(n_selected_clusters),
raft::util::FastIntDiv<value_idx>(static_cast<value_idx>(n_selected_clusters)),
selected_clusters);

// fetch the max lambda of the cluster to which the nearest MR neighbor belongs in the condensed
// hierarchy
Expand Down
Loading