Function for computing the slow RHS for the evolution equations for the scalars other than density or potential temperature
99 BL_PROFILE_REGION(
"erf_slow_rhs_post()");
101 const BCRec* bc_ptr_d = domain_bcs_type_d.data();
102 const BCRec* bc_ptr_h = domain_bcs_type_h.data();
108 const MultiFab* t_mean_mf =
nullptr;
109 if (most) t_mean_mf = most->get_mac_avg(level,2);
111 const bool l_use_terrain = (solverChoice.
mesh_type != MeshType::ConstantDz);
112 const bool l_moving_terrain = (solverChoice.
terrain_type == TerrainType::MovingFittedMesh);
113 const bool l_reflux = (solverChoice.
coupling_type != CouplingType::OneWay);
114 if (l_moving_terrain) AMREX_ALWAYS_ASSERT(l_use_terrain);
117 const bool l_use_ddorf = (tc.
les_type == LESType::Deardorff);
118 const bool l_use_KE = ( (tc.
les_type == LESType::Deardorff) ||
124 const bool l_use_turb = ( tc.
les_type == LESType::Smagorinsky ||
125 tc.
les_type == LESType::Deardorff ||
131 const Box& domain = geom.Domain();
133 const GpuArray<Real, AMREX_SPACEDIM> dxInv = geom.InvCellSizeArray();
134 const Real* dx = geom.CellSize();
139 const Array<Real,AMREX_SPACEDIM> grav{0.0, 0.0, -solverChoice.
gravity};
140 const GpuArray<Real,AMREX_SPACEDIM> grav_gpu{grav[0], grav[1], grav[2]};
147 const DistributionMapping& dm = S_data[
IntVars::cons].DistributionMap();
149 std::unique_ptr<MultiFab> dflux_x;
150 std::unique_ptr<MultiFab> dflux_y;
151 std::unique_ptr<MultiFab> dflux_z;
154 dflux_x = std::make_unique<MultiFab>(convert(ba,IntVect(1,0,0)), dm,
nvars, 0);
155 dflux_y = std::make_unique<MultiFab>(convert(ba,IntVect(0,1,0)), dm,
nvars, 0);
156 dflux_z = std::make_unique<MultiFab>(convert(ba,IntVect(0,0,1)), dm,
nvars, 0);
164 Vector<int> is_valid_slow_var; is_valid_slow_var.resize(
RhoQ1_comp+1,0);
165 if (l_use_KE) {is_valid_slow_var[
RhoKE_comp] = 1;}
175 Vector<Real> max_scal(nvar, 1.0e34); Gpu::DeviceVector<Real> max_scal_d(nvar);
176 Vector<Real> min_scal(nvar,-1.0e34); Gpu::DeviceVector<Real> min_scal_d(nvar);
177 if (l_use_mono_adv) {
179 for (
int ivar(
RhoKE_comp); ivar<nvar; ++ivar) {
180 GpuTuple<Real,Real> mm = ParReduce(TypeList<ReduceOpMax,ReduceOpMin>{},
181 TypeList<Real, Real>{},
183 [=] AMREX_GPU_DEVICE (
int box_no,
int i,
int j,
int k) noexcept
184 -> GpuTuple<Real,Real>
186 return { ma_s_arr[box_no](i,j,k,ivar), ma_s_arr[box_no](i,j,k,ivar) };
188 max_scal[ivar] = get<0>(mm);
189 min_scal[ivar] = get<1>(mm);
192 Gpu::copy(Gpu::hostToDevice, max_scal.begin(), max_scal.end(), max_scal_d.begin());
193 Gpu::copy(Gpu::hostToDevice, min_scal.begin(), min_scal.end(), min_scal_d.begin());
194 Real* max_s_ptr = max_scal_d.data();
195 Real* min_s_ptr = min_scal_d.data();
213 #pragma omp parallel if (Gpu::notInLaunchRegion())
216 std::array<FArrayBox,AMREX_SPACEDIM> flux;
217 std::array<FArrayBox,AMREX_SPACEDIM> flux_tmp;
222 for ( MFIter mfi(S_data[
IntVars::cons],TilingIfNotGPU()); mfi.isValid(); ++mfi) {
224 Box tbx = mfi.tilebox();
229 for (
int dir = 0; dir < AMREX_SPACEDIM; ++dir) {
230 flux[dir].resize(surroundingNodes(tbx,dir),
nvars);
231 flux[dir].setVal<RunOn::Device>(0.);
232 if (l_use_mono_adv) {
233 flux_tmp[dir].resize(surroundingNodes(tbx,dir),
nvars);
234 flux_tmp[dir].setVal<RunOn::Device>(0.);
237 const GpuArray<const Array4<Real>, AMREX_SPACEDIM>
238 flx_arr{{AMREX_D_DECL(flux[0].array(), flux[1].array(), flux[2].array())}};
239 Array4<Real> tmpx = (l_use_mono_adv) ? flux_tmp[0].array() : Array4<Real>{};
240 Array4<Real> tmpy = (l_use_mono_adv) ? flux_tmp[1].array() : Array4<Real>{};
241 Array4<Real> tmpz = (l_use_mono_adv) ? flux_tmp[2].array() : Array4<Real>{};
242 const GpuArray<Array4<Real>, AMREX_SPACEDIM> flx_tmp_arr{{AMREX_D_DECL(tmpx,tmpy,tmpz)}};
247 const Array4<const Real> & old_cons = S_old[
IntVars::cons].array(mfi);
248 const Array4< Real> & cell_rhs = S_rhs[
IntVars::cons].array(mfi);
250 const Array4< Real> & new_cons = S_new[
IntVars::cons].array(mfi);
251 const Array4< Real> & new_xmom = S_new[
IntVars::xmom].array(mfi);
252 const Array4< Real> & new_ymom = S_new[
IntVars::ymom].array(mfi);
253 const Array4< Real> & new_zmom = S_new[
IntVars::zmom].array(mfi);
255 const Array4< Real> & cur_cons = S_data[
IntVars::cons].array(mfi);
256 const Array4<const Real> & cur_prim = S_prim.array(mfi);
257 const Array4< Real> & cur_xmom = S_data[
IntVars::xmom].array(mfi);
258 const Array4< Real> & cur_ymom = S_data[
IntVars::ymom].array(mfi);
259 const Array4< Real> & cur_zmom = S_data[
IntVars::zmom].array(mfi);
265 const Array4<const Real> & u =
xvel.array(mfi);
266 const Array4<const Real> & v =
yvel.array(mfi);
268 const Array4<Real const>& mu_turb = l_use_turb ? eddyDiffs->const_array(mfi) : Array4<const Real>{};
270 const Array4<const Real>& z_nd = l_use_terrain ? z_phys_nd->const_array(mfi) : Array4<const Real>{};
271 const Array4<const Real>& detJ_new_arr = l_moving_terrain ? detJ_new->const_array(mfi) : Array4<const Real>{};
274 const Array4<const Real>& mf_m = mapfac_m->const_array(mfi);
275 const Array4<const Real>& mf_u = mapfac_u->const_array(mfi);
276 const Array4<const Real>& mf_v = mapfac_v->const_array(mfi);
279 const Array4<const Real>& SmnSmn_a = l_use_ddorf ? SmnSmn->const_array(mfi) : Array4<const Real>{};
285 const GpuArray<int, IntVars::NumTypes> scomp_slow = { 2,0,0,0};
286 const GpuArray<int, IntVars::NumTypes> ncomp_slow = {nsv,0,0,0};
292 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k,
int nn) {
294 cur_cons(i,j,k,n) = new_cons(i,j,k,n);
301 Box tbx_inc = mfi.nodaltilebox(0);
302 Box tby_inc = mfi.nodaltilebox(1);
303 Box tbz_inc = mfi.nodaltilebox(2);
305 ParallelFor(tbx_inc, tby_inc, tbz_inc,
306 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
307 avg_xmom(i,j,k) = cur_xmom(i,j,k);
309 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
310 avg_ymom(i,j,k) = cur_ymom(i,j,k);
312 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
313 avg_zmom(i,j,k) = cur_zmom(i,j,k);
321 Array4<const Real> ax_arr;
322 Array4<const Real> ay_arr;
323 Array4<const Real> az_arr;
324 Array4<const Real> detJ_arr;
326 ax_arr = ebfact.getAreaFrac()[0]->const_array(mfi);
327 ay_arr = ebfact.getAreaFrac()[1]->const_array(mfi);
328 az_arr = ebfact.getAreaFrac()[2]->const_array(mfi);
329 detJ_arr = ebfact.getVolFrac().const_array(mfi);
331 ax_arr = ax->const_array(mfi);
332 ay_arr = ay->const_array(mfi);
333 az_arr = az->const_array(mfi);
334 detJ_arr = detJ->const_array(mfi);
337 AdvType horiz_adv_type, vert_adv_type;
338 Real horiz_upw_frac, vert_upw_frac;
340 Array4<Real> diffflux_x, diffflux_y, diffflux_z;
341 Array4<Real> hfx_x, hfx_y, hfx_z, diss;
342 Array4<Real> q1fx_x, q1fx_y, q1fx_z, q2fx_z;
343 const bool use_most = (most !=
nullptr);
346 diffflux_x = dflux_x->array(mfi);
347 diffflux_y = dflux_y->array(mfi);
348 diffflux_z = dflux_z->array(mfi);
350 hfx_x = Hfx1->array(mfi);
351 hfx_y = Hfx2->array(mfi);
352 hfx_z = Hfx3->array(mfi);
353 diss = Diss->array(mfi);
355 if (Q1fx1) q1fx_x = Q1fx1->array(mfi);
356 if (Q1fx2) q1fx_y = Q1fx2->array(mfi);
357 if (Q1fx3) q1fx_z = Q1fx3->array(mfi);
358 if (Q2fx3) q2fx_z = Q2fx3->array(mfi);
366 if (is_valid_slow_var[ivar])
400 cur_cons, cur_prim, cell_rhs,
401 l_use_mono_adv, max_s_ptr, min_s_ptr,
402 detJ_arr, dxInv, mf_m,
403 horiz_adv_type, vert_adv_type,
404 horiz_upw_frac, vert_upw_frac,
405 flx_arr, flx_tmp_arr, domain, bc_ptr_h);
409 const Array4<const Real> tm_arr = t_mean_mf ? t_mean_mf->const_array(mfi) : Array4<const Real>{};
412 new_cons, cur_prim, cell_rhs,
413 diffflux_x, diffflux_y, diffflux_z,
414 z_nd, ax_arr, ay_arr, az_arr, detJ_arr,
415 dxInv, SmnSmn_a, mf_m, mf_u, mf_v,
416 hfx_x, hfx_y, hfx_z, q1fx_x, q1fx_y, q1fx_z,q2fx_z, diss,
417 mu_turb, solverChoice, level,
418 tm_arr, grav_gpu, bc_ptr_d, use_most);
421 new_cons, cur_prim, cell_rhs,
422 diffflux_x, diffflux_y, diffflux_z,
423 dxInv, SmnSmn_a, mf_m, mf_u, mf_v,
424 hfx_z, q1fx_z, q2fx_z, diss,
425 mu_turb, solverChoice, level,
426 tm_arr, grav_gpu, bc_ptr_d, use_most);
432 #if defined(ERF_USE_NETCDF)
433 if (moist_set_rhs_bool)
435 Box gtbx_moist = mfi.tilebox(IntVect(0),IntVect(2,2,0));
436 const Array4<const Real> & old_cons_const = S_old[
IntVars::cons].const_array(mfi);
437 const Array4<const Real> & new_cons_const = S_new[
IntVars::cons].const_array(mfi);
438 moist_set_rhs(tbx, gtbx_moist, old_cons_const, new_cons_const, cell_rhs,
439 bdy_time_interval, start_bdy_time, new_stage_time, dt, width, set_width, domain,
440 bdy_data_xlo, bdy_data_xhi, bdy_data_ylo, bdy_data_yhi);
446 BL_PROFILE(
"rhs_post_8");
448 const Real eps = std::numeric_limits<Real>::epsilon();
450 auto const& src_arr = source.const_array(mfi);
454 if (is_valid_slow_var[ivar])
464 if (l_moving_terrain)
466 ParallelFor(tbx, num_comp,
467 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k,
int nn) noexcept {
468 const int n = start_comp + nn;
469 cell_rhs(i,j,k,n) += src_arr(i,j,k,n);
470 Real temp_val = detJ_arr(i,j,k) * old_cons(i,j,k,n) + dt * detJ_arr(i,j,k) * cell_rhs(i,j,k,n);
471 cur_cons(i,j,k,n) = temp_val / detJ_new_arr(i,j,k);
473 cur_cons(i,j,k,n) = amrex::max(cur_cons(i,j,k,n), eps);
479 ParallelFor(tbx, num_comp,
480 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k,
int nn) noexcept {
481 const int n = start_comp + nn;
482 cell_rhs(i,j,k,n) += src_arr(i,j,k,n);
483 cur_cons(i,j,k,n) = old_cons(i,j,k,n) + dt * cell_rhs(i,j,k,n);
485 cur_cons(i,j,k,n) = amrex::max(cur_cons(i,j,k,n), eps);
487 cur_cons(i,j,k,n) = amrex::max(cur_cons(i,j,k,n), 0.0);
498 BL_PROFILE(
"rhs_post_9");
501 ParallelFor(tbx, num_comp_all,
502 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k,
int n) noexcept {
503 new_cons(i,j,k,n) = cur_cons(i,j,k,n);
507 Box xtbx = mfi.nodaltilebox(0);
508 Box ytbx = mfi.nodaltilebox(1);
509 Box ztbx = mfi.nodaltilebox(2);
512 BL_PROFILE(
"rhs_post_10()");
513 ParallelFor(xtbx, ytbx, ztbx,
514 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
515 new_xmom(i,j,k) = cur_xmom(i,j,k);
517 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
518 new_ymom(i,j,k) = cur_ymom(i,j,k);
520 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
521 new_zmom(i,j,k) = cur_zmom(i,j,k);
526 BL_PROFILE(
"rhs_post_10");
528 if (l_reflux && nrk == 2) {
530 int num_comp_reflux =
nvars - strt_comp_reflux;
531 if (level < finest_level) {
532 fr_as_crse->CrseAdd(mfi,
533 {{AMREX_D_DECL(&(flux[0]), &(flux[1]), &(flux[2]))}},
534 dx, dt, strt_comp_reflux, strt_comp_reflux, num_comp_reflux, RunOn::Device);
537 fr_as_fine->FineAdd(mfi,
538 {{AMREX_D_DECL(&(flux[0]), &(flux[1]), &(flux[2]))}},
539 dx, dt, strt_comp_reflux, strt_comp_reflux, num_comp_reflux, RunOn::Device);
545 Gpu::streamSynchronize();
void AdvectionSrcForScalars(const amrex::Real &dt, const amrex::Box &bx, const int icomp, const int ncomp, const amrex::Array4< const amrex::Real > &avg_xmom, const amrex::Array4< const amrex::Real > &avg_ymom, const amrex::Array4< const amrex::Real > &avg_zmom, const amrex::Array4< const amrex::Real > &cur_cons, const amrex::Array4< const amrex::Real > &cell_prim, const amrex::Array4< amrex::Real > &src, const bool &use_mono_adv, amrex::Real *max_s_ptr, amrex::Real *min_s_ptr, const amrex::Array4< const amrex::Real > &vf_arr, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &cellSizeInv, const amrex::Array4< const amrex::Real > &mf_m, const AdvType horiz_adv_type, const AdvType vert_adv_type, const amrex::Real horiz_upw_frac, const amrex::Real vert_upw_frac, const amrex::GpuArray< const amrex::Array4< amrex::Real >, AMREX_SPACEDIM > &flx_arr, const amrex::GpuArray< amrex::Array4< amrex::Real >, AMREX_SPACEDIM > &flx_tmp_arr, const amrex::Box &domain, const amrex::BCRec *bc_ptr_h)
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE AdvType EfficientAdvType(int nrk, AdvType adv_type)
Definition: ERF_Advection.H:172
@ nvars
Definition: ERF_DataStruct.H:70
void DiffusionSrcForState_N(const amrex::Box &bx, const amrex::Box &domain, int start_comp, int num_comp, const bool &exp_most, const amrex::Array4< const amrex::Real > &u, const amrex::Array4< const amrex::Real > &v, const amrex::Array4< const amrex::Real > &cell_data, const amrex::Array4< const amrex::Real > &cell_prim, const amrex::Array4< amrex::Real > &cell_rhs, const amrex::Array4< amrex::Real > &xflux, const amrex::Array4< amrex::Real > &yflux, const amrex::Array4< amrex::Real > &zflux, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &cellSizeInv, const amrex::Array4< const amrex::Real > &SmnSmn_a, const amrex::Array4< const amrex::Real > &mf_m, const amrex::Array4< const amrex::Real > &mf_u, const amrex::Array4< const amrex::Real > &mf_v, amrex::Array4< amrex::Real > &hfx_z, amrex::Array4< amrex::Real > &qfx1_z, amrex::Array4< amrex::Real > &qfx2_z, amrex::Array4< amrex::Real > &diss, const amrex::Array4< const amrex::Real > &mu_turb, const SolverChoice &solverChoice, const int level, const amrex::Array4< const amrex::Real > &tm_arr, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > grav_gpu, const amrex::BCRec *bc_ptr, const bool use_most)
void DiffusionSrcForState_T(const amrex::Box &bx, const amrex::Box &domain, int start_comp, int num_comp, const bool &exp_most, const bool &rot_most, const amrex::Array4< const amrex::Real > &u, const amrex::Array4< const amrex::Real > &v, const amrex::Array4< const amrex::Real > &cell_data, const amrex::Array4< const amrex::Real > &cell_prim, const amrex::Array4< amrex::Real > &cell_rhs, const amrex::Array4< amrex::Real > &xflux, const amrex::Array4< amrex::Real > &yflux, const amrex::Array4< amrex::Real > &zflux, const amrex::Array4< const amrex::Real > &z_nd, const amrex::Array4< const amrex::Real > &ax, const amrex::Array4< const amrex::Real > &ay, const amrex::Array4< const amrex::Real > &az, const amrex::Array4< const amrex::Real > &detJ, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &dxInv, const amrex::Array4< const amrex::Real > &SmnSmn_a, const amrex::Array4< const amrex::Real > &mf_m, const amrex::Array4< const amrex::Real > &mf_u, const amrex::Array4< const amrex::Real > &mf_v, amrex::Array4< amrex::Real > &hfx_x, amrex::Array4< amrex::Real > &hfx_y, amrex::Array4< amrex::Real > &hfx_z, amrex::Array4< amrex::Real > &qfx1_x, amrex::Array4< amrex::Real > &qfx1_y, amrex::Array4< amrex::Real > &qfx1_z, amrex::Array4< amrex::Real > &qfx2_z, amrex::Array4< amrex::Real > &diss, const amrex::Array4< const amrex::Real > &mu_turb, const SolverChoice &solverChoice, const int level, const amrex::Array4< const amrex::Real > &tm_arr, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > grav_gpu, const amrex::BCRec *bc_ptr, const bool use_most)
#define RhoScalar_comp
Definition: ERF_IndexDefines.H:40
#define RhoTheta_comp
Definition: ERF_IndexDefines.H:37
#define RhoQ1_comp
Definition: ERF_IndexDefines.H:42
AdvType
Definition: ERF_IndexDefines.H:191
#define RhoKE_comp
Definition: ERF_IndexDefines.H:38
@ ymom
Definition: ERF_IndexDefines.H:141
@ cons
Definition: ERF_IndexDefines.H:139
@ zmom
Definition: ERF_IndexDefines.H:142
@ xmom
Definition: ERF_IndexDefines.H:140
@ xvel
Definition: ERF_IndexDefines.H:130
@ yvel
Definition: ERF_IndexDefines.H:131
Definition: ERF_AdvStruct.H:19
amrex::Real dryscal_vert_upw_frac
Definition: ERF_AdvStruct.H:295
AdvType moistscal_horiz_adv_type
Definition: ERF_AdvStruct.H:286
AdvType moistscal_vert_adv_type
Definition: ERF_AdvStruct.H:287
amrex::Real moistscal_vert_upw_frac
Definition: ERF_AdvStruct.H:297
bool use_efficient_advection
Definition: ERF_AdvStruct.H:281
amrex::Real moistscal_horiz_upw_frac
Definition: ERF_AdvStruct.H:296
AdvType dryscal_horiz_adv_type
Definition: ERF_AdvStruct.H:284
AdvType dryscal_vert_adv_type
Definition: ERF_AdvStruct.H:285
amrex::Real dryscal_horiz_upw_frac
Definition: ERF_AdvStruct.H:294
Definition: ERF_DiffStruct.H:19
MolecDiffType molec_diff_type
Definition: ERF_DiffStruct.H:84
bool use_explicit_most
Definition: ERF_DataStruct.H:654
static MeshType mesh_type
Definition: ERF_DataStruct.H:579
bool use_mono_adv
Definition: ERF_DataStruct.H:670
DiffChoice diffChoice
Definition: ERF_DataStruct.H:588
amrex::Real gravity
Definition: ERF_DataStruct.H:626
amrex::Vector< TurbChoice > turbChoice
Definition: ERF_DataStruct.H:590
amrex::Vector< int > anelastic
Definition: ERF_DataStruct.H:598
AdvChoice advChoice
Definition: ERF_DataStruct.H:587
MoistureType moisture_type
Definition: ERF_DataStruct.H:673
static TerrainType terrain_type
Definition: ERF_DataStruct.H:576
bool use_rotate_most
Definition: ERF_DataStruct.H:657
CouplingType coupling_type
Definition: ERF_DataStruct.H:672
Definition: ERF_TurbStruct.H:31
PBLType pbl_type
Definition: ERF_TurbStruct.H:205
bool use_KE
Definition: ERF_TurbStruct.H:219
LESType les_type
Definition: ERF_TurbStruct.H:176
bool advect_KE
Definition: ERF_TurbStruct.H:221