diff options
Diffstat (limited to 'Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp')
| -rw-r--r-- | Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp | 40 |
1 files changed, 33 insertions, 7 deletions
diff --git a/Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp b/Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp index 3762d3f71..cd1e3bdf9 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp @@ -5,6 +5,7 @@ * License: BSD-3-Clause-LBNL */ +#include "WarpX.H" #include "Utils/WarpXAlgorithmSelection.H" #include "FiniteDifferenceSolver.H" #ifdef WARPX_DIM_RZ @@ -28,27 +29,27 @@ void FiniteDifferenceSolver::EvolveE ( std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Jfield, std::unique_ptr<amrex::MultiFab> const& Ffield, - amrex::Real const dt ) { + int lev, amrex::Real const dt ) { // Select algorithm (The choice of algorithm is a runtime option, // but we compile code for each algorithm, using templates) #ifdef WARPX_DIM_RZ if (m_fdtd_algo == MaxwellSolverAlgo::Yee){ - EvolveECylindrical <CylindricalYeeAlgorithm> ( Efield, Bfield, Jfield, Ffield, dt ); + EvolveECylindrical <CylindricalYeeAlgorithm> ( Efield, Bfield, Jfield, Ffield, lev, dt ); #else if (m_do_nodal) { - EvolveECartesian <CartesianNodalAlgorithm> ( Efield, Bfield, Jfield, Ffield, dt ); + EvolveECartesian <CartesianNodalAlgorithm> ( Efield, Bfield, Jfield, Ffield, lev, dt ); } else if (m_fdtd_algo == MaxwellSolverAlgo::Yee) { - EvolveECartesian <CartesianYeeAlgorithm> ( Efield, Bfield, Jfield, Ffield, dt ); + EvolveECartesian <CartesianYeeAlgorithm> ( Efield, Bfield, Jfield, Ffield, lev, dt ); } else if (m_fdtd_algo == MaxwellSolverAlgo::CKC) { - EvolveECartesian <CartesianCKCAlgorithm> ( Efield, Bfield, Jfield, Ffield, dt ); + EvolveECartesian <CartesianCKCAlgorithm> ( Efield, Bfield, Jfield, Ffield, lev, dt ); #endif } else { @@ -66,8 +67,9 @@ void FiniteDifferenceSolver::EvolveECartesian ( std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Jfield, std::unique_ptr<amrex::MultiFab> const& Ffield, - amrex::Real const dt ) { + int lev, amrex::Real const dt ) { + amrex::LayoutData<amrex::Real>* cost = WarpX::getCosts(lev); Real constexpr c2 = PhysConst::c * PhysConst::c; // Loop through the grids, and over the tiles within each grid @@ -75,6 +77,11 @@ void FiniteDifferenceSolver::EvolveECartesian ( #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*Efield[0], TilingIfNotGPU()); mfi.isValid(); ++mfi ) { + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } + Real wt = amrex::second(); // Extract field data for this grid/tile Array4<Real> const& Ex = Efield[0]->array(mfi); @@ -150,6 +157,12 @@ void FiniteDifferenceSolver::EvolveECartesian ( } + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + wt = amrex::second() - wt; + amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); + } } } @@ -162,13 +175,20 @@ void FiniteDifferenceSolver::EvolveECylindrical ( std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Jfield, std::unique_ptr<amrex::MultiFab> const& Ffield, - amrex::Real const dt ) { + int lev, amrex::Real const dt ) { + + amrex::LayoutData<amrex::Real>* cost = WarpX::getCosts(lev); // Loop through the grids, and over the tiles within each grid #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif for ( MFIter mfi(*Efield[0], TilingIfNotGPU()); mfi.isValid(); ++mfi ) { + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } + Real wt = amrex::second(); // Extract field data for this grid/tile Array4<Real> const& Er = Efield[0]->array(mfi); @@ -339,6 +359,12 @@ void FiniteDifferenceSolver::EvolveECylindrical ( } // end of if condition for F + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + wt = amrex::second() - wt; + amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); + } } // end of loop over grid/tiles } |