| 1472 | } |
| 1473 | |
| 1474 | std::unique_ptr<MultiFab> |
| 1475 | MultiFab::OverlapMask (const Periodicity& period) const |
| 1476 | { |
| 1477 | BL_PROFILE("MultiFab::OverlapMask()"); |
| 1478 | |
| 1479 | const BoxArray& ba = boxArray(); |
| 1480 | const DistributionMapping& dm = DistributionMap(); |
| 1481 | |
| 1482 | auto p = std::make_unique<MultiFab>(ba,dm,1,0, MFInfo(), Factory()); |
| 1483 | |
| 1484 | const std::vector<IntVect>& pshifts = period.shiftIntVect(); |
| 1485 | |
| 1486 | Vector<Array4BoxTag<Real> > tags; |
| 1487 | |
| 1488 | bool run_on_gpu = Gpu::inLaunchRegion(); |
| 1489 | amrex::ignore_unused(run_on_gpu, tags); |
| 1490 | #ifdef AMREX_USE_OMP |
| 1491 | #pragma omp parallel if (!run_on_gpu) |
| 1492 | #endif |
| 1493 | { |
| 1494 | std::vector< std::pair<int,Box> > isects; |
| 1495 | |
| 1496 | for (MFIter mfi(*p); mfi.isValid(); ++mfi) |
| 1497 | { |
| 1498 | const Box& bx = (*p)[mfi].box(); |
| 1499 | Array4<Real> const& arr = p->array(mfi); |
| 1500 | |
| 1501 | AMREX_HOST_DEVICE_PARALLEL_FOR_3D(bx, i, j, k, |
| 1502 | { |
| 1503 | arr(i,j,k) = Real(0.0); |
| 1504 | }); |
| 1505 | |
| 1506 | for (const auto& iv : pshifts) |
| 1507 | { |
| 1508 | ba.intersections(bx+iv, isects); |
| 1509 | for (const auto& is : isects) |
| 1510 | { |
| 1511 | Box const& b = is.second-iv; |
| 1512 | #ifdef AMREX_USE_GPU |
| 1513 | if (run_on_gpu) { |
| 1514 | tags.push_back(Array4BoxTag<Real>{.dfab = arr, .dbox = b}); |
| 1515 | } else |
| 1516 | #endif |
| 1517 | { |
| 1518 | amrex::LoopConcurrentOnCpu(b, [=] (int i, int j, int k) noexcept |
| 1519 | { |
| 1520 | arr(i,j,k) += Real(1.0); |
| 1521 | }); |
| 1522 | } |
| 1523 | } |
| 1524 | } |
| 1525 | } |
| 1526 | } |
| 1527 | |
| 1528 | #ifdef AMREX_USE_GPU |
| 1529 | amrex::ParallelFor(tags, 1, |
| 1530 | [=] AMREX_GPU_DEVICE (int i, int j, int k, int n, Array4BoxTag<Real> const& tag) noexcept |
| 1531 | { |
no test coverage detected