| 851 | } |
| 852 | |
| 853 | Real |
| 854 | MultiFab::max (int comp, int nghost, bool local) const |
| 855 | { |
| 856 | BL_ASSERT(nghost >= 0 && n_grow.allGE(nghost)); |
| 857 | |
| 858 | BL_PROFILE("MultiFab::max()"); |
| 859 | |
| 860 | Real mx = std::numeric_limits<Real>::lowest(); |
| 861 | |
| 862 | #ifdef AMREX_USE_EB |
| 863 | if ( this->hasEBFabFactory() ) |
| 864 | { |
| 865 | const auto& ebfactory = dynamic_cast<EBFArrayBoxFactory const&>(this->Factory()); |
| 866 | auto const& flags = ebfactory.getMultiEBCellFlagFab(); |
| 867 | #ifdef AMREX_USE_GPU |
| 868 | if (Gpu::inLaunchRegion()) { |
| 869 | auto const& flagsma = flags.const_arrays(); |
| 870 | auto const& ma = this->const_arrays(); |
| 871 | mx = ParReduce(TypeList<ReduceOpMax>{}, TypeList<Real>{}, *this, IntVect(nghost), |
| 872 | [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept -> GpuTuple<Real> |
| 873 | { |
| 874 | if (flagsma[box_no](i,j,k).isCovered()) { |
| 875 | return AMREX_REAL_LOWEST; |
| 876 | } else { |
| 877 | return ma[box_no](i,j,k,comp); |
| 878 | } |
| 879 | }); |
| 880 | } else |
| 881 | #endif |
| 882 | { |
| 883 | #ifdef AMREX_USE_OMP |
| 884 | #pragma omp parallel reduction(max:mx) |
| 885 | #endif |
| 886 | for (MFIter mfi(*this,true); mfi.isValid(); ++mfi) { |
| 887 | Box const& bx = mfi.growntilebox(nghost); |
| 888 | if (flags[mfi].getType(bx) != FabType::covered) { |
| 889 | auto const& flag = flags.const_array(mfi); |
| 890 | auto const& a = this->const_array(mfi); |
| 891 | AMREX_LOOP_3D(bx, i, j, k, |
| 892 | { |
| 893 | if (!flag(i,j,k).isCovered()) { |
| 894 | mx = std::max(mx, a(i,j,k,comp)); |
| 895 | } |
| 896 | }); |
| 897 | } |
| 898 | } |
| 899 | } |
| 900 | } |
| 901 | else |
| 902 | #endif |
| 903 | { |
| 904 | #ifdef AMREX_USE_GPU |
| 905 | if (Gpu::inLaunchRegion()) { |
| 906 | auto const& ma = this->const_arrays(); |
| 907 | mx = ParReduce(TypeList<ReduceOpMax>{}, TypeList<Real>{}, *this, IntVect(nghost), |
| 908 | [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept -> GpuTuple<Real> |
| 909 | { |
| 910 | return ma[box_no](i,j,k,comp); |
no test coverage detected