Add hipsolver bindings and enable lobpcg kernel tests for HIP

Author: nbeams

cmake/hip.cmake

@@ -33,6 +33,9 @@ find_package(hipsparse REQUIRED)
 find_package(rocrand REQUIRED)
 find_package(rocthrust REQUIRED)
 find_package(ROCTX)
+if(GINKGO_BUILD_LAPACK)
+    find_package(hipsolver REQUIRED)
+endif()
 
 if(GINKGO_HIP_AMD_UNSAFE_ATOMIC AND GINKGO_HIP_VERSION VERSION_GREATER_EQUAL 5)
     set(CMAKE_HIP_FLAGS

common/cuda_hip/base/dev_lapack_bindings.hpp

@@ -8,8 +8,12 @@
 
 #if defined(GKO_COMPILING_CUDA)
 #include "cuda/base/cusolver_bindings.hpp"
+#define GKO_DEV_LAPACK_ERROR GKO_CUSOLVER_ERROR
+#define DEV_LAPACK_INTERNAL_ERROR CUSOLVER_STATUS_INTERNAL_ERROR
 #elif defined(GKO_COMPILING_HIP)
 #include "hip/base/hipsolver_bindings.hip.hpp"
+#define GKO_DEV_LAPACK_ERROR GKO_HIPSOLVER_ERROR
+#define DEV_LAPACK_INTERNAL_ERROR HIPSOLVER_STATUS_INTERNAL_ERROR
 #else
 #error "Executor definition missing"
 #endif

common/cuda_hip/eigensolver/lobpcg_kernels.cpp

@@ -125,7 +125,7 @@ void symm_eig(std::shared_ptr<const DefaultExecutor> exec,
 
         int32 host_info = exec->copy_val_to_host(dev_info.get_data());
         if (host_info != 0) {
-            throw GKO_CUSOLVER_ERROR(CUSOLVER_STATUS_INTERNAL_ERROR);
+            throw GKO_DEV_LAPACK_ERROR(DEV_LAPACK_INTERNAL_ERROR);
         }
     } catch (std::exception& e) {
         std::cout << e.what() << std::endl;
@@ -196,7 +196,7 @@ void symm_generalized_eig(std::shared_ptr<const DefaultExecutor> exec,
 
         int32 host_info = exec->copy_val_to_host(dev_info.get_data());
         if (host_info != 0) {
-            throw GKO_CUSOLVER_ERROR(CUSOLVER_STATUS_INTERNAL_ERROR);
+            throw GKO_DEV_LAPACK_ERROR(DEV_LAPACK_INTERNAL_ERROR);
         }
     } catch (std::exception& e) {
         std::cout << e.what() << std::endl;
@@ -260,7 +260,7 @@ void b_orthonormalize(std::shared_ptr<const DefaultExecutor> exec,
 
         int32 host_info = exec->copy_val_to_host(dev_info.get_data());
         if (host_info != 0) {
-            throw GKO_CUSOLVER_ERROR(CUSOLVER_STATUS_INTERNAL_ERROR);
+            throw GKO_DEV_LAPACK_ERROR(DEV_LAPACK_INTERNAL_ERROR);
         }
     } catch (std::exception& e) {
         std::cout << e.what() << std::endl;
@@ -287,7 +287,7 @@ void b_orthonormalize(std::shared_ptr<const DefaultExecutor> exec,
 
         int32 host_info = exec->copy_val_to_host(dev_info.get_data());
         if (host_info != 0) {
-            throw GKO_CUSOLVER_ERROR(CUSOLVER_STATUS_INTERNAL_ERROR);
+            throw GKO_DEV_LAPACK_ERROR(DEV_LAPACK_INTERNAL_ERROR);
         }
     } catch (std::exception& e) {
         std::cout << e.what() << std::endl;

core/device_hooks/hip_hooks.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -174,6 +174,12 @@ std::string HipsparseError::get_error(int64)
 }
 
 
+std::string HipsolverError::get_error(int64)
+{
+    return "ginkgo HIP module is not compiled";
+}
+
+
 std::string HipfftError::get_error(int64)
 {
     return "ginkgo HIP module is not compiled";

core/test/base/exception.cpp

@@ -127,6 +127,16 @@ TEST(ExceptionClasses, HipsparseErrorReturnsCorrectWhatMessage)
 }
 
 
+#if GKO_HAVE_LAPACK
+TEST(ExceptionClasses, HipsolverErrorReturnsCorrectWhatMessage)
+{
+    gko::HipsolverError error("test_file.cpp", 123, "test_func", 1);
+    std::string expected = "test_file.cpp:123: test_func: ";
+    ASSERT_EQ(expected, std::string(error.what()).substr(0, expected.size()));
+}
+#endif
+
+
 TEST(ExceptionClasses, HipfftErrorReturnsCorrectWhatMessage)
 {
     gko::HipfftError error("test_file.cpp", 123, "test_func", 1);

core/test/base/exception_helpers.cpp

@@ -153,6 +153,16 @@ TEST(HipError, ReturnsHipsparseError)
 }
 
 
+#if GKO_HAVE_LAPACK
+void throws_hipsolver_error() { throw GKO_HIPSOLVER_ERROR(0); }
+
+TEST(HipError, ReturnsHipsolverError)
+{
+    ASSERT_THROW(throws_hipsolver_error(), gko::HipsolverError);
+}
+#endif
+
+
 void throws_hipfft_error() { throw GKO_HIPFFT_ERROR(0); }
 
 TEST(HipError, ReturnsHipfftError)

hip/CMakeLists.txt

@@ -120,6 +120,9 @@ endif()
 if(GINKGO_HAVE_ROCTX)
     target_link_libraries(ginkgo_hip PRIVATE roc::roctx)
 endif()
+if(GINKGO_BUILD_LAPACK)
+    target_link_libraries(ginkgo_hip PRIVATE roc::hipsolver)
+endif()
 
 target_compile_options(
     ginkgo_hip

hip/base/exception.hip.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -11,10 +11,16 @@
 #include <hipblas/hipblas.h>
 #include <hiprand/hiprand.h>
 #include <hipsparse/hipsparse.h>
+#if GKO_HAVE_LAPACK
+#include <hipsolver/hipsolver.h>
+#endif
 #else
 #include <hipblas.h>
 #include <hiprand.h>
 #include <hipsparse.h>
+#if GKO_HAVE_LAPACK
+#include <hipsolver.h>
+#endif
 #endif
 
 
@@ -107,4 +113,29 @@ std::string HipsparseError::get_error(int64 error_code)
 }
 
 
+std::string HipsolverError::get_error(int64 error_code)
+{
+#if GKO_HAVE_LAPACK
+#define GKO_REGISTER_HIPSOLVER_ERROR(error_name)        \
+    if (error_code == static_cast<int64>(error_name)) { \
+        return #error_name;                             \
+    }
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_SUCCESS);
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_NOT_INITIALIZED);
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_ALLOC_FAILED);
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_INVALID_VALUE);
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_ARCH_MISMATCH);
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_EXECUTION_FAILED);
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_INTERNAL_ERROR);
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_MATRIX_TYPE_NOT_SUPPORTED);
+    GKO_REGISTER_HIPSOLVER_ERROR(HIPSOLVER_STATUS_NOT_SUPPORTED);
+    return "Unknown error";
+
+#undef GKO_REGISTER_HIPSOLVER_ERROR
+#else
+    return "Ginkgo must be built with LAPACK support to enable hipSOLVER";
+#endif
+}
+
+
 }  // namespace gko

hip/base/executor.hip.cpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -14,6 +14,9 @@
 #include "common/cuda_hip/base/runtime.hpp"
 #include "hip/base/device.hpp"
 #include "hip/base/hipblas_handle.hpp"
+#if GKO_HAVE_LAPACK
+#include "hip/base/hipsolver_handle.hpp"
+#endif
 #include "hip/base/hipsparse_handle.hpp"
 #include "hip/base/scoped_device_id.hip.hpp"
 
@@ -260,6 +263,14 @@ void HipExecutor::init_handles()
                 detail::hip_scoped_device_id_guard g(id);
                 kernels::hip::hipsparse::destroy_hipsparse_handle(handle);
             });
+#if GKO_HAVE_LAPACK
+        this->hipsolver_handle_ = handle_manager<hipsolverDnContext>(
+            kernels::hip::hipsolver::init(this->get_stream()),
+            [id](hipsolverDnContext* handle) {
+                detail::hip_scoped_device_id_guard g(id);
+                kernels::hip::hipsolver::destroy_hipsolver_handle(handle);
+            });
+#endif
     }
 }
 

hip/base/hipblas_bindings.hip.hpp

@@ -1,4 +1,4 @@
-// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+// SPDX-FileCopyrightText: 2017 - 2025 The Ginkgo authors
 //
 // SPDX-License-Identifier: BSD-3-Clause
 
@@ -128,6 +128,32 @@ GKO_BIND_HIPBLAS_GEAM(ValueType, detail::not_implemented);
 #undef GKO_BIND_HIPBLAS_GEAM
 
 
+#define GKO_BIND_HIPBLAS_TRMM(ValueType, HipblasName)                        \
+    inline void trmm(hipblasHandle_t handle, hipblasSideMode_t side,         \
+                     hipblasFillMode_t uplo, hipblasOperation_t trans,       \
+                     hipblasDiagType_t diag, int m, int n,                   \
+                     const ValueType* alpha, const ValueType* a, int lda,    \
+                     const ValueType* b, int ldb, ValueType* c, int ldc)     \
+    {                                                                        \
+        GKO_ASSERT_NO_HIPBLAS_ERRORS(                                        \
+            HipblasName(handle, side, uplo, trans, diag, m, n,               \
+                        as_hipblas_type(alpha), as_hipblas_type(a), lda,     \
+                        as_hipblas_type(b), ldb, as_hipblas_type(c), ldc));  \
+    }                                                                        \
+    static_assert(true,                                                      \
+                  "This assert is used to counter the false positive extra " \
+                  "semi-colon warnings")
+
+GKO_BIND_HIPBLAS_TRMM(float, hipblasStrmm);
+GKO_BIND_HIPBLAS_TRMM(double, hipblasDtrmm);
+GKO_BIND_HIPBLAS_TRMM(std::complex<float>, hipblasCtrmm);
+GKO_BIND_HIPBLAS_TRMM(std::complex<double>, hipblasZtrmm);
+template <typename ValueType>
+GKO_BIND_HIPBLAS_TRMM(ValueType, detail::not_implemented);
+
+#undef GKO_BIND_HIPBLAS_TRMM
+
+
 #define GKO_BIND_HIPBLAS_SCAL(ValueType, HipblasName)                        \
     inline void scal(hipblasHandle_t handle, int n, const ValueType* alpha,  \
                      ValueType* x, int incx)                                 \
@@ -255,6 +281,12 @@ using namespace hipblas;
 #define BLAS_OP_T HIPBLAS_OP_T
 #define BLAS_OP_C HIPBLAS_OP_C
 
+#define BLAS_SIDE_LEFT HIPBLAS_SIDE_LEFT
+#define BLAS_SIDE_RIGHT HIPBLAS_SIDE_RIGHT
+
+#define BLAS_DIAG_UNIT HIPBLAS_DIAG_UNIT
+#define BLAS_DIAG_NONUNIT HIPBLAS_DIAG_NON_UNIT
+
 
 }  // namespace blas
 }  // namespace hip