| 660 | namespace { |
| 661 | template <typename FAB> |
| 662 | void makeFineMask_doit (FabArray<FAB>& mask, const BoxArray& fba, |
| 663 | const IntVect& ratio, Periodicity const& period, |
| 664 | typename FAB::value_type crse_value, |
| 665 | typename FAB::value_type fine_value) |
| 666 | { |
| 667 | using value_type = typename FAB::value_type; |
| 668 | |
| 669 | Vector<Array4BoxTag<value_type> > tags; |
| 670 | |
| 671 | bool run_on_gpu = Gpu::inLaunchRegion(); |
| 672 | amrex::ignore_unused(run_on_gpu, tags); |
| 673 | |
| 674 | const BoxArray& cfba = amrex::coarsen(fba,ratio); |
| 675 | const std::vector<IntVect>& pshifts = period.shiftIntVect(); |
| 676 | #ifdef AMREX_USE_OMP |
| 677 | #pragma omp parallel if (!run_on_gpu) |
| 678 | #endif |
| 679 | { |
| 680 | std::vector <std::pair<int,Box> > isects; |
| 681 | for (MFIter mfi(mask); mfi.isValid(); ++mfi) |
| 682 | { |
| 683 | const Box& bx = mfi.fabbox(); |
| 684 | Array4<value_type> const& arr = mask.array(mfi); |
| 685 | auto& fab = mask[mfi]; |
| 686 | |
| 687 | AMREX_HOST_DEVICE_PARALLEL_FOR_3D(bx, i, j, k, |
| 688 | { |
| 689 | arr(i,j,k) = crse_value; |
| 690 | }); |
| 691 | |
| 692 | for (const auto& iv : pshifts) { |
| 693 | cfba.intersections(bx+iv, isects); |
| 694 | for (const auto& is : isects) { |
| 695 | Box const& b = is.second-iv; |
| 696 | #ifdef AMREX_USE_GPU |
| 697 | if (run_on_gpu) { |
| 698 | tags.push_back(Array4BoxTag<value_type>{.dfab = arr, .dbox = b}); |
| 699 | } else |
| 700 | #endif |
| 701 | { |
| 702 | fab.template setVal<RunOn::Host>(fine_value, b); |
| 703 | } |
| 704 | } |
| 705 | } |
| 706 | } |
| 707 | } |
| 708 | |
| 709 | #ifdef AMREX_USE_GPU |
| 710 | amrex::ParallelFor(tags, 1, |
| 711 | [=] AMREX_GPU_DEVICE (int i, int j, int k, int n, |
| 712 | Array4BoxTag<value_type> const& tag) noexcept |
| 713 | { |
| 714 | tag.dfab(i,j,k,n) = fine_value; |
| 715 | }); |
| 716 | #endif |
| 717 | } |
| 718 | } |
| 719 |
no test coverage detected