diff options
Diffstat (limited to 'Source/Particles')
4 files changed, 62 insertions, 3 deletions
diff --git a/Source/Particles/Collision/CollisionHandler.cpp b/Source/Particles/Collision/CollisionHandler.cpp index 588e8d008..c4ba94a4d 100644 --- a/Source/Particles/Collision/CollisionHandler.cpp +++ b/Source/Particles/Collision/CollisionHandler.cpp @@ -46,7 +46,7 @@ void CollisionHandler::doCollisions ( amrex::Real cur_time, MultiParticleContain { for (auto& collision : allcollisions) { - collision->doCollisions( cur_time, mypc); + collision->doCollisions(cur_time, mypc); } } diff --git a/Source/Particles/Collision/PairWiseCoulombCollision.H b/Source/Particles/Collision/PairWiseCoulombCollision.H index eea91d7f9..15ae329f6 100644 --- a/Source/Particles/Collision/PairWiseCoulombCollision.H +++ b/Source/Particles/Collision/PairWiseCoulombCollision.H @@ -20,6 +20,7 @@ public: /** Perform the collisions * + * @param lev AMR level of the tile * @param cur_time Current time * @param mypc Container of species involved * @@ -28,7 +29,6 @@ public: /** Perform all binary collisions within a tile * - * @param lev AMR level of the tile * @param mfi iterator for multifab * @param species1/2 pointer to species container * diff --git a/Source/Particles/Collision/PairWiseCoulombCollision.cpp b/Source/Particles/Collision/PairWiseCoulombCollision.cpp index 0bdaf7cc8..efa1c04f9 100644 --- a/Source/Particles/Collision/PairWiseCoulombCollision.cpp +++ b/Source/Particles/Collision/PairWiseCoulombCollision.cpp @@ -36,7 +36,6 @@ PairWiseCoulombCollision::PairWiseCoulombCollision (std::string const collision_ void PairWiseCoulombCollision::doCollisions (amrex::Real cur_time, MultiParticleContainer* mypc) { - const amrex::Real dt = WarpX::GetInstance().getdt(0); if ( int(std::floor(cur_time/dt)) % m_ndt != 0 ) return; @@ -50,13 +49,28 @@ PairWiseCoulombCollision::doCollisions (amrex::Real cur_time, MultiParticleConta // Loop over refinement levels for (int lev = 0; lev <= species1.finestLevel(); ++lev){ + amrex::LayoutData<amrex::Real>* cost = WarpX::getCosts(lev); + // Loop over all grids/tiles at this level #ifdef AMREX_USE_OMP info.SetDynamic(true); #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif for (amrex::MFIter mfi = species1.MakeMFIter(lev, info); mfi.isValid(); ++mfi){ + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } + amrex::Real wt = amrex::second(); + doCoulombCollisionsWithinTile( lev, mfi, species1, species2 ); + + 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); + } } } } diff --git a/Source/Particles/MultiParticleContainer.cpp b/Source/Particles/MultiParticleContainer.cpp index 658dbf23f..5501d43d8 100644 --- a/Source/Particles/MultiParticleContainer.cpp +++ b/Source/Particles/MultiParticleContainer.cpp @@ -687,6 +687,8 @@ MultiParticleContainer::doFieldIonization (int lev, { WARPX_PROFILE("MultiParticleContainer::doFieldIonization()"); + amrex::LayoutData<amrex::Real>* cost = WarpX::getCosts(lev); + // Loop over all species. // Ionized particles in pc_source create particles in pc_product for (auto& pc_source : allcontainers) @@ -711,6 +713,12 @@ MultiParticleContainer::doFieldIonization (int lev, #endif for (WarpXParIter pti(*pc_source, lev, info); pti.isValid(); ++pti) { + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } + Real wt = amrex::second(); + auto& src_tile = pc_source ->ParticlesAt(lev, pti); auto& dst_tile = pc_product->ParticlesAt(lev, pti); @@ -723,6 +731,13 @@ MultiParticleContainer::doFieldIonization (int lev, Filter, Copy, Transform); setNewParticleIDs(dst_tile, np_dst, num_added); + + 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); + } } } } @@ -1252,6 +1267,8 @@ void MultiParticleContainer::doQedBreitWheeler (int lev, { WARPX_PROFILE("MultiParticleContainer::doQedBreitWheeler()"); + amrex::LayoutData<amrex::Real>* cost = WarpX::getCosts(lev); + // Loop over all species. // Photons undergoing Breit Wheeler process create electrons // in pc_product_ele and positrons in pc_product_pos @@ -1286,6 +1303,12 @@ void MultiParticleContainer::doQedBreitWheeler (int lev, #endif for (WarpXParIter pti(*pc_source, lev, info); pti.isValid(); ++pti) { + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } + Real wt = amrex::second(); + auto Transform = PairGenerationTransformFunc(pair_gen_functor, pti, lev, Ex.nGrow(), Ex[pti], Ey[pti], Ez[pti], @@ -1305,6 +1328,13 @@ void MultiParticleContainer::doQedBreitWheeler (int lev, setNewParticleIDs(dst_ele_tile, np_dst_ele, num_added); setNewParticleIDs(dst_pos_tile, np_dst_pos, num_added); + + 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); + } } } } @@ -1319,6 +1349,8 @@ void MultiParticleContainer::doQedQuantumSync (int lev, { WARPX_PROFILE("MultiParticleContainer::doQedQuantumSync()"); + amrex::LayoutData<amrex::Real>* cost = WarpX::getCosts(lev); + // Loop over all species. // Electrons or positrons undergoing Quantum photon emission process // create photons in pc_product_phot @@ -1347,6 +1379,12 @@ void MultiParticleContainer::doQedQuantumSync (int lev, #endif for (WarpXParIter pti(*pc_source, lev, info); pti.isValid(); ++pti) { + if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) + { + amrex::Gpu::synchronize(); + } + Real wt = amrex::second(); + auto Transform = PhotonEmissionTransformFunc( m_shr_p_qs_engine->build_optical_depth_functor(), pc_source->particle_runtime_comps["optical_depth_QSR"], @@ -1370,6 +1408,13 @@ void MultiParticleContainer::doQedQuantumSync (int lev, cleanLowEnergyPhotons( dst_tile, np_dst, num_added, m_quantum_sync_photon_creation_energy_threshold); + + 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); + } } } } |