Avoid accidental FP64 promotion in UMAP, t-SNE and HDBSCAN float kernels

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,10 +143,11 @@ 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_t lambda_value = distance > 0.0 ? 1.0 / distance : std::numeric_limits<value_t>::max();
+    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<value_t>::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;

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
  */
 
@@ -16,7 +16,7 @@ template <typename value_idx, typename value_t>
 __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<value_t>::max();
+  return delta > value_t(0.0) ? value_t(1.0) / delta : std::numeric_limits<value_t>::max();
 }
 
 /**

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<value_t>::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<value_t>::max() / n_selected_clusters;
         });
     }
@@ -197,7 +197,7 @@ void all_points_outlier_membership_vector(
     static_cast<value_idx>(n_selected_clusters),
     static_cast<value_idx>(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,

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
  */
 
@@ -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;

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;

cpp/src/tsne/fft_tsne.cuh

@@ -55,7 +55,7 @@ struct FunctionalSqrt {
   template <typename value_t>
   __host__ __device__ float operator()(const value_t& x) const
   {
-    return pow(x, 0.5);
+    return sqrtf(static_cast<float>(x));
   }
 };
 struct FunctionalSquare {

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
  */
 
@@ -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<value_t>& 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);

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
  */
 
@@ -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<T, TPB_X><<<grid, blk>>>(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<T, TPB_X><<<grid, blk, 0, stream>>>(
     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<T, TPB_X><<<grid, blk, 0, stream>>>(
     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));
     });
 
   /**

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
  */
 
@@ -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;
         }
       }
     }