Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions Source/Particles/Collision/BackgroundMCCCollision.H
Original file line number Diff line number Diff line change
Expand Up @@ -52,6 +52,7 @@ public:
*/
void doBackgroundIonization (
int lev,
amrex::LayoutData<amrex::Real>* cost,
WarpXParticleContainer& species1,
WarpXParticleContainer& species2
);
Expand Down
35 changes: 32 additions & 3 deletions Source/Particles/Collision/BackgroundMCCCollision.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -209,18 +209,33 @@ BackgroundMCCCollision::doCollisions (amrex::Real cur_time, MultiParticleContain
auto const flvl = species1.finestLevel();
for (int lev = 0; lev <= flvl; ++lev) {

auto cost = WarpX::getCosts(lev);

// firstly loop over particles box by box and do all particle conserving
// scattering
#ifdef _OPENMP
#pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
#endif
for (WarpXParIter pti(species1, lev); pti.isValid(); ++pti) {
if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers)
{
amrex::Gpu::synchronize();
}
amrex::Real wt = amrex::second();

doBackgroundCollisionsWithinTile(pti);

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);
}
}

// secondly perform ionization through the SmartCopyFactory if needed
if (ionization_flag) {
doBackgroundIonization(lev, species1, species2);
doBackgroundIonization(lev, cost, species1, species2);
}
}
}
Expand Down Expand Up @@ -348,8 +363,8 @@ void BackgroundMCCCollision::doBackgroundCollisionsWithinTile


void BackgroundMCCCollision::doBackgroundIonization
( int lev, WarpXParticleContainer& species1,
WarpXParticleContainer& species2)
( int lev, amrex::LayoutData<amrex::Real>* cost,
WarpXParticleContainer& species1, WarpXParticleContainer& species2)
{
WARPX_PROFILE("BackgroundMCCCollision::doBackgroundIonization()");

Expand All @@ -372,6 +387,13 @@ void BackgroundMCCCollision::doBackgroundIonization
#pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
#endif
for (WarpXParIter pti(species1, lev); pti.isValid(); ++pti) {

if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers)
{
amrex::Gpu::synchronize();
}
amrex::Real wt = amrex::second();

auto& elec_tile = species1.ParticlesAt(lev, pti);
auto& ion_tile = species2.ParticlesAt(lev, pti);

Expand All @@ -389,5 +411,12 @@ void BackgroundMCCCollision::doBackgroundIonization

setNewParticleIDs(elec_tile, np_elec, num_added);
setNewParticleIDs(ion_tile, np_ion, 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);
}
}
}