diff --git a/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp b/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp index b70af6fb4ed..65d18ed7a13 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp @@ -187,6 +187,8 @@ void FiniteDifferenceSolver::EvolveBCartesian ( } ); } + + cost_tracker.add(); } } @@ -340,6 +342,7 @@ void FiniteDifferenceSolver::EvolveBCartesianECT ( }); + cost_tracker.add(); } } #else @@ -449,6 +452,8 @@ void FiniteDifferenceSolver::EvolveBCylindrical ( } ); + + cost_tracker.add(); } } diff --git a/Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp b/Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp index 5d9a0ac2a96..71d565a38ba 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/EvolveE.cpp @@ -212,6 +212,8 @@ void FiniteDifferenceSolver::EvolveECartesian ( ); } + + cost_tracker.add(); } } @@ -426,6 +428,8 @@ void FiniteDifferenceSolver::EvolveECylindrical ( } // end of if condition for F + cost_tracker.add(); + } // end of loop over grid/tiles } diff --git a/Source/FieldSolver/FiniteDifferenceSolver/EvolveECTRho.cpp b/Source/FieldSolver/FiniteDifferenceSolver/EvolveECTRho.cpp index 02c78c029a7..471a479259e 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/EvolveECTRho.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/EvolveECTRho.cpp @@ -146,6 +146,8 @@ void FiniteDifferenceSolver::EvolveRhoCartesianECT ( #ifdef WARPX_DIM_XZ amrex::ignore_unused(Ey, Rhox, Rhoz, ly); #endif + + cost_tracker.add(); } #else amrex::ignore_unused(Efield, edge_lengths, face_areas, ECTRhofield, lev); diff --git a/Source/FieldSolver/FiniteDifferenceSolver/HybridPICSolveE.cpp b/Source/FieldSolver/FiniteDifferenceSolver/HybridPICSolveE.cpp index 09b559c661e..4ced40ad9b5 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/HybridPICSolveE.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/HybridPICSolveE.cpp @@ -232,6 +232,8 @@ void FiniteDifferenceSolver::CalculateCurrentAmpereCylindrical ( } } ); + + cost_tracker.add(); } } @@ -338,6 +340,8 @@ void FiniteDifferenceSolver::CalculateCurrentAmpereCartesian ( ); } ); + + cost_tracker.add(); } } #endif @@ -501,6 +505,8 @@ void FiniteDifferenceSolver::HybridPICSolveECylindrical ( - (jt_interp - jit_interp - Jextt(i, j, 0)) * Br_interp ); }); + + cost_tracker.add(); } // Loop through the grids, and over the tiles within each grid again @@ -670,6 +676,8 @@ void FiniteDifferenceSolver::HybridPICSolveECylindrical ( } } ); + + cost_tracker.add(); } } @@ -790,6 +798,8 @@ void FiniteDifferenceSolver::HybridPICSolveECartesian ( - (jy_interp - jiy_interp - Jexty(i, j, k)) * Bx_interp ); }); + + cost_tracker.add(); } // Loop through the grids, and over the tiles within each grid again @@ -953,6 +963,8 @@ void FiniteDifferenceSolver::HybridPICSolveECartesian ( } } ); + + cost_tracker.add(); } } #endif diff --git a/Source/FieldSolver/SpectralSolver/SpectralFieldData.cpp b/Source/FieldSolver/SpectralSolver/SpectralFieldData.cpp index 20c97f4b5d4..0dc7a980b7b 100644 --- a/Source/FieldSolver/SpectralSolver/SpectralFieldData.cpp +++ b/Source/FieldSolver/SpectralSolver/SpectralFieldData.cpp @@ -7,6 +7,7 @@ */ #include "SpectralFieldData.H" +#include "LoadBalance/LoadBalance.H" #include "Utils/WarpXAlgorithmSelection.H" #include "Utils/WarpXUtil.H" #include "WarpX.H" @@ -125,9 +126,6 @@ SpectralFieldData::SpectralFieldData( const int lev, const bool periodic_single_box): m_periodic_single_box{periodic_single_box} { - amrex::LayoutData* cost = WarpX::getCosts(lev); - const bool do_costs = WarpXUtilLoadBalance::doCosts(cost, realspace_ba, dm); - const BoxArray& spectralspace_ba = k_space.spectralspace_ba; // Allocate the arrays that contain the fields in spectral space @@ -168,11 +166,8 @@ SpectralFieldData::SpectralFieldData( const int lev, // Loop over boxes and allocate the corresponding plan // for each box owned by the local MPI proc for ( MFIter mfi(spectralspace_ba, dm); mfi.isValid(); ++mfi ){ - if (do_costs) - { - amrex::Gpu::synchronize(); - } - auto wt = static_cast(amrex::second()); + + const auto cost_tracker = warpx::LoadBalance::CostTracker(lev, mfi); // Note: the size of the real-space box and spectral-space box // differ when using real-to-complex FFT. When initializing @@ -189,12 +184,7 @@ SpectralFieldData::SpectralFieldData( const int lev, reinterpret_cast( tmpSpectralField[mfi].dataPtr()), ablastr::math::anyfft::direction::C2R, AMREX_SPACEDIM); - if (do_costs) - { - amrex::Gpu::synchronize(); - wt = static_cast(amrex::second()) - wt; - amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); - } + cost_tracker.add(); } } @@ -217,9 +207,6 @@ SpectralFieldData::ForwardTransform (const int lev, const MultiFab& mf, const int field_index, const int i_comp) { - amrex::LayoutData* cost = WarpX::getCosts(lev); - const bool do_costs = WarpXUtilLoadBalance::doCosts(cost, mf.boxArray(), mf.DistributionMap()); - // Check field index type, in order to apply proper shift in spectral space #if (AMREX_SPACEDIM >= 2) const bool is_nodal_x = mf.is_nodal(0); @@ -237,11 +224,8 @@ SpectralFieldData::ForwardTransform (const int lev, // Note: we do NOT OpenMP parallelize here, since we use OpenMP threads for // the FFTs on each box! for ( MFIter mfi(mf); mfi.isValid(); ++mfi ){ - if (do_costs) - { - amrex::Gpu::synchronize(); - } - auto wt = static_cast(amrex::second()); + + const auto cost_tracker = warpx::LoadBalance::CostTracker(lev, mfi); // Copy the real-space field `mf` to the temporary field `tmpRealField` // This ensures that all fields have the same number of points @@ -305,12 +289,7 @@ SpectralFieldData::ForwardTransform (const int lev, }); } - if (do_costs) - { - amrex::Gpu::synchronize(); - wt = static_cast(amrex::second()) - wt; - amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); - } + cost_tracker.add(); } } @@ -324,9 +303,6 @@ SpectralFieldData::BackwardTransform (const int lev, const amrex::IntVect& fill_guards, const int i_comp) { - amrex::LayoutData* cost = WarpX::getCosts(lev); - const bool do_costs = WarpXUtilLoadBalance::doCosts(cost, mf.boxArray(), mf.DistributionMap()); - // Check field index type, in order to apply proper shift in spectral space #if (AMREX_SPACEDIM >= 2) const bool is_nodal_x = mf.is_nodal(0); @@ -362,11 +338,8 @@ SpectralFieldData::BackwardTransform (const int lev, // Note: we do NOT OpenMP parallelize here, since we use OpenMP threads for // the iFFTs on each box! for ( MFIter mfi(mf); mfi.isValid(); ++mfi ){ - if (do_costs) - { - amrex::Gpu::synchronize(); - } - auto wt = static_cast(amrex::second()); + + const auto cost_tracker = warpx::LoadBalance::CostTracker(lev, mfi); // Copy the spectral-space field `tmpSpectralField` to the appropriate // field (specified by the input argument field_index) @@ -466,12 +439,8 @@ SpectralFieldData::BackwardTransform (const int lev, }); } - if (do_costs) - { - amrex::Gpu::synchronize(); - wt = static_cast(amrex::second()) - wt; - amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); - } + cost_tracker.add(); + } } diff --git a/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp b/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp index d295d61b43a..7021d597143 100644 --- a/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp +++ b/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp @@ -6,6 +6,7 @@ */ #include "SpectralFieldDataRZ.H" +#include "LoadBalance/LoadBalance.H" #include "Utils/WarpXUtil.H" #include "WarpX.H" @@ -452,9 +453,6 @@ SpectralFieldDataRZ::ForwardTransform (const int lev, amrex::MultiFab const & field_mf, int const field_index, int const i_comp) { - amrex::LayoutData* cost = WarpX::getCosts(lev); - const bool do_costs = WarpXUtilLoadBalance::doCosts(cost, field_mf.boxArray(), field_mf.DistributionMap()); - // Check field index type, in order to apply proper shift in spectral space. // Only cell centered in r is supported. bool const is_nodal_z = field_mf.is_nodal(1); @@ -474,11 +472,7 @@ SpectralFieldDataRZ::ForwardTransform (const int lev, // Loop over boxes. for (amrex::MFIter mfi(field_mf); mfi.isValid(); ++mfi){ - if (do_costs) - { - amrex::Gpu::synchronize(); - } - auto wt = static_cast(amrex::second()); + const auto cost_tracker = warpx::LoadBalance::CostTracker(lev, mfi); // Perform the Hankel transform first. // tempHTransformedSplit includes the imaginary component of mode 0. @@ -495,12 +489,7 @@ SpectralFieldDataRZ::ForwardTransform (const int lev, FABZForwardTransform(mfi, realspace_bx, tempHTransformedSplit, field_index, is_nodal_z); - if (do_costs) - { - amrex::Gpu::synchronize(); - wt = static_cast(amrex::second()) - wt; - amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); - } + cost_tracker.add(); } } @@ -512,8 +501,6 @@ SpectralFieldDataRZ::ForwardTransform (const int lev, amrex::MultiFab const & field_mf_r, int const field_index_r, amrex::MultiFab const & field_mf_t, int const field_index_t) { - amrex::LayoutData* cost = WarpX::getCosts(lev); - const bool do_costs = WarpXUtilLoadBalance::doCosts(cost, field_mf_r.boxArray(), field_mf_r.DistributionMap()); // Check field index type, in order to apply proper shift in spectral space. // Only cell centered in r is supported. @@ -531,11 +518,7 @@ SpectralFieldDataRZ::ForwardTransform (const int lev, // Loop over boxes. for (amrex::MFIter mfi(field_mf_r); mfi.isValid(); ++mfi){ - if (do_costs) - { - amrex::Gpu::synchronize(); - } - auto wt = static_cast(amrex::second()); + const auto cost_tracker = warpx::LoadBalance::CostTracker(lev, mfi); amrex::Box const& realspace_bx = tempHTransformed[mfi].box(); @@ -561,12 +544,7 @@ SpectralFieldDataRZ::ForwardTransform (const int lev, FABZForwardTransform(mfi, realspace_bx, tempHTransformedSplit_p, field_index_r, is_nodal_z); FABZForwardTransform(mfi, realspace_bx, tempHTransformedSplit_m, field_index_t, is_nodal_z); - if (do_costs) - { - amrex::Gpu::synchronize(); - wt = static_cast(amrex::second()) - wt; - amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); - } + cost_tracker.add(); } } @@ -577,9 +555,6 @@ SpectralFieldDataRZ::BackwardTransform (const int lev, amrex::MultiFab& field_mf, int const field_index, int const i_comp) { - amrex::LayoutData* cost = WarpX::getCosts(lev); - const bool do_costs = WarpXUtilLoadBalance::doCosts(cost, field_mf.boxArray(), field_mf.DistributionMap()); - // Check field index type, in order to apply proper shift in spectral space. const bool is_nodal_z = field_mf.is_nodal(1); @@ -593,11 +568,7 @@ SpectralFieldDataRZ::BackwardTransform (const int lev, // Loop over boxes. for (amrex::MFIter mfi(field_mf); mfi.isValid(); ++mfi){ - if (do_costs) - { - amrex::Gpu::synchronize(); - } - auto wt = static_cast(amrex::second()); + const auto cost_tracker = warpx::LoadBalance::CostTracker(lev, mfi); amrex::Box realspace_bx = tempHTransformed[mfi].box(); @@ -645,12 +616,8 @@ SpectralFieldDataRZ::BackwardTransform (const int lev, field_mf_array(i,j,k,ic) = sign*field_mf_copy_array(ii,j,k,icomp); }); - if (do_costs) - { - amrex::Gpu::synchronize(); - wt = static_cast(amrex::second()) - wt; - amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); - } + cost_tracker.add(); + } } @@ -661,8 +628,6 @@ SpectralFieldDataRZ::BackwardTransform (const int lev, amrex::MultiFab& field_mf_r, int const field_index_r, amrex::MultiFab& field_mf_t, int const field_index_t) { - amrex::LayoutData* cost = WarpX::getCosts(lev); - const bool do_costs = WarpXUtilLoadBalance::doCosts(cost, field_mf_r.boxArray(), field_mf_r.DistributionMap()); // Check field index type, in order to apply proper shift in spectral space. bool const is_nodal_z = field_mf_r.is_nodal(1); @@ -678,11 +643,7 @@ SpectralFieldDataRZ::BackwardTransform (const int lev, // Loop over boxes. for (amrex::MFIter mfi(field_mf_r); mfi.isValid(); ++mfi){ - if (do_costs) - { - amrex::Gpu::synchronize(); - } - auto wt = static_cast(amrex::second()); + const auto cost_tracker = warpx::LoadBalance::CostTracker(lev, mfi); amrex::Box realspace_bx = tempHTransformed[mfi].box(); @@ -741,12 +702,7 @@ SpectralFieldDataRZ::BackwardTransform (const int lev, } }); - if (do_costs) - { - amrex::Gpu::synchronize(); - wt = static_cast(amrex::second()) - wt; - amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); - } + cost_tracker.add(); } } @@ -772,16 +728,9 @@ SpectralFieldDataRZ::InitFilter (amrex::IntVect const & filter_npass_each_dir, b void SpectralFieldDataRZ::ApplyFilter (const int lev, int const field_index) { - amrex::LayoutData* cost = WarpX::getCosts(lev); - const bool do_costs = WarpXUtilLoadBalance::doCosts(cost, binomialfilter.boxArray(), binomialfilter.DistributionMap()); - for (amrex::MFIter mfi(binomialfilter); mfi.isValid(); ++mfi){ - if (do_costs) - { - amrex::Gpu::synchronize(); - } - auto wt = static_cast(amrex::second()); + const auto cost_tracker = warpx::LoadBalance::CostTracker(lev, mfi); auto const & filter_r = binomialfilter[mfi].getFilterArrayR(); auto const & filter_z = binomialfilter[mfi].getFilterArrayZ(); @@ -803,12 +752,8 @@ SpectralFieldDataRZ::ApplyFilter (const int lev, int const field_index) fields_arr(i,j,k,ic) *= filter_r_arr[ir]*filter_z_arr[j]; }); - if (do_costs) - { - amrex::Gpu::synchronize(); - wt = static_cast(amrex::second()) - wt; - amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt); - } + cost_tracker.add(); + } } diff --git a/Source/FieldSolver/WarpX_QED_Field_Pushers.cpp b/Source/FieldSolver/WarpX_QED_Field_Pushers.cpp index f09ae9ba762..aaee712f3f3 100644 --- a/Source/FieldSolver/WarpX_QED_Field_Pushers.cpp +++ b/Source/FieldSolver/WarpX_QED_Field_Pushers.cpp @@ -176,5 +176,7 @@ WarpX::Hybrid_QED_Push (int lev, PatchType patch_type, amrex::Real a_dt) a_dt, xi_c2); } ); + + cost_tracker.add(); } } diff --git a/Source/Filter/Filter.cpp b/Source/Filter/Filter.cpp index 96aa190bf6e..17d09443b53 100644 --- a/Source/Filter/Filter.cpp +++ b/Source/Filter/Filter.cpp @@ -51,6 +51,8 @@ Filter::ApplyStencil (MultiFab& dstmf, const MultiFab& srcmf, const int lev, int // Apply filter DoFilter(tbx, src, dst, scomp, dcomp, ncomp); + + cost_tracker.add(); } } @@ -215,6 +217,8 @@ Filter::ApplyStencil (amrex::MultiFab& dstmf, const amrex::MultiFab& srcmf, cons tmpfab.copy(srcfab, ibx, scomp, ibx, 0, ncomp); // Apply filter DoFilter(tbx, tmpfab.array(), dstfab.array(), 0, dcomp, ncomp); + + cost_tracker.add(); } } } diff --git a/Source/LoadBalance/LoadBalance.H b/Source/LoadBalance/LoadBalance.H index 1a37995542d..cb6abe11139 100644 --- a/Source/LoadBalance/LoadBalance.H +++ b/Source/LoadBalance/LoadBalance.H @@ -30,7 +30,8 @@ namespace warpx::load_balance { public: CostTracker (int lev, std::size_t mfi_iter_index); - ~CostTracker (); + + void add () const noexcept; private: int m_lev; diff --git a/Source/LoadBalance/LoadBalance.cpp b/Source/LoadBalance/LoadBalance.cpp index 7de72920e56..4be58f7a0df 100644 --- a/Source/LoadBalance/LoadBalance.cpp +++ b/Source/LoadBalance/LoadBalance.cpp @@ -48,7 +48,7 @@ namespace } } -CostTracker::CostTracker(int lev, std::size_t mfi_iter_index): +CostTracker::CostTracker (int lev, std::size_t mfi_iter_index): m_lev{lev}, m_mfi_iter_index{mfi_iter_index}, m_wt{0.0_rt} @@ -61,15 +61,16 @@ CostTracker::CostTracker(int lev, std::size_t mfi_iter_index): m_wt = static_cast(amrex::second()); } -CostTracker::~CostTracker() +void +CostTracker::add () const noexcept { if (AllCosts::get_instance().get_update_algo() == CostsUpdateAlgo::Heuristic) return; amrex::Gpu::synchronize(); - m_wt = static_cast(amrex::second()) - m_wt; + const auto time_diff = static_cast(amrex::second()) - m_wt; auto& cost = AllCosts::get_instance().m_costs[m_lev]; - amrex::HostDevice::Atomic::Add( &(*cost)[m_mfi_iter_index], m_wt); + amrex::HostDevice::Atomic::Add( &(*cost)[m_mfi_iter_index], time_diff); } AllCosts::AllCosts () diff --git a/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.H b/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.H index 3fd8b4d6fc2..d4fea4e8f57 100644 --- a/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.H +++ b/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.H @@ -54,7 +54,6 @@ public: /** Perform MCC ionization interactions * * @param[in] lev the mesh-refinement level - * @param[in,out] cost pointer to (load balancing) cost corresponding to box where present particles are ionized. * @param[in,out] species1,species2 reference to species container used to inject * new particles * @param t current time @@ -62,7 +61,6 @@ public: */ void doBackgroundIonization ( int lev, - amrex::LayoutData* cost, WarpXParticleContainer& species1, WarpXParticleContainer& species2, amrex::Real t diff --git a/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp b/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp index 32324344203..30b682e8a19 100644 --- a/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp +++ b/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp @@ -283,8 +283,6 @@ BackgroundMCCCollision::doCollisions (amrex::Real cur_time, amrex::Real dt, Mult auto const flvl = species1.finestLevel(); for (int lev = 0; lev <= flvl; ++lev) { - auto *cost = WarpX::getCosts(lev); - // firstly loop over particles box by box and do all particle conserving // scattering #ifdef _OPENMP @@ -293,11 +291,12 @@ BackgroundMCCCollision::doCollisions (amrex::Real cur_time, amrex::Real dt, Mult for (WarpXParIter pti(species1, lev); pti.isValid(); ++pti) { const auto cost_tracker = warpx::load_balance::CostTracker(lev, pti.index()); doBackgroundCollisionsWithinTile(pti, cur_time); + cost_tracker.add(); } // secondly perform ionization through the SmartCopyFactory if needed if (ionization_flag) { - doBackgroundIonization(lev, cost, species1, species2, cur_time); + doBackgroundIonization(lev, species1, species2, cur_time); } } } @@ -456,8 +455,7 @@ void BackgroundMCCCollision::doBackgroundCollisionsWithinTile void BackgroundMCCCollision::doBackgroundIonization -( int lev, amrex::LayoutData* cost, - WarpXParticleContainer& species1, WarpXParticleContainer& species2, amrex::Real t) +( int lev, WarpXParticleContainer& species1, WarpXParticleContainer& species2, amrex::Real t) { WARPX_PROFILE("BackgroundMCCCollision::doBackgroundIonization()"); @@ -499,5 +497,7 @@ void BackgroundMCCCollision::doBackgroundIonization setNewParticleIDs(elec_tile, np_elec, num_added); setNewParticleIDs(ion_tile, np_ion, num_added); + + cost_tracker.add(); } } diff --git a/Source/Particles/Collision/BackgroundStopping/BackgroundStopping.cpp b/Source/Particles/Collision/BackgroundStopping/BackgroundStopping.cpp index bf2cfceb1c3..577a494a3b6 100644 --- a/Source/Particles/Collision/BackgroundStopping/BackgroundStopping.cpp +++ b/Source/Particles/Collision/BackgroundStopping/BackgroundStopping.cpp @@ -105,8 +105,6 @@ BackgroundStopping::doCollisions (amrex::Real cur_time, amrex::Real dt, MultiPar auto const flvl = species.finestLevel(); for (int lev = 0; lev <= flvl; ++lev) { - auto *cost = WarpX::getCosts(lev); - // loop over particles box by box #ifdef _OPENMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) @@ -119,6 +117,8 @@ BackgroundStopping::doCollisions (amrex::Real cur_time, amrex::Real dt, MultiPar } else if (background_type == BackgroundStoppingType::IONS) { doBackgroundStoppingOnIonsWithinTile(pti, dt, cur_time, species_mass, species_charge); } + + cost_tracker.add(); } } diff --git a/Source/Particles/Collision/BinaryCollision/BinaryCollision.H b/Source/Particles/Collision/BinaryCollision/BinaryCollision.H index 0c9287216df..afcc54c7cce 100644 --- a/Source/Particles/Collision/BinaryCollision/BinaryCollision.H +++ b/Source/Particles/Collision/BinaryCollision/BinaryCollision.H @@ -195,6 +195,8 @@ public: doCollisionsWithinTile( dt, lev, mfi, species1, species2, product_species_vector, copy_species1_data, copy_species2_data); + + cost_tracker.add(); } if (m_have_product_species) { diff --git a/Source/Particles/LaserParticleContainer.cpp b/Source/Particles/LaserParticleContainer.cpp index 55756292e9f..c8ea4714af5 100644 --- a/Source/Particles/LaserParticleContainer.cpp +++ b/Source/Particles/LaserParticleContainer.cpp @@ -685,6 +685,8 @@ LaserParticleContainer::Evolve (int lev, // This is necessary because of plane_Xp, plane_Yp and amplitude_E amrex::Gpu::synchronize(); + + cost_tracker.add(); } } } diff --git a/Source/Particles/MultiParticleContainer.cpp b/Source/Particles/MultiParticleContainer.cpp index 5fe705463de..fc705a7958d 100644 --- a/Source/Particles/MultiParticleContainer.cpp +++ b/Source/Particles/MultiParticleContainer.cpp @@ -911,6 +911,8 @@ MultiParticleContainer::doFieldIonization (int lev, Filter, Copy, Transform); setNewParticleIDs(dst_tile, np_dst, num_added); + + cost_tracker.add(); } } } @@ -1543,6 +1545,8 @@ void MultiParticleContainer::doQedBreitWheeler (int lev, setNewParticleIDs(dst_ele_tile, np_dst_ele, num_added); setNewParticleIDs(dst_pos_tile, np_dst_pos, num_added); + + cost_tracker.add(); } } } @@ -1611,6 +1615,8 @@ void MultiParticleContainer::doQedQuantumSync (int lev, cleanLowEnergyPhotons( dst_tile, np_dst, num_added, m_quantum_sync_photon_creation_energy_threshold); + + cost_tracker.add(); } } } diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index f6b4bea09a5..46bf128ec97 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1468,6 +1468,8 @@ PhysicalParticleContainer::AddPlasma (PlasmaInjector const& plasma_injector, int }); amrex::Gpu::synchronize(); + + cost_tracker.add(); } // Remove particles that are inside the embedded boundaries @@ -1514,8 +1516,6 @@ PhysicalParticleContainer::AddPlasmaFlux (PlasmaInjector const& plasma_injector, if (plasma_injector.flux_normal_axis == 2) { scale_fac /= dx[0]; } #endif - amrex::LayoutData* cost = WarpX::getCosts(0); - // Create temporary particle container to which particles will be added; // we will then call Redistribute on this new container and finally // add the new particles to the original container. @@ -1956,6 +1956,7 @@ PhysicalParticleContainer::AddPlasmaFlux (PlasmaInjector const& plasma_injector, amrex::Gpu::synchronize(); + cost_tracker.add(); } // Remove particles that are inside the embedded boundaries @@ -2226,6 +2227,8 @@ PhysicalParticleContainer::Evolve (int lev, } amrex::Gpu::synchronize(); + + cost_tracker.add(); } } // Split particles at the end of the timestep. diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 49993aadd24..3e43396d49f 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -1464,8 +1464,6 @@ WarpXParticleContainer::PushX (int lev, amrex::Real dt) if (do_not_push) { return; } - amrex::LayoutData* costs = WarpX::getCosts(lev); - #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif @@ -1497,6 +1495,8 @@ WarpXParticleContainer::PushX (int lev, amrex::Real dt) SetPosition(i, x, y, z); } ); + + cost_tracker.add(); } } } diff --git a/Source/Utils/WarpXMovingWindow.cpp b/Source/Utils/WarpXMovingWindow.cpp index a3a7dc8c1a1..7af284fa166 100644 --- a/Source/Utils/WarpXMovingWindow.cpp +++ b/Source/Utils/WarpXMovingWindow.cpp @@ -584,6 +584,9 @@ WarpX::shiftMF (amrex::MultiFab& mf, const amrex::Geometry& geom, { dstfab(i,j,k,n) = srcfab(i+shift.x,j+shift.y,k+shift.z,n); }) + + if (update_cost_flag) + cost_tracker.add(); } #if (defined WARPX_DIM_RZ) && (defined WARPX_USE_FFT) diff --git a/Source/WarpX.H b/Source/WarpX.H index f12874b22cb..3372dd869cc 100644 --- a/Source/WarpX.H +++ b/Source/WarpX.H @@ -183,6 +183,10 @@ public: static int max_particle_its_in_implicit_scheme; //! Relative tolerance used for self-consistent particle update in implicit particle-suppressed evolve schemes static amrex::ParticleReal particle_tol_in_implicit_scheme; + /** Records a number corresponding to the load balance cost update strategy + * being used (0 or 1 corresponding to timers or heuristic). + */ + static short load_balance_costs_update_algo; //! Integer that corresponds to electromagnetic Maxwell solver (vacuum - 0, macroscopic - 1) static int em_solver_medium; /** Integer that correspond to macroscopic Maxwell solver algorithm