From 6b94e4fd4de09f50527c172566f3433af69cb26b Mon Sep 17 00:00:00 2001 From: Ben Frederickson Date: Thu, 11 May 2023 11:31:23 -0700 Subject: [PATCH] Remove raft/matrix/matrix.cuh includes (#1498) The `raft/matrix/matrix.cuh` file has been marked as deprecated, and produces a compile warning when included. However it was still being referenced in a bunch of different spots within raft - making it hard to avoid these warnings. Remove the includes, in favour of either the newer API's or in certain cases the detail API Authors: - Ben Frederickson (https://github.com/benfred) Approvers: - Corey J. Nolet (https://github.com/cjnolet) URL: https://github.com/rapidsai/raft/pull/1498 --- .../raft/cluster/detail/kmeans_balanced.cuh | 1 - cpp/include/raft/linalg/detail/eig.cuh | 32 +++++++---- cpp/include/raft/linalg/detail/lstsq.cuh | 1 - .../raft/linalg/detail/matrix_vector_op.cuh | 53 ++++++++++++++---- cpp/include/raft/linalg/detail/qr.cuh | 7 ++- cpp/include/raft/linalg/detail/rsvd.cuh | 54 ++++++++++--------- cpp/include/raft/linalg/detail/svd.cuh | 31 +++++++---- cpp/include/raft/linalg/matrix_vector_op.cuh | 3 +- cpp/include/raft/matrix/copy.cuh | 20 ++++++- cpp/include/raft/matrix/diagonal.cuh | 17 +++--- cpp/include/raft/matrix/init.cuh | 1 - cpp/include/raft/matrix/linewise_op.cuh | 4 +- cpp/include/raft/matrix/print.cuh | 1 - cpp/include/raft/matrix/sign_flip.cuh | 1 - cpp/include/raft/matrix/slice.cuh | 2 +- cpp/include/raft/matrix/sqrt.cuh | 1 - cpp/include/raft/matrix/triangular.cuh | 13 +++-- .../raft/neighbors/detail/ivf_pq_build.cuh | 8 ++- cpp/include/raft/neighbors/refine-inl.cuh | 1 - .../raft/random/detail/make_regression.cuh | 7 ++- .../raft/sparse/neighbors/detail/knn.cuh | 1 - .../raft/spatial/knn/detail/ball_cover.cuh | 20 +++---- cpp/test/linalg/svd.cu | 1 - cpp/test/matrix/matrix.cu | 2 +- cpp/test/neighbors/ann_utils.cuh | 9 ++-- 25 files changed, 191 insertions(+), 100 deletions(-) diff --git a/cpp/include/raft/cluster/detail/kmeans_balanced.cuh b/cpp/include/raft/cluster/detail/kmeans_balanced.cuh index eb89ebe402..9e5f7a7c9a 100644 --- a/cpp/include/raft/cluster/detail/kmeans_balanced.cuh +++ b/cpp/include/raft/cluster/detail/kmeans_balanced.cuh @@ -37,7 +37,6 @@ #include #include #include -#include #include #include #include diff --git a/cpp/include/raft/linalg/detail/eig.cuh b/cpp/include/raft/linalg/detail/eig.cuh index 94493efb24..7896136631 100644 --- a/cpp/include/raft/linalg/detail/eig.cuh +++ b/cpp/include/raft/linalg/detail/eig.cuh @@ -19,7 +19,7 @@ #include "cusolver_wrappers.hpp" #include #include -#include +#include #include #include #include @@ -52,7 +52,9 @@ void eigDC_legacy(raft::device_resources const& handle, rmm::device_uvector d_work(lwork, stream); rmm::device_scalar d_dev_info(stream); - raft::matrix::copy(in, eig_vectors, n_rows, n_cols, stream); + raft::matrix::copy(handle, + make_device_matrix_view(in, n_rows, n_cols), + make_device_matrix_view(eig_vectors, n_rows, n_cols)); RAFT_CUSOLVER_TRY(cusolverDnsyevd(cusolverH, CUSOLVER_EIG_MODE_VECTOR, @@ -108,7 +110,9 @@ void eigDC(raft::device_resources const& handle, rmm::device_scalar d_dev_info(stream); std::vector h_work(workspaceHost / sizeof(math_t)); - raft::matrix::copy(in, eig_vectors, n_rows, n_cols, stream); + raft::matrix::copy(handle, + make_device_matrix_view(in, n_rows, n_cols), + make_device_matrix_view(eig_vectors, n_rows, n_cols)); RAFT_CUSOLVER_TRY(cusolverDnxsyevd(cusolverH, dn_params, @@ -191,7 +195,9 @@ void eigSelDC(raft::device_resources const& handle, stream)); } else if (memUsage == COPY_INPUT) { d_eig_vectors.resize(n_rows * n_cols, stream); - raft::matrix::copy(in, d_eig_vectors.data(), n_rows, n_cols, stream); + raft::matrix::copy(handle, + make_device_matrix_view(in, n_rows, n_cols), + make_device_matrix_view(eig_vectors, n_rows, n_cols)); RAFT_CUSOLVER_TRY(cusolverDnsyevdx(cusolverH, CUSOLVER_EIG_MODE_VECTOR, @@ -220,10 +226,16 @@ void eigSelDC(raft::device_resources const& handle, "This usually occurs when some of the features do not vary enough."); if (memUsage == OVERWRITE_INPUT) { - raft::matrix::truncZeroOrigin(in, n_rows, eig_vectors, n_rows, n_eig_vals, stream); + raft::matrix::trunc_zero_origin( + handle, + make_device_matrix_view(in, n_rows, n_eig_vals), + make_device_matrix_view(eig_vectors, n_rows, n_eig_vals)); } else if (memUsage == COPY_INPUT) { - raft::matrix::truncZeroOrigin( - d_eig_vectors.data(), n_rows, eig_vectors, n_rows, n_eig_vals, stream); + raft::matrix::trunc_zero_origin( + handle, + make_device_matrix_view( + d_eig_vectors.data(), n_rows, n_eig_vals), + make_device_matrix_view(eig_vectors, n_rows, n_eig_vals)); } } @@ -259,7 +271,9 @@ void eigJacobi(raft::device_resources const& handle, rmm::device_uvector d_work(lwork, stream); rmm::device_scalar dev_info(stream); - raft::matrix::copy(in, eig_vectors, n_rows, n_cols, stream); + raft::matrix::copy(handle, + make_device_matrix_view(in, n_rows, n_cols), + make_device_matrix_view(eig_vectors, n_rows, n_cols)); RAFT_CUSOLVER_TRY(cusolverDnsyevj(cusolverH, CUSOLVER_EIG_MODE_VECTOR, @@ -283,4 +297,4 @@ void eigJacobi(raft::device_resources const& handle, } // namespace detail } // namespace linalg -} // namespace raft \ No newline at end of file +} // namespace raft diff --git a/cpp/include/raft/linalg/detail/lstsq.cuh b/cpp/include/raft/linalg/detail/lstsq.cuh index 207bcefc32..fd6b00f9fd 100644 --- a/cpp/include/raft/linalg/detail/lstsq.cuh +++ b/cpp/include/raft/linalg/detail/lstsq.cuh @@ -28,7 +28,6 @@ #include #include #include -#include #include #include #include diff --git a/cpp/include/raft/linalg/detail/matrix_vector_op.cuh b/cpp/include/raft/linalg/detail/matrix_vector_op.cuh index 62ec9bb7a4..0c1261261c 100644 --- a/cpp/include/raft/linalg/detail/matrix_vector_op.cuh +++ b/cpp/include/raft/linalg/detail/matrix_vector_op.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2022, NVIDIA CORPORATION. + * Copyright (c) 2022-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -16,7 +16,7 @@ #pragma once -#include +#include namespace raft { namespace linalg { @@ -33,10 +33,26 @@ void matrixVectorOp(MatT* out, Lambda op, cudaStream_t stream) { - IdxType stride = rowMajor ? D : N; - IdxType nLines = rowMajor ? N : D; - return matrix::linewiseOp( - out, matrix, stride, nLines, rowMajor == bcastAlongRows, op, stream, vec); + raft::device_resources handle(stream); + + bool along_lines = rowMajor == bcastAlongRows; + if (rowMajor) { + matrix::linewise_op( + handle, + make_device_matrix_view(matrix, N, D), + make_device_matrix_view(out, N, D), + along_lines, + op, + make_device_vector_view(vec, bcastAlongRows ? N : D)); + } else { + matrix::linewise_op( + handle, + make_device_matrix_view(matrix, N, D), + make_device_matrix_view(out, N, D), + along_lines, + op, + make_device_vector_view(vec, bcastAlongRows ? N : D)); + } } template ( + handle, + make_device_matrix_view(matrix, N, D), + make_device_matrix_view(out, N, D), + along_lines, + op, + make_device_vector_view(vec1, bcastAlongRows ? N : D), + make_device_vector_view(vec2, bcastAlongRows ? N : D)); + } else { + matrix::linewise_op( + handle, + make_device_matrix_view(matrix, N, D), + make_device_matrix_view(out, N, D), + along_lines, + op, + make_device_vector_view(vec1, bcastAlongRows ? N : D), + make_device_vector_view(vec2, bcastAlongRows ? N : D)); + } } }; // end namespace detail diff --git a/cpp/include/raft/linalg/detail/qr.cuh b/cpp/include/raft/linalg/detail/qr.cuh index bc7c551d89..16a721dfd3 100644 --- a/cpp/include/raft/linalg/detail/qr.cuh +++ b/cpp/include/raft/linalg/detail/qr.cuh @@ -20,7 +20,7 @@ #include "cusolver_wrappers.hpp" #include #include -#include +#include #include #include @@ -132,7 +132,10 @@ void qrGetQR(raft::resources const& handle, devInfo.data(), stream)); - raft::matrix::copyUpperTriangular(R_full.data(), R, m, n, stream); + raft::matrix::upper_triangular( + handle, + make_device_matrix_view(R_full.data(), m, n), + make_device_matrix_view(R, std::min(m, n), std::min(m, n))); RAFT_CUDA_TRY( cudaMemcpyAsync(Q, R_full.data(), sizeof(math_t) * m * n, cudaMemcpyDeviceToDevice, stream)); diff --git a/cpp/include/raft/linalg/detail/rsvd.cuh b/cpp/include/raft/linalg/detail/rsvd.cuh index a66a23179b..48b9e1d2db 100644 --- a/cpp/include/raft/linalg/detail/rsvd.cuh +++ b/cpp/include/raft/linalg/detail/rsvd.cuh @@ -21,8 +21,11 @@ #include #include #include +#include #include -#include +#include +#include +#include #include #include @@ -202,15 +205,13 @@ void rsvdFixedRank(raft::device_resources const& handle, true, true, stream); - raft::matrix::sliceMatrix(S_vec_tmp.data(), - 1, - l, - S_vec, - 0, - 0, - 1, - k, - stream); // First k elements of S_vec + + // First k elements of S_vec + raft::matrix::slice( + handle, + make_device_matrix_view(S_vec_tmp.data(), 1, l), + make_device_matrix_view(S_vec, 1, k), + raft::matrix::slice_coordinates(0, 0, 1, k)); // Merge step 14 & 15 by calculating U = Q*Vhat[:,1:k] mxl * lxk = mxk if (gen_left_vec) { @@ -272,23 +273,26 @@ void rsvdFixedRank(raft::device_resources const& handle, RAFT_CUDA_TRY(cudaMemsetAsync(Uhat.data(), 0, sizeof(math_t) * l * l, stream)); rmm::device_uvector Uhat_dup(l * l, stream); RAFT_CUDA_TRY(cudaMemsetAsync(Uhat_dup.data(), 0, sizeof(math_t) * l * l, stream)); - raft::matrix::copyUpperTriangular(BBt.data(), Uhat_dup.data(), l, l, stream); + + raft::matrix::upper_triangular( + handle, + make_device_matrix_view(BBt.data(), l, l), + make_device_matrix_view(Uhat_dup.data(), l, l)); + if (use_jacobi) raft::linalg::eigJacobi( handle, Uhat_dup.data(), l, l, Uhat.data(), S_vec_tmp.data(), stream, tol, max_sweeps); else raft::linalg::eigDC(handle, Uhat_dup.data(), l, l, Uhat.data(), S_vec_tmp.data(), stream); raft::matrix::seqRoot(S_vec_tmp.data(), l, stream); - raft::matrix::sliceMatrix(S_vec_tmp.data(), - 1, - l, - S_vec, - 0, - p, - 1, - l, - stream); // Last k elements of S_vec - raft::matrix::colReverse(S_vec, 1, k, stream); + + auto S_vec_view = make_device_matrix_view(S_vec, 1, k); + raft::matrix::slice( + handle, + raft::make_device_matrix_view(S_vec_tmp.data(), 1, l), + S_vec_view, + raft::matrix::slice_coordinates(0, p, 1, l)); // Last k elements of S_vec + raft::matrix::col_reverse(handle, S_vec_view); // Merge step 14 & 15 by calculating U = Q*Uhat[:,(p+1):l] mxl * lxk = mxk if (gen_left_vec) { @@ -305,7 +309,7 @@ void rsvdFixedRank(raft::device_resources const& handle, alpha, beta, stream); - raft::matrix::colReverse(U, m, k, stream); + raft::matrix::col_reverse(handle, make_device_matrix_view(U, m, k)); } // Merge step 14 & 15 by calculating V = B^T Uhat[:,(p+1):l] * @@ -316,7 +320,9 @@ void rsvdFixedRank(raft::device_resources const& handle, rmm::device_uvector UhatSinv(l * k, stream); RAFT_CUDA_TRY(cudaMemsetAsync(UhatSinv.data(), 0, sizeof(math_t) * l * k, stream)); raft::matrix::reciprocal(S_vec_tmp.data(), l, stream); - raft::matrix::initializeDiagonalMatrix(S_vec_tmp.data() + p, Sinv.data(), k, k, stream); + raft::matrix::set_diagonal(handle, + make_device_vector_view(S_vec_tmp.data() + p, k), + make_device_matrix_view(Sinv.data(), k, k)); raft::linalg::gemm(handle, Uhat.data() + p * l, @@ -344,7 +350,7 @@ void rsvdFixedRank(raft::device_resources const& handle, alpha, beta, stream); - raft::matrix::colReverse(V, n, k, stream); + raft::matrix::col_reverse(handle, make_device_matrix_view(V, n, k)); } } } diff --git a/cpp/include/raft/linalg/detail/svd.cuh b/cpp/include/raft/linalg/detail/svd.cuh index 998bea5b1b..94cd9e2789 100644 --- a/cpp/include/raft/linalg/detail/svd.cuh +++ b/cpp/include/raft/linalg/detail/svd.cuh @@ -24,8 +24,10 @@ #include #include +#include #include -#include +#include +#include #include #include #include @@ -139,8 +141,10 @@ void svdEig(raft::device_resources const& handle, raft::linalg::eigDC(handle, in_cross_mult.data(), n_cols, n_cols, V, S, stream); - raft::matrix::colReverse(V, n_cols, n_cols, stream); - raft::matrix::rowReverse(S, n_cols, idx_t(1), stream); + raft::matrix::col_reverse(handle, + make_device_matrix_view(V, n_cols, n_cols)); + raft::matrix::row_reverse(handle, + make_device_matrix_view(S, n_cols, idx_t(1))); raft::matrix::seqRoot(S, S, alpha, n_cols, stream, true); @@ -285,15 +289,19 @@ bool evaluateSVDByL2Norm(raft::device_resources const& handle, RAFT_CUDA_TRY(cudaMemsetAsync(P_d.data(), 0, sizeof(math_t) * m * n, stream)); RAFT_CUDA_TRY(cudaMemsetAsync(S_mat.data(), 0, sizeof(math_t) * k * k, stream)); - raft::matrix::initializeDiagonalMatrix(S_vec, S_mat.data(), k, k, stream); + raft::matrix::set_diagonal(handle, + make_device_vector_view(S_vec, k), + make_device_matrix_view(S_mat.data(), k, k)); svdReconstruction(handle, U, S_mat.data(), V, P_d.data(), m, n, k, stream); // get norms of each - math_t normA = raft::matrix::getL2Norm(handle, A_d, m * n, stream); - math_t normU = raft::matrix::getL2Norm(handle, U, m * k, stream); - math_t normS = raft::matrix::getL2Norm(handle, S_mat.data(), k * k, stream); - math_t normV = raft::matrix::getL2Norm(handle, V, n * k, stream); - math_t normP = raft::matrix::getL2Norm(handle, P_d.data(), m * n, stream); + math_t normA = raft::matrix::l2_norm(handle, make_device_matrix_view(A_d, m, n)); + math_t normU = raft::matrix::l2_norm(handle, make_device_matrix_view(U, m, k)); + math_t normS = + raft::matrix::l2_norm(handle, make_device_matrix_view(S_mat.data(), k, k)); + math_t normV = raft::matrix::l2_norm(handle, make_device_matrix_view(V, n, k)); + math_t normP = + raft::matrix::l2_norm(handle, make_device_matrix_view(P_d.data(), m, n)); // calculate percent error const math_t alpha = 1.0, beta = -1.0; @@ -315,8 +323,9 @@ bool evaluateSVDByL2Norm(raft::device_resources const& handle, m, stream)); - math_t norm_A_minus_P = raft::matrix::getL2Norm(handle, A_minus_P.data(), m * n, stream); - math_t percent_error = 100.0 * norm_A_minus_P / normA; + math_t norm_A_minus_P = + raft::matrix::l2_norm(handle, make_device_matrix_view(A_minus_P.data(), m, n)); + math_t percent_error = 100.0 * norm_A_minus_P / normA; return (percent_error / 100.0 < tol); } diff --git a/cpp/include/raft/linalg/matrix_vector_op.cuh b/cpp/include/raft/linalg/matrix_vector_op.cuh index 6c65626ac5..e8833a2779 100644 --- a/cpp/include/raft/linalg/matrix_vector_op.cuh +++ b/cpp/include/raft/linalg/matrix_vector_op.cuh @@ -22,6 +22,7 @@ #include "linalg_types.hpp" #include +#include #include namespace raft { @@ -241,4 +242,4 @@ void matrix_vector_op(raft::device_resources const& handle, }; // end namespace linalg }; // end namespace raft -#endif \ No newline at end of file +#endif diff --git a/cpp/include/raft/matrix/copy.cuh b/cpp/include/raft/matrix/copy.cuh index 42d2562e5e..e4e5526e71 100644 --- a/cpp/include/raft/matrix/copy.cuh +++ b/cpp/include/raft/matrix/copy.cuh @@ -42,7 +42,7 @@ template void copy_rows(raft::device_resources const& handle, raft::device_matrix_view in, raft::device_matrix_view out, - raft::device_vector_view indices) + raft::device_vector_view indices) { RAFT_EXPECTS(in.extent(1) == out.extent(1), "Input and output matrices must have same number of columns"); @@ -58,6 +58,24 @@ void copy_rows(raft::device_resources const& handle, raft::is_row_major(in)); } +/** + * @brief copy matrix operation for row major matrices. + * @param[in] handle: raft handle + * @param[in] in: input matrix + * @param[out] out: output matrix + */ +template +void copy(raft::device_resources const& handle, + raft::device_matrix_view in, + raft::device_matrix_view out) +{ + RAFT_EXPECTS(in.extent(0) == out.extent(0) && in.extent(1) == out.extent(1), + "Input and output matrix shapes must match."); + + raft::copy_async( + out.data_handle(), in.data_handle(), in.extent(0) * out.extent(1), handle.get_stream()); +} + /** * @brief copy matrix operation for column major matrices. * @param[in] handle: raft handle diff --git a/cpp/include/raft/matrix/diagonal.cuh b/cpp/include/raft/matrix/diagonal.cuh index 22147e9f34..c7a3681983 100644 --- a/cpp/include/raft/matrix/diagonal.cuh +++ b/cpp/include/raft/matrix/diagonal.cuh @@ -17,8 +17,8 @@ #pragma once #include +#include #include -#include namespace raft::matrix { @@ -34,7 +34,7 @@ namespace raft::matrix { * @param[out] matrix: matrix of size n_rows x n_cols */ template -void set_diagonal(raft::device_resources const& handle, +void set_diagonal(raft::resources const& handle, raft::device_vector_view vec, raft::device_matrix_view matrix) { @@ -45,7 +45,7 @@ void set_diagonal(raft::device_resources const& handle, matrix.data_handle(), matrix.extent(0), matrix.extent(1), - handle.get_stream()); + resource::get_cuda_stream(handle)); } /** @@ -55,7 +55,7 @@ void set_diagonal(raft::device_resources const& handle, * @param[out] vec: vector of length k = min(n_rows, n_cols) */ template -void get_diagonal(raft::device_resources const& handle, +void get_diagonal(raft::resources const& handle, raft::device_matrix_view matrix, raft::device_vector_view vec) { @@ -65,7 +65,7 @@ void get_diagonal(raft::device_resources const& handle, matrix.data_handle(), matrix.extent(0), matrix.extent(1), - handle.get_stream()); + resource::get_cuda_stream(handle)); } /** @@ -74,14 +74,15 @@ void get_diagonal(raft::device_resources const& handle, * @param[inout] inout: square input matrix with size len x len */ template -void invert_diagonal(raft::device_resources const& handle, +void invert_diagonal(raft::resources const& handle, raft::device_matrix_view inout) { // TODO: Use get_diagonal for this to support rectangular RAFT_EXPECTS(inout.extent(0) == inout.extent(1), "Matrix must be square."); - detail::getDiagonalInverseMatrix(inout.data_handle(), inout.extent(0), handle.get_stream()); + detail::getDiagonalInverseMatrix( + inout.data_handle(), inout.extent(0), resource::get_cuda_stream(handle)); } /** @} */ // end of group matrix_diagonal -} // namespace raft::matrix \ No newline at end of file +} // namespace raft::matrix diff --git a/cpp/include/raft/matrix/init.cuh b/cpp/include/raft/matrix/init.cuh index ed2fb4d209..9611e044f4 100644 --- a/cpp/include/raft/matrix/init.cuh +++ b/cpp/include/raft/matrix/init.cuh @@ -20,7 +20,6 @@ #include #include #include -#include namespace raft::matrix { diff --git a/cpp/include/raft/matrix/linewise_op.cuh b/cpp/include/raft/matrix/linewise_op.cuh index 33de112a35..f8e3555d9d 100644 --- a/cpp/include/raft/matrix/linewise_op.cuh +++ b/cpp/include/raft/matrix/linewise_op.cuh @@ -17,8 +17,8 @@ #pragma once #include -#include -#include +#include +#include namespace raft::matrix { diff --git a/cpp/include/raft/matrix/print.cuh b/cpp/include/raft/matrix/print.cuh index 6a4bfbdd01..f2c2653211 100644 --- a/cpp/include/raft/matrix/print.cuh +++ b/cpp/include/raft/matrix/print.cuh @@ -19,7 +19,6 @@ #include #include #include -#include #include namespace raft::matrix { diff --git a/cpp/include/raft/matrix/sign_flip.cuh b/cpp/include/raft/matrix/sign_flip.cuh index d069c55880..93962fb67d 100644 --- a/cpp/include/raft/matrix/sign_flip.cuh +++ b/cpp/include/raft/matrix/sign_flip.cuh @@ -18,7 +18,6 @@ #include #include -#include namespace raft::matrix { diff --git a/cpp/include/raft/matrix/slice.cuh b/cpp/include/raft/matrix/slice.cuh index bb92b2b86f..071a10a847 100644 --- a/cpp/include/raft/matrix/slice.cuh +++ b/cpp/include/raft/matrix/slice.cuh @@ -76,4 +76,4 @@ void slice(raft::device_resources const& handle, /** @} */ // end group matrix_slice -} // namespace raft::matrix \ No newline at end of file +} // namespace raft::matrix diff --git a/cpp/include/raft/matrix/sqrt.cuh b/cpp/include/raft/matrix/sqrt.cuh index 9729f9b3d5..309ae3452f 100644 --- a/cpp/include/raft/matrix/sqrt.cuh +++ b/cpp/include/raft/matrix/sqrt.cuh @@ -19,7 +19,6 @@ #include #include #include -#include namespace raft::matrix { diff --git a/cpp/include/raft/matrix/triangular.cuh b/cpp/include/raft/matrix/triangular.cuh index 3c60cc362f..0c89140046 100644 --- a/cpp/include/raft/matrix/triangular.cuh +++ b/cpp/include/raft/matrix/triangular.cuh @@ -17,6 +17,7 @@ #pragma once #include +#include #include namespace raft::matrix { @@ -33,17 +34,19 @@ namespace raft::matrix { * @param[out] dst: output matrix with a size of kxk, k = min(n_rows, n_cols) */ template -void upper_triangular(raft::device_resources const& handle, +void upper_triangular(raft::resources const& handle, raft::device_matrix_view src, raft::device_matrix_view dst) { auto k = std::min(src.extent(0), src.extent(1)); RAFT_EXPECTS(k == dst.extent(0) && k == dst.extent(1), "dst should be of size kxk, k = min(n_rows, n_cols)"); - detail::copyUpperTriangular( - src.data_handle(), dst.data_handle(), src.extent(0), src.extent(1), handle.get_stream()); + detail::copyUpperTriangular(src.data_handle(), + dst.data_handle(), + src.extent(0), + src.extent(1), + resource::get_cuda_stream(handle)); } - /** @} */ // end group matrix_triangular -} // namespace raft::matrix \ No newline at end of file +} // namespace raft::matrix diff --git a/cpp/include/raft/neighbors/detail/ivf_pq_build.cuh b/cpp/include/raft/neighbors/detail/ivf_pq_build.cuh index b17b3a3559..53d8823eea 100644 --- a/cpp/include/raft/neighbors/detail/ivf_pq_build.cuh +++ b/cpp/include/raft/neighbors/detail/ivf_pq_build.cuh @@ -181,8 +181,12 @@ void select_residuals(raft::device_resources const& handle, dataset, utils::mapping{}); raft::matrix::gather(mapping_itr, (IdxT)dim, n_rows, row_ids, n_rows, tmp.data(), stream); - raft::matrix::linewiseOp( - tmp.data(), tmp.data(), IdxT(dim), n_rows, true, raft::sub_op{}, stream, center); + raft::matrix::linewise_op(handle, + make_device_matrix_view(tmp.data(), n_rows, dim), + make_device_matrix_view(tmp.data(), n_rows, dim), + true, + raft::sub_op{}, + make_device_vector_view(center, dim)); float alpha = 1.0; float beta = 0.0; diff --git a/cpp/include/raft/neighbors/refine-inl.cuh b/cpp/include/raft/neighbors/refine-inl.cuh index 4243d7e723..2c4dfb422e 100644 --- a/cpp/include/raft/neighbors/refine-inl.cuh +++ b/cpp/include/raft/neighbors/refine-inl.cuh @@ -19,7 +19,6 @@ #include #include #include -#include #include #include diff --git a/cpp/include/raft/random/detail/make_regression.cuh b/cpp/include/raft/random/detail/make_regression.cuh index 1715dcbe81..aec1a15f84 100644 --- a/cpp/include/raft/random/detail/make_regression.cuh +++ b/cpp/include/raft/random/detail/make_regression.cuh @@ -29,7 +29,7 @@ #include #include #include -#include +#include #include #include #include @@ -83,7 +83,10 @@ static void _make_low_rank_matrix(raft::resources const& handle, RAFT_CUDA_TRY(cudaPeekAtLastError()); rmm::device_uvector singular_mat(n * n, stream); RAFT_CUDA_TRY(cudaMemsetAsync(singular_mat.data(), 0, n * n * sizeof(DataT), stream)); - raft::matrix::initializeDiagonalMatrix(singular_vec.data(), singular_mat.data(), n, n, stream); + + raft::matrix::set_diagonal(handle, + make_device_vector_view(singular_vec.data(), n), + make_device_matrix_view(singular_mat.data(), n, n)); // Generate the column-major matrix rmm::device_uvector temp_q0s(n_rows * n, stream); diff --git a/cpp/include/raft/sparse/neighbors/detail/knn.cuh b/cpp/include/raft/sparse/neighbors/detail/knn.cuh index 6649c10c47..527fc14208 100644 --- a/cpp/include/raft/sparse/neighbors/detail/knn.cuh +++ b/cpp/include/raft/sparse/neighbors/detail/knn.cuh @@ -20,7 +20,6 @@ #include #include -#include #include #include diff --git a/cpp/include/raft/spatial/knn/detail/ball_cover.cuh b/cpp/include/raft/spatial/knn/detail/ball_cover.cuh index c8fc6eefda..a58847ee41 100644 --- a/cpp/include/raft/spatial/knn/detail/ball_cover.cuh +++ b/cpp/include/raft/spatial/knn/detail/ball_cover.cuh @@ -30,7 +30,7 @@ #include -#include +#include #include #include #include @@ -94,14 +94,16 @@ void sample_landmarks(raft::device_resources const& handle, (value_idx)index.n_landmarks, (value_idx)index.m); - raft::matrix::copyRows(index.get_X().data_handle(), - index.m, - index.n, - index.get_R().data_handle(), - R_1nn_cols2.data(), - index.n_landmarks, - handle.get_stream(), - true); + // index.get_X() returns the wrong indextype (uint32_t where we need value_idx), so need to + // create new device_matrix_view here + auto x = index.get_X(); + auto r = index.get_R(); + + raft::matrix::copy_rows( + handle, + make_device_matrix_view(x.data_handle(), x.extent(0), x.extent(1)), + make_device_matrix_view(r.data_handle(), r.extent(0), r.extent(1)), + make_device_vector_view(R_1nn_cols2.data(), index.n_landmarks)); } /** diff --git a/cpp/test/linalg/svd.cu b/cpp/test/linalg/svd.cu index c780476a5f..9907172956 100644 --- a/cpp/test/linalg/svd.cu +++ b/cpp/test/linalg/svd.cu @@ -18,7 +18,6 @@ #include #include #include -#include #include #include diff --git a/cpp/test/matrix/matrix.cu b/cpp/test/matrix/matrix.cu index 10105203f7..07ab3c5ce4 100644 --- a/cpp/test/matrix/matrix.cu +++ b/cpp/test/matrix/matrix.cu @@ -143,7 +143,7 @@ class MatrixCopyRowsTest : public ::testing::Test { output.data(), n_selected, n_cols); auto indices_view = - raft::make_device_vector_view(indices.data(), n_selected); + raft::make_device_vector_view(indices.data(), n_selected); raft::matrix::copy_rows(handle, input_view, output_view, indices_view); diff --git a/cpp/test/neighbors/ann_utils.cuh b/cpp/test/neighbors/ann_utils.cuh index 438c56da21..67df5f2abe 100644 --- a/cpp/test/neighbors/ann_utils.cuh +++ b/cpp/test/neighbors/ann_utils.cuh @@ -18,8 +18,8 @@ #include // raft::make_device_matrix #include +#include #include -#include #include #include @@ -188,8 +188,11 @@ auto eval_distances(raft::device_resources const& handle, auto y = raft::make_device_matrix(handle, k, n_cols); auto naive_dist = raft::make_device_matrix(handle, 1, k); - raft::matrix::copyRows( - x, k, n_cols, y.data_handle(), neighbors + i * k, k, handle.get_stream(), true); + raft::matrix::copy_rows( + handle, + make_device_matrix_view(x, k, n_cols), + y.view(), + make_device_vector_view(neighbors + i * k, k)); dim3 block_dim(16, 32, 1); auto grid_y =