From eb4080b05342edd63c09091f393638c9053b7b19 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 21 Aug 2019 11:31:10 -0700 Subject: reimplement the boosted frame diagnostics without the runtime components --- Source/Particles/PhysicalParticleContainer.cpp | 65 +++++++------------------- 1 file changed, 18 insertions(+), 47 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index d10390204..14807d0d9 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -196,18 +196,6 @@ PhysicalParticleContainer::CheckAndAddParticle(Real x, Real y, Real z, attribs[PIdx::uz] = u[2]; attribs[PIdx::w ] = weight; - if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags) - { - // need to create old values - auto& particle_tile = DefineAndReturnParticleTile(0, 0, 0); - 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); } @@ -430,15 +418,6 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) for (int ia = 0; ia < PIdx::nattribs; ++ia) { pa[ia] = soa.GetRealData(ia).data() + old_size; } - GpuArray 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; - } const GpuArray overlap_corner {AMREX_D_DECL(overlap_realbox.lo(0), @@ -589,15 +568,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; @@ -1633,21 +1603,22 @@ 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(); + tmp_particle_data.resize(finestLevel()+1); + for (int i = 0; i < 6; ++i) tmp_particle_data[lev][index][i].resize(np); + Real* AMREX_RESTRICT xpold = tmp_particle_data[lev][index][0].dataPtr(); + Real* AMREX_RESTRICT ypold = tmp_particle_data[lev][index][1].dataPtr(); + Real* AMREX_RESTRICT zpold = tmp_particle_data[lev][index][2].dataPtr(); + Real* AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][3].dataPtr(); + Real* AMREX_RESTRICT uypold = tmp_particle_data[lev][index][4].dataPtr(); + Real* AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][5].dataPtr(); ParallelFor( np, [=] AMREX_GPU_DEVICE (long i) { @@ -1733,15 +1704,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][0]; + auto& yp_old = tmp_particle_data[lev][index][1]; + auto& zp_old = tmp_particle_data[lev][index][2]; + auto& uxp_old = tmp_particle_data[lev][index][3]; + auto& uyp_old = tmp_particle_data[lev][index][4]; + auto& uzp_old = tmp_particle_data[lev][index][5]; const long np = pti.numParticles(); - + Real uzfrm = -WarpX::gamma_boost*WarpX::beta_boost*PhysConst::c; Real inv_c2 = 1.0/PhysConst::c/PhysConst::c; -- cgit v1.2.3 From d2a176f231788f1a04472f09bbe7134a09d14ac5 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 21 Aug 2019 11:43:28 -0700 Subject: some typedefs and an enum for clarity --- Source/Particles/PhysicalParticleContainer.cpp | 34 ++++++++++---------------- Source/Particles/WarpXParticleContainer.H | 18 +++++++++++--- 2 files changed, 28 insertions(+), 24 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 14807d0d9..c6d3cd9e7 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -277,9 +277,6 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) const int grid_id = mfi.index(); const int tile_id = mfi.LocalTileIndex(); GetParticles(lev)[std::make_pair(grid_id, tile_id)]; - if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags) { - DefineAndReturnParticleTile(lev, grid_id, tile_id); - } } #endif @@ -403,11 +400,6 @@ 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 (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags) { - do_boosted = true; - DefineAndReturnParticleTile(lev, grid_id, tile_id); - } auto old_size = particle_tile.GetArrayOfStructs().size(); auto new_size = old_size + max_new_particles; particle_tile.resize(new_size); @@ -1612,13 +1604,13 @@ void PhysicalParticleContainer::copy_attribs(WarpXParIter& pti,const Real* xp, const auto lev = pti.GetLevel(); const auto index = pti.GetPairIndex(); tmp_particle_data.resize(finestLevel()+1); - for (int i = 0; i < 6; ++i) tmp_particle_data[lev][index][i].resize(np); - Real* AMREX_RESTRICT xpold = tmp_particle_data[lev][index][0].dataPtr(); - Real* AMREX_RESTRICT ypold = tmp_particle_data[lev][index][1].dataPtr(); - Real* AMREX_RESTRICT zpold = tmp_particle_data[lev][index][2].dataPtr(); - Real* AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][3].dataPtr(); - Real* AMREX_RESTRICT uypold = tmp_particle_data[lev][index][4].dataPtr(); - Real* AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][5].dataPtr(); + for (int i = 0; i < TmpIdx::nattribs; ++i) tmp_particle_data[lev][index][i].resize(np); + 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) { @@ -1704,12 +1696,12 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real auto& uyp_new = attribs[PIdx::uy ]; auto& uzp_new = attribs[PIdx::uz ]; - auto& xp_old = tmp_particle_data[lev][index][0]; - auto& yp_old = tmp_particle_data[lev][index][1]; - auto& zp_old = tmp_particle_data[lev][index][2]; - auto& uxp_old = tmp_particle_data[lev][index][3]; - auto& uyp_old = tmp_particle_data[lev][index][4]; - auto& uzp_old = tmp_particle_data[lev][index][5]; + 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(); diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index cbdf0cdbc..a21506ec3 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 to_index = { @@ -297,7 +306,10 @@ protected: amrex::Vector local_jy; amrex::Vector local_jz; - amrex::Vector > m_xp, m_yp, m_zp, m_giv; + using DataContainer = amrex::Gpu::ManagedDeviceVector; + using PairIndex = std::pair; + + amrex::Vector m_xp, m_yp, m_zp, m_giv; // Whether to dump particle quantities. // If true, particle position is always dumped. @@ -307,8 +319,8 @@ protected: amrex::Vector plot_flags; // list of names of attributes to dump. amrex::Vector plot_vars; - - amrex::Vector, std::array, 6> > > tmp_particle_data; + + amrex::Vector > > tmp_particle_data; private: virtual void particlePostLocate(ParticleType& p, const amrex::ParticleLocData& pld, -- cgit v1.2.3 From d5fc8e40d0373bd4ba1e306f4ce5dc0a895ba6e7 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 21 Aug 2019 12:30:35 -0700 Subject: touch all map entries for tmp_particle_data --- Source/Particles/PhysicalParticleContainer.cpp | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index c6d3cd9e7..95909feba 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -274,9 +274,10 @@ 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); + tmp_particle_data[lev][index]; } #endif @@ -1611,7 +1612,7 @@ void PhysicalParticleContainer::copy_attribs(WarpXParIter& pti,const Real* xp, 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]; -- cgit v1.2.3 From 11d362d949107d9c51a6a9771b05d62bc8e1cd8e Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Mon, 26 Aug 2019 10:26:06 -0700 Subject: fix race condition --- Source/Particles/PhysicalParticleContainer.cpp | 15 +++++++++++++-- 1 file changed, 13 insertions(+), 2 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 79a6f88f1..df8dcf836 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -949,6 +949,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 @@ -1680,8 +1693,6 @@ void PhysicalParticleContainer::copy_attribs(WarpXParIter& pti,const Real* xp, 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); 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(); -- cgit v1.2.3 From 49dc1a600e30443c6d946689bfa80d3cd1091b1f Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Mon, 26 Aug 2019 11:06:56 -0700 Subject: add comment --- Source/Particles/PhysicalParticleContainer.cpp | 1 + 1 file changed, 1 insertion(+) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index df8dcf836..9169e6c1f 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -295,6 +295,7 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) 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, mfi.index(), mfi.LocalTileIndex()); -- cgit v1.2.3