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