5 Vector<MultiFab>& S_old,
6 Vector<MultiFab>& F_slow,
7 const Real time_for_fp,
const Real slow_dt,
10 BL_PROFILE(
"no_substep_fun");
11 amrex::ignore_unused(nrk);
14 const auto& dxInv = fine_geom.InvCellSizeArray();
16 const amrex::GpuArray<int, IntVars::NumTypes> scomp_fast = {0,0,0,0};
17 const amrex::GpuArray<int, IntVars::NumTypes> ncomp_fast = {2,1,1,1};
19 if (verbose) amrex::Print() <<
" No-substepping time integration at level " << level
20 << std::setprecision(timeprecision)
21 <<
" to " << time_for_fp
22 <<
" with dt = " << slow_dt << std::endl;
26 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
29 for ( MFIter mfi(S_sum[
IntVars::cons],TilingIfNotGPU()); mfi.isValid(); ++mfi)
31 const Box bx = mfi.tilebox();
32 Box tbx = mfi.nodaltilebox(0);
33 Box tby = mfi.nodaltilebox(1);
34 Box tbz = mfi.nodaltilebox(2);
36 Vector<Array4<Real> > ssum_h(n_data);
37 Vector<Array4<Real> > sold_h(n_data);
38 Vector<Array4<Real> > fslow_h(n_data);
40 for (
int i = 0; i < n_data; ++i) {
41 ssum_h[i] = S_sum[i].array(mfi);
42 sold_h[i] = S_old[i].array(mfi);
43 fslow_h[i] = F_slow[i].array(mfi);
46 Gpu::AsyncVector<Array4<Real> > sold_d(n_data);
47 Gpu::AsyncVector<Array4<Real> > ssum_d(n_data);
48 Gpu::AsyncVector<Array4<Real> > fslow_d(n_data);
50 Gpu::copy(Gpu::hostToDevice, sold_h.begin(), sold_h.end(), sold_d.begin());
51 Gpu::copy(Gpu::hostToDevice, ssum_h.begin(), ssum_h.end(), ssum_d.begin());
52 Gpu::copy(Gpu::hostToDevice, fslow_h.begin(), fslow_h.end(), fslow_d.begin());
54 Array4<Real>* sold = sold_d.dataPtr();
55 Array4<Real>* ssum = ssum_d.dataPtr();
56 Array4<Real>* fslow = fslow_d.dataPtr();
59 if ( solverChoice.terrain_type == TerrainType::MovingFittedMesh )
61 const Array4<const Real>& dJ_old = detJ_cc[level]->const_array(mfi);
62 const Array4<const Real>& dJ_new = detJ_cc_new[level]->const_array(mfi);
64 const Array4<const Real>& z_nd_old = z_phys_nd[level]->const_array(mfi);
65 const Array4<const Real>& z_nd_new = z_phys_nd_new[level]->const_array(mfi);
67 const Array4<const Real>& mf_u = mapfac[level][
MapFacType::u_x]->const_array(mfi);
68 const Array4<const Real>& mf_v = mapfac[level][
MapFacType::v_y]->const_array(mfi);
70 const Array4<Real >& z_t_arr = z_t_rk[level]->array(mfi);
74 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k,
int nn) {
83 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
89 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
96 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
103 mf_u, mf_v, z_nd_new,dxInv);
105 Real dJ_old_kface = 0.5 * (dJ_old(i,j,k) + dJ_old(i,j,k-1));
106 Real dJ_new_kface = 0.5 * (dJ_new(i,j,k) + dJ_new(i,j,k-1));
113 const Array4<const Real>& dJ_old = detJ_cc[level]->const_array(mfi);
115 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k,
int nn) {
117 if (dJ_old(i,j,k) > 0.0) {
126 if (solverChoice.terrain_type == TerrainType::EB)
128 const Array4<const Real>& vfrac_u = (get_eb(level).get_u_const_factory())->getVolFrac().const_array(mfi);
129 const Array4<const Real>& vfrac_v = (get_eb(level).get_v_const_factory())->getVolFrac().const_array(mfi);
130 const Array4<const Real>& vfrac_w = (get_eb(level).get_w_const_factory())->getVolFrac().const_array(mfi);
132 ParallelFor(tbx, tby, tbz,
133 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
134 if (vfrac_u(i,j,k) > 0.0) {
142 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
143 if (vfrac_v(i,j,k) > 0.0) {
151 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
152 if (vfrac_w(i,j,k) > 0.0) {
162 ParallelFor(tbx, tby, tbz,
163 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
167 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
171 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
179 if (solverChoice.terrain_type == TerrainType::EB)
181 BL_PROFILE(
"no_substep_fun_redistribute");
186 for (
int i = 1; i <= AMREX_SPACEDIM; ++i) {
187 dUdt[i].define(F_slow[i].boxArray(), F_slow[i].DistributionMap(), F_slow[i].nComp(), F_slow[i].nGrow(), MFInfo());
189 for (
int i = 0; i <= AMREX_SPACEDIM; ++i) {
190 dUdt[i].setVal(0.0, 0, ncomp_fast[i], dUdt[i].nGrow());
191 MultiFab::Copy(dUdt[i], F_slow[i], 0, 0, F_slow[i].nComp(), 0);
192 dUdt[i].FillBoundary(fine_geom.periodicity());
193 dUdt[i].setDomainBndry(1.234e10, 0, ncomp_fast[i], fine_geom);
196 BCRec
const* bc_ptr_d = domain_bcs_type_d.data();
211 for ( MFIter mfi(S_sum[
IntVars::cons],TilingIfNotGPU()); mfi.isValid(); ++mfi)
213 const Box bx = mfi.tilebox();
214 Box tbx = mfi.nodaltilebox(0);
215 Box tby = mfi.nodaltilebox(1);
216 Box tbz = mfi.nodaltilebox(2);
218 Vector<Array4<Real> > ssum_h(n_data);
219 Vector<Array4<Real> > sold_h(n_data);
220 Vector<Array4<Real> > fslow_h(n_data);
222 for (
int i = 0; i < n_data; ++i) {
223 ssum_h[i] = S_sum[i].array(mfi);
224 sold_h[i] = S_old[i].array(mfi);
225 fslow_h[i] = F_slow[i].array(mfi);
228 Gpu::AsyncVector<Array4<Real> > sold_d(n_data);
229 Gpu::AsyncVector<Array4<Real> > ssum_d(n_data);
230 Gpu::AsyncVector<Array4<Real> > fslow_d(n_data);
232 Gpu::copy(Gpu::hostToDevice, sold_h.begin(), sold_h.end(), sold_d.begin());
233 Gpu::copy(Gpu::hostToDevice, ssum_h.begin(), ssum_h.end(), ssum_d.begin());
234 Gpu::copy(Gpu::hostToDevice, fslow_h.begin(), fslow_h.end(), fslow_d.begin());
236 Array4<Real>* sold = sold_d.dataPtr();
237 Array4<Real>* ssum = ssum_d.dataPtr();
238 Array4<Real>* fslow = fslow_d.dataPtr();
241 const Array4<const Real>& vfrac_c = (get_eb(level).get_const_factory())->getVolFrac().const_array(mfi);
243 ParallelFor(bx, ncomp_fast[
IntVars::cons], [=] AMREX_GPU_DEVICE (
int i,
int j,
int k,
int nn)
246 if (vfrac_c(i,j,k) > 0.0) {
252 const Array4<const Real>& vfrac_u = (get_eb(level).get_u_const_factory())->getVolFrac().const_array(mfi);
253 const Array4<const Real>& vfrac_v = (get_eb(level).get_v_const_factory())->getVolFrac().const_array(mfi);
254 const Array4<const Real>& vfrac_w = (get_eb(level).get_w_const_factory())->getVolFrac().const_array(mfi);
256 ParallelFor(tbx, tby, tbz,
257 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
258 if (vfrac_u(i,j,k) > 0.0) {
263 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
264 if (vfrac_v(i,j,k) > 0.0) {
269 [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) noexcept {
270 if (vfrac_w(i,j,k) > 0.0) {
284 fast_only=
true, vel_and_mom_synced=
false);
286 if (solverChoice.anelastic[level]) {
287 bool have_tb = (thin_xforce[0] || thin_yforce[0] || thin_zforce[0]);
289 project_velocity_tb(level, slow_dt, S_sum);
291 project_momenta(level, slow_dt, S_sum);
@ v_y
Definition: ERF_DataStruct.H:23
@ u_x
Definition: ERF_DataStruct.H:22
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:13
#define Rho_comp
Definition: ERF_IndexDefines.H:36
amrex::Real Real
Definition: ERF_ShocInterface.H:19
auto no_substep_fun
Definition: ERF_TI_no_substep_fun.H:4
auto apply_bcs
Definition: ERF_TI_utils.H:71
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:96
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:139
AMREX_GPU_DEVICE AMREX_FORCE_INLINE amrex::Real WFromOmega(int &i, int &j, int &k, amrex::Real omega, const amrex::Array4< const amrex::Real > &u_arr, const amrex::Array4< const amrex::Real > &v_arr, const amrex::Array4< const amrex::Real > &mf_u, const amrex::Array4< const amrex::Real > &mf_v, const amrex::Array4< const amrex::Real > &z_nd, const amrex::GpuArray< amrex::Real, AMREX_SPACEDIM > &dxInv)
Definition: ERF_TerrainMetrics.H:465
@ NumTypes
Definition: ERF_IndexDefines.H:162
@ ymom
Definition: ERF_IndexDefines.H:160
@ cons
Definition: ERF_IndexDefines.H:158
@ zmom
Definition: ERF_IndexDefines.H:161
@ xmom
Definition: ERF_IndexDefines.H:159