From 24c73c8f73ec3ac0c663b61cb2130ca6908d4cbe Mon Sep 17 00:00:00 2001 From: Luc Berger-Vergiat Date: Thu, 9 Nov 2023 12:57:08 -0700 Subject: [PATCH] LAPACK: adding rocsolver TPL Adding the necessary CMake logic and TPL layer to support rocsolver for LAPACK. Enabling the TPL in gesv and updating gesv test to run by default the more common configurations and only run specific ones when the associated TPL (MAGMA) is enabled. --- CMakeLists.txt | 1 + cmake/Dependencies.cmake | 2 +- cmake/KokkosKernels_config.h.in | 2 + cmake/Modules/FindTPLROCSOLVER.cmake | 9 + cmake/kokkoskernels_features.cmake | 12 + cmake/kokkoskernels_tpls.cmake | 8 +- .../src/KokkosKernels_PrintConfiguration.hpp | 7 + lapack/CMakeLists.txt | 7 - lapack/impl/KokkosLapack_gesv_spec.hpp | 44 +- lapack/src/KokkosLapack_gesv.hpp | 44 +- .../tpls/KokkosLapack_gesv_tpl_spec_avail.hpp | 86 ++-- .../tpls/KokkosLapack_gesv_tpl_spec_decl.hpp | 442 ++++++++---------- lapack/unit_test/Test_Lapack_gesv.hpp | 164 ++++--- scripts/cm_test_all_sandia | 4 +- 14 files changed, 428 insertions(+), 404 deletions(-) create mode 100644 cmake/Modules/FindTPLROCSOLVER.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index 812640374b..fc41d40452 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -377,6 +377,7 @@ ELSE() KOKKOSKERNELS_LINK_TPL(kokkoskernels PUBLIC CUSPARSE) KOKKOSKERNELS_LINK_TPL(kokkoskernels PUBLIC ROCBLAS) KOKKOSKERNELS_LINK_TPL(kokkoskernels PUBLIC ROCSPARSE) + KOKKOSKERNELS_LINK_TPL(kokkoskernels PUBLIC ROCSOLVER) KOKKOSKERNELS_LINK_TPL(kokkoskernels PUBLIC METIS) KOKKOSKERNELS_LINK_TPL(kokkoskernels PUBLIC ARMPL) KOKKOSKERNELS_LINK_TPL(kokkoskernels PUBLIC MAGMA) diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index 777d4445b3..104d153347 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -1,6 +1,6 @@ TRIBITS_PACKAGE_DEFINE_DEPENDENCIES( LIB_REQUIRED_PACKAGES Kokkos - LIB_OPTIONAL_TPLS quadmath MKL BLAS LAPACK CUSPARSE METIS SuperLU Cholmod CUBLAS ROCBLAS ROCSPARSE + LIB_OPTIONAL_TPLS quadmath MKL BLAS LAPACK METIS SuperLU Cholmod CUBLAS CUSPARSE ROCBLAS ROCSPARSE ROCSOLVER TEST_OPTIONAL_TPLS yaml-cpp ) # NOTE: If you update names in LIB_OPTIONAL_TPLS above, make sure to map those names in diff --git a/cmake/KokkosKernels_config.h.in b/cmake/KokkosKernels_config.h.in index d94860e380..c40a2b18a7 100644 --- a/cmake/KokkosKernels_config.h.in +++ b/cmake/KokkosKernels_config.h.in @@ -138,6 +138,8 @@ #cmakedefine KOKKOSKERNELS_ENABLE_TPL_ROCBLAS /* ROCSPARSE */ #cmakedefine KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE +/* ROCSOLVER */ +#cmakedefine KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER #cmakedefine KOKKOSKERNELS_ENABLE_SUPERNODAL_SPTRSV diff --git a/cmake/Modules/FindTPLROCSOLVER.cmake b/cmake/Modules/FindTPLROCSOLVER.cmake new file mode 100644 index 0000000000..8f2a92cfda --- /dev/null +++ b/cmake/Modules/FindTPLROCSOLVER.cmake @@ -0,0 +1,9 @@ +# LBV: 11/08/2023: This file follows the partern of FindTPLROCBLAS.cmake/FindTPLROCSPARSE.cmake +FIND_PACKAGE(ROCSOLVER) +if(TARGET roc::rocsolver) + SET(TPL_ROCSOLVER_IMPORTED_NAME roc::rocsolver) + SET(TPL_IMPORTED_NAME roc::rocsolver) + ADD_LIBRARY(KokkosKernels::ROCSOLVER ALIAS roc::rocsolver) +ELSE() + MESSAGE(FATAL_ERROR "Package ROCSOLVER requested but not found") +ENDIF() diff --git a/cmake/kokkoskernels_features.cmake b/cmake/kokkoskernels_features.cmake index aacc1c8451..3ecc95d6b5 100644 --- a/cmake/kokkoskernels_features.cmake +++ b/cmake/kokkoskernels_features.cmake @@ -27,3 +27,15 @@ IF (KOKKOSKERNELS_ENABLE_TPL_BLAS OR KOKKOSKERNELS_ENABLE_TPL_MKL OR KOKKOSKERNE INCLUDE(CheckHostBlasReturnComplex.cmake) CHECK_HOST_BLAS_RETURN_COMPLEX(KOKKOSKERNELS_TPL_BLAS_RETURN_COMPLEX) ENDIF() + +# ================================================================== +# Lapack requirements +# ================================================================== + +IF (KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER AND NOT KOKKOSKERNELS_ENABLE_TPL_ROCBLAS AND NOT KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE) + MESSAGE(FATAL_ERROR "rocSOLVER requires rocBLAS and rocSPARSE, please reconfigure with KOKKOSKERNELS_ENABLE_TPL_ROCBLAS:BOOL=ON and KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE:BOOL=ON.") +ELSEIF (KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER AND NOT KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE) + MESSAGE(FATAL_ERROR "rocSOLVER requires rocSPARSE, please reconfigure with KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE:BOOL=ON.") +ELSEIF (KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER AND NOT KOKKOSKERNELS_ENABLE_TPL_ROCBLAS) + MESSAGE(FATAL_ERROR "rocSOLVER requires rocBLAS, please reconfigure with KOKKOSKERNELS_ENABLE_TPL_ROCBLAS:BOOL=ON.") +ENDIF() diff --git a/cmake/kokkoskernels_tpls.cmake b/cmake/kokkoskernels_tpls.cmake index 08c7158148..2f54278d1b 100644 --- a/cmake/kokkoskernels_tpls.cmake +++ b/cmake/kokkoskernels_tpls.cmake @@ -460,15 +460,18 @@ KOKKOSKERNELS_ADD_OPTION(NO_DEFAULT_ROCM_TPLS OFF BOOL "Whether ROCM TPLs should # Unlike CUDA, ROCm does not automatically install these TPLs SET(ROCBLAS_DEFAULT OFF) SET(ROCSPARSE_DEFAULT OFF) +SET(ROCSOLVER_DEFAULT OFF) # Since the default is OFF we do not really need this piece of logic here. # IF(KOKKOSKERNELS_NO_DEFAULT_ROCM_TPLS) # SET(ROCBLAS_DEFAULT OFF) # SET(ROCSPARSE_DEFAULT OFF) # ENDIF() KOKKOSKERNELS_ADD_TPL_OPTION(ROCBLAS ${ROCBLAS_DEFAULT} "Whether to enable ROCBLAS" - DEFAULT_DOCSTRING "ON if HIP-enabled Kokkos, otherwise OFF") + DEFAULT_DOCSTRING "OFF even if HIP-enabled Kokkos") KOKKOSKERNELS_ADD_TPL_OPTION(ROCSPARSE ${ROCSPARSE_DEFAULT} "Whether to enable ROCSPARSE" - DEFAULT_DOCSTRING "ON if HIP-enabled Kokkos, otherwise OFF") + DEFAULT_DOCSTRING "OFF even if HIP-enabled Kokkos") +KOKKOSKERNELS_ADD_TPL_OPTION(ROCSOLVER ${ROCSOLVER_DEFAULT} "Whether to enable ROCSOLVER" + DEFAULT_DOCSTRING "OFF even if HIP-enabled Kokkos") IF (KOKKOSKERNELS_ENABLE_TPL_MAGMA) IF (F77_BLAS_MANGLE STREQUAL "(name,NAME) name ## _") @@ -507,6 +510,7 @@ IF (NOT KOKKOSKERNELS_HAS_TRILINOS) KOKKOSKERNELS_IMPORT_TPL(MAGMA) KOKKOSKERNELS_IMPORT_TPL(ROCBLAS) KOKKOSKERNELS_IMPORT_TPL(ROCSPARSE) + KOKKOSKERNELS_IMPORT_TPL(ROCSOLVER) ELSE () IF (Trilinos_ENABLE_SuperLU5_API) SET(HAVE_KOKKOSKERNELS_SUPERLU5_API TRUE) diff --git a/common/src/KokkosKernels_PrintConfiguration.hpp b/common/src/KokkosKernels_PrintConfiguration.hpp index cd2333b3ec..55e7285ed2 100644 --- a/common/src/KokkosKernels_PrintConfiguration.hpp +++ b/common/src/KokkosKernels_PrintConfiguration.hpp @@ -110,6 +110,13 @@ inline void print_enabled_tpls(std::ostream& os) { os << " " << "KOKKOSKERNELS_ENABLE_TPL_ROCSPARSE: no\n"; #endif +#ifdef KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER + os << " " + << "KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER: yes\n"; +#else + os << " " + << "KOKKOSKERNELS_ENABLE_TPL_ROCOLVER: no\n"; +#endif #ifdef KOKKOSKERNELS_ENABLE_TPL_METIS os << "KOKKOSKERNELS_ENABLE_TPL_METIS: yes\n"; #else diff --git a/lapack/CMakeLists.txt b/lapack/CMakeLists.txt index 8ab784a325..ee91079378 100644 --- a/lapack/CMakeLists.txt +++ b/lapack/CMakeLists.txt @@ -34,13 +34,6 @@ IF (KOKKOSKERNELS_ENABLE_TPL_CUSOLVER) ) ENDIF() -# Include rocm lapack TPL source file -IF (KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER) - LIST(APPEND SOURCES - lapack/tpls/KokkosLapack_Rocm_tpl.cpp - ) -ENDIF() - ################## # # # ETI generation # diff --git a/lapack/impl/KokkosLapack_gesv_spec.hpp b/lapack/impl/KokkosLapack_gesv_spec.hpp index b9f8549311..57098f75fc 100644 --- a/lapack/impl/KokkosLapack_gesv_spec.hpp +++ b/lapack/impl/KokkosLapack_gesv_spec.hpp @@ -28,7 +28,7 @@ namespace KokkosLapack { namespace Impl { // Specialization struct which defines whether a specialization exists -template +template struct gesv_eti_spec_avail { enum : bool { value = false }; }; @@ -46,12 +46,16 @@ struct gesv_eti_spec_avail { EXEC_SPACE_TYPE, MEM_SPACE_TYPE) \ template <> \ struct gesv_eti_spec_avail< \ + EXEC_SPACE_TYPE, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits > > { \ + Kokkos::MemoryTraits>, \ + Kokkos::View, \ + Kokkos::MemoryTraits>> { \ enum : bool { value = true }; \ }; @@ -65,20 +69,24 @@ namespace Impl { // Unification layer /// \brief Implementation of KokkosLapack::gesv. -template ::value, - bool eti_spec_avail = gesv_eti_spec_avail::value> +template ::value, + bool eti_spec_avail = + gesv_eti_spec_avail::value> struct GESV { - static void gesv(const AMatrix &A, const BXMV &B, const IPIVV &IPIV); + static void gesv(const ExecutionSpace &space, const AMatrix &A, const BXMV &B, + const IPIVV &IPIV); }; #if !defined(KOKKOSKERNELS_ETI_ONLY) || KOKKOSKERNELS_IMPL_COMPILE_LIBRARY //! Full specialization of gesv for multi vectors. // Unification layer -template -struct GESV { - static void gesv(const AMatrix & /* A */, const BXMV & /* B */, - const IPIVV & /* IPIV */) { +template +struct GESV { + static void gesv(const ExecutionSpace & /* space */, const AMatrix & /* A */, + const BXMV & /* B */, const IPIVV & /* IPIV */) { // NOTE: Might add the implementation of KokkosLapack::gesv later throw std::runtime_error( "No fallback implementation of GESV (general LU factorization & solve) " @@ -100,31 +108,33 @@ struct GESV { #define KOKKOSLAPACK_GESV_ETI_SPEC_DECL(SCALAR_TYPE, LAYOUT_TYPE, \ EXEC_SPACE_TYPE, MEM_SPACE_TYPE) \ extern template struct GESV< \ + EXEC_SPACE_TYPE, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ false, true>; #define KOKKOSLAPACK_GESV_ETI_SPEC_INST(SCALAR_TYPE, LAYOUT_TYPE, \ EXEC_SPACE_TYPE, MEM_SPACE_TYPE) \ template struct GESV< \ + EXEC_SPACE_TYPE, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ false, true>; #include diff --git a/lapack/src/KokkosLapack_gesv.hpp b/lapack/src/KokkosLapack_gesv.hpp index 4c9058f8ab..74d2e01cf9 100644 --- a/lapack/src/KokkosLapack_gesv.hpp +++ b/lapack/src/KokkosLapack_gesv.hpp @@ -34,28 +34,40 @@ namespace KokkosLapack { /// \brief Solve the dense linear equation system A*X = B. /// +/// \tparam ExecutionSpace the space where the kernel will run. /// \tparam AMatrix Input matrix/Output LU, as a 2-D Kokkos::View. /// \tparam BXMV Input (right-hand side)/Output (solution) (multi)vector, as a -/// 1-D or 2-D Kokkos::View. \tparam IPIVV Output pivot indices, as a 1-D -/// Kokkos::View +/// 1-D or 2-D Kokkos::View. +/// \tparam IPIVV Output pivot indices, as a 1-D Kokkos::View /// /// \param A [in,out] On entry, the N-by-N matrix to be solved. On exit, the /// factors L and U from /// the factorization A = P*L*U; the unit diagonal elements of L are not /// stored. /// \param B [in,out] On entry, the right hand side (multi)vector B. On exit, -/// the solution (multi)vector X. \param IPIV [out] On exit, the pivot indices -/// (for partial pivoting). If the View extents are zero and -/// its data pointer is NULL, pivoting is not used. +/// the solution (multi)vector X. +/// \param IPIV [out] On exit, the pivot indices (for partial pivoting). +/// If the View extents are zero and its data pointer is NULL, pivoting is not +/// used. /// -template -void gesv(const AMatrix& A, const BXMV& B, const IPIVV& IPIV) { - // NOTE: Currently, KokkosLapack::gesv only supports for MAGMA TPL and LAPACK - // TPL. - // MAGMA TPL should be enabled to call the MAGMA GPU interface for - // device views LAPACK TPL should be enabled to call the LAPACK - // interface for host views +template +void gesv(const ExecutionSpace& space, const AMatrix& A, const BXMV& B, + const IPIVV& IPIV) { + // NOTE: Currently, KokkosLapack::gesv only supports LAPACK, MAGMA and + // rocSOLVER TPLs. + // MAGMA/rocSOLVER TPL should be enabled to call the MAGMA/rocSOLVER GPU + // interface for device views LAPACK TPL should be enabled to call the + // LAPACK interface for host views + static_assert( + Kokkos::SpaceAccessibility::accessible); + static_assert( + Kokkos::SpaceAccessibility::accessible); + static_assert( + Kokkos::SpaceAccessibility::accessible); static_assert(Kokkos::is_view::value, "KokkosLapack::gesv: A must be a Kokkos::View."); static_assert(Kokkos::is_view::value, @@ -137,12 +149,12 @@ void gesv(const AMatrix& A, const BXMV& B, const IPIVV& IPIV) { if (BXMV::rank == 1) { auto B_i = BXMV_Internal(B.data(), B.extent(0), 1); - KokkosLapack::Impl::GESV::gesv(A_i, B_i, IPIV_i); + KokkosLapack::Impl::GESV::gesv(space, A_i, B_i, IPIV_i); } else { // BXMV::rank == 2 auto B_i = BXMV_Internal(B.data(), B.extent(0), B.extent(1)); - KokkosLapack::Impl::GESV::gesv(A_i, B_i, IPIV_i); + KokkosLapack::Impl::GESV::gesv(space, A_i, B_i, IPIV_i); } } diff --git a/lapack/tpls/KokkosLapack_gesv_tpl_spec_avail.hpp b/lapack/tpls/KokkosLapack_gesv_tpl_spec_avail.hpp index a3d8bb6ee9..fc8f634078 100644 --- a/lapack/tpls/KokkosLapack_gesv_tpl_spec_avail.hpp +++ b/lapack/tpls/KokkosLapack_gesv_tpl_spec_avail.hpp @@ -20,7 +20,7 @@ namespace KokkosLapack { namespace Impl { // Specialization struct which defines whether a specialization exists -template +template struct gesv_tpl_spec_avail { enum : bool { value = false }; }; @@ -31,9 +31,12 @@ struct gesv_tpl_spec_avail { #define KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK(SCALAR, LAYOUT, MEMSPACE) \ template \ struct gesv_tpl_spec_avail< \ + ExecSpace, \ Kokkos::View, \ Kokkos::MemoryTraits >, \ Kokkos::View, \ + Kokkos::MemoryTraits >, \ + Kokkos::View, \ Kokkos::MemoryTraits > > { \ enum : bool { value = true }; \ }; @@ -46,37 +49,22 @@ KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK(Kokkos::complex, Kokkos::LayoutLeft, Kokkos::HostSpace) KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK(Kokkos::complex, Kokkos::LayoutLeft, Kokkos::HostSpace) -/* -#if defined (KOKKOSKERNELS_INST_DOUBLE) \ - && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT) - KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK( double, Kokkos::LayoutRight, -Kokkos::HostSpace) #endif -#if defined (KOKKOSKERNELS_INST_FLOAT) \ - && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT) - KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK( float, Kokkos::LayoutRight, -Kokkos::HostSpace) #endif -#if defined (KOKKOSKERNELS_INST_KOKKOS_COMPLEX_DOUBLE_) \ - && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT) - KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK( Kokkos::complex, -Kokkos::LayoutRight, Kokkos::HostSpace) #endif -#if defined (KOKKOSKERNELS_INST_KOKKOS_COMPLEX_FLOAT_) \ - && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT) - KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK( Kokkos::complex, -Kokkos::LayoutRight, Kokkos::HostSpace) #endif -*/ #endif // MAGMA #ifdef KOKKOSKERNELS_ENABLE_TPL_MAGMA -#define KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA(SCALAR, LAYOUT, MEMSPACE) \ - template \ - struct gesv_tpl_spec_avail< \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - Kokkos::View, \ - Kokkos::MemoryTraits > > { \ - enum : bool { value = true }; \ +#define KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA(SCALAR, LAYOUT, MEMSPACE) \ + template <> \ + struct gesv_tpl_spec_avail< \ + Kokkos::Cuda, \ + Kokkos::View, \ + Kokkos::MemoryTraits >, \ + Kokkos::View, \ + Kokkos::MemoryTraits >, \ + Kokkos::View > > { \ + enum : bool { value = true }; \ }; KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA(double, Kokkos::LayoutLeft, @@ -88,24 +76,32 @@ KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA(Kokkos::complex, KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA(Kokkos::complex, Kokkos::LayoutLeft, Kokkos::CudaSpace) -/* -#if defined (KOKKOSKERNELS_INST_DOUBLE) \ - && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT) - KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA( double, Kokkos::LayoutRight, -Kokkos::CudaSpace) #endif -#if defined (KOKKOSKERNELS_INST_FLOAT) \ - && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT) - KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA( float, Kokkos::LayoutRight, -Kokkos::CudaSpace) #endif -#if defined (KOKKOSKERNELS_INST_KOKKOS_COMPLEX_DOUBLE_) \ - && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT) - KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA( -Kokkos::complex,Kokkos::LayoutRight, Kokkos::CudaSpace) #endif -#if defined (KOKKOSKERNELS_INST_KOKKOS_COMPLEX_FLOAT_) \ - && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT) - KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA( Kokkos::complex, -Kokkos::LayoutRight, Kokkos::CudaSpace) #endif -*/ +#endif + +#ifdef KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER + +#define KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_ROCSOLVER(SCALAR, LAYOUT, MEMSPACE) \ + template <> \ + struct gesv_tpl_spec_avail< \ + Kokkos::HIP, \ + Kokkos::View, \ + Kokkos::MemoryTraits >, \ + Kokkos::View, \ + Kokkos::MemoryTraits >, \ + Kokkos::View, \ + Kokkos::MemoryTraits > > { \ + enum : bool { value = true }; \ + }; + +KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_ROCSOLVER(double, Kokkos::LayoutLeft, + Kokkos::HIPSpace) +KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_ROCSOLVER(float, Kokkos::LayoutLeft, + Kokkos::HIPSpace) +KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_ROCSOLVER(Kokkos::complex, + Kokkos::LayoutLeft, Kokkos::HIPSpace) +KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_ROCSOLVER(Kokkos::complex, + Kokkos::LayoutLeft, Kokkos::HIPSpace) + #endif } // namespace Impl diff --git a/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp b/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp index 2baa76a132..957ac7c138 100644 --- a/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp +++ b/lapack/tpls/KokkosLapack_gesv_tpl_spec_decl.hpp @@ -45,229 +45,84 @@ inline void gesv_print_specialization() { namespace KokkosLapack { namespace Impl { -#define KOKKOSLAPACK_DGESV_LAPACK(LAYOUT, MEM_SPACE, ETI_SPEC_AVAIL) \ - template \ - struct GESV< \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - true, ETI_SPEC_AVAIL> { \ - typedef double SCALAR; \ - typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ - AViewType; \ - typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ - BViewType; \ - typedef Kokkos::View< \ - int*, LAYOUT, \ - Kokkos::Device, \ - Kokkos::MemoryTraits > \ - PViewType; \ - \ - static void gesv(const AViewType& A, const BViewType& B, \ - const PViewType& IPIV) { \ - Kokkos::Profiling::pushRegion("KokkosLapack::gesv[TPL_LAPACK,double]"); \ - gesv_print_specialization(); \ - const bool with_pivot = \ - !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr)); \ - \ - const int N = static_cast(A.extent(1)); \ - const int AST = static_cast(A.stride(1)); \ - const int LDA = (AST == 0) ? 1 : AST; \ - const int BST = static_cast(B.stride(1)); \ - const int LDB = (BST == 0) ? 1 : BST; \ - const int NRHS = static_cast(B.extent(1)); \ - \ - int info = 0; \ - \ - if (with_pivot) { \ - HostLapack::gesv(N, NRHS, A.data(), LDA, IPIV.data(), \ - B.data(), LDB, info); \ - } \ - Kokkos::Profiling::popRegion(); \ - } \ - }; +template +void lapackGesvWrapper(const AViewType& A, const BViewType& B, + const IPIVViewType& IPIV) { + using Scalar = typename AViewType::non_const_value_type; -#define KOKKOSLAPACK_SGESV_LAPACK(LAYOUT, MEM_SPACE, ETI_SPEC_AVAIL) \ - template \ - struct GESV< \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - true, ETI_SPEC_AVAIL> { \ - typedef float SCALAR; \ - typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ - AViewType; \ - typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ - BViewType; \ - typedef Kokkos::View< \ - int*, LAYOUT, \ - Kokkos::Device, \ - Kokkos::MemoryTraits > \ - PViewType; \ - \ - static void gesv(const AViewType& A, const BViewType& B, \ - const PViewType& IPIV) { \ - Kokkos::Profiling::pushRegion("KokkosLapack::gesv[TPL_LAPACK,float]"); \ - gesv_print_specialization(); \ - const bool with_pivot = \ - !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr)); \ - \ - const int N = static_cast(A.extent(1)); \ - const int AST = static_cast(A.stride(1)); \ - const int LDA = (AST == 0) ? 1 : AST; \ - const int BST = static_cast(B.stride(1)); \ - const int LDB = (BST == 0) ? 1 : BST; \ - const int NRHS = static_cast(B.extent(1)); \ - \ - int info = 0; \ - \ - if (with_pivot) { \ - HostLapack::gesv(N, NRHS, A.data(), LDA, IPIV.data(), B.data(), \ - LDB, info); \ - } \ - Kokkos::Profiling::popRegion(); \ - } \ - }; + const bool with_pivot = !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr)); -#define KOKKOSLAPACK_ZGESV_LAPACK(LAYOUT, MEM_SPACE, ETI_SPEC_AVAIL) \ - template \ - struct GESV**, LAYOUT, \ - Kokkos::Device, \ - Kokkos::MemoryTraits >, \ - Kokkos::View**, LAYOUT, \ - Kokkos::Device, \ - Kokkos::MemoryTraits >, \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - true, ETI_SPEC_AVAIL> { \ - typedef Kokkos::complex SCALAR; \ - typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ - AViewType; \ - typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ - BViewType; \ - typedef Kokkos::View< \ - int*, LAYOUT, \ - Kokkos::Device, \ - Kokkos::MemoryTraits > \ - PViewType; \ - \ - static void gesv(const AViewType& A, const BViewType& B, \ - const PViewType& IPIV) { \ - Kokkos::Profiling::pushRegion( \ - "KokkosLapack::gesv[TPL_LAPACK,complex]"); \ - gesv_print_specialization(); \ - const bool with_pivot = \ - !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr)); \ - \ - const int N = static_cast(A.extent(1)); \ - const int AST = static_cast(A.stride(1)); \ - const int LDA = (AST == 0) ? 1 : AST; \ - const int BST = static_cast(B.stride(1)); \ - const int LDB = (BST == 0) ? 1 : BST; \ - const int NRHS = static_cast(B.extent(1)); \ - \ - int info = 0; \ - \ - if (with_pivot) { \ - HostLapack >::gesv( \ - N, NRHS, reinterpret_cast*>(A.data()), LDA, \ - IPIV.data(), reinterpret_cast*>(B.data()), \ - LDB, info); \ - } \ - Kokkos::Profiling::popRegion(); \ - } \ - }; + const int N = static_cast(A.extent(1)); + const int AST = static_cast(A.stride(1)); + const int LDA = (AST == 0) ? 1 : AST; + const int BST = static_cast(B.stride(1)); + const int LDB = (BST == 0) ? 1 : BST; + const int NRHS = static_cast(B.extent(1)); -#define KOKKOSLAPACK_CGESV_LAPACK(LAYOUT, MEM_SPACE, ETI_SPEC_AVAIL) \ - template \ - struct GESV**, LAYOUT, \ - Kokkos::Device, \ - Kokkos::MemoryTraits >, \ - Kokkos::View**, LAYOUT, \ - Kokkos::Device, \ - Kokkos::MemoryTraits >, \ - Kokkos::View, \ - Kokkos::MemoryTraits >, \ - true, ETI_SPEC_AVAIL> { \ - typedef Kokkos::complex SCALAR; \ - typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ - AViewType; \ - typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ - BViewType; \ - typedef Kokkos::View< \ - int*, LAYOUT, \ - Kokkos::Device, \ - Kokkos::MemoryTraits > \ - PViewType; \ - \ - static void gesv(const AViewType& A, const BViewType& B, \ - const PViewType& IPIV) { \ - Kokkos::Profiling::pushRegion( \ - "KokkosLapack::gesv[TPL_LAPACK,complex]"); \ - gesv_print_specialization(); \ - const bool with_pivot = \ - !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr)); \ - \ - const int N = static_cast(A.extent(1)); \ - const int AST = static_cast(A.stride(1)); \ - const int LDA = (AST == 0) ? 1 : AST; \ - const int BST = static_cast(B.stride(1)); \ - const int LDB = (BST == 0) ? 1 : BST; \ - const int NRHS = static_cast(B.extent(1)); \ - \ - int info = 0; \ - \ - if (with_pivot) { \ - HostLapack >::gesv( \ - N, NRHS, reinterpret_cast*>(A.data()), LDA, \ - IPIV.data(), reinterpret_cast*>(B.data()), \ - LDB, info); \ - } \ - Kokkos::Profiling::popRegion(); \ - } \ + int info = 0; + + if (with_pivot) { + if constexpr (Kokkos::ArithTraits::is_complex) { + using MagType = typename Kokkos::ArithTraits::mag_type; + + HostLapack>::gesv( + N, NRHS, reinterpret_cast*>(A.data()), LDA, + IPIV.data(), reinterpret_cast*>(B.data()), LDB, + info); + } else { + HostLapack::gesv(N, NRHS, A.data(), LDA, IPIV.data(), B.data(), + LDB, info); + } + } +} + +#define KOKKOSLAPACK_GESV_LAPACK(SCALAR, LAYOUT, MEM_SPACE, ETI_SPEC_AVAIL) \ + template \ + struct GESV< \ + ExecSpace, \ + Kokkos::View, \ + Kokkos::MemoryTraits>, \ + Kokkos::View, \ + Kokkos::MemoryTraits>, \ + Kokkos::View, \ + Kokkos::MemoryTraits>, \ + true, ETI_SPEC_AVAIL> { \ + using AViewType = \ + Kokkos::View, \ + Kokkos::MemoryTraits>; \ + using BViewType = \ + Kokkos::View, \ + Kokkos::MemoryTraits>; \ + using PViewType = Kokkos::View>; \ + \ + static void gesv(const ExecSpace& /* space */, const AViewType& A, \ + const BViewType& B, const PViewType& IPIV) { \ + Kokkos::Profiling::pushRegion("KokkosLapack::gesv[TPL_LAPACK," #SCALAR \ + "]"); \ + gesv_print_specialization(); \ + lapackGesvWrapper(A, B, IPIV); \ + Kokkos::Profiling::popRegion(); \ + } \ }; -KOKKOSLAPACK_DGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, true) -KOKKOSLAPACK_DGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, false) +KOKKOSLAPACK_GESV_LAPACK(float, Kokkos::LayoutLeft, Kokkos::HostSpace, true) +KOKKOSLAPACK_GESV_LAPACK(float, Kokkos::LayoutLeft, Kokkos::HostSpace, false) -KOKKOSLAPACK_SGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, true) -KOKKOSLAPACK_SGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, false) +KOKKOSLAPACK_GESV_LAPACK(double, Kokkos::LayoutLeft, Kokkos::HostSpace, true) +KOKKOSLAPACK_GESV_LAPACK(double, Kokkos::LayoutLeft, Kokkos::HostSpace, false) -KOKKOSLAPACK_ZGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, true) -KOKKOSLAPACK_ZGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, false) +KOKKOSLAPACK_GESV_LAPACK(Kokkos::complex, Kokkos::LayoutLeft, + Kokkos::HostSpace, true) +KOKKOSLAPACK_GESV_LAPACK(Kokkos::complex, Kokkos::LayoutLeft, + Kokkos::HostSpace, false) -KOKKOSLAPACK_CGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, true) -KOKKOSLAPACK_CGESV_LAPACK(Kokkos::LayoutLeft, Kokkos::HostSpace, false) +KOKKOSLAPACK_GESV_LAPACK(Kokkos::complex, Kokkos::LayoutLeft, + Kokkos::HostSpace, true) +KOKKOSLAPACK_GESV_LAPACK(Kokkos::complex, Kokkos::LayoutLeft, + Kokkos::HostSpace, false) } // namespace Impl } // namespace KokkosLapack @@ -284,27 +139,27 @@ namespace Impl { template \ struct GESV< \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ true, ETI_SPEC_AVAIL> { \ typedef double SCALAR; \ typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ AViewType; \ typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ BViewType; \ typedef Kokkos::View< \ magma_int_t*, LAYOUT, \ Kokkos::Device, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ PViewType; \ \ static void gesv(const AViewType& A, const BViewType& B, \ @@ -343,27 +198,27 @@ namespace Impl { template \ struct GESV< \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ true, ETI_SPEC_AVAIL> { \ typedef float SCALAR; \ typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ AViewType; \ typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ BViewType; \ typedef Kokkos::View< \ magma_int_t*, LAYOUT, \ Kokkos::Device, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ PViewType; \ \ static void gesv(const AViewType& A, const BViewType& B, \ @@ -402,28 +257,28 @@ namespace Impl { template \ struct GESV**, LAYOUT, \ Kokkos::Device, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View**, LAYOUT, \ Kokkos::Device, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ true, ETI_SPEC_AVAIL> { \ typedef Kokkos::complex SCALAR; \ typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ AViewType; \ typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ BViewType; \ typedef Kokkos::View< \ magma_int_t*, LAYOUT, \ Kokkos::Device, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ PViewType; \ \ static void gesv(const AViewType& A, const BViewType& B, \ @@ -463,28 +318,28 @@ namespace Impl { template \ struct GESV**, LAYOUT, \ Kokkos::Device, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View**, LAYOUT, \ Kokkos::Device, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ Kokkos::View, \ - Kokkos::MemoryTraits >, \ + Kokkos::MemoryTraits>, \ true, ETI_SPEC_AVAIL> { \ typedef Kokkos::complex SCALAR; \ typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ AViewType; \ typedef Kokkos::View, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ BViewType; \ typedef Kokkos::View< \ magma_int_t*, LAYOUT, \ Kokkos::Device, \ - Kokkos::MemoryTraits > \ + Kokkos::MemoryTraits> \ PViewType; \ \ static void gesv(const AViewType& A, const BViewType& B, \ @@ -536,4 +391,113 @@ KOKKOSLAPACK_CGESV_MAGMA(Kokkos::LayoutLeft, Kokkos::CudaSpace, false) } // namespace KokkosLapack #endif // KOKKOSKERNELS_ENABLE_TPL_MAGMA +// ROCSOLVER +#ifdef KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER +#include +#include + +namespace KokkosLapack { +namespace Impl { + +template +void rocsolverGesvWrapper(const ExecutionSpace& space, const IPIVViewType& IPIV, + const AViewType& A, const BViewType& B) { + using Scalar = typename BViewType::non_const_value_type; + using ALayout_t = typename AViewType::array_layout; + using BLayout_t = typename BViewType::array_layout; + + const rocblas_int N = static_cast(A.extent(0)); + const rocblas_int nrhs = static_cast(B.extent(1)); + const rocblas_int lda = std::is_same_v + ? A.stride(0) + : A.stride(1); + const rocblas_int ldb = std::is_same_v + ? B.stride(0) + : B.stride(1); + Kokkos::View info("rocsolver info"); + + KokkosBlas::Impl::RocBlasSingleton& s = + KokkosBlas::Impl::RocBlasSingleton::singleton(); + KOKKOS_ROCBLAS_SAFE_CALL_IMPL( + rocblas_set_stream(s.handle, space.hip_stream())); + if constexpr (std::is_same_v) { + KOKKOS_ROCBLAS_SAFE_CALL_IMPL( + rocsolver_sgesv(s.handle, N, nrhs, A.data(), lda, + reinterpret_cast(IPIV.data()), B.data(), + ldb, info.data())); + } + if constexpr (std::is_same_v) { + KOKKOS_ROCBLAS_SAFE_CALL_IMPL( + rocsolver_dgesv(s.handle, N, nrhs, A.data(), lda, + reinterpret_cast(IPIV.data()), B.data(), + ldb, info.data())); + } + if constexpr (std::is_same_v>) { + KOKKOS_ROCBLAS_SAFE_CALL_IMPL(rocsolver_cgesv( + s.handle, N, nrhs, reinterpret_cast(A.data()), + lda, reinterpret_cast(IPIV.data()), + reinterpret_cast(B.data()), ldb, info.data())); + } + if constexpr (std::is_same_v>) { + KOKKOS_ROCBLAS_SAFE_CALL_IMPL(rocsolver_zgesv( + s.handle, N, nrhs, reinterpret_cast(A.data()), + lda, reinterpret_cast(IPIV.data()), + reinterpret_cast(B.data()), ldb, info.data())); + } + KOKKOS_ROCBLAS_SAFE_CALL_IMPL(rocblas_set_stream(s.handle, NULL)); +} + +#define KOKKOSLAPACK_GESV_ROCSOLVER(SCALAR, LAYOUT, MEM_SPACE, ETI_SPEC_AVAIL) \ + template \ + struct GESV< \ + ExecSpace, \ + Kokkos::View, \ + Kokkos::MemoryTraits>, \ + Kokkos::View, \ + Kokkos::MemoryTraits>, \ + Kokkos::View, \ + Kokkos::MemoryTraits>, \ + true, ETI_SPEC_AVAIL> { \ + using AViewType = \ + Kokkos::View, \ + Kokkos::MemoryTraits>; \ + using BViewType = \ + Kokkos::View, \ + Kokkos::MemoryTraits>; \ + using PViewType = Kokkos::View, \ + Kokkos::MemoryTraits>; \ + \ + static void gesv(const ExecSpace& space, const AViewType& A, \ + const BViewType& B, const PViewType& IPIV) { \ + Kokkos::Profiling::pushRegion( \ + "KokkosLapack::gesv[TPL_ROCSOLVER," #SCALAR "]"); \ + gesv_print_specialization(); \ + \ + rocsolverGesvWrapper(space, IPIV, A, B); \ + Kokkos::Profiling::popRegion(); \ + } \ + }; + +KOKKOSLAPACK_GESV_ROCSOLVER(float, Kokkos::LayoutLeft, Kokkos::HIPSpace, true) +KOKKOSLAPACK_GESV_ROCSOLVER(float, Kokkos::LayoutLeft, Kokkos::HIPSpace, false) + +KOKKOSLAPACK_GESV_ROCSOLVER(double, Kokkos::LayoutLeft, Kokkos::HIPSpace, true) +KOKKOSLAPACK_GESV_ROCSOLVER(double, Kokkos::LayoutLeft, Kokkos::HIPSpace, false) + +KOKKOSLAPACK_GESV_ROCSOLVER(Kokkos::complex, Kokkos::LayoutLeft, + Kokkos::HIPSpace, true) +KOKKOSLAPACK_GESV_ROCSOLVER(Kokkos::complex, Kokkos::LayoutLeft, + Kokkos::HIPSpace, false) + +KOKKOSLAPACK_GESV_ROCSOLVER(Kokkos::complex, Kokkos::LayoutLeft, + Kokkos::HIPSpace, true) +KOKKOSLAPACK_GESV_ROCSOLVER(Kokkos::complex, Kokkos::LayoutLeft, + Kokkos::HIPSpace, false) + +} // namespace Impl +} // namespace KokkosLapack +#endif // KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER + #endif diff --git a/lapack/unit_test/Test_Lapack_gesv.hpp b/lapack/unit_test/Test_Lapack_gesv.hpp index 06f51b7eb0..e1cf743f91 100644 --- a/lapack/unit_test/Test_Lapack_gesv.hpp +++ b/lapack/unit_test/Test_Lapack_gesv.hpp @@ -16,11 +16,13 @@ // only enable this test where KokkosLapack supports gesv: // CUDA+MAGMA and HOST+LAPACK -#if (defined(TEST_CUDA_LAPACK_CPP) && \ - defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA)) || \ - (defined(KOKKOSKERNELS_ENABLE_TPL_LAPACK) && \ - (defined(TEST_OPENMP_LAPACK_CPP) || \ - defined(TEST_OPENMPTARGET_LAPACK_CPP) || \ +#if (defined(TEST_CUDA_LAPACK_CPP) && \ + defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA)) || \ + (defined(TEST_HIP_LAPACK_CPP) && \ + defined(KOKKOSKERNELS_ENABLE_TPL_ROCSOLVER)) || \ + (defined(KOKKOSKERNELS_ENABLE_TPL_LAPACK) && \ + (defined(TEST_OPENMP_LAPACK_CPP) || \ + defined(TEST_OPENMPTARGET_LAPACK_CPP) || \ defined(TEST_SERIAL_LAPACK_CPP) || defined(TEST_THREADS_LAPACK_CPP))) #include @@ -34,11 +36,13 @@ namespace Test { -template +template void impl_test_gesv(const char* mode, const char* padding, int N) { - typedef typename Device::execution_space execution_space; - typedef typename ViewTypeA::value_type ScalarA; - typedef Kokkos::ArithTraits ats; + using execution_space = typename Device::execution_space; + using ScalarA = typename ViewTypeA::value_type; + using ats = Kokkos::ArithTraits; + + execution_space space{}; Kokkos::Random_XorShift64_Pool rand_pool(13718); @@ -80,7 +84,9 @@ void impl_test_gesv(const char* mode, const char* padding, int N) { Kokkos::deep_copy(h_X0, X0); // Allocate IPIV view on host - typedef Kokkos::View ViewTypeP; + using ViewTypeP = typename std::conditional< + MAGMA, Kokkos::View, + Kokkos::View>::type; ViewTypeP ipiv; int Nt = 0; if (mode[0] == 'Y') { @@ -90,7 +96,7 @@ void impl_test_gesv(const char* mode, const char* padding, int N) { // Solve. try { - KokkosLapack::gesv(A, B, ipiv); + KokkosLapack::gesv(space, A, B, ipiv); } catch (const std::runtime_error& error) { // Check for expected runtime errors due to: // no-pivoting case (note: only MAGMA supports no-pivoting interface) @@ -138,12 +144,14 @@ void impl_test_gesv(const char* mode, const char* padding, int N) { ASSERT_EQ(test_flag, true); } -template +template void impl_test_gesv_mrhs(const char* mode, const char* padding, int N, int nrhs) { - typedef typename Device::execution_space execution_space; - typedef typename ViewTypeA::value_type ScalarA; - typedef Kokkos::ArithTraits ats; + using execution_space = typename Device::execution_space; + using ScalarA = typename ViewTypeA::value_type; + using ats = Kokkos::ArithTraits; + + execution_space space{}; Kokkos::Random_XorShift64_Pool rand_pool(13718); @@ -185,7 +193,9 @@ void impl_test_gesv_mrhs(const char* mode, const char* padding, int N, Kokkos::deep_copy(h_X0, X0); // Allocate IPIV view on host - typedef Kokkos::View ViewTypeP; + using ViewTypeP = typename std::conditional< + MAGMA, Kokkos::View, + Kokkos::View>::type; ViewTypeP ipiv; int Nt = 0; if (mode[0] == 'Y') { @@ -195,7 +205,7 @@ void impl_test_gesv_mrhs(const char* mode, const char* padding, int N, // Solve. try { - KokkosLapack::gesv(A, B, ipiv); + KokkosLapack::gesv(space, A, B, ipiv); } catch (const std::runtime_error& error) { // Check for expected runtime errors due to: // no-pivoting case (note: only MAGMA supports no-pivoting interface) @@ -253,41 +263,44 @@ int test_gesv(const char* mode) { #if defined(KOKKOSKERNELS_INST_LAYOUTLEFT) || \ (!defined(KOKKOSKERNELS_ETI_ONLY) && \ !defined(KOKKOSKERNELS_IMPL_CHECK_ETI_CALLS)) - typedef Kokkos::View view_type_a_ll; - typedef Kokkos::View view_type_b_ll; - Test::impl_test_gesv( + using view_type_a_ll = Kokkos::View; + using view_type_b_ll = Kokkos::View; + + Test::impl_test_gesv( &mode[0], "N", 2); // no padding - Test::impl_test_gesv( + Test::impl_test_gesv( &mode[0], "N", 13); // no padding - Test::impl_test_gesv( + Test::impl_test_gesv( &mode[0], "N", 179); // no padding - Test::impl_test_gesv( + Test::impl_test_gesv( &mode[0], "N", 64); // no padding - Test::impl_test_gesv( + Test::impl_test_gesv( &mode[0], "N", 1024); // no padding - Test::impl_test_gesv(&mode[0], "Y", - 13); // padding - Test::impl_test_gesv(&mode[0], "Y", - 179); // padding + +#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) && defined(KOKKOS_ENABLE_CUDA) + if constexpr (std::is_same_v) { + Test::impl_test_gesv( + &mode[0], "N", 2); // no padding + Test::impl_test_gesv( + &mode[0], "N", 13); // no padding + Test::impl_test_gesv( + &mode[0], "N", 179); // no padding + Test::impl_test_gesv( + &mode[0], "N", 64); // no padding + Test::impl_test_gesv( + &mode[0], "N", 1024); // no padding + + Test::impl_test_gesv( + &mode[0], "Y", + 13); // padding + Test::impl_test_gesv( + &mode[0], "Y", + 179); // padding + } +#endif #endif - /* - #if defined(KOKKOSKERNELS_INST_LAYOUTRIGHT) || - (!defined(KOKKOSKERNELS_ETI_ONLY) && - !defined(KOKKOSKERNELS_IMPL_CHECK_ETI_CALLS)) typedef Kokkos::View view_type_a_lr; typedef Kokkos::View view_type_b_lr; - Test::impl_test_gesv(&mode[0], "N", - 2); //no padding Test::impl_test_gesv(&mode[0], "N", 13); //no padding Test::impl_test_gesv(&mode[0], "N", 179); //no padding - Test::impl_test_gesv(&mode[0], "N", - 64); //no padding Test::impl_test_gesv(&mode[0], "N", 1024);//no padding Test::impl_test_gesv(&mode[0], "Y", 13); //padding - Test::impl_test_gesv(&mode[0], "Y", - 179); //padding #endif - */ // Supress unused parameters on CUDA10 (void)mode; return 1; @@ -298,42 +311,43 @@ int test_gesv_mrhs(const char* mode) { #if defined(KOKKOSKERNELS_INST_LAYOUTLEFT) || \ (!defined(KOKKOSKERNELS_ETI_ONLY) && \ !defined(KOKKOSKERNELS_IMPL_CHECK_ETI_CALLS)) - typedef Kokkos::View view_type_a_ll; - typedef Kokkos::View view_type_b_ll; - Test::impl_test_gesv_mrhs( + using view_type_a_ll = Kokkos::View; + using view_type_b_ll = Kokkos::View; + + Test::impl_test_gesv_mrhs( &mode[0], "N", 2, 5); // no padding - Test::impl_test_gesv_mrhs( + Test::impl_test_gesv_mrhs( &mode[0], "N", 13, 5); // no padding - Test::impl_test_gesv_mrhs( + Test::impl_test_gesv_mrhs( &mode[0], "N", 179, 5); // no padding - Test::impl_test_gesv_mrhs( + Test::impl_test_gesv_mrhs( &mode[0], "N", 64, 5); // no padding - Test::impl_test_gesv_mrhs( + Test::impl_test_gesv_mrhs( &mode[0], "N", 1024, 5); // no padding - Test::impl_test_gesv_mrhs( - &mode[0], "Y", 13, 5); // padding - Test::impl_test_gesv_mrhs( - &mode[0], "Y", 179, 5); // padding + +// When appropriate run MAGMA specific tests +#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) && defined(KOKKOS_ENABLE_CUDA) + if constexpr (std::is_same_v) { + Test::impl_test_gesv_mrhs( + &mode[0], "N", 2, 5); // no padding + Test::impl_test_gesv_mrhs( + &mode[0], "N", 13, 5); // no padding + Test::impl_test_gesv_mrhs( + &mode[0], "N", 179, 5); // no padding + Test::impl_test_gesv_mrhs( + &mode[0], "N", 64, 5); // no padding + Test::impl_test_gesv_mrhs( + &mode[0], "N", 1024, 5); // no padding + + Test::impl_test_gesv_mrhs( + &mode[0], "Y", 13, 5); // padding + Test::impl_test_gesv_mrhs( + &mode[0], "Y", 179, 5); // padding + } +#endif #endif - /* - #if defined(KOKKOSKERNELS_INST_LAYOUTRIGHT) || - (!defined(KOKKOSKERNELS_ETI_ONLY) && - !defined(KOKKOSKERNELS_IMPL_CHECK_ETI_CALLS)) typedef Kokkos::View view_type_a_lr; typedef Kokkos::View view_type_b_lr; - Test::impl_test_gesv_mrhs(&mode[0], - "N", 2, 5);//no padding Test::impl_test_gesv_mrhs(&mode[0], "N", 13, 5);//no padding - Test::impl_test_gesv_mrhs(&mode[0], - "N", 179, 5);//no padding Test::impl_test_gesv_mrhs(&mode[0], "N", 64, 5);//no padding - Test::impl_test_gesv_mrhs(&mode[0], - "N", 1024,5);//no padding Test::impl_test_gesv_mrhs(&mode[0], "Y", 13, 5);//padding - Test::impl_test_gesv_mrhs(&mode[0], - "Y", 179, 5);//padding #endif - */ // Supress unused parameters on CUDA10 (void)mode; return 1; @@ -411,4 +425,4 @@ TEST_F(TestCategory, gesv_mrhs_complex_float) { } #endif -#endif // CUDA+MAGMA or LAPACK+HOST +#endif // CUDA+MAGMA or HIP+ROCSOLVER or LAPACK+HOST diff --git a/scripts/cm_test_all_sandia b/scripts/cm_test_all_sandia index 730f3c5382..3a8079dc66 100755 --- a/scripts/cm_test_all_sandia +++ b/scripts/cm_test_all_sandia @@ -91,7 +91,7 @@ print_help() { echo "--with-tpls=TPLS: set KOKKOSKERNELS_ENABLE_TPLS" echo " Provide a comma-separated list of TPLs" echo " Valid items:" - echo " blas, mkl, cublas, cusparse, magma, armpl, rocblas, rocsparse" + echo " blas, mkl, cublas, cusparse, magma, armpl, rocblas, rocsparse, rocsolver" echo "" echo "ARGS: list of expressions matching compilers to test" @@ -1087,7 +1087,7 @@ setup_env() { export KOKKOS_CUDA_OPTIONS="${KOKKOS_CUDA_OPTIONS},enable_lambda" fi if [[ "$compiler" == rocm* ]]; then - NEW_TPL_LIST="rocblas,rocsparse," + NEW_TPL_LIST="rocblas,rocsparse,rocsolver," fi # host tpls - use mkl with intel, else use host blas if [[ "$compiler" == intel* ]]; then