Profiler wrapper to allow for cudaDeviceSynchronize (#738)

* add warpx profiler wrapper

* call WX profiler wrapper instead of BL

* EOL and cleaning

* add doc on profile syncs

* do not use PROFILE macros in .H files

* replace WX_PROFILE with WARPX_PROFILE and use bool instead of int

* typo

* this file shouldn't be there

4 files changed, 30 insertions, 2 deletions

diff --git a/Source/Utils/Make.package b/Source/Utils/Make.package
index 7e814ba89..022ff286a 100644
--- a/Source/Utils/Make.package
+++ b/Source/Utils/Make.package
@@ -12,6 +12,7 @@ CEXE_sources += WarpXAlgorithmSelection.cpp
 CEXE_headers += NCIGodfreyTables.H
 CEXE_headers += WarpX_Complex.H
 CEXE_headers += IonizationEnergiesTable.H
+CEXE_headers += WarpXProfilerWrapper.H
 
 INCLUDE_LOCATIONS += $(WARPX_HOME)/Source/Utils
 VPATH_LOCATIONS   += $(WARPX_HOME)/Source/Utils
diff --git a/Source/Utils/WarpXMovingWindow.cpp b/Source/Utils/WarpXMovingWindow.cpp
index 8f1666b2f..d7bdaf45f 100644
--- a/Source/Utils/WarpXMovingWindow.cpp
+++ b/Source/Utils/WarpXMovingWindow.cpp
@@ -235,7 +235,7 @@ WarpX::shiftMF (MultiFab& mf, const Geometry& geom, int num_shift, int dir,
                 IntVect ng_extra, amrex::Real external_field, bool useparser,
                 ParserWrapper<3> *field_parser)
 {
-    BL_PROFILE("WarpX::shiftMF()");
+    WARPX_PROFILE("WarpX::shiftMF()");
     const BoxArray& ba = mf.boxArray();
     const DistributionMapping& dm = mf.DistributionMap();
     const int nc = mf.nComp();
diff --git a/Source/Utils/WarpXProfilerWrapper.H b/Source/Utils/WarpXProfilerWrapper.H
new file mode 100644
index 000000000..6d18e22ba
--- /dev/null
+++ b/Source/Utils/WarpXProfilerWrapper.H
@@ -0,0 +1,27 @@
+/* Copyright 2020 Axel Huebl, Maxence Thevenet
+ *
+ * This file is part of WarpX.
+ *
+ * License: BSD-3-Clause-LBNL
+ */
+
+#ifndef WARPX_PROFILERWRAPPER_H_
+#define WARPX_PROFILERWRAPPER_H_
+
+#include "AMReX_BLProfiler.H"
+#include "AMReX_GpuDevice.H"
+
+static void doDeviceSynchronize(int do_device_synchronize)
+{
+    if ( do_device_synchronize )
+        amrex::Gpu::synchronize();
+}
+
+#define WARPX_PROFILE(fname) doDeviceSynchronize(WarpX::do_device_synchronize_before_profile); BL_PROFILE(fname)
+#define WARPX_PROFILE_VAR(fname, vname) doDeviceSynchronize(WarpX::do_device_synchronize_before_profile); BL_PROFILE_VAR(fname, vname)
+#define WARPX_PROFILE_VAR_NS(fname, vname) doDeviceSynchronize(WarpX::do_device_synchronize_before_profile); BL_PROFILE_VAR_NS(fname, vname)
+#define WARPX_PROFILE_VAR_START(vname) doDeviceSynchronize(WarpX::do_device_synchronize_before_profile); BL_PROFILE_VAR_START(vname)
+#define WARPX_PROFILE_VAR_STOP(vname) doDeviceSynchronize(WarpX::do_device_synchronize_before_profile); BL_PROFILE_VAR_STOP(vname)
+#define WARPX_PROFILE_REGION(rname) doDeviceSynchronize(WarpX::do_device_synchronize_before_profile); BL_PROFILE_REGION(rname)
+
+#endif // WARPX_PROFILERWRAPPER_H_
diff --git a/Source/Utils/WarpXUtil.cpp b/Source/Utils/WarpXUtil.cpp
index 1805bb4ed..63cf72aa2 100644
--- a/Source/Utils/WarpXUtil.cpp
+++ b/Source/Utils/WarpXUtil.cpp
@@ -126,7 +126,7 @@ void ConvertLabParamsToBoost()
  * zmin and zmax.
  */
 void NullifyMF(amrex::MultiFab& mf, int lev, amrex::Real zmin, amrex::Real zmax){
-    BL_PROFILE("WarpX::NullifyMF()");
+    WARPX_PROFILE("WarpX::NullifyMF()");
 #ifdef _OPENMP
 #pragma omp parallel if (Gpu::notInLaunchRegion())
 #endif