| 790 | } |
| 791 | |
| 792 | void computeGradient (MultiFab& grad, const Array<MultiFab const*,AMREX_SPACEDIM>& umac, |
| 793 | const Geometry& geom) |
| 794 | { |
| 795 | AMREX_ASSERT(grad.nComp() >= AMREX_SPACEDIM); |
| 796 | |
| 797 | #if (AMREX_SPACEDIM==2) |
| 798 | const auto& ba = grad.boxArray(); |
| 799 | const auto& dm = grad.DistributionMap(); |
| 800 | MultiFab volume, areax, areay; |
| 801 | if (geom.IsRZ()) { |
| 802 | geom.GetVolume(volume, ba, dm, 0); |
| 803 | geom.GetFaceArea(areax, ba, dm, 0, 0); |
| 804 | geom.GetFaceArea(areay, ba, dm, 1, 0); |
| 805 | } |
| 806 | #endif |
| 807 | |
| 808 | const GpuArray<Real,AMREX_SPACEDIM> dxinv = geom.InvCellSizeArray(); |
| 809 | |
| 810 | #ifdef AMREX_USE_OMP |
| 811 | #pragma omp parallel if (Gpu::notInLaunchRegion()) |
| 812 | #endif |
| 813 | for (MFIter mfi(grad,TilingIfNotGPU()); mfi.isValid(); ++mfi) |
| 814 | { |
| 815 | const Box& bx = mfi.tilebox(); |
| 816 | const auto& gradfab = grad.array(mfi); |
| 817 | AMREX_D_TERM(const auto& ufab = umac[0]->const_array(mfi);, |
| 818 | const auto& vfab = umac[1]->const_array(mfi);, |
| 819 | const auto& wfab = umac[2]->const_array(mfi);); |
| 820 | #if (AMREX_SPACEDIM==2) |
| 821 | if (geom.IsRZ()) { |
| 822 | Array4<Real const> const& ax = areax.array(mfi); |
| 823 | Array4<Real const> const& ay = areay.array(mfi); |
| 824 | Array4<Real const> const& vol = volume.array(mfi); |
| 825 | |
| 826 | AMREX_LAUNCH_HOST_DEVICE_LAMBDA (bx, tbx, |
| 827 | { |
| 828 | amrex_compute_gradient_rz(tbx,gradfab,AMREX_D_DECL(ufab,vfab,wfab),ax,ay,vol); |
| 829 | }); |
| 830 | } else |
| 831 | #endif |
| 832 | { |
| 833 | AMREX_LAUNCH_HOST_DEVICE_LAMBDA (bx, tbx, |
| 834 | { |
| 835 | amrex_compute_gradient(tbx,gradfab,AMREX_D_DECL(ufab,vfab,wfab),dxinv); |
| 836 | }); |
| 837 | } |
| 838 | } |
| 839 | } |
| 840 | |
| 841 | MultiFab periodicShift (MultiFab const& mf, IntVect const& offset, |
| 842 | Periodicity const& period) |
nothing calls this directly
no test coverage detected