| 738 | } |
| 739 | |
| 740 | void computeDivergence (MultiFab& divu, const Array<MultiFab const*,AMREX_SPACEDIM>& umac, |
| 741 | const Geometry& geom) |
| 742 | { |
| 743 | AMREX_D_TERM(AMREX_ASSERT(divu.nComp()==umac[0]->nComp());, |
| 744 | AMREX_ASSERT(divu.nComp()==umac[1]->nComp());, |
| 745 | AMREX_ASSERT(divu.nComp()==umac[2]->nComp())); |
| 746 | |
| 747 | #if (AMREX_SPACEDIM==2) |
| 748 | const auto& ba = divu.boxArray(); |
| 749 | const auto& dm = divu.DistributionMap(); |
| 750 | MultiFab volume, areax, areay; |
| 751 | if (geom.IsRZ()) { |
| 752 | geom.GetVolume(volume, ba, dm, 0); |
| 753 | geom.GetFaceArea(areax, ba, dm, 0, 0); |
| 754 | geom.GetFaceArea(areay, ba, dm, 1, 0); |
| 755 | } |
| 756 | #endif |
| 757 | |
| 758 | const GpuArray<Real,AMREX_SPACEDIM> dxinv = geom.InvCellSizeArray(); |
| 759 | |
| 760 | #ifdef AMREX_USE_OMP |
| 761 | #pragma omp parallel if (Gpu::notInLaunchRegion()) |
| 762 | #endif |
| 763 | for (MFIter mfi(divu,TilingIfNotGPU()); mfi.isValid(); ++mfi) |
| 764 | { |
| 765 | const Box& bx = mfi.tilebox(); |
| 766 | Array4<Real> const& divuarr = divu.array(mfi); |
| 767 | AMREX_D_TERM(Array4<Real const> const& uarr = umac[0]->const_array(mfi);, |
| 768 | Array4<Real const> const& varr = umac[1]->const_array(mfi);, |
| 769 | Array4<Real const> const& warr = umac[2]->const_array(mfi);); |
| 770 | |
| 771 | #if (AMREX_SPACEDIM==2) |
| 772 | if (geom.IsRZ()) { |
| 773 | Array4<Real const> const& ax = areax.array(mfi); |
| 774 | Array4<Real const> const& ay = areay.array(mfi); |
| 775 | Array4<Real const> const& vol = volume.array(mfi); |
| 776 | |
| 777 | AMREX_LAUNCH_HOST_DEVICE_LAMBDA (bx, tbx, |
| 778 | { |
| 779 | amrex_compute_divergence_rz(tbx,divuarr,AMREX_D_DECL(uarr,varr,warr),ax,ay,vol); |
| 780 | }); |
| 781 | } else |
| 782 | #endif |
| 783 | { |
| 784 | AMREX_LAUNCH_HOST_DEVICE_LAMBDA (bx, tbx, |
| 785 | { |
| 786 | amrex_compute_divergence(tbx,divuarr,AMREX_D_DECL(uarr,varr,warr),dxinv); |
| 787 | }); |
| 788 | } |
| 789 | } |
| 790 | } |
| 791 | |
| 792 | void computeGradient (MultiFab& grad, const Array<MultiFab const*,AMREX_SPACEDIM>& umac, |
| 793 | const Geometry& geom) |
no test coverage detected