AMReX-Hydro
AMReX-based hydro routines for low Mach number flows
 
Loading...
Searching...
No Matches
hydro_godunov_ppm.H
Go to the documentation of this file.
1/*
2 * \file hydro_godunov_ppm.H
3 *
4 * \addtogroup Godunov
5 * @{
6 */
7
8#ifndef HYDRO_GODUNOV_PPM_H
9#define HYDRO_GODUNOV_PPM_H
10
11#include <AMReX_MultiFab.H>
12#include <AMReX_BCRec.H>
13#include <AMReX_BC_TYPES.H>
14#include <hydro_bcs_K.H>
15#include <hydro_constants.H>
16
17AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
18void
19minmod_fn (const amrex::Real sm1,
20 const amrex::Real s0,
21 const amrex::Real sp1,
22 amrex::Real& dsm,
23 amrex::Real& dsp)
24{
25 // Calculate gradients on both sides
26 dsp = sp1 - s0;
27 dsm = s0 - sm1;
28
29 if (!(dsp * dsm < amrex::Real(0.0))) {
30 // Select the smaller slope if same sign
31 if (std::abs(dsp) < std::abs(dsm)) {
32 dsm = dsp;
33 } else {
34 dsp = dsm;
35 }
36 } else {
37 // Set to zero if opposite sign
38 dsp = amrex::Real(0.0);
39 dsm = amrex::Real(0.0);
40 }
41}
42
43AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
44void
45Godunov_minmod_bc_lo (const int n,
46 const amrex::Real sm1,
47 const amrex::Real s0,
48 const amrex::Real vel_edge,
49 amrex::Real& dsm,
50 const int bclo,
51 const int domlo)
52{
53 // Ensure that left-side slope is used unchanged
54 if (n == domlo) {
55 if ( bclo == amrex::BCType::ext_dir || bclo == amrex::BCType::hoextrap ||
56 (bclo == amrex::BCType::direction_dependent && vel_edge >= amrex::Real(0.0)) )
57 {
58 dsm = s0 - sm1;
59 }
60 }
61}
62
63AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
64void
65Godunov_minmod_bc_hi (const int n,
66 const amrex::Real s0,
67 const amrex::Real sp1,
68 const amrex::Real vel_edge,
69 amrex::Real& dsp,
70 const int bchi,
71 const int domhi)
72{
73 // Ensure that the right-side slope is used unchanged
74 if (n == domhi) {
75 if ( bchi == amrex::BCType::ext_dir || bchi == amrex::BCType::hoextrap ||
76 (bchi == amrex::BCType::direction_dependent && vel_edge <= amrex::Real(0.0)) )
77 {
78 dsp = sp1 - s0;
79 }
80 }
81}
82
83namespace PPM {
84
85enum limiters {VanLeer, WENOZ, WENO_JS, NoLimiter, UPWIND, MINMOD};
86
87static constexpr int default_limiter = VanLeer;
88
89struct nolimiter {
90
91 static constexpr bool do_limiting = false;
92
93 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
94 static amrex::Real
95 sedge1(const amrex::Real sm2,
96 const amrex::Real sm1,
97 const amrex::Real s0,
98 const amrex::Real sp1,
99 const amrex::Real /*sp2*/)
100 {
101
102 amrex::Real d1 = amrex::Real(0.5) * (sp1 - sm1);
103 amrex::Real d2 = amrex::Real(0.5) * (s0 - sm2);
104
105 amrex::Real sedge = amrex::Real(0.5)*(s0 + sm1) - (amrex::Real(1)/amrex::Real(6))*(d1 - d2);
106 return sedge;
107 }
108
109 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
110 static amrex::Real
111 sedge2(const amrex::Real /*sm2*/,
112 const amrex::Real sm1,
113 const amrex::Real s0,
114 const amrex::Real sp1,
115 const amrex::Real sp2)
116 {
117 amrex::Real d1 = amrex::Real(0.5) * (sp2 - s0);
118 amrex::Real d2 = amrex::Real(0.5) * (sp1 - sm1);
119
120 amrex::Real sedge = amrex::Real(0.5)*(sp1 + s0) - (amrex::Real(1)/amrex::Real(6))*(d1 - d2);
121 return sedge;
122 }
123
124 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
125 static amrex::GpuTuple<amrex::Real,amrex::Real>
126 sm_sp(const amrex::Real /*s0*/,
127 const amrex::Real sedge1,
128 const amrex::Real sedge2)
129 {
130 return amrex::makeTuple(sedge1, sedge2);
131 }
132}; // nolimiter
133
134struct vanleer {
135
136 static constexpr bool do_limiting = true;
137
138 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
139 static amrex::Real vanLeer(const amrex::Real a, const amrex::Real b, const amrex::Real c)
140 {
141 constexpr auto small_qty_sq = amrex::Real(1.e-10) * amrex::Real(1.e-10);
142
143 amrex::Real dsc = amrex::Real(0.5)*(b - c);
144 amrex::Real dsl = amrex::Real(2.0)*(a - c);
145 amrex::Real dsr = amrex::Real(2.0)*(b - a);
146 return (dsl*dsr > small_qty_sq) ?
147 std::copysign(amrex::Real(1.0), dsc)*amrex::min(amrex::Math::abs(dsc),amrex::min(amrex::Math::abs(dsl),
148 amrex::Math::abs(dsr))) : amrex::Real(0.0);
149 }
150
151 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
152 static amrex::Real
153 sedge1(const amrex::Real sm2,
154 const amrex::Real sm1,
155 const amrex::Real s0,
156 const amrex::Real sp1,
157 const amrex::Real /*sp2*/)
158 {
159
160 amrex::Real d1 = vanLeer(s0,sp1,sm1);
161 amrex::Real d2 = vanLeer(sm1,s0,sm2);
162
163 amrex::Real sedge = amrex::Real(0.5)*(s0 + sm1) - (amrex::Real(1)/amrex::Real(6))*(d1 - d2);
164 return amrex::min(amrex::max(sedge, amrex::min(s0, sm1)),amrex::max(s0,sm1));
165
166 }
167
168 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
169 static amrex::Real
170 sedge2(const amrex::Real /*sm2*/,
171 const amrex::Real sm1,
172 const amrex::Real s0,
173 const amrex::Real sp1,
174 const amrex::Real sp2)
175 {
176
177 amrex::Real d1 = vanLeer(sp1,sp2,s0);
178 amrex::Real d2 = vanLeer(s0,sp1,sm1);
179
180 amrex::Real sedge = amrex::Real(0.5)*(sp1 + s0) - (amrex::Real(1)/amrex::Real(6))*(d1 - d2);
181 return amrex::min(amrex::max(sedge, amrex::min(s0, sp1)),amrex::max(s0,sp1));
182
183
184 }
185
186 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
187 static amrex::GpuTuple<amrex::Real,amrex::Real>
188 sm_sp(const amrex::Real s0,
189 const amrex::Real sedge1,
190 const amrex::Real sedge2)
191 {
192 amrex::Real sm_ = sedge1;
193 amrex::Real sp_ = sedge2;
194
195 if ((sedge2-s0)*(s0-sedge1) < amrex::Real(0.0)) {
196 sm_ = s0;
197 sp_ = s0;
198 } else if (amrex::Math::abs(sedge2-s0) >= amrex::Real(2.0)*amrex::Math::abs(sedge1-s0)) {
199 sp_ = amrex::Real(3)*s0 - amrex::Real(2)*sedge1;
200 } else if (amrex::Math::abs(sedge1-s0) >= amrex::Real(2.0)*amrex::Math::abs(sedge2-s0)) {
201 sm_ = amrex::Real(3)*s0 - amrex::Real(2)*sedge2;
202 }
203 return amrex::makeTuple(sm_,sp_);
204
205 // this might be the correct way to do this in case the else if are both == 0.0
206 // could also change the >= to just > for each
207// if ((sedge2-s0)*(s0-sedge1) < amrex::Real(0.0)) {
208// return amrex::makeTuple(s0,s0);
209// }
210// amrex::Real sm_ = sedge1;
211// amrex::Real sp_ = sedge2;
212// if (amrex::Math::abs(sedge2-s0) >= amrex::Real(2.0)*amrex::Math::abs(sedge1-s0)) {
213// sp_ = amrex::Real(3)*s0 - amrex::Real(2)*sedge1;
214// }
215// if (amrex::Math::abs(sedge1-s0) >= amrex::Real(2.0)*amrex::Math::abs(sedge2-s0)) {
216// sm_ = amrex::Real(3)*s0 - amrex::Real(2)*sedge2;
217// }
218// return amrex::makeTuple(sm_,sp_);
219
220 }
221}; // vanleer
222
223struct wenoz {
224
225 static constexpr bool do_limiting = false;
226
227 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
228 static amrex::Real
229 sedge2(const amrex::Real sm2,
230 const amrex::Real sm1,
231 const amrex::Real s0,
232 const amrex::Real sp1,
233 const amrex::Real sp2)
234 {
235 constexpr auto eps = amrex::Real(1.0e-6);
236
237 const amrex::Real beta1 =
238 amrex::Real(13) / amrex::Real(12) * (sm2 - amrex::Real(2) * sm1 + s0) * (sm2 - amrex::Real(2) * sm1 + s0) +
239 amrex::Real(0.25) * (sm2 - amrex::Real(4) * sm1 + amrex::Real(3) * s0) * (sm2 - amrex::Real(4) * sm1 + amrex::Real(3) * s0);
240 const amrex::Real beta2 =
241 amrex::Real(13) / amrex::Real(12) * (sm1 - amrex::Real(2) * s0 + sp1) * (sm1 - amrex::Real(2) * s0 + sp1) +
242 amrex::Real(0.25) * (sm1 - sp1) * (sm1 - sp1);
243 const amrex::Real beta3 =
244 amrex::Real(13) / amrex::Real(12) * (s0 - amrex::Real(2) * sp1 + sp2) * (s0 - amrex::Real(2) * sp1 + sp2) +
245 amrex::Real(0.25) * (amrex::Real(3) * s0 - amrex::Real(4) * sp1 + sp2) * (amrex::Real(3) * s0 - amrex::Real(4) * sp1 + sp2);
246
247 const amrex::Real t5 = amrex::Math::abs(beta3 - beta1);
248 const amrex::Real omega1 = amrex::Real(0.1) * (amrex::Real(1) + t5 / (eps + beta1));
249 const amrex::Real omega2 = amrex::Real(0.6) * (amrex::Real(1) + t5 / (eps + beta2));
250 const amrex::Real omega3 = amrex::Real(0.3) * (amrex::Real(1) + t5 / (eps + beta3));
251
252 const amrex::Real omega = omega1 + omega2 + omega3;
253
254 const amrex::Real v_1 = amrex::Real(2) * sm2 - amrex::Real(7) * sm1 + amrex::Real(11) * s0;
255 const amrex::Real v_2 = -sm1 + amrex::Real(5) * s0 + amrex::Real(2) * sp1;
256 const amrex::Real v_3 = amrex::Real(2) * s0 + amrex::Real(5) * sp1 - sp2;
257
258 return (omega1 * v_1 + omega2 * v_2 + omega3 * v_3) / (amrex::Real(6) * omega);
259 }
260
261 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
262 static amrex::Real
263 sedge1(const amrex::Real sm2,
264 const amrex::Real sm1,
265 const amrex::Real s0,
266 const amrex::Real sp1,
267 const amrex::Real sp2)
268 {
269 return sedge2(sp2,sp1,s0,sm1,sm2); // NOLINT(readability-suspicious-call-argument)
270 }
271
272 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
273 static amrex::GpuTuple<amrex::Real,amrex::Real>
274 sm_sp(const amrex::Real /*s0*/,
275 const amrex::Real sedge1,
276 const amrex::Real sedge2)
277 {
278 return amrex::makeTuple(sedge1, sedge2);
279 }
280}; // wenoz
281
282struct weno_js {
283
284 static constexpr bool do_limiting = false;
285
286 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
287 static amrex::Real
288 sedge2(const amrex::Real sm2,
289 const amrex::Real sm1,
290 const amrex::Real s0,
291 const amrex::Real sp1,
292 const amrex::Real sp2)
293 {
294 constexpr auto eps = amrex::Real(1.0e-6);
295
296 const amrex::Real beta1 =
297 amrex::Real(13) / amrex::Real(12) * (sm2 - amrex::Real(2) * sm1 + s0) *
298 (sm2 - amrex::Real(2) * sm1 + s0) +
299 amrex::Real(0.25) * (sm2 - amrex::Real(4) * sm1 + amrex::Real(3) * s0) *
300 (sm2 - amrex::Real(4) * sm1 + amrex::Real(3) * s0);
301 const amrex::Real beta2 =
302 amrex::Real(13) / amrex::Real(12) * (sm1 - amrex::Real(2) * s0 + sp1) *
303 (sm1 - amrex::Real(2) * s0 + sp1) +
304 amrex::Real(0.25) * (sm1 - sp1) * (sm1 - sp1);
305 const amrex::Real beta3 =
306 amrex::Real(13) / amrex::Real(12) * (s0 - amrex::Real(2) * sp1 + sp2) *
307 (s0 - amrex::Real(2) * sp1 + sp2) +
308 amrex::Real(0.25) * (amrex::Real(3) * s0 - amrex::Real(4) * sp1 + sp2) *
309 (amrex::Real(3) * s0 - amrex::Real(4) * sp1 + sp2);
310
311 const amrex::Real omega1 = amrex::Real(0.1) / (eps + beta1);
312 const amrex::Real omega2 = amrex::Real(0.6) / (eps + beta2);
313 const amrex::Real omega3 = amrex::Real(0.3) / (eps + beta3);
314
315 const amrex::Real omega = omega1 + omega2 + omega3;
316
317 const amrex::Real v_1 = amrex::Real(2) * sm2 - amrex::Real(7) * sm1 + amrex::Real(11) * s0;
318 const amrex::Real v_2 = -sm1 + amrex::Real(5) * s0 + amrex::Real(2) * sp1;
319 const amrex::Real v_3 = amrex::Real(2) * s0 + amrex::Real(5) * sp1 - sp2;
320
321 return (omega1 * v_1 + omega2 * v_2 + omega3 * v_3) / (amrex::Real(6) * omega);
322 }
323
324 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
325 static amrex::Real
326 sedge1(const amrex::Real sm2,
327 const amrex::Real sm1,
328 const amrex::Real s0,
329 const amrex::Real sp1,
330 const amrex::Real sp2)
331 {
332 return sedge2(sp2,sp1,s0,sm1,sm2); // NOLINT(readability-suspicious-call-argument)
333 }
334
335 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
336 static amrex::GpuTuple<amrex::Real,amrex::Real>
337 sm_sp(const amrex::Real /*s0*/,
338 const amrex::Real sedge1,
339 const amrex::Real sedge2)
340 {
341 return amrex::makeTuple(sedge1, sedge2);
342 }
343}; // weno_js
344
345struct upwind {
346
347 static constexpr bool do_limiting = false;
348
349 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
350 static amrex::Real
351 sedge1(const amrex::Real /*sm2*/,
352 const amrex::Real /*sm1*/,
353 const amrex::Real /*s0*/,
354 const amrex::Real /*sp1*/,
355 const amrex::Real /*sp2*/)
356 {
357 return 0;
358 }
359
360 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
361 static amrex::Real
362 sedge2(const amrex::Real /*sm2*/,
363 const amrex::Real /*sm1*/,
364 const amrex::Real /*s0*/,
365 const amrex::Real /*sp1*/,
366 const amrex::Real /*sp2*/)
367 {
368 return 0;
369 }
370
371 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
372 static amrex::GpuTuple<amrex::Real,amrex::Real>
373 sm_sp(const amrex::Real /*s0*/,
374 const amrex::Real sedge1,
375 const amrex::Real sedge2)
376 {
377 return amrex::makeTuple(sedge1, sedge2);
378 }
379}; // upwind
380
381struct minmod {
382
383 static constexpr bool do_limiting = false;
384
385 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
386 static amrex::Real
387 sedge1(const amrex::Real /*sm2*/,
388 const amrex::Real /*sm1*/,
389 const amrex::Real /*s0*/,
390 const amrex::Real /*sp1*/,
391 const amrex::Real /*sp2*/)
392 {
393 return 0;
394 }
395
396 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
397 static amrex::Real
398 sedge2(const amrex::Real /*sm2*/,
399 const amrex::Real /*sm1*/,
400 const amrex::Real /*s0*/,
401 const amrex::Real /*sp1*/,
402 const amrex::Real /*sp2*/)
403 {
404 return 0;
405 }
406
407 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
408 static amrex::GpuTuple<amrex::Real,amrex::Real>
409 sm_sp(const amrex::Real /*s0*/,
410 const amrex::Real sedge1,
411 const amrex::Real sedge2)
412 {
413 return amrex::makeTuple(sedge1, sedge2);
414 }
415}; // minmod
416
417template <typename Limiter>
418AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
419void SetXBCs ( const int i, const int j, const int k, const int n,
420 amrex::Real &sm, amrex::Real &sp,
421 amrex::Real &sedge1, amrex::Real &sedge2,
422 const amrex::Array4<const amrex::Real> &s,
423 const amrex::Real vel_dlo, const amrex::Real vel_dhi,
424 const int bclo, const int bchi,
425 const int domlo, const int domhi)
426{
427 using namespace amrex;
428
429 // Both the domlo and the domlo+1 branch below use s(domlo-1) as a face value.
430 // For direction_dependent that only holds when the domain face is inflow, so the
431 // gate is the domain-boundary velocity vel_dlo, not the velocity of the face
432 // bounding the cell being filled.
433 if ( (bclo == BCType::ext_dir) || (bclo == BCType::hoextrap) ||
434 (bclo == BCType::direction_dependent && vel_dlo >= Real(0)) )
435 {
436 if (i == domlo)
437 {
438 sedge2 = -amrex::Real(0.2)*s(domlo-1,j,k,n) + amrex::Real(0.75)*s(domlo,j, k, n)
439 +amrex::Real(0.5)*s(domlo+1,j,k,n) - amrex::Real(0.05)*s(domlo+2,j,k,n);
440 if constexpr (Limiter::do_limiting) {
441 sedge2 = amrex::max(sedge2, amrex::min(s(domlo+1,j,k,n), s(domlo,j,k,n)));
442 sedge2 = amrex::min(sedge2, amrex::max(s(domlo+1,j,k,n), s(domlo,j,k,n)));
443 }
444
445 sm = s(domlo-1,j,k,n);
446 sp = sedge2;
447
448 } else if (i == domlo+1) {
449
450 sedge1 = -amrex::Real(0.2)*s(domlo-1,j,k,n) + amrex::Real(0.75)*s(domlo ,j,k,n)
451 + amrex::Real(0.5)*s(domlo+1,j,k,n) - amrex::Real(0.05)*s(domlo+2,j,k,n);
452 if constexpr (Limiter::do_limiting) {
453 sedge1 = amrex::max(sedge1, amrex::min(s(domlo+1,j,k,n), s(domlo,j,k,n)));
454 sedge1 = amrex::min(sedge1, amrex::max(s(domlo+1,j,k,n), s(domlo,j,k,n)));
455 }
456
457 sp = sedge2;
458 sm = sedge1;
459
460 if constexpr (Limiter::do_limiting) {
461 if ( (sp - s(domlo+1,j,k,n))*(s(domlo+1,j,k,n) - sm) <= amrex::Real(0))
462 {
463 sp = s(domlo+1,j,k,n);
464 sm = s(domlo+1,j,k,n);
465 } else if(amrex::Math::abs(sp - s(domlo+1,j,k,n)) >= amrex::Real(2)*amrex::Math::abs(sm - s(domlo+1,j,k,n))) {
466 sp = amrex::Real(3)*s(domlo+1,j,k,n) - amrex::Real(2)*sm;
467 } else if(amrex::Math::abs(sm - s(domlo+1,j,k,n)) >= amrex::Real(2)*amrex::Math::abs(sp - s(domlo+1,j,k,n))) {
468 sm = amrex::Real(3)*s(domlo+1,j,k,n) - amrex::Real(2)*sp;
469 }
470 }
471 }
472 }
473
474 // Likewise the domhi and domhi-1 branches are gated on the hi domain face.
475 if ( (bchi == BCType::ext_dir) || (bchi == BCType::hoextrap) ||
476 (bchi == BCType::direction_dependent && vel_dhi <= Real(0)) )
477 {
478 if (i == domhi)
479 {
480 sedge1 = -amrex::Real(0.2)*s(domhi+1,j,k,n) + amrex::Real(0.75)*s(domhi,j,k, n)
481 +amrex::Real(0.5)*s(domhi-1,j,k,n) - amrex::Real(0.05)*s(domhi-2,j,k,n);
482 if constexpr (Limiter::do_limiting) {
483 sedge1 = amrex::max(sedge1, amrex::min(s(domhi-1,j,k,n), s(domhi,j,k,n)));
484 sedge1 = amrex::min(sedge1, amrex::max(s(domhi-1,j,k,n), s(domhi,j,k,n)));
485 }
486
487 sp = s(domhi+1,j,k,n);
488 sm = sedge1;
489
490 } else if (i == domhi-1) {
491
492 sedge2 = -amrex::Real(0.2)*s(domhi+1,j,k,n) + amrex::Real(0.75)*s(domhi ,j,k,n)
493 +amrex::Real(0.5)*s(domhi-1,j,k,n) - amrex::Real(0.05)*s(domhi-2,j,k,n);
494 if constexpr (Limiter::do_limiting) {
495 sedge2 = amrex::max(sedge2, amrex::min(s(domhi-1,j,k,n), s(domhi,j,k,n)));
496 sedge2 = amrex::min(sedge2, amrex::max(s(domhi-1,j,k,n), s(domhi,j,k,n)));
497 }
498
499 sp = sedge2;
500 sm = sedge1;
501
502 if constexpr (Limiter::do_limiting) {
503 if( (sp - s(domhi-1,j,k,n))*(s(domhi-1,j,k,n) - sm) <= amrex::Real(0))
504 {
505 sp = s(domhi-1,j,k,n);
506 sm = s(domhi-1,j,k,n);
507 } else if(amrex::Math::abs(sp - s(domhi-1,j,k,n)) >= amrex::Real(2.)*amrex::Math::abs(sm - s(domhi-1,j,k,n))) {
508 sp = amrex::Real(3)*s(domhi-1,j,k,n) - amrex::Real(2)*sm;
509 } else if(amrex::Math::abs(sm - s(domhi-1,j,k,n)) >= amrex::Real(2.)*amrex::Math::abs(sp - s(domhi-1,j,k,n))) {
510 sm = amrex::Real(3)*s(domhi-1,j,k,n) - amrex::Real(2)*sp;
511 }
512 }
513 }
514 }
515}
516
517template <typename Limiter>
518AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
519void SetYBCs ( const int i, const int j, const int k, const int n,
520 amrex::Real &sm, amrex::Real &sp,
521 amrex::Real &sedge1, amrex::Real &sedge2,
522 const amrex::Array4<const amrex::Real> &s,
523 const amrex::Real vel_dlo, const amrex::Real vel_dhi,
524 const int bclo, const int bchi,
525 const int domlo, const int domhi)
526{
527 using namespace amrex;
528
529 // Both the domlo and the domlo+1 branch below use s(domlo-1) as a face value.
530 // For direction_dependent that only holds when the domain face is inflow, so the
531 // gate is the domain-boundary velocity vel_dlo, not the velocity of the face
532 // bounding the cell being filled.
533 if ( (bclo == BCType::ext_dir) || (bclo == BCType::hoextrap) ||
534 (bclo == BCType::direction_dependent && vel_dlo >= Real(0)) )
535 {
536 if (j == domlo)
537 {
538 sedge2 = -amrex::Real(0.2)*s(i,domlo-1,k,n) + amrex::Real(0.75)*s(i,domlo ,k,n)
539 +amrex::Real(0.5)*s(i,domlo+1,k,n) - amrex::Real(0.05)*s(i,domlo+2,k,n);
540 if constexpr (Limiter::do_limiting) {
541 sedge2 = amrex::max(sedge2, amrex::min(s(i,domlo+1,k,n), s(i,domlo,k,n)));
542 sedge2 = amrex::min(sedge2, amrex::max(s(i,domlo+1,k,n), s(i,domlo,k,n)));
543 }
544
545 sm = s(i,domlo-1,k,n);
546 sp = sedge2;
547
548 } else if (j == domlo+1) {
549
550 sedge1 = -amrex::Real(0.2)*s(i,domlo-1,k,n) + amrex::Real(0.75)*s(i,domlo ,k,n)
551 +amrex::Real(0.5)*s(i,domlo+1,k,n) - amrex::Real(0.05)*s(i,domlo+2,k,n);
552 if constexpr (Limiter::do_limiting) {
553 sedge1 = amrex::max(sedge1, amrex::min(s(i,domlo+1,k,n), s(i,domlo,k,n)));
554 sedge1 = amrex::min(sedge1, amrex::max(s(i,domlo+1,k,n), s(i,domlo,k,n)));
555 }
556
557 sp = sedge2;
558 sm = sedge1;
559
560 if constexpr (Limiter::do_limiting) {
561 if ( (sp - s(i,domlo+1,k,n))*(s(i,domlo+1,k,n) - sm) <= amrex::Real(0))
562 {
563 sp = s(i,domlo+1,k,n);
564 sm = s(i,domlo+1,k,n);
565 } else if(amrex::Math::abs(sp - s(i,domlo+1,k,n)) >= amrex::Real(2)*amrex::Math::abs(sm - s(i,domlo+1,k,n))) {
566 sp = amrex::Real(3)*s(i,domlo+1,k,n) - amrex::Real(2)*sm;
567 } else if(amrex::Math::abs(sm - s(i,domlo+1,k,n)) >= amrex::Real(2)*amrex::Math::abs(sp - s(i,domlo+1,k,n))) {
568 sm = amrex::Real(3)*s(i,domlo+1,k,n) - amrex::Real(2)*sp;
569 }
570 }
571 }
572 }
573
574 // Likewise the domhi and domhi-1 branches are gated on the hi domain face.
575 if ( (bchi == BCType::ext_dir) || (bchi == BCType::hoextrap) ||
576 (bchi == BCType::direction_dependent && vel_dhi <= Real(0)) )
577 {
578 if (j == domhi)
579 {
580 sedge1 = -amrex::Real(0.2)*s(i,domhi+1,k,n) + amrex::Real(0.75)*s(i,domhi ,k,n)
581 +amrex::Real(0.5)*s(i,domhi-1,k,n) - amrex::Real(0.05)*s(i,domhi-2,k,n);
582 if constexpr (Limiter::do_limiting) {
583 sedge1 = amrex::max(sedge1, amrex::min(s(i,domhi-1,k,n), s(i,domhi,k,n)));
584 sedge1 = amrex::min(sedge1, amrex::max(s(i,domhi-1,k,n), s(i,domhi,k,n)));
585 }
586
587 sp = s(i,domhi+1,k, n);
588 sm = sedge1;
589
590 } else if (j == domhi-1) {
591
592 sedge2 = -amrex::Real(0.2)*s(i,domhi+1,k,n) + amrex::Real(0.75)*s(i,domhi ,k,n)
593 +amrex::Real(0.5)*s(i,domhi-1,k,n) - amrex::Real(0.05)*s(i,domhi-2,k,n);
594 if constexpr (Limiter::do_limiting) {
595 sedge2 = amrex::max(sedge2, amrex::min(s(i,domhi-1,k,n), s(i,domhi,k,n)));
596 sedge2 = amrex::min(sedge2, amrex::max(s(i,domhi-1,k,n), s(i,domhi,k,n)));
597 }
598
599 sp = sedge2;
600 sm = sedge1;
601
602 if constexpr (Limiter::do_limiting) {
603 if( (sp - s(i,domhi-1,k,n))*(s(i,domhi-1,k,n) - sm) <= amrex::Real(0)){
604 sp = s(i,domhi-1,k,n);
605 sm = s(i,domhi-1,k,n);
606 } else if(amrex::Math::abs(sp - s(i,domhi-1,k,n)) >= amrex::Real(2.)*amrex::Math::abs(sm - s(i,domhi-1,k,n))) {
607 sp = amrex::Real(3)*s(i,domhi-1,k,n) - amrex::Real(2)*sm;
608 } else if(amrex::Math::abs(sm - s(i,domhi-1,k,n)) >= amrex::Real(2.)*amrex::Math::abs(sp - s(i,domhi-1,k,n))) {
609 sm = amrex::Real(3)*s(i,domhi-1,k,n) - amrex::Real(2)*sp;
610 }
611 }
612 }
613 }
614}
615
616#if (AMREX_SPACEDIM==3)
617template <typename Limiter>
618AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
619void SetZBCs ( const int i, const int j, const int k, const int n,
620 amrex::Real &sm, amrex::Real &sp,
621 amrex::Real &sedge1, amrex::Real &sedge2,
622 const amrex::Array4<const amrex::Real> &s,
623 const amrex::Real vel_dlo, const amrex::Real vel_dhi,
624 const int bclo, const int bchi,
625 const int domlo, const int domhi)
626{
627 using namespace amrex;
628
629 // Both the domlo and the domlo+1 branch below use s(domlo-1) as a face value.
630 // For direction_dependent that only holds when the domain face is inflow, so the
631 // gate is the domain-boundary velocity vel_dlo, not the velocity of the face
632 // bounding the cell being filled.
633 if ( (bclo == BCType::ext_dir) || (bclo == BCType::hoextrap) ||
634 (bclo == BCType::direction_dependent && vel_dlo >= Real(0)) )
635 {
636 if (k == domlo)
637 {
638 sedge2 = -amrex::Real(0.2)*s(i,j,domlo-1,n) + amrex::Real(0.75)*s(i,j,domlo ,n)
639 +amrex::Real(0.5)*s(i,j,domlo+1,n) - amrex::Real(0.05)*s(i,j,domlo+2,n);
640 if constexpr (Limiter::do_limiting) {
641 sedge2 = amrex::max(sedge2, amrex::min(s(i,j,domlo+1,n), s(i,j,domlo,n)));
642 sedge2 = amrex::min(sedge2, amrex::max(s(i,j,domlo+1,n), s(i,j,domlo,n)));
643 }
644
645 sm = s(i,j,domlo-1,n);
646 sp = sedge2;
647
648 } else if (k == domlo+1) {
649
650 sedge1 = -amrex::Real(0.2)*s(i,j,domlo-1,n) + amrex::Real(0.75)*s(i,j,domlo ,n)
651 +amrex::Real(0.5)*s(i,j,domlo+1,n) - amrex::Real(0.05)*s(i,j,domlo+2,n);
652 if constexpr (Limiter::do_limiting) {
653 sedge1 = amrex::max(sedge1, amrex::min(s(i,j,domlo+1,n), s(i,j,domlo,n)));
654 sedge1 = amrex::min(sedge1, amrex::max(s(i,j,domlo+1,n), s(i,j,domlo,n)));
655 }
656
657 sp = sedge2;
658 sm = sedge1;
659
660 if constexpr (Limiter::do_limiting) {
661 if ( (sp - s(i,j,domlo+1,n))*(s(i,j,domlo+1,n) - sm) <= amrex::Real(0.) )
662 {
663 sp = s(i,j,domlo+1,n);
664 sm = s(i,j,domlo+1,n);
665 } else if(amrex::Math::abs(sp - s(i,j,domlo+1,n)) >= amrex::Real(2.)*amrex::Math::abs(sm - s(i,j,domlo+1,n))) {
666 sp = amrex::Real(3)*s(i,j,domlo+1,n) - amrex::Real(2)*sm;
667 } else if(amrex::Math::abs(sm - s(i,j,domlo+1,n)) >= amrex::Real(2.)*amrex::Math::abs(sp - s(i,j,domlo+1,n))) {
668 sm = amrex::Real(3)*s(i,j,domlo+1,n) - amrex::Real(2)*sp;
669 }
670 }
671 }
672 }
673
674 // Likewise the domhi and domhi-1 branches are gated on the hi domain face.
675 if ( (bchi == BCType::ext_dir) || (bchi == BCType::hoextrap) ||
676 (bchi == BCType::direction_dependent && vel_dhi <= Real(0)) )
677 {
678 if (k == domhi)
679 {
680 sedge1 = -amrex::Real(0.2)*s(i,j,domhi+1,n) + amrex::Real(0.75)*s(i,j,domhi ,n)
681 +amrex::Real(0.5)*s(i,j,domhi-1,n) - amrex::Real(0.05)*s(i,j,domhi-2,n);
682 if constexpr (Limiter::do_limiting) {
683 sedge1 = amrex::max(sedge1, amrex::min(s(i,j,domhi-1,n), s(i,j,domhi,n)));
684 sedge1 = amrex::min(sedge1, amrex::max(s(i,j,domhi-1,n), s(i,j,domhi,n)));
685 }
686
687 sp = s(i,j,domhi+1,n);
688 sm = sedge1;
689
690 } else if (k == domhi-1) {
691
692 sedge2 = -amrex::Real(0.2)*s(i,j,domhi+1,n) + amrex::Real(0.75)*s(i,j,domhi ,n)
693 +amrex::Real(0.5)*s(i,j,domhi-1,n) - amrex::Real(0.05)*s(i,j,domhi-2,n);
694 if constexpr (Limiter::do_limiting) {
695 sedge2 = amrex::max(sedge2, amrex::min(s(i,j,domhi-1,n), s(i,j,domhi,n)));
696 sedge2 = amrex::min(sedge2, amrex::max(s(i,j,domhi-1,n), s(i,j,domhi,n)));
697 }
698
699 sp = sedge2;
700 sm = sedge1;
701
702 if constexpr (Limiter::do_limiting) {
703 if ( (sp - s(i,j,domhi-1,n))*(s(i,j,domhi-1,n) - sm) <= amrex::Real(0.) )
704 {
705 sp = s(i,j,domhi-1,n);
706 sm = s(i,j,domhi-1,n);
707 } else if(amrex::Math::abs(sp - s(i,j,domhi-1,n)) >= amrex::Real(2.)*amrex::Math::abs(sm - s(i,j,domhi-1,n))) {
708 sp = amrex::Real(3)*s(i,j,domhi-1,n) - amrex::Real(2)*sm;
709 } else if(amrex::Math::abs(sm - s(i,j,domhi-1,n)) >= amrex::Real(2.)*amrex::Math::abs(sp - s(i,j,domhi-1,n))) {
710 sm = amrex::Real(3)*s(i,j,domhi-1,n) - amrex::Real(2)*sp;
711 }
712 }
713
714 }
715 }
716}
717#endif
718
719
720// Right now only ppm type 1 is supported on GPU
721// This version is called before the MAC projection, when we use the cell-centered velocity
722// for upwinding
723template <typename Limiter>
724AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
725void PredictVelOnXFace ( const int i, const int j, const int k, const int n,
726 const amrex::Real dtdx, const amrex::Real v_ad,
727 const amrex::Real vel_dlo, const amrex::Real vel_dhi,
728 const amrex::Array4<const amrex::Real> &S,
729 const amrex::Array4<amrex::Real> &Im,
730 const amrex::Array4<amrex::Real> &Ip,
731 const amrex::BCRec bc, const int domlo, const int domhi,
732 const Limiter& /*limiter*/)
733{
734 constexpr auto half{amrex::Real(0.5)};
735 constexpr auto one{amrex::Real(1)};
736 constexpr auto two3rds{amrex::Real(2.0)/amrex::Real(3.0)};
737
738 const amrex::Real sm2 = S(i-2,j,k,n);
739 const amrex::Real sm1 = S(i-1,j,k,n);
740 const amrex::Real s0 = S(i ,j,k,n);
741 const amrex::Real sp1 = S(i+1,j,k,n);
742 const amrex::Real sp2 = S(i+2,j,k,n);
743
744 amrex::Real sedge1 = Limiter::sedge1(sm2,sm1,s0,sp1,sp2);
745 amrex::Real sedge2 = Limiter::sedge2(sm2,sm1,s0,sp1,sp2);
746
747 auto [sm, sp] = Limiter::sm_sp(s0, sedge1, sedge2);
748
749 SetXBCs<Limiter>(i, j, k, n, sm, sp, sedge1, sedge2, S, vel_dlo, vel_dhi,
750 bc.lo(0), bc.hi(0), domlo, domhi);
751
752 const amrex::Real s6 = amrex::Real(6)*s0 - amrex::Real(3)*(sm + sp);
753
754 const amrex::Real sigma = amrex::Math::abs(v_ad)*dtdx;
755
756 if (v_ad > small_vel)
757 {
758 Ip(i,j,k,n) = sp - (half*sigma)*((sp-sm) - (one - two3rds*sigma)*s6);
759 Im(i,j,k,n) = S(i,j,k,n);
760 }
761 else if (v_ad < -small_vel)
762 {
763 Ip(i,j,k,n) = S(i,j,k,n);
764 Im(i,j,k,n) = sm + (half*sigma)*((sp-sm) + (one - two3rds*sigma)*s6);
765 } else
766 {
767 Ip(i,j,k,n) = S(i,j,k,n);
768 Im(i,j,k,n) = S(i,j,k,n);
769 }
770}
771
772template <typename Limiter>
773AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
774void PredictVelOnYFace ( const int i, const int j, const int k, const int n,
775 const amrex::Real dtdy, const amrex::Real v_ad,
776 const amrex::Real vel_dlo, const amrex::Real vel_dhi,
777 const amrex::Array4<const amrex::Real> &S,
778 const amrex::Array4<amrex::Real> &Im,
779 const amrex::Array4<amrex::Real> &Ip,
780 const amrex::BCRec bc, const int domlo, const int domhi,
781 const Limiter& /*limiter*/)
782{
783 constexpr auto half{amrex::Real(0.5)};
784 constexpr auto one{amrex::Real(1)};
785 constexpr auto two3rds{amrex::Real(2.0)/amrex::Real(3.0)};
786
787 const amrex::Real sm2 = S(i,j-2,k,n);
788 const amrex::Real sm1 = S(i,j-1,k,n);
789 const amrex::Real s0 = S(i,j ,k,n);
790 const amrex::Real sp1 = S(i,j+1,k,n);
791 const amrex::Real sp2 = S(i,j+2,k,n);
792
793 amrex::Real sedge1 = Limiter::sedge1(sm2,sm1,s0,sp1,sp2);
794 amrex::Real sedge2 = Limiter::sedge2(sm2,sm1,s0,sp1,sp2);
795
796 auto [sm, sp] = Limiter::sm_sp(s0, sedge1, sedge2);
797
798 SetYBCs<Limiter>(i, j, k, n, sm, sp, sedge1, sedge2, S, vel_dlo, vel_dhi,
799 bc.lo(1), bc.hi(1), domlo, domhi);
800
801 const amrex::Real s6 = amrex::Real(6)*s0 - amrex::Real(3)*(sm + sp);
802
803 const amrex::Real sigma = amrex::Math::abs(v_ad)*dtdy;
804
805 if (v_ad > small_vel)
806 {
807 Ip(i,j,k,n) = sp - (half*sigma)*((sp-sm) - (one - two3rds*sigma)*s6);
808 Im(i,j,k,n) = S(i,j,k,n);
809 }
810 else if (v_ad < -small_vel)
811 {
812 Ip(i,j,k,n) = S(i,j,k,n);
813 Im(i,j,k,n) = sm + (half*sigma)*((sp-sm) + (one - two3rds*sigma)*s6);
814 } else
815 {
816 Ip(i,j,k,n) = S(i,j,k,n);
817 Im(i,j,k,n) = S(i,j,k,n);
818 }
819}
820
821#if (AMREX_SPACEDIM==3)
822template <typename Limiter>
823AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
824void PredictVelOnZFace ( const int i, const int j, const int k, const int n,
825 const amrex::Real dtdz, const amrex::Real v_ad,
826 const amrex::Real vel_dlo, const amrex::Real vel_dhi,
827 const amrex::Array4<const amrex::Real> &S,
828 const amrex::Array4<amrex::Real> &Im,
829 const amrex::Array4<amrex::Real> &Ip,
830 const amrex::BCRec bc, const int domlo, const int domhi,
831 const Limiter& /*limiter*/)
832{
833 constexpr auto half{amrex::Real(0.5)};
834 constexpr auto one{amrex::Real(1)};
835 constexpr auto two3rds{amrex::Real(2.0)/amrex::Real(3.0)};
836
837 const amrex::Real sm2 = S(i,j,k-2,n);
838 const amrex::Real sm1 = S(i,j,k-1,n);
839 const amrex::Real s0 = S(i,j,k ,n);
840 const amrex::Real sp1 = S(i,j,k+1,n);
841 const amrex::Real sp2 = S(i,j,k+2,n);
842
843 amrex::Real sedge1 = Limiter::sedge1(sm2,sm1,s0,sp1,sp2);
844 amrex::Real sedge2 = Limiter::sedge2(sm2,sm1,s0,sp1,sp2);
845
846 auto [sm, sp] = Limiter::sm_sp(s0, sedge1, sedge2);
847
848 SetZBCs<Limiter>(i, j, k, n, sm, sp, sedge1, sedge2, S, vel_dlo, vel_dhi,
849 bc.lo(2), bc.hi(2), domlo, domhi);
850
851 const amrex::Real s6 = amrex::Real(6)*s0 - amrex::Real(3)*(sm + sp);
852
853 const amrex::Real sigma = amrex::Math::abs(v_ad)*dtdz;
854
855 if (v_ad > small_vel)
856 {
857 Ip(i,j,k,n) = sp - (half*sigma)*((sp-sm) - (one - two3rds*sigma)*s6);
858 Im(i,j,k,n) = S(i,j,k,n);
859 }
860 else if (v_ad < -small_vel)
861 {
862 Ip(i,j,k,n) = S(i,j,k,n);
863 Im(i,j,k,n) = sm + (half*sigma)*((sp-sm) + (one - two3rds*sigma)*s6);
864 } else
865 {
866 Ip(i,j,k,n) = S(i,j,k,n);
867 Im(i,j,k,n) = S(i,j,k,n);
868 }
869}
870
871#endif
872
873// Right now only ppm type 1 is supported on GPU
874// This version is called after the MAC projection, when we use the MAC-projected velocity
875// for upwinding
876template <typename Limiter>
877AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
878void PredictStateOnXFace ( const int i, const int j, const int k, const int n,
879 const amrex::Real dt, const amrex::Real dx,
880 amrex::Real& Im, amrex::Real& Ip,
881 const amrex::Array4<const amrex::Real> &S,
882 const amrex::Array4<const amrex::Real> &vel_edge,
883 const amrex::BCRec bc,
884 const int domlo, const int domhi,
885 const Limiter& /*limiter*/,
886 int /*unused*/)
887{
888 const amrex::Real sm1 = S(i-1,j,k,n);
889 const amrex::Real s0 = S(i ,j,k,n);
890 const amrex::Real sp1 = S(i+1,j,k,n);
891
892 constexpr auto half{amrex::Real(0.5)};
893 constexpr auto one{amrex::Real(1)};
894
895 const amrex::Real sigmap = amrex::Math::abs(vel_edge(i+1,j,k))*dt/dx;
896 const amrex::Real sigmam = amrex::Math::abs(vel_edge(i ,j,k))*dt/dx;
897
898 constexpr auto two3rds{amrex::Real(2.0)/amrex::Real(3.0)};
899
900 const amrex::Real sm2 = S(i-2,j,k,n);
901 const amrex::Real sp2 = S(i+2,j,k,n);
902
903 amrex::Real sedge1 = Limiter::sedge1(sm2,sm1,s0,sp1,sp2);
904 amrex::Real sedge2 = Limiter::sedge2(sm2,sm1,s0,sp1,sp2);
905
906 auto [sm, sp] = Limiter::sm_sp(s0, sedge1, sedge2);
907
908 // SetXBCs applies its boundary stencil to both domlo and domlo+1 (domhi-1 and
909 // domhi), so for direction_dependent the gate must be the domain face, not the
910 // face bounding this cell. vel_edge carries one ghost face, so the domain face is
911 // out of reach only for a box whose valid region starts at domlo+2 (ends at
912 // domhi-2); there we fall back to this cell's face.
913 const amrex::Real vel_dlo = (i <= domlo+1 && amrex::lbound(vel_edge).x <= domlo)
914 ? vel_edge(domlo ,j,k) : vel_edge(i,j,k);
915 const amrex::Real vel_dhi = (i >= domhi-1 && amrex::ubound(vel_edge).x >= domhi+1)
916 ? vel_edge(domhi+1,j,k) : vel_edge(i+1,j,k);
917 SetXBCs<Limiter>(i, j, k, n, sm, sp, sedge1, sedge2, S, vel_dlo, vel_dhi,
918 bc.lo(0), bc.hi(0), domlo, domhi);
919
920 const amrex::Real s6 = amrex::Real(6)*s0 - amrex::Real(3)*(sm + sp);
921
922 if (vel_edge(i+1,j,k) > small_vel) {
923 Ip = sp - (half*sigmap)*((sp - sm) - (one -two3rds*sigmap)*s6);
924 } else {
925 Ip = S(i,j,k,n);
926 }
927
928 if(vel_edge(i,j,k) < -small_vel) {
929 Im = sm + (half*sigmam)*((sp-sm) + (one - two3rds*sigmam)*s6);
930 } else {
931 Im = S(i,j,k,n);
932 }
933}
934
935template <>
936AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
937void PredictStateOnXFace<PPM::upwind> ( const int i, const int j, const int k, const int n,
938 const amrex::Real /*unused*/, const amrex::Real /*unused*/,
939 amrex::Real& Im, amrex::Real& Ip,
940 const amrex::Array4<const amrex::Real> &S,
941 const amrex::Array4<const amrex::Real> & /*unused*/,
942 const amrex::BCRec /*unused*/,
943 const int /*unused*/, const int /*unused*/,
944 const PPM::upwind& /*limiter*/,
945 int /*unused*/)
946{
947 Im = S(i,j,k,n);
948 Ip = S(i,j,k,n);
949}
950
951
952template <>
953AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
954void PredictStateOnXFace<PPM::minmod> ( const int i, const int j, const int k, const int n,
955 const amrex::Real dt, const amrex::Real dx,
956 amrex::Real& Im, amrex::Real& Ip,
957 const amrex::Array4<const amrex::Real> &S,
958 const amrex::Array4<const amrex::Real> &vel_edge,
959 const amrex::BCRec bc,
960 const int domlo, const int domhi,
961 const PPM::minmod& /*limiter*/,
962 int /*unused*/)
963{
964 const amrex::Real sm1 = S(i-1,j,k,n);
965 const amrex::Real s0 = S(i ,j,k,n);
966 const amrex::Real sp1 = S(i+1,j,k,n);
967
968 constexpr auto half{amrex::Real(0.5)};
969 constexpr auto one{amrex::Real(1.0)};
970
971 const amrex::Real sigmap = amrex::Math::abs(vel_edge(i+1,j,k))*dt/dx;
972 const amrex::Real sigmam = amrex::Math::abs(vel_edge(i ,j,k))*dt/dx;
973
974 auto dsp = amrex::Real(0.0);
975 auto dsm = amrex::Real(0.0);
976 minmod_fn(sm1, s0, sp1, dsm, dsp);
977 Godunov_minmod_bc_lo(i, sm1, s0, vel_edge(i ,j,k), dsm, bc.lo(0), domlo);
978 Godunov_minmod_bc_hi(i, s0, sp1, vel_edge(i+1,j,k), dsp, bc.hi(0), domhi);
979
980 if (vel_edge(i + 1, j, k) > small_vel) {
981 Ip = s0 + half * (one - sigmap) * dsp;
982 } else {
983 Ip = S(i,j,k,n);
984 }
985
986 if (vel_edge(i, j, k) < -small_vel) {
987 Im = s0 - half * (one - sigmam) * dsm;
988 } else {
989 Im = S(i,j,k,n);
990 }
991}
992
993template <typename Limiter>
994AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
995void PredictStateOnYFace ( const int i, const int j, const int k, const int n,
996 const amrex::Real dt, const amrex::Real dx,
997 amrex::Real& Im, amrex::Real& Ip,
998 const amrex::Array4<const amrex::Real> &S,
999 const amrex::Array4<const amrex::Real> &vel_edge,
1000 const amrex::BCRec bc,
1001 const int domlo, const int domhi,
1002 const Limiter& /*limiter*/,
1003 int /*unused*/)
1004{
1005 const amrex::Real sm1 = S(i,j-1,k,n);
1006 const amrex::Real s0 = S(i,j ,k,n);
1007 const amrex::Real sp1 = S(i,j+1,k,n);
1008
1009 constexpr auto half{amrex::Real(0.5)};
1010 constexpr auto one{amrex::Real(1.0)};
1011
1012 const amrex::Real sigmap = amrex::Math::abs(vel_edge(i,j+1,k))*dt/dx;
1013 const amrex::Real sigmam = amrex::Math::abs(vel_edge(i,j ,k))*dt/dx;
1014
1015 constexpr auto two3rds{amrex::Real(2.0)/amrex::Real(3.0)};
1016
1017 const amrex::Real sm2 = S(i,j-2,k,n);
1018 const amrex::Real sp2 = S(i,j+2,k,n);
1019
1020 amrex::Real sedge1 = Limiter::sedge1(sm2,sm1,s0,sp1,sp2);
1021 amrex::Real sedge2 = Limiter::sedge2(sm2,sm1,s0,sp1,sp2);
1022
1023 auto [sm, sp] = Limiter::sm_sp(s0, sedge1, sedge2);
1024
1025 // SetYBCs applies its boundary stencil to both domlo and domlo+1 (domhi-1 and
1026 // domhi), so for direction_dependent the gate must be the domain face, not the
1027 // face bounding this cell. vel_edge carries one ghost face, so the domain face is
1028 // out of reach only for a box whose valid region starts at domlo+2 (ends at
1029 // domhi-2); there we fall back to this cell's face.
1030 const amrex::Real vel_dlo = (j <= domlo+1 && amrex::lbound(vel_edge).y <= domlo)
1031 ? vel_edge(i,domlo ,k) : vel_edge(i,j,k);
1032 const amrex::Real vel_dhi = (j >= domhi-1 && amrex::ubound(vel_edge).y >= domhi+1)
1033 ? vel_edge(i,domhi+1,k) : vel_edge(i,j+1,k);
1034 SetYBCs<Limiter>(i, j, k, n, sm, sp, sedge1, sedge2, S, vel_dlo, vel_dhi,
1035 bc.lo(1), bc.hi(1), domlo, domhi);
1036
1037 const amrex::Real s6 = amrex::Real(6)*s0- amrex::Real(3)*(sm + sp);
1038
1039 if (vel_edge(i,j+1,k) > small_vel) {
1040 Ip = sp - (half*sigmap)*((sp - sm) - (one -two3rds*sigmap)*s6);
1041 } else {
1042 Ip = S(i,j,k,n);
1043 }
1044
1045 if (vel_edge(i,j,k) < -small_vel) {
1046 Im = sm + (half*sigmam)*((sp-sm) + (one - two3rds*sigmam)*s6);
1047 } else {
1048 Im = S(i,j,k,n);
1049 }
1050}
1051
1052
1053template <>
1054AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
1055void PredictStateOnYFace<PPM::upwind> ( const int i, const int j, const int k, const int n,
1056 const amrex::Real /*unused*/, const amrex::Real /*unused*/,
1057 amrex::Real& Im, amrex::Real& Ip,
1058 const amrex::Array4<const amrex::Real> &S,
1059 const amrex::Array4<const amrex::Real> & /*unused*/,
1060 const amrex::BCRec /*unused*/,
1061 const int /*unused*/, const int /*unused*/,
1062 const PPM::upwind& /*limiter*/,
1063 int /*unused*/)
1064{
1065 Im = S(i,j,k,n);
1066 Ip = S(i,j,k,n);
1067}
1068
1069
1070template <>
1071AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
1072void PredictStateOnYFace<PPM::minmod> ( const int i, const int j, const int k, const int n,
1073 const amrex::Real dt, const amrex::Real dx,
1074 amrex::Real& Im, amrex::Real& Ip,
1075 const amrex::Array4<const amrex::Real> &S,
1076 const amrex::Array4<const amrex::Real> &vel_edge,
1077 const amrex::BCRec bc,
1078 const int domlo, const int domhi,
1079 const PPM::minmod& /*limiter*/,
1080 int /*unused*/)
1081{
1082 const amrex::Real sm1 = S(i,j-1,k,n);
1083 const amrex::Real s0 = S(i,j ,k,n);
1084 const amrex::Real sp1 = S(i,j+1,k,n);
1085
1086 constexpr auto half{amrex::Real(0.5)};
1087 constexpr auto one{amrex::Real(1.0)};
1088
1089 const amrex::Real sigmap = amrex::Math::abs(vel_edge(i,j+1,k))*dt/dx;
1090 const amrex::Real sigmam = amrex::Math::abs(vel_edge(i,j ,k))*dt/dx;
1091
1092 auto dsp = amrex::Real(0.0);
1093 auto dsm = amrex::Real(0.0);
1094 minmod_fn(sm1, s0, sp1, dsm, dsp);
1095 Godunov_minmod_bc_lo(j, sm1, s0, vel_edge(i,j ,k), dsm, bc.lo(1), domlo);
1096 Godunov_minmod_bc_hi(j, s0, sp1, vel_edge(i,j+1,k), dsp, bc.hi(1), domhi);
1097
1098 if (vel_edge(i,j+1,k) > small_vel) {
1099 Ip = s0 + half * (one - sigmap) * dsp;
1100 } else {
1101 Ip = S(i,j,k,n);
1102 }
1103 if (vel_edge(i,j,k) < -small_vel) {
1104 Im = s0 - half * (one - sigmam) * dsm;
1105 } else {
1106 Im = S(i,j,k,n);
1107 }
1108}
1109
1110
1111#if (AMREX_SPACEDIM==3)
1112template <typename Limiter>
1113AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
1114void PredictStateOnZFace ( const int i, const int j, const int k, const int n,
1115 const amrex::Real dt, const amrex::Real dx,
1116 amrex::Real& Im, amrex::Real& Ip,
1117 const amrex::Array4<const amrex::Real> &S,
1118 const amrex::Array4<const amrex::Real> &vel_edge,
1119 const amrex::BCRec bc,
1120 const int domlo, const int domhi,
1121 const Limiter& /*limiter*/,
1122 int /*unused*/)
1123{
1124 const amrex::Real sm1 = S(i,j,k-1,n);
1125 const amrex::Real s0 = S(i,j,k ,n);
1126 const amrex::Real sp1 = S(i,j,k+1,n);
1127
1128 constexpr auto half{amrex::Real(0.5)};
1129
1130 const amrex::Real sigmap = amrex::Math::abs(vel_edge(i,j,k+1))*dt/dx;
1131 const amrex::Real sigmam = amrex::Math::abs(vel_edge(i,j,k ))*dt/dx;
1132
1133 constexpr auto two3rds{amrex::Real(2.0)/amrex::Real(3.0)};
1134
1135 const amrex::Real sm2 = S(i,j,k-2,n);
1136 const amrex::Real sp2 = S(i,j,k+2,n);
1137
1138 amrex::Real sedge1 = Limiter::sedge1(sm2,sm1,s0,sp1,sp2);
1139 amrex::Real sedge2 = Limiter::sedge2(sm2,sm1,s0,sp1,sp2);
1140
1141 auto [sm, sp] = Limiter::sm_sp(s0, sedge1, sedge2);
1142
1143 // SetZBCs applies its boundary stencil to both domlo and domlo+1 (domhi-1 and
1144 // domhi), so for direction_dependent the gate must be the domain face, not the
1145 // face bounding this cell. vel_edge carries one ghost face, so the domain face is
1146 // out of reach only for a box whose valid region starts at domlo+2 (ends at
1147 // domhi-2); there we fall back to this cell's face.
1148 const amrex::Real vel_dlo = (k <= domlo+1 && amrex::lbound(vel_edge).z <= domlo)
1149 ? vel_edge(i,j,domlo ) : vel_edge(i,j,k);
1150 const amrex::Real vel_dhi = (k >= domhi-1 && amrex::ubound(vel_edge).z >= domhi+1)
1151 ? vel_edge(i,j,domhi+1) : vel_edge(i,j,k+1);
1152 SetZBCs<Limiter>(i, j, k, n, sm, sp, sedge1, sedge2, S, vel_dlo, vel_dhi,
1153 bc.lo(2), bc.hi(2), domlo, domhi);
1154
1155 const amrex::Real s6 = amrex::Real(6)*s0- amrex::Real(3)*(sm + sp);
1156
1157 if(vel_edge(i,j,k+1) > small_vel) {
1158 Ip = sp - (half*sigmap)*((sp-sm) - (amrex::Real(1.0) -two3rds*sigmap)*s6);
1159 } else {
1160 Ip = S(i,j,k,n);
1161 }
1162
1163 if(vel_edge(i,j,k) < -small_vel) {
1164 Im = sm + (half*sigmam)*((sp-sm) + (amrex::Real(1.0) - two3rds*sigmam)*s6);
1165 } else {
1166 Im = S(i,j,k,n);
1167 }
1168}
1169
1170template <>
1171AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
1172void PredictStateOnZFace<PPM::upwind> ( const int i, const int j, const int k, const int n,
1173 const amrex::Real /*unused*/, const amrex::Real /*unused*/,
1174 amrex::Real& Im, amrex::Real& Ip,
1175 const amrex::Array4<const amrex::Real> &S,
1176 const amrex::Array4<const amrex::Real> & /*unused*/,
1177 const amrex::BCRec /*unused*/,
1178 const int /*unused*/, const int /*unused*/,
1179 const PPM::upwind& /*limiter*/,
1180 int /*unused*/)
1181{
1182 Im = S(i,j,k,n);
1183 Ip = S(i,j,k,n);
1184}
1185
1186template <>
1187AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
1188void PredictStateOnZFace<PPM::minmod> ( const int i, const int j, const int k, const int n,
1189 const amrex::Real dt, const amrex::Real dx,
1190 amrex::Real& Im, amrex::Real& Ip,
1191 const amrex::Array4<const amrex::Real> &S,
1192 const amrex::Array4<const amrex::Real> &vel_edge,
1193 const amrex::BCRec bc,
1194 const int domlo, const int domhi,
1195 const PPM::minmod& /*limiter*/,
1196 int /*unused*/)
1197{
1198 const amrex::Real sm1 = S(i,j,k-1,n);
1199 const amrex::Real s0 = S(i,j,k ,n);
1200 const amrex::Real sp1 = S(i,j,k+1,n);
1201
1202 constexpr auto half{amrex::Real(0.5)};
1203 constexpr auto one{amrex::Real(1.0)};
1204
1205 const amrex::Real sigmap = amrex::Math::abs(vel_edge(i,j,k+1))*dt/dx;
1206 const amrex::Real sigmam = amrex::Math::abs(vel_edge(i,j,k ))*dt/dx;
1207
1208 auto dsp = amrex::Real(0.0);
1209 auto dsm = amrex::Real(0.0);
1210 minmod_fn(sm1, s0, sp1, dsm, dsp);
1211 Godunov_minmod_bc_lo(k, sm1, s0, vel_edge(i,j,k ), dsm, bc.lo(2), domlo);
1212 Godunov_minmod_bc_hi(k, s0, sp1, vel_edge(i,j,k+1), dsp, bc.hi(2), domhi);
1213
1214 if (vel_edge(i, j, k + 1) > small_vel) {
1215 Ip = s0 + half * (one - sigmap) * dsp;
1216 } else {
1217 Ip = S(i,j,k,n);
1218 }
1219
1220 if (vel_edge(i, j, k) < -small_vel) {
1221 Im = s0 - half * (one - sigmam) * dsm;
1222 } else {
1223 Im = S(i,j,k,n);
1224 }
1225}
1226#endif
1227
1228template <typename Limiter>
1229void PredictVelOnFaces (amrex::Box const& bx,
1230 AMREX_D_DECL(amrex::Array4<amrex::Real> const& Imx,
1231 amrex::Array4<amrex::Real> const& Imy,
1232 amrex::Array4<amrex::Real> const& Imz),
1233 AMREX_D_DECL(amrex::Array4<amrex::Real> const& Ipx,
1234 amrex::Array4<amrex::Real> const& Ipy,
1235 amrex::Array4<amrex::Real> const& Ipz),
1236 amrex::Array4<amrex::Real const> const& q,
1237 amrex::Array4<amrex::Real const> const& vel,
1238 amrex::Geometry geom,
1239 amrex::Real dt,
1240 amrex::BCRec const* pbc,
1241 const Limiter& limiter,
1242 amrex::Array4<int const> const& bc_arr = {})
1243{
1244 const amrex::Box& domain = geom.Domain();
1245 const amrex::Dim3 dlo = amrex::lbound(domain);
1246 const amrex::Dim3 dhi = amrex::ubound(domain);
1247
1248 const auto dx = geom.CellSizeArray();
1249 AMREX_D_TERM( const amrex::Real l_dtdx = dt / dx[0];,
1250 const amrex::Real l_dtdy = dt / dx[1];,
1251 const amrex::Real l_dtdz = dt / dx[2];);
1252
1253 amrex::ParallelFor(bx, AMREX_SPACEDIM,
1254 [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept
1255 {
1256 // The boundary stencils in SetXBCs/SetYBCs/SetZBCs cover the first two cell
1257 // layers, so for direction_dependent they must be gated on the velocity at the
1258 // domain boundary rather than on the velocity of the cell being filled. Here
1259 // that is the normal velocity of the boundary cell itself. The reads are
1260 // guarded so we never touch the boundary cell from a box that is nowhere near
1261 // it; vel carries 3 ghost cells, so it is in reach wherever the gate is used.
1262 const amrex::Real xlo_v = (i <= dlo.x+1) ? vel(dlo.x,j,k,0) : amrex::Real(0.0);
1263 const amrex::Real xhi_v = (i >= dhi.x-1) ? vel(dhi.x,j,k,0) : amrex::Real(0.0);
1264 const amrex::Real ylo_v = (j <= dlo.y+1) ? vel(i,dlo.y,k,1) : amrex::Real(0.0);
1265 const amrex::Real yhi_v = (j >= dhi.y-1) ? vel(i,dhi.y,k,1) : amrex::Real(0.0);
1266 // Use the position-dependent boundary condition where one was given, as the
1267 // PLM path does; getBC falls back to pbc[n] when bc_arr is null.
1268 const amrex::BCRec bc = HydroBC::getBC(i, j, k, n, domain, pbc, bc_arr);
1269 PredictVelOnXFace(i,j,k,n,l_dtdx,vel(i,j,k,0),xlo_v,xhi_v,q,Imx,Ipx,bc,dlo.x,dhi.x,limiter);
1270 PredictVelOnYFace(i,j,k,n,l_dtdy,vel(i,j,k,1),ylo_v,yhi_v,q,Imy,Ipy,bc,dlo.y,dhi.y,limiter);
1271#if (AMREX_SPACEDIM==3)
1272 const amrex::Real zlo_v = (k <= dlo.z+1) ? vel(i,j,dlo.z,2) : amrex::Real(0.0);
1273 const amrex::Real zhi_v = (k >= dhi.z-1) ? vel(i,j,dhi.z,2) : amrex::Real(0.0);
1274 PredictVelOnZFace(i,j,k,n,l_dtdz,vel(i,j,k,2),zlo_v,zhi_v,q,Imz,Ipz,bc,dlo.z,dhi.z,limiter);
1275#endif
1276 });
1277}
1278
1279template <typename Limiter>
1280void PredictStateOnFaces (amrex::Box const& bx,
1281 AMREX_D_DECL(amrex::Array4<amrex::Real> const& Imx,
1282 amrex::Array4<amrex::Real> const& Imy,
1283 amrex::Array4<amrex::Real> const& Imz),
1284 AMREX_D_DECL(amrex::Array4<amrex::Real> const& Ipx,
1285 amrex::Array4<amrex::Real> const& Ipy,
1286 amrex::Array4<amrex::Real> const& Ipz),
1287 AMREX_D_DECL(amrex::Array4<amrex::Real const> const& umac,
1288 amrex::Array4<amrex::Real const> const& vmac,
1289 amrex::Array4<amrex::Real const> const& wmac),
1290 amrex::Array4<amrex::Real const> const& q,
1291 amrex::Geometry geom,
1292 amrex::Real l_dt,
1293 amrex::BCRec const* pbc,
1294 const int ncomp,
1295 const Limiter& limiter,
1296 int limiter_type,
1297 amrex::Array4<int const> const& bc_arr = {})
1298{
1299 const amrex::Box& domain = geom.Domain();
1300 const amrex::Dim3 dlo = amrex::lbound(domain);
1301 const amrex::Dim3 dhi = amrex::ubound(domain);
1302
1303 AMREX_D_TERM (const auto dx = geom.CellSize(0);,
1304 const auto dy = geom.CellSize(1);,
1305 const auto dz = geom.CellSize(2););
1306
1307 amrex::ParallelFor(bx, ncomp,
1308 [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept
1309 {
1310 // Use the position-dependent boundary condition where one was given, as the
1311 // PLM path does; getBC falls back to pbc[n] when bc_arr is null.
1312 const amrex::BCRec bc = HydroBC::getBC(i, j, k, n, domain, pbc, bc_arr);
1313 PPM::PredictStateOnXFace(i, j, k, n, l_dt, dx, Imx(i,j,k,n), Ipx(i,j,k,n),
1314 q, umac, bc, dlo.x, dhi.x, limiter, limiter_type);
1315 PPM::PredictStateOnYFace(i, j, k, n, l_dt, dy, Imy(i,j,k,n), Ipy(i,j,k,n),
1316 q, vmac, bc, dlo.y, dhi.y, limiter, limiter_type);
1317#if (AMREX_SPACEDIM==3)
1318 PPM::PredictStateOnZFace(i, j, k, n, l_dt, dz, Imz(i,j,k,n), Ipz(i,j,k,n),
1319 q, wmac, bc, dlo.z, dhi.z, limiter, limiter_type);
1320#endif
1321 });
1322}
1323} // namespace
1324
1325#endif
1326/** @} */
AMReX_BCRec.H
AMReX_BC_TYPES.H
AMREX_FORCE_INLINE
#define AMREX_FORCE_INLINE
AMREX_GPU_HOST_DEVICE
#define AMREX_GPU_HOST_DEVICE
AMReX_MultiFab.H
AMREX_D_TERM
#define AMREX_D_TERM(a, b, c)
AMREX_D_DECL
#define AMREX_D_DECL(a, b, c)
amrex::BCRec::hi
__host__ __device__ const int * hi() const &noexcept
amrex::BCRec::lo
__host__ __device__ const int * lo() const &noexcept
amrex::CoordSys::CellSize
const Real * CellSize() const noexcept
amrex::CoordSys::CellSizeArray
GpuArray< Real, 3 > CellSizeArray() const noexcept
amrex::Geometry::Domain
const Box & Domain() const noexcept
eps
constexpr amrex::Real eps
Definition hydro_bds_edge_state_2D.cpp:17
small_vel
static constexpr amrex::Real small_vel
Definition hydro_constants.H:37
amrex::Real
amrex_real Real
amrex::ubound
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
amrex::lbound
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
amrex::min
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
amrex::max
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
hydro_bcs_K.H
hydro_constants.H
Godunov_minmod_bc_lo
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void Godunov_minmod_bc_lo(const int n, const amrex::Real sm1, const amrex::Real s0, const amrex::Real vel_edge, amrex::Real &dsm, const int bclo, const int domlo)
Definition hydro_godunov_ppm.H:45
minmod_fn
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void minmod_fn(const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, amrex::Real &dsm, amrex::Real &dsp)
Definition hydro_godunov_ppm.H:19
Godunov_minmod_bc_hi
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void Godunov_minmod_bc_hi(const int n, const amrex::Real s0, const amrex::Real sp1, const amrex::Real vel_edge, amrex::Real &dsp, const int bchi, const int domhi)
Definition hydro_godunov_ppm.H:65
HydroBC::getBC
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE amrex::BCRec getBC(const int i, const int j, const int k, const int n, const amrex::Box &m_domain, const amrex::BCRec *bcr, amrex::Array4< int const > const &bca)
Definition hydro_bcs_K.H:34
PPM
Definition hydro_godunov_ppm.H:83
PPM::PredictVelOnFaces
void PredictVelOnFaces(amrex::Box const &bx, AMREX_D_DECL(amrex::Array4< amrex::Real > const &Imx, amrex::Array4< amrex::Real > const &Imy, amrex::Array4< amrex::Real > const &Imz), AMREX_D_DECL(amrex::Array4< amrex::Real > const &Ipx, amrex::Array4< amrex::Real > const &Ipy, amrex::Array4< amrex::Real > const &Ipz), amrex::Array4< amrex::Real const > const &q, amrex::Array4< amrex::Real const > const &vel, amrex::Geometry geom, amrex::Real dt, amrex::BCRec const *pbc, const Limiter &limiter, amrex::Array4< int const > const &bc_arr={})
Definition hydro_godunov_ppm.H:1229
PPM::SetXBCs
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void SetXBCs(const int i, const int j, const int k, const int n, amrex::Real &sm, amrex::Real &sp, amrex::Real &sedge1, amrex::Real &sedge2, const amrex::Array4< const amrex::Real > &s, const amrex::Real vel_dlo, const amrex::Real vel_dhi, const int bclo, const int bchi, const int domlo, const int domhi)
Definition hydro_godunov_ppm.H:419
PPM::PredictStateOnXFace
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void PredictStateOnXFace(const int i, const int j, const int k, const int n, const amrex::Real dt, const amrex::Real dx, amrex::Real &Im, amrex::Real &Ip, const amrex::Array4< const amrex::Real > &S, const amrex::Array4< const amrex::Real > &vel_edge, const amrex::BCRec bc, const int domlo, const int domhi, const Limiter &, int)
Definition hydro_godunov_ppm.H:878
PPM::default_limiter
static constexpr int default_limiter
Definition hydro_godunov_ppm.H:87
PPM::PredictVelOnYFace
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void PredictVelOnYFace(const int i, const int j, const int k, const int n, const amrex::Real dtdy, const amrex::Real v_ad, const amrex::Real vel_dlo, const amrex::Real vel_dhi, const amrex::Array4< const amrex::Real > &S, const amrex::Array4< amrex::Real > &Im, const amrex::Array4< amrex::Real > &Ip, const amrex::BCRec bc, const int domlo, const int domhi, const Limiter &)
Definition hydro_godunov_ppm.H:774
PPM::PredictStateOnXFace< PPM::minmod >
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void PredictStateOnXFace< PPM::minmod >(const int i, const int j, const int k, const int n, const amrex::Real dt, const amrex::Real dx, amrex::Real &Im, amrex::Real &Ip, const amrex::Array4< const amrex::Real > &S, const amrex::Array4< const amrex::Real > &vel_edge, const amrex::BCRec bc, const int domlo, const int domhi, const PPM::minmod &, int)
Definition hydro_godunov_ppm.H:954
PPM::PredictStateOnXFace< PPM::upwind >
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void PredictStateOnXFace< PPM::upwind >(const int i, const int j, const int k, const int n, const amrex::Real, const amrex::Real, amrex::Real &Im, amrex::Real &Ip, const amrex::Array4< const amrex::Real > &S, const amrex::Array4< const amrex::Real > &, const amrex::BCRec, const int, const int, const PPM::upwind &, int)
Definition hydro_godunov_ppm.H:937
PPM::PredictStateOnFaces
void PredictStateOnFaces(amrex::Box const &bx, AMREX_D_DECL(amrex::Array4< amrex::Real > const &Imx, amrex::Array4< amrex::Real > const &Imy, amrex::Array4< amrex::Real > const &Imz), AMREX_D_DECL(amrex::Array4< amrex::Real > const &Ipx, amrex::Array4< amrex::Real > const &Ipy, amrex::Array4< amrex::Real > const &Ipz), AMREX_D_DECL(amrex::Array4< amrex::Real const > const &umac, amrex::Array4< amrex::Real const > const &vmac, amrex::Array4< amrex::Real const > const &wmac), amrex::Array4< amrex::Real const > const &q, amrex::Geometry geom, amrex::Real l_dt, amrex::BCRec const *pbc, const int ncomp, const Limiter &limiter, int limiter_type, amrex::Array4< int const > const &bc_arr={})
Definition hydro_godunov_ppm.H:1280
PPM::SetYBCs
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void SetYBCs(const int i, const int j, const int k, const int n, amrex::Real &sm, amrex::Real &sp, amrex::Real &sedge1, amrex::Real &sedge2, const amrex::Array4< const amrex::Real > &s, const amrex::Real vel_dlo, const amrex::Real vel_dhi, const int bclo, const int bchi, const int domlo, const int domhi)
Definition hydro_godunov_ppm.H:519
PPM::PredictStateOnYFace
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void PredictStateOnYFace(const int i, const int j, const int k, const int n, const amrex::Real dt, const amrex::Real dx, amrex::Real &Im, amrex::Real &Ip, const amrex::Array4< const amrex::Real > &S, const amrex::Array4< const amrex::Real > &vel_edge, const amrex::BCRec bc, const int domlo, const int domhi, const Limiter &, int)
Definition hydro_godunov_ppm.H:995
PPM::PredictStateOnYFace< PPM::minmod >
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void PredictStateOnYFace< PPM::minmod >(const int i, const int j, const int k, const int n, const amrex::Real dt, const amrex::Real dx, amrex::Real &Im, amrex::Real &Ip, const amrex::Array4< const amrex::Real > &S, const amrex::Array4< const amrex::Real > &vel_edge, const amrex::BCRec bc, const int domlo, const int domhi, const PPM::minmod &, int)
Definition hydro_godunov_ppm.H:1072
PPM::PredictStateOnYFace< PPM::upwind >
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void PredictStateOnYFace< PPM::upwind >(const int i, const int j, const int k, const int n, const amrex::Real, const amrex::Real, amrex::Real &Im, amrex::Real &Ip, const amrex::Array4< const amrex::Real > &S, const amrex::Array4< const amrex::Real > &, const amrex::BCRec, const int, const int, const PPM::upwind &, int)
Definition hydro_godunov_ppm.H:1055
PPM::PredictVelOnXFace
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void PredictVelOnXFace(const int i, const int j, const int k, const int n, const amrex::Real dtdx, const amrex::Real v_ad, const amrex::Real vel_dlo, const amrex::Real vel_dhi, const amrex::Array4< const amrex::Real > &S, const amrex::Array4< amrex::Real > &Im, const amrex::Array4< amrex::Real > &Ip, const amrex::BCRec bc, const int domlo, const int domhi, const Limiter &)
Definition hydro_godunov_ppm.H:725
PPM::limiters
limiters
Definition hydro_godunov_ppm.H:85
PPM::MINMOD
@ MINMOD
Definition hydro_godunov_ppm.H:85
PPM::VanLeer
@ VanLeer
Definition hydro_godunov_ppm.H:85
PPM::WENOZ
@ WENOZ
Definition hydro_godunov_ppm.H:85
PPM::WENO_JS
@ WENO_JS
Definition hydro_godunov_ppm.H:85
PPM::NoLimiter
@ NoLimiter
Definition hydro_godunov_ppm.H:85
PPM::UPWIND
@ UPWIND
Definition hydro_godunov_ppm.H:85
amrex::makeTuple
__host__ __device__ constexpr GpuTuple< detail::tuple_decay_t< Ts >... > makeTuple(Ts &&... args)
amrex::ParallelFor
void ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
PPM::minmod
Definition hydro_godunov_ppm.H:381
PPM::minmod::sm_sp
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::GpuTuple< amrex::Real, amrex::Real > sm_sp(const amrex::Real, const amrex::Real sedge1, const amrex::Real sedge2)
Definition hydro_godunov_ppm.H:409
PPM::minmod::sedge1
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge1(const amrex::Real, const amrex::Real, const amrex::Real, const amrex::Real, const amrex::Real)
Definition hydro_godunov_ppm.H:387
PPM::minmod::sedge2
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge2(const amrex::Real, const amrex::Real, const amrex::Real, const amrex::Real, const amrex::Real)
Definition hydro_godunov_ppm.H:398
PPM::minmod::do_limiting
static constexpr bool do_limiting
Definition hydro_godunov_ppm.H:383
PPM::nolimiter
Definition hydro_godunov_ppm.H:89
PPM::nolimiter::sm_sp
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::GpuTuple< amrex::Real, amrex::Real > sm_sp(const amrex::Real, const amrex::Real sedge1, const amrex::Real sedge2)
Definition hydro_godunov_ppm.H:126
PPM::nolimiter::sedge2
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge2(const amrex::Real, const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, const amrex::Real sp2)
Definition hydro_godunov_ppm.H:111
PPM::nolimiter::sedge1
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge1(const amrex::Real sm2, const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, const amrex::Real)
Definition hydro_godunov_ppm.H:95
PPM::nolimiter::do_limiting
static constexpr bool do_limiting
Definition hydro_godunov_ppm.H:91
PPM::upwind
Definition hydro_godunov_ppm.H:345
PPM::upwind::sedge1
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge1(const amrex::Real, const amrex::Real, const amrex::Real, const amrex::Real, const amrex::Real)
Definition hydro_godunov_ppm.H:351
PPM::upwind::do_limiting
static constexpr bool do_limiting
Definition hydro_godunov_ppm.H:347
PPM::upwind::sm_sp
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::GpuTuple< amrex::Real, amrex::Real > sm_sp(const amrex::Real, const amrex::Real sedge1, const amrex::Real sedge2)
Definition hydro_godunov_ppm.H:373
PPM::upwind::sedge2
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge2(const amrex::Real, const amrex::Real, const amrex::Real, const amrex::Real, const amrex::Real)
Definition hydro_godunov_ppm.H:362
PPM::vanleer
Definition hydro_godunov_ppm.H:134
PPM::vanleer::sm_sp
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::GpuTuple< amrex::Real, amrex::Real > sm_sp(const amrex::Real s0, const amrex::Real sedge1, const amrex::Real sedge2)
Definition hydro_godunov_ppm.H:188
PPM::vanleer::do_limiting
static constexpr bool do_limiting
Definition hydro_godunov_ppm.H:136
PPM::vanleer::sedge2
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge2(const amrex::Real, const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, const amrex::Real sp2)
Definition hydro_godunov_ppm.H:170
PPM::vanleer::vanLeer
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real vanLeer(const amrex::Real a, const amrex::Real b, const amrex::Real c)
Definition hydro_godunov_ppm.H:139
PPM::vanleer::sedge1
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge1(const amrex::Real sm2, const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, const amrex::Real)
Definition hydro_godunov_ppm.H:153
PPM::weno_js
Definition hydro_godunov_ppm.H:282
PPM::weno_js::do_limiting
static constexpr bool do_limiting
Definition hydro_godunov_ppm.H:284
PPM::weno_js::sedge1
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge1(const amrex::Real sm2, const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, const amrex::Real sp2)
Definition hydro_godunov_ppm.H:326
PPM::weno_js::sm_sp
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::GpuTuple< amrex::Real, amrex::Real > sm_sp(const amrex::Real, const amrex::Real sedge1, const amrex::Real sedge2)
Definition hydro_godunov_ppm.H:337
PPM::weno_js::sedge2
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge2(const amrex::Real sm2, const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, const amrex::Real sp2)
Definition hydro_godunov_ppm.H:288
PPM::wenoz
Definition hydro_godunov_ppm.H:223
PPM::wenoz::do_limiting
static constexpr bool do_limiting
Definition hydro_godunov_ppm.H:225
PPM::wenoz::sedge2
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge2(const amrex::Real sm2, const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, const amrex::Real sp2)
Definition hydro_godunov_ppm.H:229
PPM::wenoz::sm_sp
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::GpuTuple< amrex::Real, amrex::Real > sm_sp(const amrex::Real, const amrex::Real sedge1, const amrex::Real sedge2)
Definition hydro_godunov_ppm.H:274
PPM::wenoz::sedge1
AMREX_GPU_HOST_DEVICE static AMREX_FORCE_INLINE amrex::Real sedge1(const amrex::Real sm2, const amrex::Real sm1, const amrex::Real s0, const amrex::Real sp1, const amrex::Real sp2)
Definition hydro_godunov_ppm.H:263