xref: /petsc/src/sys/objects/device/impls/cupm/cupmthrustutility.hpp (revision 7233ce556edaaa29ce4b396318a4ca76d3b0291b)
1a4963045SJacob Faibussowitsch #pragma once
20e6b6b59SJacob Faibussowitsch 
3f3146f24SJacob Faibussowitsch #include <petsclog.h>         // PetscLogGpuTimeBegin()/End()
4*7233ce55SJed Brown #include <petscsys.h>         // SETERRQ()
5f3146f24SJacob Faibussowitsch #include <petscdevice_cupm.h> // PETSC_USING_NVCC
6f3146f24SJacob Faibussowitsch 
7ebc16b7dSJacob Faibussowitsch #include <thrust/version.h>          // THRUST_VERSION
8f3146f24SJacob Faibussowitsch #include <thrust/system_error.h>     // thrust::system_error
9f3146f24SJacob Faibussowitsch #include <thrust/execution_policy.h> // thrust::cuda/hip::par
100e6b6b59SJacob Faibussowitsch 
11d71ae5a4SJacob Faibussowitsch namespace Petsc
12d71ae5a4SJacob Faibussowitsch {
130e6b6b59SJacob Faibussowitsch 
14d71ae5a4SJacob Faibussowitsch namespace device
15d71ae5a4SJacob Faibussowitsch {
160e6b6b59SJacob Faibussowitsch 
17d71ae5a4SJacob Faibussowitsch namespace cupm
18d71ae5a4SJacob Faibussowitsch {
190e6b6b59SJacob Faibussowitsch 
200e6b6b59SJacob Faibussowitsch #if PetscDefined(USING_NVCC)
210e6b6b59SJacob Faibussowitsch   #if !defined(THRUST_VERSION)
220e6b6b59SJacob Faibussowitsch     #error "THRUST_VERSION not defined!"
230e6b6b59SJacob Faibussowitsch   #endif
24ebc16b7dSJacob Faibussowitsch   #if THRUST_VERSION >= 101600
25ebc16b7dSJacob Faibussowitsch     #define PETSC_THRUST_HAS_ASYNC                 1
26f3146f24SJacob Faibussowitsch     #define PETSC_THRUST_CALL_PAR_ON(func, s, ...) func(thrust::cuda::par_nosync.on(s), __VA_ARGS__)
270e6b6b59SJacob Faibussowitsch   #else
28f3146f24SJacob Faibussowitsch     #define PETSC_THRUST_CALL_PAR_ON(func, s, ...) func(thrust::cuda::par.on(s), __VA_ARGS__)
290e6b6b59SJacob Faibussowitsch   #endif
300e6b6b59SJacob Faibussowitsch #elif PetscDefined(USING_HCC) // rocThrust has no par_nosync
31f3146f24SJacob Faibussowitsch   #define PETSC_THRUST_CALL_PAR_ON(func, s, ...) func(thrust::hip::par.on(s), __VA_ARGS__)
320e6b6b59SJacob Faibussowitsch #else
33f3146f24SJacob Faibussowitsch   #define PETSC_THRUST_CALL_PAR_ON(func, s, ...) func(__VA_ARGS__)
340e6b6b59SJacob Faibussowitsch #endif
350e6b6b59SJacob Faibussowitsch 
36ebc16b7dSJacob Faibussowitsch #ifndef PETSC_THRUST_HAS_ASYNC
37ebc16b7dSJacob Faibussowitsch   #define PETSC_THRUST_HAS_ASYNC 0
38ebc16b7dSJacob Faibussowitsch #endif
39ebc16b7dSJacob Faibussowitsch 
40d71ae5a4SJacob Faibussowitsch namespace detail
41d71ae5a4SJacob Faibussowitsch {
420e6b6b59SJacob Faibussowitsch 
430e6b6b59SJacob Faibussowitsch struct PetscLogGpuTimer {
44ebc16b7dSJacob Faibussowitsch   PetscLogGpuTimer() noexcept
45ebc16b7dSJacob Faibussowitsch   {
46ebc16b7dSJacob Faibussowitsch     PetscFunctionBegin;
47ebc16b7dSJacob Faibussowitsch     PetscCallAbort(PETSC_COMM_SELF, PetscLogGpuTimeBegin());
48ebc16b7dSJacob Faibussowitsch     PetscFunctionReturnVoid();
49ebc16b7dSJacob Faibussowitsch   }
50ebc16b7dSJacob Faibussowitsch 
51ebc16b7dSJacob Faibussowitsch   ~PetscLogGpuTimer() noexcept
52ebc16b7dSJacob Faibussowitsch   {
53ebc16b7dSJacob Faibussowitsch     PetscFunctionBegin;
54ebc16b7dSJacob Faibussowitsch     PetscCallAbort(PETSC_COMM_SELF, PetscLogGpuTimeEnd());
55ebc16b7dSJacob Faibussowitsch     PetscFunctionReturnVoid();
56ebc16b7dSJacob Faibussowitsch   }
570e6b6b59SJacob Faibussowitsch };
580e6b6b59SJacob Faibussowitsch 
590e6b6b59SJacob Faibussowitsch } // namespace detail
600e6b6b59SJacob Faibussowitsch 
610e6b6b59SJacob Faibussowitsch #define THRUST_CALL(...) \
620e6b6b59SJacob Faibussowitsch   [&] { \
63f3146f24SJacob Faibussowitsch     const auto timer = ::Petsc::device::cupm::detail::PetscLogGpuTimer{}; \
64f3146f24SJacob Faibussowitsch     return PETSC_THRUST_CALL_PAR_ON(__VA_ARGS__); \
650e6b6b59SJacob Faibussowitsch   }()
660e6b6b59SJacob Faibussowitsch 
670e6b6b59SJacob Faibussowitsch #define PetscCallThrust(...) \
680e6b6b59SJacob Faibussowitsch   do { \
690e6b6b59SJacob Faibussowitsch     try { \
70f3146f24SJacob Faibussowitsch       { \
710e6b6b59SJacob Faibussowitsch         __VA_ARGS__; \
72f3146f24SJacob Faibussowitsch       } \
73d71ae5a4SJacob Faibussowitsch     } catch (const thrust::system_error &ex) { \
74d71ae5a4SJacob Faibussowitsch       SETERRQ(PETSC_COMM_SELF, PETSC_ERR_LIB, "Thrust error: %s", ex.what()); \
75d71ae5a4SJacob Faibussowitsch     } \
760e6b6b59SJacob Faibussowitsch   } while (0)
770e6b6b59SJacob Faibussowitsch 
780e6b6b59SJacob Faibussowitsch } // namespace cupm
790e6b6b59SJacob Faibussowitsch 
800e6b6b59SJacob Faibussowitsch } // namespace device
810e6b6b59SJacob Faibussowitsch 
820e6b6b59SJacob Faibussowitsch } // namespace Petsc
83