From f6c64d45bea78b478be40fb73c810247af38c48b Mon Sep 17 00:00:00 2001 From: Revathi Jambunathan <41089244+RevathiJambunathan@users.noreply.github.com> Date: Fri, 19 Jun 2020 14:42:10 -0700 Subject: 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 --- Source/Particles/PhysicalParticleContainer.cpp | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') 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 FlagForPartCopy; + amrex::Gpu::ManagedDeviceVector 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 FlagForPartCopy(np); - amrex::Gpu::ManagedDeviceVector IndexForPartCopy(np); + FlagForPartCopy.resize(np); + IndexForPartCopy.resize(np); int* const AMREX_RESTRICT Flag = FlagForPartCopy.dataPtr(); int* const AMREX_RESTRICT IndexLocation = IndexForPartCopy.dataPtr(); -- cgit v1.2.3