diff options
Diffstat (limited to 'Source/WarpX.cpp')
| -rw-r--r-- | Source/WarpX.cpp | 127 |
1 files changed, 85 insertions, 42 deletions
diff --git a/Source/WarpX.cpp b/Source/WarpX.cpp index c2cf97f30..877882037 100644 --- a/Source/WarpX.cpp +++ b/Source/WarpX.cpp @@ -18,6 +18,7 @@ #include <WarpXWrappers.h> #include <WarpXUtil.H> #include <WarpXAlgorithmSelection.H> +#include <WarpX_FDTD.H> #ifdef BL_USE_SENSEI_INSITU #include <AMReX_AmrMeshInSituBridge.H> @@ -924,18 +925,32 @@ WarpX::ComputeDivB (MultiFab& divB, int dcomp, const std::array<const MultiFab*, 3>& B, const std::array<Real,3>& dx) { + Real dxinv = 1./dx[0], dyinv = 1./dx[1], dzinv = 1./dx[2]; + +#ifdef WARPX_RZ + const Real rmin = GetInstance().Geom(0).ProbLo(0); +#endif + #ifdef _OPENMP -#pragma omp parallel +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif - for (MFIter mfi(divB, true); mfi.isValid(); ++mfi) + for (MFIter mfi(divB, TilingIfNotGPU()); mfi.isValid(); ++mfi) { const Box& bx = mfi.tilebox(); - WRPX_COMPUTE_DIVB(bx.loVect(), bx.hiVect(), - BL_TO_FORTRAN_N_ANYD(divB[mfi],dcomp), - BL_TO_FORTRAN_ANYD((*B[0])[mfi]), - BL_TO_FORTRAN_ANYD((*B[1])[mfi]), - BL_TO_FORTRAN_ANYD((*B[2])[mfi]), - dx.data()); + auto const& Bxfab = B[0]->array(mfi); + auto const& Byfab = B[1]->array(mfi); + auto const& Bzfab = B[2]->array(mfi); + auto const& divBfab = divB.array(mfi); + + ParallelFor(bx, + [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept + { + warpx_computedivb(i, j, k, dcomp, divBfab, Bxfab, Byfab, Bzfab, dxinv, dyinv, dzinv +#ifdef WARPX_RZ + ,rmin +#endif + ); + }); } } @@ -944,18 +959,32 @@ WarpX::ComputeDivB (MultiFab& divB, int dcomp, const std::array<const MultiFab*, 3>& B, const std::array<Real,3>& dx, int ngrow) { + Real dxinv = 1./dx[0], dyinv = 1./dx[1], dzinv = 1./dx[2]; + +#ifdef WARPX_RZ + const Real rmin = GetInstance().Geom(0).ProbLo(0); +#endif + #ifdef _OPENMP -#pragma omp parallel +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif - for (MFIter mfi(divB, true); mfi.isValid(); ++mfi) + for (MFIter mfi(divB, TilingIfNotGPU()); mfi.isValid(); ++mfi) { Box bx = mfi.growntilebox(ngrow); - WRPX_COMPUTE_DIVB(bx.loVect(), bx.hiVect(), - BL_TO_FORTRAN_N_ANYD(divB[mfi],dcomp), - BL_TO_FORTRAN_ANYD((*B[0])[mfi]), - BL_TO_FORTRAN_ANYD((*B[1])[mfi]), - BL_TO_FORTRAN_ANYD((*B[2])[mfi]), - dx.data()); + auto const& Bxfab = B[0]->array(mfi); + auto const& Byfab = B[1]->array(mfi); + auto const& Bzfab = B[2]->array(mfi); + auto const& divBfab = divB.array(mfi); + + ParallelFor(bx, + [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept + { + warpx_computedivb(i, j, k, dcomp, divBfab, Bxfab, Byfab, Bzfab, dxinv, dyinv, dzinv +#ifdef WARPX_RZ + ,rmin +#endif + ); + }); } } @@ -964,25 +993,32 @@ WarpX::ComputeDivE (MultiFab& divE, int dcomp, const std::array<const MultiFab*, 3>& E, const std::array<Real,3>& dx) { + Real dxinv = 1./dx[0], dyinv = 1./dx[1], dzinv = 1./dx[2]; + +#ifdef WARPX_RZ + const Real rmin = GetInstance().Geom(0).ProbLo(0); +#endif + #ifdef _OPENMP -#pragma omp parallel +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif - for (MFIter mfi(divE, true); mfi.isValid(); ++mfi) + for (MFIter mfi(divE, TilingIfNotGPU()); mfi.isValid(); ++mfi) { const Box& bx = mfi.tilebox(); + auto const& Exfab = E[0]->array(mfi); + auto const& Eyfab = E[1]->array(mfi); + auto const& Ezfab = E[2]->array(mfi); + auto const& divEfab = divE.array(mfi); + + ParallelFor(bx, + [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept + { + warpx_computedive(i, j, k, dcomp, divEfab, Exfab, Eyfab, Ezfab, dxinv, dyinv, dzinv #ifdef WARPX_RZ - const Real xmin = GetInstance().Geom(0).ProbLo(0); -#endif - WRPX_COMPUTE_DIVE(bx.loVect(), bx.hiVect(), - BL_TO_FORTRAN_N_ANYD(divE[mfi],dcomp), - BL_TO_FORTRAN_ANYD((*E[0])[mfi]), - BL_TO_FORTRAN_ANYD((*E[1])[mfi]), - BL_TO_FORTRAN_ANYD((*E[2])[mfi]), - dx.data() -#ifdef WARPX_RZ - ,&xmin + ,rmin #endif - ); + ); + }); } } @@ -991,25 +1027,32 @@ WarpX::ComputeDivE (MultiFab& divE, int dcomp, const std::array<const MultiFab*, 3>& E, const std::array<Real,3>& dx, int ngrow) { + Real dxinv = 1./dx[0], dyinv = 1./dx[1], dzinv = 1./dx[2]; + +#ifdef WARPX_RZ + const Real rmin = GetInstance().Geom(0).ProbLo(0); +#endif + #ifdef _OPENMP -#pragma omp parallel +#pragma omp parallel if (Gpu::notInLaunchRegion()) #endif - for (MFIter mfi(divE, true); mfi.isValid(); ++mfi) + for (MFIter mfi(divE, TilingIfNotGPU()); mfi.isValid(); ++mfi) { Box bx = mfi.growntilebox(ngrow); + auto const& Exfab = E[0]->array(mfi); + auto const& Eyfab = E[1]->array(mfi); + auto const& Ezfab = E[2]->array(mfi); + auto const& divEfab = divE.array(mfi); + + ParallelFor(bx, + [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept + { + warpx_computedive(i, j, k, dcomp, divEfab, Exfab, Eyfab, Ezfab, dxinv, dyinv, dzinv #ifdef WARPX_RZ - const Real xmin = GetInstance().Geom(0).ProbLo(0); -#endif - WRPX_COMPUTE_DIVE(bx.loVect(), bx.hiVect(), - BL_TO_FORTRAN_N_ANYD(divE[mfi],dcomp), - BL_TO_FORTRAN_ANYD((*E[0])[mfi]), - BL_TO_FORTRAN_ANYD((*E[1])[mfi]), - BL_TO_FORTRAN_ANYD((*E[2])[mfi]), - dx.data() -#ifdef WARPX_RZ - ,&xmin + ,rmin #endif - ); + ); + }); } } |