Remove raft/matrix/matrix.cuh includes

cpp/include/raft/cluster/detail/kmeans_balanced.cuh

@@ -37,7 +37,6 @@
 #include <raft/linalg/unary_op.cuh>
 #include <raft/matrix/argmin.cuh>
 #include <raft/matrix/gather.cuh>
-#include <raft/matrix/matrix.cuh>
 #include <raft/util/cuda_utils.cuh>
 #include <raft/util/device_atomics.cuh>
 #include <raft/util/integer_utils.hpp>

cpp/include/raft/linalg/detail/eig.cuh

@@ -19,7 +19,7 @@
 #include "cusolver_wrappers.hpp"
 #include <cuda_runtime_api.h>
 #include <raft/core/device_resources.hpp>
-#include <raft/matrix/matrix.cuh>
+#include <raft/matrix/copy.cuh>
 #include <raft/util/cudart_utils.hpp>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
@@ -52,7 +52,9 @@ void eigDC_legacy(raft::device_resources const& handle,
   rmm::device_uvector<math_t> d_work(lwork, stream);
   rmm::device_scalar<int> d_dev_info(stream);
 
-  raft::matrix::copy(in, eig_vectors, n_rows, n_cols, stream);
+  raft::matrix::copy(handle,
+                     make_device_matrix_view<const math_t>(in, n_rows, n_cols),
+                     make_device_matrix_view<math_t>(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<int> d_dev_info(stream);
   std::vector<math_t> 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<const math_t>(in, n_rows, n_cols),
+                     make_device_matrix_view<math_t>(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<const math_t>(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<const math_t, size_t, col_major>(in, n_rows, n_eig_vals),
+      make_device_matrix_view<math_t, size_t, col_major>(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<const math_t, size_t, col_major>(
+        d_eig_vectors.data(), n_rows, n_eig_vals),
+      make_device_matrix_view<math_t, size_t, col_major>(eig_vectors, n_rows, n_eig_vals));
   }
 }
 
@@ -259,7 +271,9 @@ void eigJacobi(raft::device_resources const& handle,
   rmm::device_uvector<math_t> d_work(lwork, stream);
   rmm::device_scalar<int> dev_info(stream);
 
-  raft::matrix::copy(in, eig_vectors, n_rows, n_cols, stream);
+  raft::matrix::copy(handle,
+                     make_device_matrix_view<const math_t>(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
+}  // namespace raft

cpp/include/raft/linalg/detail/lstsq.cuh

@@ -28,7 +28,6 @@
 #include <raft/linalg/svd.cuh>
 #include <raft/linalg/transpose.cuh>
 #include <raft/matrix/math.cuh>
-#include <raft/matrix/matrix.cuh>
 #include <raft/util/cudart_utils.hpp>
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_scalar.hpp>

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 <raft/matrix/matrix.cuh>
+#include <raft/matrix/linewise_op.cuh>
 
 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<MatT, IdxType, row_major, Lambda>(
+      handle,
+      make_device_matrix_view<const MatT, IdxType, row_major>(matrix, N, D),
+      make_device_matrix_view<MatT, IdxType, row_major>(out, N, D),
+      along_lines,
+      op,
+      make_device_vector_view<const VecT, IdxType>(vec, bcastAlongRows ? N : D));
+  } else {
+    matrix::linewise_op<MatT, IdxType, col_major, Lambda>(
+      handle,
+      make_device_matrix_view<const MatT, IdxType, col_major>(matrix, N, D),
+      make_device_matrix_view<MatT, IdxType, col_major>(out, N, D),
+      along_lines,
+      op,
+      make_device_vector_view<const VecT, IdxType>(vec, bcastAlongRows ? N : D));
+  }
 }
 
 template <typename MatT,
@@ -56,10 +72,27 @@ 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, vec1, vec2);
+  raft::device_resources handle(stream);
+  bool along_lines = rowMajor == bcastAlongRows;
+  if (rowMajor) {
+    matrix::linewise_op<MatT, IdxType, row_major, Lambda>(
+      handle,
+      make_device_matrix_view<const MatT, IdxType, row_major>(matrix, N, D),
+      make_device_matrix_view<MatT, IdxType, row_major>(out, N, D),
+      along_lines,
+      op,
+      make_device_vector_view<const Vec1T, IdxType>(vec1, bcastAlongRows ? N : D),
+      make_device_vector_view<const Vec2T, IdxType>(vec2, bcastAlongRows ? N : D));
+  } else {
+    matrix::linewise_op<MatT, IdxType, col_major, Lambda>(
+      handle,
+      make_device_matrix_view<const MatT, IdxType, col_major>(matrix, N, D),
+      make_device_matrix_view<MatT, IdxType, col_major>(out, N, D),
+      along_lines,
+      op,
+      make_device_vector_view<const Vec1T, IdxType>(vec1, bcastAlongRows ? N : D),
+      make_device_vector_view<const Vec2T, IdxType>(vec2, bcastAlongRows ? N : D));
+  }
 }
 
 };  // end namespace detail

cpp/include/raft/linalg/detail/qr.cuh

@@ -20,7 +20,7 @@
 #include "cusolver_wrappers.hpp"
 #include <raft/core/resource/cusolver_dn_handle.hpp>
 #include <raft/core/resources.hpp>
-#include <raft/matrix/matrix.cuh>
+#include <raft/matrix/triangular.cuh>
 #include <rmm/device_scalar.hpp>
 #include <rmm/device_uvector.hpp>
 
@@ -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<math_t, int>(
+    handle,
+    make_device_matrix_view<const math_t, int, col_major>(R_full.data(), m, n),
+    make_device_matrix_view<math_t, int, col_major>(R, std::min(m, n), std::min(m, n)));
 
   RAFT_CUDA_TRY(
     cudaMemcpyAsync(Q, R_full.data(), sizeof(math_t) * m * n, cudaMemcpyDeviceToDevice, stream));

cpp/include/raft/linalg/detail/rsvd.cuh

@@ -21,8 +21,11 @@
 #include <raft/linalg/qr.cuh>
 #include <raft/linalg/svd.cuh>
 #include <raft/linalg/transpose.cuh>
+#include <raft/matrix/diagonal.cuh>
 #include <raft/matrix/math.cuh>
-#include <raft/matrix/matrix.cuh>
+#include <raft/matrix/reverse.cuh>
+#include <raft/matrix/slice.cuh>
+#include <raft/matrix/triangular.cuh>
 #include <raft/random/rng.cuh>
 #include <raft/util/cuda_utils.cuh>
 
@@ -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<const math_t, int, col_major>(S_vec_tmp.data(), 1, l),
+      make_device_matrix_view<math_t, int, col_major>(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<math_t> 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<const math_t, int, col_major>(BBt.data(), l, l),
+      make_device_matrix_view<math_t, int, col_major>(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<math_t, int, col_major>(S_vec, 1, k);
+    raft::matrix::slice(
+      handle,
+      raft::make_device_matrix_view<const math_t, int, col_major>(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<math_t, int, col_major>(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<math_t> 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<const math_t>(S_vec_tmp.data() + p, k),
+                                 make_device_matrix_view<math_t>(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<math_t, int, col_major>(V, n, k));
     }
   }
 }

cpp/include/raft/linalg/detail/svd.cuh

@@ -24,8 +24,10 @@
 
 #include <raft/common/nvtx.hpp>
 #include <raft/core/device_resources.hpp>
+#include <raft/matrix/diagonal.cuh>
 #include <raft/matrix/math.cuh>
-#include <raft/matrix/matrix.cuh>
+#include <raft/matrix/norm.cuh>
+#include <raft/matrix/reverse.cuh>
 #include <raft/util/cuda_utils.cuh>
 #include <raft/util/cudart_utils.hpp>
 #include <rmm/device_scalar.hpp>
@@ -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<math_t, idx_t, col_major>(V, n_cols, n_cols));
+  raft::matrix::row_reverse(handle,
+                            make_device_matrix_view<math_t, idx_t, col_major>(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<const math_t>(S_vec, k),
+                             make_device_matrix_view<math_t>(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<const math_t>(A_d, m, n));
+  math_t normU = raft::matrix::l2_norm(handle, make_device_matrix_view<const math_t>(U, m, k));
+  math_t normS =
+    raft::matrix::l2_norm(handle, make_device_matrix_view<const math_t>(S_mat.data(), k, k));
+  math_t normV = raft::matrix::l2_norm(handle, make_device_matrix_view<const math_t>(V, n, k));
+  math_t normP =
+    raft::matrix::l2_norm(handle, make_device_matrix_view<const math_t>(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<const math_t>(A_minus_P.data(), m, n));
+  math_t percent_error = 100.0 * norm_A_minus_P / normA;
   return (percent_error / 100.0 < tol);
 }
 

cpp/include/raft/linalg/matrix_vector_op.cuh

@@ -22,6 +22,7 @@
 #include "linalg_types.hpp"
 
 #include <raft/core/device_mdspan.hpp>
+#include <raft/core/device_resources.hpp>
 #include <raft/util/input_validation.hpp>
 
 namespace raft {
@@ -241,4 +242,4 @@ void matrix_vector_op(raft::device_resources const& handle,
 };         // end namespace linalg
 };         // end namespace raft
 
-#endif
+#endif