diff --git a/cpp/include/raft/sparse/convert/dense.cuh b/cpp/include/raft/sparse/convert/dense.cuh index 02abb590ed..dc9da25df2 100644 --- a/cpp/include/raft/sparse/convert/dense.cuh +++ b/cpp/include/raft/sparse/convert/dense.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 */ #ifndef __DENSE_H @@ -8,6 +8,8 @@ #pragma once #include +#include +#include #include namespace raft { @@ -15,39 +17,28 @@ namespace sparse { namespace convert { /** - * Convert CSR arrays to a dense matrix in either row- - * or column-major format. A custom kernel is used when - * row-major output is desired since cusparse does not - * output row-major. - * @tparam value_idx : data type of the CSR index arrays - * @tparam value_t : data type of the CSR value array - * @param[in] handle : cusparse handle for conversion - * @param[in] nrows : number of rows in CSR - * @param[in] ncols : number of columns in CSR - * @param[in] nnz : number of nonzeros in CSR - * @param[in] csr_indptr : CSR row index pointer array - * @param[in] csr_indices : CSR column indices array - * @param[in] csr_data : CSR data array - * @param[in] lda : Leading dimension (used for col-major only) - * @param[out] out : Dense output array of size nrows * ncols - * @param[in] stream : Cuda stream for ordering events - * @param[in] row_major : Is row-major output desired? + * Convert a sparse matrix view to a dense matrix view. + * + * Supports both COO and CSR sparse matrix views and row- or column-major dense output. + * + * @param[in] handle RAFT resources + * @param[in] sparse Sparse COO or CSR matrix view + * @param[out] dense Dense matrix view */ -template -void csr_to_dense(cusparseHandle_t handle, - value_idx nrows, - value_idx ncols, - value_idx nnz, - const value_idx* csr_indptr, - const value_idx* csr_indices, - const value_t* csr_data, - value_idx lda, - value_t* out, - cudaStream_t stream, - bool row_major = true) +template +void sparse_to_dense(raft::resources const& handle, + SparseMatrixViewType sparse, + raft::device_matrix_view dense) { - detail::csr_to_dense( - handle, nrows, ncols, nnz, csr_indptr, csr_indices, csr_data, lda, out, stream, row_major); + auto structure = sparse.structure_view(); + RAFT_EXPECTS(dense.extent(0) == static_cast(structure.get_n_rows()) && + dense.extent(1) == static_cast(structure.get_n_cols()), + "Sparse and dense matrix dimensions must match"); + + detail::sparse_to_dense(handle, sparse, dense); } }; // end NAMESPACE convert diff --git a/cpp/include/raft/sparse/convert/detail/dense.cuh b/cpp/include/raft/sparse/convert/detail/dense.cuh index fa22902e89..31ee3e027e 100644 --- a/cpp/include/raft/sparse/convert/detail/dense.cuh +++ b/cpp/include/raft/sparse/convert/detail/dense.cuh @@ -6,134 +6,52 @@ #pragma once #include +#include +#include +#include +#include +#include #include -#include -#include -#include -#include - -#include +#include #include -#include -#include #include -#include - -#include -#include namespace raft { namespace sparse { namespace convert { namespace detail { -template -RAFT_KERNEL csr_to_dense_warp_per_row_kernel( - int n_cols, const value_t* csrVal, const int* csrRowPtr, const int* csrColInd, value_t* a) +template +void sparse_to_dense(raft::resources const& handle, + SparseMatrixViewType sparse, + raft::device_matrix_view dense) { - int row = blockIdx.x; - int tid = threadIdx.x; - - int colStart = csrRowPtr[row]; - int colEnd = csrRowPtr[row + 1]; - int rowNnz = colEnd - colStart; - - for (int i = tid; i < rowNnz; i += blockDim.x) { - int colIdx = colStart + i; - if (colIdx < colEnd) { - int col = csrColInd[colIdx]; - a[row * n_cols + col] = csrVal[colIdx]; - } - } -} - -/** - * Convert CSR arrays to a dense matrix in either row- - * or column-major format. A custom kernel is used when - * row-major output is desired since cusparse does not - * output row-major. - * @tparam value_idx : data type of the CSR index arrays - * @tparam value_t : data type of the CSR value array - * @param[in] handle : cusparse handle for conversion - * @param[in] nrows : number of rows in CSR - * @param[in] ncols : number of columns in CSR - * @param[in] nnz : the number of nonzeros in CSR - * @param[in] csr_indptr : CSR row index pointer array - * @param[in] csr_indices : CSR column indices array - * @param[in] csr_data : CSR data array - * @param[in] lda : Leading dimension (used for col-major only) - * @param[out] out : Dense output array of size nrows * ncols - * @param[in] stream : Cuda stream for ordering events - * @param[in] row_major : Is row-major output desired? - */ -template -void csr_to_dense(cusparseHandle_t handle, - value_idx nrows, - value_idx ncols, - value_idx nnz, - const value_idx* csr_indptr, - const value_idx* csr_indices, - const value_t* csr_data, - value_idx lda, - value_t* out, - cudaStream_t stream, - bool row_major = true) -{ - if (!row_major) { - /** - * If we need col-major, use cusparse. - */ - cusparseMatDescr_t out_mat; - RAFT_CUSPARSE_TRY(cusparseCreateMatDescr(&out_mat)); - RAFT_CUSPARSE_TRY(cusparseSetMatIndexBase(out_mat, CUSPARSE_INDEX_BASE_ZERO)); - RAFT_CUSPARSE_TRY(cusparseSetMatType(out_mat, CUSPARSE_MATRIX_TYPE_GENERAL)); - - size_t buffer_size; - RAFT_CUSPARSE_TRY(raft::sparse::detail::cusparsecsr2dense_buffersize(handle, - nrows, - ncols, - nnz, - out_mat, - csr_data, - csr_indptr, - csr_indices, - out, - lda, - &buffer_size, - stream)); - - rmm::device_uvector buffer(buffer_size, stream); - - RAFT_CUSPARSE_TRY(raft::sparse::detail::cusparsecsr2dense(handle, - nrows, - ncols, - nnz, - out_mat, - csr_data, - csr_indptr, - csr_indices, - out, - lda, - buffer.data(), - stream)); - - RAFT_CUSPARSE_TRY_NO_THROW(cusparseDestroyMatDescr(out_mat)); - - } else { - int blockdim = block_dim(ncols); - RAFT_CUDA_TRY(cudaMemsetAsync(out, 0, nrows * ncols * sizeof(value_t), stream)); - raft::launch_kernel(stream, - nrows, - blockdim, - csr_to_dense_warp_per_row_kernel, - ncols, - csr_data, - csr_indptr, - csr_indices, - out); - } + auto stream = raft::resource::get_cuda_stream(handle); + auto cusparse_handle = raft::resource::get_cusparse_handle(handle); + auto sparse_descriptor = raft::sparse::linalg::detail::create_descriptor(sparse); + auto dense_descriptor = raft::sparse::linalg::detail::create_descriptor(dense); + + RAFT_CUSPARSE_TRY(cusparseSetStream(cusparse_handle, stream)); + std::size_t buffer_size; + RAFT_CUSPARSE_TRY(cusparseSparseToDense_bufferSize(cusparse_handle, + sparse_descriptor, + dense_descriptor, + CUSPARSE_SPARSETODENSE_ALG_DEFAULT, + &buffer_size)); + auto buffer = raft::make_device_vector(handle, buffer_size); + RAFT_CUSPARSE_TRY(cusparseSparseToDense(cusparse_handle, + sparse_descriptor, + dense_descriptor, + CUSPARSE_SPARSETODENSE_ALG_DEFAULT, + buffer.data_handle())); + + RAFT_CUSPARSE_TRY_NO_THROW(cusparseDestroySpMat(sparse_descriptor)); + RAFT_CUSPARSE_TRY_NO_THROW(cusparseDestroyDnMat(dense_descriptor)); } }; // namespace detail diff --git a/cpp/include/raft/sparse/solver/detail/lanczos.cuh b/cpp/include/raft/sparse/solver/detail/lanczos.cuh index b7c24124f7..0f05644fc1 100644 --- a/cpp/include/raft/sparse/solver/detail/lanczos.cuh +++ b/cpp/include/raft/sparse/solver/detail/lanczos.cuh @@ -18,6 +18,8 @@ #include #include #include +#include +#include #include #include #include @@ -44,6 +46,7 @@ #include #include #include +#include #include #include #include @@ -129,9 +132,7 @@ RAFT_KERNEL kernel_clamp_down_vector(T* vec, T threshold, int size) template void lanczos_solve_ritz( raft::resources const& handle, - raft::device_matrix_view alpha, - raft::device_matrix_view beta, - std::optional> beta_k, + raft::device_matrix_view matrix, IndexTypeT k, LANCZOS_WHICH which, int ncv, @@ -145,46 +146,7 @@ void lanczos_solve_ritz( { auto stream = resource::get_cuda_stream(handle); - ValueTypeT zero = 0; - auto triangular_matrix = - raft::make_device_matrix(handle, ncv, ncv); - raft::matrix::fill(handle, triangular_matrix.view(), zero); - - raft::device_vector_view alphaVec = - raft::make_device_vector_view(alpha.data_handle(), ncv); - raft::matrix::set_diagonal(handle, alphaVec, triangular_matrix.view()); - - // raft::matrix::initializeDiagonalMatrix( - // alpha.data_handle(), triangular_matrix.data_handle(), ncv, ncv, stream); - - int blockSize = 256; - int numBlocks = raft::div_rounding_up_safe(ncv, blockSize); - raft::launch_kernel(handle, - numBlocks, - blockSize, - kernel_triangular_populate, - triangular_matrix.data_handle(), - beta.data_handle(), - ncv); - - if (beta_k) { - int threadsPerBlock = 256; - int blocksPerGrid = raft::div_rounding_up_safe(k, threadsPerBlock); - raft::launch_kernel(handle, - blocksPerGrid, - threadsPerBlock, - kernel_triangular_beta_k, - triangular_matrix.data_handle(), - beta_k.value().data_handle(), - k, - ncv); - } - - auto triangular_matrix_view = - raft::make_device_matrix_view( - triangular_matrix.data_handle(), ncv, ncv); - - raft::linalg::eig_dc(handle, triangular_matrix_view, eigenvectors, eigenvalues); + raft::linalg::eig_dc(handle, matrix, eigenvectors, eigenvalues); IndexTypeT nEigVecs = k; @@ -256,6 +218,76 @@ void lanczos_solve_ritz( } } +template +void lanczos_solve_ritz( + raft::resources const& handle, + raft::device_matrix_view alpha, + raft::device_matrix_view beta, + std::optional> beta_k, + IndexTypeT k, + LANCZOS_WHICH which, + int ncv, + raft::device_matrix_view eigenvectors, + raft::device_vector_view eigenvalues, + raft::device_matrix_view& eigenvectors_k, + raft::device_vector_view& eigenvalues_k, + raft::device_matrix_view& eigenvectors_k_slice, + raft::device_vector_view sm_eigenvalues, + raft::device_matrix_view sm_eigenvectors) +{ + ValueTypeT zero = 0; + auto triangular_matrix = + raft::make_device_matrix(handle, ncv, ncv); + raft::matrix::fill(handle, triangular_matrix.view(), zero); + + raft::device_vector_view alphaVec = + raft::make_device_vector_view(alpha.data_handle(), ncv); + raft::matrix::set_diagonal(handle, alphaVec, triangular_matrix.view()); + + // raft::matrix::initializeDiagonalMatrix( + // alpha.data_handle(), triangular_matrix.data_handle(), ncv, ncv, stream); + + int blockSize = 256; + int numBlocks = raft::div_rounding_up_safe(ncv, blockSize); + raft::launch_kernel(handle, + numBlocks, + blockSize, + kernel_triangular_populate, + triangular_matrix.data_handle(), + beta.data_handle(), + ncv); + + if (beta_k) { + int threadsPerBlock = 256; + int blocksPerGrid = raft::div_rounding_up_safe(k, threadsPerBlock); + raft::launch_kernel(handle, + blocksPerGrid, + threadsPerBlock, + kernel_triangular_beta_k, + triangular_matrix.data_handle(), + beta_k.value().data_handle(), + k, + ncv); + } + + auto triangular_matrix_view = + raft::make_device_matrix_view( + triangular_matrix.data_handle(), ncv, ncv); + + lanczos_solve_ritz(handle, + triangular_matrix_view, + k, + which, + ncv, + eigenvectors, + eigenvalues, + eigenvectors_k, + eigenvalues_k, + eigenvectors_k_slice, + sm_eigenvalues, + sm_eigenvectors); +} + template void lanczos_aux(raft::resources const& handle, AType A, @@ -797,6 +829,46 @@ auto lanczos_compute_eigenpairs( raft::device_vector_view eigenvalues, raft::device_matrix_view eigenvectors) -> int { + auto stream = resource::get_cuda_stream(handle); + auto n = A.structure_view().get_n_rows(); + RAFT_EXPECTS(config.n_components > 0 && config.n_components < n, + "n_components must satisfy 0 < n_components < n"); + + if (n <= 3) { + auto k = config.n_components; + auto dense_matrix = + raft::make_device_matrix(handle, n, n); + auto all_eigenvectors = + raft::make_device_matrix(handle, n, n); + auto all_eigenvalues = raft::make_device_vector(handle, n); + auto sm_eigenvectors = + raft::make_device_matrix(handle, n, k); + auto sm_eigenvalues = raft::make_device_vector(handle, k); + raft::device_matrix_view eigenvectors_k; + raft::device_vector_view eigenvalues_k; + raft::device_matrix_view eigenvectors_k_slice; + + raft::sparse::convert::sparse_to_dense(handle, A, dense_matrix.view()); + lanczos_solve_ritz(handle, + raft::make_const_mdspan(dense_matrix.view()), + k, + config.which, + n, + all_eigenvectors.view(), + all_eigenvalues.view(), + eigenvectors_k, + eigenvalues_k, + eigenvectors_k_slice, + sm_eigenvalues.view(), + sm_eigenvectors.view()); + raft::copy(eigenvalues.data_handle(), eigenvalues_k.data_handle(), k, stream); + raft::copy(eigenvectors.data_handle(), eigenvectors_k.data_handle(), n * k, stream); + return 0; + } + + RAFT_EXPECTS(config.ncv > config.n_components + 1 && config.ncv < n, + "ncv must satisfy n_components + 1 < ncv < n"); + if (v0.has_value()) { return lanczos_smallest(handle, A, @@ -811,7 +883,6 @@ auto lanczos_compute_eigenpairs( config.seed); } else { // Handle the optional v0 initial Lanczos vector if nullopt is used - auto n = A.structure_view().get_n_rows(); auto temp_v0 = raft::make_device_vector(handle, n); uint64_t seed = config.seed.value_or(std::random_device{}()); raft::random::RngState rng_state(seed); diff --git a/cpp/tests/sparse/csr_to_dense.cu b/cpp/tests/sparse/csr_to_dense.cu index 4d4690e5ad..fd2748a866 100644 --- a/cpp/tests/sparse/csr_to_dense.cu +++ b/cpp/tests/sparse/csr_to_dense.cu @@ -1,19 +1,18 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2018-2024, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2018-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ #include "../test_utils.cuh" +#include #include #include #include -#include #include #include -#include #include namespace raft { @@ -79,24 +78,17 @@ class CSRToDenseTest : public ::testing::TestWithParam( + indptr.data(), indices.data(), params.nrows, params.ncols, params.nnz); + auto csr = raft::make_device_csr_matrix_view( + data.data(), structure); + auto dense = raft::make_device_matrix_view( + out.data(), params.nrows, params.ncols); + convert::sparse_to_dense(raft_handle, csr, dense); RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); - RAFT_CUSPARSE_TRY(cusparseDestroy(handle)); } void compare() @@ -109,8 +101,6 @@ class CSRToDenseTest : public ::testing::TestWithParam indptr, indices; rmm::device_uvector data; diff --git a/cpp/tests/util/preprocess_utils.cu b/cpp/tests/util/preprocess_utils.cu index c71316d9d9..0a0929ccca 100644 --- a/cpp/tests/util/preprocess_utils.cu +++ b/cpp/tests/util/preprocess_utils.cu @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2025-2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * SPDX-License-Identifier: Apache-2.0 */ @@ -135,32 +135,14 @@ void calc_tfidf_bm25(raft::resources& handle, cudaStream_t stream = raft::resource::get_cuda_stream(handle); int num_rows = csr_in.structure_view().get_n_rows(); int num_cols = csr_in.structure_view().get_n_cols(); - int rows_size = csr_in.structure_view().get_indptr().size(); - int cols_size = csr_in.structure_view().get_indices().size(); - int elements_size = csr_in.get_elements().size(); - - auto indptr = raft::make_device_vector_view( - csr_in.structure_view().get_indptr().data(), rows_size); - auto indices = raft::make_device_vector_view( - csr_in.structure_view().get_indices().data(), cols_size); - auto values = - raft::make_device_vector_view(csr_in.get_elements().data(), elements_size); + auto dense_values = raft::make_device_vector(handle, num_rows * num_cols); - cusparseHandle_t cu_handle; - RAFT_CUSPARSE_TRY(cusparseCreate(&cu_handle)); - - raft::sparse::convert::csr_to_dense(cu_handle, - num_rows, - num_cols, - elements_size, - indptr.data_handle(), - indices.data_handle(), - values.data_handle(), - num_rows, - dense_values.data_handle(), - stream, - true); + raft::sparse::convert::sparse_to_dense( + handle, + csr_in, + raft::make_device_matrix_view( + dense_values.data_handle(), num_rows, num_cols)); RAFT_CUDA_TRY(cudaStreamSynchronize(stream)); preproc(handle, dense_values.view(), results, num_rows, num_cols, tf_idf);