diff options
Diffstat (limited to 'Source/FieldSolver/FiniteDifferenceSolver')
3 files changed, 73 insertions, 20 deletions
diff --git a/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp b/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp index 338cc31f9..2f66796f3 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp @@ -5,6 +5,7 @@ * License: BSD-3-Clause-LBNL */ +#include "WarpX.H" #include "Utils/WarpXAlgorithmSelection.H" #include "FiniteDifferenceSolver.H" #ifdef WARPX_DIM_RZ @@ -24,27 +25,27 @@ using namespace amrex; void FiniteDifferenceSolver::EvolveB ( std::array< std::unique_ptr<amrex::MultiFab>, 3 >& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Efield, - 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){ - EvolveBCylindrical <CylindricalYeeAlgorithm> ( Bfield, Efield, dt ); + EvolveBCylindrical <CylindricalYeeAlgorithm> ( Bfield, Efield, lev, dt ); #else if (m_do_nodal) { - EvolveBCartesian <CartesianNodalAlgorithm> ( Bfield, Efield, dt ); + EvolveBCartesian <CartesianNodalAlgorithm> ( Bfield, Efield, lev, dt ); } else if (m_fdtd_algo == MaxwellSolverAlgo::Yee) { - EvolveBCartesian <CartesianYeeAlgorithm> ( Bfield, Efield, dt ); + EvolveBCartesian <CartesianYeeAlgorithm> ( Bfield, Efield, lev, dt ); } else if (m_fdtd_algo == MaxwellSolverAlgo::CKC) { - EvolveBCartesian <CartesianCKCAlgorithm> ( Bfield, Efield, dt ); + EvolveBCartesian <CartesianCKCAlgorithm> ( Bfield, Efield, lev, dt ); #endif } else { @@ -60,13 +61,20 @@ template<typename T_Algo> void FiniteDifferenceSolver::EvolveBCartesian ( std::array< std::unique_ptr<amrex::MultiFab>, 3 >& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Efield, - 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(*Bfield[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& Bx = Bfield[0]->array(mfi); @@ -109,8 +117,13 @@ void FiniteDifferenceSolver::EvolveBCartesian ( ); + 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); + } } - } #else // corresponds to ifndef WARPX_DIM_RZ @@ -119,13 +132,20 @@ template<typename T_Algo> void FiniteDifferenceSolver::EvolveBCylindrical ( std::array< std::unique_ptr<amrex::MultiFab>, 3 >& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Efield, - 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(*Bfield[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& Br = Bfield[0]->array(mfi); @@ -214,8 +234,13 @@ void FiniteDifferenceSolver::EvolveBCylindrical ( ); + 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); + } } - } #endif // corresponds to ifndef WARPX_DIM_RZ 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 } diff --git a/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H b/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H index bb96dee45..72fabb54e 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H +++ b/Source/FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H @@ -38,13 +38,13 @@ class FiniteDifferenceSolver void EvolveB ( std::array< std::unique_ptr<amrex::MultiFab>, 3 >& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Efield, - amrex::Real const dt ); + int lev, amrex::Real const dt ); void EvolveE ( std::array< std::unique_ptr<amrex::MultiFab>, 3 >& Efield, 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 ); void EvolveF ( std::unique_ptr<amrex::MultiFab>& Ffield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Efield, @@ -114,6 +114,7 @@ class FiniteDifferenceSolver void EvolveBCylindrical ( std::array< std::unique_ptr<amrex::MultiFab>, 3 >& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Efield, + const int lev, amrex::Real const dt ); template< typename T_Algo > @@ -122,6 +123,7 @@ class FiniteDifferenceSolver 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, + const int lev, amrex::Real const dt ); template< typename T_Algo > @@ -142,7 +144,7 @@ class FiniteDifferenceSolver void EvolveBCartesian ( std::array< std::unique_ptr<amrex::MultiFab>, 3 >& Bfield, std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Efield, - amrex::Real const dt ); + int lev, amrex::Real const dt ); template< typename T_Algo > void EvolveECartesian ( @@ -150,7 +152,7 @@ class FiniteDifferenceSolver 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 ); template< typename T_Algo > void EvolveFCartesian ( |