From 60917ab3d92cf5f325468ffd15a9c35b101699a6 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Thu, 19 Sep 2019 13:22:24 -0700 Subject: do not store old attribs in second particle push when subcycling --- Source/Particles/PhysicalParticleContainer.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 4bc0ee16e..2d0b90902 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -942,7 +942,7 @@ PhysicalParticleContainer::Evolve (int lev, MultiFab* rho, MultiFab* crho, const MultiFab* cEx, const MultiFab* cEy, const MultiFab* cEz, const MultiFab* cBx, const MultiFab* cBy, const MultiFab* cBz, - Real t, Real dt) + Real t, Real dt, DtType a_dt_type) { BL_PROFILE("PPC::Evolve()"); BL_PROFILE_VAR_NS("PPC::Evolve::Copy", blp_copy); @@ -1301,7 +1301,7 @@ PhysicalParticleContainer::Evolve (int lev, // Particle Push // BL_PROFILE_VAR_START(blp_ppc_pp); - PushPX(pti, m_xp[thread_num], m_yp[thread_num], m_zp[thread_num], dt); + PushPX(pti, m_xp[thread_num], m_yp[thread_num], m_zp[thread_num], dt, a_dt_type); BL_PROFILE_VAR_STOP(blp_ppc_pp); // @@ -1526,7 +1526,7 @@ PhysicalParticleContainer::PushPX(WarpXParIter& pti, Cuda::ManagedDeviceVector& xp, Cuda::ManagedDeviceVector& yp, Cuda::ManagedDeviceVector& zp, - Real dt) + Real dt, DtType a_dt_type) { // This wraps the momentum and position advance so that inheritors can modify the call. @@ -1545,8 +1545,10 @@ PhysicalParticleContainer::PushPX(WarpXParIter& pti, const Real* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); const Real* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); - if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags) + Print()<<"a_dt_type "< Date: Thu, 19 Sep 2019 13:33:32 -0700 Subject: cleaning --- Source/Particles/PhysicalParticleContainer.cpp | 2 -- 1 file changed, 2 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 2d0b90902..8a30b3acb 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1545,10 +1545,8 @@ PhysicalParticleContainer::PushPX(WarpXParIter& pti, const Real* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); const Real* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); - Print()<<"a_dt_type "< Date: Fri, 20 Sep 2019 20:47:20 -0700 Subject: fix particle splitting --- Source/Particles/PhysicalParticleContainer.cpp | 11 ++++++----- Source/Particles/WarpXParticleContainer.H | 2 +- 2 files changed, 7 insertions(+), 6 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 8a30b3acb..c3d6a35fd 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1446,15 +1446,16 @@ PhysicalParticleContainer::SplitParticles(int lev) } #elif (AMREX_SPACEDIM==3) if (split_type==0){ - // Split particle in two along each axis - // 6 particles in 2d + Print()<<"split_type==0\n"; + // Split particle in two along each diagonals + // 8 particles in 3d for (int ishift = -1; ishift < 2; ishift +=2 ){ for (int jshift = -1; jshift < 2; jshift +=2 ){ for (int kshift = -1; kshift < 2; kshift +=2 ){ // Add one particle with offset in x, y and z psplit_x.push_back( xp[i] + ishift*dx[0]/2 ); psplit_y.push_back( yp[i] + jshift*dx[1]/2 ); - psplit_z.push_back( zp[i] + jshift*dx[2]/2 ); + psplit_z.push_back( zp[i] + kshift*dx[2]/2 ); psplit_ux.push_back( uxp[i] ); psplit_uy.push_back( uyp[i] ); psplit_uz.push_back( uzp[i] ); @@ -1463,8 +1464,8 @@ PhysicalParticleContainer::SplitParticles(int lev) } } } else { - // Split particle in two along each diagonal - // 8 particles in 3d + // Split particle in two along each axis + // 6 particles in 3d for (int ishift = -1; ishift < 2; ishift +=2 ){ // Add one particle with offset in x psplit_x.push_back( xp[i] + ishift*dx[0]/2 ); diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index 877b42f10..ee75bc511 100644 --- a/Source/Particles/WarpXParticleContainer.H +++ b/Source/Particles/WarpXParticleContainer.H @@ -258,7 +258,7 @@ public: bool do_splitting = false; - // split along axes (0) or diagonals (1) + // split along diagonals (0) or axes (1) int split_type = 0; using amrex::ParticleContainer<0, 0, PIdx::nattribs>::AddRealComp; -- cgit v1.2.3 From 897d2fb86dae052903501afb75e2f5ca7c2bb3fc Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Fri, 20 Sep 2019 20:55:23 -0700 Subject: cleaning --- Source/Particles/PhysicalParticleContainer.cpp | 1 - 1 file changed, 1 deletion(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index c3d6a35fd..d73874650 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1446,7 +1446,6 @@ PhysicalParticleContainer::SplitParticles(int lev) } #elif (AMREX_SPACEDIM==3) if (split_type==0){ - Print()<<"split_type==0\n"; // Split particle in two along each diagonals // 8 particles in 3d for (int ishift = -1; ishift < 2; ishift +=2 ){ -- cgit v1.2.3 From 05fc79ae166b44930a865af348c5ec765694a1c0 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Fri, 20 Sep 2019 20:58:57 -0700 Subject: typo --- Source/Particles/PhysicalParticleContainer.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index d73874650..d5f133332 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1408,7 +1408,7 @@ PhysicalParticleContainer::SplitParticles(int lev) np_split_to_add += np_split; #if (AMREX_SPACEDIM==2) if (split_type==0){ - // Split particle in two along each axis + // Split particle in two along each diagonals // 4 particles in 2d for (int ishift = -1; ishift < 2; ishift +=2 ){ for (int kshift = -1; kshift < 2; kshift +=2 ){ @@ -1423,7 +1423,7 @@ PhysicalParticleContainer::SplitParticles(int lev) } } } else { - // Split particle in two along each diagonal + // Split particle in two along each axis // 4 particles in 2d for (int ishift = -1; ishift < 2; ishift +=2 ){ // Add one particle with offset in x -- cgit v1.2.3 From 5f3db4cd214b27297e3ea69c8582e64bf190d3e6 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Mon, 23 Sep 2019 13:10:31 -0700 Subject: split particles only during second half-push when subcycling --- Source/Particles/PhysicalParticleContainer.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index d5f133332..c878ad7b3 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1364,7 +1364,9 @@ PhysicalParticleContainer::Evolve (int lev, } } // Split particles - if (do_splitting){ SplitParticles(lev); } + if (do_splitting && a_dt_type == DtType::SecondHalf){ + SplitParticles(lev); + } } // Loop over all particles in the particle container and -- cgit v1.2.3 From 97782d2ee0929b1485923e8dc98ccf0192031649 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Mon, 23 Sep 2019 14:10:38 -0700 Subject: make sure splitting works without subcycling --- Source/Particles/PhysicalParticleContainer.cpp | 2 +- Source/WarpX.H | 4 ++-- Source/WarpX.cpp | 2 ++ 3 files changed, 5 insertions(+), 3 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index c878ad7b3..8eec9ae41 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1364,7 +1364,7 @@ PhysicalParticleContainer::Evolve (int lev, } } // Split particles - if (do_splitting && a_dt_type == DtType::SecondHalf){ + if (do_splitting && a_dt_type == DtType::SecondHalf && WarpX::do_subcycling == 1){ SplitParticles(lev); } } diff --git a/Source/WarpX.H b/Source/WarpX.H index d4e1e6813..7a5344f4f 100644 --- a/Source/WarpX.H +++ b/Source/WarpX.H @@ -118,6 +118,8 @@ public: static int sort_int; + static int do_subcycling; + // buffers static int n_field_gather_buffer; //! in number of cells from the edge (identical for each dimension) static int n_current_deposition_buffer; //! in number of cells from the edge (identical for each dimension) @@ -527,8 +529,6 @@ private: // Other runtime parameters int verbose = 1; - int do_subcycling = 0; - int max_step = std::numeric_limits::max(); amrex::Real stop_time = std::numeric_limits::max(); diff --git a/Source/WarpX.cpp b/Source/WarpX.cpp index 6c6bb8107..a844e7aa0 100644 --- a/Source/WarpX.cpp +++ b/Source/WarpX.cpp @@ -69,6 +69,8 @@ bool WarpX::do_boosted_frame_particles = true; bool WarpX::do_dynamic_scheduling = true; +int WarpX::do_subcycling = 0; + #if (AMREX_SPACEDIM == 3) IntVect WarpX::Bx_nodal_flag(1,0,0); IntVect WarpX::By_nodal_flag(0,1,0); -- cgit v1.2.3 From b4c4315c77ade3ebb4b11e137ddfe354f1c95d23 Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Mon, 23 Sep 2019 16:24:57 -0700 Subject: Create dedicated method for partition in buffers --- Source/Particles/Make.package | 1 + Source/Particles/PhysicalParticleContainer.H | 13 +++ Source/Particles/PhysicalParticleContainer.cpp | 100 +-------------------- Source/Particles/Sorting/Make.package | 3 + Source/Particles/Sorting/Partition.cpp | 115 +++++++++++++++++++++++++ 5 files changed, 135 insertions(+), 97 deletions(-) create mode 100644 Source/Particles/Sorting/Make.package create mode 100644 Source/Particles/Sorting/Partition.cpp (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/Make.package b/Source/Particles/Make.package index 95f36cf65..c9d520873 100644 --- a/Source/Particles/Make.package +++ b/Source/Particles/Make.package @@ -18,6 +18,7 @@ CEXE_headers += ShapeFactors.H include $(WARPX_HOME)/Source/Particles/Pusher/Make.package include $(WARPX_HOME)/Source/Particles/Deposition/Make.package include $(WARPX_HOME)/Source/Particles/Gather/Make.package +include $(WARPX_HOME)/Source/Particles/Sorting/Make.package INCLUDE_LOCATIONS += $(WARPX_HOME)/Source/Particles VPATH_LOCATIONS += $(WARPX_HOME)/Source/Particles diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index ace1ec7f8..b809acc45 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -100,6 +100,19 @@ public: const amrex::MultiFab& By, const amrex::MultiFab& Bz) override; + void PartitionParticlesInBuffers( + long& nfine_current, + long& nfine_gather, + long const np, + WarpXParIter& pti, + int const lev, + amrex::iMultiFab const* current_masks, + amrex::iMultiFab const* gather_masks, + RealVector& uxp, + RealVector& uyp, + RealVector& uzp, + RealVector& wp ); + void copy_attribs(WarpXParIter& pti,const amrex::Real* xp, const amrex::Real* yp, const amrex::Real* zp); diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index d5f133332..e1046ac7a 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -948,7 +948,6 @@ PhysicalParticleContainer::Evolve (int lev, BL_PROFILE_VAR_NS("PPC::Evolve::Copy", blp_copy); BL_PROFILE_VAR_NS("PPC::FieldGather", blp_fg); BL_PROFILE_VAR_NS("PPC::ParticlePush", blp_ppc_pp); - BL_PROFILE_VAR_NS("PPC::Evolve::partition", blp_partition); const std::array& dx = WarpX::CellSize(lev); const std::array& cdx = WarpX::CellSize(std::max(lev-1,0)); @@ -960,7 +959,6 @@ PhysicalParticleContainer::Evolve (int lev, BL_ASSERT(OnSameGrids(lev,jx)); MultiFab* cost = WarpX::getCosts(lev); - const iMultiFab* current_masks = WarpX::CurrentBufferMasks(lev); const iMultiFab* gather_masks = WarpX::GatherBufferMasks(lev); @@ -991,10 +989,6 @@ PhysicalParticleContainer::Evolve (int lev, FArrayBox filtered_Ex, filtered_Ey, filtered_Ez; FArrayBox filtered_Bx, filtered_By, filtered_Bz; - std::vector inexflag; - Vector pid; - RealVector tmp; - ParticleVector particle_tmp; for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { @@ -1087,97 +1081,9 @@ PhysicalParticleContainer::Evolve (int lev, long nfine_current = np; //! number of particles depositing to fine grid long nfine_gather = np; //! number of particles gathering from fine grid - if (has_buffer && !do_not_push) - { - BL_PROFILE_VAR_START(blp_partition); - inexflag.resize(np); - auto& aos = pti.GetArrayOfStructs(); - // We need to partition the large buffer first - iMultiFab const* bmasks = (WarpX::n_field_gather_buffer >= WarpX::n_current_deposition_buffer) ? - gather_masks : current_masks; - int i = 0; - const auto& msk = (*bmasks)[pti]; - for (const auto& p : aos) { - const IntVect& iv = Index(p, lev); - inexflag[i++] = msk(iv); - } - - pid.resize(np); - std::iota(pid.begin(), pid.end(), 0L); - - auto sep = std::stable_partition(pid.begin(), pid.end(), - [&inexflag](long id) { return inexflag[id]; }); - - if (WarpX::n_current_deposition_buffer == WarpX::n_field_gather_buffer) { - nfine_current = nfine_gather = std::distance(pid.begin(), sep); - } else if (sep != pid.end()) { - int n_buf; - if (bmasks == gather_masks) { - nfine_gather = std::distance(pid.begin(), sep); - bmasks = current_masks; - n_buf = WarpX::n_current_deposition_buffer; - } else { - nfine_current = std::distance(pid.begin(), sep); - bmasks = gather_masks; - n_buf = WarpX::n_field_gather_buffer; - } - if (n_buf > 0) - { - const auto& msk2 = (*bmasks)[pti]; - for (auto it = sep; it != pid.end(); ++it) { - const long id = *it; - const IntVect& iv = Index(aos[id], lev); - inexflag[id] = msk2(iv); - } - - auto sep2 = std::stable_partition(sep, pid.end(), - [&inexflag](long id) {return inexflag[id];}); - if (bmasks == gather_masks) { - nfine_gather = std::distance(pid.begin(), sep2); - } else { - nfine_current = std::distance(pid.begin(), sep2); - } - } - } - - // only deposit / gather to coarsest grid - if (m_deposit_on_main_grid && lev > 0) { - nfine_current = 0; - } - if (m_gather_from_main_grid && lev > 0) { - nfine_gather = 0; - } - - if (nfine_current != np || nfine_gather != np) - { - particle_tmp.resize(np); - for (long ip = 0; ip < np; ++ip) { - particle_tmp[ip] = aos[pid[ip]]; - } - std::swap(aos(), particle_tmp); - - tmp.resize(np); - for (long ip = 0; ip < np; ++ip) { - tmp[ip] = wp[pid[ip]]; - } - std::swap(wp, tmp); - - for (long ip = 0; ip < np; ++ip) { - tmp[ip] = uxp[pid[ip]]; - } - std::swap(uxp, tmp); - - for (long ip = 0; ip < np; ++ip) { - tmp[ip] = uyp[pid[ip]]; - } - std::swap(uyp, tmp); - - for (long ip = 0; ip < np; ++ip) { - tmp[ip] = uzp[pid[ip]]; - } - std::swap(uzp, tmp); - } - BL_PROFILE_VAR_STOP(blp_partition); + if (has_buffer && !do_not_push) { + PartitionParticlesInBuffers( nfine_current, nfine_gather, np, + pti, lev, current_masks, gather_masks, uxp, uyp, uzp, wp); } const long np_current = (cjx) ? nfine_current : np; diff --git a/Source/Particles/Sorting/Make.package b/Source/Particles/Sorting/Make.package new file mode 100644 index 000000000..750d2607e --- /dev/null +++ b/Source/Particles/Sorting/Make.package @@ -0,0 +1,3 @@ +CEXE_sources += Partition.cpp +INCLUDE_LOCATIONS += $(WARPX_HOME)/Source/Particles/Sorting +VPATH_LOCATIONS += $(WARPX_HOME)/Source/Particles/Sorting diff --git a/Source/Particles/Sorting/Partition.cpp b/Source/Particles/Sorting/Partition.cpp new file mode 100644 index 000000000..d000a2c04 --- /dev/null +++ b/Source/Particles/Sorting/Partition.cpp @@ -0,0 +1,115 @@ +#include +#include +#include + +using namespace amrex; + +/* \brief +*/ +void +PhysicalParticleContainer::PartitionParticlesInBuffers( + long& nfine_current, long& nfine_gather, long const np, + WarpXParIter& pti, int const lev, + iMultiFab const* current_masks, + iMultiFab const* gather_masks, + RealVector& uxp, RealVector& uyp, RealVector& uzp, RealVector& wp) +{ + BL_PROFILE("PPC::Evolve::partition"); + + std::vector inexflag; + inexflag.resize(np); + + auto& aos = pti.GetArrayOfStructs(); + + // We need to partition the large buffer first + iMultiFab const* bmasks = + (WarpX::n_field_gather_buffer >= WarpX::n_current_deposition_buffer) ? + gather_masks : current_masks; + int i = 0; + const auto& msk = (*bmasks)[pti]; + for (const auto& p : aos) { + const IntVect& iv = Index(p, lev); + inexflag[i++] = msk(iv); + } + + Vector pid; + pid.resize(np); + std::iota(pid.begin(), pid.end(), 0L); + auto sep = std::stable_partition(pid.begin(), pid.end(), + [&inexflag](long id) { return inexflag[id]; }); + + if (WarpX::n_current_deposition_buffer == WarpX::n_field_gather_buffer) { + nfine_current = nfine_gather = std::distance(pid.begin(), sep); + } else if (sep != pid.end()) { + int n_buf; + if (bmasks == gather_masks) { + nfine_gather = std::distance(pid.begin(), sep); + bmasks = current_masks; + n_buf = WarpX::n_current_deposition_buffer; + } else { + nfine_current = std::distance(pid.begin(), sep); + bmasks = gather_masks; + n_buf = WarpX::n_field_gather_buffer; + } + if (n_buf > 0) + { + const auto& msk2 = (*bmasks)[pti]; + for (auto it = sep; it != pid.end(); ++it) { + const long id = *it; + const IntVect& iv = Index(aos[id], lev); + inexflag[id] = msk2(iv); + } + + auto sep2 = std::stable_partition(sep, pid.end(), + [&inexflag](long id) {return inexflag[id];}); + if (bmasks == gather_masks) { + nfine_gather = std::distance(pid.begin(), sep2); + } else { + nfine_current = std::distance(pid.begin(), sep2); + } + } + } + + // only deposit / gather to coarsest grid + if (m_deposit_on_main_grid && lev > 0) { + nfine_current = 0; + } + if (m_gather_from_main_grid && lev > 0) { + nfine_gather = 0; + } + + if (nfine_current != np || nfine_gather != np) + { + RealVector tmp; + ParticleVector particle_tmp; + + particle_tmp.resize(np); + for (long ip = 0; ip < np; ++ip) { + particle_tmp[ip] = aos[pid[ip]]; + } + std::swap(aos(), particle_tmp); + + tmp.resize(np); + + for (long ip = 0; ip < np; ++ip) { + tmp[ip] = wp[pid[ip]]; + } + std::swap(wp, tmp); + + for (long ip = 0; ip < np; ++ip) { + tmp[ip] = uxp[pid[ip]]; + } + std::swap(uxp, tmp); + + for (long ip = 0; ip < np; ++ip) { + tmp[ip] = uyp[pid[ip]]; + } + std::swap(uyp, tmp); + + for (long ip = 0; ip < np; ++ip) { + tmp[ip] = uzp[pid[ip]]; + } + std::swap(uzp, tmp); + } + +} -- cgit v1.2.3 From 3b44103b7cd689bb30e5e379f67c8a1c8ad9c618 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Mon, 23 Sep 2019 16:48:45 -0700 Subject: fix and add comments --- Source/Particles/PhysicalParticleContainer.cpp | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 8eec9ae41..8dfddc833 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1363,8 +1363,14 @@ PhysicalParticleContainer::Evolve (int lev, } } } - // Split particles - if (do_splitting && a_dt_type == DtType::SecondHalf && WarpX::do_subcycling == 1){ + // Split particles at the end of the timestep. + // When subcycling is ON, the splitting is done on the last call to + // PhysicalParticleContainer::Evolve on the finest level, i.e., at the + // end of the large timestep. Otherwise, the pushes on different levels + // are not consistent, and the call to Redistribute (inside + // SplitParticles) may result in split particles to deposit twice on the + // coarse level. + if (do_splitting && (a_dt_type == DtType::SecondHalf || a_dt_type == DtType::Full) ){ SplitParticles(lev); } } -- cgit v1.2.3 From c51f53a2dc60076995e99bcf91ca8167edb1818f Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Mon, 23 Sep 2019 20:09:09 -0700 Subject: Remove trailing space --- Source/Particles/PhysicalParticleContainer.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 8dfddc833..b1e83d652 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1367,7 +1367,7 @@ PhysicalParticleContainer::Evolve (int lev, // When subcycling is ON, the splitting is done on the last call to // PhysicalParticleContainer::Evolve on the finest level, i.e., at the // end of the large timestep. Otherwise, the pushes on different levels - // are not consistent, and the call to Redistribute (inside + // are not consistent, and the call to Redistribute (inside // SplitParticles) may result in split particles to deposit twice on the // coarse level. if (do_splitting && (a_dt_type == DtType::SecondHalf || a_dt_type == DtType::Full) ){ -- cgit v1.2.3 From 37e62ed5e384256699d442fc34c9e5e3c3b9a50a Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Tue, 24 Sep 2019 13:46:17 -0700 Subject: Add comments --- Source/Particles/PhysicalParticleContainer.H | 4 ++- Source/Particles/PhysicalParticleContainer.cpp | 20 +++++++++++--- Source/Particles/Sorting/Partition.cpp | 37 +++++++++++++++++++++----- 3 files changed, 51 insertions(+), 10 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index b809acc45..c0fdce983 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -111,7 +111,9 @@ public: RealVector& uxp, RealVector& uyp, RealVector& uzp, - RealVector& wp ); + RealVector& wp, + RealVector& tmp, + ParticleVector& particle_tmp ); void copy_attribs(WarpXParIter& pti,const amrex::Real* xp, const amrex::Real* yp, const amrex::Real* zp); diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index e1046ac7a..dae723b3f 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -989,6 +989,8 @@ PhysicalParticleContainer::Evolve (int lev, FArrayBox filtered_Ex, filtered_Ey, filtered_Ez; FArrayBox filtered_Bx, filtered_By, filtered_Bz; + RealVector tmp; + ParticleVector particle_tmp; for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { @@ -1020,6 +1022,7 @@ PhysicalParticleContainer::Evolve (int lev, FArrayBox const* bzfab = &(Bz[pti]); Elixir exeli, eyeli, ezeli, bxeli, byeli, bzeli; + if (WarpX::use_fdtd_nci_corr) { #if (AMREX_SPACEDIM == 2) @@ -1079,11 +1082,22 @@ PhysicalParticleContainer::Evolve (int lev, Byp.assign(np,WarpX::B_external[1]); Bzp.assign(np,WarpX::B_external[2]); - long nfine_current = np; //! number of particles depositing to fine grid - long nfine_gather = np; //! number of particles gathering from fine grid + // Determine which particles deposit/gather in the buffer, and + // which particles deposit/gather in the fine patch + long nfine_current = np; + long nfine_gather = np; if (has_buffer && !do_not_push) { + // - Modify `nfine_current` and `nfine_gather` (in place) + // so that they correspond to the number of particles + // that deposit/gather in the fine patch respectively. + // - Reorder the particle arrays, + // so that the `nfine_current`/`nfine_gather` first particles + // deposit/gather in the fine patch + // and (thus) the `np-nfine_current`/`np-nfine_gather` last particles + // deposit/gather in the buffer PartitionParticlesInBuffers( nfine_current, nfine_gather, np, - pti, lev, current_masks, gather_masks, uxp, uyp, uzp, wp); + pti, lev, current_masks, gather_masks, uxp, uyp, uzp, wp, + tmp, particle_tmp); } const long np_current = (cjx) ? nfine_current : np; diff --git a/Source/Particles/Sorting/Partition.cpp b/Source/Particles/Sorting/Partition.cpp index d000a2c04..4c8e35f38 100644 --- a/Source/Particles/Sorting/Partition.cpp +++ b/Source/Particles/Sorting/Partition.cpp @@ -4,15 +4,43 @@ using namespace amrex; -/* \brief -*/ +/* \brief Determine which particles deposit/gather in the buffer, and + * and reorder the particle arrays accordingly + * + * More specifically: + * - Modify `nfine_current` and `nfine_gather` (in place) + * so that they correspond to the number of particles + * that deposit/gather in the fine patch respectively. + * - Reorder the particle arrays, + * so that the `nfine_current`/`nfine_gather` first particles + * deposit/gather in the fine patch + * and (thus) the `np-nfine_current`/`np-nfine_gather` last particles + * deposit/gather in the buffer + * + * \param nfine_current number of particles that deposit to the fine patch + * (modified by this function) + * \param nfine_gather number of particles that gather into the fine patch + * (modified by this function) + * \param np total number of particles in this tile + * \param pti object that holds the particle information for this tile + * \param lev current refinement level + * \param current_masks indicates, for each cell, whether that cell is + * in the deposition buffers or in the interior of the fine patch + * \param gather_masks indicates, for each cell, whether that cell is + * in the gather buffers or in the interior of the fine patch + * \param uxp, uyp, uzp, wp references to the particle momenta and weight + * (modified by this function) + * \param tmp temporary vector, used in reference swapping + * \param particle_tmp temporary vector, used in reference swapping + */ void PhysicalParticleContainer::PartitionParticlesInBuffers( long& nfine_current, long& nfine_gather, long const np, WarpXParIter& pti, int const lev, iMultiFab const* current_masks, iMultiFab const* gather_masks, - RealVector& uxp, RealVector& uyp, RealVector& uzp, RealVector& wp) + RealVector& uxp, RealVector& uyp, RealVector& uzp, RealVector& wp, + RealVector& tmp, ParticleVector& particle_tmp) { BL_PROFILE("PPC::Evolve::partition"); @@ -80,9 +108,6 @@ PhysicalParticleContainer::PartitionParticlesInBuffers( if (nfine_current != np || nfine_gather != np) { - RealVector tmp; - ParticleVector particle_tmp; - particle_tmp.resize(np); for (long ip = 0; ip < np; ++ip) { particle_tmp[ip] = aos[pid[ip]]; -- cgit v1.2.3 From f241d5e67d3b975b39aa58851bda814662130592 Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Tue, 24 Sep 2019 15:17:12 -0700 Subject: Move temporary variables into function --- Source/Particles/PhysicalParticleContainer.H | 4 +--- Source/Particles/PhysicalParticleContainer.cpp | 5 +---- Source/Particles/Sorting/Partition.cpp | 9 +++++---- 3 files changed, 7 insertions(+), 11 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index c0fdce983..b809acc45 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -111,9 +111,7 @@ public: RealVector& uxp, RealVector& uyp, RealVector& uzp, - RealVector& wp, - RealVector& tmp, - ParticleVector& particle_tmp ); + RealVector& wp ); void copy_attribs(WarpXParIter& pti,const amrex::Real* xp, const amrex::Real* yp, const amrex::Real* zp); diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index dae723b3f..7e1dbe74a 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -989,8 +989,6 @@ PhysicalParticleContainer::Evolve (int lev, FArrayBox filtered_Ex, filtered_Ey, filtered_Ez; FArrayBox filtered_Bx, filtered_By, filtered_Bz; - RealVector tmp; - ParticleVector particle_tmp; for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { @@ -1096,8 +1094,7 @@ PhysicalParticleContainer::Evolve (int lev, // and (thus) the `np-nfine_current`/`np-nfine_gather` last particles // deposit/gather in the buffer PartitionParticlesInBuffers( nfine_current, nfine_gather, np, - pti, lev, current_masks, gather_masks, uxp, uyp, uzp, wp, - tmp, particle_tmp); + pti, lev, current_masks, gather_masks, uxp, uyp, uzp, wp ); } const long np_current = (cjx) ? nfine_current : np; diff --git a/Source/Particles/Sorting/Partition.cpp b/Source/Particles/Sorting/Partition.cpp index 4c8e35f38..683dbbd04 100644 --- a/Source/Particles/Sorting/Partition.cpp +++ b/Source/Particles/Sorting/Partition.cpp @@ -30,8 +30,6 @@ using namespace amrex; * in the gather buffers or in the interior of the fine patch * \param uxp, uyp, uzp, wp references to the particle momenta and weight * (modified by this function) - * \param tmp temporary vector, used in reference swapping - * \param particle_tmp temporary vector, used in reference swapping */ void PhysicalParticleContainer::PartitionParticlesInBuffers( @@ -39,8 +37,7 @@ PhysicalParticleContainer::PartitionParticlesInBuffers( WarpXParIter& pti, int const lev, iMultiFab const* current_masks, iMultiFab const* gather_masks, - RealVector& uxp, RealVector& uyp, RealVector& uzp, RealVector& wp, - RealVector& tmp, ParticleVector& particle_tmp) + RealVector& uxp, RealVector& uyp, RealVector& uzp, RealVector& wp) { BL_PROFILE("PPC::Evolve::partition"); @@ -108,12 +105,16 @@ PhysicalParticleContainer::PartitionParticlesInBuffers( if (nfine_current != np || nfine_gather != np) { + + ParticleVector particle_tmp; particle_tmp.resize(np); + for (long ip = 0; ip < np; ++ip) { particle_tmp[ip] = aos[pid[ip]]; } std::swap(aos(), particle_tmp); + RealVector tmp; tmp.resize(np); for (long ip = 0; ip < np; ++ip) { -- cgit v1.2.3 From e78ecdb5b6cfe7cb6693007eef8ea9f410adc418 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Wed, 25 Sep 2019 09:14:36 -0700 Subject: spacing between split particles depends on PPC --- Source/Particles/PhysicalParticleContainer.cpp | 27 ++++++++++++++++---------- 1 file changed, 17 insertions(+), 10 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 617d31868..ca633fc41 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -31,8 +31,11 @@ PhysicalParticleContainer::PhysicalParticleContainer (AmrCore* amr_core, int isp pp.query("boost_adjust_transverse_positions", boost_adjust_transverse_positions); pp.query("do_backward_propagation", do_backward_propagation); + + // Initialize splitting pp.query("do_splitting", do_splitting); pp.query("split_type", split_type); + pp.query("do_continuous_injection", do_continuous_injection); // Whether to plot back-transformed (lab-frame) diagnostics // for this species. @@ -1315,7 +1318,11 @@ PhysicalParticleContainer::SplitParticles(int lev) for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { pti.GetPosition(xp, yp, zp); + + const amrex::Vector ppc_nd = plasma_injector->num_particles_per_cell_each_dim; const std::array& dx = WarpX::CellSize(lev); + amrex::Vector split_offset = {dx[0]/2/ppc_nd[0], dx[1]/2/ppc_nd[1], dx[2]/2/ppc_nd[2]}; + // particle Array Of Structs data auto& particles = pti.GetArrayOfStructs(); // particle Struct Of Arrays data @@ -1338,9 +1345,9 @@ PhysicalParticleContainer::SplitParticles(int lev) for (int ishift = -1; ishift < 2; ishift +=2 ){ for (int kshift = -1; kshift < 2; kshift +=2 ){ // Add one particle with offset in x and z - psplit_x.push_back( xp[i] + ishift*dx[0]/2 ); + psplit_x.push_back( xp[i] + ishift*split_offset[0] ); psplit_y.push_back( yp[i] ); - psplit_z.push_back( zp[i] + kshift*dx[2]/2 ); + psplit_z.push_back( zp[i] + kshift*split_offset[2] ); psplit_ux.push_back( uxp[i] ); psplit_uy.push_back( uyp[i] ); psplit_uz.push_back( uzp[i] ); @@ -1352,7 +1359,7 @@ PhysicalParticleContainer::SplitParticles(int lev) // 4 particles in 2d for (int ishift = -1; ishift < 2; ishift +=2 ){ // Add one particle with offset in x - psplit_x.push_back( xp[i] + ishift*dx[0]/2 ); + psplit_x.push_back( xp[i] + ishift*split_offset[0] ); psplit_y.push_back( yp[i] ); psplit_z.push_back( zp[i] ); psplit_ux.push_back( uxp[i] ); @@ -1362,7 +1369,7 @@ PhysicalParticleContainer::SplitParticles(int lev) // Add one particle with offset in z psplit_x.push_back( xp[i] ); psplit_y.push_back( yp[i] ); - psplit_z.push_back( zp[i] + ishift*dx[2]/2 ); + psplit_z.push_back( zp[i] + ishift*split_offset[2] ); psplit_ux.push_back( uxp[i] ); psplit_uy.push_back( uyp[i] ); psplit_uz.push_back( uzp[i] ); @@ -1377,9 +1384,9 @@ PhysicalParticleContainer::SplitParticles(int lev) for (int jshift = -1; jshift < 2; jshift +=2 ){ for (int kshift = -1; kshift < 2; kshift +=2 ){ // Add one particle with offset in x, y and z - psplit_x.push_back( xp[i] + ishift*dx[0]/2 ); - psplit_y.push_back( yp[i] + jshift*dx[1]/2 ); - psplit_z.push_back( zp[i] + kshift*dx[2]/2 ); + psplit_x.push_back( xp[i] + ishift*split_offset[0] ); + psplit_y.push_back( yp[i] + jshift*split_offset[1] ); + psplit_z.push_back( zp[i] + kshift*split_offset[2] ); psplit_ux.push_back( uxp[i] ); psplit_uy.push_back( uyp[i] ); psplit_uz.push_back( uzp[i] ); @@ -1392,7 +1399,7 @@ PhysicalParticleContainer::SplitParticles(int lev) // 6 particles in 3d for (int ishift = -1; ishift < 2; ishift +=2 ){ // Add one particle with offset in x - psplit_x.push_back( xp[i] + ishift*dx[0]/2 ); + psplit_x.push_back( xp[i] + ishift*split_offset[0] ); psplit_y.push_back( yp[i] ); psplit_z.push_back( zp[i] ); psplit_ux.push_back( uxp[i] ); @@ -1401,7 +1408,7 @@ PhysicalParticleContainer::SplitParticles(int lev) psplit_w.push_back( wp[i]/np_split ); // Add one particle with offset in y psplit_x.push_back( xp[i] ); - psplit_y.push_back( yp[i] + ishift*dx[1]/2 ); + psplit_y.push_back( yp[i] + ishift*split_offset[1] ); psplit_z.push_back( zp[i] ); psplit_ux.push_back( uxp[i] ); psplit_uy.push_back( uyp[i] ); @@ -1410,7 +1417,7 @@ PhysicalParticleContainer::SplitParticles(int lev) // Add one particle with offset in z psplit_x.push_back( xp[i] ); psplit_y.push_back( yp[i] ); - psplit_z.push_back( zp[i] + ishift*dx[2]/2 ); + psplit_z.push_back( zp[i] + ishift*split_offset[2] ); psplit_ux.push_back( uxp[i] ); psplit_uy.push_back( uyp[i] ); psplit_uz.push_back( uzp[i] ); -- cgit v1.2.3 From 6069ed791ca2f71e00b81c32aca4e785d7735290 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Wed, 25 Sep 2019 09:16:41 -0700 Subject: comments --- Source/Particles/PhysicalParticleContainer.cpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index ca633fc41..f2d3c0d75 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1319,9 +1319,14 @@ PhysicalParticleContainer::SplitParticles(int lev) { pti.GetPosition(xp, yp, zp); + // offset for split particles is computed as a function of cell size + // and number of particles per cell, so that a uniform distribution + // before splitting results in a uniform distribution after splitting const amrex::Vector ppc_nd = plasma_injector->num_particles_per_cell_each_dim; const std::array& dx = WarpX::CellSize(lev); - amrex::Vector split_offset = {dx[0]/2/ppc_nd[0], dx[1]/2/ppc_nd[1], dx[2]/2/ppc_nd[2]}; + amrex::Vector split_offset = {dx[0]/2/ppc_nd[0], + dx[1]/2/ppc_nd[1], + dx[2]/2/ppc_nd[2]}; // particle Array Of Structs data auto& particles = pti.GetArrayOfStructs(); -- cgit v1.2.3 From a07da5a3c4abc320d2cefff5c790f670c8e8e772 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Wed, 25 Sep 2019 09:21:29 -0700 Subject: fix EOL whitespaces --- Source/Particles/PhysicalParticleContainer.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index f2d3c0d75..5292dd3ff 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1327,7 +1327,7 @@ PhysicalParticleContainer::SplitParticles(int lev) amrex::Vector split_offset = {dx[0]/2/ppc_nd[0], dx[1]/2/ppc_nd[1], dx[2]/2/ppc_nd[2]}; - + // particle Array Of Structs data auto& particles = pti.GetArrayOfStructs(); // particle Struct Of Arrays data -- cgit v1.2.3 From eee301aa9c25cc9efa859728df20a9e5bf47ae95 Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Wed, 25 Sep 2019 10:29:18 -0700 Subject: Use explicit floating point operation --- Source/Particles/PhysicalParticleContainer.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 5292dd3ff..e276fd5ef 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1324,9 +1324,9 @@ PhysicalParticleContainer::SplitParticles(int lev) // before splitting results in a uniform distribution after splitting const amrex::Vector ppc_nd = plasma_injector->num_particles_per_cell_each_dim; const std::array& dx = WarpX::CellSize(lev); - amrex::Vector split_offset = {dx[0]/2/ppc_nd[0], - dx[1]/2/ppc_nd[1], - dx[2]/2/ppc_nd[2]}; + amrex::Vector split_offset = {dx[0]/2./ppc_nd[0], + dx[1]/2./ppc_nd[1], + dx[2]/2./ppc_nd[2]}; // particle Array Of Structs data auto& particles = pti.GetArrayOfStructs(); -- cgit v1.2.3 From 867353e82b862591a2d309d923df973598cd6178 Mon Sep 17 00:00:00 2001 From: Weiqun Zhang Date: Thu, 26 Sep 2019 13:28:39 -0700 Subject: fix types, amrex::Real and amrex::ParticleReal --- Source/Diagnostics/BoostedFrameDiagnostic.cpp | 14 +-- Source/Diagnostics/ParticleIO.cpp | 6 +- Source/FortranInterface/WarpX_f.H | 20 ++-- Source/Initialization/InjectorPosition.H | 4 +- Source/Initialization/PlasmaInjector.H | 6 +- Source/Laser/LaserParticleContainer.H | 8 +- Source/Laser/LaserParticleContainer.cpp | 16 ++-- Source/Laser/LaserProfiles.cpp | 14 +-- Source/Parser/GpuParser.H | 8 +- Source/Parser/WarpXParser.H | 26 +++--- Source/Parser/WarpXParser.cpp | 4 +- Source/Parser/wp_parser.tab.c | 2 +- Source/Parser/wp_parser.tab.h | 2 +- Source/Parser/wp_parser.y | 2 +- Source/Parser/wp_parser_c.h | 9 +- Source/Parser/wp_parser_y.c | 82 ++++++++-------- Source/Parser/wp_parser_y.h | 21 +++-- Source/Particles/Deposition/ChargeDeposition.H | 8 +- Source/Particles/Deposition/CurrentDeposition.H | 28 +++--- Source/Particles/Gather/FieldGather.H | 12 +-- Source/Particles/MultiParticleContainer.cpp | 8 +- Source/Particles/PhotonParticleContainer.H | 6 +- Source/Particles/PhotonParticleContainer.cpp | 30 +++--- Source/Particles/PhysicalParticleContainer.H | 10 +- Source/Particles/PhysicalParticleContainer.cpp | 104 ++++++++++----------- Source/Particles/Pusher/GetAndSetPosition.H | 16 ++-- Source/Particles/Pusher/UpdateMomentumBoris.H | 6 +- Source/Particles/Pusher/UpdateMomentumVay.H | 6 +- Source/Particles/Pusher/UpdatePosition.H | 4 +- Source/Particles/Pusher/UpdatePositionPhoton.H | 4 +- Source/Particles/RigidInjectedParticleContainer.H | 6 +- .../Particles/RigidInjectedParticleContainer.cpp | 76 +++++++-------- Source/Particles/WarpXParticleContainer.H | 30 +++--- Source/Particles/WarpXParticleContainer.cpp | 58 ++++++------ Source/Particles/deposit_cic.F90 | 10 +- Source/Particles/interpolate_cic.F90 | 10 +- Source/Particles/push_particles_ES.F90 | 18 ++-- Source/WarpX.cpp | 8 +- 38 files changed, 355 insertions(+), 347 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Diagnostics/BoostedFrameDiagnostic.cpp b/Source/Diagnostics/BoostedFrameDiagnostic.cpp index abd9e99fe..45343a0cb 100644 --- a/Source/Diagnostics/BoostedFrameDiagnostic.cpp +++ b/Source/Diagnostics/BoostedFrameDiagnostic.cpp @@ -876,37 +876,37 @@ writeParticleData(const WarpXParticleContainer::DiagnosticParticleData& pdata, field_name = name + Concatenate("w_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::w).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::w).data(), np, ofs); ofs.close(); field_name = name + Concatenate("x_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::x).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::x).data(), np, ofs); ofs.close(); field_name = name + Concatenate("y_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::y).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::y).data(), np, ofs); ofs.close(); field_name = name + Concatenate("z_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::z).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::z).data(), np, ofs); ofs.close(); field_name = name + Concatenate("ux_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::ux).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::ux).data(), np, ofs); ofs.close(); field_name = name + Concatenate("uy_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::uy).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::uy).data(), np, ofs); ofs.close(); field_name = name + Concatenate("uz_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::uz).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::uz).data(), np, ofs); ofs.close(); } diff --git a/Source/Diagnostics/ParticleIO.cpp b/Source/Diagnostics/ParticleIO.cpp index 5cf3f3047..2a9c16aa8 100644 --- a/Source/Diagnostics/ParticleIO.cpp +++ b/Source/Diagnostics/ParticleIO.cpp @@ -174,9 +174,9 @@ PhysicalParticleContainer::ConvertUnits(ConvertDirection convert_direction) { // - momenta are stored as a struct of array, in `attribs` auto& attribs = pti.GetAttribs(); - Real* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + ParticleReal* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); // Loop over the particles and convert momentum const long np = pti.numParticles(); ParallelFor( np, diff --git a/Source/FortranInterface/WarpX_f.H b/Source/FortranInterface/WarpX_f.H index 00212aec9..8b9fc2afd 100644 --- a/Source/FortranInterface/WarpX_f.H +++ b/Source/FortranInterface/WarpX_f.H @@ -86,14 +86,14 @@ extern "C" #endif const amrex::Real* dx); - void WRPX_DEPOSIT_CIC(const amrex::Real* particles, int ns, int np, - const amrex::Real* weights, + void WRPX_DEPOSIT_CIC(const amrex::ParticleReal* particles, int ns, int np, + const amrex::ParticleReal* weights, const amrex::Real* charge, amrex::Real* rho, const int* lo, const int* hi, const amrex::Real* plo, const amrex::Real* dx, const int* ng); - void WRPX_INTERPOLATE_CIC_TWO_LEVELS(const amrex::Real* particles, int ns, int np, + void WRPX_INTERPOLATE_CIC_TWO_LEVELS(const amrex::ParticleReal* particles, int ns, int np, amrex::Real* Ex_p, amrex::Real* Ey_p, #if (AMREX_SPACEDIM == 3) amrex::Real* Ez_p, @@ -111,7 +111,7 @@ extern "C" const int* clo, const int* chi, const amrex::Real* cdx, const amrex::Real* plo, const int* ng, const int* lev); - void WRPX_INTERPOLATE_CIC(const amrex::Real* particles, int ns, int np, + void WRPX_INTERPOLATE_CIC(const amrex::ParticleReal* particles, int ns, int np, amrex::Real* Ex_p, amrex::Real* Ey_p, #if (AMREX_SPACEDIM == 3) amrex::Real* Ez_p, @@ -124,10 +124,10 @@ extern "C" const amrex::Real* plo, const amrex::Real* dx, const int* ng); - void WRPX_PUSH_LEAPFROG(amrex::Real* particles, int ns, int np, - amrex::Real* vx_p, amrex::Real* vy_p, + void WRPX_PUSH_LEAPFROG(amrex::ParticleReal* particles, int ns, int np, + amrex::ParticleReal* vx_p, amrex::ParticleReal* vy_p, #if (AMREX_SPACEDIM == 3) - amrex::Real* vz_p, + amrex::ParticleReal* vz_p, #endif const amrex::Real* Ex_p, const amrex::Real* Ey_p, #if (AMREX_SPACEDIM == 3) @@ -136,10 +136,10 @@ extern "C" const amrex::Real* charge, const amrex::Real* mass, const amrex::Real* dt, const amrex::Real* prob_lo, const amrex::Real* prob_hi); - void WRPX_PUSH_LEAPFROG_POSITIONS(amrex::Real* particles, int ns, int np, - amrex::Real* vx_p, amrex::Real* vy_p, + void WRPX_PUSH_LEAPFROG_POSITIONS(amrex::ParticleReal* particles, int ns, int np, + amrex::ParticleReal* vx_p, amrex::ParticleReal* vy_p, #if (AMREX_SPACEDIM == 3) - amrex::Real* vz_p, + amrex::ParticleReal* vz_p, #endif const amrex::Real* dt, const amrex::Real* prob_lo, const amrex::Real* prob_hi); diff --git a/Source/Initialization/InjectorPosition.H b/Source/Initialization/InjectorPosition.H index f8f16746c..6ecae93e0 100644 --- a/Source/Initialization/InjectorPosition.H +++ b/Source/Initialization/InjectorPosition.H @@ -41,7 +41,9 @@ struct InjectorPositionRegular int ix_part = i_part/(ny*nz); // written this way backward compatibility int iz_part = (i_part-ix_part*(ny*nz)) / ny; int iy_part = (i_part-ix_part*(ny*nz)) - ny*iz_part; - return amrex::XDim3{(0.5+ix_part)/nx, (0.5+iy_part)/ny, (0.5+iz_part) / nz}; + return amrex::XDim3{(amrex::Real(0.5)+ix_part)/nx, + (amrex::Real(0.5)+iy_part)/ny, + (amrex::Real(0.5)+iz_part) / nz}; } private: amrex::Dim3 ppc; diff --git a/Source/Initialization/PlasmaInjector.H b/Source/Initialization/PlasmaInjector.H index a944165d6..56b32c827 100644 --- a/Source/Initialization/PlasmaInjector.H +++ b/Source/Initialization/PlasmaInjector.H @@ -43,9 +43,9 @@ public: bool doInjection () const noexcept { return inj_pos != NULL;} bool add_single_particle = false; - amrex::Vector single_particle_pos; - amrex::Vector single_particle_vel; - amrex::Real single_particle_weight; + amrex::Vector single_particle_pos; + amrex::Vector single_particle_vel; + amrex::ParticleReal single_particle_weight; bool gaussian_beam = false; amrex::Real x_m; diff --git a/Source/Laser/LaserParticleContainer.H b/Source/Laser/LaserParticleContainer.H index 3176b78f8..e2a0743bc 100644 --- a/Source/Laser/LaserParticleContainer.H +++ b/Source/Laser/LaserParticleContainer.H @@ -56,10 +56,10 @@ public: amrex::Real * AMREX_RESTRICT const pplane_Xp, amrex::Real * AMREX_RESTRICT const pplane_Yp); - void update_laser_particle (const int np, amrex::Real * AMREX_RESTRICT const puxp, - amrex::Real * AMREX_RESTRICT const puyp, - amrex::Real * AMREX_RESTRICT const puzp, - amrex::Real const * AMREX_RESTRICT const pwp, + void update_laser_particle (const int np, amrex::ParticleReal * AMREX_RESTRICT const puxp, + amrex::ParticleReal * AMREX_RESTRICT const puyp, + amrex::ParticleReal * AMREX_RESTRICT const puzp, + amrex::ParticleReal const * AMREX_RESTRICT const pwp, amrex::Real const * AMREX_RESTRICT const amplitude, const amrex::Real dt, const int thread_num); diff --git a/Source/Laser/LaserParticleContainer.cpp b/Source/Laser/LaserParticleContainer.cpp index a7be7101c..8571c74ad 100644 --- a/Source/Laser/LaserParticleContainer.cpp +++ b/Source/Laser/LaserParticleContainer.cpp @@ -643,9 +643,9 @@ LaserParticleContainer::calculate_laser_plane_coordinates ( Real * AMREX_RESTRICT const pplane_Xp, Real * AMREX_RESTRICT const pplane_Yp) { - Real const * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr(); - Real const * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr(); - Real const * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); + ParticleReal const * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr(); + ParticleReal const * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr(); + ParticleReal const * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); Real tmp_u_X_0 = u_X[0]; Real tmp_u_X_2 = u_X[2]; Real tmp_position_0 = position[0]; @@ -691,14 +691,14 @@ LaserParticleContainer::calculate_laser_plane_coordinates ( */ void LaserParticleContainer::update_laser_particle( - const int np, Real * AMREX_RESTRICT const puxp, Real * AMREX_RESTRICT const puyp, - Real * AMREX_RESTRICT const puzp, Real const * AMREX_RESTRICT const pwp, + const int np, ParticleReal * AMREX_RESTRICT const puxp, ParticleReal * AMREX_RESTRICT const puyp, + ParticleReal * AMREX_RESTRICT const puzp, ParticleReal const * AMREX_RESTRICT const pwp, Real const * AMREX_RESTRICT const amplitude, const Real dt, const int thread_num) { - Real * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr(); - Real * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr(); - Real * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); + ParticleReal * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr(); + ParticleReal * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr(); + ParticleReal * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); Real tmp_p_X_0 = p_X[0]; Real tmp_p_X_1 = p_X[1]; Real tmp_p_X_2 = p_X[2]; diff --git a/Source/Laser/LaserProfiles.cpp b/Source/Laser/LaserProfiles.cpp index 69804b17c..281ab2101 100644 --- a/Source/Laser/LaserProfiles.cpp +++ b/Source/Laser/LaserProfiles.cpp @@ -28,16 +28,16 @@ LaserParticleContainer::gaussian_laser_profile ( const Real oscillation_phase = k0 * PhysConst::c * ( t - profile_t_peak ); // The coefficients below contain info about Gouy phase, // laser diffraction, and phase front curvature - const Complex diffract_factor = 1. + I * profile_focal_distance - * 2./( k0 * profile_waist * profile_waist ); - const Complex inv_complex_waist_2 = 1./( profile_waist*profile_waist * diffract_factor ); + const Complex diffract_factor = Real(1.) + I * profile_focal_distance + * Real(2.)/( k0 * profile_waist * profile_waist ); + const Complex inv_complex_waist_2 = Real(1.)/( profile_waist*profile_waist * diffract_factor ); // Time stretching due to STCs and phi2 complex envelope // (1 if zeta=0, beta=0, phi2=0) - const Complex stretch_factor = 1. + 4. * + const Complex stretch_factor = Real(1.) + Real(4.) * (zeta+beta*profile_focal_distance) * (zeta+beta*profile_focal_distance) * (inv_tau2*inv_complex_waist_2) + - 2.*I*(phi2 - beta*beta*k0*profile_focal_distance) * inv_tau2; + Real(2.)*I*(phi2 - beta*beta*k0*profile_focal_distance) * inv_tau2; // Amplitude and monochromatic oscillations Complex prefactor = e_max * MathFunc::exp( I * oscillation_phase ); @@ -61,10 +61,10 @@ LaserParticleContainer::gaussian_laser_profile ( amrex::ParallelFor( np, [=] AMREX_GPU_DEVICE (int i) { - const Complex stc_exponent = 1./stretch_factor * inv_tau2 * + const Complex stc_exponent = Real(1.)/stretch_factor * inv_tau2 * MathFunc::pow((t - tmp_profile_t_peak - tmp_beta*k0*(Xp[i]*std::cos(tmp_theta_stc) + Yp[i]*std::sin(tmp_theta_stc)) - - 2.*I*(Xp[i]*std::cos(tmp_theta_stc) + Yp[i]*std::sin(tmp_theta_stc)) + Real(2.)*I*(Xp[i]*std::cos(tmp_theta_stc) + Yp[i]*std::sin(tmp_theta_stc)) *( tmp_zeta - tmp_beta*tmp_profile_focal_distance ) * inv_complex_waist_2),2); // stcfactor = everything but complex transverse envelope const Complex stcfactor = prefactor * MathFunc::exp( - stc_exponent ); diff --git a/Source/Parser/GpuParser.H b/Source/Parser/GpuParser.H index e49671e06..c158ee314 100644 --- a/Source/Parser/GpuParser.H +++ b/Source/Parser/GpuParser.H @@ -16,16 +16,16 @@ public: void clear (); AMREX_GPU_HOST_DEVICE - double - operator() (double x, double y, double z) const noexcept + amrex::Real + operator() (amrex::Real x, amrex::Real y, amrex::Real z) const noexcept { #ifdef AMREX_USE_GPU #ifdef AMREX_DEVICE_COMPILE // WarpX compiled for GPU, function compiled for __device__ // the 3D position of each particle is stored in shared memory. - amrex::Gpu::SharedMemory gsm; - double* p = gsm.dataPtr(); + amrex::Gpu::SharedMemory gsm; + amrex::Real* p = gsm.dataPtr(); int tid = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*(blockDim.x*blockDim.y); p[tid*3] = x; p[tid*3+1] = y; diff --git a/Source/Parser/WarpXParser.H b/Source/Parser/WarpXParser.H index ffa61e457..8c1d854d8 100644 --- a/Source/Parser/WarpXParser.H +++ b/Source/Parser/WarpXParser.H @@ -6,6 +6,8 @@ #include #include +#include + #include "wp_parser_c.h" #include "wp_parser_y.h" @@ -23,15 +25,15 @@ public: ~WarpXParser (); void define (std::string const& func_body); - void setConstant (std::string const& name, double c); + void setConstant (std::string const& name, amrex::Real c); // // Option 1: Register every variable to an address provided. // Assign values to external variables. // Call eval(). - void registerVariable (std::string const& name, double& var); + void registerVariable (std::string const& name, amrex::Real& var); // - inline double eval () const noexcept; + inline amrex::Real eval () const noexcept; // // Option 2: Register all variables at once. Parser will create @@ -40,7 +42,7 @@ public: void registerVariables (std::vector const& names); // template inline - double eval (T x, Ts... yz) const noexcept; + amrex::Real eval (T x, Ts... yz) const noexcept; void print () const; @@ -54,23 +56,23 @@ private: void clear (); template inline - void unpack (double* p, T x) const noexcept; + void unpack (amrex::Real* p, T x) const noexcept; template inline - void unpack (double* p, T x, Ts... yz) const noexcept; + void unpack (amrex::Real* p, T x, Ts... yz) const noexcept; std::string m_expression; #ifdef _OPENMP std::vector m_parser; - mutable std::vector > m_variables; + mutable std::vector > m_variables; #else struct wp_parser* m_parser = nullptr; - mutable std::array m_variables; + mutable std::array m_variables; #endif }; inline -double +amrex::Real WarpXParser::eval () const noexcept { #ifdef _OPENMP @@ -82,7 +84,7 @@ WarpXParser::eval () const noexcept template inline -double +amrex::Real WarpXParser::eval (T x, Ts... yz) const noexcept { #ifdef _OPENMP @@ -96,7 +98,7 @@ WarpXParser::eval (T x, Ts... yz) const noexcept template inline void -WarpXParser::unpack (double* p, T x) const noexcept +WarpXParser::unpack (amrex::Real* p, T x) const noexcept { *p = x; } @@ -104,7 +106,7 @@ WarpXParser::unpack (double* p, T x) const noexcept template inline void -WarpXParser::unpack (double* p, T x, Ts... yz) const noexcept +WarpXParser::unpack (amrex::Real* p, T x, Ts... yz) const noexcept { *p++ = x; unpack(p, yz...); diff --git a/Source/Parser/WarpXParser.cpp b/Source/Parser/WarpXParser.cpp index 3237086f2..ced536327 100644 --- a/Source/Parser/WarpXParser.cpp +++ b/Source/Parser/WarpXParser.cpp @@ -69,7 +69,7 @@ WarpXParser::clear () } void -WarpXParser::registerVariable (std::string const& name, double& var) +WarpXParser::registerVariable (std::string const& name, amrex::Real& var) { // We assume this is called inside OMP parallel region #ifdef _OPENMP @@ -105,7 +105,7 @@ WarpXParser::registerVariables (std::vector const& names) } void -WarpXParser::setConstant (std::string const& name, double c) +WarpXParser::setConstant (std::string const& name, amrex::Real c) { #ifdef _OPENMP diff --git a/Source/Parser/wp_parser.tab.c b/Source/Parser/wp_parser.tab.c index 3981894a5..0f7c2403d 100644 --- a/Source/Parser/wp_parser.tab.c +++ b/Source/Parser/wp_parser.tab.c @@ -138,7 +138,7 @@ union YYSTYPE #line 19 "wp_parser.y" /* yacc.c:352 */ struct wp_node* n; - double d; + amrex_real d; struct wp_symbol* s; enum wp_f1_t f1; enum wp_f2_t f2; diff --git a/Source/Parser/wp_parser.tab.h b/Source/Parser/wp_parser.tab.h index b50516808..0c859fc03 100644 --- a/Source/Parser/wp_parser.tab.h +++ b/Source/Parser/wp_parser.tab.h @@ -75,7 +75,7 @@ union YYSTYPE #line 19 "wp_parser.y" /* yacc.c:1921 */ struct wp_node* n; - double d; + amrex_real d; struct wp_symbol* s; enum wp_f1_t f1; enum wp_f2_t f2; diff --git a/Source/Parser/wp_parser.y b/Source/Parser/wp_parser.y index 453eda1cd..809dbfa5e 100644 --- a/Source/Parser/wp_parser.y +++ b/Source/Parser/wp_parser.y @@ -18,7 +18,7 @@ */ %union { struct wp_node* n; - double d; + amrex_real d; struct wp_symbol* s; enum wp_f1_t f1; enum wp_f2_t f2; diff --git a/Source/Parser/wp_parser_c.h b/Source/Parser/wp_parser_c.h index 3aafdec65..970d6b355 100644 --- a/Source/Parser/wp_parser_c.h +++ b/Source/Parser/wp_parser_c.h @@ -4,6 +4,7 @@ #include "wp_parser_y.h" #include #include +#include #ifdef __cplusplus extern "C" { @@ -21,15 +22,15 @@ extern "C" { #include AMREX_GPU_HOST_DEVICE -inline double +inline amrex_real wp_ast_eval (struct wp_node* node) { - double result; + amrex_real result; #ifdef AMREX_DEVICE_COMPILE - extern __shared__ double extern_xyz[]; + extern __shared__ amrex_real extern_xyz[]; int tid = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*(blockDim.x*blockDim.y); - double* x = extern_xyz + tid*3; + amrex_real* x = extern_xyz + tid*3; #endif switch (node->type) diff --git a/Source/Parser/wp_parser_y.c b/Source/Parser/wp_parser_y.c index 259f9368b..b71b42638 100644 --- a/Source/Parser/wp_parser_y.c +++ b/Source/Parser/wp_parser_y.c @@ -6,8 +6,6 @@ #include "wp_parser_y.h" #include "wp_parser.tab.h" -#include - static struct wp_node* wp_root = NULL; /* This is called by a bison rule to store the original AST in a @@ -21,7 +19,7 @@ wp_defexpr (struct wp_node* body) } struct wp_node* -wp_newnumber (double d) +wp_newnumber (amrex_real d) { struct wp_number* r = (struct wp_number*) malloc(sizeof(struct wp_number)); r->type = WP_NUMBER; @@ -154,8 +152,8 @@ wp_parser_dup (struct wp_parser* source) } AMREX_GPU_HOST_DEVICE -double -wp_call_f1 (enum wp_f1_t type, double a) +amrex_real +wp_call_f1 (enum wp_f1_t type, amrex_real a) { switch (type) { case WP_SQRT: return sqrt(a); @@ -185,8 +183,8 @@ wp_call_f1 (enum wp_f1_t type, double a) } AMREX_GPU_HOST_DEVICE -double -wp_call_f2 (enum wp_f2_t type, double a, double b) +amrex_real +wp_call_f2 (enum wp_f2_t type, amrex_real a, amrex_real b) { switch (type) { case WP_POW: @@ -356,13 +354,13 @@ wp_parser_ast_dup (struct wp_parser* my_parser, struct wp_node* node, int move) #define WP_MOVEUP_R(node, v) \ struct wp_node* n = node->r->r; \ - double* p = node->r->rip.p; \ + amrex_real* p = node->r->rip.p; \ node->r = n; \ node->lvp.v = v; \ node->rip.p = p; #define WP_MOVEUP_L(node, v) \ struct wp_node* n = node->l->r; \ - double* p = node->l->rip.p; \ + amrex_real* p = node->l->rip.p; \ node->r = n; \ node->lvp.v = v; \ node->rip.p = p; @@ -392,7 +390,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = ((struct wp_number*)(node->l))->value + amrex_real v = ((struct wp_number*)(node->l))->value + ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; @@ -422,28 +420,28 @@ wp_ast_optimize (struct wp_node* node) else if (node->l->type == WP_NUMBER && node->r->type == WP_ADD_VP) { - double v = ((struct wp_number*)(node->l))->value + WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value + WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_ADD_VP; } else if (node->l->type == WP_NUMBER && node->r->type == WP_SUB_VP) { - double v = ((struct wp_number*)(node->l))->value + WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value + WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_SUB_VP; } else if (node->l->type == WP_ADD_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) + ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) + ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_ADD_VP; } else if (node->l->type == WP_SUB_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) + ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) + ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_SUB_VP; } @@ -455,7 +453,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = ((struct wp_number*)(node->l))->value + amrex_real v = ((struct wp_number*)(node->l))->value - ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; @@ -485,28 +483,28 @@ wp_ast_optimize (struct wp_node* node) else if (node->l->type == WP_NUMBER && node->r->type == WP_ADD_VP) { - double v = ((struct wp_number*)(node->l))->value - WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value - WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_SUB_VP; } else if (node->l->type == WP_NUMBER && node->r->type == WP_SUB_VP) { - double v = ((struct wp_number*)(node->l))->value - WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value - WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_ADD_VP; } else if (node->l->type == WP_ADD_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) - ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) - ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_ADD_VP; } else if (node->l->type == WP_SUB_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) - ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) - ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_SUB_VP; } @@ -518,7 +516,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = ((struct wp_number*)(node->l))->value + amrex_real v = ((struct wp_number*)(node->l))->value * ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; @@ -548,28 +546,28 @@ wp_ast_optimize (struct wp_node* node) else if (node->l->type == WP_NUMBER && node->r->type == WP_MUL_VP) { - double v = ((struct wp_number*)(node->l))->value * WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value * WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_MUL_VP; } else if (node->l->type == WP_NUMBER && node->r->type == WP_DIV_VP) { - double v = ((struct wp_number*)(node->l))->value * WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value * WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_DIV_VP; } else if (node->l->type == WP_MUL_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) * ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) * ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_MUL_VP; } else if (node->l->type == WP_DIV_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) * ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) * ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_DIV_VP; } @@ -581,7 +579,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = ((struct wp_number*)(node->l))->value + amrex_real v = ((struct wp_number*)(node->l))->value / ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; @@ -611,28 +609,28 @@ wp_ast_optimize (struct wp_node* node) else if (node->l->type == WP_NUMBER && node->r->type == WP_MUL_VP) { - double v = ((struct wp_number*)(node->l))->value / WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value / WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_DIV_VP; } else if (node->l->type == WP_NUMBER && node->r->type == WP_DIV_VP) { - double v = ((struct wp_number*)(node->l))->value / WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value / WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_MUL_VP; } else if (node->l->type == WP_MUL_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) / ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) / ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_MUL_VP; } else if (node->l->type == WP_DIV_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) / ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) / ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_DIV_VP; } @@ -641,7 +639,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->l); if (node->l->type == WP_NUMBER) { - double v = -((struct wp_number*)(node->l))->value; + amrex_real v = -((struct wp_number*)(node->l))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -675,7 +673,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->l); if (node->l->type == WP_NUMBER) { - double v = wp_call_f1 + amrex_real v = wp_call_f1 (((struct wp_f1*)node)->ftype, ((struct wp_number*)(((struct wp_f1*)node)->l))->value); ((struct wp_number*)node)->type = WP_NUMBER; @@ -688,7 +686,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = wp_call_f2 + amrex_real v = wp_call_f2 (((struct wp_f2*)node)->ftype, ((struct wp_number*)(((struct wp_f2*)node)->l))->value, ((struct wp_number*)(((struct wp_f2*)node)->r))->value); @@ -698,7 +696,7 @@ wp_ast_optimize (struct wp_node* node) else if (node->r->type == WP_NUMBER && ((struct wp_f2*)node)->ftype == WP_POW) { struct wp_node* n = node->l; - double v = ((struct wp_number*)(node->r))->value; + amrex_real v = ((struct wp_number*)(node->r))->value; if (-3.0 == v) { ((struct wp_f1*)node)->type = WP_F1; ((struct wp_f1*)node)->l = n; @@ -733,7 +731,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->r); if (node->r->type == WP_NUMBER) { - double v = node->lvp.v + ((struct wp_number*)(node->r))->value; + amrex_real v = node->lvp.v + ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -742,7 +740,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->r); if (node->r->type == WP_NUMBER) { - double v = node->lvp.v - ((struct wp_number*)(node->r))->value; + amrex_real v = node->lvp.v - ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -751,7 +749,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->r); if (node->r->type == WP_NUMBER) { - double v = node->lvp.v * ((struct wp_number*)(node->r))->value; + amrex_real v = node->lvp.v * ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -760,7 +758,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->r); if (node->r->type == WP_NUMBER) { - double v = node->lvp.v / ((struct wp_number*)(node->r))->value; + amrex_real v = node->lvp.v / ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -769,7 +767,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->l); if (node->l->type == WP_NUMBER) { - double v = -((struct wp_number*)(node->l))->value; + amrex_real v = -((struct wp_number*)(node->l))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -938,7 +936,7 @@ wp_ast_print (struct wp_node* node) } void -wp_ast_regvar (struct wp_node* node, char const* name, double* p) +wp_ast_regvar (struct wp_node* node, char const* name, amrex_real* p) { switch (node->type) { @@ -1047,7 +1045,7 @@ wp_ast_regvar_gpu (struct wp_node* node, char const* name, int i) } } -void wp_ast_setconst (struct wp_node* node, char const* name, double c) +void wp_ast_setconst (struct wp_node* node, char const* name, amrex_real c) { switch (node->type) { @@ -1099,7 +1097,7 @@ void wp_ast_setconst (struct wp_node* node, char const* name, double c) } void -wp_parser_regvar (struct wp_parser* parser, char const* name, double* p) +wp_parser_regvar (struct wp_parser* parser, char const* name, amrex_real* p) { wp_ast_regvar(parser->ast, name, p); } @@ -1111,7 +1109,7 @@ wp_parser_regvar_gpu (struct wp_parser* parser, char const* name, int i) } void -wp_parser_setconst (struct wp_parser* parser, char const* name, double c) +wp_parser_setconst (struct wp_parser* parser, char const* name, amrex_real c) { wp_ast_setconst(parser->ast, name, c); wp_ast_optimize(parser->ast); diff --git a/Source/Parser/wp_parser_y.h b/Source/Parser/wp_parser_y.h index 8c9f8e4e4..d83815090 100644 --- a/Source/Parser/wp_parser_y.h +++ b/Source/Parser/wp_parser_y.h @@ -2,6 +2,7 @@ #define WP_PARSER_Y_H_ #include +#include #ifdef __cplusplus #include @@ -77,11 +78,11 @@ enum wp_node_t { union wp_ip { int i; - double* p; + amrex_real* p; }; union wp_vp { - double v; + amrex_real v; union wp_ip ip; }; @@ -95,7 +96,7 @@ struct wp_node { struct wp_number { enum wp_node_t type; - double value; + amrex_real value; }; struct wp_symbol { @@ -122,7 +123,7 @@ struct wp_f2 { /* Builtin functions with two arguments */ /* These functions are used in bison rules to generate the original * AST. */ void wp_defexpr (struct wp_node* body); -struct wp_node* wp_newnumber (double d); +struct wp_node* wp_newnumber (amrex_real d); struct wp_symbol* wp_makesymbol (char* name); struct wp_node* wp_newsymbol (struct wp_symbol* sym); struct wp_node* wp_newnode (enum wp_node_t type, struct wp_node* l, @@ -153,20 +154,20 @@ void wp_parser_delete (struct wp_parser* parser); struct wp_parser* wp_parser_dup (struct wp_parser* source); struct wp_node* wp_parser_ast_dup (struct wp_parser* parser, struct wp_node* src, int move); -void wp_parser_regvar (struct wp_parser* parser, char const* name, double* p); +void wp_parser_regvar (struct wp_parser* parser, char const* name, amrex_real* p); void wp_parser_regvar_gpu (struct wp_parser* parser, char const* name, int i); -void wp_parser_setconst (struct wp_parser* parser, char const* name, double c); +void wp_parser_setconst (struct wp_parser* parser, char const* name, amrex_real c); /* We need to walk the tree in these functions */ void wp_ast_optimize (struct wp_node* node); size_t wp_ast_size (struct wp_node* node); void wp_ast_print (struct wp_node* node); -void wp_ast_regvar (struct wp_node* node, char const* name, double* p); +void wp_ast_regvar (struct wp_node* node, char const* name, amrex_real* p); void wp_ast_regvar_gpu (struct wp_node* node, char const* name, int i); -void wp_ast_setconst (struct wp_node* node, char const* name, double c); +void wp_ast_setconst (struct wp_node* node, char const* name, amrex_real c); -AMREX_GPU_HOST_DEVICE double wp_call_f1 (enum wp_f1_t type, double a); -AMREX_GPU_HOST_DEVICE double wp_call_f2 (enum wp_f2_t type, double a, double b); +AMREX_GPU_HOST_DEVICE amrex_real wp_call_f1 (enum wp_f1_t type, amrex_real a); +AMREX_GPU_HOST_DEVICE amrex_real wp_call_f2 (enum wp_f2_t type, amrex_real a, amrex_real b); #ifdef __cplusplus } diff --git a/Source/Particles/Deposition/ChargeDeposition.H b/Source/Particles/Deposition/ChargeDeposition.H index b9210e67c..eec407a2b 100755 --- a/Source/Particles/Deposition/ChargeDeposition.H +++ b/Source/Particles/Deposition/ChargeDeposition.H @@ -18,10 +18,10 @@ * /param q : species charge. */ template -void doChargeDepositionShapeN(const amrex::Real * const xp, - const amrex::Real * const yp, - const amrex::Real * const zp, - const amrex::Real * const wp, +void doChargeDepositionShapeN(const amrex::ParticleReal * const xp, + const amrex::ParticleReal * const yp, + const amrex::ParticleReal * const zp, + const amrex::ParticleReal * const wp, const int * const ion_lev, const amrex::Array4& rho_arr, const long np_to_depose, diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H index 7a96dab9a..f809ef1ec 100644 --- a/Source/Particles/Deposition/CurrentDeposition.H +++ b/Source/Particles/Deposition/CurrentDeposition.H @@ -24,13 +24,13 @@ * /param q : species charge. */ template -void doDepositionShapeN(const amrex::Real * const xp, - const amrex::Real * const yp, - const amrex::Real * const zp, - const amrex::Real * const wp, - const amrex::Real * const uxp, - const amrex::Real * const uyp, - const amrex::Real * const uzp, +void doDepositionShapeN(const amrex::ParticleReal * const xp, + const amrex::ParticleReal * const yp, + const amrex::ParticleReal * const zp, + const amrex::ParticleReal * const wp, + const amrex::ParticleReal * const uxp, + const amrex::ParticleReal * const uyp, + const amrex::ParticleReal * const uzp, const int * const ion_lev, const amrex::Array4& jx_arr, const amrex::Array4& jy_arr, @@ -189,13 +189,13 @@ void doDepositionShapeN(const amrex::Real * const xp, * \param n_rz_azimuthal_modes: Number of azimuthal modes when using RZ geometry */ template -void doEsirkepovDepositionShapeN (const amrex::Real * const xp, - const amrex::Real * const yp, - const amrex::Real * const zp, - const amrex::Real * const wp, - const amrex::Real * const uxp, - const amrex::Real * const uyp, - const amrex::Real * const uzp, +void doEsirkepovDepositionShapeN (const amrex::ParticleReal * const xp, + const amrex::ParticleReal * const yp, + const amrex::ParticleReal * const zp, + const amrex::ParticleReal * const wp, + const amrex::ParticleReal * const uxp, + const amrex::ParticleReal * const uyp, + const amrex::ParticleReal * const uzp, const int * ion_lev, const amrex::Array4& Jx_arr, const amrex::Array4& Jy_arr, diff --git a/Source/Particles/Gather/FieldGather.H b/Source/Particles/Gather/FieldGather.H index 6727b0aa9..c5ec6fb5b 100644 --- a/Source/Particles/Gather/FieldGather.H +++ b/Source/Particles/Gather/FieldGather.H @@ -19,12 +19,12 @@ * \param n_rz_azimuthal_modes: Number of azimuthal modes when using RZ geometry */ template -void doGatherShapeN(const amrex::Real * const xp, - const amrex::Real * const yp, - const amrex::Real * const zp, - amrex::Real * const Exp, amrex::Real * const Eyp, - amrex::Real * const Ezp, amrex::Real * const Bxp, - amrex::Real * const Byp, amrex::Real * const Bzp, +void doGatherShapeN(const amrex::ParticleReal * const xp, + const amrex::ParticleReal * const yp, + const amrex::ParticleReal * const zp, + amrex::ParticleReal * const Exp, amrex::ParticleReal * const Eyp, + amrex::ParticleReal * const Ezp, amrex::ParticleReal * const Bxp, + amrex::ParticleReal * const Byp, amrex::ParticleReal * const Bzp, const amrex::Array4& ex_arr, const amrex::Array4& ey_arr, const amrex::Array4& ez_arr, diff --git a/Source/Particles/MultiParticleContainer.cpp b/Source/Particles/MultiParticleContainer.cpp index bb795465e..715c97b99 100644 --- a/Source/Particles/MultiParticleContainer.cpp +++ b/Source/Particles/MultiParticleContainer.cpp @@ -545,12 +545,12 @@ namespace WarpXParticleContainer::ParticleType* particles_source = ptile_source.GetArrayOfStructs()().data(); // --- source SoA particle data auto& soa_source = ptile_source.GetStructOfArrays(); - GpuArray attribs_source; + GpuArray attribs_source; for (int ia = 0; ia < PIdx::nattribs; ++ia) { attribs_source[ia] = soa_source.GetRealData(ia).data(); } // --- source runtime attribs - GpuArray runtime_uold_source; + GpuArray runtime_uold_source; // Prepare arrays for boosted frame diagnostics. runtime_uold_source[0] = soa_source.GetRealData(PIdx::ux).data(); runtime_uold_source[1] = soa_source.GetRealData(PIdx::uy).data(); @@ -590,13 +590,13 @@ namespace WarpXParticleContainer::ParticleType* particles_product = ptile_product.GetArrayOfStructs()().data() + np_product_old; // --- product SoA particle data auto& soa_product = ptile_product.GetStructOfArrays(); - GpuArray attribs_product; + GpuArray attribs_product; for (int ia = 0; ia < PIdx::nattribs; ++ia) { // First element is the first newly-created product particle attribs_product[ia] = soa_product.GetRealData(ia).data() + np_product_old; } // --- product runtime attribs - GpuArray runtime_attribs_product; + GpuArray runtime_attribs_product; bool do_boosted_product = WarpX::do_boosted_frame_diagnostic && pc_product->DoBoostedFrameDiags(); if (do_boosted_product) { diff --git a/Source/Particles/PhotonParticleContainer.H b/Source/Particles/PhotonParticleContainer.H index 3ac304bdc..a6ffd1d76 100644 --- a/Source/Particles/PhotonParticleContainer.H +++ b/Source/Particles/PhotonParticleContainer.H @@ -41,9 +41,9 @@ public: DtType a_dt_type=DtType::Full) override; virtual void PushPX(WarpXParIter& pti, - amrex::Cuda::ManagedDeviceVector& xp, - amrex::Cuda::ManagedDeviceVector& yp, - amrex::Cuda::ManagedDeviceVector& zp, + amrex::Cuda::ManagedDeviceVector& xp, + amrex::Cuda::ManagedDeviceVector& yp, + amrex::Cuda::ManagedDeviceVector& zp, amrex::Real dt, DtType a_dt_type=DtType::Full) override; diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index 9de441e5c..4a75ec9f3 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -34,27 +34,27 @@ void PhotonParticleContainer::InitData() void PhotonParticleContainer::PushPX(WarpXParIter& pti, - Cuda::ManagedDeviceVector& xp, - Cuda::ManagedDeviceVector& yp, - Cuda::ManagedDeviceVector& zp, + Cuda::ManagedDeviceVector& xp, + Cuda::ManagedDeviceVector& yp, + Cuda::ManagedDeviceVector& zp, Real dt, DtType a_dt_type) { // This wraps the momentum and position advance so that inheritors can modify the call. auto& attribs = pti.GetAttribs(); // Extract pointers to the different particle quantities - Real* const AMREX_RESTRICT x = xp.dataPtr(); - Real* const AMREX_RESTRICT y = yp.dataPtr(); - Real* const AMREX_RESTRICT z = zp.dataPtr(); - Real* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); - const Real* const AMREX_RESTRICT Ex = attribs[PIdx::Ex].dataPtr(); - const Real* const AMREX_RESTRICT Ey = attribs[PIdx::Ey].dataPtr(); - const Real* const AMREX_RESTRICT Ez = attribs[PIdx::Ez].dataPtr(); - const Real* const AMREX_RESTRICT Bx = attribs[PIdx::Bx].dataPtr(); - const Real* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); - const Real* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); + ParticleReal* const AMREX_RESTRICT x = xp.dataPtr(); + ParticleReal* const AMREX_RESTRICT y = yp.dataPtr(); + ParticleReal* const AMREX_RESTRICT z = zp.dataPtr(); + ParticleReal* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ex = attribs[PIdx::Ex].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ey = attribs[PIdx::Ey].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ez = attribs[PIdx::Ez].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bx = attribs[PIdx::Bx].dataPtr(); + const ParticleReal* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags) { diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index b809acc45..1ee3407ff 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -87,9 +87,9 @@ public: DtType a_dt_type=DtType::Full) override; virtual void PushPX(WarpXParIter& pti, - amrex::Cuda::ManagedDeviceVector& xp, - amrex::Cuda::ManagedDeviceVector& yp, - amrex::Cuda::ManagedDeviceVector& zp, + amrex::Cuda::ManagedDeviceVector& xp, + amrex::Cuda::ManagedDeviceVector& yp, + amrex::Cuda::ManagedDeviceVector& zp, amrex::Real dt, DtType a_dt_type=DtType::Full); virtual void PushP (int lev, amrex::Real dt, @@ -113,8 +113,8 @@ public: RealVector& uzp, RealVector& wp ); - void copy_attribs(WarpXParIter& pti,const amrex::Real* xp, - const amrex::Real* yp, const amrex::Real* zp); + void copy_attribs(WarpXParIter& pti,const amrex::ParticleReal* xp, + const amrex::ParticleReal* yp, const amrex::ParticleReal* zp); virtual void PostRestart () final {} diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index e276fd5ef..0a1e40953 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -199,7 +199,7 @@ PhysicalParticleContainer::CheckAndAddParticle(Real x, Real y, Real z, std::array u, Real weight) { - std::array attribs; + std::array attribs; attribs.fill(0.0); // update attribs with input arguments @@ -364,13 +364,13 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) for (int dir=0; dir= part_realbox.lo(dir) ) { Real ncells_adjust = std::floor( (part_realbox.hi(dir) - tile_realbox.hi(dir))/dx[dir] ); - overlap_realbox.setHi( dir, part_realbox.hi(dir) - std::max(ncells_adjust, 0.) * dx[dir]); + overlap_realbox.setHi( dir, part_realbox.hi(dir) - std::max(ncells_adjust, Real(0.)) * dx[dir]); } else { no_overlap = true; break; } @@ -440,7 +440,7 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) ParticleType* pp = particle_tile.GetArrayOfStructs()().data() + old_size; auto& soa = particle_tile.GetStructOfArrays(); - GpuArray pa; + GpuArray pa; for (int ia = 0; ia < PIdx::nattribs; ++ia) { pa[ia] = soa.GetRealData(ia).data() + old_size; } @@ -1302,7 +1302,7 @@ PhysicalParticleContainer::SplitParticles(int lev) { auto& mypc = WarpX::GetInstance().GetPartContainer(); auto& pctmp_split = mypc.GetPCtmp(); - Cuda::ManagedDeviceVector xp, yp, zp; + Cuda::ManagedDeviceVector xp, yp, zp; RealVector psplit_x, psplit_y, psplit_z, psplit_w; RealVector psplit_ux, psplit_uy, psplit_uz; long np_split_to_add = 0; @@ -1460,27 +1460,27 @@ PhysicalParticleContainer::SplitParticles(int lev) void PhysicalParticleContainer::PushPX(WarpXParIter& pti, - Cuda::ManagedDeviceVector& xp, - Cuda::ManagedDeviceVector& yp, - Cuda::ManagedDeviceVector& zp, + Cuda::ManagedDeviceVector& xp, + Cuda::ManagedDeviceVector& yp, + Cuda::ManagedDeviceVector& zp, Real dt, DtType a_dt_type) { // This wraps the momentum and position advance so that inheritors can modify the call. auto& attribs = pti.GetAttribs(); // Extract pointers to the different particle quantities - Real* const AMREX_RESTRICT x = xp.dataPtr(); - Real* const AMREX_RESTRICT y = yp.dataPtr(); - Real* const AMREX_RESTRICT z = zp.dataPtr(); - Real* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); - const Real* const AMREX_RESTRICT Ex = attribs[PIdx::Ex].dataPtr(); - const Real* const AMREX_RESTRICT Ey = attribs[PIdx::Ey].dataPtr(); - const Real* const AMREX_RESTRICT Ez = attribs[PIdx::Ez].dataPtr(); - const Real* const AMREX_RESTRICT Bx = attribs[PIdx::Bx].dataPtr(); - const Real* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); - const Real* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); + ParticleReal* const AMREX_RESTRICT x = xp.dataPtr(); + ParticleReal* const AMREX_RESTRICT y = yp.dataPtr(); + ParticleReal* const AMREX_RESTRICT z = zp.dataPtr(); + ParticleReal* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ex = attribs[PIdx::Ex].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ey = attribs[PIdx::Ey].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ez = attribs[PIdx::Ez].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bx = attribs[PIdx::Bx].dataPtr(); + const ParticleReal* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags && (a_dt_type!=DtType::SecondHalf)) { @@ -1589,15 +1589,15 @@ PhysicalParticleContainer::PushP (int lev, Real dt, // This wraps the momentum advance so that inheritors can modify the call. // Extract pointers to the different particle quantities - Real* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); - const Real* const AMREX_RESTRICT Expp = Exp.dataPtr(); - const Real* const AMREX_RESTRICT Eypp = Eyp.dataPtr(); - const Real* const AMREX_RESTRICT Ezpp = Ezp.dataPtr(); - const Real* const AMREX_RESTRICT Bxpp = Bxp.dataPtr(); - const Real* const AMREX_RESTRICT Bypp = Byp.dataPtr(); - const Real* const AMREX_RESTRICT Bzpp = Bzp.dataPtr(); + ParticleReal* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Expp = Exp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Eypp = Eyp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ezpp = Ezp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bxpp = Bxp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bypp = Byp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bzpp = Bzp.dataPtr(); // Loop over the particles and update their momentum const Real q = this->charge; @@ -1623,23 +1623,23 @@ PhysicalParticleContainer::PushP (int lev, Real dt, } } -void PhysicalParticleContainer::copy_attribs(WarpXParIter& pti,const Real* xp, - const Real* yp, const Real* zp) +void PhysicalParticleContainer::copy_attribs(WarpXParIter& pti,const ParticleReal* xp, + const ParticleReal* yp, const ParticleReal* 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(); + ParticleReal* AMREX_RESTRICT uxp = attribs[PIdx::ux].dataPtr(); + ParticleReal* AMREX_RESTRICT uyp = attribs[PIdx::uy].dataPtr(); + ParticleReal* AMREX_RESTRICT uzp = attribs[PIdx::uz].dataPtr(); 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(); + ParticleReal* AMREX_RESTRICT xpold = tmp_particle_data[lev][index][TmpIdx::xold ].dataPtr(); + ParticleReal* AMREX_RESTRICT ypold = tmp_particle_data[lev][index][TmpIdx::yold ].dataPtr(); + ParticleReal* AMREX_RESTRICT zpold = tmp_particle_data[lev][index][TmpIdx::zold ].dataPtr(); + ParticleReal* AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][TmpIdx::uxold].dataPtr(); + ParticleReal* AMREX_RESTRICT uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr(); + ParticleReal* AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); ParallelFor( np, [=] AMREX_GPU_DEVICE (long i) { @@ -1858,9 +1858,9 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, const Array4& by_arr = byfab->array(); const Array4& bz_arr = bzfab->array(); - const Real * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; - const Real * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; - const Real * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; + const ParticleReal * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; + const ParticleReal * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; + const ParticleReal * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; // Lower corner of tile box physical domain const std::array& xyzmin = WarpX::LowerCorner(box, gather_lev); @@ -2007,15 +2007,15 @@ PhysicalParticleContainer::buildIonizationMask (const amrex::MFIter& mfi, const // Otherwise, resize ionization_mask, and get poiters to attribs arrays. ionization_mask.resize(np); int * const AMREX_RESTRICT p_ionization_mask = ionization_mask.data(); - const Real * const AMREX_RESTRICT ux = soa.GetRealData(PIdx::ux).data(); - const Real * const AMREX_RESTRICT uy = soa.GetRealData(PIdx::uy).data(); - const Real * const AMREX_RESTRICT uz = soa.GetRealData(PIdx::uz).data(); - const Real * const AMREX_RESTRICT ex = soa.GetRealData(PIdx::Ex).data(); - const Real * const AMREX_RESTRICT ey = soa.GetRealData(PIdx::Ey).data(); - const Real * const AMREX_RESTRICT ez = soa.GetRealData(PIdx::Ez).data(); - const Real * const AMREX_RESTRICT bx = soa.GetRealData(PIdx::Bx).data(); - const Real * const AMREX_RESTRICT by = soa.GetRealData(PIdx::By).data(); - const Real * const AMREX_RESTRICT bz = soa.GetRealData(PIdx::Bz).data(); + const ParticleReal * const AMREX_RESTRICT ux = soa.GetRealData(PIdx::ux).data(); + const ParticleReal * const AMREX_RESTRICT uy = soa.GetRealData(PIdx::uy).data(); + const ParticleReal * const AMREX_RESTRICT uz = soa.GetRealData(PIdx::uz).data(); + const ParticleReal * const AMREX_RESTRICT ex = soa.GetRealData(PIdx::Ex).data(); + const ParticleReal * const AMREX_RESTRICT ey = soa.GetRealData(PIdx::Ey).data(); + const ParticleReal * const AMREX_RESTRICT ez = soa.GetRealData(PIdx::Ez).data(); + const ParticleReal * const AMREX_RESTRICT bx = soa.GetRealData(PIdx::Bx).data(); + const ParticleReal * const AMREX_RESTRICT by = soa.GetRealData(PIdx::By).data(); + const ParticleReal * const AMREX_RESTRICT bz = soa.GetRealData(PIdx::Bz).data(); int* ion_lev = soa.GetIntData(particle_icomps["ionization_level"]).data(); Real c = PhysConst::c; diff --git a/Source/Particles/Pusher/GetAndSetPosition.H b/Source/Particles/Pusher/GetAndSetPosition.H index 3c74baeb2..f0dfa4c83 100644 --- a/Source/Particles/Pusher/GetAndSetPosition.H +++ b/Source/Particles/Pusher/GetAndSetPosition.H @@ -11,7 +11,7 @@ * and stores them in the variables `x`, `y`, `z`. */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void GetPosition( - amrex::Real& x, amrex::Real& y, amrex::Real& z, + amrex::ParticleReal& x, amrex::ParticleReal& y, amrex::ParticleReal& z, const WarpXParticleContainer::ParticleType& p) { #if (AMREX_SPACEDIM==3) @@ -20,7 +20,7 @@ void GetPosition( z = p.pos(2); #else x = p.pos(0); - y = std::numeric_limits::quiet_NaN(); + y = std::numeric_limits::quiet_NaN(); z = p.pos(1); #endif } @@ -30,7 +30,7 @@ void GetPosition( AMREX_GPU_HOST_DEVICE AMREX_INLINE void SetPosition( WarpXParticleContainer::ParticleType& p, - const amrex::Real x, const amrex::Real y, const amrex::Real z) + const amrex::ParticleReal x, const amrex::ParticleReal y, const amrex::ParticleReal z) { #if (AMREX_SPACEDIM==3) p.pos(0) = x; @@ -49,10 +49,10 @@ void SetPosition( * and store them in the variables `x`, `y`, `z` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void GetCartesianPositionFromCylindrical( - amrex::Real& x, amrex::Real& y, amrex::Real& z, - const WarpXParticleContainer::ParticleType& p, const amrex::Real theta) + amrex::ParticleReal& x, amrex::ParticleReal& y, amrex::ParticleReal& z, + const WarpXParticleContainer::ParticleType& p, const amrex::ParticleReal theta) { - const amrex::Real r = p.pos(0); + const amrex::ParticleReal r = p.pos(0); x = r*std::cos(theta); y = r*std::sin(theta); z = p.pos(1); @@ -63,8 +63,8 @@ void GetCartesianPositionFromCylindrical( * from the values of `x`, `y`, `z` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void SetCylindricalPositionFromCartesian( - WarpXParticleContainer::ParticleType& p, amrex::Real& theta, - const amrex::Real x, const amrex::Real y, const amrex::Real z ) + WarpXParticleContainer::ParticleType& p, amrex::ParticleReal& theta, + const amrex::ParticleReal x, const amrex::ParticleReal y, const amrex::ParticleReal z ) { theta = std::atan2(y, x); p.pos(0) = std::sqrt(x*x + y*y); diff --git a/Source/Particles/Pusher/UpdateMomentumBoris.H b/Source/Particles/Pusher/UpdateMomentumBoris.H index a33058347..205cc9a71 100644 --- a/Source/Particles/Pusher/UpdateMomentumBoris.H +++ b/Source/Particles/Pusher/UpdateMomentumBoris.H @@ -7,9 +7,9 @@ * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdateMomentumBoris( - amrex::Real& ux, amrex::Real& uy, amrex::Real& uz, - const amrex::Real Ex, const amrex::Real Ey, const amrex::Real Ez, - const amrex::Real Bx, const amrex::Real By, const amrex::Real Bz, + amrex::ParticleReal& ux, amrex::ParticleReal& uy, amrex::ParticleReal& uz, + const amrex::ParticleReal Ex, const amrex::ParticleReal Ey, const amrex::ParticleReal Ez, + const amrex::ParticleReal Bx, const amrex::ParticleReal By, const amrex::ParticleReal Bz, const amrex::Real q, const amrex::Real m, const amrex::Real dt ) { const amrex::Real econst = 0.5*q*dt/m; diff --git a/Source/Particles/Pusher/UpdateMomentumVay.H b/Source/Particles/Pusher/UpdateMomentumVay.H index 1f0f19e63..433a891c5 100644 --- a/Source/Particles/Pusher/UpdateMomentumVay.H +++ b/Source/Particles/Pusher/UpdateMomentumVay.H @@ -9,9 +9,9 @@ * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdateMomentumVay( - amrex::Real& ux, amrex::Real& uy, amrex::Real& uz, - const amrex::Real Ex, const amrex::Real Ey, const amrex::Real Ez, - const amrex::Real Bx, const amrex::Real By, const amrex::Real Bz, + amrex::ParticleReal& ux, amrex::ParticleReal& uy, amrex::ParticleReal& uz, + const amrex::ParticleReal Ex, const amrex::ParticleReal Ey, const amrex::ParticleReal Ez, + const amrex::ParticleReal Bx, const amrex::ParticleReal By, const amrex::ParticleReal Bz, const amrex::Real q, const amrex::Real m, const amrex::Real dt ) { // Constants diff --git a/Source/Particles/Pusher/UpdatePosition.H b/Source/Particles/Pusher/UpdatePosition.H index a9df63a30..da0e9cdf9 100644 --- a/Source/Particles/Pusher/UpdatePosition.H +++ b/Source/Particles/Pusher/UpdatePosition.H @@ -9,8 +9,8 @@ * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdatePosition( - amrex::Real& x, amrex::Real& y, amrex::Real& z, - const amrex::Real ux, const amrex::Real uy, const amrex::Real uz, + amrex::ParticleReal& x, amrex::ParticleReal& y, amrex::ParticleReal& z, + const amrex::ParticleReal ux, const amrex::ParticleReal uy, const amrex::ParticleReal uz, const amrex::Real dt ) { diff --git a/Source/Particles/Pusher/UpdatePositionPhoton.H b/Source/Particles/Pusher/UpdatePositionPhoton.H index bd6e6cd21..f95c2b09d 100644 --- a/Source/Particles/Pusher/UpdatePositionPhoton.H +++ b/Source/Particles/Pusher/UpdatePositionPhoton.H @@ -10,8 +10,8 @@ * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdatePositionPhoton( - amrex::Real& x, amrex::Real& y, amrex::Real& z, - const amrex::Real ux, const amrex::Real uy, const amrex::Real uz, + amrex::ParticleReal& x, amrex::ParticleReal& y, amrex::ParticleReal& z, + const amrex::ParticleReal ux, const amrex::ParticleReal uy, const amrex::ParticleReal uz, const amrex::Real dt ) { // Compute speed of light over inverse of momentum modulus diff --git a/Source/Particles/RigidInjectedParticleContainer.H b/Source/Particles/RigidInjectedParticleContainer.H index be3dd21f9..3abbb4afe 100644 --- a/Source/Particles/RigidInjectedParticleContainer.H +++ b/Source/Particles/RigidInjectedParticleContainer.H @@ -44,9 +44,9 @@ public: DtType a_dt_type=DtType::Full) override; virtual void PushPX(WarpXParIter& pti, - amrex::Cuda::ManagedDeviceVector& xp, - amrex::Cuda::ManagedDeviceVector& yp, - amrex::Cuda::ManagedDeviceVector& zp, + amrex::Cuda::ManagedDeviceVector& xp, + amrex::Cuda::ManagedDeviceVector& yp, + amrex::Cuda::ManagedDeviceVector& zp, amrex::Real dt, DtType a_dt_type=DtType::Full) override; virtual void PushP (int lev, amrex::Real dt, diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp index 7d129fc01..891ade76d 100644 --- a/Source/Particles/RigidInjectedParticleContainer.cpp +++ b/Source/Particles/RigidInjectedParticleContainer.cpp @@ -76,7 +76,7 @@ RigidInjectedParticleContainer::RemapParticles() // Note that the particles are already in the boosted frame. // This value is saved to advance the particles not injected yet - Cuda::ManagedDeviceVector xp, yp, zp; + Cuda::ManagedDeviceVector xp, yp, zp; for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { @@ -136,7 +136,7 @@ RigidInjectedParticleContainer::BoostandRemapParticles() #pragma omp parallel #endif { - Cuda::ManagedDeviceVector xp, yp, zp; + Cuda::ManagedDeviceVector xp, yp, zp; for (WarpXParIter pti(*this, 0); pti.isValid(); ++pti) { @@ -207,9 +207,9 @@ RigidInjectedParticleContainer::BoostandRemapParticles() void RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, - Cuda::ManagedDeviceVector& xp, - Cuda::ManagedDeviceVector& yp, - Cuda::ManagedDeviceVector& zp, + Cuda::ManagedDeviceVector& xp, + Cuda::ManagedDeviceVector& yp, + Cuda::ManagedDeviceVector& zp, Real dt, DtType a_dt_type) { @@ -220,21 +220,21 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, auto& uzp = attribs[PIdx::uz]; // Save the position and momenta, making copies - Cuda::ManagedDeviceVector xp_save, yp_save, zp_save; + Cuda::ManagedDeviceVector xp_save, yp_save, zp_save; RealVector uxp_save, uyp_save, uzp_save; - Real* const AMREX_RESTRICT x = xp.dataPtr(); - Real* const AMREX_RESTRICT y = yp.dataPtr(); - Real* const AMREX_RESTRICT z = zp.dataPtr(); - Real* const AMREX_RESTRICT ux = uxp.dataPtr(); - Real* const AMREX_RESTRICT uy = uyp.dataPtr(); - Real* const AMREX_RESTRICT uz = uzp.dataPtr(); - Real* const AMREX_RESTRICT Exp = attribs[PIdx::Ex].dataPtr(); - Real* const AMREX_RESTRICT Eyp = attribs[PIdx::Ey].dataPtr(); - Real* const AMREX_RESTRICT Ezp = attribs[PIdx::Ez].dataPtr(); - Real* const AMREX_RESTRICT Bxp = attribs[PIdx::Bx].dataPtr(); - Real* const AMREX_RESTRICT Byp = attribs[PIdx::By].dataPtr(); - Real* const AMREX_RESTRICT Bzp = attribs[PIdx::Bz].dataPtr(); + ParticleReal* const AMREX_RESTRICT x = xp.dataPtr(); + ParticleReal* const AMREX_RESTRICT y = yp.dataPtr(); + ParticleReal* const AMREX_RESTRICT z = zp.dataPtr(); + ParticleReal* const AMREX_RESTRICT ux = uxp.dataPtr(); + ParticleReal* const AMREX_RESTRICT uy = uyp.dataPtr(); + ParticleReal* const AMREX_RESTRICT uz = uzp.dataPtr(); + ParticleReal* const AMREX_RESTRICT Exp = attribs[PIdx::Ex].dataPtr(); + ParticleReal* const AMREX_RESTRICT Eyp = attribs[PIdx::Ey].dataPtr(); + ParticleReal* const AMREX_RESTRICT Ezp = attribs[PIdx::Ez].dataPtr(); + ParticleReal* const AMREX_RESTRICT Bxp = attribs[PIdx::Bx].dataPtr(); + ParticleReal* const AMREX_RESTRICT Byp = attribs[PIdx::By].dataPtr(); + ParticleReal* const AMREX_RESTRICT Bzp = attribs[PIdx::Bz].dataPtr(); if (!done_injecting_lev) { // If the old values are not already saved, create copies here. @@ -271,12 +271,12 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, if (!done_injecting_lev) { - 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(); + ParticleReal* AMREX_RESTRICT x_save = xp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT y_save = yp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT z_save = zp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT ux_save = uxp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT uy_save = uyp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT uz_save = uzp_save.dataPtr(); // Undo the push for particles not injected yet. // The zp are advanced a fixed amount. @@ -415,16 +415,16 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, // This wraps the momentum advance so that inheritors can modify the call. // Extract pointers to the different particle quantities - const Real* const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); - Real* const AMREX_RESTRICT uxpp = uxp.dataPtr(); - Real* const AMREX_RESTRICT uypp = uyp.dataPtr(); - Real* const AMREX_RESTRICT uzpp = uzp.dataPtr(); - const Real* const AMREX_RESTRICT Expp = Exp.dataPtr(); - const Real* const AMREX_RESTRICT Eypp = Eyp.dataPtr(); - const Real* const AMREX_RESTRICT Ezpp = Ezp.dataPtr(); - const Real* const AMREX_RESTRICT Bxpp = Bxp.dataPtr(); - const Real* const AMREX_RESTRICT Bypp = Byp.dataPtr(); - const Real* const AMREX_RESTRICT Bzpp = Bzp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); + ParticleReal* const AMREX_RESTRICT uxpp = uxp.dataPtr(); + ParticleReal* const AMREX_RESTRICT uypp = uyp.dataPtr(); + ParticleReal* const AMREX_RESTRICT uzpp = uzp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Expp = Exp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Eypp = Eyp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ezpp = Ezp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bxpp = Bxp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bypp = Byp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bzpp = Bzp.dataPtr(); // Loop over the particles and update their momentum const Real q = this->charge; @@ -450,10 +450,10 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, // Undo the push for particles not injected yet. // It is assumed that PushP will only be called on the first and last steps // and that no particles will cross zinject_plane. - const Real* const AMREX_RESTRICT ux_save = uxp_save.dataPtr(); - const Real* const AMREX_RESTRICT uy_save = uyp_save.dataPtr(); - const Real* const AMREX_RESTRICT uz_save = uzp_save.dataPtr(); - const Real zz = zinject_plane_levels[lev]; + const ParticleReal* const AMREX_RESTRICT ux_save = uxp_save.dataPtr(); + const ParticleReal* const AMREX_RESTRICT uy_save = uyp_save.dataPtr(); + const ParticleReal* const AMREX_RESTRICT uz_save = uzp_save.dataPtr(); + const ParticleReal zz = zinject_plane_levels[lev]; amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { if (zp[i] <= zz) { diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index ee75bc511..7b0d2d1d0 100644 --- a/Source/Particles/WarpXParticleContainer.H +++ b/Source/Particles/WarpXParticleContainer.H @@ -68,12 +68,12 @@ public: WarpXParIter (ContainerType& pc, int level); #if (AMREX_SPACEDIM == 2) - void GetPosition (amrex::Cuda::ManagedDeviceVector& x, - amrex::Cuda::ManagedDeviceVector& y, - amrex::Cuda::ManagedDeviceVector& z) const; - void SetPosition (const amrex::Cuda::ManagedDeviceVector& x, - const amrex::Cuda::ManagedDeviceVector& y, - const amrex::Cuda::ManagedDeviceVector& z); + void GetPosition (amrex::Cuda::ManagedDeviceVector& x, + amrex::Cuda::ManagedDeviceVector& y, + amrex::Cuda::ManagedDeviceVector& z) const; + void SetPosition (const amrex::Cuda::ManagedDeviceVector& x, + const amrex::Cuda::ManagedDeviceVector& y, + const amrex::Cuda::ManagedDeviceVector& z); #endif const std::array& GetAttribs () const { return GetStructOfArrays().GetRealData(); @@ -104,7 +104,7 @@ class WarpXParticleContainer public: friend MultiParticleContainer; - // amrex::StructOfArrays with DiagIdx::nattribs amrex::Real components + // amrex::StructOfArrays with DiagIdx::nattribs amrex::ParticleReal components // and 0 int components for the particle data. using DiagnosticParticleData = amrex::StructOfArrays; // DiagnosticParticles is a vector, with one element per MR level. @@ -232,17 +232,17 @@ public: amrex::Real maxParticleVelocity(bool local = false); void AddNParticles (int lev, - int n, const amrex::Real* x, const amrex::Real* y, const amrex::Real* z, - const amrex::Real* vx, const amrex::Real* vy, const amrex::Real* vz, - int nattr, const amrex::Real* attr, int uniqueparticles, int id=-1); + int n, const amrex::ParticleReal* x, const amrex::ParticleReal* y, const amrex::ParticleReal* z, + const amrex::ParticleReal* vx, const amrex::ParticleReal* vy, const amrex::ParticleReal* vz, + int nattr, const amrex::ParticleReal* attr, int uniqueparticles, int id=-1); void AddOneParticle (int lev, int grid, int tile, - amrex::Real x, amrex::Real y, amrex::Real z, - std::array& attribs); + amrex::ParticleReal x, amrex::ParticleReal y, amrex::ParticleReal z, + std::array& attribs); void AddOneParticle (ParticleTileType& particle_tile, - amrex::Real x, amrex::Real y, amrex::Real z, - std::array& attribs); + amrex::ParticleReal x, amrex::ParticleReal y, amrex::ParticleReal z, + std::array& attribs); virtual void ReadHeader (std::istream& is); @@ -326,7 +326,7 @@ protected: amrex::Vector local_jy; amrex::Vector local_jz; - using DataContainer = amrex::Gpu::ManagedDeviceVector; + using DataContainer = amrex::Gpu::ManagedDeviceVector; using PairIndex = std::pair; amrex::Vector m_xp, m_yp, m_zp; diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 4a8684a6a..65a82f233 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -25,7 +25,9 @@ WarpXParIter::WarpXParIter (ContainerType& pc, int level) #if (AMREX_SPACEDIM == 2) void -WarpXParIter::GetPosition (Cuda::ManagedDeviceVector& x, Cuda::ManagedDeviceVector& y, Cuda::ManagedDeviceVector& z) const +WarpXParIter::GetPosition (Gpu::ManagedDeviceVector& x, + Gpu::ManagedDeviceVector& y, + Gpu::ManagedDeviceVector& z) const { amrex::ParIter<0,0,PIdx::nattribs>::GetPosition(x, z); #ifdef WARPX_DIM_RZ @@ -38,17 +40,19 @@ WarpXParIter::GetPosition (Cuda::ManagedDeviceVector& x, Cuda::ManagedDevi x[i] = x[i]*std::cos(theta[i]); } #else - y.resize(x.size(), std::numeric_limits::quiet_NaN()); + y.resize(x.size(), std::numeric_limits::quiet_NaN()); #endif } void -WarpXParIter::SetPosition (const Cuda::ManagedDeviceVector& x, const Cuda::ManagedDeviceVector& y, const Cuda::ManagedDeviceVector& z) +WarpXParIter::SetPosition (const Gpu::ManagedDeviceVector& x, + const Gpu::ManagedDeviceVector& y, + const Gpu::ManagedDeviceVector& z) { #ifdef WARPX_DIM_RZ auto& attribs = GetAttribs(); auto& theta = attribs[PIdx::theta]; - Cuda::ManagedDeviceVector r(x.size()); + Gpu::ManagedDeviceVector r(x.size()); for (unsigned int i=0 ; i < x.size() ; i++) { theta[i] = std::atan2(y[i], x[i]); r[i] = std::sqrt(x[i]*x[i] + y[i]*y[i]); @@ -132,8 +136,8 @@ WarpXParticleContainer::AllocData () void WarpXParticleContainer::AddOneParticle (int lev, int grid, int tile, - Real x, Real y, Real z, - std::array& attribs) + ParticleReal x, ParticleReal y, ParticleReal z, + std::array& attribs) { auto& particle_tile = DefineAndReturnParticleTile(lev, grid, tile); AddOneParticle(particle_tile, x, y, z, attribs); @@ -141,8 +145,8 @@ WarpXParticleContainer::AddOneParticle (int lev, int grid, int tile, void WarpXParticleContainer::AddOneParticle (ParticleTileType& particle_tile, - Real x, Real y, Real z, - std::array& attribs) + ParticleReal x, ParticleReal y, ParticleReal z, + std::array& attribs) { ParticleType p; p.id() = ParticleType::NextID(); @@ -171,12 +175,12 @@ WarpXParticleContainer::AddOneParticle (ParticleTileType& particle_tile, void WarpXParticleContainer::AddNParticles (int lev, - int n, const Real* x, const Real* y, const Real* z, - const Real* vx, const Real* vy, const Real* vz, - int nattr, const Real* attr, int uniqueparticles, int id) + int n, const ParticleReal* x, const ParticleReal* y, const ParticleReal* z, + const ParticleReal* vx, const ParticleReal* vy, const ParticleReal* vz, + int nattr, const ParticleReal* attr, int uniqueparticles, int id) { BL_ASSERT(nattr == 1); - const Real* weight = attr; + const ParticleReal* weight = attr; int ibegin, iend; if (uniqueparticles) { @@ -204,7 +208,7 @@ WarpXParticleContainer::AddNParticles (int lev, std::size_t np = iend-ibegin; #ifdef WARPX_DIM_RZ - Vector theta(np); + Vector theta(np); #endif for (int i = ibegin; i < iend; ++i) @@ -362,9 +366,9 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, // CPU, tiling: deposit into local_jx // (same for jx and jz) - Real* AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; - Real* AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; - Real* AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow @@ -503,9 +507,9 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector& wp, // GPU, no tiling: deposit directly in rho // CPU, tiling: deposit into local_rho - Real* AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; - Real* AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; - Real* AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow @@ -731,7 +735,7 @@ std::array WarpXParticleContainer::meanParticleVelocity(bool local) { Real WarpXParticleContainer::maxParticleVelocity(bool local) { - amrex::Real max_v = 0.0; + amrex::ParticleReal max_v = 0.0; for (int lev = 0; lev <= finestLevel(); ++lev) { @@ -745,12 +749,12 @@ Real WarpXParticleContainer::maxParticleVelocity(bool local) { auto& uy = pti.GetAttribs(PIdx::uy); auto& uz = pti.GetAttribs(PIdx::uz); for (unsigned long i = 0; i < ux.size(); i++) { - max_v = std::max(max_v, sqrt(ux[i]*ux[i] + uy[i]*uy[i] + uz[i]*uz[i])); + max_v = std::max(max_v, std::sqrt(ux[i]*ux[i] + uy[i]*uy[i] + uz[i]*uz[i])); } } } - if (!local) ParallelDescriptor::ReduceRealMax(max_v); + if (!local) ParallelAllReduce::Max(max_v, ParallelDescriptor::Communicator()); return max_v; } @@ -819,17 +823,17 @@ WarpXParticleContainer::PushX (int lev, Real dt) ParticleType * AMREX_RESTRICT pstructs = &(pti.GetArrayOfStructs()[0]); // - momenta are stored as a struct of array, in `attribs` auto& attribs = pti.GetAttribs(); - Real* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + ParticleReal* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); #ifdef WARPX_DIM_RZ - Real* AMREX_RESTRICT theta = attribs[PIdx::theta].dataPtr(); + ParticleReal* AMREX_RESTRICT theta = attribs[PIdx::theta].dataPtr(); #endif // Loop over the particles and update their position amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleType& p = pstructs[i]; // Particle object that gets updated - Real x, y, z; // Temporary variables + ParticleReal x, y, z; // Temporary variables #ifndef WARPX_DIM_RZ GetPosition( x, y, z, p ); // Initialize x, y, z UpdatePosition( x, y, z, ux[i], uy[i], uz[i], dt); diff --git a/Source/Particles/deposit_cic.F90 b/Source/Particles/deposit_cic.F90 index 1fe80016f..2831ce96b 100644 --- a/Source/Particles/deposit_cic.F90 +++ b/Source/Particles/deposit_cic.F90 @@ -1,7 +1,7 @@ module warpx_ES_deposit_cic use iso_c_binding - use amrex_fort_module, only : amrex_real + use amrex_fort_module, only : amrex_real, amrex_particle_real implicit none @@ -28,8 +28,8 @@ contains ng) & bind(c,name='warpx_deposit_cic_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) - real(amrex_real), intent(in) :: weights(np) + real(amrex_particle_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: weights(np) real(amrex_real), intent(in) :: charge integer, intent(in) :: lo(3) integer, intent(in) :: hi(3) @@ -86,8 +86,8 @@ contains ng) & bind(c,name='warpx_deposit_cic_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) - real(amrex_real), intent(in) :: weights(np) + real(amrex_particle_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: weights(np) real(amrex_real), intent(in) :: charge integer, intent(in) :: lo(2) integer, intent(in) :: hi(2) diff --git a/Source/Particles/interpolate_cic.F90 b/Source/Particles/interpolate_cic.F90 index 005ab35f4..3eb361d2f 100644 --- a/Source/Particles/interpolate_cic.F90 +++ b/Source/Particles/interpolate_cic.F90 @@ -1,7 +1,7 @@ module warpx_ES_interpolate_cic use iso_c_binding - use amrex_fort_module, only : amrex_real + use amrex_fort_module, only : amrex_real, amrex_particle_real implicit none @@ -31,7 +31,7 @@ contains lo, hi, plo, dx, ng) & bind(c,name='warpx_interpolate_cic_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: particles(ns,np) real(amrex_real), intent(inout) :: Ex_p(np), Ey_p(np), Ez_p(np) integer, intent(in) :: ng integer, intent(in) :: lo(3) @@ -103,7 +103,7 @@ contains lo, hi, plo, dx, ng) & bind(c,name='warpx_interpolate_cic_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: particles(ns,np) real(amrex_real), intent(inout) :: Ex_p(np), Ey_p(np) integer, intent(in) :: ng integer, intent(in) :: lo(2) @@ -157,7 +157,7 @@ contains plo, ng, lev) & bind(c,name='warpx_interpolate_cic_two_levels_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: particles(ns,np) real(amrex_real), intent(inout) :: Ex_p(np), Ey_p(np), Ez_p(np) integer, intent(in) :: ng, lev integer, intent(in) :: lo(3), hi(3) @@ -290,7 +290,7 @@ contains plo, ng, lev) & bind(c,name='warpx_interpolate_cic_two_levels_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: particles(ns,np) real(amrex_real), intent(inout) :: Ex_p(np), Ey_p(np) integer, intent(in) :: ng, lev integer, intent(in) :: lo(2), hi(2) diff --git a/Source/Particles/push_particles_ES.F90 b/Source/Particles/push_particles_ES.F90 index 53dd3c181..b84f48d5f 100644 --- a/Source/Particles/push_particles_ES.F90 +++ b/Source/Particles/push_particles_ES.F90 @@ -1,7 +1,7 @@ module warpx_ES_push_particles use iso_c_binding - use amrex_fort_module, only : amrex_real + use amrex_fort_module, only : amrex_real, amrex_particle_real implicit none @@ -36,8 +36,8 @@ contains prob_lo, prob_hi) & bind(c,name='warpx_push_leapfrog_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(inout) :: particles(ns,np) - real(amrex_real), intent(inout) :: vx_p(np), vy_p(np), vz_p(np) + real(amrex_particle_real), intent(inout) :: particles(ns,np) + real(amrex_particle_real), intent(inout) :: vx_p(np), vy_p(np), vz_p(np) real(amrex_real), intent(in) :: Ex_p(np), Ey_p(np), Ez_p(np) real(amrex_real), intent(in) :: charge real(amrex_real), intent(in) :: mass @@ -100,8 +100,8 @@ contains prob_lo, prob_hi) & bind(c,name='warpx_push_leapfrog_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(inout) :: particles(ns,np) - real(amrex_real), intent(inout) :: vx_p(np), vy_p(np) + real(amrex_particle_real), intent(inout) :: particles(ns,np) + real(amrex_particle_real), intent(inout) :: vx_p(np), vy_p(np) real(amrex_real), intent(in) :: Ex_p(np), Ey_p(np) real(amrex_real), intent(in) :: charge real(amrex_real), intent(in) :: mass @@ -167,8 +167,8 @@ contains prob_lo, prob_hi) & bind(c,name='warpx_push_leapfrog_positions_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(inout) :: particles(ns,np) - real(amrex_real), intent(inout) :: vx_p(np), vy_p(np), vz_p(np) + real(amrex_particle_real), intent(inout) :: particles(ns,np) + real(amrex_particle_real), intent(inout) :: vx_p(np), vy_p(np), vz_p(np) real(amrex_real), intent(in) :: dt real(amrex_real), intent(in) :: prob_lo(3), prob_hi(3) @@ -219,8 +219,8 @@ contains prob_lo, prob_hi) & bind(c,name='warpx_push_leapfrog_positions_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(inout) :: particles(ns,np) - real(amrex_real), intent(inout) :: vx_p(np), vy_p(np) + real(amrex_particle_real), intent(inout) :: particles(ns,np) + real(amrex_particle_real), intent(inout) :: vx_p(np), vy_p(np) real(amrex_real), intent(in) :: dt real(amrex_real), intent(in) :: prob_lo(2), prob_hi(2) diff --git a/Source/WarpX.cpp b/Source/WarpX.cpp index a844e7aa0..62bbfc44e 100644 --- a/Source/WarpX.cpp +++ b/Source/WarpX.cpp @@ -149,8 +149,8 @@ WarpX::WarpX () #endif t_new.resize(nlevs_max, 0.0); - t_old.resize(nlevs_max, -1.e100); - dt.resize(nlevs_max, 1.e100); + t_old.resize(nlevs_max, std::numeric_limits::lowest()); + dt.resize(nlevs_max, std::numeric_limits::max()); // Particle Container mypc = std::unique_ptr (new MultiParticleContainer(this)); @@ -988,7 +988,7 @@ WarpX::LowerCorner(const Box& bx, int lev) #if (AMREX_SPACEDIM == 3) return { xyzmin[0], xyzmin[1], xyzmin[2] }; #elif (AMREX_SPACEDIM == 2) - return { xyzmin[0], -1.e100, xyzmin[1] }; + return { xyzmin[0], std::numeric_limits::lowest(), xyzmin[1] }; #endif } @@ -1000,7 +1000,7 @@ WarpX::UpperCorner(const Box& bx, int lev) #if (AMREX_SPACEDIM == 3) return { xyzmax[0], xyzmax[1], xyzmax[2] }; #elif (AMREX_SPACEDIM == 2) - return { xyzmax[0], 1.e100, xyzmax[1] }; + return { xyzmax[0], std::numeric_limits::max(), xyzmax[1] }; #endif } -- cgit v1.2.3 From bf1c8d22fd01aa35927b13cd42ebc6377df7e814 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Thu, 26 Sep 2019 13:57:07 -0700 Subject: put NCI filter into a separate function --- Source/Particles/PhysicalParticleContainer.H | 18 +++- Source/Particles/PhysicalParticleContainer.cpp | 121 +++++++++++++++---------- 2 files changed, 88 insertions(+), 51 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index b809acc45..e175e3bbe 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -1,12 +1,13 @@ #ifndef WARPX_PhysicalParticleContainer_H_ #define WARPX_PhysicalParticleContainer_H_ -#include +#include +#include +#include #include -#include -#include +#include class PhysicalParticleContainer : public WarpXParticleContainer @@ -144,6 +145,17 @@ public: virtual void ConvertUnits (ConvertDirection convert_dir) override; +void applyNCIFilter ( + int lev, const amrex::Box& box, + amrex::Elixir& exeli, amrex::Elixir& eyeli, amrex::Elixir& ezeli, + amrex::Elixir& bxeli, amrex::Elixir& byeli, amrex::Elixir& bzeli, + amrex::FArrayBox& filtered_Ex, amrex::FArrayBox& filtered_Ey, amrex::FArrayBox& filtered_Ez, + amrex::FArrayBox& filtered_Bx, amrex::FArrayBox& filtered_By, amrex::FArrayBox& filtered_Bz, + const amrex::FArrayBox& Ex_pti, const amrex::FArrayBox& Ey_pti, const amrex::FArrayBox& Ez_pti, + const amrex::FArrayBox& Bx_pti, const amrex::FArrayBox& By_pti, const amrex::FArrayBox& Bz_pti, + amrex::FArrayBox const * & exfab, amrex::FArrayBox const * & eyfab, amrex::FArrayBox const * & ezfab, + amrex::FArrayBox const * & bxfab, amrex::FArrayBox const * & byfab, amrex::FArrayBox const * & bzfab); + protected: std::string species_name; diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index e276fd5ef..da1ed506e 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -936,6 +936,73 @@ PhysicalParticleContainer::FieldGather (int lev, } } +void +PhysicalParticleContainer::applyNCIFilter ( + int lev, const Box& box, + Elixir& exeli, Elixir& eyeli, Elixir& ezeli, + Elixir& bxeli, Elixir& byeli, Elixir& bzeli, + FArrayBox& filtered_Ex, FArrayBox& filtered_Ey, FArrayBox& filtered_Ez, + FArrayBox& filtered_Bx, FArrayBox& filtered_By, FArrayBox& filtered_Bz, + const FArrayBox& Ex_pti, const FArrayBox& Ey_pti, const FArrayBox& Ez_pti, + const FArrayBox& Bx_pti, const FArrayBox& By_pti, const FArrayBox& Bz_pti, + FArrayBox const * & exfab, FArrayBox const * & eyfab, FArrayBox const * & ezfab, + FArrayBox const * & bxfab, FArrayBox const * & byfab, FArrayBox const * & bzfab) +{ + + // Get instances of NCI Godfrey filters + const auto& nci_godfrey_filter_exeybz = WarpX::GetInstance().nci_godfrey_filter_exeybz; + const auto& nci_godfrey_filter_bxbyez = WarpX::GetInstance().nci_godfrey_filter_bxbyez; + +#if (AMREX_SPACEDIM == 2) + const Box& tbox = amrex::grow(box,{static_cast(WarpX::nox), + static_cast(WarpX::noz)}); +#else + const Box& tbox = amrex::grow(box,{static_cast(WarpX::nox), + static_cast(WarpX::noy), + static_cast(WarpX::noz)}); +#endif + + // Filter Ex (Both 2D and 3D) + filtered_Ex.resize(amrex::convert(tbox,WarpX::Ex_nodal_flag)); + // Safeguard for GPU + exeli = filtered_Ex.elixir(); + // Apply filter on Ex, result stored in filtered_Ex + nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Ex, Ex_pti, filtered_Ex.box()); + // Update exfab reference + exfab = &filtered_Ex; + + // Filter Ez + filtered_Ez.resize(amrex::convert(tbox,WarpX::Ez_nodal_flag)); + ezeli = filtered_Ez.elixir(); + nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_Ez, Ez_pti, filtered_Ez.box()); + ezfab = &filtered_Ez; + + // Filter By + filtered_By.resize(amrex::convert(tbox,WarpX::By_nodal_flag)); + byeli = filtered_By.elixir(); + nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_By, By_pti, filtered_By.box()); + byfab = &filtered_By; +#if (AMREX_SPACEDIM == 3) + // Filter Ey + filtered_Ey.resize(amrex::convert(tbox,WarpX::Ey_nodal_flag)); + eyeli = filtered_Ey.elixir(); + nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Ey, Ey_pti, filtered_Ey.box()); + eyfab = &filtered_Ey; + + // Filter Bx + filtered_Bx.resize(amrex::convert(tbox,WarpX::Bx_nodal_flag)); + bxeli = filtered_Bx.elixir(); + nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_Bx, Bx_pti, filtered_Bx.box()); + bxfab = &filtered_Bx; + + // Filter Bz + filtered_Bz.resize(amrex::convert(tbox,WarpX::Bz_nodal_flag)); + bzeli = filtered_Bz.elixir(); + nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Bz, Bz_pti, filtered_Bz.box()); + bzfab = &filtered_Bz; +#endif +} + void PhysicalParticleContainer::Evolve (int lev, const MultiFab& Ex, const MultiFab& Ey, const MultiFab& Ez, @@ -958,7 +1025,7 @@ PhysicalParticleContainer::Evolve (int lev, // Get instances of NCI Godfrey filters const auto& nci_godfrey_filter_exeybz = WarpX::GetInstance().nci_godfrey_filter_exeybz; const auto& nci_godfrey_filter_bxbyez = WarpX::GetInstance().nci_godfrey_filter_bxbyez; - + BL_ASSERT(OnSameGrids(lev,jx)); MultiFab* cost = WarpX::getCosts(lev); @@ -1026,54 +1093,12 @@ PhysicalParticleContainer::Evolve (int lev, if (WarpX::use_fdtd_nci_corr) { -#if (AMREX_SPACEDIM == 2) - const Box& tbox = amrex::grow(pti.tilebox(),{static_cast(WarpX::nox), - static_cast(WarpX::noz)}); -#else - const Box& tbox = amrex::grow(pti.tilebox(),{static_cast(WarpX::nox), - static_cast(WarpX::noy), - static_cast(WarpX::noz)}); -#endif + applyNCIFilter(lev, pti.tilebox(), exeli, eyeli, ezeli, bxeli, byeli, bzeli, + filtered_Ex, filtered_Ey, filtered_Ez, + filtered_Bx, filtered_By, filtered_Bz, + Ex[pti], Ey[pti], Ez[pti], Bx[pti], By[pti], Bz[pti], + exfab, eyfab, ezfab, bxfab, byfab, bzfab); - // Filter Ex (Both 2D and 3D) - filtered_Ex.resize(amrex::convert(tbox,WarpX::Ex_nodal_flag)); - // Safeguard for GPU - exeli = filtered_Ex.elixir(); - // Apply filter on Ex, result stored in filtered_Ex - nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Ex, Ex[pti], filtered_Ex.box()); - // Update exfab reference - exfab = &filtered_Ex; - - // Filter Ez - filtered_Ez.resize(amrex::convert(tbox,WarpX::Ez_nodal_flag)); - ezeli = filtered_Ez.elixir(); - nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_Ez, Ez[pti], filtered_Ez.box()); - ezfab = &filtered_Ez; - - // Filter By - filtered_By.resize(amrex::convert(tbox,WarpX::By_nodal_flag)); - byeli = filtered_By.elixir(); - nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_By, By[pti], filtered_By.box()); - byfab = &filtered_By; -#if (AMREX_SPACEDIM == 3) - // Filter Ey - filtered_Ey.resize(amrex::convert(tbox,WarpX::Ey_nodal_flag)); - eyeli = filtered_Ey.elixir(); - nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Ey, Ey[pti], filtered_Ey.box()); - eyfab = &filtered_Ey; - - // Filter Bx - filtered_Bx.resize(amrex::convert(tbox,WarpX::Bx_nodal_flag)); - bxeli = filtered_Bx.elixir(); - nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_Bx, Bx[pti], filtered_Bx.box()); - bxfab = &filtered_Bx; - - // Filter Bz - filtered_Bz.resize(amrex::convert(tbox,WarpX::Bz_nodal_flag)); - bzeli = filtered_Bz.elixir(); - nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Bz, Bz[pti], filtered_Bz.box()); - bzfab = &filtered_Bz; -#endif } Exp.assign(np,0.0); -- cgit v1.2.3 From effa03e8d28fed99c8afac392538548afacfd915 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Thu, 26 Sep 2019 14:32:32 -0700 Subject: use clearer argument names --- Source/Particles/PhysicalParticleContainer.H | 23 +-- Source/Particles/PhysicalParticleContainer.cpp | 189 ++++++++++--------------- 2 files changed, 87 insertions(+), 125 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index e175e3bbe..392f9df1c 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -145,16 +145,19 @@ public: virtual void ConvertUnits (ConvertDirection convert_dir) override; -void applyNCIFilter ( - int lev, const amrex::Box& box, - amrex::Elixir& exeli, amrex::Elixir& eyeli, amrex::Elixir& ezeli, - amrex::Elixir& bxeli, amrex::Elixir& byeli, amrex::Elixir& bzeli, - amrex::FArrayBox& filtered_Ex, amrex::FArrayBox& filtered_Ey, amrex::FArrayBox& filtered_Ez, - amrex::FArrayBox& filtered_Bx, amrex::FArrayBox& filtered_By, amrex::FArrayBox& filtered_Bz, - const amrex::FArrayBox& Ex_pti, const amrex::FArrayBox& Ey_pti, const amrex::FArrayBox& Ez_pti, - const amrex::FArrayBox& Bx_pti, const amrex::FArrayBox& By_pti, const amrex::FArrayBox& Bz_pti, - amrex::FArrayBox const * & exfab, amrex::FArrayBox const * & eyfab, amrex::FArrayBox const * & ezfab, - amrex::FArrayBox const * & bxfab, amrex::FArrayBox const * & byfab, amrex::FArrayBox const * & bzfab); + void applyNCIFilter ( + int lev, const amrex::Box& box, + amrex::Elixir& exeli, amrex::Elixir& eyeli, amrex::Elixir& ezeli, + amrex::Elixir& bxeli, amrex::Elixir& byeli, amrex::Elixir& bzeli, + amrex::FArrayBox& filtered_Ex, amrex::FArrayBox& filtered_Ey, + amrex::FArrayBox& filtered_Ez, amrex::FArrayBox& filtered_Bx, + amrex::FArrayBox& filtered_By, amrex::FArrayBox& filtered_Bz, + const amrex::FArrayBox& Ex, const amrex::FArrayBox& Ey, + const amrex::FArrayBox& Ez, const amrex::FArrayBox& Bx, + const amrex::FArrayBox& By, const amrex::FArrayBox& Bz, + amrex::FArrayBox const * & exfab, amrex::FArrayBox const * & eyfab, + amrex::FArrayBox const * & ezfab, amrex::FArrayBox const * & bxfab, + amrex::FArrayBox const * & byfab, amrex::FArrayBox const * & bzfab); protected: diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index da1ed506e..b0da0e887 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -936,73 +936,6 @@ PhysicalParticleContainer::FieldGather (int lev, } } -void -PhysicalParticleContainer::applyNCIFilter ( - int lev, const Box& box, - Elixir& exeli, Elixir& eyeli, Elixir& ezeli, - Elixir& bxeli, Elixir& byeli, Elixir& bzeli, - FArrayBox& filtered_Ex, FArrayBox& filtered_Ey, FArrayBox& filtered_Ez, - FArrayBox& filtered_Bx, FArrayBox& filtered_By, FArrayBox& filtered_Bz, - const FArrayBox& Ex_pti, const FArrayBox& Ey_pti, const FArrayBox& Ez_pti, - const FArrayBox& Bx_pti, const FArrayBox& By_pti, const FArrayBox& Bz_pti, - FArrayBox const * & exfab, FArrayBox const * & eyfab, FArrayBox const * & ezfab, - FArrayBox const * & bxfab, FArrayBox const * & byfab, FArrayBox const * & bzfab) -{ - - // Get instances of NCI Godfrey filters - const auto& nci_godfrey_filter_exeybz = WarpX::GetInstance().nci_godfrey_filter_exeybz; - const auto& nci_godfrey_filter_bxbyez = WarpX::GetInstance().nci_godfrey_filter_bxbyez; - -#if (AMREX_SPACEDIM == 2) - const Box& tbox = amrex::grow(box,{static_cast(WarpX::nox), - static_cast(WarpX::noz)}); -#else - const Box& tbox = amrex::grow(box,{static_cast(WarpX::nox), - static_cast(WarpX::noy), - static_cast(WarpX::noz)}); -#endif - - // Filter Ex (Both 2D and 3D) - filtered_Ex.resize(amrex::convert(tbox,WarpX::Ex_nodal_flag)); - // Safeguard for GPU - exeli = filtered_Ex.elixir(); - // Apply filter on Ex, result stored in filtered_Ex - nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Ex, Ex_pti, filtered_Ex.box()); - // Update exfab reference - exfab = &filtered_Ex; - - // Filter Ez - filtered_Ez.resize(amrex::convert(tbox,WarpX::Ez_nodal_flag)); - ezeli = filtered_Ez.elixir(); - nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_Ez, Ez_pti, filtered_Ez.box()); - ezfab = &filtered_Ez; - - // Filter By - filtered_By.resize(amrex::convert(tbox,WarpX::By_nodal_flag)); - byeli = filtered_By.elixir(); - nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_By, By_pti, filtered_By.box()); - byfab = &filtered_By; -#if (AMREX_SPACEDIM == 3) - // Filter Ey - filtered_Ey.resize(amrex::convert(tbox,WarpX::Ey_nodal_flag)); - eyeli = filtered_Ey.elixir(); - nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Ey, Ey_pti, filtered_Ey.box()); - eyfab = &filtered_Ey; - - // Filter Bx - filtered_Bx.resize(amrex::convert(tbox,WarpX::Bx_nodal_flag)); - bxeli = filtered_Bx.elixir(); - nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_Bx, Bx_pti, filtered_Bx.box()); - bxfab = &filtered_Bx; - - // Filter Bz - filtered_Bz.resize(amrex::convert(tbox,WarpX::Bz_nodal_flag)); - bzeli = filtered_Bz.elixir(); - nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Bz, Bz_pti, filtered_Bz.box()); - bzfab = &filtered_Bz; -#endif -} - void PhysicalParticleContainer::Evolve (int lev, const MultiFab& Ex, const MultiFab& Ey, const MultiFab& Ez, @@ -1180,54 +1113,12 @@ PhysicalParticleContainer::Evolve (int lev, if (WarpX::use_fdtd_nci_corr) { -#if (AMREX_SPACEDIM == 2) - const Box& tbox = amrex::grow(cbox,{static_cast(WarpX::nox), - static_cast(WarpX::noz)}); -#else - const Box& tbox = amrex::grow(cbox,{static_cast(WarpX::nox), - static_cast(WarpX::noy), - static_cast(WarpX::noz)}); -#endif - - // Filter Ex (both 2D and 3D) - filtered_Ex.resize(amrex::convert(tbox,WarpX::Ex_nodal_flag)); - // Safeguard for GPU - exeli = filtered_Ex.elixir(); - // Apply filter on Ex, result stored in filtered_Ex - nci_godfrey_filter_exeybz[lev-1]->ApplyStencil(filtered_Ex, (*cEx)[pti], filtered_Ex.box()); - // Update exfab reference - cexfab = &filtered_Ex; - - // Filter Ez - filtered_Ez.resize(amrex::convert(tbox,WarpX::Ez_nodal_flag)); - ezeli = filtered_Ez.elixir(); - nci_godfrey_filter_bxbyez[lev-1]->ApplyStencil(filtered_Ez, (*cEz)[pti], filtered_Ez.box()); - cezfab = &filtered_Ez; - - // Filter By - filtered_By.resize(amrex::convert(tbox,WarpX::By_nodal_flag)); - byeli = filtered_By.elixir(); - nci_godfrey_filter_bxbyez[lev-1]->ApplyStencil(filtered_By, (*cBy)[pti], filtered_By.box()); - cbyfab = &filtered_By; -#if (AMREX_SPACEDIM == 3) - // Filter Ey - filtered_Ey.resize(amrex::convert(tbox,WarpX::Ey_nodal_flag)); - eyeli = filtered_Ey.elixir(); - nci_godfrey_filter_exeybz[lev-1]->ApplyStencil(filtered_Ey, (*cEy)[pti], filtered_Ey.box()); - ceyfab = &filtered_Ey; - - // Filter Bx - filtered_Bx.resize(amrex::convert(tbox,WarpX::Bx_nodal_flag)); - bxeli = filtered_Bx.elixir(); - nci_godfrey_filter_bxbyez[lev-1]->ApplyStencil(filtered_Bx, (*cBx)[pti], filtered_Bx.box()); - cbxfab = &filtered_Bx; - - // Filter Bz - filtered_Bz.resize(amrex::convert(tbox,WarpX::Bz_nodal_flag)); - bzeli = filtered_Bz.elixir(); - nci_godfrey_filter_exeybz[lev-1]->ApplyStencil(filtered_Bz, (*cBz)[pti], filtered_Bz.box()); - cbzfab = &filtered_Bz; -#endif + applyNCIFilter(lev-1, cbox, exeli, eyeli, ezeli, bxeli, byeli, bzeli, + filtered_Ex, filtered_Ey, filtered_Ez, + filtered_Bx, filtered_By, filtered_Bz, + (*cEx)[pti], (*cEy)[pti], (*cEz)[pti], + (*cBx)[pti], (*cBy)[pti], (*cBz)[pti], + cexfab, ceyfab, cezfab, cbxfab, cbyfab, cbzfab); } // Field gather for particles in gather buffers @@ -1320,6 +1211,74 @@ PhysicalParticleContainer::Evolve (int lev, } } +void +PhysicalParticleContainer::applyNCIFilter ( + int lev, const Box& box, + Elixir& exeli, Elixir& eyeli, Elixir& ezeli, + Elixir& bxeli, Elixir& byeli, Elixir& bzeli, + FArrayBox& filtered_Ex, FArrayBox& filtered_Ey, FArrayBox& filtered_Ez, + FArrayBox& filtered_Bx, FArrayBox& filtered_By, FArrayBox& filtered_Bz, + const FArrayBox& Ex, const FArrayBox& Ey, const FArrayBox& Ez, + const FArrayBox& Bx, const FArrayBox& By, const FArrayBox& Bz, + FArrayBox const * & ex_ptr, FArrayBox const * & ey_ptr, + FArrayBox const * & ez_ptr, FArrayBox const * & bx_ptr, + FArrayBox const * & by_ptr, FArrayBox const * & bz_ptr) +{ + + // Get instances of NCI Godfrey filters + const auto& nci_godfrey_filter_exeybz = WarpX::GetInstance().nci_godfrey_filter_exeybz; + const auto& nci_godfrey_filter_bxbyez = WarpX::GetInstance().nci_godfrey_filter_bxbyez; + +#if (AMREX_SPACEDIM == 2) + const Box& tbox = amrex::grow(box,{static_cast(WarpX::nox), + static_cast(WarpX::noz)}); +#else + const Box& tbox = amrex::grow(box,{static_cast(WarpX::nox), + static_cast(WarpX::noy), + static_cast(WarpX::noz)}); +#endif + + // Filter Ex (Both 2D and 3D) + filtered_Ex.resize(amrex::convert(tbox,WarpX::Ex_nodal_flag)); + // Safeguard for GPU + exeli = filtered_Ex.elixir(); + // Apply filter on Ex, result stored in filtered_Ex + nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Ex, Ex, filtered_Ex.box()); + // Update ex_ptr reference + ex_ptr = &filtered_Ex; + + // Filter Ez + filtered_Ez.resize(amrex::convert(tbox,WarpX::Ez_nodal_flag)); + ezeli = filtered_Ez.elixir(); + nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_Ez, Ez, filtered_Ez.box()); + ez_ptr = &filtered_Ez; + + // Filter By + filtered_By.resize(amrex::convert(tbox,WarpX::By_nodal_flag)); + byeli = filtered_By.elixir(); + nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_By, By, filtered_By.box()); + by_ptr = &filtered_By; +#if (AMREX_SPACEDIM == 3) + // Filter Ey + filtered_Ey.resize(amrex::convert(tbox,WarpX::Ey_nodal_flag)); + eyeli = filtered_Ey.elixir(); + nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Ey, Ey, filtered_Ey.box()); + ey_ptr = &filtered_Ey; + + // Filter Bx + filtered_Bx.resize(amrex::convert(tbox,WarpX::Bx_nodal_flag)); + bxeli = filtered_Bx.elixir(); + nci_godfrey_filter_bxbyez[lev]->ApplyStencil(filtered_Bx, Bx, filtered_Bx.box()); + bx_ptr = &filtered_Bx; + + // Filter Bz + filtered_Bz.resize(amrex::convert(tbox,WarpX::Bz_nodal_flag)); + bzeli = filtered_Bz.elixir(); + nci_godfrey_filter_exeybz[lev]->ApplyStencil(filtered_Bz, Bz, filtered_Bz.box()); + bz_ptr = &filtered_Bz; +#endif +} + // Loop over all particles in the particle container and // split particles tagged with p.id()=DoSplitParticleID void -- cgit v1.2.3 From 153eabf2e6742a63013f4bcbc56f951998104ae3 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Thu, 26 Sep 2019 15:00:07 -0700 Subject: add comments --- Source/Particles/PhysicalParticleContainer.H | 12 ++++++++++++ Source/Particles/PhysicalParticleContainer.cpp | 9 ++++++++- 2 files changed, 20 insertions(+), 1 deletion(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index 392f9df1c..4075bfa4d 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -145,6 +145,18 @@ public: virtual void ConvertUnits (ConvertDirection convert_dir) override; +/* \brief Apply NCI Godfrey filter to all components of E and B before gather + * \param lev MR level + * \param box box onto which the filter is applied + * \param exeli safeguard to avoid destructing arrays between ParIter iterations on GPU + * \param filtered_Ex Array containing filtered value + * \param Ex Field array before filtering (not modified) + * \param ex_ptr pointer to the array used for field gather. + * + * The NCI Godfrey filter is applied on Ex, the result is stored in filtered_Ex + * and the pointer is modified (before this function is called, it points to Ex + * and after this function is called, it points to Ex_filtered) + */ void applyNCIFilter ( int lev, const amrex::Box& box, amrex::Elixir& exeli, amrex::Elixir& eyeli, amrex::Elixir& ezeli, diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index b0da0e887..b2bccf2ed 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1026,12 +1026,15 @@ PhysicalParticleContainer::Evolve (int lev, if (WarpX::use_fdtd_nci_corr) { + // Filter arrays Ex[pti], store the result in + // filtered_Ex and update pointer exfab so that it + // points to filtered_Ex (and do the same for all + // components of E and B). applyNCIFilter(lev, pti.tilebox(), exeli, eyeli, ezeli, bxeli, byeli, bzeli, filtered_Ex, filtered_Ey, filtered_Ez, filtered_Bx, filtered_By, filtered_Bz, Ex[pti], Ey[pti], Ez[pti], Bx[pti], By[pti], Bz[pti], exfab, eyfab, ezfab, bxfab, byfab, bzfab); - } Exp.assign(np,0.0); @@ -1113,6 +1116,10 @@ PhysicalParticleContainer::Evolve (int lev, if (WarpX::use_fdtd_nci_corr) { + // Filter arrays (*cEx)[pti], store the result in + // filtered_Ex and update pointer cexfab so that it + // points to filtered_Ex (and do the same for all + // components of E and B) applyNCIFilter(lev-1, cbox, exeli, eyeli, ezeli, bxeli, byeli, bzeli, filtered_Ex, filtered_Ey, filtered_Ez, filtered_Bx, filtered_By, filtered_Bz, -- cgit v1.2.3 From a8f9fcafac08b9e51b894dc56b1e8a999b3a923e Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Thu, 26 Sep 2019 16:30:59 -0700 Subject: eol whitespace --- Source/Particles/PhysicalParticleContainer.H | 22 +++++++++++----------- Source/Particles/PhysicalParticleContainer.cpp | 16 ++++++++-------- 2 files changed, 19 insertions(+), 19 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index 4075bfa4d..35398ae4c 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -151,24 +151,24 @@ public: * \param exeli safeguard to avoid destructing arrays between ParIter iterations on GPU * \param filtered_Ex Array containing filtered value * \param Ex Field array before filtering (not modified) - * \param ex_ptr pointer to the array used for field gather. - * + * \param ex_ptr pointer to the array used for field gather. + * * The NCI Godfrey filter is applied on Ex, the result is stored in filtered_Ex * and the pointer is modified (before this function is called, it points to Ex * and after this function is called, it points to Ex_filtered) */ void applyNCIFilter ( - int lev, const amrex::Box& box, + int lev, const amrex::Box& box, amrex::Elixir& exeli, amrex::Elixir& eyeli, amrex::Elixir& ezeli, amrex::Elixir& bxeli, amrex::Elixir& byeli, amrex::Elixir& bzeli, - amrex::FArrayBox& filtered_Ex, amrex::FArrayBox& filtered_Ey, - amrex::FArrayBox& filtered_Ez, amrex::FArrayBox& filtered_Bx, - amrex::FArrayBox& filtered_By, amrex::FArrayBox& filtered_Bz, - const amrex::FArrayBox& Ex, const amrex::FArrayBox& Ey, - const amrex::FArrayBox& Ez, const amrex::FArrayBox& Bx, - const amrex::FArrayBox& By, const amrex::FArrayBox& Bz, - amrex::FArrayBox const * & exfab, amrex::FArrayBox const * & eyfab, - amrex::FArrayBox const * & ezfab, amrex::FArrayBox const * & bxfab, + amrex::FArrayBox& filtered_Ex, amrex::FArrayBox& filtered_Ey, + amrex::FArrayBox& filtered_Ez, amrex::FArrayBox& filtered_Bx, + amrex::FArrayBox& filtered_By, amrex::FArrayBox& filtered_Bz, + const amrex::FArrayBox& Ex, const amrex::FArrayBox& Ey, + const amrex::FArrayBox& Ez, const amrex::FArrayBox& Bx, + const amrex::FArrayBox& By, const amrex::FArrayBox& Bz, + amrex::FArrayBox const * & exfab, amrex::FArrayBox const * & eyfab, + amrex::FArrayBox const * & ezfab, amrex::FArrayBox const * & bxfab, amrex::FArrayBox const * & byfab, amrex::FArrayBox const * & bzfab); protected: diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index b2bccf2ed..647386a51 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -958,7 +958,7 @@ PhysicalParticleContainer::Evolve (int lev, // Get instances of NCI Godfrey filters const auto& nci_godfrey_filter_exeybz = WarpX::GetInstance().nci_godfrey_filter_exeybz; const auto& nci_godfrey_filter_bxbyez = WarpX::GetInstance().nci_godfrey_filter_bxbyez; - + BL_ASSERT(OnSameGrids(lev,jx)); MultiFab* cost = WarpX::getCosts(lev); @@ -1220,15 +1220,15 @@ PhysicalParticleContainer::Evolve (int lev, void PhysicalParticleContainer::applyNCIFilter ( - int lev, const Box& box, + int lev, const Box& box, Elixir& exeli, Elixir& eyeli, Elixir& ezeli, Elixir& bxeli, Elixir& byeli, Elixir& bzeli, - FArrayBox& filtered_Ex, FArrayBox& filtered_Ey, FArrayBox& filtered_Ez, - FArrayBox& filtered_Bx, FArrayBox& filtered_By, FArrayBox& filtered_Bz, - const FArrayBox& Ex, const FArrayBox& Ey, const FArrayBox& Ez, - const FArrayBox& Bx, const FArrayBox& By, const FArrayBox& Bz, - FArrayBox const * & ex_ptr, FArrayBox const * & ey_ptr, - FArrayBox const * & ez_ptr, FArrayBox const * & bx_ptr, + FArrayBox& filtered_Ex, FArrayBox& filtered_Ey, FArrayBox& filtered_Ez, + FArrayBox& filtered_Bx, FArrayBox& filtered_By, FArrayBox& filtered_Bz, + const FArrayBox& Ex, const FArrayBox& Ey, const FArrayBox& Ez, + const FArrayBox& Bx, const FArrayBox& By, const FArrayBox& Bz, + FArrayBox const * & ex_ptr, FArrayBox const * & ey_ptr, + FArrayBox const * & ez_ptr, FArrayBox const * & bx_ptr, FArrayBox const * & by_ptr, FArrayBox const * & bz_ptr) { -- cgit v1.2.3