32 BL_PROFILE_REGION(
"erf_make_tau_terms()");
34 const BCRec* bc_ptr_h = domain_bcs_type_h.data();
39 const bool l_use_terrain = (solverChoice.
terrain_type != TerrainType::None);
40 const bool l_moving_terrain = (solverChoice.
terrain_type == TerrainType::Moving);
41 if (l_moving_terrain) AMREX_ALWAYS_ASSERT (l_use_terrain);
48 const bool l_use_turb = ( tc.
les_type == LESType::Smagorinsky ||
52 const bool use_ddorf = (tc.
les_type == LESType::Deardorff);
54 const bool use_most = (most !=
nullptr);
58 const Box& domain = geom.Domain();
59 const int domlo_z = domain.smallEnd(2);
60 const int domhi_z = domain.bigEnd(2);
62 const GpuArray<Real, AMREX_SPACEDIM> dxInv = geom.InvCellSizeArray();
69 const DistributionMapping& dm = S_data[
IntVars::cons].DistributionMap();
71 std::unique_ptr<MultiFab> expr;
72 std::unique_ptr<MultiFab> dflux_x;
73 std::unique_ptr<MultiFab> dflux_y;
74 std::unique_ptr<MultiFab> dflux_z;
77 expr = std::make_unique<MultiFab>(ba , dm, 1, IntVect(1,1,1));
78 dflux_x = std::make_unique<MultiFab>(convert(ba,IntVect(1,0,0)), dm,
nvars, 0);
79 dflux_y = std::make_unique<MultiFab>(convert(ba,IntVect(0,1,0)), dm,
nvars, 0);
80 dflux_z = std::make_unique<MultiFab>(convert(ba,IntVect(0,0,1)), dm,
nvars, 0);
89 #pragma omp parallel if (Gpu::notInLaunchRegion())
93 const Box& bx = mfi.tilebox();
94 const Box& valid_bx = mfi.validbox();
97 const Array4<const Real> & u =
xvel.array(mfi);
98 const Array4<const Real> & v =
yvel.array(mfi);
99 const Array4<const Real> & w =
zvel.array(mfi);
102 const Array4<const Real>& mf_m = mapfac_m->const_array(mfi);
103 const Array4<const Real>& mf_u = mapfac_u->const_array(mfi);
104 const Array4<const Real>& mf_v = mapfac_v->const_array(mfi);
107 const Array4<Real const>& mu_turb = l_use_turb ? eddyDiffs->const_array(mfi) : Array4<const Real>{};
108 const Array4<Real const>& cell_data = l_use_constAlpha ? S_data[
IntVars::cons].const_array(mfi) : Array4<const Real>{};
111 const Array4<const Real>& z_nd = l_use_terrain ? z_phys_nd->const_array(mfi) : Array4<const Real>{};
112 const Array4<const Real>& detJ_arr = detJ->const_array(mfi);
131 Box bxcc = mfi.tilebox();
132 Box tbxxy = mfi.tilebox(IntVect(1,1,0));
133 Box tbxxz = mfi.tilebox(IntVect(1,0,1));
134 Box tbxyz = mfi.tilebox(IntVect(0,1,1));
137 bxcc.grow(IntVect(1,1,0));
138 tbxxy.grow(IntVect(1,1,0));
139 tbxxz.grow(IntVect(1,1,0));
140 tbxyz.grow(IntVect(1,1,0));
142 if (bxcc.smallEnd(2) != domain.smallEnd(2)) {
149 if (bxcc.bigEnd(2) != domain.bigEnd(2)) {
157 Array4<Real> er_arr = expr->array(mfi);
160 FArrayBox S11,S22,S33;
161 FArrayBox S12,S13,S23;
162 S11.resize( bxcc,1,The_Async_Arena()); S22.resize(bxcc,1,The_Async_Arena()); S33.resize(bxcc,1,The_Async_Arena());
163 S12.resize(tbxxy,1,The_Async_Arena()); S13.resize(tbxxz,1,The_Async_Arena()); S23.resize(tbxyz,1,The_Async_Arena());
164 Array4<Real> s11 = S11.array(); Array4<Real> s22 = S22.array(); Array4<Real> s33 = S33.array();
165 Array4<Real> s12 = S12.array(); Array4<Real> s13 = S13.array(); Array4<Real> s23 = S23.array();
168 Array4<Real> tau11 = Tau11->array(mfi); Array4<Real> tau22 = Tau22->array(mfi); Array4<Real> tau33 = Tau33->array(mfi);
169 Array4<Real> tau12 = Tau12->array(mfi); Array4<Real> tau13 = Tau13->array(mfi); Array4<Real> tau23 = Tau23->array(mfi);
172 Array4<Real> SmnSmn_a;
176 FArrayBox S21,S31,S32;
177 S21.resize(tbxxy,1,The_Async_Arena()); S31.resize(tbxxz,1,The_Async_Arena()); S32.resize(tbxyz,1,The_Async_Arena());
178 Array4<Real> s21 = S21.array(); Array4<Real> s31 = S31.array(); Array4<Real> s32 = S32.array();
179 Array4<Real> tau21 = Tau21->array(mfi); Array4<Real> tau31 = Tau31->array(mfi); Array4<Real> tau32 = Tau32->array(mfi);
186 BL_PROFILE(
"slow_rhs_making_er_T");
187 Box gbxo = surroundingNodes(bxcc,2);
192 Omega.resize(gbxo,1,The_Async_Arena());
195 Array4<Real> omega_arr = Omega.array();
196 ParallelFor(gbxo, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept
198 omega_arr(i,j,k) = (k == 0) ? 0. :
OmegaFromW(i,j,k,w(i,j,k),u,v,z_nd,dxInv);
201 ParallelFor(bxcc, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept
210 Real Omega_hi = omega_arr(i,j,k+1);
211 Real Omega_lo = omega_arr(i,j,k );
213 Real mfsq = mf_m(i,j,0)*mf_m(i,j,0);
215 Real expansionRate = (u(i+1,j ,k)/mf_u(i+1,j,0)*met_u_h_zeta_hi - u(i,j,k)/mf_u(i,j,0)*met_u_h_zeta_lo)*dxInv[0]*mfsq +
216 (v(i ,j+1,k)/mf_v(i,j+1,0)*met_v_h_zeta_hi - v(i,j,k)/mf_v(i,j,0)*met_v_h_zeta_lo)*dxInv[1]*mfsq +
217 (Omega_hi - Omega_lo)*dxInv[2];
219 er_arr(i,j,k) = expansionRate / detJ_arr(i,j,k);
227 BL_PROFILE(
"slow_rhs_making_strain_T");
234 z_nd, detJ_arr, bc_ptr_h, dxInv,
240 if ((nrk==0) && (use_ddorf)) {
241 SmnSmn_a = SmnSmn->array(mfi);
242 ParallelFor(bx, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept
247 domlo_z,use_most,exp_most);
255 if (use_most && exp_most) {
269 BL_PROFILE(
"slow_rhs_making_stress_T");
272 tbxxy.grow(IntVect(-1,-1,0));
273 tbxxz.grow(IntVect(-1,-1,0));
274 tbxyz.grow(IntVect(-1,-1,0));
283 er_arr, z_nd, detJ_arr, dxInv);
291 er_arr, z_nd, detJ_arr, dxInv);
295 bxcc.grow(IntVect(-1,-1,0));
296 if (bxcc.smallEnd(0) == valid_bx.smallEnd(0)) bxcc.growLo(0, 1);
297 if (bxcc.bigEnd(0) == valid_bx.bigEnd(0)) bxcc.growHi(0, 1);
298 if (bxcc.smallEnd(1) == valid_bx.smallEnd(1)) bxcc.growLo(1, 1);
299 if (bxcc.bigEnd(1) == valid_bx.bigEnd(1)) bxcc.growHi(1, 1);
303 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
304 tau11(i,j,k) = s11(i,j,k);
305 tau22(i,j,k) = s22(i,j,k);
306 tau33(i,j,k) = s33(i,j,k);
309 ParallelFor(tbxxy, tbxxz, tbxyz,
310 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
311 tau12(i,j,k) = s12(i,j,k);
312 tau21(i,j,k) = s21(i,j,k);
314 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
315 tau13(i,j,k) = s13(i,j,k);
316 tau31(i,j,k) = s31(i,j,k);
318 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
319 tau23(i,j,k) = s23(i,j,k);
320 tau32(i,j,k) = s32(i,j,k);
330 BL_PROFILE(
"slow_rhs_making_er_N");
331 ParallelFor(bxcc, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
332 Real mfsq = mf_m(i,j,0)*mf_m(i,j,0);
333 er_arr(i,j,k) = (u(i+1, j , k )/mf_u(i+1,j,0) - u(i, j, k)/mf_u(i,j,0))*dxInv[0]*mfsq +
334 (v(i , j+1, k )/mf_v(i,j+1,0) - v(i, j, k)/mf_v(i,j,0))*dxInv[1]*mfsq +
335 (w(i , j , k+1) - w(i, j, k))*dxInv[2];
344 BL_PROFILE(
"slow_rhs_making_strain_N");
355 if ((nrk==0) && (use_ddorf)) {
356 SmnSmn_a = SmnSmn->array(mfi);
357 ParallelFor(bx, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept
362 domlo_z,use_most,exp_most);
370 if (use_most && exp_most) {
380 BL_PROFILE(
"slow_rhs_making_stress_N");
383 tbxxy.grow(IntVect(-1,-1,0));
384 tbxxz.grow(IntVect(-1,-1,0));
385 tbxyz.grow(IntVect(-1,-1,0));
386 if (tbxxy.smallEnd(2) > domlo_z) {
391 if (tbxxy.bigEnd(2) < domhi_z) {
412 bxcc.grow(IntVect(-1,-1,0));
413 if (bxcc.smallEnd(0) == valid_bx.smallEnd(0)) bxcc.growLo(0, 1);
414 if (bxcc.bigEnd(0) == valid_bx.bigEnd(0)) bxcc.growHi(0, 1);
415 if (bxcc.smallEnd(1) == valid_bx.smallEnd(1)) bxcc.growLo(1, 1);
416 if (bxcc.bigEnd(1) == valid_bx.bigEnd(1)) bxcc.growHi(1, 1);
420 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
421 tau11(i,j,k) = s11(i,j,k);
422 tau22(i,j,k) = s22(i,j,k);
423 tau33(i,j,k) = s33(i,j,k);
425 ParallelFor(tbxxy, tbxxz, tbxyz,
426 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
427 tau12(i,j,k) = s12(i,j,k);
429 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
430 tau13(i,j,k) = s13(i,j,k);
432 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
433 tau23(i,j,k) = s23(i,j,k);
void ComputeStrain_N(Box bxcc, Box tbxxy, Box tbxxz, Box tbxyz, Box domain, const Array4< const Real > &u, const Array4< const Real > &v, const Array4< const Real > &w, Array4< Real > &tau11, Array4< Real > &tau22, Array4< Real > &tau33, Array4< Real > &tau12, Array4< Real > &tau13, Array4< Real > &tau23, const BCRec *bc_ptr, const GpuArray< Real, AMREX_SPACEDIM > &dxInv, const Array4< const Real > &mf_m, const Array4< const Real > &mf_u, const Array4< const Real > &mf_v)
Definition: ERF_ComputeStrain_N.cpp:28
void ComputeStrain_T(Box bxcc, Box tbxxy, Box tbxxz, Box tbxyz, Box domain, const Array4< const Real > &u, const Array4< const Real > &v, const Array4< const Real > &w, Array4< Real > &tau11, Array4< Real > &tau22, Array4< Real > &tau33, Array4< Real > &tau12, Array4< Real > &tau13, Array4< Real > &tau21, Array4< Real > &tau23, Array4< Real > &tau31, Array4< Real > &tau32, const Array4< const Real > &z_nd, const Array4< const Real > &detJ, const BCRec *bc_ptr, const GpuArray< Real, AMREX_SPACEDIM > &dxInv, const Array4< const Real > &mf_m, const Array4< const Real > &mf_u, const Array4< const Real > &mf_v)
Definition: ERF_ComputeStrain_T.cpp:33
void ComputeStressVarVisc_N(Box bxcc, Box tbxxy, Box tbxxz, Box tbxyz, Real mu_eff, const Array4< const Real > &mu_turb, const Array4< const Real > &cell_data, Array4< Real > &tau11, Array4< Real > &tau22, Array4< Real > &tau33, Array4< Real > &tau12, Array4< Real > &tau13, Array4< Real > &tau23, const Array4< const Real > &er_arr)
Definition: ERF_ComputeStress_N.cpp:105
void ComputeStressConsVisc_N(Box bxcc, Box tbxxy, Box tbxxz, Box tbxyz, Real mu_eff, const Array4< const Real > &cell_data, Array4< Real > &tau11, Array4< Real > &tau22, Array4< Real > &tau33, Array4< Real > &tau12, Array4< Real > &tau13, Array4< Real > &tau23, const Array4< const Real > &er_arr)
Definition: ERF_ComputeStress_N.cpp:23
void ComputeStressVarVisc_T(Box bxcc, Box tbxxy, Box tbxxz, Box tbxyz, Real mu_eff, const Array4< const Real > &mu_turb, const Array4< const Real > &cell_data, Array4< Real > &tau11, Array4< Real > &tau22, Array4< Real > &tau33, Array4< Real > &tau12, Array4< Real > &tau13, Array4< Real > &tau21, Array4< Real > &tau23, Array4< Real > &tau31, Array4< Real > &tau32, const Array4< const Real > &er_arr, const Array4< const Real > &z_nd, const Array4< const Real > &detJ, const GpuArray< Real, AMREX_SPACEDIM > &dxInv)
Definition: ERF_ComputeStress_T.cpp:303
void ComputeStressConsVisc_T(Box bxcc, Box tbxxy, Box tbxxz, Box tbxyz, Real mu_eff, const Array4< const Real > &cell_data, Array4< Real > &tau11, Array4< Real > &tau22, Array4< Real > &tau33, Array4< Real > &tau12, Array4< Real > &tau13, Array4< Real > &tau21, Array4< Real > &tau23, Array4< Real > &tau31, Array4< Real > &tau32, const Array4< const Real > &er_arr, const Array4< const Real > &z_nd, const Array4< const Real > &detJ, const GpuArray< Real, AMREX_SPACEDIM > &dxInv)
Definition: ERF_ComputeStress_T.cpp:29
@ nvars
Definition: ERF_DataStruct.H:70
AMREX_GPU_DEVICE AMREX_FORCE_INLINE amrex::Real ComputeSmnSmn(int &i, int &j, int &k, const amrex::Array4< amrex::Real const > &tau11, const amrex::Array4< amrex::Real const > &tau22, const amrex::Array4< amrex::Real const > &tau33, const amrex::Array4< amrex::Real const > &tau12, const amrex::Array4< amrex::Real const > &tau13, const amrex::Array4< amrex::Real const > &tau23, const int &klo, const bool &use_most, const bool &exp_most)
Definition: ERF_EddyViscosity.H:31
AMREX_GPU_DEVICE AMREX_FORCE_INLINE amrex::Real OmegaFromW(int i, int j, int k, amrex::Real w, const amrex::Array4< const amrex::Real > u_arr, const amrex::Array4< const amrex::Real > v_arr, const amrex::Array4< const amrex::Real > z_nd, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &dxInv)
Definition: ERF_TerrainMetrics.H:374
AMREX_GPU_DEVICE AMREX_FORCE_INLINE amrex::Real Compute_h_zeta_AtIface(const int &i, const int &j, const int &k, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &cellSizeInv, const amrex::Array4< const amrex::Real > &z_nd)
Definition: ERF_TerrainMetrics.H:87
AMREX_GPU_DEVICE AMREX_FORCE_INLINE amrex::Real Compute_h_zeta_AtJface(const int &i, const int &j, const int &k, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &cellSizeInv, const amrex::Array4< const amrex::Real > &z_nd)
Definition: ERF_TerrainMetrics.H:130
AMREX_FORCE_INLINE amrex::IntVect TileNoZ()
Definition: ERF_TileNoZ.H:11
@ cons
Definition: ERF_IndexDefines.H:139
@ xvel
Definition: ERF_IndexDefines.H:130
@ zvel
Definition: ERF_IndexDefines.H:132
@ yvel
Definition: ERF_IndexDefines.H:131
Definition: ERF_DiffStruct.H:19
amrex::Real rho0_trans
Definition: ERF_DiffStruct.H:91
MolecDiffType molec_diff_type
Definition: ERF_DiffStruct.H:84
amrex::Real dynamic_viscosity
Definition: ERF_DiffStruct.H:96
bool use_explicit_most
Definition: ERF_DataStruct.H:645
DiffChoice diffChoice
Definition: ERF_DataStruct.H:579
amrex::Vector< TurbChoice > turbChoice
Definition: ERF_DataStruct.H:581
static TerrainType terrain_type
Definition: ERF_DataStruct.H:567
bool use_rotate_most
Definition: ERF_DataStruct.H:648
Definition: ERF_TurbStruct.H:31
PBLType pbl_type
Definition: ERF_TurbStruct.H:201
LESType les_type
Definition: ERF_TurbStruct.H:175