Actual source code: cupmthrustutility.hpp
1: #ifndef PETSC_CUPM_THRUST_UTILITY_HPP
2: #define PETSC_CUPM_THRUST_UTILITY_HPP
4: #if defined(__cplusplus)
5: #include <petsclog.h>
6: #include <petscerror.h>
7: #include <petscdevice_cupm.h>
9: #include <thrust/system_error.h> // thrust::system_error
10: #include <thrust/execution_policy.h> // thrust::cuda/hip::par
12: namespace Petsc
13: {
15: namespace device
16: {
18: namespace cupm
19: {
21: #if PetscDefined(USING_NVCC)
22: #if !defined(THRUST_VERSION)
23: #error "THRUST_VERSION not defined!"
24: #endif
25: #if !PetscDefined(USE_DEBUG) && (THRUST_VERSION >= 101600)
26: #define PETSC_THRUST_CALL_PAR_ON(func, s, ...) func(thrust::cuda::par_nosync.on(s), __VA_ARGS__)
27: #else
28: #define PETSC_THRUST_CALL_PAR_ON(func, s, ...) func(thrust::cuda::par.on(s), __VA_ARGS__)
29: #endif
30: #elif PetscDefined(USING_HCC) // rocThrust has no par_nosync
31: #define PETSC_THRUST_CALL_PAR_ON(func, s, ...) func(thrust::hip::par.on(s), __VA_ARGS__)
32: #else
33: #define PETSC_THRUST_CALL_PAR_ON(func, s, ...) func(__VA_ARGS__)
34: #endif
36: namespace detail
37: {
39: struct PetscLogGpuTimer {
40: PetscLogGpuTimer() noexcept { PetscCallAbort(PETSC_COMM_SELF, PetscLogGpuTimeBegin()); }
41: ~PetscLogGpuTimer() noexcept { PetscCallAbort(PETSC_COMM_SELF, PetscLogGpuTimeEnd()); }
42: };
44: } // namespace detail
46: #define THRUST_CALL(...) \
47: [&] { \
48: const auto timer = ::Petsc::device::cupm::detail::PetscLogGpuTimer{}; \
49: return PETSC_THRUST_CALL_PAR_ON(__VA_ARGS__); \
50: }()
52: #define PetscCallThrust(...) \
53: do { \
54: try { \
55: { \
56: __VA_ARGS__; \
57: } \
58: } catch (const thrust::system_error &ex) { \
59: SETERRQ(PETSC_COMM_SELF, PETSC_ERR_LIB, "Thrust error: %s", ex.what()); \
60: } \
61: } while (0)
63: } // namespace cupm
65: } // namespace device
67: } // namespace Petsc
69: #endif // __cplusplus
71: #endif // PETSC_CUPM_THRUST_UTILITY_HPP