diff options
author | 2022-02-03 07:06:03 +0100 | |
---|---|---|
committer | 2022-02-03 06:06:03 +0000 | |
commit | c3c88dfacf1e4fd4fb2148531e20034eed27b836 (patch) | |
tree | e4541c59c5fb68765afe9e7c48cbf7029cdf096e /Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp | |
parent | ec072594fb1bddb4631c55fb3018050cbf461243 (diff) | |
download | WarpX-c3c88dfacf1e4fd4fb2148531e20034eed27b836.tar.gz WarpX-c3c88dfacf1e4fd4fb2148531e20034eed27b836.tar.zst WarpX-c3c88dfacf1e4fd4fb2148531e20034eed27b836.zip |
Fixing the computation of ECT Rho Field (#2711)
* Fixed computation of ECT Rho
* [pre-commit.ci] auto fixes from pre-commit.com hooks
for more information, see https://pre-commit.ci
* Adding a missing preprocessor directive
* Updated contributors list
* Apply suggestions from code review
Co-authored-by: Axel Huebl <axel.huebl@plasma.ninja>
* Apply suggestions from code review
Co-authored-by: Edoardo Zoni <59625522+EZoni@users.noreply.github.com>
* Fixed an abort message
* Fix in Source/Initialization/WarpXInitData.cpp
Co-authored-by: Edoardo Zoni <59625522+EZoni@users.noreply.github.com>
* Added some comments
* Made EvolveRhoCartesianECT not static
Co-authored-by: Axel Huebl <axel.huebl@plasma.ninja>
Co-authored-by: lgiacome <lorenzo.giacome@cern.ch>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
Co-authored-by: Axel Huebl <axel.huebl@plasma.ninja>
Co-authored-by: Edoardo Zoni <59625522+EZoni@users.noreply.github.com>
Diffstat (limited to 'Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp')
-rw-r--r-- | Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp | 91 |
1 files changed, 0 insertions, 91 deletions
diff --git a/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp b/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp index 9de1e8f8b..5e0a944b2 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/EvolveB.cpp @@ -197,97 +197,6 @@ void FiniteDifferenceSolver::EvolveBCartesian ( } } -void FiniteDifferenceSolver::EvolveRhoCartesianECT ( - std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& Efield, - std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& edge_lengths, - std::array< std::unique_ptr<amrex::MultiFab>, 3 > const& face_areas, - std::array< std::unique_ptr<amrex::MultiFab>, 3 >& ECTRhofield, const int lev ) { -#ifdef AMREX_USE_EB - -#if !(defined(WARPX_DIM_3D) || defined(WARPX_DIM_XZ)) - amrex::Abort("EvolveRhoCartesianECT: Embedded Boundaries are only implemented in 2D3V and 3D3V"); -#endif - - amrex::LayoutData<amrex::Real>* cost = WarpX::getCosts(lev); - - // Loop through the grids, and over the tiles within each grid -#ifdef AMREX_USE_OMP -#pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) -#endif - for (MFIter mfi(*ECTRhofield[0], TilingIfNotGPU()); mfi.isValid(); ++mfi ) { - if (cost && WarpX::load_balance_costs_update_algo == LoadBalanceCostsUpdateAlgo::Timers) { - amrex::Gpu::synchronize(); - } - Real wt = amrex::second(); - - // Extract field data for this grid/tile - Array4<Real> const &Ex = Efield[0]->array(mfi); - Array4<Real> const &Ey = Efield[1]->array(mfi); - Array4<Real> const &Ez = Efield[2]->array(mfi); - Array4<Real> const &Rhox = ECTRhofield[0]->array(mfi); - Array4<Real> const &Rhoy = ECTRhofield[1]->array(mfi); - Array4<Real> const &Rhoz = ECTRhofield[2]->array(mfi); - amrex::Array4<amrex::Real> const &lx = edge_lengths[0]->array(mfi); - amrex::Array4<amrex::Real> const &ly = edge_lengths[1]->array(mfi); - amrex::Array4<amrex::Real> const &lz = edge_lengths[2]->array(mfi); - amrex::Array4<amrex::Real> const &Sx = face_areas[0]->array(mfi); - amrex::Array4<amrex::Real> const &Sy = face_areas[1]->array(mfi); - amrex::Array4<amrex::Real> const &Sz = face_areas[2]->array(mfi); - - // Extract tileboxes for which to loop - Box const &trhox = mfi.tilebox(ECTRhofield[0]->ixType().toIntVect()); - Box const &trhoy = mfi.tilebox(ECTRhofield[1]->ixType().toIntVect()); - Box const &trhoz = mfi.tilebox(ECTRhofield[2]->ixType().toIntVect()); - - amrex::ParallelFor(trhox, trhoy, trhoz, - - [=] AMREX_GPU_DEVICE(int i, int j, int k) { - if (Sx(i, j, k) <= 0) return; - -#ifndef WARPX_DIM_XZ - Rhox(i, j, k) = (Ey(i, j, k) * ly(i, j, k) - Ey(i, j, k + 1) * ly(i, j, k + 1) + - Ez(i, j + 1, k) * lz(i, j + 1, k) - Ez(i, j, k) * lz(i, j, k)) / Sx(i, j, k); -#endif - }, - - [=] AMREX_GPU_DEVICE(int i, int j, int k) { - if (Sy(i, j, k) <= 0) return; - -#ifdef WARPX_DIM_XZ - Rhoy(i, j, k) = (Ez(i, j, k) * lz(i, j, k) - Ez(i + 1, j, k) * lz(i + 1, j, k) + - Ex(i, j + 1, k) * lx(i, j + 1, k) - Ex(i, j, k) * lx(i, j, k)) / Sy(i, j, k); -#elif defined(WARPX_DIM_3D) - Rhoy(i, j, k) = (Ez(i, j, k) * lz(i, j, k) - Ez(i + 1, j, k) * lz(i + 1, j, k) + - Ex(i, j, k + 1) * lx(i, j, k + 1) - Ex(i, j, k) * lx(i, j, k)) / Sy(i, j, k); -#else - amrex::Abort("EvolveRhoCartesianECT: Embedded Boundaries are only implemented in 2D3V and 3D3V"); -#endif - }, - - [=] AMREX_GPU_DEVICE(int i, int j, int k) { - if (Sz(i, j, k) <= 0) return; - -#ifndef WARPX_DIM_XZ - Rhoz(i, j, k) = (Ex(i, j, k) * lx(i, j, k) - Ex(i, j + 1, k) * lx(i, j + 1, k) + - Ey(i + 1, j, k) * ly(i + 1, j, k) - Ey(i, j, k) * ly(i, j, k)) / Sz(i, j, k); -#endif - } - ); - - 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); - } -#ifdef WARPX_DIM_XZ - amrex::ignore_unused(Ey, Rhox, Rhoz, ly); -#endif - } -#else - amrex::ignore_unused(Efield, edge_lengths, face_areas, ECTRhofield, lev); -#endif -} void FiniteDifferenceSolver::EvolveBCartesianECT ( std::array< std::unique_ptr<amrex::MultiFab>, 3 >& Bfield, |