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 <class AVT, class BVT>
+template <class ExecutionSpace, class AVT, class BVT, class IPIVV>
 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<SCALAR_TYPE **, LAYOUT_TYPE,                           \
                    Kokkos::Device<EXEC_SPACE_TYPE, MEM_SPACE_TYPE>,       \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,             \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,              \
       Kokkos::View<SCALAR_TYPE **, LAYOUT_TYPE,                           \
                    Kokkos::Device<EXEC_SPACE_TYPE, MEM_SPACE_TYPE>,       \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> > > {          \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,              \
+      Kokkos::View<int *, LAYOUT_TYPE,                                    \
+                   Kokkos::Device<EXEC_SPACE_TYPE, MEM_SPACE_TYPE>,       \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>> {            \
     enum : bool { value = true };                                         \
   };
 
@@ -65,20 +69,24 @@ namespace Impl {
 // Unification layer
 /// \brief Implementation of KokkosLapack::gesv.
 
-template <class AMatrix, class BXMV, class IPIVV,
-          bool tpl_spec_avail = gesv_tpl_spec_avail<AMatrix, BXMV>::value,
-          bool eti_spec_avail = gesv_eti_spec_avail<AMatrix, BXMV>::value>
+template <class ExecutionSpace, class AMatrix, class BXMV, class IPIVV,
+          bool tpl_spec_avail =
+              gesv_tpl_spec_avail<ExecutionSpace, AMatrix, BXMV, IPIVV>::value,
+          bool eti_spec_avail =
+              gesv_eti_spec_avail<ExecutionSpace, AMatrix, BXMV, IPIVV>::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 <class AMatrix, class BXMV, class IPIVV>
-struct GESV<AMatrix, BXMV, IPIVV, false, KOKKOSKERNELS_IMPL_COMPILE_LIBRARY> {
-  static void gesv(const AMatrix & /* A */, const BXMV & /* B */,
-                   const IPIVV & /* IPIV */) {
+template <class ExecutionSpace, class AMatrix, class BXMV, class IPIVV>
+struct GESV<ExecutionSpace, AMatrix, BXMV, IPIVV, false,
+            KOKKOSKERNELS_IMPL_COMPILE_LIBRARY> {
+  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<AMatrix, BXMV, IPIVV, false, KOKKOSKERNELS_IMPL_COMPILE_LIBRARY> {
 #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<SCALAR_TYPE **, LAYOUT_TYPE,                          \
                    Kokkos::Device<EXEC_SPACE_TYPE, MEM_SPACE_TYPE>,      \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,            \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,             \
       Kokkos::View<SCALAR_TYPE **, LAYOUT_TYPE,                          \
                    Kokkos::Device<EXEC_SPACE_TYPE, MEM_SPACE_TYPE>,      \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,            \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,             \
       Kokkos::View<int *, LAYOUT_TYPE,                                   \
                    Kokkos::Device<Kokkos::DefaultHostExecutionSpace,     \
                                   Kokkos::HostSpace>,                    \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,            \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,             \
       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<SCALAR_TYPE **, LAYOUT_TYPE,                          \
                    Kokkos::Device<EXEC_SPACE_TYPE, MEM_SPACE_TYPE>,      \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,            \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,             \
       Kokkos::View<SCALAR_TYPE **, LAYOUT_TYPE,                          \
                    Kokkos::Device<EXEC_SPACE_TYPE, MEM_SPACE_TYPE>,      \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,            \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,             \
       Kokkos::View<int *, LAYOUT_TYPE,                                   \
                    Kokkos::Device<Kokkos::DefaultHostExecutionSpace,     \
                                   Kokkos::HostSpace>,                    \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,            \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,             \
       false, true>;
 
 #include <KokkosLapack_gesv_tpl_spec_decl.hpp>
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 <class AMatrix, class BXMV, class IPIVV>
-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 <class ExecutionSpace, class AMatrix, class BXMV, class IPIVV>
+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<ExecutionSpace,
+                                 typename AMatrix::memory_space>::accessible);
+  static_assert(
+      Kokkos::SpaceAccessibility<ExecutionSpace,
+                                 typename BXMV::memory_space>::accessible);
+  static_assert(
+      Kokkos::SpaceAccessibility<ExecutionSpace,
+                                 typename IPIVV::memory_space>::accessible);
   static_assert(Kokkos::is_view<AMatrix>::value,
                 "KokkosLapack::gesv: A must be a Kokkos::View.");
   static_assert(Kokkos::is_view<BXMV>::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<AMatrix_Internal, BXMV_Internal,
-                             IPIVV_Internal>::gesv(A_i, B_i, IPIV_i);
+    KokkosLapack::Impl::GESV<ExecutionSpace, AMatrix_Internal, BXMV_Internal,
+                             IPIVV_Internal>::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<AMatrix_Internal, BXMV_Internal,
-                             IPIVV_Internal>::gesv(A_i, B_i, IPIV_i);
+    KokkosLapack::Impl::GESV<ExecutionSpace, AMatrix_Internal, BXMV_Internal,
+                             IPIVV_Internal>::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 <class AMatrix, class BXMV>
+template <class ExecutionSpace, class AMatrix, class BXMV, class IPIVV>
 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 <class ExecSpace>                                              \
   struct gesv_tpl_spec_avail<                                             \
+      ExecSpace,                                                          \
       Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEMSPACE>, \
                    Kokkos::MemoryTraits<Kokkos::Unmanaged> >,             \
       Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEMSPACE>, \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,             \
+      Kokkos::View<int*, LAYOUT, Kokkos::Device<ExecSpace, MEMSPACE>,     \
                    Kokkos::MemoryTraits<Kokkos::Unmanaged> > > {          \
     enum : bool { value = true };                                         \
   };
@@ -46,37 +49,22 @@ KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK(Kokkos::complex<double>,
                                         Kokkos::LayoutLeft, Kokkos::HostSpace)
 KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK(Kokkos::complex<float>,
                                         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<double>,
-Kokkos::LayoutRight, Kokkos::HostSpace) #endif
-#if defined (KOKKOSKERNELS_INST_KOKKOS_COMPLEX_FLOAT_) \
- && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT)
- KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_LAPACK( Kokkos::complex<float>,
-Kokkos::LayoutRight, Kokkos::HostSpace) #endif
-*/
 #endif
 
 // MAGMA
 #ifdef KOKKOSKERNELS_ENABLE_TPL_MAGMA
 
-#define KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA(SCALAR, LAYOUT, MEMSPACE)  \
-  template <class ExecSpace>                                              \
-  struct gesv_tpl_spec_avail<                                             \
-      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEMSPACE>, \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,             \
-      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEMSPACE>, \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> > > {          \
-    enum : bool { value = true };                                         \
+#define KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA(SCALAR, LAYOUT, MEMSPACE)     \
+  template <>                                                                \
+  struct gesv_tpl_spec_avail<                                                \
+      Kokkos::Cuda,                                                          \
+      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<Kokkos::Cuda, MEMSPACE>, \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                \
+      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<Kokkos::Cuda, MEMSPACE>, \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                \
+      Kokkos::View<int*, LAYOUT, Kokkos::HostSpace,                          \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged> > > {             \
+    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<double>,
 KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA(Kokkos::complex<float>,
                                        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<double>,Kokkos::LayoutRight, Kokkos::CudaSpace) #endif
-#if defined (KOKKOSKERNELS_INST_KOKKOS_COMPLEX_FLOAT_) \
- && defined (KOKKOSKERNELS_INST_LAYOUTRIGHT)
- KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_MAGMA( Kokkos::complex<float>,
-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<SCALAR**, LAYOUT, Kokkos::Device<Kokkos::HIP, MEMSPACE>,  \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                \
+      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<Kokkos::HIP, MEMSPACE>,  \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                \
+      Kokkos::View<int*, LAYOUT, Kokkos::Device<Kokkos::HIP, MEMSPACE>,      \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged> > > {             \
+    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<double>,
+                                           Kokkos::LayoutLeft, Kokkos::HIPSpace)
+KOKKOSLAPACK_GESV_TPL_SPEC_AVAIL_ROCSOLVER(Kokkos::complex<float>,
+                                           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 <class ExecSpace>                                                  \
-  struct GESV<                                                                \
-      Kokkos::View<double**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,    \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
-      Kokkos::View<double**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,    \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
-      Kokkos::View<int*, LAYOUT,                                              \
-                   Kokkos::Device<Kokkos::DefaultHostExecutionSpace,          \
-                                  Kokkos::HostSpace>,                         \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
-      true, ETI_SPEC_AVAIL> {                                                 \
-    typedef double SCALAR;                                                    \
-    typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
-                         Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
-        AViewType;                                                            \
-    typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
-                         Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
-        BViewType;                                                            \
-    typedef Kokkos::View<                                                     \
-        int*, LAYOUT,                                                         \
-        Kokkos::Device<Kokkos::DefaultHostExecutionSpace, Kokkos::HostSpace>, \
-        Kokkos::MemoryTraits<Kokkos::Unmanaged> >                             \
-        PViewType;                                                            \
-                                                                              \
-    static void gesv(const AViewType& A, const BViewType& B,                  \
-                     const PViewType& IPIV) {                                 \
-      Kokkos::Profiling::pushRegion("KokkosLapack::gesv[TPL_LAPACK,double]"); \
-      gesv_print_specialization<AViewType, BViewType, PViewType>();           \
-      const bool with_pivot =                                                 \
-          !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr));               \
-                                                                              \
-      const int N    = static_cast<int>(A.extent(1));                         \
-      const int AST  = static_cast<int>(A.stride(1));                         \
-      const int LDA  = (AST == 0) ? 1 : AST;                                  \
-      const int BST  = static_cast<int>(B.stride(1));                         \
-      const int LDB  = (BST == 0) ? 1 : BST;                                  \
-      const int NRHS = static_cast<int>(B.extent(1));                         \
-                                                                              \
-      int info = 0;                                                           \
-                                                                              \
-      if (with_pivot) {                                                       \
-        HostLapack<double>::gesv(N, NRHS, A.data(), LDA, IPIV.data(),         \
-                                 B.data(), LDB, info);                        \
-      }                                                                       \
-      Kokkos::Profiling::popRegion();                                         \
-    }                                                                         \
-  };
+template <class AViewType, class BViewType, class IPIVViewType>
+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 <class ExecSpace>                                                   \
-  struct GESV<                                                                 \
-      Kokkos::View<float**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,      \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                  \
-      Kokkos::View<float**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,      \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                  \
-      Kokkos::View<int*, LAYOUT,                                               \
-                   Kokkos::Device<Kokkos::DefaultHostExecutionSpace,           \
-                                  Kokkos::HostSpace>,                          \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                  \
-      true, ETI_SPEC_AVAIL> {                                                  \
-    typedef float SCALAR;                                                      \
-    typedef Kokkos::View<SCALAR**, LAYOUT,                                     \
-                         Kokkos::Device<ExecSpace, MEM_SPACE>,                 \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >             \
-        AViewType;                                                             \
-    typedef Kokkos::View<SCALAR**, LAYOUT,                                     \
-                         Kokkos::Device<ExecSpace, MEM_SPACE>,                 \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >             \
-        BViewType;                                                             \
-    typedef Kokkos::View<                                                      \
-        int*, LAYOUT,                                                          \
-        Kokkos::Device<Kokkos::DefaultHostExecutionSpace, Kokkos::HostSpace>,  \
-        Kokkos::MemoryTraits<Kokkos::Unmanaged> >                              \
-        PViewType;                                                             \
-                                                                               \
-    static void gesv(const AViewType& A, const BViewType& B,                   \
-                     const PViewType& IPIV) {                                  \
-      Kokkos::Profiling::pushRegion("KokkosLapack::gesv[TPL_LAPACK,float]");   \
-      gesv_print_specialization<AViewType, BViewType, PViewType>();            \
-      const bool with_pivot =                                                  \
-          !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr));                \
-                                                                               \
-      const int N    = static_cast<int>(A.extent(1));                          \
-      const int AST  = static_cast<int>(A.stride(1));                          \
-      const int LDA  = (AST == 0) ? 1 : AST;                                   \
-      const int BST  = static_cast<int>(B.stride(1));                          \
-      const int LDB  = (BST == 0) ? 1 : BST;                                   \
-      const int NRHS = static_cast<int>(B.extent(1));                          \
-                                                                               \
-      int info = 0;                                                            \
-                                                                               \
-      if (with_pivot) {                                                        \
-        HostLapack<float>::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 <class ExecSpace>                                                  \
-  struct GESV<Kokkos::View<Kokkos::complex<double>**, LAYOUT,                 \
-                           Kokkos::Device<ExecSpace, MEM_SPACE>,              \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
-              Kokkos::View<Kokkos::complex<double>**, LAYOUT,                 \
-                           Kokkos::Device<ExecSpace, MEM_SPACE>,              \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
-              Kokkos::View<int*, LAYOUT,                                      \
-                           Kokkos::Device<Kokkos::DefaultHostExecutionSpace,  \
-                                          Kokkos::HostSpace>,                 \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
-              true, ETI_SPEC_AVAIL> {                                         \
-    typedef Kokkos::complex<double> SCALAR;                                   \
-    typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
-                         Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
-        AViewType;                                                            \
-    typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
-                         Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
-        BViewType;                                                            \
-    typedef Kokkos::View<                                                     \
-        int*, LAYOUT,                                                         \
-        Kokkos::Device<Kokkos::DefaultHostExecutionSpace, Kokkos::HostSpace>, \
-        Kokkos::MemoryTraits<Kokkos::Unmanaged> >                             \
-        PViewType;                                                            \
-                                                                              \
-    static void gesv(const AViewType& A, const BViewType& B,                  \
-                     const PViewType& IPIV) {                                 \
-      Kokkos::Profiling::pushRegion(                                          \
-          "KokkosLapack::gesv[TPL_LAPACK,complex<double>]");                  \
-      gesv_print_specialization<AViewType, BViewType, PViewType>();           \
-      const bool with_pivot =                                                 \
-          !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr));               \
-                                                                              \
-      const int N    = static_cast<int>(A.extent(1));                         \
-      const int AST  = static_cast<int>(A.stride(1));                         \
-      const int LDA  = (AST == 0) ? 1 : AST;                                  \
-      const int BST  = static_cast<int>(B.stride(1));                         \
-      const int LDB  = (BST == 0) ? 1 : BST;                                  \
-      const int NRHS = static_cast<int>(B.extent(1));                         \
-                                                                              \
-      int info = 0;                                                           \
-                                                                              \
-      if (with_pivot) {                                                       \
-        HostLapack<std::complex<double> >::gesv(                              \
-            N, NRHS, reinterpret_cast<std::complex<double>*>(A.data()), LDA,  \
-            IPIV.data(), reinterpret_cast<std::complex<double>*>(B.data()),   \
-            LDB, info);                                                       \
-      }                                                                       \
-      Kokkos::Profiling::popRegion();                                         \
-    }                                                                         \
-  };
+  const int N    = static_cast<int>(A.extent(1));
+  const int AST  = static_cast<int>(A.stride(1));
+  const int LDA  = (AST == 0) ? 1 : AST;
+  const int BST  = static_cast<int>(B.stride(1));
+  const int LDB  = (BST == 0) ? 1 : BST;
+  const int NRHS = static_cast<int>(B.extent(1));
 
-#define KOKKOSLAPACK_CGESV_LAPACK(LAYOUT, MEM_SPACE, ETI_SPEC_AVAIL)          \
-  template <class ExecSpace>                                                  \
-  struct GESV<Kokkos::View<Kokkos::complex<float>**, LAYOUT,                  \
-                           Kokkos::Device<ExecSpace, MEM_SPACE>,              \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
-              Kokkos::View<Kokkos::complex<float>**, LAYOUT,                  \
-                           Kokkos::Device<ExecSpace, MEM_SPACE>,              \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
-              Kokkos::View<int*, LAYOUT,                                      \
-                           Kokkos::Device<Kokkos::DefaultHostExecutionSpace,  \
-                                          Kokkos::HostSpace>,                 \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
-              true, ETI_SPEC_AVAIL> {                                         \
-    typedef Kokkos::complex<float> SCALAR;                                    \
-    typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
-                         Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
-        AViewType;                                                            \
-    typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
-                         Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
-        BViewType;                                                            \
-    typedef Kokkos::View<                                                     \
-        int*, LAYOUT,                                                         \
-        Kokkos::Device<Kokkos::DefaultHostExecutionSpace, Kokkos::HostSpace>, \
-        Kokkos::MemoryTraits<Kokkos::Unmanaged> >                             \
-        PViewType;                                                            \
-                                                                              \
-    static void gesv(const AViewType& A, const BViewType& B,                  \
-                     const PViewType& IPIV) {                                 \
-      Kokkos::Profiling::pushRegion(                                          \
-          "KokkosLapack::gesv[TPL_LAPACK,complex<float>]");                   \
-      gesv_print_specialization<AViewType, BViewType, PViewType>();           \
-      const bool with_pivot =                                                 \
-          !((IPIV.extent(0) == 0) && (IPIV.data() == nullptr));               \
-                                                                              \
-      const int N    = static_cast<int>(A.extent(1));                         \
-      const int AST  = static_cast<int>(A.stride(1));                         \
-      const int LDA  = (AST == 0) ? 1 : AST;                                  \
-      const int BST  = static_cast<int>(B.stride(1));                         \
-      const int LDB  = (BST == 0) ? 1 : BST;                                  \
-      const int NRHS = static_cast<int>(B.extent(1));                         \
-                                                                              \
-      int info = 0;                                                           \
-                                                                              \
-      if (with_pivot) {                                                       \
-        HostLapack<std::complex<float> >::gesv(                               \
-            N, NRHS, reinterpret_cast<std::complex<float>*>(A.data()), LDA,   \
-            IPIV.data(), reinterpret_cast<std::complex<float>*>(B.data()),    \
-            LDB, info);                                                       \
-      }                                                                       \
-      Kokkos::Profiling::popRegion();                                         \
-    }                                                                         \
+  int info = 0;
+
+  if (with_pivot) {
+    if constexpr (Kokkos::ArithTraits<Scalar>::is_complex) {
+      using MagType = typename Kokkos::ArithTraits<Scalar>::mag_type;
+
+      HostLapack<std::complex<MagType>>::gesv(
+          N, NRHS, reinterpret_cast<std::complex<MagType>*>(A.data()), LDA,
+          IPIV.data(), reinterpret_cast<std::complex<MagType>*>(B.data()), LDB,
+          info);
+    } else {
+      HostLapack<Scalar>::gesv(N, NRHS, A.data(), LDA, IPIV.data(), B.data(),
+                               LDB, info);
+    }
+  }
+}
+
+#define KOKKOSLAPACK_GESV_LAPACK(SCALAR, LAYOUT, MEM_SPACE, ETI_SPEC_AVAIL)  \
+  template <class ExecSpace>                                                 \
+  struct GESV<                                                               \
+      ExecSpace,                                                             \
+      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,   \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                 \
+      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,   \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                 \
+      Kokkos::View<int*, LAYOUT,                                             \
+                   Kokkos::Device<Kokkos::DefaultHostExecutionSpace,         \
+                                  Kokkos::HostSpace>,                        \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                 \
+      true, ETI_SPEC_AVAIL> {                                                \
+    using AViewType =                                                        \
+        Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>, \
+                     Kokkos::MemoryTraits<Kokkos::Unmanaged>>;               \
+    using BViewType =                                                        \
+        Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>, \
+                     Kokkos::MemoryTraits<Kokkos::Unmanaged>>;               \
+    using PViewType = Kokkos::View<int*, LAYOUT, Kokkos::HostSpace,          \
+                                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>; \
+                                                                             \
+    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<AViewType, BViewType, PViewType>();          \
+      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<float>, Kokkos::LayoutLeft,
+                         Kokkos::HostSpace, true)
+KOKKOSLAPACK_GESV_LAPACK(Kokkos::complex<float>, 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<double>, Kokkos::LayoutLeft,
+                         Kokkos::HostSpace, true)
+KOKKOSLAPACK_GESV_LAPACK(Kokkos::complex<double>, Kokkos::LayoutLeft,
+                         Kokkos::HostSpace, false)
 
 }  // namespace Impl
 }  // namespace KokkosLapack
@@ -284,27 +139,27 @@ namespace Impl {
   template <class ExecSpace>                                                  \
   struct GESV<                                                                \
       Kokkos::View<double**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,    \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                  \
       Kokkos::View<double**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,    \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                  \
       Kokkos::View<magma_int_t*, LAYOUT,                                      \
                    Kokkos::Device<Kokkos::DefaultHostExecutionSpace,          \
                                   Kokkos::HostSpace>,                         \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                  \
       true, ETI_SPEC_AVAIL> {                                                 \
     typedef double SCALAR;                                                    \
     typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
                          Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
+                         Kokkos::MemoryTraits<Kokkos::Unmanaged>>             \
         AViewType;                                                            \
     typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
                          Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
+                         Kokkos::MemoryTraits<Kokkos::Unmanaged>>             \
         BViewType;                                                            \
     typedef Kokkos::View<                                                     \
         magma_int_t*, LAYOUT,                                                 \
         Kokkos::Device<Kokkos::DefaultHostExecutionSpace, Kokkos::HostSpace>, \
-        Kokkos::MemoryTraits<Kokkos::Unmanaged> >                             \
+        Kokkos::MemoryTraits<Kokkos::Unmanaged>>                              \
         PViewType;                                                            \
                                                                               \
     static void gesv(const AViewType& A, const BViewType& B,                  \
@@ -343,27 +198,27 @@ namespace Impl {
   template <class ExecSpace>                                                  \
   struct GESV<                                                                \
       Kokkos::View<float**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,     \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                  \
       Kokkos::View<float**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,     \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                  \
       Kokkos::View<magma_int_t*, LAYOUT,                                      \
                    Kokkos::Device<Kokkos::DefaultHostExecutionSpace,          \
                                   Kokkos::HostSpace>,                         \
-                   Kokkos::MemoryTraits<Kokkos::Unmanaged> >,                 \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                  \
       true, ETI_SPEC_AVAIL> {                                                 \
     typedef float SCALAR;                                                     \
     typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
                          Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
+                         Kokkos::MemoryTraits<Kokkos::Unmanaged>>             \
         AViewType;                                                            \
     typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
                          Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
+                         Kokkos::MemoryTraits<Kokkos::Unmanaged>>             \
         BViewType;                                                            \
     typedef Kokkos::View<                                                     \
         magma_int_t*, LAYOUT,                                                 \
         Kokkos::Device<Kokkos::DefaultHostExecutionSpace, Kokkos::HostSpace>, \
-        Kokkos::MemoryTraits<Kokkos::Unmanaged> >                             \
+        Kokkos::MemoryTraits<Kokkos::Unmanaged>>                              \
         PViewType;                                                            \
                                                                               \
     static void gesv(const AViewType& A, const BViewType& B,                  \
@@ -402,28 +257,28 @@ namespace Impl {
   template <class ExecSpace>                                                  \
   struct GESV<Kokkos::View<Kokkos::complex<double>**, LAYOUT,                 \
                            Kokkos::Device<ExecSpace, MEM_SPACE>,              \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
+                           Kokkos::MemoryTraits<Kokkos::Unmanaged>>,          \
               Kokkos::View<Kokkos::complex<double>**, LAYOUT,                 \
                            Kokkos::Device<ExecSpace, MEM_SPACE>,              \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
+                           Kokkos::MemoryTraits<Kokkos::Unmanaged>>,          \
               Kokkos::View<magma_int_t*, LAYOUT,                              \
                            Kokkos::Device<Kokkos::DefaultHostExecutionSpace,  \
                                           Kokkos::HostSpace>,                 \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
+                           Kokkos::MemoryTraits<Kokkos::Unmanaged>>,          \
               true, ETI_SPEC_AVAIL> {                                         \
     typedef Kokkos::complex<double> SCALAR;                                   \
     typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
                          Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
+                         Kokkos::MemoryTraits<Kokkos::Unmanaged>>             \
         AViewType;                                                            \
     typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
                          Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
+                         Kokkos::MemoryTraits<Kokkos::Unmanaged>>             \
         BViewType;                                                            \
     typedef Kokkos::View<                                                     \
         magma_int_t*, LAYOUT,                                                 \
         Kokkos::Device<Kokkos::DefaultHostExecutionSpace, Kokkos::HostSpace>, \
-        Kokkos::MemoryTraits<Kokkos::Unmanaged> >                             \
+        Kokkos::MemoryTraits<Kokkos::Unmanaged>>                              \
         PViewType;                                                            \
                                                                               \
     static void gesv(const AViewType& A, const BViewType& B,                  \
@@ -463,28 +318,28 @@ namespace Impl {
   template <class ExecSpace>                                                  \
   struct GESV<Kokkos::View<Kokkos::complex<float>**, LAYOUT,                  \
                            Kokkos::Device<ExecSpace, MEM_SPACE>,              \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
+                           Kokkos::MemoryTraits<Kokkos::Unmanaged>>,          \
               Kokkos::View<Kokkos::complex<float>**, LAYOUT,                  \
                            Kokkos::Device<ExecSpace, MEM_SPACE>,              \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
+                           Kokkos::MemoryTraits<Kokkos::Unmanaged>>,          \
               Kokkos::View<magma_int_t*, LAYOUT,                              \
                            Kokkos::Device<Kokkos::DefaultHostExecutionSpace,  \
                                           Kokkos::HostSpace>,                 \
-                           Kokkos::MemoryTraits<Kokkos::Unmanaged> >,         \
+                           Kokkos::MemoryTraits<Kokkos::Unmanaged>>,          \
               true, ETI_SPEC_AVAIL> {                                         \
     typedef Kokkos::complex<float> SCALAR;                                    \
     typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
                          Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
+                         Kokkos::MemoryTraits<Kokkos::Unmanaged>>             \
         AViewType;                                                            \
     typedef Kokkos::View<SCALAR**, LAYOUT,                                    \
                          Kokkos::Device<ExecSpace, MEM_SPACE>,                \
-                         Kokkos::MemoryTraits<Kokkos::Unmanaged> >            \
+                         Kokkos::MemoryTraits<Kokkos::Unmanaged>>             \
         BViewType;                                                            \
     typedef Kokkos::View<                                                     \
         magma_int_t*, LAYOUT,                                                 \
         Kokkos::Device<Kokkos::DefaultHostExecutionSpace, Kokkos::HostSpace>, \
-        Kokkos::MemoryTraits<Kokkos::Unmanaged> >                             \
+        Kokkos::MemoryTraits<Kokkos::Unmanaged>>                              \
         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 <KokkosBlas_tpl_spec.hpp>
+#include <rocsolver/rocsolver.h>
+
+namespace KokkosLapack {
+namespace Impl {
+
+template <class ExecutionSpace, class IPIVViewType, class AViewType,
+          class BViewType>
+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<rocblas_int>(A.extent(0));
+  const rocblas_int nrhs = static_cast<rocblas_int>(B.extent(1));
+  const rocblas_int lda  = std::is_same_v<ALayout_t, Kokkos::LayoutRight>
+                              ? A.stride(0)
+                              : A.stride(1);
+  const rocblas_int ldb = std::is_same_v<BLayout_t, Kokkos::LayoutRight>
+                              ? B.stride(0)
+                              : B.stride(1);
+  Kokkos::View<rocblas_int, ExecutionSpace> 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<Scalar, float>) {
+    KOKKOS_ROCBLAS_SAFE_CALL_IMPL(
+        rocsolver_sgesv(s.handle, N, nrhs, A.data(), lda,
+                        reinterpret_cast<rocblas_int*>(IPIV.data()), B.data(),
+                        ldb, info.data()));
+  }
+  if constexpr (std::is_same_v<Scalar, double>) {
+    KOKKOS_ROCBLAS_SAFE_CALL_IMPL(
+        rocsolver_dgesv(s.handle, N, nrhs, A.data(), lda,
+                        reinterpret_cast<rocblas_int*>(IPIV.data()), B.data(),
+                        ldb, info.data()));
+  }
+  if constexpr (std::is_same_v<Scalar, Kokkos::complex<float>>) {
+    KOKKOS_ROCBLAS_SAFE_CALL_IMPL(rocsolver_cgesv(
+        s.handle, N, nrhs, reinterpret_cast<rocblas_float_complex*>(A.data()),
+        lda, reinterpret_cast<rocblas_int*>(IPIV.data()),
+        reinterpret_cast<rocblas_float_complex*>(B.data()), ldb, info.data()));
+  }
+  if constexpr (std::is_same_v<Scalar, Kokkos::complex<double>>) {
+    KOKKOS_ROCBLAS_SAFE_CALL_IMPL(rocsolver_zgesv(
+        s.handle, N, nrhs, reinterpret_cast<rocblas_double_complex*>(A.data()),
+        lda, reinterpret_cast<rocblas_int*>(IPIV.data()),
+        reinterpret_cast<rocblas_double_complex*>(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 <class ExecSpace>                                                   \
+  struct GESV<                                                                 \
+      ExecSpace,                                                               \
+      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,     \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                   \
+      Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,     \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                   \
+      Kokkos::View<rocblas_int*, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>, \
+                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>,                   \
+      true, ETI_SPEC_AVAIL> {                                                  \
+    using AViewType =                                                          \
+        Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,   \
+                     Kokkos::MemoryTraits<Kokkos::Unmanaged>>;                 \
+    using BViewType =                                                          \
+        Kokkos::View<SCALAR**, LAYOUT, Kokkos::Device<ExecSpace, MEM_SPACE>,   \
+                     Kokkos::MemoryTraits<Kokkos::Unmanaged>>;                 \
+    using PViewType = Kokkos::View<rocblas_int*, LAYOUT,                       \
+                                   Kokkos::Device<ExecSpace, MEM_SPACE>,       \
+                                   Kokkos::MemoryTraits<Kokkos::Unmanaged>>;   \
+                                                                               \
+    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<AViewType, BViewType, PViewType>();            \
+                                                                               \
+      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<float>, Kokkos::LayoutLeft,
+                            Kokkos::HIPSpace, true)
+KOKKOSLAPACK_GESV_ROCSOLVER(Kokkos::complex<float>, Kokkos::LayoutLeft,
+                            Kokkos::HIPSpace, false)
+
+KOKKOSLAPACK_GESV_ROCSOLVER(Kokkos::complex<double>, Kokkos::LayoutLeft,
+                            Kokkos::HIPSpace, true)
+KOKKOSLAPACK_GESV_ROCSOLVER(Kokkos::complex<double>, 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 <gtest/gtest.h>
@@ -34,11 +36,13 @@
 
 namespace Test {
 
-template <class ViewTypeA, class ViewTypeB, class Device>
+template <class ViewTypeA, class ViewTypeB, class Device, bool MAGMA>
 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<ScalarA> ats;
+  using execution_space = typename Device::execution_space;
+  using ScalarA         = typename ViewTypeA::value_type;
+  using ats             = Kokkos::ArithTraits<ScalarA>;
+
+  execution_space space{};
 
   Kokkos::Random_XorShift64_Pool<execution_space> 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<int*, Kokkos::LayoutLeft, Kokkos::HostSpace> ViewTypeP;
+  using ViewTypeP = typename std::conditional<
+      MAGMA, Kokkos::View<int*, Kokkos::LayoutLeft, Kokkos::HostSpace>,
+      Kokkos::View<int*, Kokkos::LayoutLeft, execution_space>>::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 <class ViewTypeA, class ViewTypeB, class Device>
+template <class ViewTypeA, class ViewTypeB, class Device, bool MAGMA>
 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<ScalarA> ats;
+  using execution_space = typename Device::execution_space;
+  using ScalarA         = typename ViewTypeA::value_type;
+  using ats             = Kokkos::ArithTraits<ScalarA>;
+
+  execution_space space{};
 
   Kokkos::Random_XorShift64_Pool<execution_space> 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<int*, Kokkos::LayoutLeft, Kokkos::HostSpace> ViewTypeP;
+  using ViewTypeP = typename std::conditional<
+      MAGMA, Kokkos::View<int*, Kokkos::LayoutLeft, Kokkos::HostSpace>,
+      Kokkos::View<int*, Kokkos::LayoutLeft, execution_space>>::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<Scalar**, Kokkos::LayoutLeft, Device> view_type_a_ll;
-  typedef Kokkos::View<Scalar*, Kokkos::LayoutLeft, Device> view_type_b_ll;
-  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device>(
+  using view_type_a_ll = Kokkos::View<Scalar**, Kokkos::LayoutLeft, Device>;
+  using view_type_b_ll = Kokkos::View<Scalar*, Kokkos::LayoutLeft, Device>;
+
+  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 2);  // no padding
-  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device>(
+  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 13);  // no padding
-  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device>(
+  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 179);  // no padding
-  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device>(
+  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 64);  // no padding
-  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device>(
+  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 1024);  // no padding
-  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device>(&mode[0], "Y",
-                                                               13);  // padding
-  Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device>(&mode[0], "Y",
-                                                               179);  // padding
+
+#if defined(KOKKOSKERNELS_ENABLE_TPL_MAGMA) && defined(KOKKOS_ENABLE_CUDA)
+  if constexpr (std::is_same_v<Kokkos::Cuda,
+                               typename Device::execution_space>) {
+    Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 2);  // no padding
+    Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 13);  // no padding
+    Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 179);  // no padding
+    Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 64);  // no padding
+    Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 1024);  // no padding
+
+    Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "Y",
+        13);  // padding
+    Test::impl_test_gesv<view_type_a_ll, view_type_b_ll, Device, true>(
+        &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<ScalarA**,
-  Kokkos::LayoutRight, Device> view_type_a_lr; typedef Kokkos::View<ScalarB*,
-  Kokkos::LayoutRight, Device> view_type_b_lr;
-    Test::impl_test_gesv<view_type_a_lr, view_type_b_lr, Device>(&mode[0], "N",
-  2);   //no padding Test::impl_test_gesv<view_type_a_lr, view_type_b_lr,
-  Device>(&mode[0], "N", 13);  //no padding Test::impl_test_gesv<view_type_a_lr,
-  view_type_b_lr, Device>(&mode[0], "N", 179); //no padding
-    Test::impl_test_gesv<view_type_a_lr, view_type_b_lr, Device>(&mode[0], "N",
-  64);  //no padding Test::impl_test_gesv<view_type_a_lr, view_type_b_lr,
-  Device>(&mode[0], "N", 1024);//no padding Test::impl_test_gesv<view_type_a_lr,
-  view_type_b_lr, Device>(&mode[0], "Y", 13);  //padding
-    Test::impl_test_gesv<view_type_a_lr, view_type_b_lr, Device>(&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<Scalar**, Kokkos::LayoutLeft, Device> view_type_a_ll;
-  typedef Kokkos::View<Scalar**, Kokkos::LayoutLeft, Device> view_type_b_ll;
-  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device>(
+  using view_type_a_ll = Kokkos::View<Scalar**, Kokkos::LayoutLeft, Device>;
+  using view_type_b_ll = Kokkos::View<Scalar**, Kokkos::LayoutLeft, Device>;
+
+  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 2, 5);  // no padding
-  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device>(
+  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 13, 5);  // no padding
-  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device>(
+  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 179, 5);  // no padding
-  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device>(
+  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 64, 5);  // no padding
-  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device>(
+  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, false>(
       &mode[0], "N", 1024, 5);  // no padding
-  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device>(
-      &mode[0], "Y", 13, 5);  // padding
-  Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device>(
-      &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<Kokkos::Cuda,
+                               typename Device::execution_space>) {
+    Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 2, 5);  // no padding
+    Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 13, 5);  // no padding
+    Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 179, 5);  // no padding
+    Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 64, 5);  // no padding
+    Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "N", 1024, 5);  // no padding
+
+    Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, true>(
+        &mode[0], "Y", 13, 5);  // padding
+    Test::impl_test_gesv_mrhs<view_type_a_ll, view_type_b_ll, Device, true>(
+        &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<ScalarA**,
-  Kokkos::LayoutRight, Device> view_type_a_lr; typedef Kokkos::View<ScalarB**,
-  Kokkos::LayoutRight, Device> view_type_b_lr;
-    Test::impl_test_gesv_mrhs<view_type_a_lr, view_type_b_lr, Device>(&mode[0],
-  "N", 2,   5);//no padding Test::impl_test_gesv_mrhs<view_type_a_lr,
-  view_type_b_lr, Device>(&mode[0], "N", 13,  5);//no padding
-    Test::impl_test_gesv_mrhs<view_type_a_lr, view_type_b_lr, Device>(&mode[0],
-  "N", 179, 5);//no padding Test::impl_test_gesv_mrhs<view_type_a_lr,
-  view_type_b_lr, Device>(&mode[0], "N", 64,  5);//no padding
-    Test::impl_test_gesv_mrhs<view_type_a_lr, view_type_b_lr, Device>(&mode[0],
-  "N", 1024,5);//no padding Test::impl_test_gesv_mrhs<view_type_a_lr,
-  view_type_b_lr, Device>(&mode[0], "Y", 13,  5);//padding
-    Test::impl_test_gesv_mrhs<view_type_a_lr, view_type_b_lr, Device>(&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