Actual source code: petscdevice_cuda.h
1: #pragma once
3: #include <petscdevice.h>
4: #include <petscpkg_version.h>
6: /* MANSEC = Sys */
8: #if defined(__NVCC__) || defined(__CUDACC__)
9: #define PETSC_USING_NVCC 1
10: #endif
12: #if PetscDefined(HAVE_CUDA)
13: #include <cuda.h>
14: #include <cuda_runtime.h>
15: #include <cublas_v2.h>
16: #define DISABLE_CUSPARSE_DEPRECATED
17: #include <cusparse.h>
18: #include <cusolverDn.h>
19: #include <cusolverSp.h>
20: #include <cufft.h>
21: #include <curand.h>
22: #if PetscDefined(HAVE_NVML)
23: #include <nvml.h> // NVML comes with the NVIDIA GPU driver
24: #endif
26: /* cuBLAS does not have cublasGetErrorName(). We create one on our own. */
27: PETSC_EXTERN const char *PetscCUBLASGetErrorName(cublasStatus_t); /* PETSC_EXTERN since it is exposed by the CHKERRCUBLAS macro */
28: PETSC_EXTERN const char *PetscCUSolverGetErrorName(cusolverStatus_t);
29: PETSC_EXTERN const char *PetscCUFFTGetErrorName(cufftResult);
31: /*MC
32: WaitForCUDA - Block the calling host thread until all previously queued work on the current CUDA device has completed
34: Synopsis:
35: #include <petscdevice_cuda.h>
36: cudaError_t WaitForCUDA(void)
38: Not Collective; No Fortran Support
40: Level: developer
42: Note:
43: Thin convenience wrapper around `cudaDeviceSynchronize()`. Marked for removal in favour of
44: explicit `PetscDeviceContext` synchronization.
46: .seealso: `PetscDeviceContext`, `PetscDeviceContextSynchronize()`, `WaitForHIP()`
47: M*/
48: /* REMOVE ME */
49: #define WaitForCUDA() cudaDeviceSynchronize()
51: /* CUDART_VERSION = 1000 x major + 10 x minor version */
53: /* Could not find exactly which CUDART_VERSION introduced cudaGetErrorName. At least it was in CUDA 8.0 (Sep. 2016) */
54: #if PETSC_PKG_CUDA_VERSION_GE(8, 0, 0)
55: #define PetscCallCUDAVoid(...) \
56: do { \
57: const cudaError_t _p_cuda_err__ = __VA_ARGS__; \
58: PetscCheckAbort(_p_cuda_err__ == cudaSuccess, PETSC_COMM_SELF, PETSC_ERR_GPU, "cuda error %d (%s) : %s", (PetscErrorCode)_p_cuda_err__, cudaGetErrorName(_p_cuda_err__), cudaGetErrorString(_p_cuda_err__)); \
59: } while (0)
61: #define PetscCallCUDA(...) \
62: do { \
63: const cudaError_t _p_cuda_err__ = __VA_ARGS__; \
64: PetscCheck(_p_cuda_err__ == cudaSuccess, PETSC_COMM_SELF, PETSC_ERR_GPU, "cuda error %d (%s) : %s", (PetscErrorCode)_p_cuda_err__, cudaGetErrorName(_p_cuda_err__), cudaGetErrorString(_p_cuda_err__)); \
65: } while (0)
66: #else /* PETSC_PKG_CUDA_VERSION_GE(8,0,0) */
67: #define PetscCallCUDA(...) \
68: do { \
69: const cudaError_t _p_cuda_err__ = __VA_ARGS__; \
70: PetscCheck(_p_cuda_err__ == cudaSuccess, PETSC_COMM_SELF, PETSC_ERR_GPU, "cuda error %d", (PetscErrorCode)_p_cuda_err__); \
71: } while (0)
73: #define PetscCallCUDAVoid(...) \
74: do { \
75: const cudaError_t _p_cuda_err__ = __VA_ARGS__; \
76: PetscCheckAbort(_p_cuda_err__ == cudaSuccess, PETSC_COMM_SELF, PETSC_ERR_GPU, "cuda error %d", (PetscErrorCode)_p_cuda_err__); \
77: } while (0)
78: #endif /* PETSC_PKG_CUDA_VERSION_GE(8,0,0) */
79: #define CHKERRCUDA(...) PetscCallCUDA(__VA_ARGS__)
81: #define PetscCUDACheckLaunch \
82: do { \
83: /* Check synchronous errors, i.e. pre-launch */ \
84: PetscCallCUDA(cudaGetLastError()); \
85: /* Check asynchronous errors, i.e. kernel failed (ULF) */ \
86: PetscCallCUDA(cudaDeviceSynchronize()); \
87: } while (0)
89: #define PetscCallCUBLAS(...) \
90: do { \
91: const cublasStatus_t _p_cublas_stat__ = __VA_ARGS__; \
92: if (PetscUnlikely(_p_cublas_stat__ != CUBLAS_STATUS_SUCCESS)) { \
93: const char *name = PetscCUBLASGetErrorName(_p_cublas_stat__); \
94: if ((_p_cublas_stat__ == CUBLAS_STATUS_NOT_INITIALIZED || _p_cublas_stat__ == CUBLAS_STATUS_ALLOC_FAILED) && PetscDeviceInitialized(PETSC_DEVICE_CUDA)) { \
95: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
96: "cuBLAS error %d (%s). " \
97: "Reports not initialized or alloc failed; " \
98: "this indicates the GPU may have run out resources", \
99: (PetscErrorCode)_p_cublas_stat__, name); \
100: } else { \
101: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuBLAS error %d (%s)", (PetscErrorCode)_p_cublas_stat__, name); \
102: } \
103: } \
104: } while (0)
105: #define CHKERRCUBLAS(...) PetscCallCUBLAS(__VA_ARGS__)
107: #if (CUSPARSE_VER_MAJOR > 10 || CUSPARSE_VER_MAJOR == 10 && CUSPARSE_VER_MINOR >= 2) /* According to cuda/10.1.168 on OLCF Summit */
108: #define PetscCallCUSPARSE(...) \
109: do { \
110: const cusparseStatus_t _p_cusparse_stat__ = __VA_ARGS__; \
111: if (PetscUnlikely(_p_cusparse_stat__)) { \
112: const char *name = cusparseGetErrorName(_p_cusparse_stat__); \
113: const char *descr = cusparseGetErrorString(_p_cusparse_stat__); \
114: PetscCheck((_p_cusparse_stat__ != CUSPARSE_STATUS_NOT_INITIALIZED) && (_p_cusparse_stat__ != CUSPARSE_STATUS_ALLOC_FAILED), PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
115: "cuSPARSE errorcode %d (%s) : %s.; " \
116: "this indicates the GPU has run out resources", \
117: (int)_p_cusparse_stat__, name, descr); \
118: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuSPARSE errorcode %d (%s) : %s", (int)_p_cusparse_stat__, name, descr); \
119: } \
120: } while (0)
121: #else /* (CUSPARSE_VER_MAJOR > 10 || CUSPARSE_VER_MAJOR == 10 && CUSPARSE_VER_MINOR >= 2) */
122: #define PetscCallCUSPARSE(...) \
123: do { \
124: const cusparseStatus_t _p_cusparse_stat__ = __VA_ARGS__; \
125: PetscCheck(_p_cusparse_stat__ == CUSPARSE_STATUS_SUCCESS, PETSC_COMM_SELF, PETSC_ERR_GPU, "cuSPARSE errorcode %d", (PetscErrorCode)_p_cusparse_stat__); \
126: } while (0)
127: #endif /* (CUSPARSE_VER_MAJOR > 10 || CUSPARSE_VER_MAJOR == 10 && CUSPARSE_VER_MINOR >= 2) */
128: #define CHKERRCUSPARSE(...) PetscCallCUSPARSE(__VA_ARGS__)
130: #define PetscCallCUSOLVER(...) \
131: do { \
132: const cusolverStatus_t _p_cusolver_stat__ = __VA_ARGS__; \
133: if (PetscUnlikely(_p_cusolver_stat__ != CUSOLVER_STATUS_SUCCESS)) { \
134: const char *name = PetscCUSolverGetErrorName(_p_cusolver_stat__); \
135: if ((_p_cusolver_stat__ == CUSOLVER_STATUS_NOT_INITIALIZED || _p_cusolver_stat__ == CUSOLVER_STATUS_ALLOC_FAILED || _p_cusolver_stat__ == CUSOLVER_STATUS_INTERNAL_ERROR) && PetscDeviceInitialized(PETSC_DEVICE_CUDA)) { \
136: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
137: "cuSolver error %d (%s). " \
138: "This indicates the GPU may have run out resources", \
139: (PetscErrorCode)_p_cusolver_stat__, name); \
140: } else { \
141: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuSolver error %d (%s)", (PetscErrorCode)_p_cusolver_stat__, name); \
142: } \
143: } \
144: } while (0)
145: #define CHKERRCUSOLVER(...) PetscCallCUSOLVER(__VA_ARGS__)
147: #define PetscCallCUFFT(...) \
148: do { \
149: const cufftResult_t _p_cufft_stat__ = __VA_ARGS__; \
150: if (PetscUnlikely(_p_cufft_stat__ != CUFFT_SUCCESS)) { \
151: const char *name = PetscCUFFTGetErrorName(_p_cufft_stat__); \
152: if ((_p_cufft_stat__ == CUFFT_SETUP_FAILED || _p_cufft_stat__ == CUFFT_ALLOC_FAILED) && PetscDeviceInitialized(PETSC_DEVICE_CUDA)) { \
153: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
154: "cuFFT error %d (%s). " \
155: "Reports not initialized or alloc failed; " \
156: "this indicates the GPU has run out resources", \
157: (PetscErrorCode)_p_cufft_stat__, name); \
158: } else { \
159: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuFFT error %d (%s)", (PetscErrorCode)_p_cufft_stat__, name); \
160: } \
161: } \
162: } while (0)
163: #define CHKERRCUFFT(...) PetscCallCUFFT(__VA_ARGS__)
165: #define PetscCallCURAND(...) \
166: do { \
167: const curandStatus_t _p_curand_stat__ = __VA_ARGS__; \
168: if (PetscUnlikely(_p_curand_stat__ != CURAND_STATUS_SUCCESS)) { \
169: if ((_p_curand_stat__ == CURAND_STATUS_INITIALIZATION_FAILED || _p_curand_stat__ == CURAND_STATUS_ALLOCATION_FAILED) && PetscDeviceInitialized(PETSC_DEVICE_CUDA)) { \
170: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
171: "cuRAND error %d. " \
172: "Reports not initialized or alloc failed; " \
173: "this indicates the GPU has run out resources", \
174: (PetscErrorCode)_p_curand_stat__); \
175: } else { \
176: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuRand error %d", (PetscErrorCode)_p_curand_stat__); \
177: } \
178: } \
179: } while (0)
180: #define CHKERRCURAND(...) PetscCallCURAND(__VA_ARGS__)
182: PETSC_EXTERN cudaStream_t PetscDefaultCudaStream; // The default stream used by PETSc
183: PETSC_EXTERN PetscErrorCode PetscCUBLASGetHandle(cublasHandle_t *);
184: PETSC_EXTERN PetscErrorCode PetscCUSOLVERDnGetHandle(cusolverDnHandle_t *);
185: PETSC_EXTERN PetscErrorCode PetscGetCurrentCUDAStream(cudaStream_t *);
187: #endif // PETSC_HAVE_CUDA
189: // these can also be defined in petscdevice_hip.h so we undef and define them *only* if the
190: // current compiler is NVCC. In this case if petscdevice_hip.h is included first, the macros
191: // would already be defined, but they would be empty since we cannot be using HCC at the same
192: // time.
193: #if PetscDefined(USING_NVCC)
194: #undef PETSC_HOST_DECL
195: #undef PETSC_DEVICE_DECL
196: #undef PETSC_KERNEL_DECL
197: #undef PETSC_SHAREDMEM_DECL
198: #undef PETSC_FORCEINLINE
199: #undef PETSC_CONSTMEM_DECL
201: #define PETSC_HOST_DECL __host__
202: #define PETSC_DEVICE_DECL __device__
203: #define PETSC_KERNEL_DECL __global__
204: #define PETSC_SHAREDMEM_DECL __shared__
205: #define PETSC_FORCEINLINE __forceinline__
206: #define PETSC_CONSTMEM_DECL __constant__
207: #endif
209: #if !defined(PETSC_HOST_DECL) // use HOST_DECL as canary
210: #define PETSC_HOST_DECL
211: #define PETSC_DEVICE_DECL
212: #define PETSC_KERNEL_DECL
213: #define PETSC_SHAREDMEM_DECL
214: #define PETSC_FORCEINLINE inline
215: #define PETSC_CONSTMEM_DECL
216: #endif
218: #if !PetscDefined(DEVICE_DEFINED_DECLS_PRIVATE)
219: #define PETSC_DEVICE_DEFINED_DECLS_PRIVATE
220: #define PETSC_HOSTDEVICE_DECL PETSC_HOST_DECL PETSC_DEVICE_DECL
221: #define PETSC_DEVICE_INLINE_DECL PETSC_DEVICE_DECL PETSC_FORCEINLINE
222: #define PETSC_HOSTDEVICE_INLINE_DECL PETSC_HOSTDEVICE_DECL PETSC_FORCEINLINE
223: #endif