Actual source code: petscdevice_hip.h
1: #pragma once
3: #include <petscdevice.h>
4: #include <petscpkg_version.h>
6: /* MANSEC = Sys */
7: /* SUBMANSEC = Device */
9: #if defined(__HCC__) || (defined(__clang__) && defined(__HIP__))
10: #define PETSC_USING_HCC 1
11: #endif
13: #if PetscDefined(HAVE_HIP)
14: #include <hip/hip_runtime.h>
16: #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
18: // cupmScalarPtrCast() returns hip{Float,Double}Complex while hipBLAS uses hipBlas{Float,Double}Complex, causing many VecCUPM errors like
19: // error: no matching function for call to 'cupmBlasXdot'.
20: // Before rocm-6.0, one can define ROCM_MATHLIBS_API_USE_HIP_COMPLEX to force rocm to 'typedef hipDoubleComplex hipBlasDoubleComplex' for example.
21: // Since then, ROCM_MATHLIBS_API_USE_HIP_COMPLEX is deprecated, and one can define HIPBLAS_V2 to use version 2 of hipBLAS that directly use hipDoubleComplex etc.
22: // Per AMD, HIPBLAS_V2 will be removed in the future so that hipBLAS only provides updated APIs (but not yet in 6.2.2 as of Sep. 27, 2024).
23: //
24: // see https://rocm.docs.amd.com/projects/hipBLAS/en/docs-6.0.0/functions.html#complex-datatypes
25: // and https://rocm.docs.amd.com/projects/hipBLAS/en/docs-6.2.2/functions.html#hipblas-v2-and-deprecations
26: #if PETSC_PKG_HIP_VERSION_GE(6, 0, 0)
27: #define HIPBLAS_V2
28: #else
29: #define ROCM_MATHLIBS_API_USE_HIP_COMPLEX
30: #endif
31: #include <hipblas/hipblas.h>
32: #include <hipsparse/hipsparse.h>
33: #else
34: #include <hipblas.h>
35: #include <hipsparse.h>
36: #endif
38: #if PETSC_PKG_HIP_VERSION_LT(5, 4, 0)
39: #define HIPSPARSE_ORDER_COL HIPSPARSE_ORDER_COLUMN
40: #endif
42: #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
43: #include <hipsolver/hipsolver.h>
44: #else
45: #include <hipsolver.h>
46: #endif
47: #include <hip/hip_complex.h> // for hipComplex, hipDoubleComplex
49: /*MC
50: WaitForHIP - Block the calling host thread until all previously queued work on the current HIP device has completed
52: Synopsis:
53: #include <petscdevice_hip.h>
54: hipError_t WaitForHIP(void)
56: Not Collective; No Fortran Support
58: Level: developer
60: Note:
61: Thin convenience wrapper around `hipDeviceSynchronize()`. Marked for removal in favour of
62: explicit `PetscDeviceContext` synchronization.
64: .seealso: `PetscDeviceContext`, `PetscDeviceContextSynchronize()`, `WaitForCUDA()`
65: M*/
66: // REMOVE ME
67: #define WaitForHIP() hipDeviceSynchronize()
69: /* hipBLAS, hipSPARSE and hipSolver does not have hip*GetErrorName(). We create one on our own. */
70: PETSC_EXTERN const char *PetscHIPBLASGetErrorName(hipblasStatus_t); /* PETSC_EXTERN since it is exposed by the CHKERRHIPBLAS macro */
71: PETSC_EXTERN const char *PetscHIPSPARSEGetErrorName(hipsparseStatus_t); /* PETSC_EXTERN since it is exposed by the CHKERRHIPSPARSE macro */
72: PETSC_EXTERN const char *PetscHIPSolverGetErrorName(hipsolverStatus_t); /* PETSC_EXTERN since it is exposed by the CHKERRHIPSOLVER macro */
74: #define PetscCallHIP(...) \
75: do { \
76: const hipError_t _p_hip_err__ = __VA_ARGS__; \
77: if (PetscUnlikely(_p_hip_err__ != hipSuccess)) { \
78: const char *name = hipGetErrorName(_p_hip_err__); \
79: const char *descr = hipGetErrorString(_p_hip_err__); \
80: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "hip error %d (%s) : %s", (PetscErrorCode)_p_hip_err__, name, descr); \
81: } \
82: } while (0)
83: #define CHKERRHIP(...) PetscCallHIP(__VA_ARGS__)
85: #define PetscHIPCheckLaunch \
86: do { \
87: /* Check synchronous errors, i.e. pre-launch */ \
88: PetscCallHIP(hipGetLastError()); \
89: /* Check asynchronous errors, i.e. kernel failed (ULF) */ \
90: PetscCallHIP(hipDeviceSynchronize()); \
91: } while (0)
93: #define PetscCallHIPBLAS(...) \
94: do { \
95: const hipblasStatus_t _p_hipblas_stat__ = __VA_ARGS__; \
96: if (PetscUnlikely(_p_hipblas_stat__ != HIPBLAS_STATUS_SUCCESS)) { \
97: const char *name = PetscHIPBLASGetErrorName(_p_hipblas_stat__); \
98: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "hipBLAS error %d (%s)", (PetscErrorCode)_p_hipblas_stat__, name); \
99: } \
100: } while (0)
101: #define CHKERRHIPBLAS(...) PetscCallHIPBLAS(__VA_ARGS__)
103: #if PETSC_PKG_HIP_VERSION_GE(4, 5, 0)
104: /* HIPSPARSE & HIPSOLVER have better functionality with ROCm-4.5 or newer */
105: #define PetscCallHIPSPARSE(...) \
106: do { \
107: const hipsparseStatus_t _p_hipsparse_stat__ = __VA_ARGS__; \
108: if (PetscUnlikely(_p_hipsparse_stat__ != HIPSPARSE_STATUS_SUCCESS)) { \
109: const char *name = PetscHIPSPARSEGetErrorName(_p_hipsparse_stat__); \
110: 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); \
111: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "hipSPARSE errorcode %d (%s)", (int)_p_hipsparse_stat__, name); \
112: } \
113: } while (0)
114: #define CHKERRHIPSPARSE(...) PetscCallHIPSPARSE(__VA_ARGS__)
116: #define PetscCallHIPSOLVER(...) \
117: do { \
118: const hipsolverStatus_t _p_hipsolver_stat__ = __VA_ARGS__; \
119: if (PetscUnlikely(_p_hipsolver_stat__ != HIPSOLVER_STATUS_SUCCESS)) { \
120: const char *name = PetscHIPSolverGetErrorName(_p_hipsolver_stat__); \
121: 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)) { \
122: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
123: "hipSolver error %d (%s). " \
124: "This indicates the GPU may have run out resources", \
125: (PetscErrorCode)_p_hipsolver_stat__, name); \
126: } else { \
127: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "hipSolver error %d (%s)", (PetscErrorCode)_p_hipsolver_stat__, name); \
128: } \
129: } \
130: } while (0)
131: #define CHKERRHIPSOLVER(...) PetscCallHIPSOLVER(__VA_ARGS__)
133: #else /* PETSC_PKG_HIP_VERSION_GE(4,5,0) */
134: /* hipSolver does not exist yet so we work around it
135: rocSOLVER users rocBLAS for the handle
136: * */
137: #if defined(__HIP_PLATFORM_NVCC__)
138: #include <cusolverDn.h>
139: /*MC
140: hipsolverHandle_t - Opaque handle to a hipSolver context used by PETSc to call into the hipSolver dense linear algebra library
142: Level: developer
144: Note:
145: On modern ROCm releases (>= 4.5) this is the native hipSolver handle. On older releases PETSc provides a compatibility shim that aliases the handle to `cusolverDnHandle_t` on the NVIDIA HIP platform or to `rocblas_handle` on the AMD HIP platform.
147: .seealso: `PetscHIPSOLVERGetHandle()`, `hipsolverStatus_t`, `PetscCallHIPSOLVER()`
148: M*/
149: typedef cusolverDnHandle_t hipsolverHandle_t;
151: /*MC
152: hipsolverStatus_t - Return-status enumeration used by hipSolver to report success or failure of a hipSolver call
154: Level: developer
156: Note:
157: On modern ROCm releases (>= 4.5) this is the native hipSolver status type. On older releases PETSc provides a compatibility shim that aliases the status to `cusolverStatus_t` on the NVIDIA HIP platform or to `rocblas_status` on the AMD HIP platform. Use `PetscCallHIPSOLVER()` to check the return value.
159: .seealso: `hipsolverHandle_t`, `PetscCallHIPSOLVER()`, `PetscHIPSolverGetErrorName()`
160: M*/
161: typedef cusolverStatus_t hipsolverStatus_t;
163: /* Alias hipsolverDestroy to cusolverDnDestroy */
164: static inline hipsolverStatus_t hipsolverDestroy(hipsolverHandle_t hipsolverhandle)
165: {
166: return cusolverDnDestroy(hipsolverhandle);
167: }
169: /* Alias hipsolverCreate to cusolverDnCreate */
170: static inline hipsolverStatus_t hipsolverCreate(hipsolverHandle_t *hipsolverhandle)
171: {
172: return cusolverDnCreate(hipsolverhandle);
173: }
175: /* Alias hipsolverGetStream to cusolverDnGetStream */
176: static inline hipsolverStatus_t hipsolverGetStream(hipsolverHandle_t handle, hipStream_t *stream)
177: {
178: return cusolverDnGetStream(handle, stream);
179: }
181: /* Alias hipsolverSetStream to cusolverDnSetStream */
182: static inline hipsolverStatus_t hipsolverSetStream(hipsolverHandle_t handle, hipStream_t stream)
183: {
184: return cusolverDnSetStream(handle, stream);
185: }
186: #else /* __HIP_PLATFORM_HCC__ */
187: #include <rocsolver.h>
188: #include <rocblas.h>
189: typedef rocblas_handle hipsolverHandle_t;
190: typedef rocblas_status hipsolverStatus_t;
192: /* Alias hipsolverDestroy to rocblas_destroy_handle */
193: static inline hipsolverStatus_t hipsolverDestroy(hipsolverHandle_t hipsolverhandle)
194: {
195: return rocblas_destroy_handle(hipsolverhandle);
196: }
198: /* Alias hipsolverCreate to rocblas_destroy_handle */
199: static inline hipsolverStatus_t hipsolverCreate(hipsolverHandle_t *hipsolverhandle)
200: {
201: return rocblas_create_handle(hipsolverhandle);
202: }
204: // Alias hipsolverGetStream to rocblas_get_stream
205: static inline hipsolverStatus_t hipsolverGetStream(hipsolverHandle_t handle, hipStream_t *stream)
206: {
207: return rocblas_get_stream(handle, stream);
208: }
210: // Alias hipsolverSetStream to rocblas_set_stream
211: static inline hipsolverStatus_t hipsolverSetStream(hipsolverHandle_t handle, hipStream_t stream)
212: {
213: return rocblas_set_stream(handle, stream);
214: }
215: #endif // __HIP_PLATFORM_NVCC__
216: #endif /* PETSC_PKG_HIP_VERSION_GE(4,5,0) */
217: // REMOVE ME
218: PETSC_EXTERN hipStream_t PetscDefaultHipStream; // The default stream used by PETSc
219: PETSC_EXTERN PetscErrorCode PetscHIPBLASGetHandle(hipblasHandle_t *);
220: PETSC_EXTERN PetscErrorCode PetscHIPSOLVERGetHandle(hipsolverHandle_t *);
221: PETSC_EXTERN PetscErrorCode PetscGetCurrentHIPStream(hipStream_t *);
223: #endif // PETSC_HAVE_HIP
225: // these can also be defined in petscdevice_cuda.h so we undef and define them *only* if the
226: // current compiler is HCC. In this case if petscdevice_cuda.h is included first, the macros
227: // would already be defined, but they would be empty since we cannot be using NVCC at the same
228: // time.
229: #if PetscDefined(USING_HCC)
230: #undef PETSC_HOST_DECL
231: #undef PETSC_DEVICE_DECL
232: #undef PETSC_KERNEL_DECL
233: #undef PETSC_SHAREDMEM_DECL
234: #undef PETSC_FORCEINLINE
235: #undef PETSC_CONSTMEM_DECL
237: #define PETSC_HOST_DECL __host__
238: #define PETSC_DEVICE_DECL __device__
239: #define PETSC_KERNEL_DECL __global__
240: #define PETSC_SHAREDMEM_DECL __shared__
241: #define PETSC_FORCEINLINE __forceinline__
242: #define PETSC_CONSTMEM_DECL __constant__
243: #endif
245: #if !defined(PETSC_HOST_DECL) // use HOST_DECL as canary
246: #define PETSC_HOST_DECL
247: #define PETSC_DEVICE_DECL
248: #define PETSC_KERNEL_DECL
249: #define PETSC_SHAREDMEM_DECL
250: #define PETSC_FORCEINLINE inline
251: #define PETSC_CONSTMEM_DECL
252: #endif
254: #if !PetscDefined(DEVICE_DEFINED_DECLS_PRIVATE)
255: #define PETSC_DEVICE_DEFINED_DECLS_PRIVATE
256: #define PETSC_HOSTDEVICE_DECL PETSC_HOST_DECL PETSC_DEVICE_DECL
257: #define PETSC_DEVICE_INLINE_DECL PETSC_DEVICE_DECL PETSC_FORCEINLINE
258: #define PETSC_HOSTDEVICE_INLINE_DECL PETSC_HOSTDEVICE_DECL PETSC_FORCEINLINE
259: #endif