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