Actual source code: cupminterface.hpp

  1: #pragma once

  3: #include <petscdevice_cupm.h>

  5: #include <petsc/private/cpputil.hpp>
  6: #include <petsc/private/petscadvancedmacros.h>

  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<int,char,bool>());
 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 <typename... T>
 66: // static constexpr auto cupmMalloc(T&&... args) *noexcept and trailing return type deduction*
 67: // {
 68: //   return cudaMalloc(std::forward<T>(args)...);
 69: // }
 70: //
 71: // PETSC_CUPM_ALIAS_FUNCTION(cupmMalloc, hipMalloc) ->
 72: // template <typename... T>
 73: // static constexpr auto cupmMalloc(T&&... args) *noexcept and trailing return type deduction*
 74: // {
 75: //   return hipMalloc(std::forward<T>(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 <typename... T, typename Tend>
 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<T>(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 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 cupmComplex_t           = util::conditional_t<PetscDefined(USE_REAL_SINGLE), cuComplex, cuDoubleComplex>;
150:   using cupmPointerAttributes_t = cudaPointerAttributes;
151:   using cupmMemoryType_t        = enum cudaMemoryType;
152:   using cupmDim3                = dim3;
153:   using cupmHostFn_t            = cudaHostFn_t;
154:   #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
155:   using cupmMemPool_t   = cudaMemPool_t;
156:   using cupmMemPoolAttr = cudaMemPoolAttr;
157:   #else
158:   using cupmMemPool_t   = void *;
159:   using cupmMemPoolAttr = unsigned int;
160:   #endif

162:   // values
163:   static const auto cupmSuccess                 = cudaSuccess;
164:   static const auto cupmErrorNotReady           = cudaErrorNotReady;
165:   static const auto cupmErrorDeviceAlreadyInUse = cudaErrorDeviceAlreadyInUse;
166:   static const auto cupmErrorSetOnActiveProcess = cudaErrorSetOnActiveProcess;
167:   static const auto cupmErrorStubLibrary =
168:   #if PETSC_PKG_CUDA_VERSION_GE(11, 1, 0)
169:     cudaErrorStubLibrary;
170:   #else
171:     cudaErrorInsufficientDriver;
172:   #endif

174:   static const auto cupmErrorNoDevice          = cudaErrorNoDevice;
175:   static const auto cupmStreamDefault          = cudaStreamDefault;
176:   static const auto cupmStreamNonBlocking      = cudaStreamNonBlocking;
177:   static const auto cupmDeviceMapHost          = cudaDeviceMapHost;
178:   static const auto cupmMemcpyHostToDevice     = cudaMemcpyHostToDevice;
179:   static const auto cupmMemcpyDeviceToHost     = cudaMemcpyDeviceToHost;
180:   static const auto cupmMemcpyDeviceToDevice   = cudaMemcpyDeviceToDevice;
181:   static const auto cupmMemcpyHostToHost       = cudaMemcpyHostToHost;
182:   static const auto cupmMemcpyDefault          = cudaMemcpyDefault;
183:   static const auto cupmMemoryTypeHost         = cudaMemoryTypeHost;
184:   static const auto cupmMemoryTypeDevice       = cudaMemoryTypeDevice;
185:   static const auto cupmMemoryTypeManaged      = cudaMemoryTypeManaged;
186:   static const auto cupmEventDisableTiming     = cudaEventDisableTiming;
187:   static const auto cupmHostAllocDefault       = cudaHostAllocDefault;
188:   static const auto cupmHostAllocWriteCombined = cudaHostAllocWriteCombined;
189:   static const auto cupmMemPoolAttrReleaseThreshold =
190:   #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
191:     cudaMemPoolAttrReleaseThreshold;
192:   #else
193:     cupmMemPoolAttr{0};
194:   #endif

196:   // error functions
197:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetErrorName, cudaGetErrorName)
198:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetErrorString, cudaGetErrorString)
199:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetLastError, cudaGetLastError)

201:   // device management
202:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceCount, cudaGetDeviceCount)
203:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceProperties, cudaGetDeviceProperties)
204:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetDevice, cudaGetDevice)
205:   PETSC_CUPM_ALIAS_FUNCTION(cupmSetDevice, cudaSetDevice)
206:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceFlags, cudaGetDeviceFlags)
207:   PETSC_CUPM_ALIAS_FUNCTION(cupmSetDeviceFlags, cudaSetDeviceFlags)
208:   PETSC_CUPM_ALIAS_FUNCTION(cupmPointerGetAttributes, cudaPointerGetAttributes)
209:   #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
210:   PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceGetMemPool, cudaDeviceGetMemPool)
211:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemPoolSetAttribute, cudaMemPoolSetAttribute)
212:   #else
213:   PETSC_NODISCARD static cupmError_t cupmDeviceGetMemPool(cupmMemPool_t *pool, int) noexcept
214:   {
215:     *pool = nullptr;
216:     return cupmSuccess;
217:   }

219:   PETSC_NODISCARD static cupmError_t cupmMemPoolSetAttribute(cupmMemPool_t, cupmMemPoolAttr, void *) noexcept { return cupmSuccess; }
220:   #endif
221:   // CUDA has no cudaInit() to match hipInit()
222:   PETSC_NODISCARD static cupmError_t cupmInit(unsigned int) noexcept { return cudaFree(nullptr); }

224:   // stream management
225:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventCreate, cudaEventCreate)
226:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventCreateWithFlags, cudaEventCreateWithFlags)
227:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventDestroy, cudaEventDestroy)
228:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventRecord, cudaEventRecord)
229:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventSynchronize, cudaEventSynchronize)
230:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventElapsedTime, cudaEventElapsedTime)
231:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventQuery, cudaEventQuery)
232:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamCreate, cudaStreamCreate)
233:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamCreateWithFlags, cudaStreamCreateWithFlags)
234:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamGetFlags, cudaStreamGetFlags)
235:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamDestroy, cudaStreamDestroy)
236:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamWaitEvent, cudaStreamWaitEvent)
237:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamQuery, cudaStreamQuery)
238:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamSynchronize, cudaStreamSynchronize)
239:   PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceSynchronize, cudaDeviceSynchronize)
240:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetSymbolAddress, cudaGetSymbolAddress)

242:   // memory management
243:   PETSC_CUPM_ALIAS_FUNCTION(cupmFree, cudaFree)
244:   PETSC_CUPM_ALIAS_FUNCTION(cupmMalloc, cudaMalloc)
245:   #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
246:   PETSC_CUPM_ALIAS_FUNCTION(cupmFreeAsync, cudaFreeAsync)
247:   PETSC_CUPM_ALIAS_FUNCTION(cupmMallocAsync, cudaMallocAsync)
248:   #else
249:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmFreeAsync, cudaFree, 1)
250:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMallocAsync, cudaMalloc, 1)
251:   #endif
252:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy, cudaMemcpy)
253:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpyAsync, cudaMemcpyAsync)
254:   PETSC_CUPM_ALIAS_FUNCTION(cupmMallocHost, cudaMallocHost)
255:   PETSC_CUPM_ALIAS_FUNCTION(cupmFreeHost, cudaFreeHost)
256:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemset, cudaMemset)
257:   #if PETSC_PKG_CUDA_VERSION_GE(11, 2, 0)
258:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemsetAsync, cudaMemsetAsync)
259:   #else
260:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMemsetAsync, cudaMemset, 1)
261:   #endif
262:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy2D, cudaMemcpy2D)
263:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMemcpy2DAsync, cudaMemcpy2DAsync, 1)
264:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemset2D, cudaMemset2D)
265:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMemset2DAsync, cudaMemset2DAsync, 1)

267:   // launch control
268:   PETSC_CUPM_ALIAS_FUNCTION(cupmLaunchHostFunc, cudaLaunchHostFunc)
269:   template <typename FunctionT, typename... KernelArgsT>
270:   PETSC_NODISCARD static cudaError_t cupmLaunchKernel(FunctionT &&func, dim3 gridDim, dim3 blockDim, std::size_t sharedMem, cudaStream_t stream, KernelArgsT &&...kernelArgs) noexcept
271:   {
272:     static_assert(!std::is_pointer<FunctionT>::value, "kernel function must not be passed by pointer");
273:     void *args[] = {(void *)std::addressof(kernelArgs)...};

275:     return cudaLaunchKernel<util::remove_reference_t<FunctionT>>(std::addressof(func), std::move(gridDim), std::move(blockDim), args, sharedMem, std::move(stream));
276:   }
277: };
278: #endif // PetscDefined(HAVE_CUDA)

280: #if PetscDefined(HAVE_HIP)
281: template <>
282: struct InterfaceImpl<DeviceType::HIP> : InterfaceBase<DeviceType::HIP> {
283:   PETSC_CUPM_BASE_CLASS_HEADER(DeviceType::HIP);

285:   // typedefs
286:   using cupmError_t             = hipError_t;
287:   using cupmEvent_t             = hipEvent_t;
288:   using cupmStream_t            = hipStream_t;
289:   using cupmDeviceProp_t        = hipDeviceProp_t;
290:   using cupmMemcpyKind_t        = hipMemcpyKind;
291:   using cupmComplex_t           = util::conditional_t<PetscDefined(USE_REAL_SINGLE), hipComplex, hipDoubleComplex>;
292:   using cupmPointerAttributes_t = hipPointerAttribute_t;
293:   using cupmMemoryType_t        = enum hipMemoryType;
294:   using cupmDim3                = dim3;
295:   #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
296:   using cupmHostFn_t    = hipHostFn_t;
297:   using cupmMemPool_t   = hipMemPool_t;
298:   using cupmMemPoolAttr = hipMemPoolAttr;
299:   #else
300:   using cupmHostFn_t    = void (*)(void *);
301:   using cupmMemPool_t   = void *;
302:   using cupmMemPoolAttr = unsigned int;
303:   #endif

305:   // values
306:   static const auto cupmSuccess       = hipSuccess;
307:   static const auto cupmErrorNotReady = hipErrorNotReady;
308:   // see https://github.com/ROCm-Developer-Tools/HIP/blob/develop/bin/hipify-perl
309:   static const auto cupmErrorDeviceAlreadyInUse = hipErrorContextAlreadyInUse;
310:   static const auto cupmErrorSetOnActiveProcess = hipErrorSetOnActiveProcess;
311:   // as of HIP v4.2 cudaErrorStubLibrary has no HIP equivalent
312:   static const auto cupmErrorStubLibrary     = hipErrorInsufficientDriver;
313:   static const auto cupmErrorNoDevice        = hipErrorNoDevice;
314:   static const auto cupmStreamDefault        = hipStreamDefault;
315:   static const auto cupmStreamNonBlocking    = hipStreamNonBlocking;
316:   static const auto cupmDeviceMapHost        = hipDeviceMapHost;
317:   static const auto cupmMemcpyHostToDevice   = hipMemcpyHostToDevice;
318:   static const auto cupmMemcpyDeviceToHost   = hipMemcpyDeviceToHost;
319:   static const auto cupmMemcpyDeviceToDevice = hipMemcpyDeviceToDevice;
320:   static const auto cupmMemcpyHostToHost     = hipMemcpyHostToHost;
321:   static const auto cupmMemcpyDefault        = hipMemcpyDefault;
322:   static const auto cupmMemoryTypeHost       = hipMemoryTypeHost;
323:   static const auto cupmMemoryTypeDevice     = hipMemoryTypeDevice;
324:   // see
325:   // https://github.com/ROCm-Developer-Tools/HIP/blob/develop/include/hip/hip_runtime_api.h#L156
326:   static const auto cupmMemoryTypeManaged      = hipMemoryTypeUnified;
327:   static const auto cupmEventDisableTiming     = hipEventDisableTiming;
328:   static const auto cupmHostAllocDefault       = hipHostMallocDefault;
329:   static const auto cupmHostAllocWriteCombined = hipHostMallocWriteCombined;
330:   static const auto cupmMemPoolAttrReleaseThreshold =
331:   #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
332:     hipMemPoolAttrReleaseThreshold;
333:   #else
334:     cupmMemPoolAttr{0};
335:   #endif

337:   // error functions
338:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetErrorName, hipGetErrorName)
339:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetErrorString, hipGetErrorString)
340:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetLastError, hipGetLastError)

342:   // device management
343:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceCount, hipGetDeviceCount)
344:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceProperties, hipGetDeviceProperties)
345:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetDevice, hipGetDevice)
346:   PETSC_CUPM_ALIAS_FUNCTION(cupmSetDevice, hipSetDevice)
347:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetDeviceFlags, hipGetDeviceFlags)
348:   PETSC_CUPM_ALIAS_FUNCTION(cupmSetDeviceFlags, hipSetDeviceFlags)
349:   PETSC_CUPM_ALIAS_FUNCTION(cupmPointerGetAttributes, hipPointerGetAttributes)
350:   #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
351:   PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceGetMemPool, hipDeviceGetMemPool)
352:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemPoolSetAttribute, hipMemPoolSetAttribute)
353:   #else
354:   PETSC_NODISCARD static cupmError_t cupmDeviceGetMemPool(cupmMemPool_t *pool, int) noexcept
355:   {
356:     *pool = nullptr;
357:     return cupmSuccess;
358:   }

360:   PETSC_NODISCARD static cupmError_t cupmMemPoolSetAttribute(cupmMemPool_t, cupmMemPoolAttr, void *) noexcept { return cupmSuccess; }
361:   #endif
362:   PETSC_CUPM_ALIAS_FUNCTION(cupmInit, hipInit)

364:   // stream management
365:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventCreate, hipEventCreate)
366:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventCreateWithFlags, hipEventCreateWithFlags)
367:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventDestroy, hipEventDestroy)
368:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventRecord, hipEventRecord)
369:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventSynchronize, hipEventSynchronize)
370:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventElapsedTime, hipEventElapsedTime)
371:   PETSC_CUPM_ALIAS_FUNCTION(cupmEventQuery, hipEventQuery)
372:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamCreate, hipStreamCreate)
373:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamCreateWithFlags, hipStreamCreateWithFlags)
374:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamGetFlags, hipStreamGetFlags)
375:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamDestroy, hipStreamDestroy)
376:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamWaitEvent, hipStreamWaitEvent)
377:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamQuery, hipStreamQuery)
378:   PETSC_CUPM_ALIAS_FUNCTION(cupmStreamSynchronize, hipStreamSynchronize)
379:   PETSC_CUPM_ALIAS_FUNCTION(cupmDeviceSynchronize, hipDeviceSynchronize)
380:   PETSC_CUPM_ALIAS_FUNCTION(cupmGetSymbolAddress, hipGetSymbolAddress)

382:   // memory management
383:   PETSC_CUPM_ALIAS_FUNCTION(cupmFree, hipFree)
384:   PETSC_CUPM_ALIAS_FUNCTION(cupmMalloc, hipMalloc)
385:   #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
386:   PETSC_CUPM_ALIAS_FUNCTION(cupmMallocAsync, hipMallocAsync)
387:   PETSC_CUPM_ALIAS_FUNCTION(cupmFreeAsync, hipFreeAsync)
388:   #else
389:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMallocAsync, hipMalloc, 1)
390:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmFreeAsync, hipFree, 1)
391:   #endif
392:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy, hipMemcpy)
393:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpyAsync, hipMemcpyAsync)
394:   // hipMallocHost is deprecated
395:   PETSC_CUPM_ALIAS_FUNCTION(cupmMallocHost, hipHostMalloc)
396:   // hipFreeHost is deprecated
397:   PETSC_CUPM_ALIAS_FUNCTION(cupmFreeHost, hipHostFree)
398:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemset, hipMemset)
399:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemsetAsync, hipMemsetAsync)
400:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemcpy2D, hipMemcpy2D)
401:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMemcpy2DAsync, hipMemcpy2DAsync, 1)
402:   PETSC_CUPM_ALIAS_FUNCTION(cupmMemset2D, hipMemset2D)
403:   PETSC_CUPM_ALIAS_FUNCTION_GOBBLE(cupmMemset2DAsync, hipMemset2DAsync, 1)

405:   // launch control
406:   // HIP appears to only have hipLaunchHostFunc from 5.2.0 onwards
407:   // https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/doc/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md#7-execution-control=
408:   #if PETSC_PKG_HIP_VERSION_GE(5, 2, 0)
409:   PETSC_CUPM_ALIAS_FUNCTION(cupmLaunchHostFunc, hipLaunchHostFunc)
410:   #else
411:   PETSC_NODISCARD static hipError_t cupmLaunchHostFunc(hipStream_t stream, cupmHostFn_t fn, void *ctx) noexcept
412:   {
413:     // the only correct way to spoof this function is to do it synchronously...
414:     auto herr = hipStreamSynchronize(stream);
415:     if (PetscUnlikely(herr != hipSuccess)) return herr;
416:     fn(ctx);
417:     return herr;
418:   }
419:   #endif

421:   template <typename FunctionT, typename... KernelArgsT>
422:   PETSC_NODISCARD static hipError_t cupmLaunchKernel(FunctionT &&func, dim3 gridDim, dim3 blockDim, std::size_t sharedMem, hipStream_t stream, KernelArgsT &&...kernelArgs) noexcept
423:   {
424:     void *args[] = {(void *)std::addressof(kernelArgs)...};

426:     return hipLaunchKernel((void *)func, std::move(gridDim), std::move(blockDim), args, sharedMem, std::move(stream));
427:   }
428: };
429: #endif // PetscDefined(HAVE_HIP)

431: // shorthand for bringing all of the typedefs from the base Interface class into your own,
432: // it's annoying that c++ doesn't have a way to do this automatically
433: #define PETSC_CUPM_IMPL_CLASS_HEADER(T) \
434:   PETSC_CUPM_BASE_CLASS_HEADER(T); \
435:   /* types */ \
436:   using cupmComplex_t           = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmComplex_t; \
437:   using cupmError_t             = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmError_t; \
438:   using cupmEvent_t             = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEvent_t; \
439:   using cupmStream_t            = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStream_t; \
440:   using cupmDeviceProp_t        = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceProp_t; \
441:   using cupmMemcpyKind_t        = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyKind_t; \
442:   using cupmPointerAttributes_t = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmPointerAttributes_t; \
443:   using cupmMemoryType_t        = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemoryType_t; \
444:   using cupmDim3                = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDim3; \
445:   using cupmMemPool_t           = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemPool_t; \
446:   using cupmMemPoolAttr         = typename ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemPoolAttr; \
447:   /* variables */ \
448:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmSuccess; \
449:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorNotReady; \
450:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorDeviceAlreadyInUse; \
451:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorSetOnActiveProcess; \
452:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorStubLibrary; \
453:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmErrorNoDevice; \
454:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamDefault; \
455:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamNonBlocking; \
456:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceMapHost; \
457:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyHostToDevice; \
458:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyDeviceToHost; \
459:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyDeviceToDevice; \
460:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyHostToHost; \
461:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyDefault; \
462:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemoryTypeHost; \
463:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemoryTypeDevice; \
464:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemoryTypeManaged; \
465:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventDisableTiming; \
466:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmHostAllocDefault; \
467:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmHostAllocWriteCombined; \
468:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemPoolAttrReleaseThreshold; \
469:   /* functions */ \
470:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetErrorName; \
471:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetErrorString; \
472:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetLastError; \
473:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetDeviceCount; \
474:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetDeviceProperties; \
475:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetDevice; \
476:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmSetDevice; \
477:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetDeviceFlags; \
478:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmSetDeviceFlags; \
479:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmPointerGetAttributes; \
480:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceGetMemPool; \
481:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemPoolSetAttribute; \
482:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmInit; \
483:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventCreate; \
484:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventCreateWithFlags; \
485:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventDestroy; \
486:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventRecord; \
487:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventSynchronize; \
488:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventElapsedTime; \
489:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmEventQuery; \
490:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamCreate; \
491:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamCreateWithFlags; \
492:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamGetFlags; \
493:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamDestroy; \
494:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamWaitEvent; \
495:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamQuery; \
496:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmStreamSynchronize; \
497:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmDeviceSynchronize; \
498:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmGetSymbolAddress; \
499:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMalloc; \
500:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMallocAsync; \
501:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpy; \
502:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpyAsync; \
503:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMallocHost; \
504:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemset; \
505:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemsetAsync; \
506:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpy2D; \
507:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemcpy2DAsync; \
508:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemset2D; \
509:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmMemset2DAsync; \
510:   using ::Petsc::device::cupm::impl::InterfaceImpl<T>::cupmLaunchHostFunc

512: #if PetscHasAttribute(always_inline)
513:   // https://gcc.gnu.org/bugzilla//show_bug.cgi?id=109464
514:   #define PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND __attribute__((always_inline))
515: #else
516:   #define PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND
517: #endif

519: // The actual interface class
520: template <DeviceType T>
521: struct Interface : InterfaceImpl<T> {
522: private:
523:   using interface_type = InterfaceImpl<T>;

525: public:
526:   PETSC_CUPM_IMPL_CLASS_HEADER(T);

528:   using cupmReal_t   = util::conditional_t<PetscDefined(USE_REAL_SINGLE), float, double>;
529:   using cupmScalar_t = util::conditional_t<PetscDefined(USE_COMPLEX), cupmComplex_t, cupmReal_t>;

531:   PETSC_NODISCARD PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND static constexpr cupmScalar_t cupmScalarCast(PetscScalar s) noexcept
532:   {
533: #if PetscDefined(USE_COMPLEX)
534:     return cupmComplex_t{PetscRealPart(s), PetscImaginaryPart(s)};
535: #else
536:     return static_cast<cupmScalar_t>(s);
537: #endif
538:   }

540:   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); }

542:   PETSC_NODISCARD PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND static constexpr cupmScalar_t *cupmScalarPtrCast(PetscScalar *s) noexcept { return reinterpret_cast<cupmScalar_t *>(s); }

544:   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); }

546:   PETSC_NODISCARD PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND static constexpr cupmReal_t *cupmRealPtrCast(PetscReal *s) noexcept { return reinterpret_cast<cupmReal_t *>(s); }

548: #if !defined(PETSC_PKG_CUDA_VERSION_GE)
549:   #define PETSC_PKG_CUDA_VERSION_GE(...) 0
550:   #define CUPM_DEFINED_PETSC_PKG_CUDA_VERSION_GE
551: #endif
552:   static PetscErrorCode PetscCUPMGetMemType(const void *data, PetscMemType *type, PetscBool *registered = nullptr, PetscBool *managed = nullptr) noexcept
553:   {
554:     cupmPointerAttributes_t attr;
555:     cupmError_t             cerr;

557:     PetscFunctionBegin;
558:     if (type) PetscAssertPointer(type, 2);
559:     if (registered) {
560:       PetscAssertPointer(registered, 3);
561:       *registered = PETSC_FALSE;
562:     }
563:     if (managed) {
564:       PetscAssertPointer(managed, 4);
565:       *managed = PETSC_FALSE;
566:     }
567:     // Do not check error, instead reset it via GetLastError() since before CUDA 11.0, passing
568:     // a host pointer returns cudaErrorInvalidValue
569:     cerr = cupmPointerGetAttributes(&attr, data);
570:     cerr = cupmGetLastError();
571:     // HIP seems to always have used memoryType though
572: #if (defined(CUDART_VERSION) && (CUDART_VERSION < 10000)) || defined(__HIP_PLATFORM_HCC__)
573:     const auto mtype = attr.memoryType;
574:     if (managed) *managed = static_cast<PetscBool>((cerr == cupmSuccess) && attr.isManaged);
575: #else
576:     if (PETSC_PKG_CUDA_VERSION_GE(11, 0, 0) && (T == DeviceType::CUDA)) PetscCallCUPM(cerr);
577:     const auto mtype = attr.type;
578:     if (managed) *managed = static_cast<PetscBool>(mtype == cupmMemoryTypeManaged);
579: #endif // CUDART_VERSION && CUDART_VERSION < 10000 || __HIP_PLATFORM_HCC__
580:     if (type) *type = ((cerr == cupmSuccess) && (mtype == cupmMemoryTypeDevice)) ? PETSC_MEMTYPE_CUPM() : PETSC_MEMTYPE_HOST;
581:     if (registered && (cerr == cupmSuccess) && (mtype == cupmMemoryTypeHost)) *registered = PETSC_TRUE;
582:     PetscFunctionReturn(PETSC_SUCCESS);
583:   }
584: #if defined(CUPM_DEFINED_PETSC_PKG_CUDA_VERSION_GE)
585:   #undef PETSC_PKG_CUDA_VERSION_GE
586: #endif

588:   PETSC_NODISCARD static PETSC_CONSTEXPR_14 cupmMemcpyKind_t PetscDeviceCopyModeToCUPMMemcpyKind(PetscDeviceCopyMode mode) noexcept
589:   {
590:     switch (mode) {
591:     case PETSC_DEVICE_COPY_HTOH:
592:       return cupmMemcpyHostToHost;
593:     case PETSC_DEVICE_COPY_HTOD:
594:       return cupmMemcpyHostToDevice;
595:     case PETSC_DEVICE_COPY_DTOD:
596:       return cupmMemcpyDeviceToDevice;
597:     case PETSC_DEVICE_COPY_DTOH:
598:       return cupmMemcpyDeviceToHost;
599:     case PETSC_DEVICE_COPY_AUTO:
600:       return cupmMemcpyDefault;
601:     }
602:     PetscUnreachable();
603:     return cupmMemcpyDefault;
604:   }

606:   // these change what the arguments mean, so need to namespace these
607:   template <typename M>
608:   static PetscErrorCode PetscCUPMMallocAsync(M **ptr, std::size_t n, cupmStream_t stream = nullptr) noexcept
609:   {
610:     static_assert(!std::is_void<M>::value, "");

612:     PetscFunctionBegin;
613:     PetscAssertPointer(ptr, 1);
614:     *ptr = nullptr;
615:     if (n) {
616:       const auto bytes = n * sizeof(M);
617:       // https://developer.nvidia.com/blog/using-cuda-stream-ordered-memory-allocator-part-2/
618:       //
619:       // TLD;DR: cudaMallocAsync() does not work with NVIDIA GPUDirect which OPENMPI uses to
620:       // underpin its cuda-aware MPI implementation, so we cannot just async allocate
621:       // blindly...
622:       if (stream) {
623:         PetscCallCUPM(cupmMallocAsync(reinterpret_cast<void **>(ptr), bytes, stream));
624:       } else {
625:         PetscCallCUPM(cupmMalloc(reinterpret_cast<void **>(ptr), bytes));
626:       }
627:     }
628:     PetscFunctionReturn(PETSC_SUCCESS);
629:   }

631:   template <typename M>
632:   static PetscErrorCode PetscCUPMMalloc(M **ptr, std::size_t n) noexcept
633:   {
634:     PetscFunctionBegin;
635:     PetscCall(PetscCUPMMallocAsync(ptr, n));
636:     PetscFunctionReturn(PETSC_SUCCESS);
637:   }

639:   template <typename M>
640:   static PetscErrorCode PetscCUPMMallocHost(M **ptr, std::size_t n, unsigned int flags = cupmHostAllocDefault) noexcept
641:   {
642:     static_assert(!std::is_void<M>::value, "");

644:     PetscFunctionBegin;
645:     PetscAssertPointer(ptr, 1);
646:     *ptr = nullptr;
647:     if (n) PetscCallCUPM(cupmMallocHost(reinterpret_cast<void **>(ptr), n * sizeof(M), flags));
648:     PetscFunctionReturn(PETSC_SUCCESS);
649:   }

651:   template <typename D>
652:   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
653:   {
654:     static_assert(!std::is_void<D>::value, "");
655:     const auto size = n * sizeof(D);

657:     PetscFunctionBegin;
658:     if (PetscUnlikely(!n)) PetscFunctionReturn(PETSC_SUCCESS);
659:     // cannot dereference (i.e. cannot call PetscAssertPointer() here)
660:     PetscCheck(dest, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to copy to a NULL pointer");
661:     PetscCheck(src, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to copy from a NULL pointer");
662:     // do early return after nullptr check since we need to check that they are not both nullptrs
663:     if (PetscUnlikely(dest == src)) PetscFunctionReturn(PETSC_SUCCESS);
664:     if (kind == cupmMemcpyHostToHost) {
665:       // If we are HTOH it is cheaper to check if the stream is idle and do a basic mempcy()
666:       // than it is to just call the vendor functions. This assumes of course that the stream
667:       // accounts for both memory regions being "idle"
668:       if (cupmStreamQuery(stream) == cupmSuccess) {
669:         PetscCall(PetscMemcpy(dest, src, size));
670:         PetscFunctionReturn(PETSC_SUCCESS);
671:       }
672:       // need to clear the potential cupmErrorNotReady generated by query above...
673:       auto cerr = cupmGetLastError();

675:       if (PetscUnlikely(cerr != cupmErrorNotReady)) PetscCallCUPM(cerr);
676:     }
677:     if (use_async || stream || (kind != cupmMemcpyDeviceToHost)) {
678:       PetscCallCUPM(cupmMemcpyAsync(dest, src, size, kind, stream));
679:     } else {
680:       PetscCallCUPM(cupmMemcpy(dest, src, size, kind));
681:     }
682:     PetscCall(PetscLogCUPMMemcpyTransfer(kind, size));
683:     PetscFunctionReturn(PETSC_SUCCESS);
684:   }

686:   template <typename D>
687:   static PetscErrorCode PetscCUPMMemcpy(D *dest, const util::type_identity_t<D> *src, std::size_t n, cupmMemcpyKind_t kind) noexcept
688:   {
689:     PetscFunctionBegin;
690:     PetscCall(PetscCUPMMemcpyAsync(dest, src, n, kind));
691:     PetscFunctionReturn(PETSC_SUCCESS);
692:   }

694:   template <typename D>
695:   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)
696:   {
697:     static_assert(!std::is_void<D>::value, "");
698:     const auto dest_pitch_bytes = dest_pitch * sizeof(D);
699:     const auto src_pitch_bytes  = src_pitch * sizeof(D);
700:     const auto width_bytes      = width * sizeof(D);
701:     const auto size             = height * width_bytes;

703:     PetscFunctionBegin;
704:     if (PetscUnlikely(!size)) PetscFunctionReturn(PETSC_SUCCESS);
705:     PetscCheck(dest, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to copy to a NULL pointer");
706:     PetscCheck(src, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to copy from a NULL pointer");
707:     if (stream || (kind != cupmMemcpyDeviceToHost)) {
708:       PetscCallCUPM(cupmMemcpy2DAsync(dest, dest_pitch_bytes, src, src_pitch_bytes, width_bytes, height, kind, stream));
709:     } else {
710:       PetscCallCUPM(cupmMemcpy2D(dest, dest_pitch_bytes, src, src_pitch_bytes, width_bytes, height, kind));
711:     }
712:     PetscCall(PetscLogCUPMMemcpyTransfer(kind, size));
713:     PetscFunctionReturn(PETSC_SUCCESS);
714:   }

716:   template <typename D>
717:   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)
718:   {
719:     PetscFunctionBegin;
720:     PetscCall(PetscCUPMMemcpy2DAsync(dest, dest_pitch, src, src_pitch, width, height, kind));
721:     PetscFunctionReturn(PETSC_SUCCESS);
722:   }

724:   template <typename M>
725:   static PetscErrorCode PetscCUPMMemsetAsync(M *ptr, int value, std::size_t n, cupmStream_t stream = nullptr, bool use_async = false) noexcept
726:   {
727:     static_assert(!std::is_void<M>::value, "");

729:     PetscFunctionBegin;
730:     if (PetscLikely(n)) {
731:       const auto bytes = n * sizeof(M);

733:       PetscCheck(ptr, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to memset a NULL pointer with size %zu != 0", n);
734:       if (stream || use_async) {
735:         PetscCallCUPM(cupmMemsetAsync(ptr, value, bytes, stream));
736:       } else {
737:         PetscCallCUPM(cupmMemset(ptr, value, bytes));
738:       }
739:     }
740:     PetscFunctionReturn(PETSC_SUCCESS);
741:   }

743:   template <typename M>
744:   static PetscErrorCode PetscCUPMMemset(M *ptr, int value, std::size_t n) noexcept
745:   {
746:     PetscFunctionBegin;
747:     PetscCall(PetscCUPMMemsetAsync(ptr, value, n));
748:     PetscFunctionReturn(PETSC_SUCCESS);
749:   }

751:   template <typename D>
752:   static PetscErrorCode PetscCUPMMemset2DAsync(D *ptr, std::size_t pitch, int value, std::size_t width, std::size_t height, cupmStream_t stream = nullptr)
753:   {
754:     static_assert(!std::is_void<D>::value, "");
755:     const auto pitch_bytes = pitch * sizeof(D);
756:     const auto width_bytes = width * sizeof(D);
757:     const auto size        = width_bytes * height;

759:     PetscFunctionBegin;
760:     if (PetscUnlikely(!size)) PetscFunctionReturn(PETSC_SUCCESS);
761:     PetscAssert(ptr, PETSC_COMM_SELF, PETSC_ERR_POINTER, "Trying to memset a NULL pointer with size %zu != 0", size);
762:     if (stream) {
763:       PetscCallCUPM(cupmMemset2DAsync(ptr, pitch_bytes, value, width_bytes, height, stream));
764:     } else {
765:       PetscCallCUPM(cupmMemset2D(ptr, pitch_bytes, value, width_bytes, height));
766:     }
767:     PetscFunctionReturn(PETSC_SUCCESS);
768:   }

770:   // these we can transparently wrap, no need to namespace it to Petsc
771:   template <typename M>
772:   PETSC_NODISCARD static cupmError_t cupmFreeAsync(M &ptr, cupmStream_t stream = nullptr) noexcept
773:   {
774:     static_assert(std::is_pointer<util::decay_t<M>>::value, "");
775:     static_assert(!std::is_const<M>::value, "");

777:     if (ptr) {
778:       auto cerr = interface_type::cupmFreeAsync(std::forward<M>(ptr), stream);

780:       ptr = nullptr;
781:       if (PetscUnlikely(cerr != cupmSuccess)) return cerr;
782:     }
783:     return cupmSuccess;
784:   }

786:   PETSC_NODISCARD static cupmError_t cupmFreeAsync(std::nullptr_t ptr, cupmStream_t stream = nullptr) { return interface_type::cupmFreeAsync(ptr, stream); }

788:   template <typename M>
789:   PETSC_NODISCARD static cupmError_t cupmFree(M &ptr) noexcept
790:   {
791:     return cupmFreeAsync(ptr);
792:   }

794:   PETSC_NODISCARD static cupmError_t cupmFree(std::nullptr_t ptr) { return cupmFreeAsync(ptr); }

796:   template <typename M>
797:   PETSC_NODISCARD static cupmError_t cupmFreeHost(M &ptr) noexcept
798:   {
799:     static_assert(std::is_pointer<util::decay_t<M>>::value, "");
800:     const auto cerr = interface_type::cupmFreeHost(std::forward<M>(ptr));
801:     ptr             = nullptr;
802:     return cerr;
803:   }

805:   PETSC_NODISCARD static cupmError_t cupmFreeHost(std::nullptr_t ptr) { return interface_type::cupmFreeHost(ptr); }

807:   // specific wrapper for device launch function, as the real function is a C routine and
808:   // doesn't have variable arguments. The actual mechanics of this are a bit complicated but
809:   // boils down to the fact that ultimately we pass a
810:   //
811:   // void *args[] = {(void*)&kernel_args...};
812:   //
813:   // to the kernel launcher. Since we pass void* this means implicit conversion does **not**
814:   // happen to the kernel arguments so we must do it ourselves here. This function does this in
815:   // 3 stages:
816:   // 1. Enumerate the kernel arguments (cupmLaunchKernel)
817:   // 2. Deduce the signature of func() and static_cast the kernel arguments to the type
818:   //    expected by func() using the enumeration above (deduceKernelCall)
819:   // 3. Form the void* array with the converted arguments and call cuda/hipLaunchKernel with
820:   //    it. (interface_type::cupmLaunchKernel)
821:   template <typename F, typename... Args>
822:   PETSC_NODISCARD static cupmError_t cupmLaunchKernel(F &&func, cupmDim3 gridDim, cupmDim3 blockDim, std::size_t sharedMem, cupmStream_t stream, Args &&...kernelArgs) noexcept
823:   {
824:     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)...);
825:   }

827:   template <std::size_t block_size = 256, std::size_t warp_size = 32, typename F, typename... Args>
828:   static PetscErrorCode PetscCUPMLaunchKernel1D(std::size_t n, std::size_t sharedMem, cupmStream_t stream, F &&func, Args &&...kernelArgs) noexcept
829:   {
830:     static_assert(block_size > 0, "");
831:     static_assert(warp_size > 0, "");
832:     // want block_size to be a multiple of the warp_size
833:     static_assert(block_size % warp_size == 0, "");
834:     const auto nthread = std::min(n, block_size);
835:     const auto nblock  = (n + block_size - 1) / block_size;

837:     PetscFunctionBegin;
838:     // if n = 0 then nthread = 0, which is not allowed. rather than letting the user try to
839:     // decipher cryptic 'cuda/hipErrorLaunchFailure' we explicitly check for zero here
840:     PetscAssert(nthread, PETSC_COMM_SELF, PETSC_ERR_ARG_WRONG, "Trying to launch kernel with grid/block size 0");
841:     PetscCallCUPM(cupmLaunchKernel(std::forward<F>(func), nblock, nthread, sharedMem, stream, std::forward<Args>(kernelArgs)...));
842:     PetscFunctionReturn(PETSC_SUCCESS);
843:   }

845: private:
846:   template <typename S, typename D, typename = void>
847:   struct is_static_castable : std::false_type { };

849:   template <typename S, typename D>
850:   struct is_static_castable<S, D, util::void_t<decltype(static_cast<D>(std::declval<S>()))>> : std::true_type { };

852:   template <typename D, typename S>
853:   static constexpr util::enable_if_t<is_static_castable<S, D>::value, D> cast_to(S &&src) noexcept
854:   {
855:     return static_cast<D>(std::forward<S>(src));
856:   }

858:   template <typename D, typename S>
859:   static constexpr util::enable_if_t<!is_static_castable<S, D>::value, D> cast_to(S &&src) noexcept
860:   {
861:     return const_cast<D>(std::forward<S>(src));
862:   }

864:   template <typename F, typename... Args, std::size_t... Idx>
865:   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
866:   {
867:     // clang-format off
868:     return interface_type::template cupmLaunchKernel(
869:       std::forward<F>(func),
870:       std::move(gridDim), std::move(blockDim), std::move(sharedMem), std::move(stream),
871:       // can't static_cast() here since the function argument type may be cv-qualified, in
872:       // which case we would need to const_cast(). But you can only const_cast() indirect types
873:       // (pointers, references). So we need a SFINAE monster that is a static_cast() if
874:       // possible, and a const_cast() if not. We could just use a C-style cast which *would*
875:       // work here since it tries the following and uses the first one that succeeds:
876:       //
877:       // 1. const_cast()
878:       // 2. static_cast()
879:       // 3. static_cast() then const_cast()
880:       // 4. reinterpret_cast()...
881:       //
882:       // the issue however is the final reinterpret_cast(). We absolutely cannot get there
883:       // because doing so would silently hide a ton of bugs, for example casting a PetscScalar
884:       // * to double * in complex builds, a PetscInt * to int * in 64idx builds, etc.
885:       cast_to<typename util::func_traits<F>::template arg<Idx>::type>(std::forward<Args>(kernelArgs))...
886:     );
887:     // clang-format on
888:   }

890:   static PetscErrorCode PetscLogCUPMMemcpyTransfer(cupmMemcpyKind_t kind, std::size_t size) noexcept
891:   {
892:     PetscFunctionBegin;
893:     // only the explicit HTOD or DTOH are handled, since we either don't log the other cases
894:     // (yet) or don't know the direction
895:     if (kind == cupmMemcpyDeviceToHost) PetscCall(PetscLogGpuToCpu(static_cast<PetscLogDouble>(size)));
896:     else if (kind == cupmMemcpyHostToDevice) PetscCall(PetscLogCpuToGpu(static_cast<PetscLogDouble>(size)));
897:     else (void)size;
898:     PetscFunctionReturn(PETSC_SUCCESS);
899:   }
900: };

902: #undef PETSC_GCC_LINKER_UNDEFINED_REFERENCE_BUG_WORKAROUND

904: #define PETSC_CUPM_INHERIT_INTERFACE_TYPEDEFS_USING(T) \
905:   PETSC_CUPM_IMPL_CLASS_HEADER(T); \
906:   using cupmReal_t   = typename ::Petsc::device::cupm::impl::Interface<T>::cupmReal_t; \
907:   using cupmScalar_t = typename ::Petsc::device::cupm::impl::Interface<T>::cupmScalar_t; \
908:   using ::Petsc::device::cupm::impl::Interface<T>::cupmScalarCast; \
909:   using ::Petsc::device::cupm::impl::Interface<T>::cupmScalarPtrCast; \
910:   using ::Petsc::device::cupm::impl::Interface<T>::cupmRealPtrCast; \
911:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMGetMemType; \
912:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemset; \
913:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemsetAsync; \
914:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMalloc; \
915:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMallocAsync; \
916:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMallocHost; \
917:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemcpy; \
918:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemcpyAsync; \
919:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemcpy2D; \
920:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemcpy2DAsync; \
921:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMMemset2DAsync; \
922:   using ::Petsc::device::cupm::impl::Interface<T>::cupmFree; \
923:   using ::Petsc::device::cupm::impl::Interface<T>::cupmFreeAsync; \
924:   using ::Petsc::device::cupm::impl::Interface<T>::cupmFreeHost; \
925:   using ::Petsc::device::cupm::impl::Interface<T>::cupmLaunchKernel; \
926:   using ::Petsc::device::cupm::impl::Interface<T>::PetscCUPMLaunchKernel1D; \
927:   using ::Petsc::device::cupm::impl::Interface<T>::PetscDeviceCopyModeToCUPMMemcpyKind

929: #if PetscDefined(HAVE_CUDA)
930: extern template struct PETSC_SINGLE_LIBRARY_VISIBILITY_INTERNAL Interface<DeviceType::CUDA>;
931: #endif

933: #if PetscDefined(HAVE_HIP)
934: extern template struct PETSC_SINGLE_LIBRARY_VISIBILITY_INTERNAL Interface<DeviceType::HIP>;
935: #endif

937: } // namespace impl

939: } // namespace cupm

941: } // namespace device

943: } // namespace Petsc