Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[WIP] Add collisional diagnostic buffers #5543

Draft
wants to merge 10 commits into
base: development
Choose a base branch
from
12 changes: 12 additions & 0 deletions Python/pywarpx/fields.py
Original file line number Diff line number Diff line change
Expand Up @@ -1030,3 +1030,15 @@ def GCPPMLWrapper(level=0, include_ghosts=False):
return _MultiFABWrapper(
mf_name="pml_G_cp", level=level, include_ghosts=include_ghosts
)


def MCCEnergyWrapper(level=0, include_ghosts=False):
return _MultiFABWrapper(
mf_name="collision_energy_change", level=level, include_ghosts=include_ghosts
)


def MCCIonizWrapper(level=0, include_ghosts=False):
return _MultiFABWrapper(
mf_name="collision_ionization", level=level, include_ghosts=include_ghosts
)
4 changes: 3 additions & 1 deletion Source/Fields.H
Original file line number Diff line number Diff line change
Expand Up @@ -81,7 +81,9 @@ namespace warpx::fields
Bfield_avg_cp,
B_old, /**< Stores the value of B at the beginning of the timestep, for the implicit solver */
ECTRhofield,
Venl
Venl,
collision_energy_change, /**< Stores collisional energy transfer data for Monte Carlo Collisions */
collision_ionization /**< Stores ionization data for Monte Carlo Collisions */
Comment on lines +85 to +86
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For clarity it would be good to add mcc to these multifab names (to make it clear they are tied to the MCC routine). Maybe something like mcc_collision_energy_change, etc.

);

/** these are vector fields */
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -45,11 +45,13 @@ public:

/** Perform particle conserving MCC collisions within a tile
*
* @param lev the mesh-refinement level
* @param species1 reference to colliding species container
* @param pti particle iterator
* @param t current time
*
*/
void doBackgroundCollisionsWithinTile ( WarpXParIter& pti, amrex::Real t);
void doBackgroundCollisionsWithinTile ( int lev, WarpXParticleContainer& species1, WarpXParIter& pti, amrex::Real t);

/** Perform MCC ionization interactions
*
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -296,7 +296,7 @@ BackgroundMCCCollision::doCollisions (amrex::Real cur_time, amrex::Real dt, Mult
}
auto wt = static_cast<amrex::Real>(amrex::second());

doBackgroundCollisionsWithinTile(pti, cur_time);
doBackgroundCollisionsWithinTile(lev, species1, pti, cur_time);

if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers)
{
Expand All @@ -315,13 +315,17 @@ BackgroundMCCCollision::doCollisions (amrex::Real cur_time, amrex::Real dt, Mult


void BackgroundMCCCollision::doBackgroundCollisionsWithinTile
( WarpXParIter& pti, amrex::Real t )
( int lev, WarpXParticleContainer& species1, WarpXParIter& pti, amrex::Real t )
{
using namespace amrex::literals;
using warpx::fields::FieldType;

// So that CUDA code gets its intrinsic, not the host-only C++ library version
using std::sqrt;

// get pointer to warpx instance, for multifab retrieval
auto & warpx = WarpX::GetInstance();

// get particle count
const long np = pti.numParticles();

Expand Down Expand Up @@ -353,6 +357,17 @@ void BackgroundMCCCollision::doBackgroundCollisionsWithinTile
amrex::ParticleReal* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr();
amrex::ParticleReal* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr();
amrex::ParticleReal* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr();
amrex::ParticleReal* const AMREX_RESTRICT w = attribs[PIdx::w].dataPtr();

// get multi-fab for collisional energy transfer (does this need to be moved out of the pti loop?)
const amrex::MultiFab &coll_E_change = *warpx.m_fields.get(FieldType::collision_energy_change, lev);
const auto& coll_E_change_arr = coll_E_change[pti].array();

// get parameters for getParticleCell (method used here is similar to TemperatureFunctor.cpp)
auto& tile = pti.GetParticleTile();
auto ptd = tile.getParticleTileData();
const auto plo = species1.Geom(lev).ProbLoArray();
const auto dxi = species1.Geom(lev).InvCellSizeArray();

amrex::ParallelForRNG(np,
[=] AMREX_GPU_HOST_DEVICE (long ip, amrex::RandomEngine const& engine)
Expand All @@ -366,8 +381,8 @@ void BackgroundMCCCollision::doBackgroundCollisionsWithinTile
const amrex::ParticleReal n_a = n_a_func(x, y, z, t);
const amrex::ParticleReal T_a = T_a_func(x, y, z, t);

amrex::ParticleReal v_coll, v_coll2, sigma_E, nu_i = 0;
double gamma, E_coll;
amrex::ParticleReal v_coll, v_coll2, sigma_E, nu_i = 0, v_lost2;
double gamma, E_coll, E_lost;
amrex::ParticleReal ua_x, ua_y, ua_z, vx, vy, vz;
amrex::ParticleReal uCOM_x, uCOM_y, uCOM_z;
const amrex::ParticleReal col_select = amrex::Random(engine);
Expand Down Expand Up @@ -406,6 +421,11 @@ void BackgroundMCCCollision::doBackgroundCollisionsWithinTile
// check if this collision should be performed
if (col_select > nu_i) { continue; }

// get the initial particle energy for tracking
if (WarpX::do_MCC_energy_tracking) {
v_lost2 = ux[ip]*ux[ip] + uy[ip]*uy[ip] + uz[ip]*uz[ip];
}

// charge exchange is implemented as a simple swap of the projectile
// and target velocities which doesn't require any of the Lorentz
// transformations below; note that if the projectile and target
Expand All @@ -414,6 +434,20 @@ void BackgroundMCCCollision::doBackgroundCollisionsWithinTile
ux[ip] = ua_x;
uy[ip] = ua_y;
uz[ip] = ua_z;

if (WarpX::do_MCC_energy_tracking) {
// Calculate the energy lost for tracking
v_lost2 -= ux[ip]*ux[ip] + uy[ip]*uy[ip] + uz[ip]*uz[ip];
ParticleUtils::getEnergy(v_lost2, m, E_lost);

// store the energy (add E_lost * w[ip] to the buffer of energy sent to the background)
// get the particle's cell index (method used here is similar to TemperatureFunctor.cpp)
const auto p = WarpXParticleContainer::ParticleType(ptd, ip);
const auto [ii, jj, kk] = amrex::getParticleCell(p , plo, dxi).dim3();

// store energy lost in the buffer
coll_E_change_arr(ii, jj, kk, 0) += E_lost * w[ip];
}
Comment on lines +438 to +450
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It would be nice if this code does not need to be duplicated as it is currently. Perhaps the break can be removed, and a else used for the non-charge exchange processes?

break;
}

Expand Down Expand Up @@ -459,6 +493,20 @@ void BackgroundMCCCollision::doBackgroundCollisionsWithinTile
ux[ip] = vx + ua_x;
uy[ip] = vy + ua_y;
uz[ip] = vz + ua_z;

if (WarpX::do_MCC_energy_tracking) {
// Calculate difference in energy for tracking
v_lost2 -= ux[ip]*ux[ip] + uy[ip]*uy[ip] + uz[ip]*uz[ip];
ParticleUtils::getEnergy(v_lost2, m, E_lost);

// store the energy (add E_lost * w[ip] to the buffer of energy sent to the background)
// get the particle's cell index (method used here is similar to TemperatureFunctor.cpp)
const auto p = WarpXParticleContainer::ParticleType(ptd, ip);
const auto [ii, jj, kk] = amrex::getParticleCell(p , plo, dxi).dim3();

// store energy lost in the buffer
coll_E_change_arr(ii, jj, kk, 0) += E_lost * w[ip];
}
break;
}
}
Expand All @@ -470,8 +518,13 @@ void BackgroundMCCCollision::doBackgroundIonization
( int lev, amrex::LayoutData<amrex::Real>* cost,
WarpXParticleContainer& species1, WarpXParticleContainer& species2, amrex::Real t)
{
using warpx::fields::FieldType;

WARPX_PROFILE("BackgroundMCCCollision::doBackgroundIonization()");

// get pointer to warpx instance, for multifab retrieval
auto & warpx = WarpX::GetInstance();

const SmartCopyFactory copy_factory_elec(species1, species1);
const SmartCopyFactory copy_factory_ion(species1, species2);
const auto CopyElec = copy_factory_elec.getSmartCopy();
Expand All @@ -485,6 +538,13 @@ void BackgroundMCCCollision::doBackgroundIonization

const amrex::ParticleReal sqrt_kb_m = std::sqrt(PhysConst::kb / m_background_mass);

// get multi-fab for ionization creation tracking
const amrex::MultiFab &coll_ionization = *warpx.m_fields.get(FieldType::collision_ionization, lev);

// get parameters for getParticleCell
const auto plo = species1.Geom(lev).ProbLoArray();
const auto dxi = species1.Geom(lev).InvCellSizeArray();

#ifdef AMREX_USE_OMP
#pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
#endif
Expand Down Expand Up @@ -515,6 +575,36 @@ void BackgroundMCCCollision::doBackgroundIonization
setNewParticleIDs(elec_tile, np_elec, num_added);
setNewParticleIDs(ion_tile, np_ion, num_added);

if (WarpX::do_MCC_ionization_tracking)
{
// get parameters for getParticleCell (method used here is similar to TemperatureFunctor.cpp)
// should these lines be removed and we just do ptd = elec_tile.getParticleTileData()?
auto& tile = pti.GetParticleTile();
auto ptd = tile.getParticleTileData();

// get multi-fab array for ionization creation tracking
const auto& coll_ionization_arr = coll_ionization[pti].array();

// get Struct-Of-Array particle data
// as long as w is set after the copy, will it include newly created particles (be of length np_elec + num_added)?
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, at this point the length will include the new particles (since this happens after setNewParticleIDs.

auto& attribs = pti.GetAttribs();
amrex::ParticleReal* const AMREX_RESTRICT w = attribs[PIdx::w].dataPtr();

// add each created particle to tracking buffer
amrex::ParallelFor(num_added,
[=] AMREX_GPU_HOST_DEVICE (long ip)
{
// store the energy (add E_lost * w[ip] to the buffer of energy sent to the background)
// get the particle's cell index (method used here is similar to TemperatureFunctor.cpp)
const auto p = WarpXParticleContainer::ParticleType(ptd, ip + np_elec);
const auto [ii, jj, kk] = amrex::getParticleCell(p , plo, dxi).dim3();

// store energy lost in the buffer
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This buffer is tracking number of electrons created, not energy, right?

coll_ionization_arr(ii, jj, kk, 0) += w[ip + np_elec];
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I've added everything that I think I'll need for this diagnostic, but am running into the following compilation error:

/home/bj8080/src/warpx/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp:603:67: error:
assignment of read-only location 'coll_ionization_arr.amrex::Array4<const double>::operator()<>(((int)ii), ((int)jj), ((int)kk), 0)'
  603 |                                coll_ionization_arr(ii, jj, kk, 0) += w[ip + np_elec];
      |                                ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~

since coll_ionization_tracking is a constant. When I correct for this by removing const from the initialization of coll_ionization_arr and the earlier line initializing coll_ionization, I get an error:

/home/bj8080/src/warpx/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp:586:67:
error: cannot bind non-const lvalue reference of type 'amrex::Array4<double>&' to an rvalue of type 'amrex::Array4<double>'
  586 |             auto& coll_ionization_arr = coll_ionization[pti].array();
      |                                         ~~~~~~~~~~~~~~~~~~~~~~~~~~^~

I think to get around this I should create a multifab to sum up the ionization events, and then add that multifab to coll_ionization_arr. How would I do this? Or, is there a better approach?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pinging this with @roelof-groenewald, since we've talked about this feature in discussion 5381. Would you be able to take a look?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think you should do

const auto& coll_ionization_arr = coll_ionization.array(pti);

Does that help?

}
);
}

if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers)
{
amrex::Gpu::synchronize();
Expand Down
4 changes: 4 additions & 0 deletions Source/WarpX.H
Original file line number Diff line number Diff line change
Expand Up @@ -868,6 +868,10 @@ public:
static amrex::Real moving_window_v;
static bool fft_do_time_averaging;

// Control whether Monte Carlo Collision statistics are collected
static bool do_MCC_energy_tracking;
static bool do_MCC_ionization_tracking;

Comment on lines +870 to +873
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These should instead be made member variables of the BackgroundMCCCollision class.

// these should be private, but can't due to Cuda limitations
static void ComputeDivB (amrex::MultiFab& divB, int dcomp,
ablastr::fields::VectorField const & B,
Expand Down
69 changes: 69 additions & 0 deletions Source/WarpX.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -104,6 +104,9 @@ Real WarpX::moving_window_v = std::numeric_limits<amrex::Real>::max();

bool WarpX::fft_do_time_averaging = false;

bool WarpX::do_MCC_energy_tracking = false;
bool WarpX::do_MCC_ionization_tracking = false;

amrex::IntVect WarpX::m_fill_guards_fields = amrex::IntVect(0);
amrex::IntVect WarpX::m_fill_guards_current = amrex::IntVect(0);

Expand Down Expand Up @@ -1744,6 +1747,53 @@ WarpX::ReadParameters ()
}
}
}

// set flags for MCC collision tracking
// TODO: set an option for only tracking certain collisions
{
const ParmParse pp_collisions("collisions");

// Check if MCC tracking is turned off
int MCC_tracking = 1;
pp_collisions.query("mcc_tracking", MCC_tracking);

if (MCC_tracking) {
std::vector<std::string> collision_names;
pp_collisions.queryarr("collision_names", collision_names);

// Look through collision types for MCC
auto const ncollisions = collision_names.size();
for (int i = 0; i < static_cast<int>(ncollisions); ++i) {
const ParmParse pp_collision_name(collision_names[i]);

// For legacy, pairwisecoulomb is the default
std::string type = "pairwisecoulomb";

pp_collision_name.query("type", type);

// Check any MCC scattering processes
if (type == "background_mcc") {
std::vector<std::string> scattering_processes;
pp_collision_name.queryarr("scattering_processes", scattering_processes);

// Look through scattering processes for ionization
auto const nprocesses = scattering_processes.size();
for (int j = 0; j < static_cast<int>(nprocesses); ++j) {
// Set ionization tracking to true
if (scattering_processes[j] == "ionization") {
do_MCC_ionization_tracking = true;
}
}

// If there are other processes set energy tracking to true
if ((do_MCC_ionization_tracking && nprocesses > 1) ||
(!do_MCC_ionization_tracking && nprocesses > 0)) {
do_MCC_energy_tracking = true;
}
}
}
}
}
}

void
Expand Down Expand Up @@ -2225,6 +2275,25 @@ WarpX::AllocLevelMFs (int lev, const BoxArray& ba, const DistributionMapping& dm
m_fields.alloc_init(FieldType::current_fp, Direction{1}, lev, amrex::convert(ba, jy_nodal_flag), dm, ncomps, ngJ, 0.0_rt);
m_fields.alloc_init(FieldType::current_fp, Direction{2}, lev, amrex::convert(ba, jz_nodal_flag), dm, ncomps, ngJ, 0.0_rt);

// Allocate extra multifabs needed for MCC collision tracking
if (do_MCC_ionization_tracking || do_MCC_energy_tracking) {
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These multifabs could also be initialized in the constructor for BackgroundMCCCollision.

// 1 component per cell (either energy change or ionization)
int MCC_ncomps = 1;

// Zero guard cells
IntVect ngMCC = IntVect::TheZeroVector();

// Prepare to initialize a cell-centered MultiFab
IntVect MCC_nodal_flag = IntVect::TheCellVector();

if (do_MCC_energy_tracking) {
m_fields.alloc_init(FieldType::collision_energy_change, lev, amrex::convert(ba, MCC_nodal_flag), dm, MCC_ncomps, ngMCC, 0.0_rt);
}
if (do_MCC_ionization_tracking) {
m_fields.alloc_init(FieldType::collision_ionization, lev, amrex::convert(ba, MCC_nodal_flag), dm, MCC_ncomps, ngMCC, 0.0_rt);
}
}

if (do_current_centering)
{
amrex::BoxArray const& nodal_ba = amrex::convert(ba, amrex::IntVect::TheNodeVector());
Expand Down