From 3ce76b4f34be02e30ac43db33d128c382f2c4fe3 Mon Sep 17 00:00:00 2001 From: Max Buckley Date: Tue, 1 Sep 2026 15:44:48 +0200 Subject: [PATCH 1/2] enh: avoid accidental FP64 promotion in UMAP, t-SNE and HDBSCAN kernels Several float kernels in these three algorithms contain unsuffixed double-precision literals (`0.5`, `1e-8`, `pow(x, 2.0 * b)`) or `double` locals. In C++ these silently promote the surrounding float expression to double, so the compiler emits FP64 instructions in kernels that are otherwise entirely single-precision. This costs little on datacenter parts, but consumer GPUs have heavily reduced FP64 throughput (1/64 of FP32 on GeForce Blackwell), so the promoted arithmetic is disproportionately expensive there. Verified by compiling to sm_120 SASS and counting FP64-class opcodes (DADD/DMUL/DFMA/DSETP/MUFU.RCP64H plus F2F/I2F conversions) in kernels whose demangled signature contains no `double`. Across umap.cu, tsne.cu and the hdbscan translation units this drops from 2173 to 12; the remaining 12 are in upstream raft::random Box-Muller code. The `double` instantiations of these same templates are unaffected -- no kernel with `double` in its signature changed its FP64 count. Results are unchanged to float round-off. UMAP and HDBSCAN outputs match to ~1e-8 relative (HDBSCAN membership vectors are bit-identical); t-SNE embeddings have the same scale and structure but are not bit-reproducible against the previous build, since the changed rounding diverges over a long iterative optimisation. One change is a deliberate reduction in working precision: compute_membership_strength_kernel previously accumulated in `double` regardless of `value_t`. It now follows `value_t`, which is what the rest of the fuzzy simplicial set construction already uses. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_01NbvBYReympEzkFon5cCnFM --- cpp/src/hdbscan/detail/condense.cuh | 3 ++- cpp/src/hdbscan/detail/kernels/condense.cuh | 2 +- cpp/src/hdbscan/detail/soft_clustering.cuh | 14 ++++++++------ cpp/src/tsne/barnes_hut_kernels.cuh | 2 +- cpp/src/tsne/fft_kernels.cuh | 8 ++++---- cpp/src/tsne/fft_tsne.cuh | 2 +- cpp/src/umap/fuzzy_simpl_set/naive.cuh | 16 ++++++++-------- cpp/src/umap/optimize.cuh | 9 +++++---- cpp/src/umap/supervised.cuh | 6 +++--- 9 files changed, 33 insertions(+), 29 deletions(-) diff --git a/cpp/src/hdbscan/detail/condense.cuh b/cpp/src/hdbscan/detail/condense.cuh index 8f4148d382..5adb3a38b2 100644 --- a/cpp/src/hdbscan/detail/condense.cuh +++ b/cpp/src/hdbscan/detail/condense.cuh @@ -146,7 +146,8 @@ void _build_condensed_hierarchy(const raft::handle_t& handle, value_idx left = h_children[(node - n_samples) * 2]; value_idx right = h_children[(node - n_samples) * 2 + 1]; value_t distance = h_delta[node - n_samples]; - value_t lambda_value = distance > 0.0 ? 1.0 / distance : std::numeric_limits::max(); + value_t lambda_value = + distance > value_t(0.0) ? value_t(1.0) / distance : std::numeric_limits::max(); value_idx left_count = left >= n_samples ? h_sizes[left - n_samples] : 1; value_idx right_count = right >= n_samples ? h_sizes[right - n_samples] : 1; diff --git a/cpp/src/hdbscan/detail/kernels/condense.cuh b/cpp/src/hdbscan/detail/kernels/condense.cuh index 88f6065fe0..272cdc383f 100644 --- a/cpp/src/hdbscan/detail/kernels/condense.cuh +++ b/cpp/src/hdbscan/detail/kernels/condense.cuh @@ -16,7 +16,7 @@ template __device__ inline value_t get_lambda(value_idx node, value_idx num_points, const value_t* deltas) { value_t delta = deltas[node - num_points]; - return delta > 0.0 ? 1.0 / delta : std::numeric_limits::max(); + return delta > value_t(0.0) ? value_t(1.0) / delta : std::numeric_limits::max(); } /** diff --git a/cpp/src/hdbscan/detail/soft_clustering.cuh b/cpp/src/hdbscan/detail/soft_clustering.cuh index 75f1165142..1ae04d9df5 100644 --- a/cpp/src/hdbscan/detail/soft_clustering.cuh +++ b/cpp/src/hdbscan/detail/soft_clustering.cuh @@ -125,7 +125,7 @@ void dist_membership_vector(const raft::handle_t& handle, samples_per_batch * n_selected_clusters), [min_dist = min_dist.data_handle()] __device__(auto idx) { value_t val = min_dist[idx]; - if (val != 0) { return value_t(exp(1.0 / val)); } + if (val != 0) { return exp(value_t(1.0) / val); } return std::numeric_limits::max(); }); } @@ -139,7 +139,7 @@ void dist_membership_vector(const raft::handle_t& handle, samples_per_batch * n_selected_clusters), [min_dist = min_dist.data_handle(), n_selected_clusters] __device__(auto idx) { value_t val = min_dist[idx]; - if (val > 0) { return value_t(1.0 / val); } + if (val > 0) { return value_t(1.0) / val; } return std::numeric_limits::max() / n_selected_clusters; }); } @@ -197,7 +197,7 @@ void all_points_outlier_membership_vector( static_cast(n_selected_clusters), static_cast(m), [] __device__(value_t mat_in, value_t vec_in) { - return exp(-(vec_in + 1e-8) / mat_in); + return exp(-(vec_in + value_t(1e-8)) / mat_in); }, //+ 1e-8 to avoid zero lambda stream); @@ -310,7 +310,7 @@ void outlier_membership_vector(const raft::handle_t& handle, n_prediction_points, [] __device__(value_t mat_in, value_t vec_in) { value_t denominator = vec_in - mat_in; - if (denominator <= 0) { denominator = 1e-8; } + if (denominator <= 0) { denominator = value_t(1e-8); } return vec_in / denominator; }, stream); @@ -359,7 +359,8 @@ void prob_in_some_cluster(const raft::handle_t& handle, n_selected_clusters] __device__(auto idx) { value_idx nearest_cluster = height_argmax[idx]; value_t max_lambda = - max(prediction_lambdas[idx], deaths[selected_clusters[nearest_cluster] - n_leaves]) + 1e-8; + max(prediction_lambdas[idx], deaths[selected_clusters[nearest_cluster] - n_leaves]) + + value_t(1e-8); return merge_heights[idx * n_selected_clusters + nearest_cluster] / max_lambda; }; raft::linalg::map_offset( @@ -589,7 +590,8 @@ void membership_vector(const raft::handle_t& handle, auto combine_op = [membership_vec, dist_membership_vec = dist_membership_vec.data()] __device__(auto idx) { - return pow(membership_vec[idx], 2) * pow(dist_membership_vec[idx], 0.5); + value_t m = membership_vec[idx]; + return m * m * sqrt(dist_membership_vec[idx]); }; raft::linalg::map_offset(handle, diff --git a/cpp/src/tsne/barnes_hut_kernels.cuh b/cpp/src/tsne/barnes_hut_kernels.cuh index 9eeb3dba13..7b3e7288e0 100644 --- a/cpp/src/tsne/barnes_hut_kernels.cuh +++ b/cpp/src/tsne/barnes_hut_kernels.cuh @@ -296,7 +296,7 @@ CUML_KERNEL __launch_bounds__(THREADS2) void TreeBuildingKernel(/* int *restrict y += ((y < py) ? (j |= 2, r) : (-r)); ch = childd[n * 4 + j]; - if (r <= 1e-10) { break; } + if (r <= value_t(1e-10)) { break; } } childd[n * 4 + j] = i; diff --git a/cpp/src/tsne/fft_kernels.cuh b/cpp/src/tsne/fft_kernels.cuh index 146b8cd5e2..b4a9c77c5a 100644 --- a/cpp/src/tsne/fft_kernels.cuh +++ b/cpp/src/tsne/fft_kernels.cuh @@ -527,10 +527,10 @@ CUML_KERNEL void IntegrationKernel(volatile value_t* __restrict__ points, value_t dy = exaggeration * attr_forces[i + num_points] - (rep_forces[i + num_points] / normalization); - gx = signbit(dx) != signbit(ux) ? gx + 0.2 : gx * 0.8; - gy = signbit(dy) != signbit(uy) ? gy + 0.2 : gy * 0.8; - gx = gx < 0.01 ? 0.01 : gx; - gy = gy < 0.01 ? 0.01 : gy; + gx = signbit(dx) != signbit(ux) ? gx + value_t(0.2) : gx * value_t(0.8); + gy = signbit(dy) != signbit(uy) ? gy + value_t(0.2) : gy * value_t(0.8); + gx = gx < value_t(0.01) ? value_t(0.01) : gx; + gy = gy < value_t(0.01) ? value_t(0.01) : gy; ux = momentum * ux - eta * gx * dx; uy = momentum * uy - eta * gy * dy; diff --git a/cpp/src/tsne/fft_tsne.cuh b/cpp/src/tsne/fft_tsne.cuh index 34adb310a8..08cdb53210 100644 --- a/cpp/src/tsne/fft_tsne.cuh +++ b/cpp/src/tsne/fft_tsne.cuh @@ -55,7 +55,7 @@ struct FunctionalSqrt { template __host__ __device__ float operator()(const value_t& x) const { - return pow(x, 0.5); + return sqrtf(static_cast(x)); } }; struct FunctionalSquare { diff --git a/cpp/src/umap/fuzzy_simpl_set/naive.cuh b/cpp/src/umap/fuzzy_simpl_set/naive.cuh index 84d3646f7d..c0df2df46f 100644 --- a/cpp/src/umap/fuzzy_simpl_set/naive.cuh +++ b/cpp/src/umap/fuzzy_simpl_set/naive.cuh @@ -207,18 +207,18 @@ CUML_KERNEL void compute_membership_strength_kernel( if (idx < to_process) { int row = idx / n_neighbors; // one neighbor per thread - double cur_rho = rhos[row]; - double cur_sigma = sigmas[row]; + value_t cur_rho = rhos[row]; + value_t cur_sigma = sigmas[row]; value_idx cur_knn_ind = knn_indices[idx]; - double cur_knn_dist = knn_dists[idx]; + value_t cur_knn_dist = knn_dists[idx]; if (cur_knn_ind != -1) { - double val = 0.0; + value_t val = value_t(0.0); if (cur_knn_ind == row) - val = 0.0; - else if (cur_knn_dist - cur_rho <= 0.0 || cur_sigma == 0.0) - val = 1.0; + val = value_t(0.0); + else if (cur_knn_dist - cur_rho <= value_t(0.0) || cur_sigma == value_t(0.0)) + val = value_t(1.0); else { val = exp(-((cur_knn_dist - cur_rho) / (cur_sigma))); @@ -355,7 +355,7 @@ void symmetrize(raft::sparse::COO& in, [set_op_mix_ratio] __device__(int row, int col, value_t result, value_t transpose) { value_t prod_matrix = result * transpose; value_t res = set_op_mix_ratio * (result + transpose - prod_matrix) + - (1.0 - set_op_mix_ratio) * prod_matrix; + (value_t(1.0) - set_op_mix_ratio) * prod_matrix; return res; }, stream); diff --git a/cpp/src/umap/optimize.cuh b/cpp/src/umap/optimize.cuh index 2ef1f6f6ea..5d8ddb8f5b 100644 --- a/cpp/src/umap/optimize.cuh +++ b/cpp/src/umap/optimize.cuh @@ -36,7 +36,7 @@ CUML_KERNEL void map_kernel(T* output, T* X, int n_rows, T* coef, Lambda grad) T a = coef[0]; T b = coef[1]; output[row] = grad(x, a, b); - if (isnan(output[row])) output[row] = 0.0; + if (isnan(output[row])) output[row] = T(0.0); } } @@ -52,7 +52,7 @@ void f(T* input, int n_rows, T* coef, T* preds) // Function: 1/1+ax^(2b) map_kernel<<>>(preds, input, n_rows, coef, [] __device__(T x, T a, T b) { - return 1.0 / (1 + a * pow(x, 2.0 * b)); + return T(1.0) / (T(1.0) + a * pow(x, T(2.0) * b)); }); } @@ -83,7 +83,7 @@ void abLossGrads( raft::copy(a_deriv.data(), input, n_rows, stream); map_kernel<<>>( a_deriv.data(), a_deriv.data(), n_rows, coef, [] __device__ __host__(T x, T a, T b) { - return -(pow(x, 2.0 * b)) / pow((1.0 + a * pow(x, 2.0 * b)), 2.0); + return -(pow(x, T(2.0) * b)) / pow((T(1.0) + a * pow(x, T(2.0) * b)), T(2.0)); }); raft::linalg::eltwiseMultiply(a_deriv.data(), a_deriv.data(), residuals.data(), n_rows, stream); @@ -96,7 +96,8 @@ void abLossGrads( raft::copy(b_deriv.data(), input, n_rows, stream); map_kernel<<>>( b_deriv.data(), b_deriv.data(), n_rows, coef, [] __device__ __host__(T x, T a, T b) { - return -(2.0 * a * pow(x, 2.0 * b) * log(x)) / pow(1 + a * pow(x, 2.0 * b), 2.0); + return -(T(2.0) * a * pow(x, T(2.0) * b) * log(x)) / + pow(T(1.0) + a * pow(x, T(2.0) * b), T(2.0)); }); /** diff --git a/cpp/src/umap/supervised.cuh b/cpp/src/umap/supervised.cuh index 4fceb52fab..7de90a9689 100644 --- a/cpp/src/umap/supervised.cuh +++ b/cpp/src/umap/supervised.cuh @@ -152,10 +152,10 @@ CUML_KERNEL void sset_intersection_kernel(int* row_ind1, } if (left_val > left_min || right_val > right_min) { - if (mix_weight < 0.5) { - result_vals[j] = left_val * powf(right_val, mix_weight / (1.0 - mix_weight)); + if (mix_weight < 0.5f) { + result_vals[j] = left_val * powf(right_val, mix_weight / (1.0f - mix_weight)); } else { - result_vals[j] = powf(left_val, (1.0 - mix_weight) / mix_weight) * right_val; + result_vals[j] = powf(left_val, (1.0f - mix_weight) / mix_weight) * right_val; } } } From 5a6d02d65c95cff2f9f74888d7ec2dfdbb5ccf87 Mon Sep 17 00:00:00 2001 From: Max Buckley Date: Thu, 3 Sep 2026 10:11:58 +0200 Subject: [PATCH 2/2] Apply pre-commit: copyright headers and clang-format `pre-commit run --all-files` flagged six of the files this PR touches: five stale/non-canonical copyright notices and one clang-format realignment in condense.cuh caused by this PR's own edit. --- cpp/src/hdbscan/detail/condense.cuh | 8 ++++---- cpp/src/hdbscan/detail/kernels/condense.cuh | 2 +- cpp/src/tsne/barnes_hut_kernels.cuh | 2 +- cpp/src/umap/fuzzy_simpl_set/naive.cuh | 2 +- cpp/src/umap/optimize.cuh | 2 +- cpp/src/umap/supervised.cuh | 2 +- 6 files changed, 9 insertions(+), 9 deletions(-) diff --git a/cpp/src/hdbscan/detail/condense.cuh b/cpp/src/hdbscan/detail/condense.cuh index 5adb3a38b2..1c37535812 100644 --- a/cpp/src/hdbscan/detail/condense.cuh +++ b/cpp/src/hdbscan/detail/condense.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -143,9 +143,9 @@ void _build_condensed_hierarchy(const raft::handle_t& handle, // Skip if already processed or is a leaf if (ignore[node] || node < n_samples) { continue; } - value_idx left = h_children[(node - n_samples) * 2]; - value_idx right = h_children[(node - n_samples) * 2 + 1]; - value_t distance = h_delta[node - n_samples]; + value_idx left = h_children[(node - n_samples) * 2]; + value_idx right = h_children[(node - n_samples) * 2 + 1]; + value_t distance = h_delta[node - n_samples]; value_t lambda_value = distance > value_t(0.0) ? value_t(1.0) / distance : std::numeric_limits::max(); diff --git a/cpp/src/hdbscan/detail/kernels/condense.cuh b/cpp/src/hdbscan/detail/kernels/condense.cuh index 272cdc383f..c86daa0270 100644 --- a/cpp/src/hdbscan/detail/kernels/condense.cuh +++ b/cpp/src/hdbscan/detail/kernels/condense.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/tsne/barnes_hut_kernels.cuh b/cpp/src/tsne/barnes_hut_kernels.cuh index 7b3e7288e0..09dcee10e7 100644 --- a/cpp/src/tsne/barnes_hut_kernels.cuh +++ b/cpp/src/tsne/barnes_hut_kernels.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2021-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2021-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/umap/fuzzy_simpl_set/naive.cuh b/cpp/src/umap/fuzzy_simpl_set/naive.cuh index c0df2df46f..b18810c20b 100644 --- a/cpp/src/umap/fuzzy_simpl_set/naive.cuh +++ b/cpp/src/umap/fuzzy_simpl_set/naive.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/umap/optimize.cuh b/cpp/src/umap/optimize.cuh index 5d8ddb8f5b..1ee5a1fd96 100644 --- a/cpp/src/umap/optimize.cuh +++ b/cpp/src/umap/optimize.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ diff --git a/cpp/src/umap/supervised.cuh b/cpp/src/umap/supervised.cuh index 7de90a9689..0b3b8d6988 100644 --- a/cpp/src/umap/supervised.cuh +++ b/cpp/src/umap/supervised.cuh @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */