aboutsummaryrefslogtreecommitdiff
path: root/Source/Particles/RigidInjectedParticleContainer.cpp
diff options
context:
space:
mode:
authorGravatar Luca Fedeli <luca.fedeli@cea.fr> 2019-09-30 09:50:27 +0200
committerGravatar Luca Fedeli <luca.fedeli@cea.fr> 2019-09-30 09:50:27 +0200
commit36175661fd27c812ac9dbdc605407a25108334cb (patch)
treed9e374c7d8e69e9abdf08eba442211dd26b7de86 /Source/Particles/RigidInjectedParticleContainer.cpp
parent8f375c0c07f17b27115b61b95d21cf9bb8f554e7 (diff)
parentc636e0f71c6ca6a854ed45bf90f90b943b34fc58 (diff)
downloadWarpX-36175661fd27c812ac9dbdc605407a25108334cb.tar.gz
WarpX-36175661fd27c812ac9dbdc605407a25108334cb.tar.zst
WarpX-36175661fd27c812ac9dbdc605407a25108334cb.zip
Merge remote-tracking branch 'upstream/dev' into classical_radiation_reaction
Diffstat (limited to 'Source/Particles/RigidInjectedParticleContainer.cpp')
-rw-r--r--Source/Particles/RigidInjectedParticleContainer.cpp76
1 files changed, 38 insertions, 38 deletions
diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp
index 006a0024b..a3ae5f4e8 100644
--- a/Source/Particles/RigidInjectedParticleContainer.cpp
+++ b/Source/Particles/RigidInjectedParticleContainer.cpp
@@ -77,7 +77,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<Real> xp, yp, zp;
+ Cuda::ManagedDeviceVector<ParticleReal> xp, yp, zp;
for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti)
{
@@ -137,7 +137,7 @@ RigidInjectedParticleContainer::BoostandRemapParticles()
#pragma omp parallel
#endif
{
- Cuda::ManagedDeviceVector<Real> xp, yp, zp;
+ Cuda::ManagedDeviceVector<ParticleReal> xp, yp, zp;
for (WarpXParIter pti(*this, 0); pti.isValid(); ++pti)
{
@@ -208,9 +208,9 @@ RigidInjectedParticleContainer::BoostandRemapParticles()
void
RigidInjectedParticleContainer::PushPX(WarpXParIter& pti,
- Cuda::ManagedDeviceVector<Real>& xp,
- Cuda::ManagedDeviceVector<Real>& yp,
- Cuda::ManagedDeviceVector<Real>& zp,
+ Cuda::ManagedDeviceVector<ParticleReal>& xp,
+ Cuda::ManagedDeviceVector<ParticleReal>& yp,
+ Cuda::ManagedDeviceVector<ParticleReal>& zp,
Real dt, DtType a_dt_type)
{
@@ -221,21 +221,21 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti,
auto& uzp = attribs[PIdx::uz];
// Save the position and momenta, making copies
- Cuda::ManagedDeviceVector<Real> xp_save, yp_save, zp_save;
+ Cuda::ManagedDeviceVector<ParticleReal> 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.
@@ -272,12 +272,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.
@@ -416,16 +416,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;
@@ -462,10 +462,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) {