PaddlePaddle
diff --git a/‎cmake/hip.cmake
Copy file name to clipboard
+1Lines changed: 1 addition & 0 deletions b/‎cmake/hip.cmake
Copy file name to clipboard
+1Lines changed: 1 addition & 0 deletions
diff --git a/‎paddle/phi/backends/dynload/CMakeLists.txt
Copy file name to clipboard
+2-1Lines changed: 2 additions & 1 deletion b/‎paddle/phi/backends/dynload/CMakeLists.txt
Copy file name to clipboard
+2-1Lines changed: 2 additions & 1 deletion
diff --git a/‎paddle/phi/backends/dynload/dynamic_loader.cc
Copy file name to clipboard
+2Lines changed: 2 additions & 0 deletions b/‎paddle/phi/backends/dynload/dynamic_loader.cc
Copy file name to clipboard
+2Lines changed: 2 additions & 0 deletions
diff --git a/‎paddle/phi/backends/dynload/rocblas.h
Copy file name to clipboard
+7Lines changed: 7 additions & 0 deletions b/‎paddle/phi/backends/dynload/rocblas.h
Copy file name to clipboard
+7Lines changed: 7 additions & 0 deletions
diff --git a/‎paddle/phi/backends/dynload/rocsolver.cc
Copy file name to clipboard
+26Lines changed: 26 additions & 0 deletions b/‎paddle/phi/backends/dynload/rocsolver.cc
Copy file name to clipboard
+26Lines changed: 26 additions & 0 deletions
diff --git a/‎paddle/phi/backends/dynload/rocsolver.h
Copy file name to clipboard
+79Lines changed: 79 additions & 0 deletions b/‎paddle/phi/backends/dynload/rocsolver.h
Copy file name to clipboard
+79Lines changed: 79 additions & 0 deletions
diff --git a/‎paddle/phi/backends/gpu/gpu_resources.cc
Copy file name to clipboard
+10-2Lines changed: 10 additions & 2 deletions b/‎paddle/phi/backends/gpu/gpu_resources.cc
Copy file name to clipboard
+10-2Lines changed: 10 additions & 2 deletions
diff --git a/‎paddle/phi/kernels/CMakeLists.txt
Copy file name to clipboard
-4Lines changed: 0 additions & 4 deletions b/‎paddle/phi/kernels/CMakeLists.txt
Copy file name to clipboard
-4Lines changed: 0 additions & 4 deletions
diff --git a/‎paddle/phi/kernels/funcs/blas/blas_impl.hip.h
Copy file name to clipboard
+8Lines changed: 8 additions & 0 deletions b/‎paddle/phi/kernels/funcs/blas/blas_impl.hip.h
Copy file name to clipboard
+8Lines changed: 8 additions & 0 deletions
diff --git a/‎paddle/phi/kernels/funcs/values_vectors_functor.h
Copy file name to clipboard
+138-2Lines changed: 138 additions & 2 deletions b/‎paddle/phi/kernels/funcs/values_vectors_functor.h
Copy file name to clipboard
+138-2Lines changed: 138 additions & 2 deletions
@@ -85,6 +85,7 @@ find_package_and_include(rocprim)
 find_package_and_include(hipsparse)
 find_package_and_include(rocsparse)
 find_package_and_include(rocfft)
+find_package_and_include(rocsolver)
 
 if(CCACHE_PATH)
   set(HIP_HIPCC_EXECUTABLE ${CCACHE_PATH} ${HIP_HIPCC_EXECUTABLE})
 
@@ -29,7 +29,8 @@ if(WITH_ROCM)
     miopen.cc
     hiprand.cc
     hipfft.cc
-    rocsparse.cc)
+    rocsparse.cc
+    rocsolver.cc)
 endif()
 
 # There is no macOS version of NCCL.
 
@@ -593,6 +593,8 @@ void* GetCusolverDsoHandle() {
   return GetDsoHandleFromSearchPath(
       FLAGS_cuda_dir, win_cusolver_lib, true, {cuda_lib_path});
 #endif
+#elif defined(PADDLE_WITH_HIP)
+  return GetDsoHandleFromSearchPath(FLAGS_rocm_dir, "librocsolver.so");
 #else
 #ifdef PADDLE_WITH_PIP_CUDA_LIBRARIES
   return GetDsoHandleFromSearchPath(FLAGS_cuda_dir, "libcusolver.so.11");
 
@@ -108,6 +108,13 @@ ROCBLAS_BLAS_ROUTINE_EACH_R2(DECLARE_DYNAMIC_LOAD_ROCBLAS_WRAP)
 
 ROCBLAS_BLAS_ROUTINE_EACH_R4(DECLARE_DYNAMIC_LOAD_ROCBLAS_WRAP)
 
+#if HIP_VERSION >= 30000000
+#define ROCBLAS_BLAS_ROUTINE_EACH_R5(__macro) \
+  __macro(rocblas_strsm_batched);             \
+  __macro(rocblas_dtrsm_batched);
+
+ROCBLAS_BLAS_ROUTINE_EACH_R5(DECLARE_DYNAMIC_LOAD_ROCBLAS_WRAP)
+#endif
 #undef DECLARE_DYNAMIC_LOAD_ROCBLAS_WRAP
 }  // namespace dynload
 }  // namespace phi
@@ -0,0 +1,26 @@
+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    https://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+
+#include "paddle/phi/backends/dynload/rocsolver.h"
+
+namespace phi::dynload {
+
+std::once_flag rocsolver_dso_flag;
+void *rocsolver_dso_handle;
+
+#define DEFINE_WRAP(__name) DynLoad__##__name __name
+
+ROCSOLVER_ROUTINE_EACH(DEFINE_WRAP);
+
+}  // namespace phi::dynload
@@ -0,0 +1,79 @@
+/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    https://www.apache.org/licenses/LICENSE-2.0
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License. */
+#pragma once
+
+#include <hip/hip_runtime.h>
+#include <rocsolver/rocsolver.h>
+
+#include <mutex>  // NOLINT
+
+#include "paddle/phi/backends/dynload/dynamic_loader.h"
+#include "paddle/phi/common/port.h"
+
+namespace phi {
+namespace dynload {
+extern std::once_flag rocsolver_dso_flag;
+extern void *rocsolver_dso_handle;
+
+#define DECLARE_DYNAMIC_LOAD_ROCSOLVER_WRAP(__name)                   \
+  struct DynLoad__##__name {                                          \
+    template <typename... Args>                                       \
+    rocblas_status operator()(Args... args) {                         \
+      using rocsolverFunc = decltype(&::__name);                      \
+      std::call_once(rocsolver_dso_flag, []() {                       \
+        rocsolver_dso_handle = phi::dynload::GetCusolverDsoHandle();  \
+      });                                                             \
+      static void *p_##__name = dlsym(rocsolver_dso_handle, #__name); \
+      return reinterpret_cast<rocsolverFunc>(p_##__name)(args...);    \
+    }                                                                 \
+  };                                                                  \
+  extern DynLoad__##__name __name
+
+#define ROCSOLVER_ROUTINE_EACH(__macro) \
+  __macro(rocsolver_spotrs);            \
+  __macro(rocsolver_dpotrs);            \
+  __macro(rocsolver_cpotrs);            \
+  __macro(rocsolver_zpotrs);            \
+  __macro(rocsolver_sgetrf);            \
+  __macro(rocsolver_dgetrf);            \
+  __macro(rocsolver_cgetrf);            \
+  __macro(rocsolver_zgetrf);            \
+  __macro(rocsolver_spotrf);            \
+  __macro(rocsolver_dpotrf);            \
+  __macro(rocsolver_cpotrf);            \
+  __macro(rocsolver_zpotrf);            \
+  __macro(rocsolver_spotrf_batched);    \
+  __macro(rocsolver_dpotrf_batched);    \
+  __macro(rocsolver_cpotrf_batched);    \
+  __macro(rocsolver_zpotrf_batched);    \
+  __macro(rocsolver_sgeqrf);            \
+  __macro(rocsolver_dgeqrf);            \
+  __macro(rocsolver_sorgqr);            \
+  __macro(rocsolver_dorgqr);            \
+  __macro(rocsolver_dormqr);            \
+  __macro(rocsolver_sormqr);
+
+ROCSOLVER_ROUTINE_EACH(DECLARE_DYNAMIC_LOAD_ROCSOLVER_WRAP);
+
+#if HIP_VERSION >= 50300000
+#define ROCSOLVER_ROUTINE_EACH1(__macro) \
+  __macro(rocsolver_ssyevj_batched);     \
+  __macro(rocsolver_dsyevj_batched);
+
+ROCSOLVER_ROUTINE_EACH1(DECLARE_DYNAMIC_LOAD_ROCSOLVER_WRAP);
+#endif
+
+#undef DECLARE_DYNAMIC_LOAD_ROCSOLVER_WRAP
+}  // namespace dynload
+}  // namespace phi
@@ -335,18 +335,26 @@ void DestroyDnnHandle(dnnHandle_t handle) {
 }
 
 void InitSolverHandle(solverHandle_t* handle, gpuStream_t stream) {
-#ifndef PADDLE_WITH_HIP
+#if defined(PADDLE_WITH_CUDA)
   PADDLE_RETRY_CUDA_SUCCESS(phi::dynload::cusolverDnCreate(handle));
   PADDLE_RETRY_CUDA_SUCCESS(phi::dynload::cusolverDnSetStream(*handle, stream));
+#elif defined(PADDLE_WITH_HIP)
+  phi::dynload::rocblas_create_handle(handle);
+  phi::dynload::rocblas_set_stream(*handle, stream);
 #endif
 }
 
 void DestroySolverHandle(solverHandle_t solver_handle) {
-#ifndef PADDLE_WITH_HIP
+#if defined(PADDLE_WITH_CUDA)
   if (solver_handle != nullptr) {
     PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::cusolverDnDestroy(solver_handle));
     solver_handle = nullptr;
   }
+#elif defined(PADDLE_WITH_HIP)
+  if (solver_handle != nullptr) {
+    phi::dynload::rocblas_destroy_handle(solver_handle);
+    solver_handle = nullptr;
+  }
 #endif
 }
 
 
@@ -232,16 +232,12 @@ if(WITH_ROCM)
     "gpu/affine_grid_grad_kernel.cu"
     "gpu/apply_per_channel_scale_kernel.cu"
     "gpu/calc_reduced_attn_kernel.cu"
-    "gpu/cholesky_solve_kernel.cu"
-    "gpu/eigh_kernel.cu"
     "gpu/eigvalsh_kernel.cu"
-    "gpu/lstsq_kernel.cu"
     "gpu/lu_kernel.cu"
     "gpu/matrix_rank_kernel.cu"
     "gpu/matrix_rank_tol_kernel.cu"
     "gpu/put_along_axis_grad_kernel.cu"
     "gpu/put_along_axis_kernel.cu"
-    "gpu/qr_kernel.cu"
     "gpu/svd_kernel.cu"
     "gpu/cuda_gemm_kernel.cu")
 endif()
 
@@ -95,8 +95,12 @@ struct CUBlas<float> {
 
   template <typename... ARGS>
   static void TRSM_BATCH(ARGS... args) {
+#if HIP_VERSION >= 30000000
+    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_strsm_batched(args...));
+#else
     PADDLE_THROW(common::errors::Unimplemented(
         "cublasStrsmBatched is not supported on HIP platform."));
+#endif
   }
 };
 
@@ -164,8 +168,12 @@ struct CUBlas<double> {
 
   template <typename... ARGS>
   static void TRSM_BATCH(ARGS... args) {
+#if HIP_VERSION >= 30000000
+    PADDLE_ENFORCE_GPU_SUCCESS(phi::dynload::rocblas_dtrsm_batched(args...));
+#else
     PADDLE_THROW(common::errors::Unimplemented(
         "cublasDtrsmBatched is not supported on HIP platform."));
+#endif
   }
 };
 
 
@@ -14,9 +14,15 @@
 
 #pragma once
 #ifdef PADDLE_WITH_CUDA
-#include "paddle/common/errors.h"
 #include "paddle/phi/backends/dynload/cusolver.h"
 #endif  // PADDLE_WITH_CUDA
+#ifdef PADDLE_WITH_HIP
+#include <thrust/device_vector.h>
+#include "paddle/phi/backends/dynload/rocsolver.h"
+#endif  // PADDLE_WITH_HIP
+#if defined(PADDLE_WITH_CUDA) || defined(PADDLE_WITH_HIP)
+#include "paddle/common/errors.h"
+#endif
 #include "paddle/phi/backends/cpu/cpu_context.h"
 #include "paddle/phi/backends/gpu/gpu_context.h"
 #include "paddle/phi/common/memory_utils.h"
@@ -185,7 +191,7 @@ inline void syevjBatched<phi::dtype::complex<double>, double>(
 }
 #endif
 
-#ifdef PADDLE_WITH_CUDA
+#if defined(PADDLE_WITH_CUDA) || defined(PADDLE_WITH_HIP)
 static void CheckEighResult(const GPUContext &dev_ctx,
                             const int64_t batch_size,
                             int *info) {
@@ -323,6 +329,136 @@ struct MatrixEighFunctor<CPUContext, T> {
   }
 };
 
+#ifdef PADDLE_WITH_HIP
+#define ROCSOLVER_SYEVJ_BATCHED_ARGTYPES(scalar_t, value_t)            \
+  solverHandle_t handle, rocblas_esort esort, rocblas_evect evect,     \
+      rocblas_fill uplo, int n, scalar_t *const A[], int lda,          \
+      const scalar_t abstol, scalar_t *residual, const int max_sweeps, \
+      int *n_sweeps, value_t *W, const int strideW, int *info,         \
+      const int batch_count
+
+template <class scalar_t, class value_t = scalar_t>
+void syevjBatched(ROCSOLVER_SYEVJ_BATCHED_ARGTYPES(scalar_t, value_t)) {
+  PADDLE_THROW(common::errors::InvalidArgument(
+      "syevjBatched: not implemented for %s", typeid(scalar_t).name()));
+}
+
+template <>
+inline void syevjBatched<float>(ROCSOLVER_SYEVJ_BATCHED_ARGTYPES(float,
+                                                                 float)) {
+  PADDLE_ENFORCE_GPU_SUCCESS(dynload::rocsolver_ssyevj_batched(handle,
+                                                               esort,
+                                                               evect,
+                                                               uplo,
+                                                               n,
+                                                               A,
+                                                               lda,
+                                                               abstol,
+                                                               residual,
+                                                               max_sweeps,
+                                                               n_sweeps,
+                                                               W,
+                                                               strideW,
+                                                               info,
+                                                               batch_count));
+}
+
+template <>
+inline void syevjBatched<double>(ROCSOLVER_SYEVJ_BATCHED_ARGTYPES(double,
+                                                                  double)) {
+  PADDLE_ENFORCE_GPU_SUCCESS(dynload::rocsolver_dsyevj_batched(handle,
+                                                               esort,
+                                                               evect,
+                                                               uplo,
+                                                               n,
+                                                               A,
+                                                               lda,
+                                                               abstol,
+                                                               residual,
+                                                               max_sweeps,
+                                                               n_sweeps,
+                                                               W,
+                                                               strideW,
+                                                               info,
+                                                               batch_count));
+}
+
+template <typename T>
+struct MatrixEighFunctor<GPUContext, T> {
+ public:
+  void operator()(const GPUContext &dev_ctx,
+                  const DenseTensor &input,
+                  DenseTensor *eigen_values,
+                  DenseTensor *eigen_vectors,
+                  bool is_lower,
+                  bool has_vectors) {
+    using ValueType = phi::dtype::Real<T>;
+
+    auto &dims = input.dims();
+    int dim_size = dims.size();
+    int64_t batch_size = GetBatchSize(dims);
+    int last_dim = dims[dim_size - 1];
+    int lda = std::max<int>(1, last_dim);
+    auto vector_stride = dims[dim_size - 1] * dims[dim_size - 2];
+    auto values_stride = dims[dim_size - 1];
+
+    rocblas_fill uplo = is_lower ? rocblas_fill_lower : rocblas_fill_upper;
+    rocblas_evect evect =
+        has_vectors ? rocblas_evect_original : rocblas_evect_none;
+
+    ValueType *out_value = dev_ctx.template Alloc<ValueType>(eigen_values);
+    DenseTensor input_trans = phi::TransposeLast2Dim<T>(dev_ctx, input);
+    T *input_vector = input_trans.data<T>();
+
+    auto handle = dev_ctx.cusolver_dn_handle();
+
+    size_t total_bytes = sizeof(T) * batch_size + sizeof(int) * batch_size * 2;
+    auto info = phi::memory_utils::Alloc(
+        dev_ctx.GetPlace(),
+        total_bytes,
+        phi::Stream(reinterpret_cast<phi::StreamId>(dev_ctx.stream())));
+    auto *residual_ptr = reinterpret_cast<T *>(info->ptr());
+    auto *info_ptr = reinterpret_cast<int *>(residual_ptr + batch_size);
+    auto *n_sweeps_ptr = reinterpret_cast<int *>(info_ptr + batch_size);
+
+    std::vector<T *> output_ptrs;
+    for (int i = 0; i < batch_size; i++) {
+      output_ptrs.emplace_back(input_vector + i * vector_stride);
+    }
+    thrust::device_vector<T *> dev_output_ptrs(output_ptrs.begin(),
+                                               output_ptrs.end());
+
+    syevjBatched<T>(handle,
+                    rocblas_esort_ascending,
+                    evect,
+                    uplo,
+                    last_dim,
+                    thrust::raw_pointer_cast(dev_output_ptrs.data()),
+                    lda,
+                    0,
+                    residual_ptr,
+                    100,  // 100 max_sweeps default
+                    n_sweeps_ptr,
+                    out_value,
+                    values_stride,
+                    info_ptr,
+                    batch_size);
+
+    CheckEighResult(dev_ctx, batch_size, info_ptr);
+
+    if (has_vectors) {
+      PADDLE_ENFORCE_NOT_NULL(eigen_vectors,
+                              common::errors::InvalidArgument(
+                                  "When has_vectors is true,"
+                                  "the eigenvectors needs to be calculated,"
+                                  "so the eigenvectors must be provided."));
+      input_trans = phi::TransposeLast2Dim<T>(dev_ctx, input_trans);
+      eigen_vectors->ShareDataWith(input_trans);
+    }
+  }
+};
+#endif
+
 #ifdef PADDLE_WITH_CUDA
 
 // Calculates the eigenvalues and eigenvectors of Hermitian or real