Actual source code: petscdevice_hip.h

  1: #ifndef PETSCDEVICE_HIP_H
  2: #define PETSCDEVICE_HIP_H

  4: #include <petscdevice.h>
  5: #include <petscpkg_version.h>

  7: #if defined(__HCC__) || (defined(__clang__) && defined(__HIP__))
  8:   #define PETSC_USING_HCC 1
  9: #endif

 11: #if PetscDefined(HAVE_HIP)
 12:   #include <hip/hip_runtime.h>

 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

 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

 33:   // REMOVE ME
 34:   #define WaitForHIP() hipDeviceSynchronize()

 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 */

 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__)

 52:   #define PetscHIPCheckLaunch \
 53:     do { \
 54:       /* Check synchronous errors, i.e. pre-launch */ \
 55:       hipGetLastError(); \
 56:       /* Check asynchronous errors, i.e. kernel failed (ULF) */ \
 57:       hipDeviceSynchronize(); \
 58:     } while (0)

 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__)

 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__); \
 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__)

 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__)

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;

109: /* Alias hipsolverDestroy to cusolverDnDestroy */
110: static inline hipsolverStatus_t hipsolverDestroy(hipsolverHandle_t *hipsolverhandle)
111: {
112:   return cusolverDnDestroy(hipsolverhandle);
113: }

115: /* Alias hipsolverCreate to cusolverDnCreate */
116: static inline hipsolverStatus_t hipsolverCreate(hipsolverHandle_t *hipsolverhandle)
117: {
118:   return cusolverDnCreate(hipsolverhandle);
119: }

121: /* Alias hipsolverGetStream to cusolverDnGetStream */
122: static inline hipsolverStatus_t hipsolverGetStream(hipsolverHandle_t handle, hipStream_t *stream)
123: {
124:   return cusolverDnGetStream(handle, stream);
125: }

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;

138: /* Alias hipsolverDestroy to rocblas_destroy_handle */
139: static inline hipsolverStatus_t hipsolverDestroy(hipsolverHandle_t hipsolverhandle)
140: {
141:   return rocblas_destroy_handle(hipsolverhandle);
142: }

144: /* Alias hipsolverCreate to rocblas_destroy_handle */
145: static inline hipsolverStatus_t hipsolverCreate(hipsolverHandle_t *hipsolverhandle)
146: {
147:   return rocblas_create_handle(hipsolverhandle);
148: }

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: }

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 *);

168: #endif // PETSC_HAVE_HIP

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

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

194: #endif // PETSCDEVICE_HIP_H