Merge pull request #348 from atmyers/pml_exchange_gpu

Pml exchange gpu

1 files changed, 8 insertions, 2 deletions

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);
+                                       });
                 }
             }
         }