Actual source code: cupminterface.hpp
1: #pragma once
3: #include <petscdevice_cupm.h>
4: #include <petscpkg_version.h>
6: #include <petsc/private/cpputil.hpp>
8: #include <petsc/private/cpp/array.hpp>
10: namespace Petsc
11: {
13: namespace device
14: {
16: namespace cupm
17: {
19: // enum describing available cupm devices, this is used as the template parameter to any
20: // class subclassing the Interface or using it as a member variable
21: enum class DeviceType : int {
22: CUDA,
23: HIP
24: };
26: // clang-format off
27: static constexpr std::array<const char *const, 5> DeviceTypes = {
28: "cuda",
29: "hip",
30: "Petsc::device::cupm::DeviceType",
31: "Petsc::device::cupm::DeviceType::",
32: nullptr
33: };
34: // clang-format on
36: namespace impl
37: {
39: #define PetscCallCUPM_(__abort_fn__, __comm__, ...) \
40: do { \
41: PetscStackUpdateLine; \
42: const cupmError_t cerr_p_ = __VA_ARGS__; \
43: __abort_fn__(cerr_p_ == cupmSuccess, __comm__, PETSC_ERR_GPU, "%s error %d (%s) : %s", cupmName(), static_cast<PetscErrorCode>(cerr_p_), cupmGetErrorName(cerr_p_), cupmGetErrorString(cerr_p_)); \
44: } while (0)
46: // A backend agnostic PetscCallCUPM() function, this will only work inside the member
47: // functions of a class inheriting from CUPM::Interface. Thanks to __VA_ARGS__ templated
48: // functions can also be wrapped inline:
49: //
50: // PetscCallCUPM(foo());
51: #define PetscCallCUPM(...) PetscCallCUPM_(PetscCheck, PETSC_COMM_SELF, __VA_ARGS__)
52: #define PetscCallCUPMAbort(comm_, ...) PetscCallCUPM_(PetscCheckAbort, comm_, __VA_ARGS__)
54: // PETSC_CUPM_ALIAS_FUNCTION() - declaration to alias a cuda/hip function
55: //
56: // input params:
57: // our_name - the name of the alias
58: // their_name - the name of the function being aliased
59: //
60: // notes:
61: // see PETSC_ALIAS_FUNCTION() for the exact nature of the expansion
62: //
63: // example usage:
64: // PETSC_CUPM_ALIAS_FUNCTION(cupmMalloc, cudaMalloc) ->
65: // template
66: // static constexpr auto cupmMalloc(T&&... args) *noexcept and trailing return type deduction*
67: // {
68: // return cudaMalloc(std::forward(args)...);
69: // }
70: //
71: // PETSC_CUPM_ALIAS_FUNCTION(cupmMalloc, hipMalloc) ->
72: // template
73: // static constexpr auto cupmMalloc(T&&... args) *noexcept and trailing return type deduction*
74: // {
75: // return hipMalloc(std::forward(args)...);
76: // }
77: #define PETSC_CUPM_ALIAS_FUNCTION(our_name, their_name) PETSC_ALIAS_FUNCTION(static our_name, their_name)
79: // PETSC_CUPM_ALIAS_FUNCTION_GOBBLE() - declaration to alias a cuda/hip function but
80: // discard the last N arguments
81: //
82: // input params:
83: // our_name - the name of the alias
84: // their_name - the name of the function being aliased
85: // N - integer constant [0, INT_MAX) dictating how many arguments to chop off the end
86: //
87: // notes:
88: // see PETSC_ALIAS_FUNCTION_GOBBLE_NTH_LAST_ARGS() for the exact nature of the expansion
89: //
90: // example use:
91: // PETSC_CUPM_ALIAS_FUNCTION_GOBBLE_COMMON(cupmMallocAsync, cudaMalloc, 1) ->
92: // template
93: // static constexpr auto cupmMallocAsync(T&&... args, Tend argend) *noexcept and trailing
94: // return type deduction*
95: // {
96: // (void)argend;
97: // return cudaMalloc(std::forward(args)...);
98: // }
99: #define PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(our_name, their_name, N) PETSC_ALIAS_FUNCTION_GOBBLE_NTH_LAST_ARGS(static our_name, their_name, N)
101: // Base class that holds functions and variables that don't require CUDA or HIP to be present
102: // on the system
103: template <DeviceType T>
104: struct InterfaceBase {
105: static const DeviceType type = T;
107: PETSC_NODISCARD static constexpr const char *cupmName() noexcept
108: {
109: static_assert(util::to_underlying(DeviceType::CUDA) == 0, "");
110: static_assert(util::to_underlying(DeviceType::HIP) == 1, "");
111: return std::get<util::to_underlying(T)>(DeviceTypes);
112: }
114: PETSC_NODISCARD static constexpr const char *cupmNAME() noexcept { return T == DeviceType::CUDA ? "CUDA" : "HIP"; }
116: PETSC_NODISCARD static constexpr PetscDeviceType PETSC_DEVICE_CUPM() noexcept { return T == DeviceType::CUDA ? PETSC_DEVICE_CUDA : PETSC_DEVICE_HIP; }
118: PETSC_NODISCARD static constexpr PetscMemType PETSC_MEMTYPE_CUPM() noexcept { return T == DeviceType::CUDA ? PETSC_MEMTYPE_CUDA : PETSC_MEMTYPE_HIP; }
119: };
121: // declare the base class static member variables
122: template <DeviceType T>
123: const DeviceType InterfaceBase<T>::type;
125: #define PETSC_CUPM_BASE_CLASS_HEADER(T) \
126: using ::Petsc::device::cupm::impl::InterfaceBase<T>::type; \
127: using ::Petsc::device::cupm::impl::InterfaceBase<T>::cupmName; \
128: using ::Petsc::device::cupm::impl::InterfaceBase<T>::cupmNAME; \
129: using ::Petsc::device::cupm::impl::InterfaceBase<T>::PETSC_DEVICE_CUPM; \
130: using ::Petsc::device::cupm::impl::InterfaceBase<T>::PETSC_MEMTYPE_CUPM
132: // A templated C++ struct that defines the entire CUPM interface. Use of templating vs
133: // preprocessor macros allows us to use both interfaces simultaneously as well as easily
134: // import them into classes.
135: template <DeviceType>
136: struct InterfaceImpl;
138: #if PetscDefined(HAVE_CUDA)
139: template <>
140: struct PETSC_SINGLE_LIBRARY_VISIBILITY_INTERNAL InterfaceImpl<DeviceType::CUDA> : InterfaceBase<DeviceType::CUDA> {
141: PETSC_CUPM_BASE_CLASS_HEADER(DeviceType::CUDA);
143: // typedefs
144: using cupmError_t = cudaError_t;
145: using cupmEvent_t = cudaEvent_t;
146: using cupmStream_t = cudaStream_t;
147: using cupmDeviceProp_t = cudaDeviceProp;
148: using cupmMemcpyKind_t = cudaMemcpyKind;
149: using cupmDeviceAttr_t = cudaDeviceAttr;
150: using cupmComplex_t = util::conditional_t<PetscDefined(USE_REAL_SINGLE), cuComplex, cuDoubleComplex>;
151: using cupmPointerAttributes_t = cudaPointerAttributes;
152: using cupmMemoryType_t = enum cudaMemoryType;
153: using cupmDim3 = dim3;
154: using cupmHostFn_t = cudaHostFn_t;
155: #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
156: using cupmMemPool_t = cudaMemPool_t;
157: using cupmMemPoolAttr = cudaMemPoolAttr;
158: #else
159: using cupmMemPool_t = void *;
160: using cupmMemPoolAttr = unsigned int;
161: #endif
163: // values
164: static const auto cupmSuccess = cudaSuccess;
165: static const auto cupmErrorNotReady = cudaErrorNotReady;
166: static const auto cupmErrorDeviceAlreadyInUse = cudaErrorDeviceAlreadyInUse;
167: static const auto cupmErrorSetOnActiveProcess = cudaErrorSetOnActiveProcess;
168: static const auto cupmErrorStubLibrary =
169: #if PETSC_PKG_CUDA_VERSION_GE(11, 1, 0)
170: cudaErrorStubLibrary;
171: #else
172: cudaErrorInsufficientDriver;
173: #endif
175: static const auto cupmErrorNoDevice = cudaErrorNoDevice;
176: static const auto cupmStreamDefault = cudaStreamDefault;
177: static const auto cupmStreamNonBlocking = cudaStreamNonBlocking;
178: static const auto cupmDeviceMapHost = cudaDeviceMapHost;
179: static const auto cupmMemcpyHostToDevice = cudaMemcpyHostToDevice;
180: static const auto cupmMemcpyDeviceToHost = cudaMemcpyDeviceToHost;
181: static const auto cupmMemcpyDeviceToDevice = cudaMemcpyDeviceToDevice;
182: static const auto cupmMemcpyHostToHost = cudaMemcpyHostToHost;
183: static const auto cupmMemcpyDefault = cudaMemcpyDefault;
184: static const auto cupmMemoryTypeHost = cudaMemoryTypeHost;
185: static const auto cupmMemoryTypeDevice = cudaMemoryTypeDevice;
186: static const auto cupmMemoryTypeManaged = cudaMemoryTypeManaged;
187: static const auto cupmEventDisableTiming = cudaEventDisableTiming;
188: static const auto cupmHostAllocDefault = cudaHostAllocDefault;
189: static const auto cupmHostAllocWriteCombined = cudaHostAllocWriteCombined;
190: static const auto cupmMemPoolAttrReleaseThreshold =
191: #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
192: cudaMemPoolAttrReleaseThreshold;
193: #else
194: cupmMemPoolAttr{0};
195: #endif
196: static const auto cupmDevAttrClockRate = cudaDevAttrClockRate;
197: static const auto cupmDevAttrMemoryClockRate = cudaDevAttrMemoryClockRate;
199: // error functions
200: PETSC_CUPM_ALIAS_FUNCTION(cupmGetErrorName, cudaGetErrorName)
201: PETSC_CUPM_ALIAS_FUNCTION(cupmGetErrorString, cudaGetErrorString)
202: PETSC_CUPM_ALIAS_FUNCTION(cupmGetLastError, cudaGetLastError)
204: // device management
205: PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceCount, cudaGetDeviceCount)
206: PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceProperties, cudaGetDeviceProperties)
207: PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceGetAttribute, cudaDeviceGetAttribute)
208: PETSC_CUPM_ALIAS_FUNCTION(cupmGetDevice, cudaGetDevice)
209: PETSC_CUPM_ALIAS_FUNCTION(cupmSetDevice, cudaSetDevice)
210: PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceFlags, cudaGetDeviceFlags)
211: PETSC_CUPM_ALIAS_FUNCTION(cupmSetDeviceFlags, cudaSetDeviceFlags)
212: PETSC_CUPM_ALIAS_FUNCTION(cupmPointerGetAttributes, cudaPointerGetAttributes)
213: PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceGetPCIBusId, cudaDeviceGetPCIBusId)
215: PETSC_NODISCARD static cupmMemoryType_t cupmPointerGetMemType(const cupmPointerAttributes_t &attr, PetscBool *managed) noexcept
216: {
217: #if CUDART_VERSION < 10000
218: if (managed) *managed = static_cast<PetscBool>(attr.isManaged);
219: return attr.memoryType;
220: #else
221: if (managed) *managed = static_cast<PetscBool>(attr.type == cupmMemoryTypeManaged);
222: return attr.type;
223: #endif
224: }
226: #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
227: PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceGetMemPool, cudaDeviceGetMemPool)
228: PETSC_CUPM_ALIAS_FUNCTION(cupmMemPoolSetAttribute, cudaMemPoolSetAttribute)
229: #else
230: PETSC_NODISCARD static cupmError_t cupmDeviceGetMemPool(cupmMemPool_t *pool, int) noexcept
231: {
232: *pool = nullptr;
233: return cupmSuccess;
234: }
236: PETSC_NODISCARD static cupmError_t cupmMemPoolSetAttribute(cupmMemPool_t, cupmMemPoolAttr, void *) noexcept { return cupmSuccess; }
237: #endif
238: // CUDA has no cudaInit() to match hipInit()
239: PETSC_NODISCARD static cupmError_t cupmInit(unsigned int) noexcept { return cudaFree(nullptr); }
241: // stream management
242: PETSC_CUPM_ALIAS_FUNCTION(cupmEventCreate, cudaEventCreate)
243: PETSC_CUPM_ALIAS_FUNCTION(cupmEventCreateWithFlags, cudaEventCreateWithFlags)
244: PETSC_CUPM_ALIAS_FUNCTION(cupmEventDestroy, cudaEventDestroy)
245: PETSC_CUPM_ALIAS_FUNCTION(cupmEventRecord, cudaEventRecord)
246: PETSC_CUPM_ALIAS_FUNCTION(cupmEventSynchronize, cudaEventSynchronize)
247: PETSC_CUPM_ALIAS_FUNCTION(cupmEventElapsedTime, cudaEventElapsedTime)
248: PETSC_CUPM_ALIAS_FUNCTION(cupmEventQuery, cudaEventQuery)
249: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamCreate, cudaStreamCreate)
250: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamCreateWithFlags, cudaStreamCreateWithFlags)
251: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamGetFlags, cudaStreamGetFlags)
252: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamDestroy, cudaStreamDestroy)
253: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamWaitEvent, cudaStreamWaitEvent)
254: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamQuery, cudaStreamQuery)
255: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamSynchronize, cudaStreamSynchronize)
256: PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceSynchronize, cudaDeviceSynchronize)
257: PETSC_CUPM_ALIAS_FUNCTION(cupmGetSymbolAddress, cudaGetSymbolAddress)
259: // memory management
260: PETSC_CUPM_ALIAS_FUNCTION(cupmFree, cudaFree)
261: PETSC_CUPM_ALIAS_FUNCTION(cupmMalloc, cudaMalloc)
262: #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
263: PETSC_CUPM_ALIAS_FUNCTION(cupmFreeAsync, cudaFreeAsync)
264: PETSC_CUPM_ALIAS_FUNCTION(cupmMallocAsync, cudaMallocAsync)
265: #else
266: PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmFreeAsync, cudaFree, 1)
267: PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMallocAsync, cudaMalloc, 1)
268: #endif
269: PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy, cudaMemcpy)
270: PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpyAsync, cudaMemcpyAsync)
271: PETSC_CUPM_ALIAS_FUNCTION(cupmMallocHost, cudaMallocHost)
272: PETSC_CUPM_ALIAS_FUNCTION(cupmFreeHost, cudaFreeHost)
273: PETSC_CUPM_ALIAS_FUNCTION(cupmMemset, cudaMemset)
274: #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
275: PETSC_CUPM_ALIAS_FUNCTION(cupmMemsetAsync, cudaMemsetAsync)
276: #else
277: PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMemsetAsync, cudaMemset, 1)
278: #endif
279: PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy2D, cudaMemcpy2D)
280: PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy2DAsync, cudaMemcpy2DAsync)
281: PETSC_CUPM_ALIAS_FUNCTION(cupmMemset2D, cudaMemset2D)
282: PETSC_CUPM_ALIAS_FUNCTION(cupmMemset2DAsync, cudaMemset2DAsync)
284: // launch control
285: PETSC_CUPM_ALIAS_FUNCTION(cupmLaunchHostFunc, cudaLaunchHostFunc)
286: template <typename FunctionT, typename... KernelArgsT>
287: PETSC_NODISCARD static cudaError_t cupmLaunchKernel(FunctionT &&func, dim3 gridDim, dim3 blockDim, std::size_t sharedMem, cudaStream_t stream, KernelArgsT &&...kernelArgs) noexcept
288: {
289: static_assert(!std::is_pointer<FunctionT>::value, "kernel function must not be passed by pointer");
290: void *args[] = {(void *)std::addressof(kernelArgs)...};
292: return cudaLaunchKernel<util::remove_reference_t<FunctionT>>(std::addressof(func), std::move(gridDim), std::move(blockDim), args, sharedMem, std::move(stream));
293: }
294: };
295: #endif // PetscDefined(HAVE_CUDA)
297: #if PetscDefined(HAVE_HIP)
298: template <>
299: struct PETSC_SINGLE_LIBRARY_VISIBILITY_INTERNAL InterfaceImpl<DeviceType::HIP> : InterfaceBase<DeviceType::HIP> {
300: PETSC_CUPM_BASE_CLASS_HEADER(DeviceType::HIP);
302: // typedefs
303: using cupmError_t = hipError_t;
304: using cupmEvent_t = hipEvent_t;
305: using cupmStream_t = hipStream_t;
306: using cupmDeviceProp_t = hipDeviceProp_t;
307: using cupmMemcpyKind_t = hipMemcpyKind;
308: using cupmDeviceAttr_t = hipDeviceAttribute_t;
309: using cupmComplex_t = util::conditional_t<PetscDefined(USE_REAL_SINGLE), hipComplex, hipDoubleComplex>;
310: using cupmPointerAttributes_t = hipPointerAttribute_t;
311: using cupmMemoryType_t = enum hipMemoryType;
312: using cupmDim3 = dim3;
313: #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
314: using cupmHostFn_t = hipHostFn_t;
315: using cupmMemPool_t = hipMemPool_t;
316: using cupmMemPoolAttr = hipMemPoolAttr;
317: #else
318: using cupmHostFn_t = void (*)(void *);
319: using cupmMemPool_t = void *;
320: using cupmMemPoolAttr = unsigned int;
321: #endif
323: // values
324: static const auto cupmSuccess = hipSuccess;
325: static const auto cupmErrorNotReady = hipErrorNotReady;
326: // see https://github.com/ROCm-Developer-Tools/HIP/blob/develop/bin/hipify-perl
327: static const auto cupmErrorDeviceAlreadyInUse = hipErrorContextAlreadyInUse;
328: static const auto cupmErrorSetOnActiveProcess = hipErrorSetOnActiveProcess;
329: // as of HIP v4.2 cudaErrorStubLibrary has no HIP equivalent
330: static const auto cupmErrorStubLibrary = hipErrorInsufficientDriver;
331: static const auto cupmErrorNoDevice = hipErrorNoDevice;
332: static const auto cupmStreamDefault = hipStreamDefault;
333: static const auto cupmStreamNonBlocking = hipStreamNonBlocking;
334: static const auto cupmDeviceMapHost = hipDeviceMapHost;
335: static const auto cupmMemcpyHostToDevice = hipMemcpyHostToDevice;
336: static const auto cupmMemcpyDeviceToHost = hipMemcpyDeviceToHost;
337: static const auto cupmMemcpyDeviceToDevice = hipMemcpyDeviceToDevice;
338: static const auto cupmMemcpyHostToHost = hipMemcpyHostToHost;
339: static const auto cupmMemcpyDefault = hipMemcpyDefault;
340: static const auto cupmMemoryTypeHost = hipMemoryTypeHost;
341: static const auto cupmMemoryTypeDevice = hipMemoryTypeDevice;
342: // see
343: // https://github.com/ROCm-Developer-Tools/HIP/blob/develop/include/hip/hip_runtime_api.h#L156
344: static const auto cupmMemoryTypeManaged = hipMemoryTypeUnified;
345: static const auto cupmEventDisableTiming = hipEventDisableTiming;
346: static const auto cupmHostAllocDefault = hipHostMallocDefault;
347: static const auto cupmHostAllocWriteCombined = hipHostMallocWriteCombined;
348: static const auto cupmMemPoolAttrReleaseThreshold =
349: #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
350: hipMemPoolAttrReleaseThreshold;
351: #else
352: cupmMemPoolAttr{0};
353: #endif
354: static const auto cupmDevAttrClockRate = hipDeviceAttributeClockRate;
355: static const auto cupmDevAttrMemoryClockRate = hipDeviceAttributeMemoryClockRate;
357: // error functions
358: PETSC_CUPM_ALIAS_FUNCTION(cupmGetErrorName, hipGetErrorName)
359: PETSC_CUPM_ALIAS_FUNCTION(cupmGetErrorString, hipGetErrorString)
360: PETSC_CUPM_ALIAS_FUNCTION(cupmGetLastError, hipGetLastError)
362: // device management
363: PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceCount, hipGetDeviceCount)
364: PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceProperties, hipGetDeviceProperties)
365: PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceGetAttribute, hipDeviceGetAttribute)
366: PETSC_CUPM_ALIAS_FUNCTION(cupmGetDevice, hipGetDevice)
367: PETSC_CUPM_ALIAS_FUNCTION(cupmSetDevice, hipSetDevice)
368: PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceFlags, hipGetDeviceFlags)
369: PETSC_CUPM_ALIAS_FUNCTION(cupmSetDeviceFlags, hipSetDeviceFlags)
370: PETSC_CUPM_ALIAS_FUNCTION(cupmPointerGetAttributes, hipPointerGetAttributes)
371: PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceGetPCIBusId, hipDeviceGetPCIBusId)
373: PETSC_NODISCARD static cupmMemoryType_t cupmPointerGetMemType(const cupmPointerAttributes_t &attr, PetscBool *managed) noexcept
374: {
375: // HIP used memoryType on both backends before 5.5.
376: #if PETSC_PKG_HIP_VERSION_LT(5, 5, 0)
377: if (managed) *managed = static_cast<PetscBool>(attr.isManaged);
378: return attr.memoryType;
379: #else
380: if (managed) *managed = static_cast<PetscBool>(attr.type == cupmMemoryTypeManaged);
381: return attr.type;
382: #endif
383: }
385: #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
386: PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceGetMemPool, hipDeviceGetMemPool)
387: PETSC_CUPM_ALIAS_FUNCTION(cupmMemPoolSetAttribute, hipMemPoolSetAttribute)
388: #else
389: PETSC_NODISCARD static cupmError_t cupmDeviceGetMemPool(cupmMemPool_t *pool, int) noexcept
390: {
391: *pool = nullptr;
392: return cupmSuccess;
393: }
395: PETSC_NODISCARD static cupmError_t cupmMemPoolSetAttribute(cupmMemPool_t, cupmMemPoolAttr, void *) noexcept { return cupmSuccess; }
396: #endif
397: PETSC_CUPM_ALIAS_FUNCTION(cupmInit, hipInit)
399: // stream management
400: PETSC_CUPM_ALIAS_FUNCTION(cupmEventCreate, hipEventCreate)
401: PETSC_CUPM_ALIAS_FUNCTION(cupmEventCreateWithFlags, hipEventCreateWithFlags)
402: PETSC_CUPM_ALIAS_FUNCTION(cupmEventDestroy, hipEventDestroy)
403: PETSC_CUPM_ALIAS_FUNCTION(cupmEventRecord, hipEventRecord)
404: PETSC_CUPM_ALIAS_FUNCTION(cupmEventSynchronize, hipEventSynchronize)
405: PETSC_CUPM_ALIAS_FUNCTION(cupmEventElapsedTime, hipEventElapsedTime)
406: PETSC_CUPM_ALIAS_FUNCTION(cupmEventQuery, hipEventQuery)
407: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamCreate, hipStreamCreate)
408: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamCreateWithFlags, hipStreamCreateWithFlags)
409: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamGetFlags, hipStreamGetFlags)
410: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamDestroy, hipStreamDestroy)
411: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamWaitEvent, hipStreamWaitEvent)
412: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamQuery, hipStreamQuery)
413: PETSC_CUPM_ALIAS_FUNCTION(cupmStreamSynchronize, hipStreamSynchronize)
414: PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceSynchronize, hipDeviceSynchronize)
415: PETSC_CUPM_ALIAS_FUNCTION(cupmGetSymbolAddress, hipGetSymbolAddress)
417: // memory management
418: PETSC_CUPM_ALIAS_FUNCTION(cupmFree, hipFree)
419: PETSC_CUPM_ALIAS_FUNCTION(cupmMalloc, hipMalloc)
420: #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
421: PETSC_CUPM_ALIAS_FUNCTION(cupmMallocAsync, hipMallocAsync)
422: PETSC_CUPM_ALIAS_FUNCTION(cupmFreeAsync, hipFreeAsync)
423: #else
424: PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMallocAsync, hipMalloc, 1)
425: PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmFreeAsync, hipFree, 1)
426: #endif
427: PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy, hipMemcpy)
428: PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpyAsync, hipMemcpyAsync)
429: // hipMallocHost is deprecated
430: PETSC_CUPM_ALIAS_FUNCTION(cupmMallocHost, hipHostMalloc)
431: // hipFreeHost is deprecated
432: PETSC_CUPM_ALIAS_FUNCTION(cupmFreeHost, hipHostFree)
433: PETSC_CUPM_ALIAS_FUNCTION(cupmMemset, hipMemset)
434: PETSC_CUPM_ALIAS_FUNCTION(cupmMemsetAsync, hipMemsetAsync)
435: PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy2D, hipMemcpy2D)
436: PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy2DAsync, hipMemcpy2DAsync)
437: PETSC_CUPM_ALIAS_FUNCTION(cupmMemset2D, hipMemset2D)
438: PETSC_CUPM_ALIAS_FUNCTION(cupmMemset2DAsync, hipMemset2DAsync)
440: // launch control
441: // HIP appears to only have hipLaunchHostFunc from 5.2.0 onwards
442: // https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md#7-execution-control=
443: #if PetscDefined(HAVE_HIPCUDA)
444: PETSC_NODISCARD static hipError_t cupmLaunchHostFunc(cupmStream_t stream, void (*fn)(void *), void *ctx) noexcept { return hipCUDAErrorTohipError(cudaLaunchHostFunc(stream, fn, ctx)); }
445: #elif PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
446: PETSC_CUPM_ALIAS_FUNCTION(cupmLaunchHostFunc, hipLaunchHostFunc)
447: #else
448: PETSC_NODISCARD static hipError_t cupmLaunchHostFunc(hipStream_t stream, cupmHostFn_t fn, void *ctx) noexcept
449: {
450: // the only correct way to spoof this function is to do it synchronously...
451: auto herr = hipStreamSynchronize(stream);
452: if (PetscUnlikely(herr != hipSuccess)) return herr;
453: fn(ctx);
454: return herr;
455: }
456: #endif
458: template <typename FunctionT, typename... KernelArgsT>
459: PETSC_NODISCARD static hipError_t cupmLaunchKernel(FunctionT &&func, dim3 gridDim, dim3 blockDim, std::size_t sharedMem, hipStream_t stream, KernelArgsT &&...kernelArgs) noexcept
460: {
461: void *args[] = {(void *)std::addressof(kernelArgs)...};
463: return hipLaunchKernel((void *)func, std::move(gridDim), std::move(blockDim), args, sharedMem, std::move(stream));
464: }
465: };
466: #endif // PetscDefined(HAVE_HIP)
468: // shorthand for bringing all of the typedefs from the base Interface class into your own,
469: // it's annoying that c++ doesn't have a way to do this automatically
470: #define PETSC_CUPM_IMPL_CLASS_HEADER(T) \
471: PETSC_CUPM_BASE_CLASS_HEADER(T); \
472: /* types */ \
473: using cupmError_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmError_t; \
474: using cupmEvent_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEvent_t; \
475: using cupmStream_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStream_t; \
476: using cupmDeviceProp_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceProp_t; \
477: using cupmMemcpyKind_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyKind_t; \
478: using cupmDeviceAttr_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceAttr_t; \
479: using cupmComplex_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmComplex_t; \
480: using cupmPointerAttributes_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmPointerAttributes_t; \
481: using cupmMemoryType_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemoryType_t; \
482: using cupmDim3 = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDim3; \
483: using cupmHostFn_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmHostFn_t; \
484: using cupmMemPool_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemPool_t; \
485: using cupmMemPoolAttr = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemPoolAttr; \
486: /* variables */ \
487: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmSuccess; \
488: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorNotReady; \
489: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorDeviceAlreadyInUse; \
490: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorSetOnActiveProcess; \
491: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorStubLibrary; \
492: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorNoDevice; \
493: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamDefault; \
494: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamNonBlocking; \
495: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceMapHost; \
496: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyHostToDevice; \
497: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyDeviceToHost; \
498: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyDeviceToDevice; \
499: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyHostToHost; \
500: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyDefault; \
501: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemoryTypeHost; \
502: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemoryTypeDevice; \
503: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemoryTypeManaged; \
504: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventDisableTiming; \
505: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmHostAllocDefault; \
506: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmHostAllocWriteCombined; \
507: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemPoolAttrReleaseThreshold; \
508: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDevAttrClockRate; \
509: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDevAttrMemoryClockRate; \
510: /* functions */ \
511: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetErrorName; \
512: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetErrorString; \
513: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetLastError; \
514: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetDeviceCount; \
515: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetDeviceProperties; \
516: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceGetAttribute; \
517: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetDevice; \
518: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmSetDevice; \
519: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetDeviceFlags; \
520: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmSetDeviceFlags; \
521: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmPointerGetAttributes; \
522: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceGetPCIBusId; \
523: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceGetMemPool; \
524: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemPoolSetAttribute; \
525: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmInit; \
526: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventCreate; \
527: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventCreateWithFlags; \
528: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventDestroy; \
529: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventRecord; \
530: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventSynchronize; \
531: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventElapsedTime; \
532: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventQuery; \
533: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamCreate; \
534: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamCreateWithFlags; \
535: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamGetFlags; \
536: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamDestroy; \
537: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamWaitEvent; \
538: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamQuery; \
539: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamSynchronize; \
540: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceSynchronize; \
541: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetSymbolAddress; \
542: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMalloc; \
543: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMallocAsync; \
544: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpy; \
545: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyAsync; \
546: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMallocHost; \
547: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemset; \
548: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemsetAsync; \
549: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpy2D; \
550: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpy2DAsync; \
551: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemset2D; \
552: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemset2DAsync; \
553: using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmLaunchHostFunc
555: #if PetscHasAttribute(always_inline)
556: // https://gcc.gnu.org/bugzilla//show_bug.cgi?id=109464
557: #define PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND __attribute__((always_inline))
558: #else
559: #define PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND
560: #endif
562: // The actual interface class
563: template <DeviceType T>
564: struct Interface : InterfaceImpl<T> {
565: private:
566: using interface_type = InterfaceImpl<T>;
568: public:
569: PETSC_CUPM_IMPL_CLASS_HEADER(T);
571: using cupmReal_t = util::conditional_t<PetscDefined(USE_REAL_SINGLE), float, double>;
572: using cupmScalar_t = util::conditional_t<PetscDefined(USE_COMPLEX), cupmComplex_t, cupmReal_t>;
574: PETSC_NODISCARD PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND static constexpr cupmScalar_t cupmScalarCast(PetscScalar s) noexcept
575: {
576: #if PetscDefined(USE_COMPLEX)
577: return cupmComplex_t{PetscRealPart(s), PetscImaginaryPart(s)};
578: #else
579: return static_cast<cupmScalar_t>(s);
580: #endif
581: }
583: PETSC_NODISCARD PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND static constexpr const cupmScalar_t *cupmScalarPtrCast(const PetscScalar *s) noexcept { return reinterpret_cast<const cupmScalar_t *>(s); }
585: PETSC_NODISCARD PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND static constexpr cupmScalar_t *cupmScalarPtrCast(PetscScalar *s) noexcept { return reinterpret_cast<cupmScalar_t *>(s); }
587: PETSC_NODISCARD PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND static constexpr const cupmReal_t *cupmRealPtrCast(const PetscReal *s) noexcept { return reinterpret_cast<const cupmReal_t *>(s); }
589: PETSC_NODISCARD PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND static constexpr cupmReal_t *cupmRealPtrCast(PetscReal *s) noexcept { return reinterpret_cast<cupmReal_t *>(s); }
591: #if !defined(PETSC_PKG_CUDA_VERSION_GE)
592: #define PETSC_PKG_CUDA_VERSION_GE(...) 0
593: #define CUPM_DEFINED_PETSC_PKG_CUDA_VERSION_GE
594: #endif
596: static PetscErrorCode PetscCUPMGetMemType(const void *data, PetscMemType *type, PetscBool *registered = nullptr, PetscBool *managed = nullptr) noexcept
597: {
598: cupmPointerAttributes_t attr;
599: cupmError_t cerr;
601: PetscFunctionBegin;
602: if (type) PetscAssertPointer(type, 2);
603: if (registered) {
604: PetscAssertPointer(registered, 3);
605: *registered = PETSC_FALSE;
606: }
607: if (managed) {
608: PetscAssertPointer(managed, 4);
609: *managed = PETSC_FALSE;
610: }
611: // Do not check error, instead reset it via GetLastError() since before CUDA 11.0, passing
612: // a host pointer returns cudaErrorInvalidValue
613: cerr = cupmPointerGetAttributes(&attr, data);
614: cerr = cupmGetLastError();
615: if (PETSC_PKG_CUDA_VERSION_GE(11, 0, 0) && (T == DeviceType::CUDA)) PetscCallCUPM(cerr);
616: if (type) *type = PETSC_MEMTYPE_HOST;
617: if (cerr == cupmSuccess) {
618: const auto mtype = interface_type::cupmPointerGetMemType(attr, managed);
620: // Managed memory is device accessible, so treat it as device memory.
621: if (type && (mtype == cupmMemoryTypeDevice || mtype == cupmMemoryTypeManaged)) *type = PETSC_MEMTYPE_CUPM();
622: if (registered && (mtype == cupmMemoryTypeHost)) *registered = PETSC_TRUE;
623: }
624: PetscFunctionReturn(PETSC_SUCCESS);
625: }
626: #if defined(CUPM_DEFINED_PETSC_PKG_CUDA_VERSION_GE)
627: #undef PETSC_PKG_CUDA_VERSION_GE
628: #endif
630: PETSC_NODISCARD static PETSC_CONSTEXPR_14 cupmMemcpyKind_t PetscDeviceCopyModeToCUPMMemcpyKind(PetscDeviceCopyMode mode) noexcept
631: {
632: switch (mode) {
633: case PETSC_DEVICE_COPY_HTOH:
634: return cupmMemcpyHostToHost;
635: case PETSC_DEVICE_COPY_HTOD:
636: return cupmMemcpyHostToDevice;
637: case PETSC_DEVICE_COPY_DTOD:
638: return cupmMemcpyDeviceToDevice;
639: case PETSC_DEVICE_COPY_DTOH:
640: return cupmMemcpyDeviceToHost;
641: case PETSC_DEVICE_COPY_AUTO:
642: return cupmMemcpyDefault;
643: }
644: PetscUnreachable();
645: return cupmMemcpyDefault;
646: }
648: // these change what the arguments mean, so need to namespace these
649: template <typename M>
650: static PetscErrorCode PetscCUPMMallocAsync(M **ptr, std::size_t n, cupmStream_t stream = nullptr) noexcept
651: {
652: static_assert(!std::is_void<M>::value, "");
654: PetscFunctionBegin;
655: PetscAssertPointer(ptr, 1);
656: *ptr = nullptr;
657: if (n) {
658: const auto bytes = n * sizeof(M);
659: // https://developer.nvidia.com/blog/using-cuda-stream-ordered-memory-allocator-part-2/
660: //
661: // TLD;DR: cudaMallocAsync() does not work with NVIDIA GPUDirect which OPENMPI uses to
662: // underpin its cuda-aware MPI implementation, so we cannot just async allocate
663: // blindly...
664: if (stream) {
665: PetscCallCUPM(cupmMallocAsync(reinterpret_cast<void **>(ptr), bytes, stream));
666: } else {
667: PetscCallCUPM(cupmMalloc(reinterpret_cast<void **>(ptr), bytes));
668: }
669: }
670: PetscFunctionReturn(PETSC_SUCCESS);
671: }
673: template <typename M>
674: static PetscErrorCode PetscCUPMMalloc(M **ptr, std::size_t n) noexcept
675: {
676: PetscFunctionBegin;
677: PetscCall(PetscCUPMMallocAsync(ptr, n));
678: PetscFunctionReturn(PETSC_SUCCESS);
679: }
681: template <typename M>
682: static PetscErrorCode PetscCUPMMallocHost(M **ptr, std::size_t n, unsigned int flags = cupmHostAllocDefault) noexcept
683: {
684: static_assert(!std::is_void<M>::value, "");
686: PetscFunctionBegin;
687: PetscAssertPointer(ptr, 1);
688: *ptr = nullptr;
689: if (n) PetscCallCUPM(cupmMallocHost(reinterpret_cast<void **>(ptr), n * sizeof(M), flags));
690: PetscFunctionReturn(PETSC_SUCCESS);
691: }
693: template <typename D>
694: static PetscErrorCode PetscCUPMMemcpyAsync(D *dest, const util::type_identity_t<D> *src, std::size_t n, cupmMemcpyKind_t kind, cupmStream_t stream = nullptr, bool use_async = false) noexcept
695: {
696: static_assert(!std::is_void<D>::value, "");
697: const auto size = n * sizeof(D);
699: PetscFunctionBegin;
700: if (PetscUnlikely(!n)) PetscFunctionReturn(PETSC_SUCCESS);
701: // cannot dereference (i.e. cannot call PetscAssertPointer() here)
702: PetscCheck(dest, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to copy to a NULL pointer");
703: PetscCheck(src, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to copy from a NULL pointer");
704: if (kind == cupmMemcpyHostToHost) {
705: // If we are HTOH it is cheaper to check if the stream is idle and do a basic mempcy()
706: // than it is to just call the vendor functions. This assumes of course that the stream
707: // accounts for both memory regions being "idle"
708: if (cupmStreamQuery(stream) == cupmSuccess) {
709: PetscCall(PetscMemcpy(dest, src, size));
710: PetscFunctionReturn(PETSC_SUCCESS);
711: }
712: // need to clear the potential cupmErrorNotReady generated by query above...
713: auto cerr = cupmGetLastError();
715: if (PetscUnlikely(cerr != cupmErrorNotReady)) PetscCallCUPM(cerr);
716: }
717: if (use_async || stream || (kind != cupmMemcpyDeviceToHost)) {
718: PetscCallCUPM(cupmMemcpyAsync(dest, src, size, kind, stream));
719: } else {
720: PetscCallCUPM(cupmMemcpy(dest, src, size, kind));
721: }
722: PetscCall(PetscLogCUPMMemcpyTransfer(kind, size));
723: PetscFunctionReturn(PETSC_SUCCESS);
724: }
726: template <typename D>
727: static PetscErrorCode PetscCUPMMemcpy(D *dest, const util::type_identity_t<D> *src, std::size_t n, cupmMemcpyKind_t kind) noexcept
728: {
729: PetscFunctionBegin;
730: PetscCall(PetscCUPMMemcpyAsync(dest, src, n, kind));
731: PetscFunctionReturn(PETSC_SUCCESS);
732: }
734: template <typename D>
735: static PetscErrorCode PetscCUPMMemcpy2DAsync(D *dest, std::size_t dest_pitch, const util::type_identity_t<D> *src, std::size_t src_pitch, std::size_t width, std::size_t height, cupmMemcpyKind_t kind, cupmStream_t stream = nullptr)
736: {
737: static_assert(!std::is_void<D>::value, "");
738: const auto dest_pitch_bytes = dest_pitch * sizeof(D);
739: const auto src_pitch_bytes = src_pitch * sizeof(D);
740: const auto width_bytes = width * sizeof(D);
741: const auto size = height * width_bytes;
743: PetscFunctionBegin;
744: if (PetscUnlikely(!size)) PetscFunctionReturn(PETSC_SUCCESS);
745: PetscCheck(dest, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to copy to a NULL pointer");
746: PetscCheck(src, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to copy from a NULL pointer");
747: if (stream || (kind != cupmMemcpyDeviceToHost)) {
748: PetscCallCUPM(cupmMemcpy2DAsync(dest, dest_pitch_bytes, src, src_pitch_bytes, width_bytes, height, kind, stream));
749: } else {
750: PetscCallCUPM(cupmMemcpy2D(dest, dest_pitch_bytes, src, src_pitch_bytes, width_bytes, height, kind));
751: }
752: PetscCall(PetscLogCUPMMemcpyTransfer(kind, size));
753: PetscFunctionReturn(PETSC_SUCCESS);
754: }
756: template <typename D>
757: static PetscErrorCode PetscCUPMMemcpy2D(D *dest, std::size_t dest_pitch, const util::type_identity_t<D> *src, std::size_t src_pitch, std::size_t width, std::size_t height, cupmMemcpyKind_t kind)
758: {
759: PetscFunctionBegin;
760: PetscCall(PetscCUPMMemcpy2DAsync(dest, dest_pitch, src, src_pitch, width, height, kind));
761: PetscFunctionReturn(PETSC_SUCCESS);
762: }
764: template <typename M>
765: static PetscErrorCode PetscCUPMMemsetAsync(M *ptr, int value, std::size_t n, cupmStream_t stream = nullptr, bool use_async = false) noexcept
766: {
767: static_assert(!std::is_void<M>::value, "");
769: PetscFunctionBegin;
770: if (PetscLikely(n)) {
771: const auto bytes = n * sizeof(M);
773: PetscCheck(ptr, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to memset a NULL pointer with size %zu != 0", n);
774: if (stream || use_async) {
775: PetscCallCUPM(cupmMemsetAsync(ptr, value, bytes, stream));
776: } else {
777: PetscCallCUPM(cupmMemset(ptr, value, bytes));
778: }
779: }
780: PetscFunctionReturn(PETSC_SUCCESS);
781: }
783: template <typename M>
784: static PetscErrorCode PetscCUPMMemset(M *ptr, int value, std::size_t n) noexcept
785: {
786: PetscFunctionBegin;
787: PetscCall(PetscCUPMMemsetAsync(ptr, value, n));
788: PetscFunctionReturn(PETSC_SUCCESS);
789: }
791: template <typename D>
792: static PetscErrorCode PetscCUPMMemset2DAsync(D *ptr, std::size_t pitch, int value, std::size_t width, std::size_t height, cupmStream_t stream = nullptr)
793: {
794: static_assert(!std::is_void<D>::value, "");
795: const auto pitch_bytes = pitch * sizeof(D);
796: const auto width_bytes = width * sizeof(D);
797: const auto size = width_bytes * height;
799: PetscFunctionBegin;
800: if (PetscUnlikely(!size)) PetscFunctionReturn(PETSC_SUCCESS);
801: PetscAssert(ptr, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to memset a NULL pointer with size %zu != 0", size);
802: if (stream) {
803: PetscCallCUPM(cupmMemset2DAsync(ptr, pitch_bytes, value, width_bytes, height, stream));
804: } else {
805: PetscCallCUPM(cupmMemset2D(ptr, pitch_bytes, value, width_bytes, height));
806: }
807: PetscFunctionReturn(PETSC_SUCCESS);
808: }
810: // these we can transparently wrap, no need to namespace it to Petsc
811: template <typename M>
812: PETSC_NODISCARD static cupmError_t cupmFreeAsync(M &ptr, cupmStream_t stream = nullptr) noexcept
813: {
814: static_assert(std::is_pointer<util::decay_t<M>>::value, "");
815: static_assert(!std::is_const<M>::value, "");
817: if (ptr) {
818: auto cerr = interface_type::cupmFreeAsync(std::forward<M>(ptr), stream);
820: ptr = nullptr;
821: if (PetscUnlikely(cerr != cupmSuccess)) return cerr;
822: }
823: return cupmSuccess;
824: }
826: PETSC_NODISCARD static cupmError_t cupmFreeAsync(std::nullptr_t ptr, cupmStream_t stream = nullptr) { return interface_type::cupmFreeAsync(ptr, stream); }
828: template <typename M>
829: PETSC_NODISCARD static cupmError_t cupmFree(M &ptr) noexcept
830: {
831: return cupmFreeAsync(ptr);
832: }
834: PETSC_NODISCARD static cupmError_t cupmFree(std::nullptr_t ptr) { return cupmFreeAsync(ptr); }
836: template <typename M>
837: PETSC_NODISCARD static cupmError_t cupmFreeHost(M &ptr) noexcept
838: {
839: static_assert(std::is_pointer<util::decay_t<M>>::value, "");
840: const auto cerr = interface_type::cupmFreeHost(std::forward<M>(ptr));
841: ptr = nullptr;
842: return cerr;
843: }
845: PETSC_NODISCARD static cupmError_t cupmFreeHost(std::nullptr_t ptr) { return interface_type::cupmFreeHost(ptr); }
847: // specific wrapper for device launch function, as the real function is a C routine and
848: // doesn't have variable arguments. The actual mechanics of this are a bit complicated but
849: // boils down to the fact that ultimately we pass a
850: //
851: // void *args[] = {(void*)&kernel_args...};
852: //
853: // to the kernel launcher. Since we pass void* this means implicit conversion does **not**
854: // happen to the kernel arguments so we must do it ourselves here. This function does this in
855: // 3 stages:
856: // 1. Enumerate the kernel arguments (cupmLaunchKernel)
857: // 2. Deduce the signature of func() and static_cast the kernel arguments to the type
858: // expected by func() using the enumeration above (deduceKernelCall)
859: // 3. Form the void* array with the converted arguments and call cuda/hipLaunchKernel with
860: // it. (interface_type::cupmLaunchKernel)
861: template <typename F, typename... Args>
862: PETSC_NODISCARD static cupmError_t cupmLaunchKernel(F &&func, cupmDim3 gridDim, cupmDim3 blockDim, std::size_t sharedMem, cupmStream_t stream, Args &&...kernelArgs) noexcept
863: {
864: return deduceKernelCall(util::index_sequence_for<Args...>{}, std::forward<F>(func), std::move(gridDim), std::move(blockDim), std::move(sharedMem), std::move(stream), std::forward<Args>(kernelArgs)...);
865: }
867: template <std::size_t block_size = 256, std::size_t warp_size = 32, typename F, typename... Args>
868: static PetscErrorCode PetscCUPMLaunchKernel1D(std::size_t n, std::size_t sharedMem, cupmStream_t stream, F &&func, Args &&...kernelArgs) noexcept
869: {
870: static_assert(block_size > 0, "");
871: static_assert(warp_size > 0, "");
872: // want block_size to be a multiple of the warp_size
873: static_assert(block_size % warp_size == 0, "");
874: const auto nthread = std::min(n, block_size);
875: const auto nblock = (n + block_size - 1) / block_size;
877: PetscFunctionBegin;
878: // if n = 0 then nthread = 0, which is not allowed. rather than letting the user try to
879: // decipher cryptic 'cuda/hipErrorLaunchFailure' we explicitly check for zero here
880: PetscAssert(nthread, PETSC_COMM_SELF, PETSC_ERR_ARG_WRONG, "Trying to launch kernel with grid/block size 0");
881: PetscCallCUPM(cupmLaunchKernel(std::forward<F>(func), (unsigned int)nblock, (unsigned int)nthread, sharedMem, stream, std::forward<Args>(kernelArgs)...));
882: PetscFunctionReturn(PETSC_SUCCESS);
883: }
885: private:
886: template <typename S, typename D, typename = void>
887: struct is_static_castable : std::false_type { };
889: template <typename S, typename D>
890: struct is_static_castable<S, D, util::void_t<decltype(static_cast<D>(std::declval<S>()))>> : std::true_type { };
892: template <typename D, typename S>
893: static constexpr util::enable_if_t<is_static_castable<S, D>::value, D> cast_to(S &&src) noexcept
894: {
895: return static_cast<D>(std::forward<S>(src));
896: }
898: template <typename D, typename S>
899: static constexpr util::enable_if_t<!is_static_castable<S, D>::value, D> cast_to(S &&src) noexcept
900: {
901: return const_cast<D>(std::forward<S>(src));
902: }
904: template <typename F, typename... Args, std::size_t... Idx>
905: PETSC_NODISCARD static cupmError_t deduceKernelCall(util::index_sequence<Idx...>, F &&func, cupmDim3 gridDim, cupmDim3 blockDim, std::size_t sharedMem, cupmStream_t stream, Args &&...kernelArgs) noexcept
906: {
907: // clang-format off
908: return interface_type::cupmLaunchKernel(
909: std::forward<F>(func),
910: std::move(gridDim), std::move(blockDim), std::move(sharedMem), std::move(stream),
911: // can't static_cast() here since the function argument type may be cv-qualified, in
912: // which case we would need to const_cast(). But you can only const_cast() indirect types
913: // (pointers, references). So we need a SFINAE monster that is a static_cast() if
914: // possible, and a const_cast() if not. We could just use a C-style cast which *would*
915: // work here since it tries the following and uses the first one that succeeds:
916: //
917: // 1. const_cast()
918: // 2. static_cast()
919: // 3. static_cast() then const_cast()
920: // 4. reinterpret_cast()...
921: //
922: // the issue however is the final reinterpret_cast(). We absolutely cannot get there
923: // because doing so would silently hide a ton of bugs, for example casting a PetscScalar
924: // * to double * in complex builds, a PetscInt * to int * in 64idx builds, etc.
925: cast_to<typename util::func_traits<F>::template arg<Idx>::type>(std::forward<Args>(kernelArgs))...
926: );
927: // clang-format on
928: }
930: static PetscErrorCode PetscLogCUPMMemcpyTransfer(cupmMemcpyKind_t kind, std::size_t size) noexcept
931: {
932: PetscFunctionBegin;
933: // only the explicit HTOD or DTOH are handled, since we either don't log the other cases
934: // (yet) or don't know the direction
935: if (kind == cupmMemcpyDeviceToHost) PetscCall(PetscLogGpuToCpu(static_cast<PetscLogDouble>(size)));
936: else if (kind == cupmMemcpyHostToDevice) PetscCall(PetscLogCpuToGpu(static_cast<PetscLogDouble>(size)));
937: else (void)size;
938: PetscFunctionReturn(PETSC_SUCCESS);
939: }
940: };
942: #undef PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND
944: #define PETSC_CUPM_INHERIT_INTERFACE_TYPEDEFS_USING(T) \
945: PETSC_CUPM_IMPL_CLASS_HEADER(T); \
946: using cupmReal_t = typename ::Petsc::device::cupm::impl::Interface<T>::cupmReal_t; \
947: using cupmScalar_t = typename ::Petsc::device::cupm::impl::Interface<T>::cupmScalar_t; \
948: using ::Petsc::device::cupm::impl::Interface<T>::cupmScalarCast; \
949: using ::Petsc::device::cupm::impl::Interface<T>::cupmScalarPtrCast; \
950: using ::Petsc::device::cupm::impl::Interface<T>::cupmRealPtrCast; \
951: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMGetMemType; \
952: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemset; \
953: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemsetAsync; \
954: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMalloc; \
955: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMallocAsync; \
956: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMallocHost; \
957: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemcpy; \
958: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemcpyAsync; \
959: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemcpy2D; \
960: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemcpy2DAsync; \
961: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemset2DAsync; \
962: using ::Petsc::device::cupm::impl::Interface<T>::cupmFree; \
963: using ::Petsc::device::cupm::impl::Interface<T>::cupmFreeAsync; \
964: using ::Petsc::device::cupm::impl::Interface<T>::cupmFreeHost; \
965: using ::Petsc::device::cupm::impl::Interface<T>::cupmLaunchKernel; \
966: using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMLaunchKernel1D; \
967: using ::Petsc::device::cupm::impl::Interface<T>::PetscDeviceCopyModeToCUPMMemcpyKind
969: #if PetscDefined(HAVE_CUDA)
970: extern template struct PETSC_SINGLE_LIBRARY_VISIBILITY_INTERNAL Interface<DeviceType::CUDA>;
971: #endif
973: #if PetscDefined(HAVE_HIP)
974: extern template struct PETSC_SINGLE_LIBRARY_VISIBILITY_INTERNAL Interface<DeviceType::HIP>;
975: #endif
977: } // namespace impl
979: } // namespace cupm
981: } // namespace device
983: } // namespace Petsc