aboutsummaryrefslogtreecommitdiff
path: root/Source/Particles/RigidInjectedParticleContainer.cpp
diff options
context:
space:
mode:
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 7d129fc01..891ade76d 100644
--- a/Source/Particles/RigidInjectedParticleContainer.cpp
+++ b/Source/Particles/RigidInjectedParticleContainer.cpp
@@ -76,7 +76,7 @@ RigidInjectedParticleContainer::RemapParticles()
// Note that the particles are already in the boosted frame.
// This value is saved to advance the particles not injected yet
- Cuda::ManagedDeviceVector<Real> xp, yp, zp;
+ Cuda::ManagedDeviceVector<ParticleReal> xp, yp, zp;
for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti)
{
@@ -136,7 +136,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)
{
@@ -207,9 +207,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)
{
@@ -220,21 +220,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.
@@ -271,12 +271,12 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti,
if (!done_injecting_lev) {
- Real* AMREX_RESTRICT x_save = xp_save.dataPtr();
- Real* AMREX_RESTRICT y_save = yp_save.dataPtr();
- Real* AMREX_RESTRICT z_save = zp_save.dataPtr();
- Real* AMREX_RESTRICT ux_save = uxp_save.dataPtr();
- Real* AMREX_RESTRICT uy_save = uyp_save.dataPtr();
- Real* AMREX_RESTRICT uz_save = uzp_save.dataPtr();
+ ParticleReal* AMREX_RESTRICT x_save = xp_save.dataPtr();
+ ParticleReal* AMREX_RESTRICT y_save = yp_save.dataPtr();
+ ParticleReal* AMREX_RESTRICT z_save = zp_save.dataPtr();
+ ParticleReal* AMREX_RESTRICT ux_save = uxp_save.dataPtr();
+ ParticleReal* AMREX_RESTRICT uy_save = uyp_save.dataPtr();
+ ParticleReal* AMREX_RESTRICT uz_save = uzp_save.dataPtr();
// Undo the push for particles not injected yet.
// The zp are advanced a fixed amount.
@@ -415,16 +415,16 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt,
// This wraps the momentum advance so that inheritors can modify the call.
// Extract pointers to the different particle quantities
- const Real* const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr();
- Real* const AMREX_RESTRICT uxpp = uxp.dataPtr();
- Real* const AMREX_RESTRICT uypp = uyp.dataPtr();
- Real* const AMREX_RESTRICT uzpp = uzp.dataPtr();
- const Real* const AMREX_RESTRICT Expp = Exp.dataPtr();
- const Real* const AMREX_RESTRICT Eypp = Eyp.dataPtr();
- const Real* const AMREX_RESTRICT Ezpp = Ezp.dataPtr();
- const Real* const AMREX_RESTRICT Bxpp = Bxp.dataPtr();
- const Real* const AMREX_RESTRICT Bypp = Byp.dataPtr();
- const Real* const AMREX_RESTRICT Bzpp = Bzp.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr();
+ ParticleReal* const AMREX_RESTRICT uxpp = uxp.dataPtr();
+ ParticleReal* const AMREX_RESTRICT uypp = uyp.dataPtr();
+ ParticleReal* const AMREX_RESTRICT uzpp = uzp.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT Expp = Exp.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT Eypp = Eyp.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT Ezpp = Ezp.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT Bxpp = Bxp.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT Bypp = Byp.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT Bzpp = Bzp.dataPtr();
// Loop over the particles and update their momentum
const Real q = this->charge;
@@ -450,10 +450,10 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt,
// Undo the push for particles not injected yet.
// It is assumed that PushP will only be called on the first and last steps
// and that no particles will cross zinject_plane.
- const Real* const AMREX_RESTRICT ux_save = uxp_save.dataPtr();
- const Real* const AMREX_RESTRICT uy_save = uyp_save.dataPtr();
- const Real* const AMREX_RESTRICT uz_save = uzp_save.dataPtr();
- const Real zz = zinject_plane_levels[lev];
+ const ParticleReal* const AMREX_RESTRICT ux_save = uxp_save.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT uy_save = uyp_save.dataPtr();
+ const ParticleReal* const AMREX_RESTRICT uz_save = uzp_save.dataPtr();
+ const ParticleReal zz = zinject_plane_levels[lev];
amrex::ParallelFor( pti.numParticles(),
[=] AMREX_GPU_DEVICE (long i) {
if (zp[i] <= zz) {