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>
23: /* cuBLAS does not have cublasGetErrorName(). We create one on our own. */
24: PETSC_EXTERN const char *PetscCUBLASGetErrorName(cublasStatus_t); /* PETSC_EXTERN since it is exposed by the CHKERRCUBLAS macro */
25: PETSC_EXTERN const char *PetscCUSolverGetErrorName(cusolverStatus_t);
26: PETSC_EXTERN const char *PetscCUFFTGetErrorName(cufftResult);
28: /* REMOVE ME */
29: #define WaitForCUDA() cudaDeviceSynchronize()
31: /* CUDART_VERSION = 1000 x major + 10 x minor version */
33: /* Could not find exactly which CUDART_VERSION introduced cudaGetErrorName. At least it was in CUDA 8.0 (Sep. 2016) */
34: #if PETSC_PKG_CUDA_VERSION_GE(8, 0, 0)
35: #define PetscCallCUDAVoid(...) \
36: do { \
37: const cudaError_t _p_cuda_err__ = __VA_ARGS__; \
38: 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__)); \
39: } while (0)
41: #define PetscCallCUDA(...) \
42: do { \
43: const cudaError_t _p_cuda_err__ = __VA_ARGS__; \
44: 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__)); \
45: } while (0)
46: #else /* PETSC_PKG_CUDA_VERSION_GE(8,0,0) */
47: #define PetscCallCUDA(...) \
48: do { \
49: const cudaError_t _p_cuda_err__ = __VA_ARGS__; \
50: PetscCheck(_p_cuda_err__ == cudaSuccess, PETSC_COMM_SELF, PETSC_ERR_GPU, "cuda error %d", (PetscErrorCode)_p_cuda_err__); \
51: } while (0)
53: #define PetscCallCUDAVoid(...) \
54: do { \
55: const cudaError_t _p_cuda_err__ = __VA_ARGS__; \
56: PetscCheckAbort(_p_cuda_err__ == cudaSuccess, PETSC_COMM_SELF, PETSC_ERR_GPU, "cuda error %d", (PetscErrorCode)_p_cuda_err__); \
57: } while (0)
58: #endif /* PETSC_PKG_CUDA_VERSION_GE(8,0,0) */
59: #define CHKERRCUDA(...) PetscCallCUDA(__VA_ARGS__)
61: #define PetscCUDACheckLaunch \
62: do { \
63: /* Check synchronous errors, i.e. pre-launch */ \
64: PetscCallCUDA(cudaGetLastError()); \
65: /* Check asynchronous errors, i.e. kernel failed (ULF) */ \
66: PetscCallCUDA(cudaDeviceSynchronize()); \
67: } while (0)
69: #define PetscCallCUBLAS(...) \
70: do { \
71: const cublasStatus_t _p_cublas_stat__ = __VA_ARGS__; \
72: if (PetscUnlikely(_p_cublas_stat__ != CUBLAS_STATUS_SUCCESS)) { \
73: const char *name = PetscCUBLASGetErrorName(_p_cublas_stat__); \
74: if (((_p_cublas_stat__ == CUBLAS_STATUS_NOT_INITIALIZED) || (_p_cublas_stat__ == CUBLAS_STATUS_ALLOC_FAILED)) && PetscDeviceInitialized(PETSC_DEVICE_CUDA)) { \
75: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
76: "cuBLAS error %d (%s). " \
77: "Reports not initialized or alloc failed; " \
78: "this indicates the GPU may have run out resources", \
79: (PetscErrorCode)_p_cublas_stat__, name); \
80: } else { \
81: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuBLAS error %d (%s)", (PetscErrorCode)_p_cublas_stat__, name); \
82: } \
83: } \
84: } while (0)
85: #define CHKERRCUBLAS(...) PetscCallCUBLAS(__VA_ARGS__)
87: #if (CUSPARSE_VER_MAJOR > 10 || CUSPARSE_VER_MAJOR == 10 && CUSPARSE_VER_MINOR >= 2) /* According to cuda/10.1.168 on OLCF Summit */
88: #define PetscCallCUSPARSE(...) \
89: do { \
90: const cusparseStatus_t _p_cusparse_stat__ = __VA_ARGS__; \
91: if (PetscUnlikely(_p_cusparse_stat__)) { \
92: const char *name = cusparseGetErrorName(_p_cusparse_stat__); \
93: const char *descr = cusparseGetErrorString(_p_cusparse_stat__); \
94: PetscCheck((_p_cusparse_stat__ != CUSPARSE_STATUS_NOT_INITIALIZED) && (_p_cusparse_stat__ != CUSPARSE_STATUS_ALLOC_FAILED), PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
95: "cuSPARSE errorcode %d (%s) : %s.; " \
96: "this indicates the GPU has run out resources", \
97: (int)_p_cusparse_stat__, name, descr); \
98: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuSPARSE errorcode %d (%s) : %s", (int)_p_cusparse_stat__, name, descr); \
99: } \
100: } while (0)
101: #else /* (CUSPARSE_VER_MAJOR > 10 || CUSPARSE_VER_MAJOR == 10 && CUSPARSE_VER_MINOR >= 2) */
102: #define PetscCallCUSPARSE(...) \
103: do { \
104: const cusparseStatus_t _p_cusparse_stat__ = __VA_ARGS__; \
105: PetscCheck(_p_cusparse_stat__ == CUSPARSE_STATUS_SUCCESS, PETSC_COMM_SELF, PETSC_ERR_GPU, "cuSPARSE errorcode %d", (PetscErrorCode)_p_cusparse_stat__); \
106: } while (0)
107: #endif /* (CUSPARSE_VER_MAJOR > 10 || CUSPARSE_VER_MAJOR == 10 && CUSPARSE_VER_MINOR >= 2) */
108: #define CHKERRCUSPARSE(...) PetscCallCUSPARSE(__VA_ARGS__)
110: #define PetscCallCUSOLVER(...) \
111: do { \
112: const cusolverStatus_t _p_cusolver_stat__ = __VA_ARGS__; \
113: if (PetscUnlikely(_p_cusolver_stat__ != CUSOLVER_STATUS_SUCCESS)) { \
114: const char *name = PetscCUSolverGetErrorName(_p_cusolver_stat__); \
115: 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)) { \
116: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
117: "cuSolver error %d (%s). " \
118: "This indicates the GPU may have run out resources", \
119: (PetscErrorCode)_p_cusolver_stat__, name); \
120: } else { \
121: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuSolver error %d (%s)", (PetscErrorCode)_p_cusolver_stat__, name); \
122: } \
123: } \
124: } while (0)
125: #define CHKERRCUSOLVER(...) PetscCallCUSOLVER(__VA_ARGS__)
127: #define PetscCallCUFFT(...) \
128: do { \
129: const cufftResult_t _p_cufft_stat__ = __VA_ARGS__; \
130: if (PetscUnlikely(_p_cufft_stat__ != CUFFT_SUCCESS)) { \
131: const char *name = PetscCUFFTGetErrorName(_p_cufft_stat__); \
132: if (((_p_cufft_stat__ == CUFFT_SETUP_FAILED) || (_p_cufft_stat__ == CUFFT_ALLOC_FAILED)) && PetscDeviceInitialized(PETSC_DEVICE_CUDA)) { \
133: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
134: "cuFFT error %d (%s). " \
135: "Reports not initialized or alloc failed; " \
136: "this indicates the GPU has run out resources", \
137: (PetscErrorCode)_p_cufft_stat__, name); \
138: } else { \
139: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuFFT error %d (%s)", (PetscErrorCode)_p_cufft_stat__, name); \
140: } \
141: } \
142: } while (0)
143: #define CHKERRCUFFT(...) PetscCallCUFFT(__VA_ARGS__)
145: #define PetscCallCURAND(...) \
146: do { \
147: const curandStatus_t _p_curand_stat__ = __VA_ARGS__; \
148: if (PetscUnlikely(_p_curand_stat__ != CURAND_STATUS_SUCCESS)) { \
149: if (((_p_curand_stat__ == CURAND_STATUS_INITIALIZATION_FAILED) || (_p_curand_stat__ == CURAND_STATUS_ALLOCATION_FAILED)) && PetscDeviceInitialized(PETSC_DEVICE_CUDA)) { \
150: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU_RESOURCE, \
151: "cuRAND error %d. " \
152: "Reports not initialized or alloc failed; " \
153: "this indicates the GPU has run out resources", \
154: (PetscErrorCode)_p_curand_stat__); \
155: } else { \
156: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_GPU, "cuRand error %d", (PetscErrorCode)_p_curand_stat__); \
157: } \
158: } \
159: } while (0)
160: #define CHKERRCURAND(...) PetscCallCURAND(__VA_ARGS__)
162: PETSC_EXTERN cudaStream_t PetscDefaultCudaStream; // The default stream used by PETSc
163: PETSC_EXTERN PetscErrorCode PetscCUBLASGetHandle(cublasHandle_t *);
164: PETSC_EXTERN PetscErrorCode PetscCUSOLVERDnGetHandle(cusolverDnHandle_t *);
165: PETSC_EXTERN PetscErrorCode PetscGetCurrentCUDAStream(cudaStream_t *);
167: #endif // PETSC_HAVE_CUDA
169: // these can also be defined in petscdevice_hip.h so we undef and define them *only* if the
170: // current compiler is NVCC. In this case if petscdevice_hip.h is included first, the macros
171: // would already be defined, but they would be empty since we cannot be using HCC at the same
172: // time.
173: #if PetscDefined(USING_NVCC)
174: #undef PETSC_HOST_DECL
175: #undef PETSC_DEVICE_DECL
176: #undef PETSC_KERNEL_DECL
177: #undef PETSC_SHAREDMEM_DECL
178: #undef PETSC_FORCEINLINE
179: #undef PETSC_CONSTMEM_DECL
181: #define PETSC_HOST_DECL __host__
182: #define PETSC_DEVICE_DECL __device__
183: #define PETSC_KERNEL_DECL __global__
184: #define PETSC_SHAREDMEM_DECL __shared__
185: #define PETSC_FORCEINLINE __forceinline__
186: #define PETSC_CONSTMEM_DECL __constant__
187: #endif
189: #ifndef PETSC_HOST_DECL // use HOST_DECL as canary
190: #define PETSC_HOST_DECL
191: #define PETSC_DEVICE_DECL
192: #define PETSC_KERNEL_DECL
193: #define PETSC_SHAREDMEM_DECL
194: #define PETSC_FORCEINLINE inline
195: #define PETSC_CONSTMEM_DECL
196: #endif
198: #ifndef PETSC_DEVICE_DEFINED_DECLS_PRIVATE
199: #define PETSC_DEVICE_DEFINED_DECLS_PRIVATE
200: #define PETSC_HOSTDEVICE_DECL PETSC_HOST_DECL PETSC_DEVICE_DECL
201: #define PETSC_DEVICE_INLINE_DECL PETSC_DEVICE_DECL PETSC_FORCEINLINE
202: #define PETSC_HOSTDEVICE_INLINE_DECL PETSC_HOSTDEVICE_DECL PETSC_FORCEINLINE
203: #endif