Fix incorrect async initialisations

cpp/src/solver/cd.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -213,8 +213,9 @@ int cdFit(const raft::handle_t& handle,
   rmm::device_uvector<ConvState<math_t>> convStateBuf(1, stream);
   auto convStateLoc = convStateBuf.data();
 
-  rmm::device_scalar<math_t> cublas_alpha(1.0, stream);
-  rmm::device_scalar<math_t> cublas_beta(0.0, stream);
+  // Passed to gemv in host pointer mode, so cuBLAS reads them during the call.
+  math_t const cublas_alpha = 1.0;
+  math_t const cublas_beta  = 0.0;
 
   int n_iter = 0;
   while (n_iter < epochs) {
@@ -238,19 +239,19 @@ int cdFit(const raft::handle_t& handle,
         handle, n_rows, coef_loc, input_col_loc, 1, residual.data(), 1, stream);
 
       // coef[ci] = dot(X[:, ci], residual[:])
-      raft::linalg::gemv<math_t, true>(handle,
-                                       false,
-                                       1,
-                                       n_rows,
-                                       cublas_alpha.data(),
-                                       input_col_loc,
-                                       1,
-                                       residual.data(),
-                                       1,
-                                       cublas_beta.data(),
-                                       coef_loc,
-                                       1,
-                                       stream);
+      raft::linalg::gemv<math_t, false>(handle,
+                                        false,
+                                        1,
+                                        n_rows,
+                                        &cublas_alpha,
+                                        input_col_loc,
+                                        1,
+                                        residual.data(),
+                                        1,
+                                        &cublas_beta,
+                                        coef_loc,
+                                        1,
+                                        stream);
 
       // Calculate the new coefficient that minimizes f along coordinate line ci
       // coef[ci] = SoftTreshold(dot(X[:, ci], residual[:]), l1_alpha) /  dot(X[:, ci], X[:, ci]))

cpp/src/tsne/fft_tsne.cuh

@@ -131,10 +131,13 @@ std::pair<value_t, value_t> min_max(const value_t* Y, const value_idx n, cudaStr
   rmm::device_scalar<value_t> min_d(stream);
   rmm::device_scalar<value_t> max_d(stream);
 
-  value_t val = std::numeric_limits<value_t>::max();
-  min_d.set_value_async(val, stream);
-  val = std::numeric_limits<value_t>::lowest();
-  max_d.set_value_async(val, stream);
+  // Each copy needs its own host source, alive and unmodified until the stream
+  // is synchronized below, because rmm may defer reading it until the stream
+  // reaches the copy.
+  value_t const min_init = std::numeric_limits<value_t>::max();
+  value_t const max_init = std::numeric_limits<value_t>::lowest();
+  min_d.set_value_async(min_init, stream);
+  max_d.set_value_async(max_init, stream);
 
   auto nthreads = 256;
   auto nblocks  = raft::ceildiv(n, (value_idx)nthreads);

cpp/src/umap/simpl_set_embed/algo.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
  */
 
@@ -174,7 +174,8 @@ bool check_outliers(const int* rows, int m, nnz_t nnz, int threshold, cudaStream
   dim3 blk(TPB_X, 1, 1);
   compute_degrees_kernel<<<grid_nnz, blk, 0, stream>>>(rows, nnz, graph_degree_head.data());
 
-  rmm::device_scalar<bool> has_outlier_d(0, stream);  // initialize to 0
+  rmm::device_scalar<bool> has_outlier_d(stream);
+  raft::linalg::zero(has_outlier_d.data(), 1, stream);
 
   dim3 grid_head_n(raft::ceildiv(static_cast<nnz_t>(m), static_cast<nnz_t>(TPB_X)), 1, 1);
   check_threshold_kernel<<<grid_head_n, blk, 0, stream>>>(