Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_EB2_GeometryShop.H
Go to the documentation of this file.
1#ifndef AMREX_EB2_GEOMETRYSHOP_H_
2#define AMREX_EB2_GEOMETRYSHOP_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_EB2_IF_Base.H>
6#include <AMReX_EB2_Graph.H>
7#include <AMReX_Geometry.H>
8#include <AMReX_BaseFab.H>
9#include <AMReX_Print.H>
10#include <AMReX_Array.H>
11#include <memory>
12#include <type_traits>
13#include <cmath>
14#include <utility>
15
16namespace amrex::EB2 {
17
18template <class F, std::enable_if_t<IsGPUable<F>::value>* FOO = nullptr>
20Real
21IF_f (F const& f, GpuArray<Real,AMREX_SPACEDIM> const& p) noexcept
22{
23 return f(AMREX_D_DECL(p[0],p[1],p[2]));
24}
25
26template <class F, std::enable_if_t<!IsGPUable<F>::value>* BAR = nullptr>
28Real
29IF_f (F const& f, GpuArray<Real,AMREX_SPACEDIM> const& p)
30{
33 amrex::Error("EB2::GeometryShop: how did this happen?");
34 return 0.0;
35 ))
36 AMREX_IF_ON_HOST((return f({AMREX_D_DECL(p[0],p[1],p[2])});))
37}
38
39template <class F>
41Real
44 int rangedir, F const& f)
45{
46#ifdef AMREX_USE_FLOAT
47 const Real tol = 1.e-4_rt;
48 const Real EPS = 1.0e-6_rt;
49#else
50 const Real tol = 1.e-12;
51 const Real EPS = 3.0e-15;
52#endif
53 const int MAXITER = 100;
54
55 Real p, q, r, s;
56
59
60 Real fa = IF_f(f, aPt);
61 Real fb = IF_f(f, bPt);
62 Real c = bPt[rangedir];
63 Real fc = fb;
64
65 if (fb*fa > 0.0_rt) {
66// amrex::AllPrint() << "fa " << fa << " fb " << fb << "\n";
67 amrex::Error("BrentRootFinder. Root must be bracketed, but instead the supplied end points have the same sign.");
68 return 0.0_rt;
69 } else if (fa == 0.0_rt) {
70 return aPt[rangedir];
71 } else if (fb == 0.0_rt) {
72 return bPt[rangedir];
73 }
74
75 Real d = 0.0_rt, e = 0.0_rt;
76 int i;
77 for (i = 0; i < MAXITER; ++i)
78 {
79 if (fb*fc > 0)
80 {
81 // Rename a, b, c and adjust bounding interval d
82 c = aPt[rangedir];
83 fc = fa;
84 d = bPt[rangedir] - aPt[rangedir];
85 e = d;
86 }
87
88 if (std::abs(fc) < std::abs(fb))
89 {
90 aPt[rangedir] = bPt[rangedir];
91 bPt[rangedir] = c;
92 c = aPt[rangedir];
93 fa = fb;
94 fb = fc;
95 fc = fa;
96 }
97
98 // Convergence check
99 Real tol1 = 2.0_rt * EPS * std::abs(bPt[rangedir]) + 0.5_rt * tol;
100 Real xm = 0.5_rt * (c - bPt[rangedir]);
101
102 if (std::abs(xm) <= tol1 || fb == 0.0_rt)
103 {
104 break;
105 }
106
107 if (std::abs(e) >= tol1 && std::abs(fa) > std::abs(fb))
108 {
109 // Attempt inverse quadratic interpolation
110 s = fb / fa;
111 if (aPt[rangedir] == c)
112 {
113 p = 2.0_rt * xm * s;
114 q = 1.0_rt - s;
115 }
116 else
117 {
118 q = fa / fc;
119 r = fb / fc;
120 p = s * (2.0_rt * xm * q * (q-r) - (bPt[rangedir]-aPt[rangedir]) * (r-1.0_rt));
121 q = (q-1.0_rt) * (r-1.0_rt) * (s-1.0_rt);
122 }
123
124 // Check whether in bounds
125 if (p > 0) { q = -q; }
126
127 p = std::abs(p);
128
129 if (2.0_rt * p < amrex::min(3.0_rt*xm*q-std::abs(tol1*q), 1.0_rt*std::abs(e*q)))
130 {
131 // Accept interpolation
132 e = d;
133 d = p / q;
134 }
135 else
136 {
137 // Interpolation failed, use bisection
138 d = xm;
139 e = d;
140 }
141 }
142 else
143 {
144 // Bounds decreasing too slowly, use bisection
145 d = xm;
146 e = d;
147 }
148
149 // Move last best guess to a
150 aPt[rangedir] = bPt[rangedir];
151 fa = fb;
152
153 // Evaluate new trial root
154 if (std::abs(d) > tol1)
155 {
156 bPt[rangedir] = bPt[rangedir] + d;
157 }
158 else
159 {
160 if (xm < 0) { bPt[rangedir] = bPt[rangedir] - tol1; }
161 else { bPt[rangedir] = bPt[rangedir] + tol1; }
162 }
163
164 fb = IF_f(f, bPt);
165 }
166
167 if (i >= MAXITER)
168 {
169 amrex::Error("BrentRootFinder: exceeding maximum iterations.");
170 }
171
172 return bPt[rangedir];
173}
174
181template <class F, class R = int>
183{
184public:
185
186 static constexpr int in_fluid = -1;
187 static constexpr int on_boundary = 0;
188 static constexpr int in_body = 1;
189 //
190 static constexpr int allregular = -1;
191 static constexpr int mixedcells = 0;
192 static constexpr int allcovered = 1;
193
195
201 explicit GeometryShop (F f)
202 : m_f(std::move(f)), m_resource()
203 {}
204
212 : m_f(std::move(f)), m_resource(std::move(r))
213 {}
214
216 [[nodiscard]] F const& GetImpFunc () const& { return m_f; }
217 F&& GetImpFunc () && { return std::move(m_f); }
218
219 [[nodiscard]] int getBoxType_Cpu (const Box& bx, Geometry const& geom) const noexcept
220 {
221 const Real* problo = geom.ProbLo();
222 const Real* dx = geom.CellSize();
223 const auto& len3 = bx.length3d();
224 const int* blo = bx.loVect();
225 int nbody = 0, nzero = 0, nfluid = 0;
226 for (int k = 0; k < len3[2]; ++k) {
227 for (int j = 0; j < len3[1]; ++j) {
228 for (int i = 0; i < len3[0]; ++i) {
229 RealArray xyz {AMREX_D_DECL(problo[0]+(i+blo[0])*dx[0],
230 problo[1]+(j+blo[1])*dx[1],
231 problo[2]+(k+blo[2])*dx[2])};
232 Real v = m_f(xyz);
233 if (v == 0.0_rt) {
234 ++nzero;
235 } else if (v > 0.0_rt) {
236 ++nbody;
237 } else {
238 ++nfluid;
239 }
240 if (nbody > 0 && nfluid > 0) { return mixedcells; }
241 }
242 }
243 }
245
246 if (nbody == 0) {
247 return allregular;
248 } else if (nfluid == 0) {
249 return allcovered;
250 } else {
251 return mixedcells;
252 }
253 }
254
262 template <class U=F, std::enable_if_t<IsGPUable<U>::value>* FOO = nullptr >
263 [[nodiscard]] int getBoxType (const Box& bx, const Geometry& geom, RunOn run_on) const noexcept
264 {
265 if (run_on == RunOn::Gpu && Gpu::inLaunchRegion())
266 {
267 const auto& problo = geom.ProbLoArray();
268 const auto& dx = geom.CellSizeArray();
269 auto const& f = m_f;
271 ReduceData<int,int,int> reduce_data(reduce_op);
272 using ReduceTuple = typename decltype(reduce_data)::Type;
273 reduce_op.eval(bx, reduce_data,
274 [=] AMREX_GPU_DEVICE (int i, int j, int k) -> ReduceTuple
275 {
277 Real v = f(AMREX_D_DECL(problo[0]+i*dx[0],
278 problo[1]+j*dx[1],
279 problo[2]+k*dx[2]));
280 int nbody = 0, nzero = 0, nfluid = 0;
281 if (v == 0.0_rt) {
282 ++nzero;
283 } else if (v > 0.0_rt) {
284 ++nbody;
285 } else {
286 ++nfluid;
287 }
288 return {nbody,nzero,nfluid};
289 });
290 ReduceTuple hv = reduce_data.value(reduce_op);
291 int nbody = amrex::get<0>(hv);
292 // int nzero = amrex::get<1>(hv);
293 int nfluid = amrex::get<2>(hv);
294
295 if (nbody == 0) {
296 return allregular;
297 } else if (nfluid == 0) {
298 return allcovered;
299 } else {
300 return mixedcells;
301 }
302 }
303 else
304 {
305 return getBoxType_Cpu(bx, geom);
306 }
307 }
308
310 template <class U=F, std::enable_if_t<!IsGPUable<U>::value>* BAR = nullptr >
311 [[nodiscard]] int getBoxType (const Box& bx, const Geometry& geom, RunOn) const noexcept
312 {
313 return getBoxType_Cpu(bx, geom);
314 }
315
316 template <class U=F, std::enable_if_t<IsGPUable<U>::value>* FOO = nullptr >
317 static constexpr bool isGPUable () noexcept { return true; }
318
319 template <class U=F, std::enable_if_t<!IsGPUable<U>::value>* BAR = nullptr >
320 static constexpr bool isGPUable () noexcept { return false; }
321
330 template <class U=F, std::enable_if_t<IsGPUable<U>::value>* FOO = nullptr >
331 void fillFab (BaseFab<Real>& levelset, const Geometry& geom, RunOn run_on,
332 Box const& bounding_box) const noexcept
333 {
334 const auto& problo = geom.ProbLoArray();
335 const auto& dx = geom.CellSizeArray();
336 const Box& bx = levelset.box();
337 const auto& a = levelset.array();
338 const auto& blo = amrex::lbound(bounding_box);
339 const auto& bhi = amrex::ubound(bounding_box);
340 const auto& f = m_f;
341 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, bx, i, j, k,
342 {
343 a(i,j,k) = f(AMREX_D_DECL(problo[0]+amrex::Clamp(i,blo.x,bhi.x)*dx[0],
344 problo[1]+amrex::Clamp(j,blo.y,bhi.y)*dx[1],
345 problo[2]+amrex::Clamp(k,blo.z,bhi.z)*dx[2]));
346 });
347 }
348
350 template <class U=F, std::enable_if_t<!IsGPUable<U>::value>* BAR = nullptr >
352 Box const& bounding_box) const noexcept
353 {
354#ifdef AMREX_USE_GPU
355 if (!levelset.arena()->isHostAccessible()) {
356 const Box& bx = levelset.box();
357 BaseFab<Real> h_levelset(bx, levelset.nComp(), The_Pinned_Arena());
358 fillFab_Cpu(h_levelset, geom, bounding_box);
359 Gpu::htod_memcpy_async(levelset.dataPtr(), h_levelset.dataPtr(),
360 levelset.nBytes(bx, levelset.nComp()));
362 }
363 else
364#endif
365 {
366 fillFab_Cpu(levelset, geom, bounding_box);
367 }
368 }
369
371 Box const& bounding_box) const noexcept
372 {
373 const auto& problo = geom.ProbLoArray();
374 const auto& dx = geom.CellSizeArray();
375 const Box& bx = levelset.box();
376 const auto& blo = amrex::lbound(bounding_box);
377 const auto& bhi = amrex::ubound(bounding_box);
378
379 const auto& a = levelset.array();
380 amrex::LoopOnCpu(bx, [&] (int i, int j, int k) noexcept
381 {
382 a(i,j,k) = m_f(RealArray{AMREX_D_DECL(problo[0]+amrex::Clamp(i,blo.x,bhi.x)*dx[0],
383 problo[1]+amrex::Clamp(j,blo.y,bhi.y)*dx[1],
384 problo[2]+amrex::Clamp(k,blo.z,bhi.z)*dx[2])});
385 });
386 }
387
388
398 template <class U=F, std::enable_if_t<IsGPUable<U>::value>* FOO = nullptr >
399 void getIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
400 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
401 Array4<Real const> const&, Geometry const& geom, RunOn run_on,
402 Box const& bounding_box) const noexcept
403 {
404 auto const& dx = geom.CellSizeArray();
405 auto const& problo = geom.ProbLoArray();
406 const auto& blo = amrex::lbound(bounding_box);
407 const auto& bhi = amrex::ubound(bounding_box);
408 auto const& f = m_f;
409 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
410 Array4<Real> const& inter = inter_arr[idim];
411 Array4<Type_t const> const& type = type_arr[idim];
412 const Box bx{inter};
413 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, bx, i, j, k,
414 {
415 if (type(i,j,k) == Type::irregular) {
416 IntVect ivlo(AMREX_D_DECL(i,j,k));
417 IntVect ivhi(AMREX_D_DECL(i,j,k));
418 ivhi[idim] += 1;
419 inter(i,j,k) = BrentRootFinder
420 ({AMREX_D_DECL(problo[0]+amrex::Clamp(ivlo[0],blo.x,bhi.x)*dx[0],
421 problo[1]+amrex::Clamp(ivlo[1],blo.y,bhi.y)*dx[1],
422 problo[2]+amrex::Clamp(ivlo[2],blo.z,bhi.z)*dx[2])},
423 {AMREX_D_DECL(problo[0]+amrex::Clamp(ivhi[0],blo.x,bhi.x)*dx[0],
424 problo[1]+amrex::Clamp(ivhi[1],blo.y,bhi.y)*dx[1],
425 problo[2]+amrex::Clamp(ivhi[2],blo.z,bhi.z)*dx[2])},
426 idim, f);
427 } else {
428 inter(i,j,k) = std::numeric_limits<Real>::quiet_NaN();
429 }
430 });
431 }
432 }
433
435 template <class U=F, std::enable_if_t<!IsGPUable<U>::value>* BAR = nullptr >
436 void getIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
437 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
438 Array4<Real const> const&, Geometry const& geom, RunOn,
439 Box const& bounding_box) const noexcept
440 {
441 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
442 Array4<Real> const& inter = inter_arr[idim];
443 Array4<Type_t const> const& type = type_arr[idim];
444 // When GShop is not GPUable, the intercept is either in managed memory or
445 // in device memory and needs to be copied to host and setup there.
446#ifdef AMREX_USE_GPU
447 if (!The_Arena()->isHostAccessible()) {
448 const Box bx{inter};
449 BaseFab<Real> h_inter_fab(bx, 1, The_Pinned_Arena());
450 BaseFab<Type_t> h_type_fab(bx, 1, The_Pinned_Arena());
451 Gpu::dtoh_memcpy_async(h_type_fab.dataPtr(), type.dataPtr(),
452 h_type_fab.nBytes(bx, 1));
454 const auto& h_inter = h_inter_fab.array();
455 const auto& h_type = h_type_fab.const_array();
456 getIntercept_Cpu(h_inter, h_type, geom, bounding_box, idim);
457 Gpu::htod_memcpy_async(inter.dataPtr(), h_inter.dataPtr(),
458 h_inter_fab.nBytes(bx, 1));
460 }
461 else
462#endif
463 {
464 getIntercept_Cpu(inter, type, geom, bounding_box, idim);
465 }
466
467 }
468 }
469
470 void getIntercept_Cpu (Array4<Real> const& inter,
471 Array4<Type_t const> const& type,
472 Geometry const& geom,
473 Box const& bounding_box,
474 const int idim) const noexcept
475 {
476 auto const& dx = geom.CellSizeArray();
477 auto const& problo = geom.ProbLoArray();
478 const auto& blo = amrex::lbound(bounding_box);
479 const auto& bhi = amrex::ubound(bounding_box);
480 const Box bx{inter};
481 amrex::LoopOnCpu(bx, [&] (int i, int j, int k) noexcept
482 {
483 if (type(i,j,k) == Type::irregular) {
484 IntVect ivlo(AMREX_D_DECL(i,j,k));
485 IntVect ivhi(AMREX_D_DECL(i,j,k));
486 ivhi[idim] += 1;
487 inter(i,j,k) = BrentRootFinder
488 ({AMREX_D_DECL(problo[0]+amrex::Clamp(ivlo[0],blo.x,bhi.x)*dx[0],
489 problo[1]+amrex::Clamp(ivlo[1],blo.y,bhi.y)*dx[1],
490 problo[2]+amrex::Clamp(ivlo[2],blo.z,bhi.z)*dx[2])},
491 {AMREX_D_DECL(problo[0]+amrex::Clamp(ivhi[0],blo.x,bhi.x)*dx[0],
492 problo[1]+amrex::Clamp(ivhi[1],blo.y,bhi.y)*dx[1],
493 problo[2]+amrex::Clamp(ivhi[2],blo.z,bhi.z)*dx[2])},
494 idim, m_f);
495 } else {
496 inter(i,j,k) = std::numeric_limits<Real>::quiet_NaN();
497 }
498 });
499 }
500
509 void updateIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
510 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
511 Array4<Real const> const& lst, Geometry const& geom) const noexcept
512 {
513 auto const& dx = geom.CellSizeArray();
514 auto const& problo = geom.ProbLoArray();
515 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
516 Array4<Real> const& inter = inter_arr[idim];
517 Array4<Type_t const> const& type = type_arr[idim];
518 const Box bx{inter};
520 {
521 if (type(i,j,k) == Type::irregular) {
522 bool is_nan = amrex::isnan(inter(i,j,k));
523 if (idim == 0) {
524 if (lst(i,j,k) == Real(0.0) ||
525 (lst(i,j,k) > Real(0.0) && is_nan))
526 {
527 // interp might still be quiet_nan because lst that
528 // was set to zero has been changed by FillBoundary
529 // at periodic boundaries.
530 inter(i,j,k) = problo[0] + i*dx[0];
531 }
532 else if (lst(i+1,j,k) == Real(0.0) ||
533 (lst(i+1,j,k) > Real(0.0) && is_nan))
534 {
535 inter(i,j,k) = problo[0] + (i+1)*dx[0];
536 }
537 } else if (idim == 1) {
538 if (lst(i,j,k) == Real(0.0) ||
539 (lst(i,j,k) > Real(0.0) && is_nan))
540 {
541 inter(i,j,k) = problo[1] + j*dx[1];
542 }
543 else if (lst(i,j+1,k) == Real(0.0) ||
544 (lst(i,j+1,k) > Real(0.0) && is_nan))
545 {
546 inter(i,j,k) = problo[1] + (j+1)*dx[1];
547 }
548 } else {
549 if (lst(i,j,k) == Real(0.0) ||
550 (lst(i,j,k) > Real(0.0) && is_nan))
551 {
552 inter(i,j,k) = problo[2] + k*dx[2];
553 }
554 else if (lst(i,j,k+1) == Real(0.0) ||
555 (lst(i,j,k+1) > Real(0.0) && is_nan))
556 {
557 inter(i,j,k) = problo[2] + (k+1)*dx[2];
558 }
559 }
560 }
561 });
562 }
563 }
564
565private:
566
567 F m_f;
568 R m_resource; // We use this to hold the ownership of resource for F if needed,
569 // because F needs to be a simply type suitable for GPU.
570};
571
577template <class F>
578GeometryShop<std::decay_t<F>>
580{
581 return GeometryShop<std::decay_t<F>>(std::forward<F>(f));
582}
583
591template <class F, class R>
592GeometryShop<std::decay_t<F>, std::decay_t<R>>
593makeShop (F&& f, R&& r)
594{
596 std::decay_t<R>>
597 (std::forward<F>(f), std::forward<R>(r));
598}
599
600}
601
602#endif
#define AMREX_HOST_DEVICE_PARALLEL_FOR_3D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:110
#define AMREX_HOST_DEVICE_FOR_3D_FLAG(where_to_run, box, i, j, k, block)
Definition AMReX_GpuLaunch.nolint.H:100
#define AMREX_IF_ON_DEVICE(CODE)
Definition AMReX_GpuQualifiers.H:56
#define AMREX_IF_ON_HOST(CODE)
Definition AMReX_GpuQualifiers.H:58
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
A FortranArrayBox(FAB)-like object.
Definition AMReX_BaseFab.H:190
Array4< T const > const_array() const noexcept
Definition AMReX_BaseFab.H:418
std::size_t nBytes() const noexcept
Returns how many bytes used.
Definition AMReX_BaseFab.H:269
Array4< T const > array() const noexcept
Definition AMReX_BaseFab.H:382
T * dataPtr(int n=0) noexcept
Returns a pointer to an object of type T that is the value of the Nth component associated with the c...
Definition AMReX_BaseFab.H:355
Samples an implicit function on boxes and fills EB data.
Definition AMReX_EB2_GeometryShop.H:183
F const & GetImpFunc() const &
Immutable access to the underlying implicit function.
Definition AMReX_EB2_GeometryShop.H:216
void getIntercept(Array< Array4< Real >, 3 > const &inter_arr, Array< Array4< Type_t const >, 3 > const &type_arr, Array4< Real const > const &, Geometry const &geom, RunOn, Box const &bounding_box) const noexcept
CPU-only overload of getIntercept().
Definition AMReX_EB2_GeometryShop.H:436
void fillFab(BaseFab< Real > &levelset, const Geometry &geom, RunOn, Box const &bounding_box) const noexcept
CPU-only overload of fillFab().
Definition AMReX_EB2_GeometryShop.H:351
static constexpr bool isGPUable() noexcept
Definition AMReX_EB2_GeometryShop.H:317
GeometryShop(F f, R r)
Construct a GeometryShop with an implicit function and resource handle.
Definition AMReX_EB2_GeometryShop.H:211
static constexpr int in_body
Definition AMReX_EB2_GeometryShop.H:188
static constexpr int allcovered
Definition AMReX_EB2_GeometryShop.H:192
int getBoxType(const Box &bx, const Geometry &geom, RunOn) const noexcept
CPU-only fallback for getBoxType().
Definition AMReX_EB2_GeometryShop.H:311
GeometryShop(F f)
Construct a GeometryShop that owns an implicit function f.
Definition AMReX_EB2_GeometryShop.H:201
static constexpr int in_fluid
Definition AMReX_EB2_GeometryShop.H:186
void fillFab_Cpu(BaseFab< Real > &levelset, const Geometry &geom, Box const &bounding_box) const noexcept
Definition AMReX_EB2_GeometryShop.H:370
F FunctionType
Definition AMReX_EB2_GeometryShop.H:194
void fillFab(BaseFab< Real > &levelset, const Geometry &geom, RunOn run_on, Box const &bounding_box) const noexcept
Fill a level-set FAB by sampling the implicit function on each point in the FAB.
Definition AMReX_EB2_GeometryShop.H:331
static constexpr int mixedcells
Definition AMReX_EB2_GeometryShop.H:191
static constexpr int allregular
Definition AMReX_EB2_GeometryShop.H:190
void getIntercept_Cpu(Array4< Real > const &inter, Array4< Type_t const > const &type, Geometry const &geom, Box const &bounding_box, const int idim) const noexcept
Definition AMReX_EB2_GeometryShop.H:470
void getIntercept(Array< Array4< Real >, 3 > const &inter_arr, Array< Array4< Type_t const >, 3 > const &type_arr, Array4< Real const > const &, Geometry const &geom, RunOn run_on, Box const &bounding_box) const noexcept
Compute EB intercepts for irregular edges (GPU-capable case).
Definition AMReX_EB2_GeometryShop.H:399
static constexpr int on_boundary
Definition AMReX_EB2_GeometryShop.H:187
void updateIntercept(Array< Array4< Real >, 3 > const &inter_arr, Array< Array4< Type_t const >, 3 > const &type_arr, Array4< Real const > const &lst, Geometry const &geom) const noexcept
Update intercepts when level-set data has been adjusted (e.g., fixing small cells).
Definition AMReX_EB2_GeometryShop.H:509
F && GetImpFunc() &&
Definition AMReX_EB2_GeometryShop.H:217
int getBoxType_Cpu(const Box &bx, Geometry const &geom) const noexcept
Definition AMReX_EB2_GeometryShop.H:219
int getBoxType(const Box &bx, const Geometry &geom, RunOn run_on) const noexcept
Classify bx as regular, covered, or mixed (GPU-capable case).
Definition AMReX_EB2_GeometryShop.H:263
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:74
Definition AMReX_Reduce.H:453
Type value()
Definition AMReX_Reduce.H:488
Definition AMReX_Reduce.H:612
std::enable_if_t< IsFabArray< MF >::value > eval(MF const &mf, IntVect const &nghost, D &reduce_data, F &&f)
Definition AMReX_Reduce.H:746
amrex_real Real
Floating Point Type for Fields.
Definition AMReX_REAL.H:79
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Return the inclusive upper bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1331
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Return the inclusive lower bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1317
std::array< T, N > Array
Definition AMReX_Array.H:26
Arena * The_Pinned_Arena()
Definition AMReX_Arena.cpp:860
Arena * The_Arena()
Definition AMReX_Arena.cpp:820
static constexpr Type_t irregular
Definition AMReX_EB2_Graph.H:51
Definition AMReX_FabArrayBase.H:33
__host__ __device__ Real BrentRootFinder(GpuArray< Real, 3 > const &lo, GpuArray< Real, 3 > const &hi, int rangedir, F const &f)
Definition AMReX_EB2_GeometryShop.H:42
__host__ __device__ Real IF_f(F const &f, GpuArray< Real, 3 > const &p) noexcept
Definition AMReX_EB2_GeometryShop.H:21
GeometryShop< std::decay_t< F > > makeShop(F &&f)
Helper that constructs a GeometryShop owning the supplied implicit function.
Definition AMReX_EB2_GeometryShop.H:579
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:310
void dtoh_memcpy_async(void *p_h, const void *p_d, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:435
bool inLaunchRegion() noexcept
Definition AMReX_GpuControl.H:92
void htod_memcpy_async(void *p_d, const void *p_h, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:421
__host__ __device__ void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:139
RunOn
Definition AMReX_GpuControl.H:69
__host__ __device__ constexpr const T & Clamp(const T &v, const T &lo, const T &hi)
Definition AMReX_Algorithm.H:106
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:24
void Error(const std::string &msg)
Print out message to cerr and exit via amrex::Abort().
Definition AMReX.cpp:234
void LoopOnCpu(Dim3 lo, Dim3 hi, F const &f) noexcept
Definition AMReX_Loop.H:365
Array< Real, 3 > RealArray
Definition AMReX_Array.H:28
A multidimensional array accessor.
Definition AMReX_Array4.H:283
__host__ __device__ constexpr T * dataPtr() const noexcept
Get raw data pointer.
Definition AMReX_Array4.H:612
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:43