aboutsummaryrefslogtreecommitdiff
path: root/Source/Particles/PhysicalParticleContainer.cpp
diff options
context:
space:
mode:
authorGravatar Revathi Jambunathan <41089244+RevathiJambunathan@users.noreply.github.com> 2020-06-19 14:42:10 -0700
committerGravatar GitHub <noreply@github.com> 2020-06-19 14:42:10 -0700
commitf6c64d45bea78b478be40fb73c810247af38c48b (patch)
tree998eb310f33663c256efc868ceb011322dd6a6c5 /Source/Particles/PhysicalParticleContainer.cpp
parent07954e866040442deaa8876b88d934b6af57a6d6 (diff)
downloadWarpX-f6c64d45bea78b478be40fb73c810247af38c48b.tar.gz
WarpX-f6c64d45bea78b478be40fb73c810247af38c48b.tar.zst
WarpX-f6c64d45bea78b478be40fb73c810247af38c48b.zip
Fix Race condition in GetParticleSlice for BTD (#1105)
* define temp array outside the ParIter loop for GetParticleSlice * Moving definition of temp arrays closer to the ParIter loop where used with comments
Diffstat (limited to '')
-rw-r--r--Source/Particles/PhysicalParticleContainer.cpp14
1 files changed, 10 insertions, 4 deletions
diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp
index e138f39a4..e0a46f43a 100644
--- a/Source/Particles/PhysicalParticleContainer.cpp
+++ b/Source/Particles/PhysicalParticleContainer.cpp
@@ -1734,6 +1734,14 @@ PhysicalParticleContainer::GetParticleSlice (
#pragma omp parallel
#endif
{
+ // Temporary arrays to store copy_flag and copy_index
+ // for particles that cross the z-slice
+ // These arrays are defined before the WarpXParIter to prevent them
+ // from going out of scope after each iteration, while the kernels
+ // may still need access to them.
+ // Note that the destructor for WarpXParIter is synchronized.
+ amrex::Gpu::ManagedDeviceVector<int> FlagForPartCopy;
+ amrex::Gpu::ManagedDeviceVector<int> IndexForPartCopy;
for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti)
{
const Box& box = pti.validbox();
@@ -1768,10 +1776,8 @@ PhysicalParticleContainer::GetParticleSlice (
Real uzfrm = -WarpX::gamma_boost*WarpX::beta_boost*PhysConst::c;
Real inv_c2 = 1.0/PhysConst::c/PhysConst::c;
- // temporary arrays to store copy_flag and copy_index
- // for particles that cross the z-slice
- amrex::Gpu::ManagedDeviceVector<int> FlagForPartCopy(np);
- amrex::Gpu::ManagedDeviceVector<int> IndexForPartCopy(np);
+ FlagForPartCopy.resize(np);
+ IndexForPartCopy.resize(np);
int* const AMREX_RESTRICT Flag = FlagForPartCopy.dataPtr();
int* const AMREX_RESTRICT IndexLocation = IndexForPartCopy.dataPtr();