diff options
author | 2020-06-19 14:42:10 -0700 | |
---|---|---|
committer | 2020-06-19 14:42:10 -0700 | |
commit | f6c64d45bea78b478be40fb73c810247af38c48b (patch) | |
tree | 998eb310f33663c256efc868ceb011322dd6a6c5 /Source/Particles/PhysicalParticleContainer.cpp | |
parent | 07954e866040442deaa8876b88d934b6af57a6d6 (diff) | |
download | WarpX-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.cpp | 14 |
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(); |