Delay some synchronize calls in addParticles (#4623)

The proposed changes:
- [ ] fix a bug or incorrect behavior in AMReX
- [ ] add new capabilities to AMReX
- [ ] changes answers in the test suite to more than roundoff level
- [ ] are likely to significantly affect the results of downstream AMReX
users
- [ ] include documentation in the code and/or rst files, if appropriate

Author: atmyers

Src/Particle/AMReX_ParticleContainerI.H

@@ -1170,22 +1170,23 @@ addParticles (const PCType& other, F const& f, bool local)
     for (int lev = 0; lev < other.numLevels(); ++lev)
     {
         const auto& plevel_other = other.GetParticles(lev);
-        for(MFIter mfi = other.MakeMFIter(lev); mfi.isValid(); ++mfi)
+        for(MFIter mfi = other.MakeMFIter(lev, MFItInfo().DisableDeviceSync()); mfi.isValid(); ++mfi)
         {
             auto index = std::make_pair(mfi.index(), mfi.LocalTileIndex());
             if(plevel_other.find(index) == plevel_other.end()) { continue; }
 
             DefineAndReturnParticleTile(lev, mfi.index(), mfi.LocalTileIndex());
         }
     }
+    Gpu::streamSynchronizeAll();
 
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (Gpu::notInLaunchRegion())
 #endif
     for (int lev = 0; lev < other.numLevels(); ++lev)
     {
         const auto& plevel_other = other.GetParticles(lev);
-        for(MFIter mfi = other.MakeMFIter(lev); mfi.isValid(); ++mfi)
+        for(MFIter mfi = other.MakeMFIter(lev, MFItInfo().DisableDeviceSync()); mfi.isValid(); ++mfi)
         {
             auto index = std::make_pair(mfi.index(), mfi.LocalTileIndex());
             if(plevel_other.find(index) == plevel_other.end()) { continue; }
@@ -1204,6 +1205,7 @@ addParticles (const PCType& other, F const& f, bool local)
             ptile.resize(dst_index + count);
         }
     }
+    Gpu::streamSynchronizeAll();
 
     if (! local) { Redistribute(); }
 }