1 #ifndef PETSCDEVICE_HIP_H 2 #define PETSCDEVICE_HIP_H 3 4 #include <petscdevice.h> 5 #include <petscpkg_version.h> 6 7 #if defined(__HCC__) || (defined(__clang__) && defined(__HIP__)) 8 #define PETSC_USING_HCC 1 9 #endif 10 11 #if PetscDefined(HAVE_HIP) 12 #include <hip/hip_runtime.h> 13 14 #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0) 15 #include <hipblas/hipblas.h> 16 #include <hipsparse/hipsparse.h> 17 #else 18 #include <hipblas.h> 19 #include <hipsparse.h> 20 #endif 21 22 #if defined(__HIP_PLATFORM_NVCC__) 23 #include <cusolverDn.h> 24 #else // __HIP_PLATFORM_HCC__ 25 #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0) 26 #include <hipsolver/hipsolver.h> 27 #else 28 #include <hipsolver.h> 29 #endif 30 #endif // __HIP_PLATFORM_NVCC__ 31 #include <hip/hip_complex.h> // for hipComplex, hipDoubleComplex 32 33 // REMOVE ME 34 #define WaitForHIP() hipDeviceSynchronize() 35 36 /* hipBLAS, hipSPARSE and hipSolver does not have hip*GetErrorName(). We create one on our own. */ 37 PETSC_EXTERN const char *PetscHIPBLASGetErrorName(hipblasStatus_t); /* PETSC_EXTERN since it is exposed by the CHKERRHIPBLAS macro */ 38 PETSC_EXTERN const char *PetscHIPSPARSEGetErrorName(hipsparseStatus_t); /* PETSC_EXTERN since it is exposed by the CHKERRHIPSPARSE macro */ 39 PETSC_EXTERN const char *PetscHIPSolverGetErrorName(hipsolverStatus_t); /* PETSC_EXTERN since it is exposed by the CHKERRHIPSOLVER macro */ 40 41 #define PetscCallHIP(...) \ 42 do { \ 43 const hipError_t _p_hip_err__ = __VA_ARGS__; \ 44 if (PetscUnlikely(_p_hip_err__ != hipSuccess)) { \ 45 const char *name = hipGetErrorName(_p_hip_err__); \ 46 const char *descr = hipGetErrorString(_p_hip_err__); \ 47 SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "hip error %d (%s) : %s", (PetscErrorCode)_p_hip_err__, name, descr); \ 48 } \ 49 } while (0) 50 #define CHKERRHIP(...) PetscCallHIP(__VA_ARGS__) 51 52 #define PetscHIPCheckLaunch \ 53 do { \ 54 /* Check synchronous errors, i.e. pre-launch */ \ 55 PetscCallHIP(hipGetLastError()); \ 56 /* Check asynchronous errors, i.e. kernel failed (ULF) */ \ 57 PetscCallHIP(hipDeviceSynchronize()); \ 58 } while (0) 59 60 #define PetscCallHIPBLAS(...) \ 61 do { \ 62 const hipblasStatus_t _p_hipblas_stat__ = __VA_ARGS__; \ 63 if (PetscUnlikely(_p_hipblas_stat__ != HIPBLAS_STATUS_SUCCESS)) { \ 64 const char *name = PetscHIPBLASGetErrorName(_p_hipblas_stat__); \ 65 SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "hipBLAS error %d (%s)", (PetscErrorCode)_p_hipblas_stat__, name); \ 66 } \ 67 } while (0) 68 #define CHKERRHIPBLAS(...) PetscCallHIPBLAS(__VA_ARGS__) 69 70 #if PETSC_PKG_HIP_VERSION_GE(4, 5, 0) 71 /* HIPSPARSE & HIPSOLVER have better functionality with ROCm-4.5 or newer */ 72 #define PetscCallHIPSPARSE(...) \ 73 do { \ 74 const hipsparseStatus_t _p_hipsparse_stat__ = __VA_ARGS__; \ 75 if (PetscUnlikely(_p_hipsparse_stat__ != HIPSPARSE_STATUS_SUCCESS)) { \ 76 const char *name = PetscHIPSPARSEGetErrorName(_p_hipsparse_stat__); \ 77 PetscCheck((_p_hipsparse_stat__ != HIPSPARSE_STATUS_NOT_INITIALIZED) && (_p_hipsparse_stat__ != HIPSPARSE_STATUS_ALLOC_FAILED), PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, "hipSPARSE errorcode %d (%s): Reports not initialized or alloc failed; this indicates the GPU has run out resources", (int)_p_hipsparse_stat__, name); \ 78 SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "hipSPARSE errorcode %d (%s)", (int)_p_hipsparse_stat__, name); \ 79 } \ 80 } while (0) 81 #define CHKERRHIPSPARSE(...) PetscCallHIPSPARSE(__VA_ARGS__) 82 83 #define PetscCallHIPSOLVER(...) \ 84 do { \ 85 const hipsolverStatus_t _p_hipsolver_stat__ = __VA_ARGS__; \ 86 if (PetscUnlikely(_p_hipsolver_stat__ != HIPSOLVER_STATUS_SUCCESS)) { \ 87 const char *name = PetscHIPSolverGetErrorName(_p_hipsolver_stat__); \ 88 if (((_p_hipsolver_stat__ == HIPSOLVER_STATUS_NOT_INITIALIZED) || (_p_hipsolver_stat__ == HIPSOLVER_STATUS_ALLOC_FAILED) || (_p_hipsolver_stat__ == HIPSOLVER_STATUS_INTERNAL_ERROR)) && PetscDeviceInitialized(PETSC_DEVICE_HIP)) { \ 89 SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \ 90 "hipSolver error %d (%s). " \ 91 "This indicates the GPU may have run out resources", \ 92 (PetscErrorCode)_p_hipsolver_stat__, name); \ 93 } else { \ 94 SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "hipSolver error %d (%s)", (PetscErrorCode)_p_hipsolver_stat__, name); \ 95 } \ 96 } \ 97 } while (0) 98 #define CHKERRHIPSOLVER(...) PetscCallHIPSOLVER(__VA_ARGS__) 99 100 #else /* PETSC_PKG_HIP_VERSION_GE(4,5,0) */ 101 /* hipSolver does not exist yet so we work around it 102 rocSOLVER users rocBLAS for the handle 103 * */ 104 #if defined(__HIP_PLATFORM_NVCC__) 105 #include <cusolverDn.h> 106 typedef cusolverDnHandle_t hipsolverHandle_t; 107 typedef cusolverStatus_t hipsolverStatus_t; 108 109 /* Alias hipsolverDestroy to cusolverDnDestroy */ 110 static inline hipsolverStatus_t hipsolverDestroy(hipsolverHandle_t *hipsolverhandle) 111 { 112 return cusolverDnDestroy(hipsolverhandle); 113 } 114 115 /* Alias hipsolverCreate to cusolverDnCreate */ 116 static inline hipsolverStatus_t hipsolverCreate(hipsolverHandle_t *hipsolverhandle) 117 { 118 return cusolverDnCreate(hipsolverhandle); 119 } 120 121 /* Alias hipsolverGetStream to cusolverDnGetStream */ 122 static inline hipsolverStatus_t hipsolverGetStream(hipsolverHandle_t handle, hipStream_t *stream) 123 { 124 return cusolverDnGetStream(handle, stream); 125 } 126 127 /* Alias hipsolverSetStream to cusolverDnSetStream */ 128 static inline hipsolverStatus_t hipsolverSetStream(hipsolverHandle_t handle, hipStream_t stream) 129 { 130 return cusolveDnSetStream(handle, stream); 131 } 132 #else /* __HIP_PLATFORM_HCC__ */ 133 #include <rocsolver.h> 134 #include <rocblas.h> 135 typedef rocblas_handle hipsolverHandle_t; 136 typedef rocblas_status hipsolverStatus_t; 137 138 /* Alias hipsolverDestroy to rocblas_destroy_handle */ 139 static inline hipsolverStatus_t hipsolverDestroy(hipsolverHandle_t hipsolverhandle) 140 { 141 return rocblas_destroy_handle(hipsolverhandle); 142 } 143 144 /* Alias hipsolverCreate to rocblas_destroy_handle */ 145 static inline hipsolverStatus_t hipsolverCreate(hipsolverHandle_t *hipsolverhandle) 146 { 147 return rocblas_create_handle(hipsolverhandle); 148 } 149 150 // Alias hipsolverGetStream to rocblas_get_stream 151 static inline hipsolverStatus_t hipsolverGetStream(hipsolverHandle_t handle, hipStream_t *stream) 152 { 153 return rocblas_get_stream(handle, stream); 154 } 155 156 // Alias hipsolverSetStream to rocblas_set_stream 157 static inline hipsolverStatus_t hipsolverSetStream(hipsolverHandle_t handle, hipStream_t stream) 158 { 159 return rocblas_set_stream(handle, stream); 160 } 161 #endif // __HIP_PLATFORM_NVCC__ 162 #endif /* PETSC_PKG_HIP_VERSION_GE(4,5,0) */ 163 // REMOVE ME 164 PETSC_EXTERN hipStream_t PetscDefaultHipStream; // The default stream used by PETSc 165 PETSC_EXTERN PetscErrorCode PetscHIPBLASGetHandle(hipblasHandle_t *); 166 PETSC_EXTERN PetscErrorCode PetscHIPSOLVERGetHandle(hipsolverHandle_t *); 167 168 #endif // PETSC_HAVE_HIP 169 170 // these can also be defined in petscdevice_cuda.h 171 #ifndef PETSC_DEVICE_DEFINED_DECLS_PRIVATE 172 #define PETSC_DEVICE_DEFINED_DECLS_PRIVATE 173 #if PetscDefined(USING_HCC) 174 #define PETSC_HOST_DECL __host__ 175 #define PETSC_DEVICE_DECL __device__ 176 #define PETSC_KERNEL_DECL __global__ 177 #define PETSC_SHAREDMEM_DECL __shared__ 178 #define PETSC_FORCEINLINE __forceinline__ 179 #define PETSC_CONSTMEM_DECL __constant__ 180 #else 181 #define PETSC_HOST_DECL 182 #define PETSC_DEVICE_DECL 183 #define PETSC_KERNEL_DECL 184 #define PETSC_SHAREDMEM_DECL 185 #define PETSC_FORCEINLINE inline 186 #define PETSC_CONSTMEM_DECL 187 #endif // PETSC_USING_NVCC 188 189 #define PETSC_HOSTDEVICE_DECL PETSC_HOST_DECL PETSC_DEVICE_DECL 190 #define PETSC_DEVICE_INLINE_DECL PETSC_DEVICE_DECL PETSC_FORCEINLINE 191 #define PETSC_HOSTDEVICE_INLINE_DECL PETSC_HOSTDEVICE_DECL PETSC_FORCEINLINE 192 #endif // PETSC_DEVICE_DEFINED_DECLS_PRIVATE 193 194 #endif // PETSCDEVICE_HIP_H 195