From 4b4f13857bd4fd623096a367b784e30fe15a8810 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Mon, 27 Jan 2020 13:42:39 -0800 Subject: remove the copies between soa and aos for the particle positions --- Source/Particles/PhotonParticleContainer.cpp | 18 +++++++----------- 1 file changed, 7 insertions(+), 11 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index c03ed00c2..bf4ed7055 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -50,19 +50,16 @@ void PhotonParticleContainer::InitData() } void -PhotonParticleContainer::PushPX(WarpXParIter& pti, - Gpu::ManagedDeviceVector& xp, - Gpu::ManagedDeviceVector& yp, - Gpu::ManagedDeviceVector& zp, - Real dt, DtType a_dt_type) +PhotonParticleContainer::PushPX(WarpXParIter& pti, 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 - ParticleReal* const AMREX_RESTRICT x = xp.dataPtr(); - ParticleReal* const AMREX_RESTRICT y = yp.dataPtr(); - ParticleReal* const AMREX_RESTRICT z = zp.dataPtr(); + auto& aos = pti.GetArrayOfStructs(); + ParticleType* AMREX_RESTRICT const pstruct = aos().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(); @@ -75,15 +72,14 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, if (WarpX::do_back_transformed_diagnostics && do_back_transformed_diagnostics) { - copy_attribs(pti, x, y, z); + copy_attribs(pti, pstruct); } amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { - UpdatePositionPhoton( x[i], y[i], z[i], - ux[i], uy[i], uz[i], dt ); + UpdatePositionPhoton( pstruct[i], ux[i], uy[i], uz[i], dt ); } ); } -- cgit v1.2.3 From cb93a9c54bd616ca01d189a0eaa910ac82a5c899 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Mon, 27 Jan 2020 13:47:28 -0800 Subject: fix EOL whitespace --- Source/Laser/LaserParticleContainer.cpp | 4 ++-- Source/Parser/WarpXParserWrapper.H | 2 +- Source/Particles/PhotonParticleContainer.cpp | 2 +- Source/Particles/PhysicalParticleContainer.cpp | 2 +- Source/Particles/RigidInjectedParticleContainer.cpp | 12 ++++++------ Source/Particles/WarpXParticleContainer.cpp | 4 ++-- 6 files changed, 13 insertions(+), 13 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Laser/LaserParticleContainer.cpp b/Source/Laser/LaserParticleContainer.cpp index 9f4374517..2520b9ec1 100644 --- a/Source/Laser/LaserParticleContainer.cpp +++ b/Source/Laser/LaserParticleContainer.cpp @@ -568,7 +568,7 @@ LaserParticleContainer::calculate_laser_plane_coordinates (const WarpXParIter& p Real * AMREX_RESTRICT const pplane_Yp) { const auto& aos = pti.GetArrayOfStructs(); - const ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); + const ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); Real tmp_u_X_0 = u_X[0]; Real tmp_u_X_2 = u_X[2]; @@ -624,7 +624,7 @@ LaserParticleContainer::update_laser_particle(WarpXParIter& pti, const Real dt) { auto& aos = pti.GetArrayOfStructs(); - ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); + ParticleType* AMREX_RESTRICT const pstruct = aos().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/Parser/WarpXParserWrapper.H b/Source/Parser/WarpXParserWrapper.H index f7bed4e45..129422e95 100644 --- a/Source/Parser/WarpXParserWrapper.H +++ b/Source/Parser/WarpXParserWrapper.H @@ -35,7 +35,7 @@ struct ParserWrapper { return m_parser(p.pos(0),p.pos(1),p.pos(2)); } - + GpuParser m_parser; }; diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index bf4ed7055..a4d738f93 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -59,7 +59,7 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, Real dt, DtType a_dt_type) // Extract pointers to the different particle quantities auto& aos = pti.GetArrayOfStructs(); ParticleType* AMREX_RESTRICT const pstruct = aos().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(); diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 36e823d5f..1430df6a0 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1468,7 +1468,7 @@ PhysicalParticleContainer::SplitParticles(int lev) { const auto& aos = pti.GetArrayOfStructs(); const ParticleType* AMREX_RESTRICT pstruct = aos().dataPtr(); - + 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._rt, diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp index ad2c89428..4d4aedb0e 100644 --- a/Source/Particles/RigidInjectedParticleContainer.cpp +++ b/Source/Particles/RigidInjectedParticleContainer.cpp @@ -87,7 +87,7 @@ RigidInjectedParticleContainer::RemapParticles() auto& aos = pti.GetArrayOfStructs(); ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); - + // Loop over particles const long np = pti.numParticles(); for (int i=0 ; i < np ; i++) { @@ -211,7 +211,7 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ auto& aos = pti.GetArrayOfStructs(); ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); - + ParticleReal* const AMREX_RESTRICT ux = uxp.dataPtr(); ParticleReal* const AMREX_RESTRICT uy = uyp.dataPtr(); ParticleReal* const AMREX_RESTRICT uz = uzp.dataPtr(); @@ -221,7 +221,7 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ 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. @@ -242,7 +242,7 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ amrex::Real* const AMREX_RESTRICT uxp_save_ptr = uxp_save.dataPtr(); amrex::Real* const AMREX_RESTRICT uyp_save_ptr = uyp_save.dataPtr(); amrex::Real* const AMREX_RESTRICT uzp_save_ptr = uzp_save.dataPtr(); - + amrex::ParallelFor( np, [=] AMREX_GPU_DEVICE (long i) { xp_save_ptr[i] = pstruct[i].pos(0); @@ -251,7 +251,7 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ uxp_save_ptr[i] = ux[i]; uyp_save_ptr[i] = uy[i]; uzp_save_ptr[i] = uz[i]; - }); + }); // Scale the fields of particles about to cross the injection plane. // This only approximates what should be happening. The particles @@ -486,7 +486,7 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, } } ); - + } } } diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index eab410a18..170770c58 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -288,7 +288,7 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, const auto& aos = pti.GetArrayOfStructs(); const ParticleType* AMREX_RESTRICT pstruct = aos().dataPtr() + offset; - + // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow const Dim3 lo = lbound(tilebox); @@ -751,7 +751,7 @@ WarpXParticleContainer::PushX (int lev, amrex::Real dt) ParticleType& p = pstructs[i]; // Particle object that gets updated #ifndef WARPX_DIM_RZ UpdatePosition( p, ux[i], uy[i], uz[i], dt); -#else +#else // For WARPX_DIM_RZ, the particles are still pushed in 3D Cartesian ParticleReal x, y, z; // Temporary variables GetCartesianPositionFromCylindrical( x, y, z, p, theta[i] ); -- cgit v1.2.3 From 80074a19c67a17a9b8169dc64004322bbd0c31b5 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Tue, 28 Jan 2020 13:24:49 -0800 Subject: switch deposition, gather, and pushers to use the get / set position functor --- Source/Particles/Deposition/ChargeDeposition.H | 17 +++--- Source/Particles/Deposition/CurrentDeposition.H | 49 ++++++++------- Source/Particles/Gather/FieldGather.H | 22 ++++--- Source/Particles/PhotonParticleContainer.cpp | 13 ++-- Source/Particles/PhysicalParticleContainer.H | 2 +- Source/Particles/PhysicalParticleContainer.cpp | 69 ++++++++++++++-------- Source/Particles/Pusher/GetAndSetPosition.H | 12 ++-- Source/Particles/Pusher/UpdatePosition.H | 8 +-- Source/Particles/Pusher/UpdatePositionPhoton.H | 9 ++- .../Particles/RigidInjectedParticleContainer.cpp | 7 +-- Source/Particles/WarpXParticleContainer.cpp | 45 +++++++------- 11 files changed, 144 insertions(+), 109 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Particles/Deposition/ChargeDeposition.H b/Source/Particles/Deposition/ChargeDeposition.H index 53861984e..9e3df600b 100755 --- a/Source/Particles/Deposition/ChargeDeposition.H +++ b/Source/Particles/Deposition/ChargeDeposition.H @@ -1,10 +1,11 @@ #ifndef CHARGEDEPOSITION_H_ #define CHARGEDEPOSITION_H_ +#include "GetAndSetPosition.H" #include "ShapeFactors.H" /* \brief Charge Deposition for thread thread_num - * /param pstruct : Pointer to array of particle structs + * /param get_position : A functor for returning the particle position. * \param wp : Pointer to array of particle weights. * \param ion_lev : Pointer to array of particle ionization level. This is required to have the charge of each macroparticle @@ -18,7 +19,7 @@ * /param q : species charge. */ template -void doChargeDepositionShapeN(const WarpXParticleContainer::ParticleType* const pstruct, +void doChargeDepositionShapeN(const GetPosition& get_position, const amrex::ParticleReal * const wp, const int * const ion_lev, const amrex::Array4& rho_arr, @@ -54,15 +55,17 @@ void doChargeDepositionShapeN(const WarpXParticleContainer::ParticleType* const wq *= ion_lev[ip]; } + amrex::Real xp, yp, zp; + get_position(ip, xp, yp, zp); + // --- Compute shape factors // x direction // Get particle position in grid coordinates #if (defined WARPX_DIM_RZ) - const amrex::Real r = std::sqrt(pstruct[ip].pos(0)*pstruct[ip].pos(0) - + pstruct[ip].pos(1)*pstruct[ip].pos(1)); + const amrex::Real r = std::sqrt(xp*xp + yp*yp); const amrex::Real x = (r - xmin)*dxi; #else - const amrex::Real x = (pstruct[ip].pos(0) - xmin)*dxi; + const amrex::Real x = (xp - xmin)*dxi; #endif // Compute shape factors for node-centered quantities amrex::Real sx[depos_order + 1]; @@ -71,12 +74,12 @@ void doChargeDepositionShapeN(const WarpXParticleContainer::ParticleType* const #if (defined WARPX_DIM_3D) // y direction - const amrex::Real y = (pstruct[ip].pos(1) - ymin)*dyi; + const amrex::Real y = (yp - ymin)*dyi; amrex::Real sy[depos_order + 1]; const int j = compute_shape_factor(sy, y); #endif // z direction - const amrex::Real z = (pstruct[ip].pos(2) - zmin)*dzi; + const amrex::Real z = (zp - zmin)*dzi; amrex::Real sz[depos_order + 1]; const int k = compute_shape_factor(sz, z); diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H index b5e5c13a6..8e7f1a552 100644 --- a/Source/Particles/Deposition/CurrentDeposition.H +++ b/Source/Particles/Deposition/CurrentDeposition.H @@ -1,6 +1,7 @@ #ifndef CURRENTDEPOSITION_H_ #define CURRENTDEPOSITION_H_ +#include "GetAndSetPosition.H" #include "ShapeFactors.H" #include @@ -9,7 +10,7 @@ /** * \brief Current Deposition for thread thread_num - * /param pstruct : Pointer to array of particle structs. + * /param get_position : A functor for returning the particle position. * \param wp : Pointer to array of particle weights. * \param uxp uyp uzp : Pointer to arrays of particle momentum. * \param ion_lev : Pointer to array of particle ionization level. This is @@ -27,7 +28,7 @@ * /param q : species charge. */ template -void doDepositionShapeN(const WarpXParticleContainer::ParticleType * const pstruct, +void doDepositionShapeN(const GetPosition& get_position, const amrex::ParticleReal * const wp, const amrex::ParticleReal * const uxp, const amrex::ParticleReal * const uyp, @@ -86,6 +87,10 @@ void doDepositionShapeN(const WarpXParticleContainer::ParticleType * const pstru if (do_ionization){ wq *= ion_lev[ip]; } + + amrex::Real xp, yp, zp; + get_position(ip, xp, yp, zp); + const amrex::Real vx = uxp[ip]*gaminv; const amrex::Real vy = uyp[ip]*gaminv; const amrex::Real vz = uzp[ip]*gaminv; @@ -93,8 +98,8 @@ void doDepositionShapeN(const WarpXParticleContainer::ParticleType * const pstru #if (defined WARPX_DIM_RZ) // In RZ, wqx is actually wqr, and wqy is wqtheta // Convert to cylinderical at the mid point - const amrex::Real xpmid = pstruct[ip].pos(0) - 0.5*dt*vx; - const amrex::Real ypmid = pstruct[ip].pos(1) - 0.5*dt*vy; + const amrex::Real xpmid = xp - 0.5*dt*vx; + const amrex::Real ypmid = yp - 0.5*dt*vy; const amrex::Real rpmid = std::sqrt(xpmid*xpmid + ypmid*ypmid); amrex::Real costheta; amrex::Real sintheta; @@ -119,7 +124,7 @@ void doDepositionShapeN(const WarpXParticleContainer::ParticleType * const pstru #if (defined WARPX_DIM_RZ) const amrex::Real xmid = (rpmid - xmin)*dxi; #else - const amrex::Real xmid = (pstruct[ip].pos(0) - xmin)*dxi - dts2dx*vx; + const amrex::Real xmid = (xp - xmin)*dxi - dts2dx*vx; #endif // j_j[xyz] leftmost grid point in x that the particle touches for the centering of each current // sx_j[xyz] shape factor along x for the centering of each current @@ -144,7 +149,7 @@ void doDepositionShapeN(const WarpXParticleContainer::ParticleType * const pstru #if (defined WARPX_DIM_3D) // y direction - const amrex::Real ymid = (pstruct[ip].pos(1) - ymin)*dyi - dts2dy*vy; + const amrex::Real ymid = (yp - ymin)*dyi - dts2dy*vy; amrex::Real sy_node[depos_order + 1]; amrex::Real sy_cell[depos_order + 1]; int k_node; @@ -164,7 +169,7 @@ void doDepositionShapeN(const WarpXParticleContainer::ParticleType * const pstru #endif // z direction - const amrex::Real zmid = (pstruct[ip].pos(2) - zmin)*dzi - dts2dz*vz; + const amrex::Real zmid = (zp - zmin)*dzi - dts2dz*vz; amrex::Real sz_node[depos_order + 1]; amrex::Real sz_cell[depos_order + 1]; int l_node; @@ -221,7 +226,7 @@ void doDepositionShapeN(const WarpXParticleContainer::ParticleType * const pstru /** * \brief Esirkepov Current Deposition for thread thread_num * - * \param pstruct : Pointer to array of particle structs. + * /param get_position : A functor for returning the particle position. * \param wp : Pointer to array of particle weights. * \param uxp uyp uzp : Pointer to arrays of particle momentum. * \param ion_lev : Pointer to array of particle ionization level. This is @@ -240,7 +245,7 @@ void doDepositionShapeN(const WarpXParticleContainer::ParticleType * const pstru * \param n_rz_azimuthal_modes: Number of azimuthal modes when using RZ geometry */ template -void doEsirkepovDepositionShapeN (const WarpXParticleContainer::ParticleType * pstruct, +void doEsirkepovDepositionShapeN (const GetPosition& get_position, const amrex::ParticleReal * const wp, const amrex::ParticleReal * const uxp, const amrex::ParticleReal * const uyp, @@ -305,6 +310,10 @@ void doEsirkepovDepositionShapeN (const WarpXParticleContainer::ParticleType * p if (do_ionization){ wq *= ion_lev[ip]; } + + Real xp, yp, zp; + get_position(ip, xp, yp, zp); + Real const wqx = wq*invdtdx; #if (defined WARPX_DIM_3D) Real const wqy = wq*invdtdy; @@ -313,18 +322,18 @@ void doEsirkepovDepositionShapeN (const WarpXParticleContainer::ParticleType * p // computes current and old position in grid units #if (defined WARPX_DIM_RZ) - Real const xp_mid = pstruct[ip].pos(0) - 0.5_rt * dt*uxp[ip]*gaminv; - Real const yp_mid = pstruct[ip].pos(1) - 0.5_rt * dt*uyp[ip]*gaminv; - Real const xp_old = pstruct[ip].pos(0) - dt*uxp[ip]*gaminv; - Real const yp_old = pstruct[ip].pos(1) - dt*uyp[ip]*gaminv; - Real const rp_new = std::sqrt(pstruct[ip].pos(0)*pstruct[ip].pos(0) - + pstruct[ip].pos(1)*pstruct[ip].pos(1)); + Real const xp_mid = xp - 0.5_rt * dt*uxp[ip]*gaminv; + Real const yp_mid = yp - 0.5_rt * dt*uyp[ip]*gaminv; + Real const xp_old = xp - dt*uxp[ip]*gaminv; + Real const yp_old = yp - dt*uyp[ip]*gaminv; + Real const rp_new = std::sqrt(xp*xp + + yp*yp); Real const rp_mid = std::sqrt(xp_mid*xp_mid + yp_mid*yp_mid); Real const rp_old = std::sqrt(xp_old*xp_old + yp_old*yp_old); Real costheta_new, sintheta_new; if (rp_new > 0._rt) { - costheta_new = pstruct[ip].pos(0)/rp_new; - sintheta_new = pstruct[ip].pos(1)/rp_new; + costheta_new = xp/rp_new; + sintheta_new = yp/rp_new; } else { costheta_new = 1.; sintheta_new = 0.; @@ -351,14 +360,14 @@ void doEsirkepovDepositionShapeN (const WarpXParticleContainer::ParticleType * p Real const x_new = (rp_new - xmin)*dxi; Real const x_old = (rp_old - xmin)*dxi; #else - Real const x_new = (pstruct[ip].pos(0) - xmin)*dxi; + Real const x_new = (xp - xmin)*dxi; Real const x_old = x_new - dtsdx0*uxp[ip]*gaminv; #endif #if (defined WARPX_DIM_3D) - Real const y_new = (pstruct[ip].pos(1) - ymin)*dyi; + Real const y_new = (yp - ymin)*dyi; Real const y_old = y_new - dtsdy0*uyp[ip]*gaminv; #endif - Real const z_new = (pstruct[ip].pos(2) - zmin)*dzi; + Real const z_new = (zp - zmin)*dzi; Real const z_old = z_new - dtsdz0*uzp[ip]*gaminv; #if (defined WARPX_DIM_RZ) diff --git a/Source/Particles/Gather/FieldGather.H b/Source/Particles/Gather/FieldGather.H index 0af1e362a..369b648c1 100644 --- a/Source/Particles/Gather/FieldGather.H +++ b/Source/Particles/Gather/FieldGather.H @@ -1,12 +1,13 @@ #ifndef FIELDGATHER_H_ #define FIELDGATHER_H_ +#include "GetAndSetPosition.H" #include "ShapeFactors.H" #include /** * \brief Field gather for particles handled by thread thread_num - * \param pstruct : Pointer to array of particle structs. + * /param get_position : A functor for returning the particle position. * \param Exp, Eyp, Ezp: Pointer to array of electric field on particles. * \param Bxp, Byp, Bzp: Pointer to array of magnetic field on particles. * \param ex_arr ey_arr: Array4 of current density, either full array or tile. @@ -20,7 +21,7 @@ * \param n_rz_azimuthal_modes: Number of azimuthal modes when using RZ geometry */ template -void doGatherShapeN(const WarpXParticleContainer::ParticleType * const pstruct, +void doGatherShapeN(const GetPosition& get_position, amrex::ParticleReal * const Exp, amrex::ParticleReal * const Eyp, amrex::ParticleReal * const Ezp, amrex::ParticleReal * const Bxp, amrex::ParticleReal * const Byp, amrex::ParticleReal * const Bzp, @@ -71,15 +72,18 @@ void doGatherShapeN(const WarpXParticleContainer::ParticleType * const pstruct, amrex::ParallelFor( np_to_gather, [=] AMREX_GPU_DEVICE (long ip) { + + amrex::Real xp, yp, zp; + get_position(ip, xp, yp, zp); + // --- Compute shape factors // x direction // Get particle position #ifdef WARPX_DIM_RZ - const amrex::Real rp = std::sqrt(pstruct[ip].pos(0)*pstruct[ip].pos(0) - + pstruct[ip].pos(1)*pstruct[ip].pos(1)); + const amrex::Real rp = std::sqrt(xp*xp + yp*yp); const amrex::Real x = (rp - xmin)*dxi; #else - const amrex::Real x = (pstruct[ip].pos(0)-xmin)*dxi; + const amrex::Real x = (xp-xmin)*dxi; #endif // j_[eb][xyz] leftmost grid point in x that the particle touches for the centering of each current @@ -121,7 +125,7 @@ void doGatherShapeN(const WarpXParticleContainer::ParticleType * const pstruct, #if (AMREX_SPACEDIM == 3) // y direction - const amrex::Real y = (pstruct[ip].pos(1)-ymin)*dyi; + const amrex::Real y = (yp-ymin)*dyi; amrex::Real sy_node[depos_order + 1]; amrex::Real sy_cell[depos_order + 1]; amrex::Real sy_node_v[depos_order + 1 - lower_in_v]; @@ -157,7 +161,7 @@ void doGatherShapeN(const WarpXParticleContainer::ParticleType * const pstruct, #endif // z direction - const amrex::Real z = (pstruct[ip].pos(2)-zmin)*dzi; + const amrex::Real z = (zp-zmin)*dzi; amrex::Real sz_node[depos_order + 1]; amrex::Real sz_cell[depos_order + 1]; amrex::Real sz_node_v[depos_order + 1 - lower_in_v]; @@ -237,8 +241,8 @@ void doGatherShapeN(const WarpXParticleContainer::ParticleType * const pstruct, amrex::Real costheta; amrex::Real sintheta; if (rp > 0.) { - costheta = pstruct[ip].pos(0)/rp; - sintheta = pstruct[ip].pos(1)/rp; + costheta = xp/rp; + sintheta = yp/rp; } else { costheta = 1.; sintheta = 0.; diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index a4d738f93..86cc206a7 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -14,7 +14,7 @@ // Import low-level single-particle kernels #include - +#include using namespace amrex; @@ -72,14 +72,19 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, Real dt, DtType a_dt_type) if (WarpX::do_back_transformed_diagnostics && do_back_transformed_diagnostics) { - copy_attribs(pti, pstruct); + copy_attribs(pti); } + const auto get_position = GetPosition(pti); + auto set_position = SetPosition(pti); + amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { - - UpdatePositionPhoton( pstruct[i], ux[i], uy[i], uz[i], dt ); + Real x, y, z; + get_position(i, x, y, z); + UpdatePositionPhoton( x, y, z, ux[i], uy[i], uz[i], dt ); + set_position(i, x, y, z); } ); } diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index 4cb71a5d2..392d13244 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -172,7 +172,7 @@ public: RealVector& uzp, RealVector& wp ); - void copy_attribs (WarpXParIter& pti, const ParticleType* pstruct); + void copy_attribs (WarpXParIter& pti); virtual void PostRestart () final {} diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index a660adadd..110d4294c 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1618,8 +1618,8 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) auto& attribs = pti.GetAttribs(); // Extract pointers to the different particle quantities - auto& aos = pti.GetArrayOfStructs(); - ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); + const auto get_position = GetPosition(pti); + auto set_position = SetPosition(pti); ParticleReal* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); ParticleReal* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); @@ -1633,7 +1633,7 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) if (WarpX::do_back_transformed_diagnostics && do_back_transformed_diagnostics && (a_dt_type!=DtType::SecondHalf)) { - copy_attribs(pti, pstruct); + copy_attribs(pti); } int* AMREX_RESTRICT ion_lev = nullptr; @@ -1666,7 +1666,10 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) Ex[i], Ey[i], Ez[i], Bx[i], By[i], Bz[i], q, m, dt); } - UpdatePosition( pstruct[i], ux[i], uy[i], uz[i], dt ); + Real x, y, z; + get_position(i, x, y, z); + UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt ); + set_position(i, x, y, z); } ); }else{ @@ -1676,7 +1679,10 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) UpdateMomentumBorisWithRadiationReaction( ux[i], uy[i], uz[i], Ex[i], Ey[i], Ez[i], Bx[i], By[i], Bz[i], q, m, dt); - UpdatePosition( pstruct[i], ux[i], uy[i], uz[i], dt ); + Real x, y, z; + get_position(i, x, y, z); + UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt ); + set_position(i, x, y, z); } ); } @@ -1690,8 +1696,10 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) UpdateMomentumBorisWithRadiationReaction( ux[i], uy[i], uz[i], Ex[i], Ey[i], Ez[i], Bx[i], By[i], Bz[i], qp, m, dt); - UpdatePosition( pstruct[i], - ux[i], uy[i], uz[i], dt ); + Real x, y, z; + get_position(i, x, y, z); + UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt ); + set_position(i, x, y, z); } ); #endif @@ -1704,8 +1712,10 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) UpdateMomentumBoris( ux[i], uy[i], uz[i], Ex[i], Ey[i], Ez[i], Bx[i], By[i], Bz[i], qp, m, dt); - UpdatePosition( pstruct[i], - ux[i], uy[i], uz[i], dt ); + Real x, y, z; + get_position(i, x, y, z); + UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt ); + set_position(i, x, y, z); } ); } else if (WarpX::particle_pusher_algo == ParticlePusherAlgo::Vay) { @@ -1717,8 +1727,10 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) UpdateMomentumVay( ux[i], uy[i], uz[i], Ex[i], Ey[i], Ez[i], Bx[i], By[i], Bz[i], qp, m, dt); - UpdatePosition( pstruct[i], - ux[i], uy[i], uz[i], dt ); + Real x, y, z; + get_position(i, x, y, z); + UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt ); + set_position(i, x, y, z); } ); } else if (WarpX::particle_pusher_algo == ParticlePusherAlgo::HigueraCary) { @@ -1730,8 +1742,10 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) UpdateMomentumHigueraCary( ux[i], uy[i], uz[i], Ex[i], Ey[i], Ez[i], Bx[i], By[i], Bz[i], qp, m, dt); - UpdatePosition( pstruct[i], - ux[i], uy[i], uz[i], dt ); + Real x, y, z; + get_position(i, x, y, z); + UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt ); + set_position(i, x, y, z); } ); } else { @@ -1901,7 +1915,7 @@ PhysicalParticleContainer::PushP (int lev, Real dt, } } -void PhysicalParticleContainer::copy_attribs (WarpXParIter& pti, const ParticleType* pstruct) +void PhysicalParticleContainer::copy_attribs (WarpXParIter& pti) { auto& attribs = pti.GetAttribs(); ParticleReal* AMREX_RESTRICT uxp = attribs[PIdx::ux].dataPtr(); @@ -1918,11 +1932,15 @@ void PhysicalParticleContainer::copy_attribs (WarpXParIter& pti, const ParticleT ParticleReal* AMREX_RESTRICT uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr(); ParticleReal* AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); + const auto get_position = GetPosition(pti); + ParallelFor( np, [=] AMREX_GPU_DEVICE (long i) { - xpold[i]=pstruct[i].pos(0); - ypold[i]=pstruct[i].pos(1); - zpold[i]=pstruct[i].pos(2); + Real x, y, z; + get_position(i, x, y, z); + xpold[i]=x; + ypold[i]=y; + zpold[i]=z; uxpold[i]=uxp[i]; uypold[i]=uyp[i]; @@ -2193,9 +2211,8 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, // Add guard cells to the box. box.grow(ngE); - const auto& aos = pti.GetArrayOfStructs(); - const ParticleType* AMREX_RESTRICT pstruct = aos().dataPtr() + offset; - + const auto get_position = GetPosition(pti, offset); + // Lower corner of tile box physical domain const std::array& xyzmin = WarpX::LowerCorner(box, gather_lev); @@ -2205,7 +2222,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, // different versions of template function doGatherShapeN if (WarpX::l_lower_order_in_v){ if (WarpX::nox == 1){ - doGatherShapeN<1,1>(pstruct, + doGatherShapeN<1,1>(get_position, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2213,7 +2230,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, np_to_gather, dx, xyzmin, lo, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ - doGatherShapeN<2,1>(pstruct, + doGatherShapeN<2,1>(get_position, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2221,7 +2238,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, np_to_gather, dx, xyzmin, lo, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ - doGatherShapeN<3,1>(pstruct, + doGatherShapeN<3,1>(get_position, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2231,7 +2248,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, } } else { if (WarpX::nox == 1){ - doGatherShapeN<1,0>(pstruct, + doGatherShapeN<1,0>(get_position, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2239,7 +2256,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, np_to_gather, dx, xyzmin, lo, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ - doGatherShapeN<2,0>(pstruct, + doGatherShapeN<2,0>(get_position, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2247,7 +2264,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, np_to_gather, dx, xyzmin, lo, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ - doGatherShapeN<3,0>(pstruct, + doGatherShapeN<3,0>(get_position, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, diff --git a/Source/Particles/Pusher/GetAndSetPosition.H b/Source/Particles/Pusher/GetAndSetPosition.H index e3c017aa6..a429a98fc 100644 --- a/Source/Particles/Pusher/GetAndSetPosition.H +++ b/Source/Particles/Pusher/GetAndSetPosition.H @@ -16,13 +16,13 @@ struct GetPosition #elif (AMREX_SPACEDIM == 2) static constexpr RType m_snan = std::numeric_limits::quiet_NaN(); #endif - GetPosition (const WarpXParIter& a_pti) noexcept + GetPosition (const WarpXParIter& a_pti, int a_offset = 0) noexcept { const auto& aos = a_pti.GetArrayOfStructs(); - m_structs = aos().dataPtr(); + m_structs = aos().dataPtr() + a_offset; #if (defined WARPX_DIM_RZ) const auto& soa = a_pti.GetStructOfArrays(); - m_theta = soa.GetRealData(PIdx::theta).dataPtr(); + m_theta = soa.GetRealData(PIdx::theta).dataPtr() + a_offset; #elif (AMREX_SPACEDIM == 2) static constexpr RType m_snan = std::numeric_limits::quiet_NaN(); #endif @@ -57,13 +57,13 @@ struct SetPosition #if (defined WARPX_DIM_RZ) RType* AMREX_RESTRICT m_theta; #endif - SetPosition (WarpXParIter& a_pti) noexcept + SetPosition (WarpXParIter& a_pti, int a_offset = 0) noexcept { auto& aos = a_pti.GetArrayOfStructs(); - m_structs = aos().dataPtr(); + m_structs = aos().dataPtr() + a_offset; #if (defined WARPX_DIM_RZ) auto& soa = a_pti.GetStructOfArrays(); - m_theta = soa.GetRealData(PIdx::theta).dataPtr(); + m_theta = soa.GetRealData(PIdx::theta).dataPtr() + a_offset; #endif } diff --git a/Source/Particles/Pusher/UpdatePosition.H b/Source/Particles/Pusher/UpdatePosition.H index 564355887..ad9f0f07e 100644 --- a/Source/Particles/Pusher/UpdatePosition.H +++ b/Source/Particles/Pusher/UpdatePosition.H @@ -8,7 +8,7 @@ /** \brief Push the particle's positions over one timestep, * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE -void UpdatePosition(WarpXParticleContainer::ParticleType& p, +void UpdatePosition(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 ) { @@ -18,11 +18,11 @@ void UpdatePosition(WarpXParticleContainer::ParticleType& p, // Compute inverse Lorentz factor const amrex::Real inv_gamma = 1./std::sqrt(1. + (ux*ux + uy*uy + uz*uz)*inv_c2); // Update positions over one time step - p.pos(0) += ux * inv_gamma * dt; + x += ux * inv_gamma * dt; #if (AMREX_SPACEDIM == 3) || (defined WARPX_DIM_RZ) // RZ pushes particles in 3D - p.pos(1) += uy * inv_gamma * dt; + y += uy * inv_gamma * dt; #endif - p.pos(2) += uz * inv_gamma * dt; + z += uz * inv_gamma * dt; } #endif // WARPX_PARTICLES_PUSHER_UPDATEPOSITION_H_ diff --git a/Source/Particles/Pusher/UpdatePositionPhoton.H b/Source/Particles/Pusher/UpdatePositionPhoton.H index fd9597b3f..325e4b748 100644 --- a/Source/Particles/Pusher/UpdatePositionPhoton.H +++ b/Source/Particles/Pusher/UpdatePositionPhoton.H @@ -12,7 +12,7 @@ */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdatePositionPhoton( - WarpXParticleContainer::ParticleType& p, + 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 ) { @@ -20,12 +20,11 @@ void UpdatePositionPhoton( const amrex::Real c_over_umod = PhysConst::c/std::sqrt(ux*ux + uy*uy + uz*uz); // Update positions over one time step - p.pos(0) += ux * c_over_umod * dt; + x += ux * c_over_umod * dt; #if (defined WARPX_DIM_3D) || (defined WARPX_DIM_RZ) // RZ pushes particles in 3D - p.pos(1) += uy * c_over_umod * dt; + y += uy * c_over_umod * dt; #endif - p.pos(2) += uz * c_over_umod * dt; - + z += uz * c_over_umod * dt; } #endif // WARPX_PARTICLES_PUSHER_UPDATEPOSITION_H_ diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp index 4d4aedb0e..fc41fd3d2 100644 --- a/Source/Particles/RigidInjectedParticleContainer.cpp +++ b/Source/Particles/RigidInjectedParticleContainer.cpp @@ -366,11 +366,6 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, #pragma omp parallel #endif { -#ifdef _OPENMP - int thread_num = omp_get_thread_num(); -#else - int thread_num = 0; -#endif for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { const Box& box = pti.validbox(); @@ -401,7 +396,7 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, FieldGather(pti, Exp, Eyp, Ezp, Bxp, Byp, Bzp, &exfab, &eyfab, &ezfab, &bxfab, &byfab, &bzfab, Ex.nGrow(), e_is_nodal, - 0, np, thread_num, lev, lev); + 0, np, lev, lev); // Save the position and momenta, making copies auto uxp_save = uxp; diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 170770c58..779771360 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -286,9 +286,8 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, // CPU, tiling: deposit into local_jx // (same for jx and jz) - const auto& aos = pti.GetArrayOfStructs(); - const ParticleType* AMREX_RESTRICT pstruct = aos().dataPtr() + offset; - + const auto get_position = GetPosition(pti, offset); + // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow const Dim3 lo = lbound(tilebox); @@ -298,19 +297,19 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, if (WarpX::current_deposition_algo == CurrentDepositionAlgo::Esirkepov) { if (WarpX::nox == 1){ doEsirkepovDepositionShapeN<1>( - pstruct, wp.dataPtr() + offset, uxp.dataPtr() + offset, + get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_arr, jy_arr, jz_arr, np_to_depose, dt, dx, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ doEsirkepovDepositionShapeN<2>( - pstruct, wp.dataPtr() + offset, uxp.dataPtr() + offset, + get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_arr, jy_arr, jz_arr, np_to_depose, dt, dx, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ doEsirkepovDepositionShapeN<3>( - pstruct, wp.dataPtr() + offset, uxp.dataPtr() + offset, + get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_arr, jy_arr, jz_arr, np_to_depose, dt, dx, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); @@ -318,19 +317,19 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, } else { if (WarpX::nox == 1){ doDepositionShapeN<1>( - pstruct, wp.dataPtr() + offset, uxp.dataPtr() + offset, + get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_fab, jy_fab, jz_fab, np_to_depose, dt, dx, xyzmin, lo, q); } else if (WarpX::nox == 2){ doDepositionShapeN<2>( - pstruct, wp.dataPtr() + offset, uxp.dataPtr() + offset, + get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_fab, jy_fab, jz_fab, np_to_depose, dt, dx, xyzmin, lo, q); } else if (WarpX::nox == 3){ doDepositionShapeN<3>( - pstruct, wp.dataPtr() + offset, uxp.dataPtr() + offset, + get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_fab, jy_fab, jz_fab, np_to_depose, dt, dx, xyzmin, lo, q); @@ -424,8 +423,7 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector& wp, // GPU, no tiling: deposit directly in rho // CPU, tiling: deposit into local_rho - const auto& aos = pti.GetArrayOfStructs(); - const ParticleType* AMREX_RESTRICT pstruct = aos().dataPtr() + offset; + const auto get_position = GetPosition(pti, offset); // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow @@ -435,13 +433,13 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector& wp, BL_PROFILE_VAR_START(blp_ppc_chd); if (WarpX::nox == 1){ - doChargeDepositionShapeN<1>(pstruct, wp.dataPtr()+offset, ion_lev, + doChargeDepositionShapeN<1>(get_position, wp.dataPtr()+offset, ion_lev, rho_arr, np_to_depose, dx, xyzmin, lo, q); } else if (WarpX::nox == 2){ - doChargeDepositionShapeN<2>(pstruct, wp.dataPtr()+offset, ion_lev, + doChargeDepositionShapeN<2>(get_position, wp.dataPtr()+offset, ion_lev, rho_arr, np_to_depose, dx, xyzmin, lo, q); } else if (WarpX::nox == 3){ - doChargeDepositionShapeN<3>(pstruct, wp.dataPtr()+offset, ion_lev, + doChargeDepositionShapeN<3>(get_position, wp.dataPtr()+offset, ion_lev, rho_arr, np_to_depose, dx, xyzmin, lo, q); } BL_PROFILE_VAR_STOP(blp_ppc_chd); @@ -734,28 +732,33 @@ WarpXParticleContainer::PushX (int lev, amrex::Real dt) // // Particle Push // - // Extract pointers to particle position and momenta, for this particle tile - // - positions are stored as an array of struct, in `ParticleType` - ParticleType * AMREX_RESTRICT pstructs = &(pti.GetArrayOfStructs()[0]); + + const auto get_position = GetPosition(pti); + auto set_position = SetPosition(pti); + // - momenta are stored as a struct of array, in `attribs` auto& attribs = pti.GetAttribs(); 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 + auto& aos = pti.GetArrayOfStructs(); + ParticleType* AMREX_RESTRICT const pstruct = aos().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; #ifndef WARPX_DIM_RZ - UpdatePosition( p, ux[i], uy[i], uz[i], dt); + get_position(i, x, y, z); + UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt); + set_position(i, x, y, z); #else // For WARPX_DIM_RZ, the particles are still pushed in 3D Cartesian - ParticleReal x, y, z; // Temporary variables + ParticleType& p = pstruct[i]; // Particle object that gets updated GetCartesianPositionFromCylindrical( x, y, z, p, theta[i] ); - UpdatePosition( p, ux[i], uy[i], uz[i], dt); + UpdatePosition( x, y, z, ux[i], uy[i], uz[i], dt); SetCylindricalPositionFromCartesian( p, theta[i], x, y, z ); #endif } -- cgit v1.2.3 From 340b8f23629c3246296b8e2287bbda49c1e39c13 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Tue, 28 Jan 2020 13:40:05 -0800 Subject: handle rigid injection and particle splitting --- Source/Particles/PhotonParticleContainer.cpp | 3 -- Source/Particles/PhysicalParticleContainer.cpp | 47 +++++++++++----------- .../Particles/RigidInjectedParticleContainer.cpp | 46 ++++++++++++--------- 3 files changed, 52 insertions(+), 44 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index 86cc206a7..7e7451ab4 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -57,9 +57,6 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, Real dt, DtType a_dt_type) auto& attribs = pti.GetAttribs(); // Extract pointers to the different particle quantities - auto& aos = pti.GetArrayOfStructs(); - ParticleType* AMREX_RESTRICT const pstruct = aos().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(); diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 110d4294c..729ab7044 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1464,8 +1464,7 @@ PhysicalParticleContainer::SplitParticles(int lev) // Loop over particle interator for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { - const auto& aos = pti.GetArrayOfStructs(); - const ParticleType* AMREX_RESTRICT pstruct = aos().dataPtr(); + const auto get_position = GetPosition(pti); const amrex::Vector ppc_nd = plasma_injector->num_particles_per_cell_each_dim; const std::array& dx = WarpX::CellSize(lev); @@ -1490,6 +1489,8 @@ PhysicalParticleContainer::SplitParticles(int lev) auto& uzp = attribs[PIdx::uz]; const long np = pti.numParticles(); for(int i=0; i #include #include +#include using namespace amrex; @@ -85,25 +86,30 @@ RigidInjectedParticleContainer::RemapParticles() auto& uyp = attribs[PIdx::uy]; auto& uzp = attribs[PIdx::uz]; - auto& aos = pti.GetArrayOfStructs(); - ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); - + const auto get_position = GetPosition(pti); + auto set_position = SetPosition(pti); + // Loop over particles const long np = pti.numParticles(); for (int i=0 ; i < np ; i++) { + Real xp, yp, zp; + get_position(i, xp, yp, zp); + const Real gammapr = std::sqrt(1. + (uxp[i]*uxp[i] + uyp[i]*uyp[i] + uzp[i]*uzp[i])/csq); const Real vzpr = uzp[i]/gammapr; // Back out the value of z_lab - const Real z_lab = (pstruct[i].pos(2) + uz_boost*t_lab + WarpX::gamma_boost*t_lab*vzpr)/(WarpX::gamma_boost + uz_boost*vzpr/csq); + const Real z_lab = (zp + uz_boost*t_lab + WarpX::gamma_boost*t_lab*vzpr)/(WarpX::gamma_boost + uz_boost*vzpr/csq); // Time of the particle in the boosted frame given its position in the lab frame at t=0. const Real tpr = WarpX::gamma_boost*t_lab - uz_boost*z_lab/csq; // Adjust the position, taking away its motion from its own velocity and adding // the motion from the average velocity - pstruct[i].pos(2) += tpr*vzpr - tpr*vzbeam_ave_boosted; + zp += tpr*vzpr - tpr*vzbeam_ave_boosted; + + set_position(i, xp, yp, zp); } } @@ -140,13 +146,16 @@ RigidInjectedParticleContainer::BoostandRemapParticles() auto& uyp = attribs[PIdx::uy]; auto& uzp = attribs[PIdx::uz]; - auto& aos = pti.GetArrayOfStructs(); - ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); - + const auto get_position = GetPosition(pti); + auto set_position = SetPosition(pti); + // Loop over particles const long np = pti.numParticles(); for (int i=0 ; i < np ; i++) { + Real xp, yp, zp; + get_position(i, xp, yp, zp); + const Real gamma_lab = std::sqrt(1. + (uxp[i]*uxp[i] + uyp[i]*uyp[i] + uzp[i]*uzp[i])/(PhysConst::c*PhysConst::c)); const Real vx_lab = uxp[i]/gamma_lab; @@ -156,24 +165,24 @@ RigidInjectedParticleContainer::BoostandRemapParticles() // t0_lab is the time in the lab frame that the particles reaches z=0 // The location and time (z=0, t=0) is a synchronization point between the // lab and boosted frames. - const Real t0_lab = -pstruct[i].pos(2)/vz_lab; + const Real t0_lab = -zp/vz_lab; if (!projected) { - pstruct[i].pos(0) += t0_lab*vx_lab; - pstruct[i].pos(1) += t0_lab*vy_lab; + xp += t0_lab*vx_lab; + yp += t0_lab*vy_lab; } if (focused) { // Correct for focusing effect from shift from z=0 to zinject const Real tfocus = -zinject_plane*WarpX::gamma_boost/vz_lab; - pstruct[i].pos(0) -= tfocus*vx_lab; - pstruct[i].pos(1) -= tfocus*vy_lab; + xp -= tfocus*vx_lab; + yp -= tfocus*vy_lab; } // Time of the particle in the boosted frame given its position in the lab frame at t=0. - const Real tpr = -WarpX::gamma_boost*WarpX::beta_boost*pstruct[i].pos(2)/PhysConst::c; + const Real tpr = -WarpX::gamma_boost*WarpX::beta_boost*zp/PhysConst::c; // Position of the particle in the boosted frame given its position in the lab frame at t=0. - const Real zpr = WarpX::gamma_boost*pstruct[i].pos(2); + const Real zpr = WarpX::gamma_boost*zp; // Momentum of the particle in the boosted frame (assuming that it is fixed). uzp[i] = WarpX::gamma_boost*(uzp[i] - WarpX::beta_boost*PhysConst::c*gamma_lab); @@ -181,15 +190,16 @@ RigidInjectedParticleContainer::BoostandRemapParticles() // Put the particle at the location in the boosted frame at boost frame t=0, if (rigid_advance) { // with the particle moving at the average velocity - pstruct[i].pos(2) = zpr - vzbeam_ave_boosted*tpr; + zp = zpr - vzbeam_ave_boosted*tpr; } else { // with the particle moving with its own velocity const Real gammapr = std::sqrt(1. + (uxp[i]*uxp[i] + uyp[i]*uyp[i] + uzp[i]*uzp[i])/(PhysConst::c*PhysConst::c)); const Real vzpr = uzp[i]/gammapr; - pstruct[i].pos(2) = zpr - vzpr*tpr; + zp = zpr - vzpr*tpr; } - + + set_position(i, xp, yp, zp); } } } -- cgit v1.2.3 From f91a7109ed9091cfd04226c247dde57ea4dd02fe Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Tue, 28 Jan 2020 13:53:47 -0800 Subject: fix EOL whitespace --- Source/Laser/LaserParticleContainer.cpp | 4 +- Source/Particles/Deposition/ChargeDeposition.H | 2 +- Source/Particles/Deposition/CurrentDeposition.H | 4 +- Source/Particles/Gather/FieldGather.H | 2 +- Source/Particles/PhotonParticleContainer.cpp | 2 +- Source/Particles/PhysicalParticleContainer.cpp | 6 +- Source/Particles/Pusher/GetAndSetPosition.H | 10 ++-- .../Particles/RigidInjectedParticleContainer.cpp | 70 ++++++++++++---------- Source/Particles/WarpXParticleContainer.cpp | 4 +- 9 files changed, 55 insertions(+), 49 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Laser/LaserParticleContainer.cpp b/Source/Laser/LaserParticleContainer.cpp index c647b520f..658e56db4 100644 --- a/Source/Laser/LaserParticleContainer.cpp +++ b/Source/Laser/LaserParticleContainer.cpp @@ -627,7 +627,7 @@ LaserParticleContainer::update_laser_particle(WarpXParIter& pti, { const auto get_position = GetPosition(pti); auto set_position = SetPosition(pti); - + Real tmp_p_X_0 = p_X[0]; Real tmp_p_X_1 = p_X[1]; Real tmp_p_X_2 = p_X[2]; @@ -660,7 +660,7 @@ LaserParticleContainer::update_laser_particle(WarpXParIter& pti, puxp[i] = gamma * vx; puyp[i] = gamma * vy; puzp[i] = gamma * vz; - + // Push the the particle positions Real x, y, z; get_position(i, x, y, z); diff --git a/Source/Particles/Deposition/ChargeDeposition.H b/Source/Particles/Deposition/ChargeDeposition.H index 9e3df600b..a2eff20b3 100755 --- a/Source/Particles/Deposition/ChargeDeposition.H +++ b/Source/Particles/Deposition/ChargeDeposition.H @@ -57,7 +57,7 @@ void doChargeDepositionShapeN(const GetPosition& get_position, amrex::Real xp, yp, zp; get_position(ip, xp, yp, zp); - + // --- Compute shape factors // x direction // Get particle position in grid coordinates diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H index 8e7f1a552..c7e79bce6 100644 --- a/Source/Particles/Deposition/CurrentDeposition.H +++ b/Source/Particles/Deposition/CurrentDeposition.H @@ -90,7 +90,7 @@ void doDepositionShapeN(const GetPosition& get_position, amrex::Real xp, yp, zp; get_position(ip, xp, yp, zp); - + const amrex::Real vx = uxp[ip]*gaminv; const amrex::Real vy = uyp[ip]*gaminv; const amrex::Real vz = uzp[ip]*gaminv; @@ -313,7 +313,7 @@ void doEsirkepovDepositionShapeN (const GetPosition& get_position, Real xp, yp, zp; get_position(ip, xp, yp, zp); - + Real const wqx = wq*invdtdx; #if (defined WARPX_DIM_3D) Real const wqy = wq*invdtdy; diff --git a/Source/Particles/Gather/FieldGather.H b/Source/Particles/Gather/FieldGather.H index 369b648c1..abdc068b4 100644 --- a/Source/Particles/Gather/FieldGather.H +++ b/Source/Particles/Gather/FieldGather.H @@ -75,7 +75,7 @@ void doGatherShapeN(const GetPosition& get_position, amrex::Real xp, yp, zp; get_position(ip, xp, yp, zp); - + // --- Compute shape factors // x direction // Get particle position diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index 7e7451ab4..590b8b9dc 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -74,7 +74,7 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, Real dt, DtType a_dt_type) const auto get_position = GetPosition(pti); auto set_position = SetPosition(pti); - + amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 729ab7044..7a14e191e 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1746,7 +1746,7 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_type) Real x, y, z; get_position(i, x, y, z); UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt ); - set_position(i, x, y, z); + set_position(i, x, y, z); } ); } else { @@ -1934,7 +1934,7 @@ void PhysicalParticleContainer::copy_attribs (WarpXParIter& pti) ParticleReal* AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); const auto get_position = GetPosition(pti); - + ParallelFor( np, [=] AMREX_GPU_DEVICE (long i) { Real x, y, z; @@ -2213,7 +2213,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, box.grow(ngE); const auto get_position = GetPosition(pti, offset); - + // Lower corner of tile box physical domain const std::array& xyzmin = WarpX::LowerCorner(box, gather_lev); diff --git a/Source/Particles/Pusher/GetAndSetPosition.H b/Source/Particles/Pusher/GetAndSetPosition.H index a429a98fc..92b37bf1c 100644 --- a/Source/Particles/Pusher/GetAndSetPosition.H +++ b/Source/Particles/Pusher/GetAndSetPosition.H @@ -9,7 +9,7 @@ struct GetPosition { using PType = WarpXParticleContainer::ParticleType; using RType = amrex::ParticleReal; - + const PType* AMREX_RESTRICT m_structs; #if (defined WARPX_DIM_RZ) const RType* m_theta; @@ -30,7 +30,7 @@ struct GetPosition AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void operator() (const int i, RType& x, RType& y, RType& z) const noexcept - { + { #ifdef WARPX_DIM_RZ RType r = m_structs[i].pos(0); x = r*std::cos(m_theta[i]); @@ -43,8 +43,8 @@ struct GetPosition #else x = m_structs[i].pos(0); y = m_snan; - z = m_structs[i].pos(2); -#endif + z = m_structs[i].pos(2); +#endif } }; @@ -66,7 +66,7 @@ struct SetPosition m_theta = soa.GetRealData(PIdx::theta).dataPtr() + a_offset; #endif } - + AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void operator() (const int i, RType x, RType y, RType z) const noexcept { diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp index 2ab967eeb..2bf4ee59d 100644 --- a/Source/Particles/RigidInjectedParticleContainer.cpp +++ b/Source/Particles/RigidInjectedParticleContainer.cpp @@ -88,14 +88,14 @@ RigidInjectedParticleContainer::RemapParticles() const auto get_position = GetPosition(pti); auto set_position = SetPosition(pti); - + // Loop over particles const long np = pti.numParticles(); for (int i=0 ; i < np ; i++) { Real xp, yp, zp; get_position(i, xp, yp, zp); - + const Real gammapr = std::sqrt(1. + (uxp[i]*uxp[i] + uyp[i]*uyp[i] + uzp[i]*uzp[i])/csq); const Real vzpr = uzp[i]/gammapr; @@ -148,14 +148,14 @@ RigidInjectedParticleContainer::BoostandRemapParticles() const auto get_position = GetPosition(pti); auto set_position = SetPosition(pti); - + // Loop over particles const long np = pti.numParticles(); for (int i=0 ; i < np ; i++) { Real xp, yp, zp; get_position(i, xp, yp, zp); - + const Real gamma_lab = std::sqrt(1. + (uxp[i]*uxp[i] + uyp[i]*uyp[i] + uzp[i]*uzp[i])/(PhysConst::c*PhysConst::c)); const Real vx_lab = uxp[i]/gamma_lab; @@ -198,7 +198,7 @@ RigidInjectedParticleContainer::BoostandRemapParticles() const Real vzpr = uzp[i]/gammapr; zp = zpr - vzpr*tpr; } - + set_position(i, xp, yp, zp); } } @@ -219,8 +219,8 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ Gpu::ManagedDeviceVector xp_save, yp_save, zp_save; RealVector uxp_save, uyp_save, uzp_save; - auto& aos = pti.GetArrayOfStructs(); - ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); + const auto get_position = GetPosition(pti); + auto set_position = SetPosition(pti); ParticleReal* const AMREX_RESTRICT ux = uxp.dataPtr(); ParticleReal* const AMREX_RESTRICT uy = uyp.dataPtr(); @@ -255,9 +255,11 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ amrex::ParallelFor( np, [=] AMREX_GPU_DEVICE (long i) { - xp_save_ptr[i] = pstruct[i].pos(0); - yp_save_ptr[i] = pstruct[i].pos(1); - zp_save_ptr[i] = pstruct[i].pos(2); + Real xp, yp, zp; + get_position(i, xp, yp, zp); + xp_save_ptr[i] = xp; + yp_save_ptr[i] = yp; + zp_save_ptr[i] = zp; uxp_save_ptr[i] = ux[i]; uyp_save_ptr[i] = uy[i]; uzp_save_ptr[i] = uz[i]; @@ -272,17 +274,19 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ const Real vz_ave_boosted = vzbeam_ave_boosted; amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { - const Real dtscale = dt - (z_plane_previous - pstruct[i].pos(2))/(vz_ave_boosted + v_boost); - if (0. < dtscale && dtscale < dt) { - Exp[i] *= dtscale; - Eyp[i] *= dtscale; - Ezp[i] *= dtscale; - Bxp[i] *= dtscale; - Byp[i] *= dtscale; - Bzp[i] *= dtscale; - } - } - ); + Real xp, yp, zp; + get_position(i, xp, yp, zp); + const Real dtscale = dt - (z_plane_previous - zp)/(vz_ave_boosted + v_boost); + if (0. < dtscale && dtscale < dt) { + Exp[i] *= dtscale; + Eyp[i] *= dtscale; + Ezp[i] *= dtscale; + Bxp[i] *= dtscale; + Byp[i] *= dtscale; + Bzp[i] *= dtscale; + } + } + ); } PhysicalParticleContainer::PushPX(pti, dt, a_dt_type); @@ -304,19 +308,22 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ const Real inv_csq = 1./(PhysConst::c*PhysConst::c); amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { - if (pstruct[i].pos(2) <= z_plane_lev) { + Real xp, yp, zp; + get_position(i, xp, yp, zp); + if (zp <= z_plane_lev) { ux[i] = ux_save[i]; uy[i] = uy_save[i]; uz[i] = uz_save[i]; - pstruct[i].pos(0) = x_save[i]; - pstruct[i].pos(1) = y_save[i]; + xp = x_save[i]; + yp = y_save[i]; if (rigid) { - pstruct[i].pos(2) = z_save[i] + dt*vz_ave_boosted; + zp = z_save[i] + dt*vz_ave_boosted; } else { const Real gi = 1./std::sqrt(1. + (ux[i]*ux[i] + uy[i]*uy[i] + uz[i]*uz[i])*inv_csq); - pstruct[i].pos(2) = z_save[i] + dt*uz[i]*gi; + zp = z_save[i] + dt*uz[i]*gi; } + set_position(i, xp, yp, zp); } }); } @@ -415,9 +422,7 @@ 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 auto& aos = pti.GetArrayOfStructs(); - const ParticleType* AMREX_RESTRICT const pstruct = aos().dataPtr(); - + const auto get_position = GetPosition(pti); ParticleReal* const AMREX_RESTRICT uxpp = uxp.dataPtr(); ParticleReal* const AMREX_RESTRICT uypp = uyp.dataPtr(); ParticleReal* const AMREX_RESTRICT uzpp = uzp.dataPtr(); @@ -484,14 +489,15 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, const ParticleReal zz = zinject_plane_levels[lev]; amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { - if (pstruct[i].pos(2) <= zz) { + ParticleReal xp, yp, zp; + get_position(i, xp, yp, zp); + if (zp <= zz) { uxpp[i] = ux_save[i]; uypp[i] = uy_save[i]; uzpp[i] = uz_save[i]; } } - ); - + ); } } } diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 779771360..db32f2bf4 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -287,7 +287,7 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, // (same for jx and jz) const auto get_position = GetPosition(pti, offset); - + // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow const Dim3 lo = lbound(tilebox); @@ -735,7 +735,7 @@ WarpXParticleContainer::PushX (int lev, amrex::Real dt) const auto get_position = GetPosition(pti); auto set_position = SetPosition(pti); - + // - momenta are stored as a struct of array, in `attribs` auto& attribs = pti.GetAttribs(); ParticleReal* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); -- cgit v1.2.3 From 2ef268b10a924644d10996daa069befea843879f Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 29 Jan 2020 12:23:59 -0800 Subject: Real -> ParticleReal --- Source/Laser/LaserParticleContainer.cpp | 4 ++-- Source/Particles/Deposition/ChargeDeposition.H | 2 +- Source/Particles/Deposition/CurrentDeposition.H | 4 ++-- Source/Particles/Gather/FieldGather.H | 2 +- Source/Particles/PhotonParticleContainer.cpp | 2 +- Source/Particles/PhysicalParticleContainer.cpp | 20 ++++++++++---------- Source/Particles/RigidInjectedParticleContainer.cpp | 10 +++++----- Source/Particles/WarpXParticleContainer.cpp | 2 +- 8 files changed, 23 insertions(+), 23 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Laser/LaserParticleContainer.cpp b/Source/Laser/LaserParticleContainer.cpp index 658e56db4..0203cc4b2 100644 --- a/Source/Laser/LaserParticleContainer.cpp +++ b/Source/Laser/LaserParticleContainer.cpp @@ -585,7 +585,7 @@ LaserParticleContainer::calculate_laser_plane_coordinates (const WarpXParIter& p amrex::ParallelFor( np, [=] AMREX_GPU_DEVICE (int i) { - Real x, y, z; + ParticleReal x, y, z; get_position(i, x, y, z); #if (defined WARPX_DIM_3D) || (defined WARPX_DIM_RZ) pplane_Xp[i] = @@ -662,7 +662,7 @@ LaserParticleContainer::update_laser_particle(WarpXParIter& pti, puzp[i] = gamma * vz; // Push the the particle positions - Real x, y, z; + ParticleReal x, y, z; get_position(i, x, y, z); x += vx * dt; #if (defined WARPX_DIM_3D) || (defined WARPX_DIM_RZ) diff --git a/Source/Particles/Deposition/ChargeDeposition.H b/Source/Particles/Deposition/ChargeDeposition.H index a2eff20b3..f97866f03 100755 --- a/Source/Particles/Deposition/ChargeDeposition.H +++ b/Source/Particles/Deposition/ChargeDeposition.H @@ -55,7 +55,7 @@ void doChargeDepositionShapeN(const GetPosition& get_position, wq *= ion_lev[ip]; } - amrex::Real xp, yp, zp; + amrex::ParticleReal xp, yp, zp; get_position(ip, xp, yp, zp); // --- Compute shape factors diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H index c7e79bce6..e4cb8c1cd 100644 --- a/Source/Particles/Deposition/CurrentDeposition.H +++ b/Source/Particles/Deposition/CurrentDeposition.H @@ -88,7 +88,7 @@ void doDepositionShapeN(const GetPosition& get_position, wq *= ion_lev[ip]; } - amrex::Real xp, yp, zp; + amrex::ParticleReal xp, yp, zp; get_position(ip, xp, yp, zp); const amrex::Real vx = uxp[ip]*gaminv; @@ -311,7 +311,7 @@ void doEsirkepovDepositionShapeN (const GetPosition& get_position, wq *= ion_lev[ip]; } - Real xp, yp, zp; + ParticleReal xp, yp, zp; get_position(ip, xp, yp, zp); Real const wqx = wq*invdtdx; diff --git a/Source/Particles/Gather/FieldGather.H b/Source/Particles/Gather/FieldGather.H index abdc068b4..2d2b85271 100644 --- a/Source/Particles/Gather/FieldGather.H +++ b/Source/Particles/Gather/FieldGather.H @@ -73,7 +73,7 @@ void doGatherShapeN(const GetPosition& get_position, np_to_gather, [=] AMREX_GPU_DEVICE (long ip) { - amrex::Real xp, yp, zp; + amrex::ParticleReal xp, yp, zp; get_position(ip, xp, yp, zp); // --- Compute shape factors diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index 590b8b9dc..1e34b4c97 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -78,7 +78,7 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, Real dt, DtType a_dt_type) amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { - Real x, y, z; + ParticleReal x, y, z; get_position(i, x, y, z); UpdatePositionPhoton( x, y, z, ux[i], uy[i], uz[i], dt ); set_position(i, x, y, z); diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index f4cdbcb92..0daa9fb5f 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1000,7 +1000,7 @@ PhysicalParticleContainer::AssignExternalFieldOnParticles(WarpXParIter& pti, Real time = warpx.gett_new(lev); amrex::ParallelFor(pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { - Real x, y, z; + ParticleReal x, y, z; get_position(i, x, y, z); Exp_data[i] = xfield_partparser->getField(x, y, z, time); Eyp_data[i] = yfield_partparser->getField(x, y, z, time); @@ -1022,7 +1022,7 @@ PhysicalParticleContainer::AssignExternalFieldOnParticles(WarpXParIter& pti, Real time = warpx.gett_new(lev); amrex::ParallelFor(pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { - Real x, y, z; + ParticleReal x, y, z; get_position(i, x, y, z); Bxp_data[i] = xfield_partparser->getField(x, y, z, time); Byp_data[i] = yfield_partparser->getField(x, y, z, time); @@ -1489,7 +1489,7 @@ PhysicalParticleContainer::SplitParticles(int lev) auto& uzp = attribs[PIdx::uz]; const long np = pti.numParticles(); for(int i=0; i Date: Wed, 29 Jan 2020 12:54:42 -0800 Subject: rename get/set particle position functors --- Source/Laser/LaserParticleContainer.cpp | 12 ++-- Source/Particles/Deposition/ChargeDeposition.H | 6 +- Source/Particles/Deposition/CurrentDeposition.H | 12 ++-- Source/Particles/Gather/FieldGather.H | 6 +- Source/Particles/PhotonParticleContainer.cpp | 8 +-- Source/Particles/PhysicalParticleContainer.cpp | 64 +++++++++++----------- Source/Particles/Pusher/GetAndSetPosition.H | 8 +-- .../Particles/RigidInjectedParticleContainer.cpp | 32 +++++------ Source/Particles/WarpXParticleContainer.cpp | 30 +++++----- 9 files changed, 89 insertions(+), 89 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Laser/LaserParticleContainer.cpp b/Source/Laser/LaserParticleContainer.cpp index 0203cc4b2..09dd8814b 100644 --- a/Source/Laser/LaserParticleContainer.cpp +++ b/Source/Laser/LaserParticleContainer.cpp @@ -568,7 +568,7 @@ LaserParticleContainer::calculate_laser_plane_coordinates (const WarpXParIter& p Real * AMREX_RESTRICT const pplane_Xp, Real * AMREX_RESTRICT const pplane_Yp) { - const auto get_position = GetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); Real tmp_u_X_0 = u_X[0]; Real tmp_u_X_2 = u_X[2]; @@ -586,7 +586,7 @@ LaserParticleContainer::calculate_laser_plane_coordinates (const WarpXParIter& p np, [=] AMREX_GPU_DEVICE (int i) { ParticleReal x, y, z; - get_position(i, x, y, z); + GetPosition(i, x, y, z); #if (defined WARPX_DIM_3D) || (defined WARPX_DIM_RZ) pplane_Xp[i] = tmp_u_X_0 * (x - tmp_position_0) + @@ -625,8 +625,8 @@ LaserParticleContainer::update_laser_particle(WarpXParIter& pti, Real const * AMREX_RESTRICT const amplitude, const Real dt) { - const auto get_position = GetPosition(pti); - auto set_position = SetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); + auto SetPosition = SetParticlePosition(pti); Real tmp_p_X_0 = p_X[0]; Real tmp_p_X_1 = p_X[1]; @@ -663,13 +663,13 @@ LaserParticleContainer::update_laser_particle(WarpXParIter& pti, // Push the the particle positions ParticleReal x, y, z; - get_position(i, x, y, z); + GetPosition(i, x, y, z); x += vx * dt; #if (defined WARPX_DIM_3D) || (defined WARPX_DIM_RZ) y += vy * dt; #endif z += vz * dt; - set_position(i, x, y, z); + SetPosition(i, x, y, z); } ); } diff --git a/Source/Particles/Deposition/ChargeDeposition.H b/Source/Particles/Deposition/ChargeDeposition.H index f97866f03..f1d737994 100755 --- a/Source/Particles/Deposition/ChargeDeposition.H +++ b/Source/Particles/Deposition/ChargeDeposition.H @@ -5,7 +5,7 @@ #include "ShapeFactors.H" /* \brief Charge Deposition for thread thread_num - * /param get_position : A functor for returning the particle position. + * /param GetPosition : A functor for returning the particle position. * \param wp : Pointer to array of particle weights. * \param ion_lev : Pointer to array of particle ionization level. This is required to have the charge of each macroparticle @@ -19,7 +19,7 @@ * /param q : species charge. */ template -void doChargeDepositionShapeN(const GetPosition& get_position, +void doChargeDepositionShapeN(const GetParticlePosition& GetPosition, const amrex::ParticleReal * const wp, const int * const ion_lev, const amrex::Array4& rho_arr, @@ -56,7 +56,7 @@ void doChargeDepositionShapeN(const GetPosition& get_position, } amrex::ParticleReal xp, yp, zp; - get_position(ip, xp, yp, zp); + GetPosition(ip, xp, yp, zp); // --- Compute shape factors // x direction diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H index e4cb8c1cd..97a707d1c 100644 --- a/Source/Particles/Deposition/CurrentDeposition.H +++ b/Source/Particles/Deposition/CurrentDeposition.H @@ -10,7 +10,7 @@ /** * \brief Current Deposition for thread thread_num - * /param get_position : A functor for returning the particle position. + * /param GetPosition : A functor for returning the particle position. * \param wp : Pointer to array of particle weights. * \param uxp uyp uzp : Pointer to arrays of particle momentum. * \param ion_lev : Pointer to array of particle ionization level. This is @@ -28,7 +28,7 @@ * /param q : species charge. */ template -void doDepositionShapeN(const GetPosition& get_position, +void doDepositionShapeN(const GetParticlePosition& GetPosition, const amrex::ParticleReal * const wp, const amrex::ParticleReal * const uxp, const amrex::ParticleReal * const uyp, @@ -89,7 +89,7 @@ void doDepositionShapeN(const GetPosition& get_position, } amrex::ParticleReal xp, yp, zp; - get_position(ip, xp, yp, zp); + GetPosition(ip, xp, yp, zp); const amrex::Real vx = uxp[ip]*gaminv; const amrex::Real vy = uyp[ip]*gaminv; @@ -226,7 +226,7 @@ void doDepositionShapeN(const GetPosition& get_position, /** * \brief Esirkepov Current Deposition for thread thread_num * - * /param get_position : A functor for returning the particle position. + * /param GetPosition : A functor for returning the particle position. * \param wp : Pointer to array of particle weights. * \param uxp uyp uzp : Pointer to arrays of particle momentum. * \param ion_lev : Pointer to array of particle ionization level. This is @@ -245,7 +245,7 @@ void doDepositionShapeN(const GetPosition& get_position, * \param n_rz_azimuthal_modes: Number of azimuthal modes when using RZ geometry */ template -void doEsirkepovDepositionShapeN (const GetPosition& get_position, +void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, const amrex::ParticleReal * const wp, const amrex::ParticleReal * const uxp, const amrex::ParticleReal * const uyp, @@ -312,7 +312,7 @@ void doEsirkepovDepositionShapeN (const GetPosition& get_position, } ParticleReal xp, yp, zp; - get_position(ip, xp, yp, zp); + GetPosition(ip, xp, yp, zp); Real const wqx = wq*invdtdx; #if (defined WARPX_DIM_3D) diff --git a/Source/Particles/Gather/FieldGather.H b/Source/Particles/Gather/FieldGather.H index 2d2b85271..1bac5d3aa 100644 --- a/Source/Particles/Gather/FieldGather.H +++ b/Source/Particles/Gather/FieldGather.H @@ -7,7 +7,7 @@ /** * \brief Field gather for particles handled by thread thread_num - * /param get_position : A functor for returning the particle position. + * /param GetPosition : A functor for returning the particle position. * \param Exp, Eyp, Ezp: Pointer to array of electric field on particles. * \param Bxp, Byp, Bzp: Pointer to array of magnetic field on particles. * \param ex_arr ey_arr: Array4 of current density, either full array or tile. @@ -21,7 +21,7 @@ * \param n_rz_azimuthal_modes: Number of azimuthal modes when using RZ geometry */ template -void doGatherShapeN(const GetPosition& get_position, +void doGatherShapeN(const GetParticlePosition& GetPosition, amrex::ParticleReal * const Exp, amrex::ParticleReal * const Eyp, amrex::ParticleReal * const Ezp, amrex::ParticleReal * const Bxp, amrex::ParticleReal * const Byp, amrex::ParticleReal * const Bzp, @@ -74,7 +74,7 @@ void doGatherShapeN(const GetPosition& get_position, [=] AMREX_GPU_DEVICE (long ip) { amrex::ParticleReal xp, yp, zp; - get_position(ip, xp, yp, zp); + GetPosition(ip, xp, yp, zp); // --- Compute shape factors // x direction diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index 1e34b4c97..92ccc17e7 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -72,16 +72,16 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, Real dt, DtType a_dt_type) copy_attribs(pti); } - const auto get_position = GetPosition(pti); - auto set_position = SetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); + auto SetPosition = SetParticlePosition(pti); amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleReal x, y, z; - get_position(i, x, y, z); + GetPosition(i, x, y, z); UpdatePositionPhoton( x, y, z, ux[i], uy[i], uz[i], dt ); - set_position(i, x, y, z); + SetPosition(i, x, y, z); } ); } diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 0daa9fb5f..55de2131d 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -990,7 +990,7 @@ PhysicalParticleContainer::AssignExternalFieldOnParticles(WarpXParIter& pti, Bzp.assign(np,mypc.m_B_external_particle[2]); } if (mypc.m_E_ext_particle_s=="parse_e_ext_particle_function") { - const auto get_position = GetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); Real* const AMREX_RESTRICT Exp_data = Exp.dataPtr(); Real* const AMREX_RESTRICT Eyp_data = Eyp.dataPtr(); Real* const AMREX_RESTRICT Ezp_data = Ezp.dataPtr(); @@ -1001,7 +1001,7 @@ PhysicalParticleContainer::AssignExternalFieldOnParticles(WarpXParIter& pti, amrex::ParallelFor(pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleReal x, y, z; - get_position(i, x, y, z); + GetPosition(i, x, y, z); Exp_data[i] = xfield_partparser->getField(x, y, z, time); Eyp_data[i] = yfield_partparser->getField(x, y, z, time); Ezp_data[i] = zfield_partparser->getField(x, y, z, time); @@ -1012,7 +1012,7 @@ PhysicalParticleContainer::AssignExternalFieldOnParticles(WarpXParIter& pti, ); } if (mypc.m_B_ext_particle_s=="parse_b_ext_particle_function") { - const auto get_position = GetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); Real* const AMREX_RESTRICT Bxp_data = Bxp.dataPtr(); Real* const AMREX_RESTRICT Byp_data = Byp.dataPtr(); Real* const AMREX_RESTRICT Bzp_data = Bzp.dataPtr(); @@ -1023,7 +1023,7 @@ PhysicalParticleContainer::AssignExternalFieldOnParticles(WarpXParIter& pti, amrex::ParallelFor(pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleReal x, y, z; - get_position(i, x, y, z); + GetPosition(i, x, y, z); Bxp_data[i] = xfield_partparser->getField(x, y, z, time); Byp_data[i] = yfield_partparser->getField(x, y, z, time); Bzp_data[i] = zfield_partparser->getField(x, y, z, time); @@ -1464,7 +1464,7 @@ PhysicalParticleContainer::SplitParticles(int lev) // Loop over particle interator for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { - const auto get_position = GetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); const amrex::Vector ppc_nd = plasma_injector->num_particles_per_cell_each_dim; const std::array& dx = WarpX::CellSize(lev); @@ -1490,7 +1490,7 @@ PhysicalParticleContainer::SplitParticles(int lev) const long np = pti.numParticles(); for(int i=0; i= z_new) && (zpold[i] <= z_old)) || ((zp <= z_new) && (zpold[i] >= z_old))) ) @@ -2090,7 +2090,7 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real [=] AMREX_GPU_DEVICE(int i) { ParticleReal xp_new, yp_new, zp_new; - get_position(i, xp_new, yp_new, zp_new); + GetPosition(i, xp_new, yp_new, zp_new); if (Flag[i] == 1) { // Lorentz Transform particles to lab-frame @@ -2215,7 +2215,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, // Add guard cells to the box. box.grow(ngE); - const auto get_position = GetPosition(pti, offset); + const auto GetPosition = GetParticlePosition(pti, offset); // Lower corner of tile box physical domain const std::array& xyzmin = WarpX::LowerCorner(box, gather_lev); @@ -2226,7 +2226,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, // different versions of template function doGatherShapeN if (WarpX::l_lower_order_in_v){ if (WarpX::nox == 1){ - doGatherShapeN<1,1>(get_position, + doGatherShapeN<1,1>(GetPosition, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2234,7 +2234,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, np_to_gather, dx, xyzmin, lo, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ - doGatherShapeN<2,1>(get_position, + doGatherShapeN<2,1>(GetPosition, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2242,7 +2242,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, np_to_gather, dx, xyzmin, lo, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ - doGatherShapeN<3,1>(get_position, + doGatherShapeN<3,1>(GetPosition, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2252,7 +2252,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, } } else { if (WarpX::nox == 1){ - doGatherShapeN<1,0>(get_position, + doGatherShapeN<1,0>(GetPosition, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2260,7 +2260,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, np_to_gather, dx, xyzmin, lo, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ - doGatherShapeN<2,0>(get_position, + doGatherShapeN<2,0>(GetPosition, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, @@ -2268,7 +2268,7 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, np_to_gather, dx, xyzmin, lo, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ - doGatherShapeN<3,0>(get_position, + doGatherShapeN<3,0>(GetPosition, Exp.dataPtr() + offset, Eyp.dataPtr() + offset, Ezp.dataPtr() + offset, Bxp.dataPtr() + offset, Byp.dataPtr() + offset, Bzp.dataPtr() + offset, diff --git a/Source/Particles/Pusher/GetAndSetPosition.H b/Source/Particles/Pusher/GetAndSetPosition.H index 785570a0c..d8a927c74 100644 --- a/Source/Particles/Pusher/GetAndSetPosition.H +++ b/Source/Particles/Pusher/GetAndSetPosition.H @@ -5,7 +5,7 @@ #include #include -struct GetPosition +struct GetParticlePosition { using PType = WarpXParticleContainer::ParticleType; using RType = amrex::ParticleReal; @@ -16,7 +16,7 @@ struct GetPosition #elif (AMREX_SPACEDIM == 2) static constexpr RType m_snan = std::numeric_limits::quiet_NaN(); #endif - GetPosition (const WarpXParIter& a_pti, int a_offset = 0) noexcept + GetParticlePosition (const WarpXParIter& a_pti, int a_offset = 0) noexcept { const auto& aos = a_pti.GetArrayOfStructs(); m_structs = aos().dataPtr() + a_offset; @@ -48,7 +48,7 @@ struct GetPosition } }; -struct SetPosition +struct SetParticlePosition { using PType = WarpXParticleContainer::ParticleType; using RType = amrex::ParticleReal; @@ -57,7 +57,7 @@ struct SetPosition #if (defined WARPX_DIM_RZ) RType* AMREX_RESTRICT m_theta; #endif - SetPosition (WarpXParIter& a_pti, int a_offset = 0) noexcept + SetParticlePosition (WarpXParIter& a_pti, int a_offset = 0) noexcept { auto& aos = a_pti.GetArrayOfStructs(); m_structs = aos().dataPtr() + a_offset; diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp index 692b1f287..eb1677194 100644 --- a/Source/Particles/RigidInjectedParticleContainer.cpp +++ b/Source/Particles/RigidInjectedParticleContainer.cpp @@ -86,15 +86,15 @@ RigidInjectedParticleContainer::RemapParticles() auto& uyp = attribs[PIdx::uy]; auto& uzp = attribs[PIdx::uz]; - const auto get_position = GetPosition(pti); - auto set_position = SetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); + auto SetPosition = SetParticlePosition(pti); // Loop over particles const long np = pti.numParticles(); for (int i=0 ; i < np ; i++) { ParticleReal xp, yp, zp; - get_position(i, xp, yp, zp); + GetPosition(i, xp, yp, zp); const Real gammapr = std::sqrt(1. + (uxp[i]*uxp[i] + uyp[i]*uyp[i] + uzp[i]*uzp[i])/csq); const Real vzpr = uzp[i]/gammapr; @@ -109,7 +109,7 @@ RigidInjectedParticleContainer::RemapParticles() // the motion from the average velocity zp += tpr*vzpr - tpr*vzbeam_ave_boosted; - set_position(i, xp, yp, zp); + SetPosition(i, xp, yp, zp); } } @@ -146,15 +146,15 @@ RigidInjectedParticleContainer::BoostandRemapParticles() auto& uyp = attribs[PIdx::uy]; auto& uzp = attribs[PIdx::uz]; - const auto get_position = GetPosition(pti); - auto set_position = SetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); + auto SetPosition = SetParticlePosition(pti); // Loop over particles const long np = pti.numParticles(); for (int i=0 ; i < np ; i++) { ParticleReal xp, yp, zp; - get_position(i, xp, yp, zp); + GetPosition(i, xp, yp, zp); const Real gamma_lab = std::sqrt(1. + (uxp[i]*uxp[i] + uyp[i]*uyp[i] + uzp[i]*uzp[i])/(PhysConst::c*PhysConst::c)); @@ -199,7 +199,7 @@ RigidInjectedParticleContainer::BoostandRemapParticles() zp = zpr - vzpr*tpr; } - set_position(i, xp, yp, zp); + SetPosition(i, xp, yp, zp); } } } @@ -219,8 +219,8 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ Gpu::ManagedDeviceVector xp_save, yp_save, zp_save; RealVector uxp_save, uyp_save, uzp_save; - const auto get_position = GetPosition(pti); - auto set_position = SetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); + auto SetPosition = SetParticlePosition(pti); ParticleReal* const AMREX_RESTRICT ux = uxp.dataPtr(); ParticleReal* const AMREX_RESTRICT uy = uyp.dataPtr(); @@ -256,7 +256,7 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ amrex::ParallelFor( np, [=] AMREX_GPU_DEVICE (long i) { ParticleReal xp, yp, zp; - get_position(i, xp, yp, zp); + GetPosition(i, xp, yp, zp); xp_save_ptr[i] = xp; yp_save_ptr[i] = yp; zp_save_ptr[i] = zp; @@ -275,7 +275,7 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleReal xp, yp, zp; - get_position(i, xp, yp, zp); + GetPosition(i, xp, yp, zp); const Real dtscale = dt - (z_plane_previous - zp)/(vz_ave_boosted + v_boost); if (0. < dtscale && dtscale < dt) { Exp[i] *= dtscale; @@ -309,7 +309,7 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleReal xp, yp, zp; - get_position(i, xp, yp, zp); + GetPosition(i, xp, yp, zp); if (zp <= z_plane_lev) { ux[i] = ux_save[i]; uy[i] = uy_save[i]; @@ -323,7 +323,7 @@ RigidInjectedParticleContainer::PushPX (WarpXParIter& pti, Real dt, DtType a_dt_ const Real gi = 1./std::sqrt(1. + (ux[i]*ux[i] + uy[i]*uy[i] + uz[i]*uz[i])*inv_csq); zp = z_save[i] + dt*uz[i]*gi; } - set_position(i, xp, yp, zp); + SetPosition(i, xp, yp, zp); } }); } @@ -422,7 +422,7 @@ 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 auto get_position = GetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); ParticleReal* const AMREX_RESTRICT uxpp = uxp.dataPtr(); ParticleReal* const AMREX_RESTRICT uypp = uyp.dataPtr(); ParticleReal* const AMREX_RESTRICT uzpp = uzp.dataPtr(); @@ -490,7 +490,7 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleReal xp, yp, zp; - get_position(i, xp, yp, zp); + GetPosition(i, xp, yp, zp); if (zp <= zz) { uxpp[i] = ux_save[i]; uypp[i] = uy_save[i]; diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 581565a5f..a45b764eb 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -286,7 +286,7 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, // CPU, tiling: deposit into local_jx // (same for jx and jz) - const auto get_position = GetPosition(pti, offset); + const auto GetPosition = GetParticlePosition(pti, offset); // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow @@ -297,19 +297,19 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, if (WarpX::current_deposition_algo == CurrentDepositionAlgo::Esirkepov) { if (WarpX::nox == 1){ doEsirkepovDepositionShapeN<1>( - get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, + GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_arr, jy_arr, jz_arr, np_to_depose, dt, dx, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ doEsirkepovDepositionShapeN<2>( - get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, + GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_arr, jy_arr, jz_arr, np_to_depose, dt, dx, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ doEsirkepovDepositionShapeN<3>( - get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, + GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_arr, jy_arr, jz_arr, np_to_depose, dt, dx, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); @@ -317,19 +317,19 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, } else { if (WarpX::nox == 1){ doDepositionShapeN<1>( - get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, + GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_fab, jy_fab, jz_fab, np_to_depose, dt, dx, xyzmin, lo, q); } else if (WarpX::nox == 2){ doDepositionShapeN<2>( - get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, + GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_fab, jy_fab, jz_fab, np_to_depose, dt, dx, xyzmin, lo, q); } else if (WarpX::nox == 3){ doDepositionShapeN<3>( - get_position, wp.dataPtr() + offset, uxp.dataPtr() + offset, + GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, jx_fab, jy_fab, jz_fab, np_to_depose, dt, dx, xyzmin, lo, q); @@ -423,7 +423,7 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector& wp, // GPU, no tiling: deposit directly in rho // CPU, tiling: deposit into local_rho - const auto get_position = GetPosition(pti, offset); + const auto GetPosition = GetParticlePosition(pti, offset); // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow @@ -433,13 +433,13 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector& wp, BL_PROFILE_VAR_START(blp_ppc_chd); if (WarpX::nox == 1){ - doChargeDepositionShapeN<1>(get_position, wp.dataPtr()+offset, ion_lev, + doChargeDepositionShapeN<1>(GetPosition, wp.dataPtr()+offset, ion_lev, rho_arr, np_to_depose, dx, xyzmin, lo, q); } else if (WarpX::nox == 2){ - doChargeDepositionShapeN<2>(get_position, wp.dataPtr()+offset, ion_lev, + doChargeDepositionShapeN<2>(GetPosition, wp.dataPtr()+offset, ion_lev, rho_arr, np_to_depose, dx, xyzmin, lo, q); } else if (WarpX::nox == 3){ - doChargeDepositionShapeN<3>(get_position, wp.dataPtr()+offset, ion_lev, + doChargeDepositionShapeN<3>(GetPosition, wp.dataPtr()+offset, ion_lev, rho_arr, np_to_depose, dx, xyzmin, lo, q); } BL_PROFILE_VAR_STOP(blp_ppc_chd); @@ -733,8 +733,8 @@ WarpXParticleContainer::PushX (int lev, amrex::Real dt) // Particle Push // - const auto get_position = GetPosition(pti); - auto set_position = SetPosition(pti); + const auto GetPosition = GetParticlePosition(pti); + auto SetPosition = SetParticlePosition(pti); // - momenta are stored as a struct of array, in `attribs` auto& attribs = pti.GetAttribs(); @@ -750,9 +750,9 @@ WarpXParticleContainer::PushX (int lev, amrex::Real dt) amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleReal x, y, z; - get_position(i, x, y, z); + GetPosition(i, x, y, z); UpdatePosition(x, y, z, ux[i], uy[i], uz[i], dt); - set_position(i, x, y, z); + SetPosition(i, x, y, z); } ); -- cgit v1.2.3