3 files changed, 17 insertions, 3 deletions
diff --git a/Source/Particles/Collision/CollisionHandler.cpp b/Source/Particles/Collision/CollisionHandler.cpp
index 588e8d008..c4ba94a4d 100644
--- a/Source/Particles/Collision/CollisionHandler.cpp
+++ b/Source/Particles/Collision/CollisionHandler.cpp
@@ -46,7 +46,7 @@ void CollisionHandler::doCollisions ( amrex::Real cur_time, MultiParticleContain
 {
 
     for (auto& collision : allcollisions) {
-        collision->doCollisions( cur_time, mypc);
+        collision->doCollisions(cur_time, mypc);
     }
 
 }
diff --git a/Source/Particles/Collision/PairWiseCoulombCollision.H b/Source/Particles/Collision/PairWiseCoulombCollision.H
index eea91d7f9..15ae329f6 100644
--- a/Source/Particles/Collision/PairWiseCoulombCollision.H
+++ b/Source/Particles/Collision/PairWiseCoulombCollision.H
@@ -20,6 +20,7 @@ public:
 
     /** Perform the collisions
      *
+     * @param lev AMR level of the tile
      * @param cur_time Current time
      * @param mypc Container of species involved
      *
@@ -28,7 +29,6 @@ public:
 
     /** Perform all binary collisions within a tile
      *
-     * @param lev AMR level of the tile
      * @param mfi iterator for multifab
      * @param species1/2 pointer to species container
      *
diff --git a/Source/Particles/Collision/PairWiseCoulombCollision.cpp b/Source/Particles/Collision/PairWiseCoulombCollision.cpp
index 0bdaf7cc8..efa1c04f9 100644
--- a/Source/Particles/Collision/PairWiseCoulombCollision.cpp
+++ b/Source/Particles/Collision/PairWiseCoulombCollision.cpp
@@ -36,7 +36,6 @@ PairWiseCoulombCollision::PairWiseCoulombCollision (std::string const collision_
 void
 PairWiseCoulombCollision::doCollisions (amrex::Real cur_time, MultiParticleContainer* mypc)
 {
-
     const amrex::Real dt = WarpX::GetInstance().getdt(0);
     if ( int(std::floor(cur_time/dt)) % m_ndt != 0 ) return;
 
@@ -50,13 +49,28 @@ PairWiseCoulombCollision::doCollisions (amrex::Real cur_time, MultiParticleConta
     // Loop over refinement levels
     for (int lev = 0; lev <= species1.finestLevel(); ++lev){
 
+    amrex::LayoutData<amrex::Real>* cost = WarpX::getCosts(lev);
+
         // Loop over all grids/tiles at this level
 #ifdef AMREX_USE_OMP
         info.SetDynamic(true);
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
         for (amrex::MFIter mfi = species1.MakeMFIter(lev, info); mfi.isValid(); ++mfi){
+            if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers)
+            {
+                amrex::Gpu::synchronize();
+            }
+            amrex::Real wt = amrex::second();
+
             doCoulombCollisionsWithinTile( lev, mfi, species1, species2 );
+
+            if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers)
+            {
+                amrex::Gpu::synchronize();
+                wt = amrex::second() - wt;
+                amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt);
+            }
         }
     }
 }