diff --git a/cpp/src/solver/cd.cuh b/cpp/src/solver/cd.cuh index 0fd646ebe0..1da8b4a764 100644 --- a/cpp/src/solver/cd.cuh +++ b/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> convStateBuf(1, stream); auto convStateLoc = convStateBuf.data(); - rmm::device_scalar cublas_alpha(1.0, stream); - rmm::device_scalar 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(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(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])) diff --git a/cpp/src/tsne/fft_tsne.cuh b/cpp/src/tsne/fft_tsne.cuh index e0a92b2df7..34adb310a8 100644 --- a/cpp/src/tsne/fft_tsne.cuh +++ b/cpp/src/tsne/fft_tsne.cuh @@ -131,10 +131,13 @@ std::pair min_max(const value_t* Y, const value_idx n, cudaStr rmm::device_scalar min_d(stream); rmm::device_scalar max_d(stream); - value_t val = std::numeric_limits::max(); - min_d.set_value_async(val, stream); - val = std::numeric_limits::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::max(); + value_t const max_init = std::numeric_limits::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); diff --git a/cpp/src/umap/simpl_set_embed/algo.cuh b/cpp/src/umap/simpl_set_embed/algo.cuh index b969af4dc9..97e0c68209 100644 --- a/cpp/src/umap/simpl_set_embed/algo.cuh +++ b/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<<>>(rows, nnz, graph_degree_head.data()); - rmm::device_scalar has_outlier_d(0, stream); // initialize to 0 + rmm::device_scalar has_outlier_d(stream); + raft::linalg::zero(has_outlier_d.data(), 1, stream); dim3 grid_head_n(raft::ceildiv(static_cast(m), static_cast(TPB_X)), 1, 1); check_threshold_kernel<<>>(