From 8bfa3488bae938d4c1c4ec9f71eababa51556324 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Fri, 8 Mar 2019 11:10:17 -0800 Subject: DeviceVector -> ManagedDeviceVector --- Source/Particles/RigidInjectedParticleContainer.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) (limited to 'Source/Particles/RigidInjectedParticleContainer.cpp') diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp index db3623705..3ee4d87e5 100644 --- a/Source/Particles/RigidInjectedParticleContainer.cpp +++ b/Source/Particles/RigidInjectedParticleContainer.cpp @@ -73,7 +73,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::DeviceVector xp, yp, zp; + Cuda::ManagedDeviceVector xp, yp, zp; for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { @@ -134,7 +134,7 @@ RigidInjectedParticleContainer::BoostandRemapParticles() #pragma omp parallel #endif { - Cuda::DeviceVector xp, yp, zp; + Cuda::ManagedDeviceVector xp, yp, zp; for (WarpXParIter pti(*this, 0); pti.isValid(); ++pti) { @@ -205,10 +205,10 @@ RigidInjectedParticleContainer::BoostandRemapParticles() void RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, - Cuda::DeviceVector& xp, - Cuda::DeviceVector& yp, - Cuda::DeviceVector& zp, - Cuda::DeviceVector& giv, + Cuda::ManagedDeviceVector& xp, + Cuda::ManagedDeviceVector& yp, + Cuda::ManagedDeviceVector& zp, + Cuda::ManagedDeviceVector& giv, Real dt) { @@ -241,7 +241,7 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, #endif // Save the position and momenta, making copies - Cuda::DeviceVector xp_save, yp_save, zp_save; + Cuda::ManagedDeviceVector xp_save, yp_save, zp_save; RealVector uxp_save, uyp_save, uzp_save; if (!done_injecting_lev) { @@ -362,7 +362,7 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, #pragma omp parallel #endif { - Cuda::DeviceVector xp, yp, zp, giv; + Cuda::ManagedDeviceVector xp, yp, zp, giv; for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { -- cgit v1.2.3 From 7e58007c9eba2b51003ff813a8bfe65ae25e8392 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Mon, 29 Apr 2019 10:41:06 -0700 Subject: replace the compile time checks used for the old particle attributes with runtime ones --- Source/Particles/PhysicalParticleContainer.cpp | 71 ++++++++++++---------- .../Particles/RigidInjectedParticleContainer.cpp | 28 ++++----- 2 files changed, 52 insertions(+), 47 deletions(-) (limited to 'Source/Particles/RigidInjectedParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 74052f842..ee219e355 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -455,16 +455,18 @@ PhysicalParticleContainer::AddPlasmaCPU (int lev, RealBox part_realbox) attribs[PIdx::ux] = u[0]; attribs[PIdx::uy] = u[1]; attribs[PIdx::uz] = u[2]; - -#ifdef WARPX_STORE_OLD_PARTICLE_ATTRIBS - attribs[PIdx::xold] = x; - attribs[PIdx::yold] = y; - attribs[PIdx::zold] = z; - - attribs[PIdx::uxold] = u[0]; - attribs[PIdx::uyold] = u[1]; - attribs[PIdx::uzold] = u[2]; -#endif + + if (WarpX::do_boosted_frame_diagnostic && WarpX::do_boosted_frame_particles) + { + auto& particle_tile = DefineAndReturnParticleTile(lev, grid_id, tile_id); + particle_tile.push_back_real(particle_comps["xold"], x); + particle_tile.push_back_real(particle_comps["yold"], y); + particle_tile.push_back_real(particle_comps["zold"], z); + + particle_tile.push_back_real(particle_comps["uxold"], u[0]); + particle_tile.push_back_real(particle_comps["uyold"], u[1]); + particle_tile.push_back_real(particle_comps["uzold"], u[2]); + } AddOneParticle(lev, grid_id, tile_id, x, y, z, attribs); } @@ -695,15 +697,18 @@ PhysicalParticleContainer::AddPlasmaGPU (int lev, RealBox part_realbox) attribs[PIdx::uy] = u[1]; attribs[PIdx::uz] = u[2]; -#ifdef WARPX_STORE_OLD_PARTICLE_ATTRIBS - attribs[PIdx::xold] = x; - attribs[PIdx::yold] = y; - attribs[PIdx::zold] = z; - - attribs[PIdx::uxold] = u[0]; - attribs[PIdx::uyold] = u[1]; - attribs[PIdx::uzold] = u[2]; -#endif + // note - this will be slow on the GPU, need to revisit + if (WarpX::do_boosted_frame_diagnostic && WarpX::do_boosted_frame_particles) + { + auto& particle_tile = DefineAndReturnParticleTile(lev, grid_id, tile_id); + particle_tile.push_back_real(particle_comps["xold"], x); + particle_tile.push_back_real(particle_comps["yold"], y); + particle_tile.push_back_real(particle_comps["zold"], z); + + particle_tile.push_back_real(particle_comps["uxold"], u[0]); + particle_tile.push_back_real(particle_comps["uyold"], u[1]); + particle_tile.push_back_real(particle_comps["uzold"], u[2]); + } ParticleType p; p.id() = ParticleType::NextID(); @@ -1666,20 +1671,20 @@ PhysicalParticleContainer::PushPX(WarpXParIter& pti, auto& Bzp = attribs[PIdx::Bz]; const long np = pti.numParticles(); -#ifdef WARPX_STORE_OLD_PARTICLE_ATTRIBS - auto& xpold = attribs[PIdx::xold]; - auto& ypold = attribs[PIdx::yold]; - auto& zpold = attribs[PIdx::zold]; - auto& uxpold = attribs[PIdx::uxold]; - auto& uypold = attribs[PIdx::uyold]; - auto& uzpold = attribs[PIdx::uzold]; - - warpx_copy_attribs(&np, xp.dataPtr(), yp.dataPtr(), zp.dataPtr(), - uxp.dataPtr(), uyp.dataPtr(), uzp.dataPtr(), - xpold.dataPtr(), ypold.dataPtr(), zpold.dataPtr(), - uxpold.dataPtr(), uypold.dataPtr(), uzpold.dataPtr()); - -#endif + if (WarpX::do_boosted_frame_diagnostic && WarpX::do_boosted_frame_particles) + { + auto& xpold = pti.GetAttribs(particle_comps["xold"]); + auto& ypold = pti.GetAttribs(particle_comps["yold"]); + auto& zpold = pti.GetAttribs(particle_comps["zold"]); + auto& uxpold = pti.GetAttribs(particle_comps["uxold"]); + auto& uypold = pti.GetAttribs(particle_comps["uyold"]); + auto& uzpold = pti.GetAttribs(particle_comps["uzold"]); + + warpx_copy_attribs(&np, xp.dataPtr(), yp.dataPtr(), zp.dataPtr(), + uxp.dataPtr(), uyp.dataPtr(), uzp.dataPtr(), + xpold.dataPtr(), ypold.dataPtr(), zpold.dataPtr(), + uxpold.dataPtr(), uypold.dataPtr(), uzpold.dataPtr()); + } warpx_particle_pusher(&np, xp.dataPtr(), diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp index 3ee4d87e5..a5acca281 100644 --- a/Source/Particles/RigidInjectedParticleContainer.cpp +++ b/Source/Particles/RigidInjectedParticleContainer.cpp @@ -225,20 +225,20 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, auto& Bzp = attribs[PIdx::Bz]; const long np = pti.numParticles(); -#ifdef WARPX_STORE_OLD_PARTICLE_ATTRIBS - auto& xpold = attribs[PIdx::xold]; - auto& ypold = attribs[PIdx::yold]; - auto& zpold = attribs[PIdx::zold]; - auto& uxpold = attribs[PIdx::uxold]; - auto& uypold = attribs[PIdx::uyold]; - auto& uzpold = attribs[PIdx::uzold]; - - warpx_copy_attribs(&np, xp.dataPtr(), yp.dataPtr(), zp.dataPtr(), - uxp.dataPtr(), uyp.dataPtr(), uzp.dataPtr(), - xpold.dataPtr(), ypold.dataPtr(), zpold.dataPtr(), - uxpold.dataPtr(), uypold.dataPtr(), uzpold.dataPtr()); - -#endif + if (WarpX::do_boosted_frame_diagnostic && WarpX::do_boosted_frame_particles) + { + auto& xpold = pti.GetAttribs(particle_comps["xold"]); + auto& ypold = pti.GetAttribs(particle_comps["yold"]); + auto& zpold = pti.GetAttribs(particle_comps["zold"]); + auto& uxpold = pti.GetAttribs(particle_comps["uxold"]); + auto& uypold = pti.GetAttribs(particle_comps["uyold"]); + auto& uzpold = pti.GetAttribs(particle_comps["uzold"]); + + warpx_copy_attribs(&np, xp.dataPtr(), yp.dataPtr(), zp.dataPtr(), + uxp.dataPtr(), uyp.dataPtr(), uzp.dataPtr(), + xpold.dataPtr(), ypold.dataPtr(), zpold.dataPtr(), + uxpold.dataPtr(), uypold.dataPtr(), uzpold.dataPtr()); + } // Save the position and momenta, making copies Cuda::ManagedDeviceVector xp_save, yp_save, zp_save; -- cgit v1.2.3