diff --git a/cpp/include/raft/sparse/linalg/detail/cusparse_utils.hpp b/cpp/include/raft/sparse/linalg/detail/cusparse_utils.hpp index c10c0de426..97ac7c45f4 100644 --- a/cpp/include/raft/sparse/linalg/detail/cusparse_utils.hpp +++ b/cpp/include/raft/sparse/linalg/detail/cusparse_utils.hpp @@ -30,7 +30,25 @@ namespace linalg { namespace detail { /** - * @brief create a cuSparse dense descriptor + * @brief create a cuSparse dense descriptor for a vector + * @tparam ValueType Data type of vector_view (float/double) + * @tparam IndexType Type of vector_view + * @param[in] vector_view input raft::device_vector_view + * @returns dense vector descriptor to be used by cuSparse API + */ +template +cusparseDnVecDescr_t create_descriptor(raft::device_vector_view vector_view) +{ + cusparseDnVecDescr_t descr; + RAFT_CUSPARSE_TRY(raft::sparse::detail::cusparsecreatednvec( + &descr, + vector_view.extent(0), + const_cast*>(vector_view.data_handle()))); + return descr; +} + +/** + * @brief create a cuSparse dense descriptor for a matrix * @tparam ValueType Data type of dense_view (float/double) * @tparam IndexType Type of dense_view * @tparam LayoutPolicy layout of dense_view diff --git a/cpp/include/raft/sparse/solver/detail/lanczos.cuh b/cpp/include/raft/sparse/solver/detail/lanczos.cuh index a0537cf970..078a751b7d 100644 --- a/cpp/include/raft/sparse/solver/detail/lanczos.cuh +++ b/cpp/include/raft/sparse/solver/detail/lanczos.cuh @@ -54,6 +54,7 @@ #include #include #include +#include #include #include #include @@ -1553,26 +1554,18 @@ void lanczos_aux(raft::resources const& handle, { auto stream = resource::get_cuda_stream(handle); - auto A_structure = A.structure_view(); - IndexTypeT n = A_structure.get_n_rows(); + IndexTypeT n = A.structure_view().get_n_rows(); + auto v_vector = raft::make_device_vector_view(v.data_handle(), n); + auto u_vector = raft::make_device_vector_view(u.data_handle(), n); raft::copy( v.data_handle(), V.data_handle() + start_idx * V.stride(0), n, stream); // V(start_idx, 0) - auto cusparse_h = resource::get_cusparse_handle(handle); - cusparseSpMatDescr_t cusparse_A; - raft::sparse::detail::cusparsecreatecsr(&cusparse_A, - A_structure.get_n_rows(), - A_structure.get_n_cols(), - A_structure.get_nnz(), - const_cast(A_structure.get_indptr().data()), - const_cast(A_structure.get_indices().data()), - const_cast(A.get_elements().data())); - - cusparseDnVecDescr_t cusparse_v; - cusparseDnVecDescr_t cusparse_u; - raft::sparse::detail::cusparsecreatednvec(&cusparse_v, n, v.data_handle()); - raft::sparse::detail::cusparsecreatednvec(&cusparse_u, n, u.data_handle()); + auto cusparse_h = resource::get_cusparse_handle(handle); + cusparseSpMatDescr_t cusparse_A = raft::sparse::linalg::detail::create_descriptor(A); + + cusparseDnVecDescr_t cusparse_v = raft::sparse::linalg::detail::create_descriptor(v_vector); + cusparseDnVecDescr_t cusparse_u = raft::sparse::linalg::detail::create_descriptor(u_vector); ValueTypeT one = 1; ValueTypeT zero = 0; @@ -1603,8 +1596,6 @@ void lanczos_aux(raft::resources const& handle, auto alpha_i = raft::make_device_scalar_view(alpha.data_handle() + i * alpha.stride(1)); // alpha(0, i) - auto v_vector = raft::make_device_vector_view(v.data_handle(), n); - auto u_vector = raft::make_device_vector_view(u.data_handle(), n); raft::linalg::dot(handle, v_vector, u_vector, alpha_i); raft::matrix::fill(handle, vv, zero); @@ -1706,17 +1697,17 @@ auto lanczos_smallest( ValueTypeT* v0, uint64_t seed) -> int { - auto A_structure = A.structure_view(); - int n = A_structure.get_n_rows(); - int ncv = restartIter; - auto stream = resource::get_cuda_stream(handle); + int n = A.structure_view().get_n_rows(); + int ncv = restartIter; + auto stream = resource::get_cuda_stream(handle); auto V = raft::make_device_matrix(handle, ncv, n); auto V_0_view = raft::make_device_matrix_view(V.data_handle(), 1, n); // First Row V[0] auto v0_view = raft::make_device_matrix_view(v0, 1, n); - auto u = raft::make_device_matrix(handle, 1, n); + auto u = raft::make_device_matrix(handle, 1, n); + auto u_vector = raft::make_device_vector_view(u.data_handle(), n); raft::copy(u.data_handle(), v0, n, stream); auto cublas_h = resource::get_cublas_handle(handle); @@ -1835,7 +1826,7 @@ auto lanczos_smallest( ValueTypeT one = 1; ValueTypeT mone = -1; - // Using raft::linalg::gemv leads to Reason=7:CUBLAS_STATUS_INVALID_VALUE + // Using raft::linalg::gemv leads to Reason=7:CUBLAS_STATUS_INVALID_VALUE (issue raft#2484) raft::linalg::detail::cublasgemv(cublas_h, CUBLAS_OP_T, nEigVecs, @@ -1866,41 +1857,28 @@ auto lanczos_smallest( auto V_0_view = raft::make_device_matrix_view(V.data_handle() + (nEigVecs * n), 1, n); - - auto unrm = raft::make_device_vector(handle, 1); - auto input = raft::make_device_matrix_view(u.data_handle(), 1, n); + auto V_0_view_vector = + raft::make_device_vector_view(V_0_view.data_handle(), n); + auto unrm = raft::make_device_vector(handle, 1); raft::linalg::norm(handle, - input, + raft::make_const_mdspan(u.view()), unrm.view(), raft::linalg::L2Norm, raft::linalg::Apply::ALONG_ROWS, raft::sqrt_op()); - auto u_vector_const = raft::make_device_vector_view(u.data_handle(), n); - raft::linalg::unary_op( - handle, u_vector_const, V_0_view, [device_scalar = unrm.data_handle()] __device__(auto y) { - return y / *device_scalar; - }); + handle, + raft::make_const_mdspan(u_vector), + V_0_view, + [device_scalar = unrm.data_handle()] __device__(auto y) { return y / *device_scalar; }); + + auto cusparse_h = resource::get_cusparse_handle(handle); + cusparseSpMatDescr_t cusparse_A = raft::sparse::linalg::detail::create_descriptor(A); - auto cusparse_h = resource::get_cusparse_handle(handle); - cusparseSpMatDescr_t cusparse_A; - // input_config.a_indptr = const_cast(x_structure.get_indptr().data()); - // input_config.a_indices = const_cast(x_structure.get_indices().data()); - // input_config.a_data = const_cast(x.get_elements().data()); - raft::sparse::detail::cusparsecreatecsr( - &cusparse_A, - A_structure.get_n_rows(), - A_structure.get_n_cols(), - A_structure.get_nnz(), - const_cast(A_structure.get_indptr().data()), - const_cast(A_structure.get_indices().data()), - const_cast(A.get_elements().data())); - - cusparseDnVecDescr_t cusparse_v; - cusparseDnVecDescr_t cusparse_u; - raft::sparse::detail::cusparsecreatednvec(&cusparse_v, n, V_0_view.data_handle()); - raft::sparse::detail::cusparsecreatednvec(&cusparse_u, n, u.data_handle()); + cusparseDnVecDescr_t cusparse_v = + raft::sparse::linalg::detail::create_descriptor(V_0_view_vector); + cusparseDnVecDescr_t cusparse_u = raft::sparse::linalg::detail::create_descriptor(u_vector); ValueTypeT zero = 0; size_t bufferSize; @@ -1928,22 +1906,20 @@ auto lanczos_smallest( stream); auto alpha_k = raft::make_device_scalar_view(alpha.data_handle() + nEigVecs); - auto V_0_view_vector = raft::make_device_vector_view(V_0_view.data_handle(), n); - auto u_view_vector = raft::make_device_vector_view(u.data_handle(), n); raft::linalg::dot( - handle, make_const_mdspan(V_0_view_vector), make_const_mdspan(u_view_vector), alpha_k); + handle, make_const_mdspan(V_0_view_vector), make_const_mdspan(u_vector), alpha_k); raft::linalg::binary_op(handle, - make_const_mdspan(u_view_vector), + make_const_mdspan(u_vector), make_const_mdspan(V_0_view_vector), - u_view_vector, + u_vector, [device_scalar_ptr = alpha_k.data_handle()] __device__( ValueTypeT u_element, ValueTypeT V_0_element) { return u_element - (*device_scalar_ptr) * V_0_element; }); - auto temp = raft::make_device_vector(handle, n); + auto temp = raft::make_device_vector(handle, n); auto V_k = raft::make_device_matrix_view( V.data_handle(), nEigVecs, n); @@ -1994,9 +1970,9 @@ auto lanczos_smallest( auto one_scalar = raft::make_device_scalar(handle, 1); raft::linalg::binary_op(handle, - make_const_mdspan(u_view_vector), + make_const_mdspan(u_vector), make_const_mdspan(temp.view()), - u_view_vector, + u_vector, [device_scalar_ptr = one_scalar.data_handle()] __device__( ValueTypeT u_element, ValueTypeT temp_element) { return u_element - (*device_scalar_ptr) * temp_element; @@ -2013,11 +1989,10 @@ auto lanczos_smallest( auto V_kplus1 = raft::make_device_vector_view(V.data_handle() + V.stride(0) * (nEigVecs + 1), n); - auto u_vector = raft::make_device_vector_view(u.data_handle(), n); raft::linalg::unary_op( handle, - u_vector, + make_const_mdspan(u_vector), V_kplus1, [device_scalar = (beta.data_handle() + beta.stride(1) * nEigVecs)] __device__(auto y) { return y / *device_scalar; diff --git a/cpp/include/raft/sparse/solver/lanczos_types.hpp b/cpp/include/raft/sparse/solver/lanczos_types.hpp index ce5bf86166..8a95f8b621 100644 --- a/cpp/include/raft/sparse/solver/lanczos_types.hpp +++ b/cpp/include/raft/sparse/solver/lanczos_types.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2024, NVIDIA CORPORATION. + * Copyright (c) 2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/include/raft_runtime/solver/lanczos.hpp b/cpp/include/raft_runtime/solver/lanczos.hpp index c43921b282..af4e396027 100644 --- a/cpp/include/raft_runtime/solver/lanczos.hpp +++ b/cpp/include/raft_runtime/solver/lanczos.hpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2024, NVIDIA CORPORATION. + * Copyright (c) 2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/raft_runtime/solver/lanczos_solver.cuh b/cpp/src/raft_runtime/solver/lanczos_solver.cuh index 958a9e8c48..39a0b65179 100644 --- a/cpp/src/raft_runtime/solver/lanczos_solver.cuh +++ b/cpp/src/raft_runtime/solver/lanczos_solver.cuh @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2024, NVIDIA CORPORATION. + * Copyright (c) 2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/raft_runtime/solver/lanczos_solver_int64_double.cu b/cpp/src/raft_runtime/solver/lanczos_solver_int64_double.cu index dbc4a0e886..f772a8a0d1 100644 --- a/cpp/src/raft_runtime/solver/lanczos_solver_int64_double.cu +++ b/cpp/src/raft_runtime/solver/lanczos_solver_int64_double.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2024, NVIDIA CORPORATION. + * Copyright (c) 2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/raft_runtime/solver/lanczos_solver_int64_float.cu b/cpp/src/raft_runtime/solver/lanczos_solver_int64_float.cu index b9bea3cf23..efaf3be565 100644 --- a/cpp/src/raft_runtime/solver/lanczos_solver_int64_float.cu +++ b/cpp/src/raft_runtime/solver/lanczos_solver_int64_float.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2024, NVIDIA CORPORATION. + * Copyright (c) 2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/raft_runtime/solver/lanczos_solver_int_double.cu b/cpp/src/raft_runtime/solver/lanczos_solver_int_double.cu index 3e716396fc..9bbc00e78a 100644 --- a/cpp/src/raft_runtime/solver/lanczos_solver_int_double.cu +++ b/cpp/src/raft_runtime/solver/lanczos_solver_int_double.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2024, NVIDIA CORPORATION. + * Copyright (c) 2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/cpp/src/raft_runtime/solver/lanczos_solver_int_float.cu b/cpp/src/raft_runtime/solver/lanczos_solver_int_float.cu index 9f1f0fc67d..316a9fb7e1 100644 --- a/cpp/src/raft_runtime/solver/lanczos_solver_int_float.cu +++ b/cpp/src/raft_runtime/solver/lanczos_solver_int_float.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2024-2024, NVIDIA CORPORATION. + * Copyright (c) 2024, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. diff --git a/python/pylibraft/CMakeLists.txt b/python/pylibraft/CMakeLists.txt index beca333bc2..5779f2ca0d 100644 --- a/python/pylibraft/CMakeLists.txt +++ b/python/pylibraft/CMakeLists.txt @@ -89,7 +89,7 @@ add_subdirectory(pylibraft/distance) add_subdirectory(pylibraft/matrix) add_subdirectory(pylibraft/neighbors) add_subdirectory(pylibraft/random) -add_subdirectory(pylibraft/solver) +add_subdirectory(pylibraft/sparse) add_subdirectory(pylibraft/cluster) if(DEFINED cython_lib_dir) diff --git a/python/pylibraft/pylibraft/sparse/CMakeLists.txt b/python/pylibraft/pylibraft/sparse/CMakeLists.txt new file mode 100644 index 0000000000..3779fd2715 --- /dev/null +++ b/python/pylibraft/pylibraft/sparse/CMakeLists.txt @@ -0,0 +1,15 @@ +# ============================================================================= +# Copyright (c) 2024, NVIDIA CORPORATION. +# +# Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except +# in compliance with the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software distributed under the License +# is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express +# or implied. See the License for the specific language governing permissions and limitations under +# the License. +# ============================================================================= + +add_subdirectory(linalg) diff --git a/python/pylibraft/pylibraft/sparse/__init__.py b/python/pylibraft/pylibraft/sparse/__init__.py new file mode 100644 index 0000000000..c77def5bb0 --- /dev/null +++ b/python/pylibraft/pylibraft/sparse/__init__.py @@ -0,0 +1,18 @@ +# Copyright (c) 2024, NVIDIA CORPORATION. +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. +# + +from pylibraft.sparse import linalg + +__all__ = ["linalg"] diff --git a/python/pylibraft/pylibraft/solver/CMakeLists.txt b/python/pylibraft/pylibraft/sparse/linalg/CMakeLists.txt similarity index 93% rename from python/pylibraft/pylibraft/solver/CMakeLists.txt rename to python/pylibraft/pylibraft/sparse/linalg/CMakeLists.txt index c9fbd5d0f0..ef16981644 100644 --- a/python/pylibraft/pylibraft/solver/CMakeLists.txt +++ b/python/pylibraft/pylibraft/sparse/linalg/CMakeLists.txt @@ -1,5 +1,5 @@ # ============================================================================= -# Copyright (c) 2022-2024, NVIDIA CORPORATION. +# Copyright (c) 2024, NVIDIA CORPORATION. # # Licensed under the Apache License, Version 2.0 (the "License"); you may not use this file except # in compliance with the License. You may obtain a copy of the License at @@ -23,5 +23,5 @@ set(linked_libraries raft::raft raft::compiled) rapids_cython_create_modules( CXX SOURCE_FILES "${cython_sources}" - LINKED_LIBRARIES "${linked_libraries}" ASSOCIATED_TARGETS raft MODULE_PREFIX solver_ + LINKED_LIBRARIES "${linked_libraries}" ASSOCIATED_TARGETS raft MODULE_PREFIX sparse_ ) diff --git a/python/pylibraft/pylibraft/solver/__init__.pxd b/python/pylibraft/pylibraft/sparse/linalg/__init__.pxd similarity index 100% rename from python/pylibraft/pylibraft/solver/__init__.pxd rename to python/pylibraft/pylibraft/sparse/linalg/__init__.pxd diff --git a/python/pylibraft/pylibraft/solver/__init__.py b/python/pylibraft/pylibraft/sparse/linalg/__init__.py similarity index 92% rename from python/pylibraft/pylibraft/solver/__init__.py rename to python/pylibraft/pylibraft/sparse/linalg/__init__.py index 30afe63e7e..04a8106496 100644 --- a/python/pylibraft/pylibraft/solver/__init__.py +++ b/python/pylibraft/pylibraft/sparse/linalg/__init__.py @@ -1,4 +1,4 @@ -# Copyright (c) 2022-2024, NVIDIA CORPORATION. +# Copyright (c) 2024, NVIDIA CORPORATION. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/pylibraft/pylibraft/solver/cpp/__init__.pxd b/python/pylibraft/pylibraft/sparse/linalg/cpp/__init__.pxd similarity index 100% rename from python/pylibraft/pylibraft/solver/cpp/__init__.pxd rename to python/pylibraft/pylibraft/sparse/linalg/cpp/__init__.pxd diff --git a/python/pylibraft/pylibraft/solver/cpp/__init__.py b/python/pylibraft/pylibraft/sparse/linalg/cpp/__init__.py similarity index 100% rename from python/pylibraft/pylibraft/solver/cpp/__init__.py rename to python/pylibraft/pylibraft/sparse/linalg/cpp/__init__.py diff --git a/python/pylibraft/pylibraft/solver/lanczos.pyx b/python/pylibraft/pylibraft/sparse/linalg/lanczos.pyx similarity index 99% rename from python/pylibraft/pylibraft/solver/lanczos.pyx rename to python/pylibraft/pylibraft/sparse/linalg/lanczos.pyx index e2819b891c..8b1cbff418 100644 --- a/python/pylibraft/pylibraft/solver/lanczos.pyx +++ b/python/pylibraft/pylibraft/sparse/linalg/lanczos.pyx @@ -1,5 +1,5 @@ # -# Copyright (c) 2024-2024, NVIDIA CORPORATION. +# Copyright (c) 2024, NVIDIA CORPORATION. # # Licensed under the Apache License, Version 2.0 (the "License"); # you may not use this file except in compliance with the License. diff --git a/python/pylibraft/pylibraft/test/test_solver.py b/python/pylibraft/pylibraft/test/test_sparse.py similarity index 99% rename from python/pylibraft/pylibraft/test/test_solver.py rename to python/pylibraft/pylibraft/test/test_sparse.py index 8cb0011975..10b261d322 100644 --- a/python/pylibraft/pylibraft/test/test_solver.py +++ b/python/pylibraft/pylibraft/test/test_sparse.py @@ -19,7 +19,7 @@ import pytest from cupyx.scipy import sparse -from pylibraft.solver import eigsh +from pylibraft.sparse.linalg import eigsh def shaped_random(