ERF
Energy Research and Forecasting: An Atmospheric Modeling Code
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Pages
ERF_SlowRhsPost.cpp File Reference
#include <AMReX.H>
#include <ERF_SrcHeaders.H>
#include <ERF_TI_slow_headers.H>
#include <ERF_EBAdvection.H>
#include <ERF_EBRedistribute.H>
Include dependency graph for ERF_SlowRhsPost.cpp:

Functions

void erf_slow_rhs_post (int level, int finest_level, int nrk, Real dt, int n_qstate, Vector< MultiFab > &S_rhs, Vector< MultiFab > &S_old, Vector< MultiFab > &S_new, Vector< MultiFab > &S_data, const MultiFab &S_prim, Vector< MultiFab > &S_scratch, const MultiFab &xvel, const MultiFab &yvel, const MultiFab &, const MultiFab &source, const MultiFab *SmnSmn, const MultiFab *eddyDiffs, MultiFab *Hfx1, MultiFab *Hfx2, MultiFab *Hfx3, MultiFab *Q1fx1, MultiFab *Q1fx2, MultiFab *Q1fx3, MultiFab *Q2fx3, MultiFab *Diss, const Geometry geom, const SolverChoice &solverChoice, std::unique_ptr< ABLMost > &most, const Gpu::DeviceVector< BCRec > &domain_bcs_type_d, const Vector< BCRec > &domain_bcs_type_h, std::unique_ptr< MultiFab > &z_phys_nd, std::unique_ptr< MultiFab > &ax, std::unique_ptr< MultiFab > &ay, std::unique_ptr< MultiFab > &az, std::unique_ptr< MultiFab > &detJ, std::unique_ptr< MultiFab > &detJ_new, std::unique_ptr< MultiFab > &mapfac_m, std::unique_ptr< MultiFab > &mapfac_u, std::unique_ptr< MultiFab > &mapfac_v, amrex::EBFArrayBoxFactory const &ebfact, YAFluxRegister *fr_as_crse, YAFluxRegister *fr_as_fine)
 

Function Documentation

◆ erf_slow_rhs_post()

void erf_slow_rhs_post ( int  level,
int  finest_level,
int  nrk,
Real  dt,
int  n_qstate,
Vector< MultiFab > &  S_rhs,
Vector< MultiFab > &  S_old,
Vector< MultiFab > &  S_new,
Vector< MultiFab > &  S_data,
const MultiFab &  S_prim,
Vector< MultiFab > &  S_scratch,
const MultiFab &  xvel,
const MultiFab &  yvel,
const MultiFab &  ,
const MultiFab &  source,
const MultiFab *  SmnSmn,
const MultiFab *  eddyDiffs,
MultiFab *  Hfx1,
MultiFab *  Hfx2,
MultiFab *  Hfx3,
MultiFab *  Q1fx1,
MultiFab *  Q1fx2,
MultiFab *  Q1fx3,
MultiFab *  Q2fx3,
MultiFab *  Diss,
const Geometry  geom,
const SolverChoice solverChoice,
std::unique_ptr< ABLMost > &  most,
const Gpu::DeviceVector< BCRec > &  domain_bcs_type_d,
const Vector< BCRec > &  domain_bcs_type_h,
std::unique_ptr< MultiFab > &  z_phys_nd,
std::unique_ptr< MultiFab > &  ax,
std::unique_ptr< MultiFab > &  ay,
std::unique_ptr< MultiFab > &  az,
std::unique_ptr< MultiFab > &  detJ,
std::unique_ptr< MultiFab > &  detJ_new,
std::unique_ptr< MultiFab > &  mapfac_m,
std::unique_ptr< MultiFab > &  mapfac_u,
std::unique_ptr< MultiFab > &  mapfac_v,
amrex::EBFArrayBoxFactory const &  ebfact,
YAFluxRegister *  fr_as_crse,
YAFluxRegister *  fr_as_fine 
)

Function for computing the slow RHS for the evolution equations for the scalars other than density or potential temperature

Parameters
[in]levellevel of resolution
[in]finest_levelfinest level of resolution
[in]nrkwhich RK stage
[in]dtslow time step
[out]S_rhsRHS computed here
[in]S_oldsolution at start of time step
[in]S_newsolution at end of current RK stage
[in]S_datacurrent solution
[in]S_primprimitive variables (i.e. conserved variables divided by density)
[in]S_scratchscratch space
[in]xvelx-component of velocity
[in]yvely-component of velocity
[in]zvelz-component of velocity
[in]sourcesource terms for conserved variables
[in]SmnSmnstrain rate magnitude
[in]eddyDiffsdiffusion coefficients for LES turbulence models
[in]Hfx3heat flux in z-dir
[in]Dissdissipation of turbulent kinetic energy
[in]geomContainer for geometric information
[in]solverChoiceContainer for solver parameters
[in]mostPointer to MOST class for Monin-Obukhov Similarity Theory boundary condition
[in]domain_bcs_type_ddevice vector for domain boundary conditions
[in]z_phys_ndheight coordinate at nodes
[in]axarea fractions on x-faces
[in]ayarea fractions on y-faces
[in]azarea fractions on z-faces
[in]detJJacobian of the metric transformation at start of time step (= 1 if use_terrain is false)
[in]detJ_newJacobian of the metric transformation at new RK stage time (= 1 if use_terrain is false)
[in]mapfac_mmap factor at cell centers
[in]mapfac_umap factor at x-faces
[in]mapfac_vmap factor at y-faces
[in,out]fr_as_crseYAFluxRegister at level l at level l / l+1 interface
[in,out]fr_as_fineYAFluxRegister at level l at level l-1 / l interface
100 {
101  BL_PROFILE_REGION("erf_slow_rhs_post()");
102 
103  const BCRec* bc_ptr_d = domain_bcs_type_d.data();
104  const BCRec* bc_ptr_h = domain_bcs_type_h.data();
105 
106  AdvChoice ac = solverChoice.advChoice;
107  DiffChoice dc = solverChoice.diffChoice;
108  TurbChoice tc = solverChoice.turbChoice[level];
109 
110  const MultiFab* t_mean_mf = nullptr;
111  if (most) t_mean_mf = most->get_mac_avg(level,2);
112 
113  const bool l_use_terrain = (solverChoice.mesh_type != MeshType::ConstantDz);
114  const bool l_moving_terrain = (solverChoice.terrain_type == TerrainType::MovingFittedMesh);
115  const bool l_reflux = (solverChoice.coupling_type != CouplingType::OneWay);
116  if (l_moving_terrain) AMREX_ALWAYS_ASSERT(l_use_terrain);
117 
118  const bool l_anelastic = solverChoice.anelastic[level];
119 
120  const bool l_use_mono_adv = solverChoice.use_mono_adv;
121  const bool l_use_KE = ( (tc.les_type == LESType::Deardorff) ||
122  (tc.rans_type == RANSType::kEqn) ||
123  (tc.pbl_type == PBLType::MYNN25) ||
124  (tc.pbl_type == PBLType::MYNNEDMF) );
125  const bool l_need_SmnSmn = ( tc.les_type == LESType::Deardorff ||
126  tc.rans_type == RANSType::kEqn );
127  const bool l_advect_KE = (tc.use_KE && tc.advect_KE);
128  const bool l_use_diff = ((dc.molec_diff_type != MolecDiffType::None) ||
129  (tc.les_type != LESType::None) ||
130  (tc.rans_type != RANSType::None) ||
131  (tc.pbl_type != PBLType::None) );
132  const bool l_use_turb = ( tc.les_type == LESType::Smagorinsky ||
133  tc.les_type == LESType::Deardorff ||
134  tc.rans_type == RANSType::kEqn ||
135  tc.pbl_type == PBLType::MYNN25 ||
136  tc.pbl_type == PBLType::MYNNEDMF ||
137  tc.pbl_type == PBLType::YSU );
138  const bool exp_most = (solverChoice.use_explicit_most);
139  const bool rot_most = (solverChoice.use_rotate_most);
140 
141  const Box& domain = geom.Domain();
142 
143  const GpuArray<Real, AMREX_SPACEDIM> dxInv = geom.InvCellSizeArray();
144  const Real* dx = geom.CellSize();
145 
146  // *************************************************************************
147  // Set gravity as a vector
148  // *************************************************************************
149  const Array<Real,AMREX_SPACEDIM> grav{0.0, 0.0, -solverChoice.gravity};
150  const GpuArray<Real,AMREX_SPACEDIM> grav_gpu{grav[0], grav[1], grav[2]};
151 
152  // *************************************************************************
153  // Pre-computed quantities
154  // *************************************************************************
155  int nvars = S_data[IntVars::cons].nComp();
156  const BoxArray& ba = S_data[IntVars::cons].boxArray();
157  const DistributionMapping& dm = S_data[IntVars::cons].DistributionMap();
158 
159  std::unique_ptr<MultiFab> dflux_x;
160  std::unique_ptr<MultiFab> dflux_y;
161  std::unique_ptr<MultiFab> dflux_z;
162 
163  if (l_use_diff) {
164  dflux_x = std::make_unique<MultiFab>(convert(ba,IntVect(1,0,0)), dm, nvars, 0);
165  dflux_y = std::make_unique<MultiFab>(convert(ba,IntVect(0,1,0)), dm, nvars, 0);
166  dflux_z = std::make_unique<MultiFab>(convert(ba,IntVect(0,0,1)), dm, nvars, 0);
167  } else {
168  dflux_x = nullptr;
169  dflux_y = nullptr;
170  dflux_z = nullptr;
171  }
172 
173  // Valid vars
174  Vector<int> is_valid_slow_var; is_valid_slow_var.resize(RhoQ1_comp+1,0);
175  if (l_use_KE) {is_valid_slow_var[ RhoKE_comp] = 1;}
176  is_valid_slow_var[RhoScalar_comp] = 1;
177  if (solverChoice.moisture_type != MoistureType::None) {
178  is_valid_slow_var[RhoQ1_comp] = 1;
179  }
180 
181  // *****************************************************************************
182  // Monotonic advection for scalars
183  // *****************************************************************************
184  int nvar = S_new[IntVars::cons].nComp();
185  Vector<Real> max_scal(nvar, 1.0e34); Gpu::DeviceVector<Real> max_scal_d(nvar);
186  Vector<Real> min_scal(nvar,-1.0e34); Gpu::DeviceVector<Real> min_scal_d(nvar);
187  if (l_use_mono_adv) {
188  auto const& ma_s_arr = S_new[IntVars::cons].const_arrays();
189  for (int ivar(RhoKE_comp); ivar<nvar; ++ivar) {
190  GpuTuple<Real,Real> mm = ParReduce(TypeList<ReduceOpMax,ReduceOpMin>{},
191  TypeList<Real, Real>{},
192  S_new[IntVars::cons], IntVect(0),
193  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept
194  -> GpuTuple<Real,Real>
195  {
196  return { ma_s_arr[box_no](i,j,k,ivar), ma_s_arr[box_no](i,j,k,ivar) };
197  });
198  max_scal[ivar] = get<0>(mm);
199  min_scal[ivar] = get<1>(mm);
200  }
201  }
202  Gpu::copy(Gpu::hostToDevice, max_scal.begin(), max_scal.end(), max_scal_d.begin());
203  Gpu::copy(Gpu::hostToDevice, min_scal.begin(), min_scal.end(), min_scal_d.begin());
204  Real* max_s_ptr = max_scal_d.data();
205  Real* min_s_ptr = min_scal_d.data();
206 
207  // *************************************************************************
208  // Calculate cell-centered eddy viscosity & diffusivities
209  //
210  // Notes -- we fill all the data in ghost cells before calling this so
211  // that we can fill the eddy viscosity in the ghost regions and
212  // not have to call a boundary filler on this data itself
213  //
214  // LES - updates both horizontal and vertical eddy viscosityS_tmp components
215  // PBL - only updates vertical eddy viscosity components so horizontal
216  // components come from the LES model or are left as zero.
217  // *************************************************************************
218 
219  // *************************************************************************
220  // Define updates and fluxes in the current RK stage
221  // *************************************************************************
222 #ifdef _OPENMP
223 #pragma omp parallel if (Gpu::notInLaunchRegion())
224 #endif
225  {
226  std::array<FArrayBox,AMREX_SPACEDIM> flux;
227  std::array<FArrayBox,AMREX_SPACEDIM> flux_tmp;
228 
229  int start_comp;
230  int num_comp;
231 
232  for ( MFIter mfi(S_data[IntVars::cons],TilingIfNotGPU()); mfi.isValid(); ++mfi) {
233 
234  Box tbx = mfi.tilebox();
235 
236  // *************************************************************************
237  // Define flux arrays for use in advection
238  // *************************************************************************
239  for (int dir = 0; dir < AMREX_SPACEDIM; ++dir) {
240  flux[dir].resize(surroundingNodes(tbx,dir),nvars);
241  flux[dir].setVal<RunOn::Device>(0.);
242  if (l_use_mono_adv) {
243  flux_tmp[dir].resize(surroundingNodes(tbx,dir),nvars);
244  flux_tmp[dir].setVal<RunOn::Device>(0.);
245  }
246  }
247  const GpuArray<const Array4<Real>, AMREX_SPACEDIM>
248  flx_arr{{AMREX_D_DECL(flux[0].array(), flux[1].array(), flux[2].array())}};
249  Array4<Real> tmpx = (l_use_mono_adv) ? flux_tmp[0].array() : Array4<Real>{};
250  Array4<Real> tmpy = (l_use_mono_adv) ? flux_tmp[1].array() : Array4<Real>{};
251  Array4<Real> tmpz = (l_use_mono_adv) ? flux_tmp[2].array() : Array4<Real>{};
252  const GpuArray<Array4<Real>, AMREX_SPACEDIM> flx_tmp_arr{{AMREX_D_DECL(tmpx,tmpy,tmpz)}};
253 
254  // *************************************************************************
255  // Define Array4's
256  // *************************************************************************
257  const Array4<const Real> & old_cons = S_old[IntVars::cons].array(mfi);
258  const Array4< Real> & cell_rhs = S_rhs[IntVars::cons].array(mfi);
259 
260  const Array4< Real> & new_cons = S_new[IntVars::cons].array(mfi);
261  const Array4< Real> & new_xmom = S_new[IntVars::xmom].array(mfi);
262  const Array4< Real> & new_ymom = S_new[IntVars::ymom].array(mfi);
263  const Array4< Real> & new_zmom = S_new[IntVars::zmom].array(mfi);
264 
265  const Array4< Real> & cur_cons = S_data[IntVars::cons].array(mfi);
266  const Array4<const Real> & cur_prim = S_prim.array(mfi);
267  const Array4< Real> & cur_xmom = S_data[IntVars::xmom].array(mfi);
268  const Array4< Real> & cur_ymom = S_data[IntVars::ymom].array(mfi);
269  const Array4< Real> & cur_zmom = S_data[IntVars::zmom].array(mfi);
270 
271  Array4<Real> avg_xmom = S_scratch[IntVars::xmom].array(mfi);
272  Array4<Real> avg_ymom = S_scratch[IntVars::ymom].array(mfi);
273  Array4<Real> avg_zmom = S_scratch[IntVars::zmom].array(mfi);
274 
275  const Array4<const Real> & u = xvel.array(mfi);
276  const Array4<const Real> & v = yvel.array(mfi);
277 
278  const Array4<Real const>& mu_turb = l_use_turb ? eddyDiffs->const_array(mfi) : Array4<const Real>{};
279 
280  const Array4<const Real>& z_nd = z_phys_nd->const_array(mfi);
281  const Array4<const Real>& detJ_new_arr = l_moving_terrain ? detJ_new->const_array(mfi) : Array4<const Real>{};
282 
283  // Map factors
284  const Array4<const Real>& mf_m = mapfac_m->const_array(mfi);
285  const Array4<const Real>& mf_u = mapfac_u->const_array(mfi);
286  const Array4<const Real>& mf_v = mapfac_v->const_array(mfi);
287 
288  // SmnSmn for KE src with Deardorff or k-eqn RANS
289  const Array4<const Real>& SmnSmn_a = l_need_SmnSmn ? SmnSmn->const_array(mfi) : Array4<const Real>{};
290 
291  // **************************************************************************
292  // Here we fill the "current" data with "new" data because that is the result of the previous RK stage
293  // **************************************************************************
294  int nsv = S_old[IntVars::cons].nComp() - 2;
295  const GpuArray<int, IntVars::NumTypes> scomp_slow = { 2,0,0,0};
296  const GpuArray<int, IntVars::NumTypes> ncomp_slow = {nsv,0,0,0};
297 
298  // **************************************************************************
299  // Note that here we do copy only the "slow" variables, not (rho) or (rho theta)
300  // **************************************************************************
301  ParallelFor(tbx, ncomp_slow[IntVars::cons],
302  [=] AMREX_GPU_DEVICE (int i, int j, int k, int nn) {
303  const int n = scomp_slow[IntVars::cons] + nn;
304  cur_cons(i,j,k,n) = new_cons(i,j,k,n);
305  });
306 
307  // We have projected the velocities stored in S_data but we will use
308  // the velocities stored in S_scratch to update the scalars, so
309  // we need to copy from S_data (projected) into S_scratch
310  if (l_anelastic) {
311  Box tbx_inc = mfi.nodaltilebox(0);
312  Box tby_inc = mfi.nodaltilebox(1);
313  Box tbz_inc = mfi.nodaltilebox(2);
314 
315  ParallelFor(tbx_inc, tby_inc, tbz_inc,
316  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept {
317  avg_xmom(i,j,k) = cur_xmom(i,j,k);
318  },
319  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept {
320  avg_ymom(i,j,k) = cur_ymom(i,j,k);
321  },
322  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept {
323  avg_zmom(i,j,k) = cur_zmom(i,j,k);
324  });
325  }
326 
327  // **************************************************************************
328  // Define updates in the RHS of continuity, temperature, and scalar equations
329  // **************************************************************************
330  // Metric terms
331  Array4<const Real> ax_arr;
332  Array4<const Real> ay_arr;
333  Array4<const Real> az_arr;
334  Array4<const Real> detJ_arr;
335  Array4<const EBCellFlag> cfg_arr;
336  if (solverChoice.terrain_type == TerrainType::EB) {
337  EBCellFlagFab const& cfg = ebfact.getMultiEBCellFlagFab()[mfi];
338  cfg_arr = cfg.const_array();
339  ax_arr = ebfact.getAreaFrac()[0]->const_array(mfi);
340  ay_arr = ebfact.getAreaFrac()[1]->const_array(mfi);
341  az_arr = ebfact.getAreaFrac()[2]->const_array(mfi);
342  detJ_arr = ebfact.getVolFrac().const_array(mfi);
343  } else {
344  ax_arr = ax->const_array(mfi);
345  ay_arr = ay->const_array(mfi);
346  az_arr = az->const_array(mfi);
347  detJ_arr = detJ->const_array(mfi);
348  }
349 
350  AdvType horiz_adv_type, vert_adv_type;
351  Real horiz_upw_frac, vert_upw_frac;
352 
353  Array4<Real> diffflux_x, diffflux_y, diffflux_z;
354  Array4<Real> hfx_x, hfx_y, hfx_z, diss;
355  Array4<Real> q1fx_x, q1fx_y, q1fx_z, q2fx_z;
356  const bool use_most = (most != nullptr);
357 
358  if (l_use_diff) {
359  diffflux_x = dflux_x->array(mfi);
360  diffflux_y = dflux_y->array(mfi);
361  diffflux_z = dflux_z->array(mfi);
362 
363  hfx_x = Hfx1->array(mfi);
364  hfx_y = Hfx2->array(mfi);
365  hfx_z = Hfx3->array(mfi);
366  diss = Diss->array(mfi);
367 
368  if (Q1fx1) q1fx_x = Q1fx1->array(mfi);
369  if (Q1fx2) q1fx_y = Q1fx2->array(mfi);
370  if (Q1fx3) q1fx_z = Q1fx3->array(mfi);
371  if (Q2fx3) q2fx_z = Q2fx3->array(mfi);
372  }
373 
374  //
375  // Note that we either advect and diffuse all or none of the moisture variables
376  //
377  for (int ivar(RhoKE_comp); ivar<= RhoQ1_comp; ++ivar)
378  {
379  if (is_valid_slow_var[ivar])
380  {
381  start_comp = ivar;
382 
383  if (ivar >= RhoQ1_comp) {
384  horiz_adv_type = ac.moistscal_horiz_adv_type;
385  vert_adv_type = ac.moistscal_vert_adv_type;
386  horiz_upw_frac = ac.moistscal_horiz_upw_frac;
387  vert_upw_frac = ac.moistscal_vert_upw_frac;
388 
389  if (ac.use_efficient_advection){
390  horiz_adv_type = EfficientAdvType(nrk,ac.moistscal_horiz_adv_type);
391  vert_adv_type = EfficientAdvType(nrk,ac.moistscal_vert_adv_type);
392  }
393 
394  num_comp = n_qstate;
395 
396  } else {
397  horiz_adv_type = ac.dryscal_horiz_adv_type;
398  vert_adv_type = ac.dryscal_vert_adv_type;
399  horiz_upw_frac = ac.dryscal_horiz_upw_frac;
400  vert_upw_frac = ac.dryscal_vert_upw_frac;
401 
402  if (ac.use_efficient_advection){
403  horiz_adv_type = EfficientAdvType(nrk,ac.dryscal_horiz_adv_type);
404  vert_adv_type = EfficientAdvType(nrk,ac.dryscal_vert_adv_type);
405  }
406  num_comp = 1;
407  }
408 
409  if (( ivar != RhoKE_comp ) ||
410  ((ivar == RhoKE_comp) && l_advect_KE))
411  {
412  if (solverChoice.terrain_type != TerrainType::EB){
413  AdvectionSrcForScalars(dt, tbx, start_comp, num_comp, avg_xmom, avg_ymom, avg_zmom,
414  cur_cons, cur_prim, cell_rhs,
415  l_use_mono_adv, max_s_ptr, min_s_ptr,
416  detJ_arr, dxInv, mf_m,
417  horiz_adv_type, vert_adv_type,
418  horiz_upw_frac, vert_upw_frac,
419  flx_arr, flx_tmp_arr, domain, bc_ptr_h);
420  } else {
421  EBAdvectionSrcForScalars(tbx, start_comp, num_comp,
422  avg_xmom, avg_ymom, avg_zmom,
423  cur_prim, cell_rhs,
424  cfg_arr, ax_arr, ay_arr, az_arr, detJ_arr, dxInv, mf_m,
425  horiz_adv_type, vert_adv_type,
426  horiz_upw_frac, vert_upw_frac,
427  flx_arr, domain, bc_ptr_h);
428  }
429  }
430 
431  if (l_use_diff) {
432  const Array4<const Real> tm_arr = t_mean_mf ? t_mean_mf->const_array(mfi) : Array4<const Real>{};
433  if (l_use_terrain) {
434  DiffusionSrcForState_T(tbx, domain, start_comp, num_comp, exp_most, rot_most, u, v,
435  new_cons, cur_prim, cell_rhs,
436  diffflux_x, diffflux_y, diffflux_z,
437  z_nd, ax_arr, ay_arr, az_arr, detJ_arr,
438  dxInv, SmnSmn_a, mf_m, mf_u, mf_v,
439  hfx_x, hfx_y, hfx_z, q1fx_x, q1fx_y, q1fx_z,q2fx_z, diss,
440  mu_turb, solverChoice, level,
441  tm_arr, grav_gpu, bc_ptr_d, use_most);
442  } else {
443  DiffusionSrcForState_N(tbx, domain, start_comp, num_comp, exp_most, u, v,
444  new_cons, cur_prim, cell_rhs,
445  diffflux_x, diffflux_y, diffflux_z,
446  dxInv, SmnSmn_a, mf_m, mf_u, mf_v,
447  hfx_z, q1fx_z, q2fx_z, diss,
448  mu_turb, solverChoice, level,
449  tm_arr, grav_gpu, bc_ptr_d, use_most);
450  }
451  } // use_diff
452  } // valid slow var
453  } // loop ivar
454 
455 #if defined(ERF_USE_NETCDF)
456  if (moist_set_rhs_bool)
457  {
458  const Array4<const Real> & old_cons_const = S_old[IntVars::cons].const_array(mfi);
459  const Array4<const Real> & new_cons_const = S_new[IntVars::cons].const_array(mfi);
460  moist_set_rhs(tbx, old_cons_const, new_cons_const, cell_rhs, bdy_time_interval,
461  start_bdy_time, new_stage_time, dt, width, set_width, domain,
462  bdy_data_xlo, bdy_data_xhi, bdy_data_ylo, bdy_data_yhi);
463  }
464 #endif
465 
466  // This updates just the "slow" conserved variables
467  {
468  BL_PROFILE("rhs_post_8");
469 
470  const Real eps = std::numeric_limits<Real>::epsilon();
471 
472  auto const& src_arr = source.const_array(mfi);
473 
474  for (int ivar(RhoKE_comp); ivar<= RhoQ1_comp; ++ivar)
475  {
476  if (is_valid_slow_var[ivar])
477  {
478  start_comp = ivar;
479 
480  if (ivar >= RhoQ1_comp) {
481  num_comp = nvars - RhoQ1_comp;
482  } else {
483  num_comp = 1;
484  }
485 
486  if (l_moving_terrain)
487  {
488  ParallelFor(tbx, num_comp,
489  [=] AMREX_GPU_DEVICE (int i, int j, int k, int nn) noexcept {
490  const int n = start_comp + nn;
491  cell_rhs(i,j,k,n) += src_arr(i,j,k,n);
492  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);
493  cur_cons(i,j,k,n) = temp_val / detJ_new_arr(i,j,k);
494  if (ivar == RhoKE_comp) {
495  cur_cons(i,j,k,n) = amrex::max(cur_cons(i,j,k,n), eps);
496  }
497  });
498 
499  } else if (l_anelastic && (nrk == 1)) { // not moving and ( (anelastic) and second RK stage) )
500 
501  ParallelFor(tbx, num_comp,
502  [=] AMREX_GPU_DEVICE (int i, int j, int k, int nn) noexcept {
503  const int n = start_comp + nn;
504  cell_rhs(i,j,k,n) += src_arr(i,j,k,n);
505 
506  // Re-construct the cell_rhs used in the first RK stage
507  Real dt_times_old_cell_rhs = cur_cons(i,j,k,n) - old_cons(i,j,k,n);
508 
509  // Add the time-averaged RHS to the old state
510  cur_cons(i,j,k,n) = old_cons(i,j,k,n) + 0.5 * (dt_times_old_cell_rhs + dt * cell_rhs(i,j,k,n));
511 
512  if (ivar == RhoKE_comp) {
513  cur_cons(i,j,k,n) = amrex::max(cur_cons(i,j,k,n), eps);
514  } else if (ivar >= RhoQ1_comp) {
515  cur_cons(i,j,k,n) = amrex::max(cur_cons(i,j,k,n), 0.0);
516  }
517  });
518 
519  } else { // not moving and ( (not anelastic) or (first RK stage) )
520 
521  ParallelFor(tbx, num_comp,
522  [=] AMREX_GPU_DEVICE (int i, int j, int k, int nn) noexcept {
523  const int n = start_comp + nn;
524  cell_rhs(i,j,k,n) += src_arr(i,j,k,n);
525  cur_cons(i,j,k,n) = old_cons(i,j,k,n) + dt * cell_rhs(i,j,k,n);
526  if (ivar == RhoKE_comp) {
527  cur_cons(i,j,k,n) = amrex::max(cur_cons(i,j,k,n), eps);
528  } else if (ivar >= RhoQ1_comp) {
529  cur_cons(i,j,k,n) = amrex::max(cur_cons(i,j,k,n), 0.0);
530  }
531  });
532 
533  } // moving, anelastic or neither?
534 
535  } // is_valid
536  } // ivar
537  } // profile
538 
539  {
540  BL_PROFILE("rhs_post_9");
541  // This updates all the conserved variables (not just the "slow" ones)
542  int num_comp_all = S_data[IntVars::cons].nComp();
543  ParallelFor(tbx, num_comp_all,
544  [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept {
545  new_cons(i,j,k,n) = cur_cons(i,j,k,n);
546  });
547  } // end profile
548 
549  Box xtbx = mfi.nodaltilebox(0);
550  Box ytbx = mfi.nodaltilebox(1);
551  Box ztbx = mfi.nodaltilebox(2);
552 
553  {
554  BL_PROFILE("rhs_post_10()");
555  ParallelFor(xtbx, ytbx, ztbx,
556  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept {
557  new_xmom(i,j,k) = cur_xmom(i,j,k);
558  },
559  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept {
560  new_ymom(i,j,k) = cur_ymom(i,j,k);
561  },
562  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept {
563  new_zmom(i,j,k) = cur_zmom(i,j,k);
564  });
565  } // end profile
566 
567  {
568  BL_PROFILE("rhs_post_10");
569  // We only add to the flux registers in the final RK step
570  if (l_reflux && nrk == 2) {
571  int strt_comp_reflux = RhoTheta_comp + 1;
572  int num_comp_reflux = nvars - strt_comp_reflux;
573  if (level < finest_level) {
574  fr_as_crse->CrseAdd(mfi,
575  {{AMREX_D_DECL(&(flux[0]), &(flux[1]), &(flux[2]))}},
576  dx, dt, strt_comp_reflux, strt_comp_reflux, num_comp_reflux, RunOn::Device);
577  }
578  if (level > 0) {
579  fr_as_fine->FineAdd(mfi,
580  {{AMREX_D_DECL(&(flux[0]), &(flux[1]), &(flux[2]))}},
581  dx, dt, strt_comp_reflux, strt_comp_reflux, num_comp_reflux, RunOn::Device);
582  }
583 
584  // This is necessary here so we don't go on to the next FArrayBox without
585  // having finished copying the fluxes into the FluxRegisters (since the fluxes
586  // are stored in temporary FArrayBox's)
587  Gpu::streamSynchronize();
588 
589  } // two-way coupling
590  } // end profile
591  } // mfi
592 
593  if (solverChoice.terrain_type == TerrainType::EB)
594  {
595  // start_comp and num_comp
596  for (int ivar(RhoKE_comp); ivar<= RhoQ1_comp; ++ivar)
597  {
598  if (is_valid_slow_var[ivar])
599  {
600  start_comp = ivar;
601  if (ivar >= RhoQ1_comp) {
602  num_comp = nvars - RhoQ1_comp;
603  } else {
604  num_comp = 1;
605  }
606  }
607  }
608 
609  // Redistribute cons states (cell-centered)
610  const int num_comp_total{S_rhs[IntVars::cons].nComp()};
611  MultiFab dUdt_tmp(ba, dm, num_comp_total, S_rhs[IntVars::cons].nGrow(), MFInfo(), ebfact);
612  dUdt_tmp.setVal(0.);
613  MultiFab::Copy(dUdt_tmp, S_rhs[IntVars::cons], start_comp, start_comp, num_comp, S_rhs[IntVars::cons].nGrow());
614  dUdt_tmp.FillBoundary(geom.periodicity());
615  dUdt_tmp.setDomainBndry(1.234e10, 0, num_comp_total, geom);
616 
617  S_old[IntVars::cons].FillBoundary(geom.periodicity());
618  S_old[IntVars::cons].setDomainBndry(1.234e10, 0, num_comp_total, geom);
619 
620  // Update S_rhs by Redistribution.
621  // To-do: Currently, redistributing all the scalar variables.
622  // This needs to be redistributed only for num_comp variables starting from ivar, for efficiency.
623  redistribute_term ( num_comp_total, geom, S_rhs[IntVars::cons], dUdt_tmp,
624  S_old[IntVars::cons], ebfact, bc_ptr_d, dt);
625 
626  // Update state using the updated S_rhs. (NOTE: redistribute_term returns RHS not state variables.)
627  for ( MFIter mfi(S_new[IntVars::cons],TilingIfNotGPU()); mfi.isValid(); ++mfi)
628  {
629  Box tbx = mfi.tilebox();
630  const Array4<Real>& snew = S_new[IntVars::cons].array(mfi);
631  const Array4<Real>& sold = S_old[IntVars::cons].array(mfi);
632  const Array4<Real>& srhs = S_rhs[IntVars::cons].array(mfi);
633  Array4<const Real> detJ_arr = ebfact.getVolFrac().const_array(mfi);
634 
635  ParallelFor(tbx, num_comp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int nn)
636  {
637  if (detJ_arr(i,j,k) > 0.0) {
638  const int n = start_comp + nn;
639  snew(i,j,k,n) = sold(i,j,k,n) + dt * srhs(i,j,k,n);
640  }
641  });
642  }
643 
644  // Redistribute momentum states (face-centered) will be added here.
645  } // EB
646 
647  } // OMP
648 }
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:254
@ nvars
Definition: ERF_DataStruct.H:74
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)
void EBAdvectionSrcForScalars(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 > &cell_prim, const amrex::Array4< amrex::Real > &src, const amrex::Array4< const amrex::EBCellFlag > &cfg_arr, const amrex::Array4< const amrex::Real > &ax_arr, const amrex::Array4< const amrex::Real > &ay_arr, const amrex::Array4< const amrex::Real > &az_arr, 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::Box &domain, const amrex::BCRec *bc_ptr_h)
void redistribute_term(int ncomp, const Geometry &geom, MultiFab &result, MultiFab &result_tmp, MultiFab const &state, EBFArrayBoxFactory const &ebfact, BCRec const *bc, Real const local_dt)
Definition: ERF_EBRedistribute.cpp:11
#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:203
#define RhoKE_comp
Definition: ERF_IndexDefines.H:38
@ ymom
Definition: ERF_IndexDefines.H:152
@ cons
Definition: ERF_IndexDefines.H:150
@ zmom
Definition: ERF_IndexDefines.H:153
@ xmom
Definition: ERF_IndexDefines.H:151
@ xvel
Definition: ERF_IndexDefines.H:141
@ yvel
Definition: ERF_IndexDefines.H:142
Definition: ERF_AdvStruct.H:19
amrex::Real dryscal_vert_upw_frac
Definition: ERF_AdvStruct.H:352
AdvType moistscal_horiz_adv_type
Definition: ERF_AdvStruct.H:343
AdvType moistscal_vert_adv_type
Definition: ERF_AdvStruct.H:344
amrex::Real moistscal_vert_upw_frac
Definition: ERF_AdvStruct.H:354
bool use_efficient_advection
Definition: ERF_AdvStruct.H:338
amrex::Real moistscal_horiz_upw_frac
Definition: ERF_AdvStruct.H:353
AdvType dryscal_horiz_adv_type
Definition: ERF_AdvStruct.H:341
AdvType dryscal_vert_adv_type
Definition: ERF_AdvStruct.H:342
amrex::Real dryscal_horiz_upw_frac
Definition: ERF_AdvStruct.H:351
Definition: ERF_DiffStruct.H:19
MolecDiffType molec_diff_type
Definition: ERF_DiffStruct.H:81
bool use_explicit_most
Definition: ERF_DataStruct.H:737
static MeshType mesh_type
Definition: ERF_DataStruct.H:665
bool use_mono_adv
Definition: ERF_DataStruct.H:753
DiffChoice diffChoice
Definition: ERF_DataStruct.H:674
amrex::Real gravity
Definition: ERF_DataStruct.H:709
amrex::Vector< TurbChoice > turbChoice
Definition: ERF_DataStruct.H:676
amrex::Vector< int > anelastic
Definition: ERF_DataStruct.H:681
AdvChoice advChoice
Definition: ERF_DataStruct.H:673
MoistureType moisture_type
Definition: ERF_DataStruct.H:759
static TerrainType terrain_type
Definition: ERF_DataStruct.H:659
bool use_rotate_most
Definition: ERF_DataStruct.H:740
CouplingType coupling_type
Definition: ERF_DataStruct.H:758
Definition: ERF_TurbStruct.H:31
PBLType pbl_type
Definition: ERF_TurbStruct.H:240
bool use_KE
Definition: ERF_TurbStruct.H:254
RANSType rans_type
Definition: ERF_TurbStruct.H:237
LESType les_type
Definition: ERF_TurbStruct.H:204
bool advect_KE
Definition: ERF_TurbStruct.H:256
Here is the call graph for this function: