WarpX
Loading...
Searching...
No Matches
KernelTimer.H
Go to the documentation of this file.
1/* Copyright 2019-2020 Michael Rowan, Axel Huebl, Kevin Gott
2 *
3 * This file is part of WarpX.
4 *
5 * License: BSD-3-Clause-LBNL
6 */
7#ifndef ABLASTR_KERNELTIMER_H_
8#define ABLASTR_KERNELTIMER_H_
9
11
12#include <AMReX.H>
13#include <AMReX_BLassert.H>
14#include <AMReX_GpuAtomic.H>
15#include <AMReX_GpuQualifiers.H>
16#include <AMReX_REAL.H>
17
18#include <climits>
19
21{
22
27{
28public:
36 (const bool do_timing, amrex::Real* cost)
37#if (defined AMREX_USE_GPU)
38 : m_do_timing(do_timing)
39#endif
40 {
41#if (defined AMREX_USE_GPU)
42 if (m_do_timing && cost) {
43 m_cost = cost;
44# if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP)
45 // Start the timer
46 m_wt = clock64();
47
48# elif defined(AMREX_USE_DPCPP)
49 // To be updated
50 ABLASTR_ALWAYS_ASSERT_WITH_MESSAGE(m_do_timing == false,
51 "KernelTimer not yet supported for this hardware.");
52# endif
53 }
54#else // AMREX_USE_GPU
55 amrex::ignore_unused(do_timing, cost);
56#endif // AMREX_USE_GPU
57 }
58
60#if (defined AMREX_USE_GPU)
61# if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP)
64 {
65 if (m_do_timing && m_cost) {
66 m_wt = clock64() - m_wt;
67 amrex::Gpu::Atomic::Add( m_cost, amrex::Real(m_wt));
68 }
69 }
70# elif defined(AMREX_USE_DPCPP)
71 // To be updated
73 ~KernelTimer () = default;
74# endif
75
76#else
77~KernelTimer () = default;
78#endif //AMREX_USE_GPU
79
80
81 KernelTimer ( KernelTimer const &) = default;
82 KernelTimer& operator= ( KernelTimer const & ) = default;
83 KernelTimer ( KernelTimer&& ) = default;
85
86#if (defined AMREX_USE_GPU)
87private:
89 bool m_do_timing;
90
92 amrex::Real* m_cost = nullptr;
93
95 long long int m_wt;
96#endif //AMREX_USE_GPU
97};
98
99} // namespace ablastr::parallelization
100
101#endif // ABLASTR_KERNELTIMER_H_
#define AMREX_GPU_DEVICE
#define ABLASTR_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition TextMsg.H:75
KernelTimer(KernelTimer &&)=default
KernelTimer & operator=(KernelTimer const &)=default
AMREX_GPU_DEVICE KernelTimer(const bool do_timing, amrex::Real *cost)
Definition KernelTimer.H:36
KernelTimer(KernelTimer const &)=default
Definition KernelTimer.H:21
__host__ __device__ AMREX_FORCE_INLINE T Add(T *sum, T value) noexcept
__host__ __device__ void ignore_unused(const Ts &...)