diff options
Diffstat (limited to 'Source/Particles/PhysicalParticleContainer.cpp')
| -rw-r--r-- | Source/Particles/PhysicalParticleContainer.cpp | 63 |
1 files changed, 32 insertions, 31 deletions
diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 188ff4cb4..fe6fe80f2 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1,8 +1,8 @@ /* Copyright 2019-2020 Andrew Myers, Aurore Blelly, Axel Huebl * David Grote, Glenn Richardson, Jean-Luc Vay * Ligia Diana Amorim, Luca Fedeli, Maxence Thevenet - * Remi Lehe, Revathi Jambunathan, Weiqun Zhang - * Yinjian Zhao + * Michael Rowan, Remi Lehe, Revathi Jambunathan + * Weiqun Zhang, Yinjian Zhao * * This file is part of WarpX. * @@ -406,7 +406,7 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) defineAllParticleTiles(); - MultiFab* cost = WarpX::getCosts(lev); + amrex::Vector<amrex::Real>* cost = WarpX::getCosts(lev); const int nlevs = numLevels(); static bool refine_injection = false; @@ -449,6 +449,10 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) #endif for (MFIter mfi = MakeMFIter(lev, info); mfi.isValid(); ++mfi) { + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } Real wt = amrex::second(); const Box& tile_box = mfi.tilebox(); @@ -789,14 +793,11 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) #endif }); - if (cost) { - wt = (amrex::second() - wt) / tile_box.d_numPts(); - Array4<Real> const& costarr = cost->array(mfi); - amrex::ParallelFor(tile_box, - [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept - { - costarr(i,j,k) += wt; - }); + 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); } } @@ -878,7 +879,7 @@ PhysicalParticleContainer::FieldGather (int lev, BL_ASSERT(OnSameGrids(lev,Ex)); - MultiFab* cost = WarpX::getCosts(lev); + amrex::Vector<amrex::Real>* cost = WarpX::getCosts(lev); #ifdef _OPENMP #pragma omp parallel @@ -886,6 +887,10 @@ PhysicalParticleContainer::FieldGather (int lev, { for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } Real wt = amrex::second(); const Box& box = pti.validbox(); @@ -918,15 +923,11 @@ PhysicalParticleContainer::FieldGather (int lev, Ex.nGrow(), e_is_nodal, 0, np, lev, lev); - if (cost) { - const Box& tbx = pti.tilebox(); - wt = (amrex::second() - wt) / tbx.d_numPts(); - Array4<Real> const& costarr = cost->array(pti); - amrex::ParallelFor(tbx, - [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept - { - costarr(i,j,k) += wt; - }); + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + wt = amrex::second() - wt; + amrex::HostDevice::Atomic::Add( &(*cost)[pti.index()], wt); } } } @@ -958,7 +959,7 @@ PhysicalParticleContainer::Evolve (int lev, BL_ASSERT(OnSameGrids(lev,jx)); - MultiFab* cost = WarpX::getCosts(lev); + amrex::Vector<amrex::Real>* cost = WarpX::getCosts(lev); const iMultiFab* current_masks = WarpX::CurrentBufferMasks(lev); const iMultiFab* gather_masks = WarpX::GatherBufferMasks(lev); @@ -993,6 +994,10 @@ PhysicalParticleContainer::Evolve (int lev, for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } Real wt = amrex::second(); const Box& box = pti.validbox(); @@ -1182,15 +1187,11 @@ PhysicalParticleContainer::Evolve (int lev, } } - if (cost) { - const Box& tbx = pti.tilebox(); - wt = (amrex::second() - wt) / tbx.d_numPts(); - Array4<Real> const& costarr = cost->array(pti); - amrex::ParallelFor(tbx, - [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept - { - costarr(i,j,k) += wt; - }); + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + wt = amrex::second() - wt; + amrex::HostDevice::Atomic::Add( &(*cost)[pti.index()], wt); } } } |