Use C++ templates for the PML field pusher (#808)

* Allow to pass component in stencil templates

* Define and use enum to address PML components

* Start implementing PML equations

* Implement EvolveEPML

* Implemented EvolveBPML

* Added interface for pml_has_particles

* Added interface for pml_has_particles

* [skip ci] Add update expressions for E

* [skip ci] Fix compilation

* Call new PML pusher for B field

* Fix compilation errors

* Fix more typos

* Abort code if `do_pml` is used in cylindrical geometry

* Add contribution from F in EvolveEPML

* Remove unused function for CKC coefficients

* Remove unneeded ExchangeF

* Add damping for J

* Apply suggestions from code review

Co-authored-by: Axel Huebl <axel.huebl@plasma.ninja>

* Revert "Apply suggestions from code review"

This reverts commit 08f262a676ba5e5b44b9118b8daba1b03c08b64b.

* Remove sanity checks for nodal

* Implement dive cleaning in PML

* Implement push F in the PML

* Clean-up unused code

Co-authored-by: Axel Huebl <axel.huebl@plasma.ninja>

1 files changed, 131 insertions, 0 deletions

diff --git a/Source/FieldSolver/FiniteDifferenceSolver/EvolveBPML.cpp b/Source/FieldSolver/FiniteDifferenceSolver/EvolveBPML.cpp
new file mode 100644
index 000000000..6282892ab
--- /dev/null
+++ b/Source/FieldSolver/FiniteDifferenceSolver/EvolveBPML.cpp
@@ -0,0 +1,131 @@
+/* Copyright 2020 Remi Lehe
+ *
+ * This file is part of WarpX.
+ *
+ * License: BSD-3-Clause-LBNL
+ */
+
+#include "Utils/WarpXAlgorithmSelection.H"
+#include "FieldSolver/FiniteDifferenceSolver/FiniteDifferenceSolver.H"
+#ifdef WARPX_DIM_RZ
+#   include "FieldSolver/FiniteDifferenceSolver/FiniteDifferenceAlgorithms/CylindricalYeeAlgorithm.H"
+#else
+#   include "FieldSolver/FiniteDifferenceSolver/FiniteDifferenceAlgorithms/CartesianYeeAlgorithm.H"
+#   include "FieldSolver/FiniteDifferenceSolver/FiniteDifferenceAlgorithms/CartesianCKCAlgorithm.H"
+#   include "FieldSolver/FiniteDifferenceSolver/FiniteDifferenceAlgorithms/CartesianNodalAlgorithm.H"
+#endif
+#include "BoundaryConditions/PMLComponent.H"
+#include <AMReX_Gpu.H>
+
+using namespace amrex;
+
+/**
+ * \brief Update the B field, over one timestep
+ */
+void FiniteDifferenceSolver::EvolveBPML (
+    std::array< amrex::MultiFab*, 3 > Bfield,
+    std::array< amrex::MultiFab*, 3 > const Efield,
+    amrex::Real const dt ) {
+
+   // Select algorithm (The choice of algorithm is a runtime option,
+   // but we compile code for each algorithm, using templates)
+#ifdef WARPX_DIM_RZ
+    amrex::Abort("PML are not implemented in cylindrical geometry.");
+#else
+    if (m_do_nodal) {
+
+        EvolveBPMLCartesian <CartesianNodalAlgorithm> ( Bfield, Efield, dt );
+
+    } else if (m_fdtd_algo == MaxwellSolverAlgo::Yee) {
+
+        EvolveBPMLCartesian <CartesianYeeAlgorithm> ( Bfield, Efield, dt );
+
+    } else if (m_fdtd_algo == MaxwellSolverAlgo::CKC) {
+
+        EvolveBPMLCartesian <CartesianCKCAlgorithm> ( Bfield, Efield, dt );
+
+    } else {
+        amrex::Abort("Unknown algorithm");
+    }
+#endif
+}
+
+
+#ifndef WARPX_DIM_RZ
+
+template<typename T_Algo>
+void FiniteDifferenceSolver::EvolveBPMLCartesian (
+    std::array< amrex::MultiFab*, 3 > Bfield,
+    std::array< amrex::MultiFab*, 3 > const Efield,
+    amrex::Real const dt ) {
+
+    // Loop through the grids, and over the tiles within each grid
+#ifdef _OPENMP
+#pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
+#endif
+    for ( MFIter mfi(*Bfield[0], TilingIfNotGPU()); mfi.isValid(); ++mfi ) {
+
+        // Extract field data for this grid/tile
+        Array4<Real> const& Bx = Bfield[0]->array(mfi);
+        Array4<Real> const& By = Bfield[1]->array(mfi);
+        Array4<Real> const& Bz = Bfield[2]->array(mfi);
+        Array4<Real> const& Ex = Efield[0]->array(mfi);
+        Array4<Real> const& Ey = Efield[1]->array(mfi);
+        Array4<Real> const& Ez = Efield[2]->array(mfi);
+
+        // Extract stencil coefficients
+        Real const * const AMREX_RESTRICT coefs_x = m_stencil_coefs_x.dataPtr();
+        int const n_coefs_x = m_stencil_coefs_x.size();
+        Real const * const AMREX_RESTRICT coefs_y = m_stencil_coefs_y.dataPtr();
+        int const n_coefs_y = m_stencil_coefs_y.size();
+        Real const * const AMREX_RESTRICT coefs_z = m_stencil_coefs_z.dataPtr();
+        int const n_coefs_z = m_stencil_coefs_z.size();
+
+        // Extract tileboxes for which to loop
+        Box const& tbx  = mfi.tilebox(Bfield[0]->ixType().ixType());
+        Box const& tby  = mfi.tilebox(Bfield[1]->ixType().ixType());
+        Box const& tbz  = mfi.tilebox(Bfield[2]->ixType().ixType());
+
+        // Loop over the cells and update the fields
+        amrex::ParallelFor(tbx, tby, tbz,
+
+            [=] AMREX_GPU_DEVICE (int i, int j, int k){
+                Bx(i, j, k, PMLComp::xz) += dt * (
+                    T_Algo::UpwardDz(Ey, coefs_z, n_coefs_z, i, j, k, PMLComp::yx)
+                  + T_Algo::UpwardDz(Ey, coefs_z, n_coefs_z, i, j, k, PMLComp::yy)
+                  + T_Algo::UpwardDz(Ey, coefs_z, n_coefs_z, i, j, k, PMLComp::yz) );
+                Bx(i, j, k, PMLComp::xy) -= dt * (
+                    T_Algo::UpwardDy(Ez, coefs_y, n_coefs_y, i, j, k, PMLComp::zx)
+                  + T_Algo::UpwardDy(Ez, coefs_y, n_coefs_y, i, j, k, PMLComp::zy)
+                  + T_Algo::UpwardDy(Ez, coefs_y, n_coefs_y, i, j, k, PMLComp::zz) );
+            },
+
+            [=] AMREX_GPU_DEVICE (int i, int j, int k){
+                By(i, j, k, PMLComp::yx) += dt * (
+                    T_Algo::UpwardDx(Ez, coefs_x, n_coefs_x, i, j, k, PMLComp::zx)
+                  + T_Algo::UpwardDx(Ez, coefs_x, n_coefs_x, i, j, k, PMLComp::zy)
+                  + T_Algo::UpwardDx(Ez, coefs_x, n_coefs_x, i, j, k, PMLComp::zz) );
+                By(i, j, k, PMLComp::yz) -= dt * (
+                    T_Algo::UpwardDz(Ex, coefs_z, n_coefs_z, i, j, k, PMLComp::xx)
+                  + T_Algo::UpwardDz(Ex, coefs_z, n_coefs_z, i, j, k, PMLComp::xy)
+                  + T_Algo::UpwardDz(Ex, coefs_z, n_coefs_z, i, j, k, PMLComp::xz) );
+            },
+
+            [=] AMREX_GPU_DEVICE (int i, int j, int k){
+                Bz(i, j, k, PMLComp::zy) += dt * (
+                    T_Algo::UpwardDy(Ex, coefs_y, n_coefs_y, i, j, k, PMLComp::xx)
+                  + T_Algo::UpwardDy(Ex, coefs_y, n_coefs_y, i, j, k, PMLComp::xy)
+                  + T_Algo::UpwardDy(Ex, coefs_y, n_coefs_y, i, j, k, PMLComp::xz) );
+                Bz(i, j, k, PMLComp::zx) -= dt * (
+                    T_Algo::UpwardDx(Ey, coefs_x, n_coefs_x, i, j, k, PMLComp::yx)
+                  + T_Algo::UpwardDx(Ey, coefs_x, n_coefs_x, i, j, k, PMLComp::yy)
+                  + T_Algo::UpwardDx(Ey, coefs_x, n_coefs_x, i, j, k, PMLComp::yz) );
+            }
+
+        );
+
+    }
+
+}
+
+#endif // corresponds to ifndef WARPX_DIM_RZ