aboutsummaryrefslogtreecommitdiff
path: root/Source
diff options
context:
space:
mode:
Diffstat (limited to 'Source')
-rw-r--r--Source/BoundaryConditions/PML.H4
-rw-r--r--Source/BoundaryConditions/PML.cpp10
2 files changed, 11 insertions, 3 deletions
diff --git a/Source/BoundaryConditions/PML.H b/Source/BoundaryConditions/PML.H
index b24b13ed0..c8c02d7e7 100644
--- a/Source/BoundaryConditions/PML.H
+++ b/Source/BoundaryConditions/PML.H
@@ -124,6 +124,7 @@ public:
const std::array<amrex::MultiFab*,3>& Bp, int do_pml_in_domain);
void ExchangeE (PatchType patch_type,
const std::array<amrex::MultiFab*,3>& Ep, int do_pml_in_domain);
+
void CopyJtoPMLs (PatchType patch_type,
const std::array<amrex::MultiFab*,3>& jp);
@@ -143,6 +144,8 @@ public:
void CheckPoint (const std::string& dir) const;
void Restart (const std::string& dir);
+ static void Exchange (amrex::MultiFab& pml, amrex::MultiFab& reg, const amrex::Geometry& geom, int do_pml_in_domain);
+
private:
bool m_ok;
@@ -174,7 +177,6 @@ private:
const amrex::IntVect do_pml_Lo = amrex::IntVect::TheUnitVector(),
const amrex::IntVect do_pml_Hi = amrex::IntVect::TheUnitVector());
- static void Exchange (amrex::MultiFab& pml, amrex::MultiFab& reg, const amrex::Geometry& geom, int do_pml_in_domain);
static void CopyToPML (amrex::MultiFab& pml, amrex::MultiFab& reg, const amrex::Geometry& geom);
};
diff --git a/Source/BoundaryConditions/PML.cpp b/Source/BoundaryConditions/PML.cpp
index 8f8a2608e..edf8c8358 100644
--- a/Source/BoundaryConditions/PML.cpp
+++ b/Source/BoundaryConditions/PML.cpp
@@ -806,16 +806,22 @@ PML::Exchange (MultiFab& pml, MultiFab& reg, const Geometry& geom,
MultiFab::Copy(tmpregmf, reg, 0, 0, 1, ngr);
tmpregmf.ParallelCopy(totpmlmf, 0, 0, 1, IntVect(0), ngr, period);
#ifdef _OPENMP
-#pragma omp parallel
+#pragma omp parallel if (Gpu::notInLaunchRegion())
#endif
for (MFIter mfi(reg); mfi.isValid(); ++mfi)
{
const FArrayBox& src = tmpregmf[mfi];
FArrayBox& dst = reg[mfi];
+ const auto srcarr = src.array();
+ auto dstarr = dst.array();
const BoxList& bl = amrex::boxDiff(dst.box(), mfi.validbox());
// boxDiff avoids the outermost valid cell
for (const Box& bx : bl) {
- dst.copy(src, bx, 0, bx, 0, 1);
+ amrex::ParallelFor(bx,
+ [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+ {
+ dstarr(i,j,k,0) = srcarr(i,j,k,0);
+ });
}
}
}