Function to compute average over a plane.
810 const auto & geom =
m_geom[lev];
825 Real d_fact_new, d_fact_old;
834 int klo =
m_geom[lev].Domain().smallEnd(2);
837 Gpu::DeviceVector<Real> pavg(plane_average.size(), 0.0);
838 Real* plane_avg = pavg.data();
841 Vector<Real> denom(plane_average.size(),0.0);
842 Vector<Real> val_old(plane_average.size(),0.0);
849 Box domain = geom.Domain();
851 Array<int,AMREX_SPACEDIM> is_per = {0,0,0};
852 for (
int idim(0); idim < AMREX_SPACEDIM-1; ++idim) {
853 if (geom.isPeriodic(idim)) is_per[idim] = 1;
857 for (
int imf(0); imf < 4; ++imf) {
860 if (!fields[imf])
continue;
862 denom[imf] = 1.0 / (
Real)ncell_plane[imf];
863 val_old[imf] = plane_average[imf]*d_fact_old;
866 #pragma omp parallel if (Gpu::notInLaunchRegion())
868 for (MFIter mfi(*fields[imf],
TileNoZ()); mfi.isValid(); ++mfi) {
869 Box vbx = mfi.validbox();
870 Box pbx = mfi.tilebox();
872 if (pbx.smallEnd(2) !=
klo) {
continue; }
879 IndexType ixt = averages[imf]->boxArray().ixType();
880 for (
int idim(0); idim < AMREX_SPACEDIM-1; ++idim) {
881 if ( ixt.nodeCentered(idim) && (pbx.bigEnd(idim) == vbx.bigEnd(idim)) ) {
882 int dom_hi = domain.bigEnd(idim)+1;
883 if (pbx.bigEnd(idim) <
dom_hi || is_per[idim]) {
889 auto mf_arr = (
m_rotate) ? rot_fields[imf]->const_array(mfi) :
890 fields[imf]->const_array(mfi);
893 const auto plo = geom.ProbLoArray();
894 const auto dxInv = geom.InvCellSizeArray();
895 const auto z_phys_arr = z_phys->const_array(mfi);
896 auto x_pos_arr = x_pos->array(mfi);
897 auto y_pos_arr = y_pos->array(mfi);
898 auto z_pos_arr = z_pos->array(mfi);
899 ParallelFor(Gpu::KernelInfo().setReduction(
true), pbx, [=]
900 AMREX_GPU_DEVICE(
int i,
int j,
int , Gpu::Handler
const& handler) noexcept
904 &interp, mf_arr, z_phys_arr, plo,
dxInv, 1);
906 Gpu::deviceReduceSum(&plane_avg[imf], val, handler);
909 auto k_arr = k_indx->const_array(mfi);
910 auto j_arr = j_indx ? j_indx->const_array(mfi) : Array4<const int> {};
911 auto i_arr = i_indx ? i_indx->const_array(mfi) : Array4<const int> {};
912 ParallelFor(Gpu::KernelInfo().setReduction(
true), pbx, [=]
913 AMREX_GPU_DEVICE(
int i,
int j,
int , Gpu::Handler
const& handler) noexcept
915 int mk = k_arr(i,j,0);
916 int mj = j_arr ? j_arr(i,j,0) : j;
917 int mi = i_arr ? i_arr(i,j,0) : i;
918 Real val = mf_arr(mi,mj,mk);
919 Gpu::deviceReduceSum(&plane_avg[imf], val, handler);
934 denom[iavg] = 1.0 / (
Real)ncell_plane[iavg];
935 val_old[iavg] = plane_average[iavg]*d_fact_old;
938 #pragma omp parallel if (Gpu::notInLaunchRegion())
940 for (MFIter mfi(*fields[3],
TileNoZ()); mfi.isValid(); ++mfi)
942 Box pbx = mfi.tilebox();
944 if (pbx.smallEnd(2) !=
klo) {
continue; }
948 const Array4<Real const>& T_mf_arr = fields[2]->const_array(mfi);
949 const Array4<Real const>& qv_mf_arr = fields[3]->const_array(mfi);
950 const Array4<Real const>& qr_mf_arr = (fields[4]) ? fields[4]->const_array(mfi) :
951 Array4<const Real> {};
954 const auto plo =
m_geom[lev].ProbLoArray();
956 const auto z_phys_arr = z_phys->const_array(mfi);
957 auto x_pos_arr = x_pos->array(mfi);
958 auto y_pos_arr = y_pos->array(mfi);
959 auto z_pos_arr = z_pos->array(mfi);
960 ParallelFor(Gpu::KernelInfo().setReduction(
true), pbx, [=]
961 AMREX_GPU_DEVICE(
int i,
int j,
int , Gpu::Handler
const& handler) noexcept
966 &T_interp, T_mf_arr, z_phys_arr, plo,
dxInv, 1);
968 &qv_interp, qv_mf_arr, z_phys_arr, plo,
dxInv, 1);
974 &qr_interp, qr_mf_arr, z_phys_arr, plo,
dxInv, 1);
975 vfac = 1.0 + 0.61*qv_interp - qr_interp;
977 vfac = 1.0 + 0.61*qv_interp;
980 Gpu::deviceReduceSum(&plane_avg[iavg], val, handler);
983 auto k_arr = k_indx->const_array(mfi);
984 auto j_arr = j_indx ? j_indx->const_array(mfi) : Array4<const int> {};
985 auto i_arr = i_indx ? i_indx->const_array(mfi) : Array4<const int> {};
986 ParallelFor(Gpu::KernelInfo().setReduction(
true), pbx, [=]
987 AMREX_GPU_DEVICE(
int i,
int j,
int , Gpu::Handler
const& handler) noexcept
989 int mk = k_arr(i,j,0);
990 int mj = j_arr ? j_arr(i,j,0) : j;
991 int mi = i_arr ? i_arr(i,j,0) : i;
995 vfac = 1.0 + 0.61*qv_mf_arr(mi,mj,mk) - qr_mf_arr(mi,mj,mk);
997 vfac = 1.0 + 0.61*qv_mf_arr(mi,mj,mk);
999 const Real val = T_mf_arr(mi,mj,mk) *
vfac;
1000 Gpu::deviceReduceSum(&plane_avg[iavg], val, handler);
1008 denom[iavg] = 1.0 / (
Real)ncell_plane[iavg];
1010 Gpu::copy(Gpu::deviceToDevice, pavg.begin() + 2, pavg.begin() + 3,
1011 pavg.begin() + iavg);
1024 denom[iavg] = 1.0 / (
Real)ncell_plane[iavg];
1025 val_old[iavg] = plane_average[iavg]*d_fact_old;
1030 #pragma omp parallel if (Gpu::notInLaunchRegion())
1032 for (MFIter mfi(*fields[imf_cc],
TileNoZ()); mfi.isValid(); ++mfi)
1034 Box pbx = mfi.tilebox();
1036 if (pbx.smallEnd(2) !=
klo) {
continue; }
1038 pbx.makeSlab(2,
klo);
1041 auto u_mf_arr = (
m_rotate) ? rot_fields[imf ]->const_array(mfi) :
1042 fields[imf ]->const_array(mfi);
1043 auto v_mf_arr = (
m_rotate) ? rot_fields[imf+1]->const_array(mfi) :
1044 fields[imf+1]->const_array(mfi);
1047 const auto plo =
m_geom[lev].ProbLoArray();
1048 const auto dxInv =
m_geom[lev].InvCellSizeArray();
1049 const auto z_phys_arr = z_phys->const_array(mfi);
1050 auto x_pos_arr = x_pos->array(mfi);
1051 auto y_pos_arr = y_pos->array(mfi);
1052 auto z_pos_arr = z_pos->array(mfi);
1053 ParallelFor(Gpu::KernelInfo().setReduction(
true), pbx, [=]
1054 AMREX_GPU_DEVICE(
int i,
int j,
int , Gpu::Handler
const& handler) noexcept
1059 &u_interp, u_mf_arr, z_phys_arr, plo,
dxInv, 1);
1061 &v_interp, v_mf_arr, z_phys_arr, plo,
dxInv, 1);
1062 const Real val = std::sqrt(u_interp*u_interp + v_interp*v_interp + Vsg*Vsg);
1063 Gpu::deviceReduceSum(&plane_avg[iavg], val, handler);
1066 auto k_arr = k_indx->const_array(mfi);
1067 auto j_arr = j_indx ? j_indx->const_array(mfi) : Array4<const int> {};
1068 auto i_arr = i_indx ? i_indx->const_array(mfi) : Array4<const int> {};
1069 ParallelFor(Gpu::KernelInfo().setReduction(
true), pbx, [=]
1070 AMREX_GPU_DEVICE(
int i,
int j,
int , Gpu::Handler
const& handler) noexcept
1072 int mk = k_arr(i,j,0);
1073 int mj = j_arr ? j_arr(i,j,0) : j;
1074 int mi = i_arr ? i_arr(i,j,0) : i;
1075 const Real u_val = 0.5 * (u_mf_arr(mi,mj,mk) + u_mf_arr(mi+1,mj ,mk));
1076 const Real v_val = 0.5 * (v_mf_arr(mi,mj,mk) + v_mf_arr(mi ,mj+1,mk));
1077 const Real val = std::sqrt(u_val*u_val + v_val*v_val + Vsg*Vsg);
1078 Gpu::deviceReduceSum(&plane_avg[iavg], val, handler);
1085 Gpu::copy(Gpu::deviceToHost, pavg.begin(), pavg.end(), plane_average.begin());
1086 ParallelDescriptor::ReduceRealSum(plane_average.data(), plane_average.size());
1089 for (
int iavg(0); iavg <
m_navg; ++iavg){
1090 plane_average[iavg] *= denom[iavg]*d_fact_new;
1091 plane_average[iavg] += val_old[iavg];
1092 averages[iavg]->setVal(plane_average[iavg]);
Real vfac
Definition: ERF_InitCustomPertVels_ABL.H:26
const int klo
Definition: ERF_InitCustomPertVels_ParticleTests.H:4
amrex::GpuArray< Real, AMREX_SPACEDIM > dxInv
Definition: ERF_InitCustomPertVels_ParticleTests.H:16
ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { const auto prob_lo=geomdata.ProbLo();const auto dx=geomdata.CellSize();const Real x=(prob_lo[0]+(i+0.5) *dx[0])/mf_m(i, j, 0);const Real z=z_cc(i, j, k);Real L=std::sqrt(std::pow((x - x_c)/x_r, 2)+std::pow((z - z_c)/z_r, 2));if(L<=1.0) { Real dT=T_pert *(std::cos(PI *L)+1.0)/2.0;Real Tbar_hse=p_hse(i, j, k)/(R_d *r_hse(i, j, k));Real theta_perturbed=(Tbar_hse+dT) *std::pow(p_0/p_hse(i, j, k), rdOcp);Real theta_0=(Tbar_hse) *std::pow(p_0/p_hse(i, j, k), rdOcp);if(const_rho) { state_pert(i, j, k, RhoTheta_comp)=r_hse(i, j, k) *(theta_perturbed - theta_0);} else { state_pert(i, j, k, Rho_comp)=getRhoThetagivenP(p_hse(i, j, k))/theta_perturbed - r_hse(i, j, k);} } })
const auto & dom_hi
Definition: ERF_SetupVertDiff.H:2
amrex::Real Real
Definition: ERF_ShocInterface.H:19
AMREX_FORCE_INLINE amrex::IntVect TileNoZ()
Definition: ERF_TileNoZ.H:11
int m_navg
Definition: ERF_MOSTAverage.H:199
amrex::Vector< amrex::Vector< std::unique_ptr< amrex::MultiFab > > > m_averages
Definition: ERF_MOSTAverage.H:210
amrex::Vector< std::unique_ptr< amrex::MultiFab > > m_y_pos
Definition: ERF_MOSTAverage.H:205
amrex::Vector< std::unique_ptr< amrex::iMultiFab > > m_i_indx
Definition: ERF_MOSTAverage.H:207
amrex::Vector< amrex::MultiFab * > m_z_phys_nd
Definition: ERF_MOSTAverage.H:191
amrex::Vector< std::unique_ptr< amrex::MultiFab > > m_x_pos
Definition: ERF_MOSTAverage.H:204
amrex::Vector< amrex::Real > m_Vsg
Definition: ERF_MOSTAverage.H:239
amrex::Vector< amrex::Vector< std::unique_ptr< amrex::MultiFab > > > m_rot_fields
Definition: ERF_MOSTAverage.H:211
amrex::Vector< amrex::Vector< amrex::Real > > m_plane_average
Definition: ERF_MOSTAverage.H:216
amrex::Vector< std::unique_ptr< amrex::MultiFab > > m_z_pos
Definition: ERF_MOSTAverage.H:206
amrex::Vector< amrex::Vector< int > > m_ncell_plane
Definition: ERF_MOSTAverage.H:215
amrex::Real m_fact_new
Definition: ERF_MOSTAverage.H:234
amrex::Vector< std::unique_ptr< amrex::iMultiFab > > m_j_indx
Definition: ERF_MOSTAverage.H:208
amrex::Vector< amrex::Vector< amrex::MultiFab * > > m_fields
Definition: ERF_MOSTAverage.H:190
amrex::Real m_fact_old
Definition: ERF_MOSTAverage.H:234
bool m_interp
Definition: ERF_MOSTAverage.H:226
const amrex::Vector< amrex::Geometry > m_geom
Definition: ERF_MOSTAverage.H:189
AMREX_GPU_HOST_DEVICE static AMREX_INLINE void trilinear_interp_T(const amrex::Real &xp, const amrex::Real &yp, const amrex::Real &zp, amrex::Real *interp_vals, amrex::Array4< amrex::Real const > const &interp_array, amrex::Array4< amrex::Real const > const &z_arr, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &plo, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &dxi, const int interp_comp)
Definition: ERF_MOSTAverage.H:120
amrex::Vector< std::unique_ptr< amrex::iMultiFab > > m_k_indx
Definition: ERF_MOSTAverage.H:209