aboutsummaryrefslogtreecommitdiff
path: root/Source/Particles
diff options
context:
space:
mode:
Diffstat (limited to 'Source/Particles')
-rwxr-xr-xSource/Particles/Deposition/ChargeDeposition.H2
-rw-r--r--Source/Particles/Deposition/CurrentDeposition.H8
-rw-r--r--Source/Particles/MultiParticleContainer.cpp20
-rw-r--r--Source/Particles/PhysicalParticleContainer.cpp110
-rw-r--r--Source/Particles/RigidInjectedParticleContainer.cpp43
-rw-r--r--Source/Particles/WarpXParticleContainer.H18
-rw-r--r--Source/Particles/WarpXParticleContainer.cpp23
7 files changed, 83 insertions, 141 deletions
diff --git a/Source/Particles/Deposition/ChargeDeposition.H b/Source/Particles/Deposition/ChargeDeposition.H
index 4253f5e76..f02eb1033 100755
--- a/Source/Particles/Deposition/ChargeDeposition.H
+++ b/Source/Particles/Deposition/ChargeDeposition.H
@@ -82,7 +82,7 @@ void doChargeDepositionShapeN(const amrex::Real * const xp,
const int k = compute_shape_factor<depos_order>(sz, z);
// Deposit charge into rho_arr
-#if (defined WARPX_DIM_2D) || (defined WARPX_DIM_RZ)
+#if (defined WARPX_DIM_XZ) || (defined WARPX_DIM_RZ)
for (int iz=0; iz<=depos_order; iz++){
for (int ix=0; ix<=depos_order; ix++){
amrex::Gpu::Atomic::Add(
diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H
index e8e7fab0a..b879970ef 100644
--- a/Source/Particles/Deposition/CurrentDeposition.H
+++ b/Source/Particles/Deposition/CurrentDeposition.H
@@ -133,7 +133,7 @@ void doDepositionShapeN(const amrex::Real * const xp,
const int l0 = compute_shape_factor<depos_order>(sz0, zmid-stagger_shift);
// Deposit current into jx_arr, jy_arr and jz_arr
-#if (defined WARPX_DIM_2D) || (defined WARPX_DIM_RZ)
+#if (defined WARPX_DIM_XZ) || (defined WARPX_DIM_RZ)
for (int iz=0; iz<=depos_order; iz++){
for (int ix=0; ix<=depos_order; ix++){
amrex::Gpu::Atomic::Add(
@@ -224,7 +224,7 @@ void doEsirkepovDepositionShapeN (const amrex::Real * const xp,
const amrex::Real invdtdx = 1.0/(dt*dx[1]*dx[2]);
const amrex::Real invdtdy = 1.0/(dt*dx[0]*dx[2]);
const amrex::Real invdtdz = 1.0/(dt*dx[0]*dx[1]);
-#elif (defined WARPX_DIM_2D) || (defined WARPX_DIM_RZ)
+#elif (defined WARPX_DIM_XZ) || (defined WARPX_DIM_RZ)
const amrex::Real invdtdx = 1.0/(dt*dx[2]);
const amrex::Real invdtdz = 1.0/(dt*dx[0]);
const amrex::Real invvol = 1.0/(dx[0]*dx[2]);
@@ -282,7 +282,7 @@ void doEsirkepovDepositionShapeN (const amrex::Real * const xp,
sintheta = 0.;
}
const amrex::Real vy = (-uxp[ip]*sintheta + uyp[ip]*costheta)*gaminv;
-#elif (defined WARPX_DIM_2D)
+#elif (defined WARPX_DIM_XZ)
const amrex::Real vy = uyp[ip]*gaminv;
#endif
@@ -357,7 +357,7 @@ void doEsirkepovDepositionShapeN (const amrex::Real * const xp,
}
}
-#elif (defined WARPX_DIM_2D) || (defined WARPX_DIM_RZ)
+#elif (defined WARPX_DIM_XZ) || (defined WARPX_DIM_RZ)
for (int k=dkl; k<=depos_order+2-dku; k++) {
amrex::Real sdxi = 0.;
diff --git a/Source/Particles/MultiParticleContainer.cpp b/Source/Particles/MultiParticleContainer.cpp
index 9e2cd097f..143f4b7f9 100644
--- a/Source/Particles/MultiParticleContainer.cpp
+++ b/Source/Particles/MultiParticleContainer.cpp
@@ -42,26 +42,6 @@ MultiParticleContainer::MultiParticleContainer (AmrCore* amr_core)
nspecies_boosted_frame_diags += 1;
}
}
-
- if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags)
- {
- for (int i = 0; i < nspecies_boosted_frame_diags; ++i)
- {
- int is = map_species_boosted_frame_diags[i];
- allcontainers[is]->AddRealComp("xold");
- allcontainers[is]->AddRealComp("yold");
- allcontainers[is]->AddRealComp("zold");
- allcontainers[is]->AddRealComp("uxold");
- allcontainers[is]->AddRealComp("uyold");
- allcontainers[is]->AddRealComp("uzold");
- }
- pc_tmp->AddRealComp("xold");
- pc_tmp->AddRealComp("yold");
- pc_tmp->AddRealComp("zold");
- pc_tmp->AddRealComp("uxold");
- pc_tmp->AddRealComp("uyold");
- pc_tmp->AddRealComp("uzold");
- }
}
void
diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp
index 9750e7e59..015482e3f 100644
--- a/Source/Particles/PhysicalParticleContainer.cpp
+++ b/Source/Particles/PhysicalParticleContainer.cpp
@@ -212,17 +212,8 @@ PhysicalParticleContainer::CheckAndAddParticle(Real x, Real y, Real z,
{
// need to create old values
auto& particle_tile = DefineAndReturnParticleTile(0, 0, 0);
- if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags)
- {
- particle_tile.push_back_real(particle_comps["xold"], x);
- particle_tile.push_back_real(particle_comps["yold"], y);
- particle_tile.push_back_real(particle_comps["zold"], z);
-
- particle_tile.push_back_real(particle_comps["uxold"], u[0]);
- particle_tile.push_back_real(particle_comps["uyold"], u[1]);
- particle_tile.push_back_real(particle_comps["uzold"], u[2]);
- }
}
+
// add particle
AddOneParticle(0, 0, 0, x, y, z, attribs);
}
@@ -301,11 +292,13 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox)
#ifdef _OPENMP
// First touch all tiles in the map in serial
for (MFIter mfi = MakeMFIter(lev); mfi.isValid(); ++mfi) {
- const int grid_id = mfi.index();
- const int tile_id = mfi.LocalTileIndex();
- GetParticles(lev)[std::make_pair(grid_id, tile_id)];
+ auto index = std::make_pair(mfi.index(), mfi.LocalTileIndex());
+ GetParticles(lev)[index];
+ tmp_particle_data.resize(finestLevel()+1);
+ // Create map entry if not there
+ tmp_particle_data[lev][index];
if ( (NumRuntimeRealComps()>0) || (NumRuntimeIntComps()>0) ) {
- DefineAndReturnParticleTile(lev, grid_id, tile_id);
+ DefineAndReturnParticleTile(lev, mfi.index(), mfi.LocalTileIndex());
}
}
#endif
@@ -405,14 +398,16 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox)
// then how many new particles will be injected is not that simple
// We have to shift fine_injection_box because overlap_box has been shifted.
Box fine_overlap_box = overlap_box & amrex::shift(fine_injection_box,shifted);
- max_new_particles += fine_overlap_box.numPts() * num_ppc
- * (AMREX_D_TERM(rrfac,*rrfac,*rrfac)-1);
- for (int icell = 0, ncells = overlap_box.numPts(); icell < ncells; ++icell) {
- IntVect iv = overlap_box.atOffset(icell);
- int r = (fine_overlap_box.contains(iv)) ? AMREX_D_TERM(rrfac,*rrfac,*rrfac) : 1;
- for (int ipart = 0; ipart < r; ++ipart) {
- cellid_v.push_back(icell);
- cellid_v.push_back(ipart);
+ if (fine_overlap_box.ok()) {
+ max_new_particles += fine_overlap_box.numPts() * num_ppc
+ * (AMREX_D_TERM(rrfac,*rrfac,*rrfac)-1);
+ for (int icell = 0, ncells = overlap_box.numPts(); icell < ncells; ++icell) {
+ IntVect iv = overlap_box.atOffset(icell);
+ int r = (fine_overlap_box.contains(iv)) ? AMREX_D_TERM(rrfac,*rrfac,*rrfac) : 1;
+ for (int ipart = 0; ipart < r; ++ipart) {
+ cellid_v.push_back(icell);
+ cellid_v.push_back(ipart);
+ }
}
}
}
@@ -430,12 +425,10 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox)
const int cpuid = ParallelDescriptor::MyProc();
auto& particle_tile = GetParticles(lev)[std::make_pair(grid_id,tile_id)];
- bool do_boosted = false;
+
if ( (NumRuntimeRealComps()>0) || (NumRuntimeIntComps()>0) ) {
DefineAndReturnParticleTile(lev, grid_id, tile_id);
}
- do_boosted = WarpX::do_boosted_frame_diagnostic
- && do_boosted_frame_diags;
auto old_size = particle_tile.GetArrayOfStructs().size();
auto new_size = old_size + max_new_particles;
@@ -447,15 +440,7 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox)
for (int ia = 0; ia < PIdx::nattribs; ++ia) {
pa[ia] = soa.GetRealData(ia).data() + old_size;
}
- GpuArray<Real*,6> pb;
- if (do_boosted) {
- pb[0] = soa.GetRealData(particle_comps[ "xold"]).data() + old_size;
- pb[1] = soa.GetRealData(particle_comps[ "yold"]).data() + old_size;
- pb[2] = soa.GetRealData(particle_comps[ "zold"]).data() + old_size;
- pb[3] = soa.GetRealData(particle_comps["uxold"]).data() + old_size;
- pb[4] = soa.GetRealData(particle_comps["uyold"]).data() + old_size;
- pb[5] = soa.GetRealData(particle_comps["uzold"]).data() + old_size;
- }
+
int* pi;
if (do_field_ionization) {
pi = soa.GetIntData(particle_icomps["ionization_level"]).data() + old_size;
@@ -617,15 +602,6 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox)
pa[PIdx::uy][ip] = u.y;
pa[PIdx::uz][ip] = u.z;
- if (do_boosted) {
- pb[0][ip] = x;
- pb[1][ip] = y;
- pb[2][ip] = z;
- pb[3][ip] = u.x;
- pb[4][ip] = u.y;
- pb[5][ip] = u.z;
- }
-
#if (AMREX_SPACEDIM == 3)
p.pos(0) = x;
p.pos(1) = y;
@@ -976,6 +952,19 @@ PhysicalParticleContainer::Evolve (int lev,
bool has_buffer = cEx || cjx;
+ if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags)
+ {
+ for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti)
+ {
+ const auto np = pti.numParticles();
+ const auto lev = pti.GetLevel();
+ const auto index = pti.GetPairIndex();
+ tmp_particle_data.resize(finestLevel()+1);
+ for (int i = 0; i < TmpIdx::nattribs; ++i)
+ tmp_particle_data[lev][index][i].resize(np);
+ }
+ }
+
#ifdef _OPENMP
#pragma omp parallel
#endif
@@ -1688,22 +1677,21 @@ PhysicalParticleContainer::PushP (int lev, Real dt,
void PhysicalParticleContainer::copy_attribs(WarpXParIter& pti,const Real* xp,
const Real* yp, const Real* zp)
{
-
auto& attribs = pti.GetAttribs();
-
Real* AMREX_RESTRICT uxp = attribs[PIdx::ux].dataPtr();
Real* AMREX_RESTRICT uyp = attribs[PIdx::uy].dataPtr();
Real* AMREX_RESTRICT uzp = attribs[PIdx::uz].dataPtr();
- Real* AMREX_RESTRICT xpold = pti.GetAttribs(particle_comps["xold"]).dataPtr();
- Real* AMREX_RESTRICT ypold = pti.GetAttribs(particle_comps["yold"]).dataPtr();
- Real* AMREX_RESTRICT zpold = pti.GetAttribs(particle_comps["zold"]).dataPtr();
- Real* AMREX_RESTRICT uxpold = pti.GetAttribs(particle_comps["uxold"]).dataPtr();
- Real* AMREX_RESTRICT uypold = pti.GetAttribs(particle_comps["uyold"]).dataPtr();
- Real* AMREX_RESTRICT uzpold = pti.GetAttribs(particle_comps["uzold"]).dataPtr();
-
- const long np = pti.numParticles();
-
+ const auto np = pti.numParticles();
+ const auto lev = pti.GetLevel();
+ const auto index = pti.GetPairIndex();
+ Real* AMREX_RESTRICT xpold = tmp_particle_data[lev][index][TmpIdx::xold ].dataPtr();
+ Real* AMREX_RESTRICT ypold = tmp_particle_data[lev][index][TmpIdx::yold ].dataPtr();
+ Real* AMREX_RESTRICT zpold = tmp_particle_data[lev][index][TmpIdx::zold ].dataPtr();
+ Real* AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][TmpIdx::uxold].dataPtr();
+ Real* AMREX_RESTRICT uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr();
+ Real* AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr();
+
ParallelFor( np,
[=] AMREX_GPU_DEVICE (long i) {
xpold[i]=xp[i];
@@ -1788,15 +1776,15 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real
auto& uyp_new = attribs[PIdx::uy ];
auto& uzp_new = attribs[PIdx::uz ];
- auto& xp_old = pti.GetAttribs(particle_comps["xold"]);
- auto& yp_old = pti.GetAttribs(particle_comps["yold"]);
- auto& zp_old = pti.GetAttribs(particle_comps["zold"]);
- auto& uxp_old = pti.GetAttribs(particle_comps["uxold"]);
- auto& uyp_old = pti.GetAttribs(particle_comps["uyold"]);
- auto& uzp_old = pti.GetAttribs(particle_comps["uzold"]);
+ auto& xp_old = tmp_particle_data[lev][index][TmpIdx::xold];
+ auto& yp_old = tmp_particle_data[lev][index][TmpIdx::yold];
+ auto& zp_old = tmp_particle_data[lev][index][TmpIdx::zold];
+ auto& uxp_old = tmp_particle_data[lev][index][TmpIdx::uxold];
+ auto& uyp_old = tmp_particle_data[lev][index][TmpIdx::uyold];
+ auto& uzp_old = tmp_particle_data[lev][index][TmpIdx::uzold];
const long np = pti.numParticles();
-
+
Real uzfrm = -WarpX::gamma_boost*WarpX::beta_boost*PhysConst::c;
Real inv_c2 = 1.0/PhysConst::c/PhysConst::c;
diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp
index 36cb9d224..0c94d1e33 100644
--- a/Source/Particles/RigidInjectedParticleContainer.cpp
+++ b/Source/Particles/RigidInjectedParticleContainer.cpp
@@ -239,15 +239,13 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti,
Real* const AMREX_RESTRICT Bzp = attribs[PIdx::Bz].dataPtr();
if (!done_injecting_lev) {
- if (!(WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags)) {
- // If the old values are not already saved, create copies here.
- xp_save = xp;
- yp_save = yp;
- zp_save = zp;
- uxp_save = uxp;
- uyp_save = uyp;
- uzp_save = uzp;
- }
+ // If the old values are not already saved, create copies here.
+ xp_save = xp;
+ yp_save = yp;
+ zp_save = zp;
+ uxp_save = uxp;
+ uyp_save = uyp;
+ uzp_save = uzp;
// Scale the fields of particles about to cross the injection plane.
// This only approximates what should be happening. The particles
@@ -275,27 +273,12 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti,
if (!done_injecting_lev) {
- Real* AMREX_RESTRICT x_save;
- Real* AMREX_RESTRICT y_save;
- Real* AMREX_RESTRICT z_save;
- Real* AMREX_RESTRICT ux_save;
- Real* AMREX_RESTRICT uy_save;
- Real* AMREX_RESTRICT uz_save;
- if (!(WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags)) {
- x_save = xp_save.dataPtr();
- y_save = yp_save.dataPtr();
- z_save = zp_save.dataPtr();
- ux_save = uxp_save.dataPtr();
- uy_save = uyp_save.dataPtr();
- uz_save = uzp_save.dataPtr();
- } else {
- x_save = pti.GetAttribs(particle_comps["xold"]).dataPtr();
- y_save = pti.GetAttribs(particle_comps["yold"]).dataPtr();
- z_save = pti.GetAttribs(particle_comps["zold"]).dataPtr();
- ux_save = pti.GetAttribs(particle_comps["uxold"]).dataPtr();
- uy_save = pti.GetAttribs(particle_comps["uyold"]).dataPtr();
- uz_save = pti.GetAttribs(particle_comps["uzold"]).dataPtr();
- }
+ Real* AMREX_RESTRICT x_save = xp_save.dataPtr();
+ Real* AMREX_RESTRICT y_save = yp_save.dataPtr();
+ Real* AMREX_RESTRICT z_save = zp_save.dataPtr();
+ Real* AMREX_RESTRICT ux_save = uxp_save.dataPtr();
+ Real* AMREX_RESTRICT uy_save = uyp_save.dataPtr();
+ Real* AMREX_RESTRICT uz_save = uzp_save.dataPtr();
// Undo the push for particles not injected yet.
// The zp are advanced a fixed amount.
diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H
index cc553f9a7..7393f7301 100644
--- a/Source/Particles/WarpXParticleContainer.H
+++ b/Source/Particles/WarpXParticleContainer.H
@@ -29,6 +29,15 @@ struct DiagIdx
};
};
+struct TmpIdx
+{
+ enum {
+ xold = 0,
+ yold, zold, uxold, uyold, uzold,
+ nattribs
+ };
+};
+
namespace ParticleStringNames
{
const std::map<std::string, int> to_index = {
@@ -308,7 +317,10 @@ protected:
amrex::Vector<amrex::FArrayBox> local_jy;
amrex::Vector<amrex::FArrayBox> local_jz;
- amrex::Vector<amrex::Cuda::ManagedDeviceVector<amrex::Real> > m_xp, m_yp, m_zp, m_giv;
+ using DataContainer = amrex::Gpu::ManagedDeviceVector<amrex::Real>;
+ using PairIndex = std::pair<int, int>;
+
+ amrex::Vector<DataContainer> m_xp, m_yp, m_zp, m_giv;
// Whether to dump particle quantities.
// If true, particle position is always dumped.
@@ -318,7 +330,9 @@ protected:
amrex::Vector<int> plot_flags;
// list of names of attributes to dump.
amrex::Vector<std::string> plot_vars;
-
+
+ amrex::Vector<std::map<PairIndex, std::array<DataContainer, TmpIdx::nattribs> > > tmp_particle_data;
+
private:
virtual void particlePostLocate(ParticleType& p, const amrex::ParticleLocData& pld,
const int lev) override;
diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp
index 646644f94..80d1caa0f 100644
--- a/Source/Particles/WarpXParticleContainer.cpp
+++ b/Source/Particles/WarpXParticleContainer.cpp
@@ -85,17 +85,6 @@ WarpXParticleContainer::WarpXParticleContainer (AmrCore* amr_core, int ispecies)
particle_comps["theta"] = PIdx::theta;
#endif
- if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags)
- {
- particle_comps["xold"] = PIdx::nattribs;
- particle_comps["yold"] = PIdx::nattribs+1;
- particle_comps["zold"] = PIdx::nattribs+2;
- particle_comps["uxold"] = PIdx::nattribs+3;
- particle_comps["uyold"] = PIdx::nattribs+4;
- particle_comps["uzold"] = PIdx::nattribs+5;
-
- }
-
// Initialize temporary local arrays for charge/current deposition
int num_threads = 1;
#ifdef _OPENMP
@@ -240,12 +229,6 @@ WarpXParticleContainer::AddNParticles (int lev,
if ( (NumRuntimeRealComps()>0) || (NumRuntimeIntComps()>0) ){
auto& ptile = DefineAndReturnParticleTile(0, 0, 0);
- if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags)
- {
- ptile.push_back_real(particle_comps["xold"], x[i]);
- ptile.push_back_real(particle_comps["yold"], y[i]);
- ptile.push_back_real(particle_comps["zold"], z[i]);
- }
}
particle_tile.push_back(p);
@@ -260,12 +243,6 @@ WarpXParticleContainer::AddNParticles (int lev,
if ( (NumRuntimeRealComps()>0) || (NumRuntimeIntComps()>0) ){
auto& ptile = DefineAndReturnParticleTile(0, 0, 0);
- if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags)
- {
- ptile.push_back_real(particle_comps["uxold"], vx + ibegin, vx + iend);
- ptile.push_back_real(particle_comps["uyold"], vy + ibegin, vy + iend);
- ptile.push_back_real(particle_comps["uzold"], vz + ibegin, vz + iend);
- }
}
for (int comp = PIdx::uz+1; comp < PIdx::nattribs; ++comp)
generated by cgit v1.2.3 (git 2.46.0) at 2025-08-19 17:00:43 +0000