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>
19requires (IsGPUable<F>::value)
21Real
22IF_f (F const& f, GpuArray<Real,AMREX_SPACEDIM> const& p) noexcept
23{
24 return f(AMREX_D_DECL(p[0],p[1],p[2]));
25}
26
27template <class F>
28requires (!IsGPUable<F>::value)
30Real
32{
35 amrex::Error("EB2::GeometryShop: how did this happen?");
36 return 0.0;
37 ))
38 AMREX_IF_ON_HOST((return f({AMREX_D_DECL(p[0],p[1],p[2])});))
39}
40
41template <class F>
43Real
46 int rangedir, F const& f)
47{
48#ifdef AMREX_USE_FLOAT
49 const Real tol = 1.e-4_rt;
50 const Real EPS = 1.0e-6_rt;
51#else
52 const Real tol = 1.e-12;
53 const Real EPS = 3.0e-15;
54#endif
55 const int MAXITER = 100;
56
57 Real p, q, r, s;
58
61
62 Real fa = IF_f(f, aPt);
63 Real fb = IF_f(f, bPt);
64 Real c = bPt[rangedir];
65 Real fc = fb;
66
67 if (fb*fa > 0.0_rt) {
68// amrex::AllPrint() << "fa " << fa << " fb " << fb << "\n";
69 amrex::Error("BrentRootFinder. Root must be bracketed, but instead the supplied end points have the same sign.");
70 return 0.0_rt;
71 } else if (fa == 0.0_rt) {
72 return aPt[rangedir];
73 } else if (fb == 0.0_rt) {
74 return bPt[rangedir];
75 }
76
77 Real d = 0.0_rt, e = 0.0_rt;
78 int i;
79 for (i = 0; i < MAXITER; ++i)
80 {
81 if (fb*fc > 0)
82 {
83 // Rename a, b, c and adjust bounding interval d
84 c = aPt[rangedir];
85 fc = fa;
86 d = bPt[rangedir] - aPt[rangedir];
87 e = d;
88 }
89
90 if (std::abs(fc) < std::abs(fb))
91 {
92 aPt[rangedir] = bPt[rangedir];
93 bPt[rangedir] = c;
94 c = aPt[rangedir];
95 fa = fb;
96 fb = fc;
97 fc = fa;
98 }
99
100 // Convergence check
101 Real tol1 = 2.0_rt * EPS * std::abs(bPt[rangedir]) + 0.5_rt * tol;
102 Real xm = 0.5_rt * (c - bPt[rangedir]);
103
104 if (std::abs(xm) <= tol1 || fb == 0.0_rt)
105 {
106 break;
107 }
108
109 if (std::abs(e) >= tol1 && std::abs(fa) > std::abs(fb))
110 {
111 // Attempt inverse quadratic interpolation
112 s = fb / fa;
113 if (aPt[rangedir] == c)
114 {
115 p = 2.0_rt * xm * s;
116 q = 1.0_rt - s;
117 }
118 else
119 {
120 q = fa / fc;
121 r = fb / fc;
122 p = s * (2.0_rt * xm * q * (q-r) - (bPt[rangedir]-aPt[rangedir]) * (r-1.0_rt));
123 q = (q-1.0_rt) * (r-1.0_rt) * (s-1.0_rt);
124 }
125
126 // Check whether in bounds
127 if (p > 0) { q = -q; }
128
129 p = std::abs(p);
130
131 if (2.0_rt * p < amrex::min(3.0_rt*xm*q-std::abs(tol1*q), 1.0_rt*std::abs(e*q)))
132 {
133 // Accept interpolation
134 e = d;
135 d = p / q;
136 }
137 else
138 {
139 // Interpolation failed, use bisection
140 d = xm;
141 e = d;
142 }
143 }
144 else
145 {
146 // Bounds decreasing too slowly, use bisection
147 d = xm;
148 e = d;
149 }
150
151 // Move last best guess to a
152 aPt[rangedir] = bPt[rangedir];
153 fa = fb;
154
155 // Evaluate new trial root
156 if (std::abs(d) > tol1)
157 {
158 bPt[rangedir] = bPt[rangedir] + d;
159 }
160 else
161 {
162 if (xm < 0) { bPt[rangedir] = bPt[rangedir] - tol1; }
163 else { bPt[rangedir] = bPt[rangedir] + tol1; }
164 }
165
166 fb = IF_f(f, bPt);
167 }
168
169 if (i >= MAXITER)
170 {
171 amrex::Error("BrentRootFinder: exceeding maximum iterations.");
172 }
173
174 return bPt[rangedir];
175}
176
183template <class F, class R = int>
185{
186public:
187
188 static constexpr int in_fluid = -1;
189 static constexpr int on_boundary = 0;
190 static constexpr int in_body = 1;
191 //
192 static constexpr int allregular = -1;
193 static constexpr int mixedcells = 0;
194 static constexpr int allcovered = 1;
195
197
203 explicit GeometryShop (F f)
204 : m_f(std::move(f)), m_resource()
205 {}
206
214 : m_f(std::move(f)), m_resource(std::move(r))
215 {}
216
218 [[nodiscard]] F const& GetImpFunc () const& { return m_f; }
219 F&& GetImpFunc () && { return std::move(m_f); }
220
221 [[nodiscard]] int getBoxType_Cpu (const Box& bx, Geometry const& geom) const noexcept
222 {
223 const Real* problo = geom.ProbLo();
224 const Real* dx = geom.CellSize();
225 const auto& len3 = bx.length3d();
226 const int* blo = bx.loVect();
227 int nbody = 0, nzero = 0, nfluid = 0;
228 for (int k = 0; k < len3[2]; ++k) {
229 for (int j = 0; j < len3[1]; ++j) {
230 for (int i = 0; i < len3[0]; ++i) {
231 RealArray xyz {AMREX_D_DECL(problo[0]+(i+blo[0])*dx[0],
232 problo[1]+(j+blo[1])*dx[1],
233 problo[2]+(k+blo[2])*dx[2])};
234 Real v = m_f(xyz);
235 if (v == 0.0_rt) {
236 ++nzero;
237 } else if (v > 0.0_rt) {
238 ++nbody;
239 } else {
240 ++nfluid;
241 }
242 if (nbody > 0 && nfluid > 0) { return mixedcells; }
243 }
244 }
245 }
247
248 if (nbody == 0) {
249 return allregular;
250 } else if (nfluid == 0) {
251 return allcovered;
252 } else {
253 return mixedcells;
254 }
255 }
256
264 [[nodiscard]] int getBoxType (const Box& bx, const Geometry& geom, RunOn run_on) const noexcept
265 requires (IsGPUable<F>::value)
266 {
267 if (run_on == RunOn::Gpu && Gpu::inLaunchRegion())
268 {
269 const auto& problo = geom.ProbLoArray();
270 const auto& dx = geom.CellSizeArray();
271 auto const& f = m_f;
273 ReduceData<int,int,int> reduce_data(reduce_op);
274 using ReduceTuple = typename decltype(reduce_data)::Type;
275 reduce_op.eval(bx, reduce_data,
276 [=] AMREX_GPU_DEVICE (int i, int j, int k) -> ReduceTuple
277 {
279 Real v = f(AMREX_D_DECL(problo[0]+i*dx[0],
280 problo[1]+j*dx[1],
281 problo[2]+k*dx[2]));
282 int nbody = 0, nzero = 0, nfluid = 0;
283 if (v == 0.0_rt) {
284 ++nzero;
285 } else if (v > 0.0_rt) {
286 ++nbody;
287 } else {
288 ++nfluid;
289 }
290 return {nbody,nzero,nfluid};
291 });
292 ReduceTuple hv = reduce_data.value(reduce_op);
293 int nbody = amrex::get<0>(hv);
294 // int nzero = amrex::get<1>(hv);
295 int nfluid = amrex::get<2>(hv);
296
297 if (nbody == 0) {
298 return allregular;
299 } else if (nfluid == 0) {
300 return allcovered;
301 } else {
302 return mixedcells;
303 }
304 }
305 else
306 {
307 return getBoxType_Cpu(bx, geom);
308 }
309 }
310
312 [[nodiscard]] int getBoxType (const Box& bx, const Geometry& geom, RunOn) const noexcept
313 requires (!IsGPUable<F>::value)
314 {
315 return getBoxType_Cpu(bx, geom);
316 }
317
318 static constexpr bool isGPUable () noexcept
319 requires (IsGPUable<F>::value)
320 { return true; }
321
322 static constexpr bool isGPUable () noexcept
323 requires (!IsGPUable<F>::value)
324 { return false; }
325
334 void fillFab (BaseFab<Real>& levelset, const Geometry& geom, RunOn run_on,
335 Box const& bounding_box) const noexcept
336 requires (IsGPUable<F>::value)
337 {
338 const auto& problo = geom.ProbLoArray();
339 const auto& dx = geom.CellSizeArray();
340 const Box& bx = levelset.box();
341 const auto& a = levelset.array();
342 const auto& blo = amrex::lbound(bounding_box);
343 const auto& bhi = amrex::ubound(bounding_box);
344 const auto& f = m_f;
345 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, bx, i, j, k,
346 {
347 a(i,j,k) = f(AMREX_D_DECL(problo[0]+amrex::Clamp(i,blo.x,bhi.x)*dx[0],
348 problo[1]+amrex::Clamp(j,blo.y,bhi.y)*dx[1],
349 problo[2]+amrex::Clamp(k,blo.z,bhi.z)*dx[2]));
350 });
351 }
352
355 Box const& bounding_box) const noexcept
356 requires (!IsGPUable<F>::value)
357 {
358#ifdef AMREX_USE_GPU
359 if (!levelset.arena()->isHostAccessible()) {
360 const Box& bx = levelset.box();
361 BaseFab<Real> h_levelset(bx, levelset.nComp(), The_Pinned_Arena());
362 fillFab_Cpu(h_levelset, geom, bounding_box);
363 Gpu::htod_memcpy_async(levelset.dataPtr(), h_levelset.dataPtr(),
364 levelset.nBytes(bx, levelset.nComp()));
366 }
367 else
368#endif
369 {
370 fillFab_Cpu(levelset, geom, bounding_box);
371 }
372 }
373
375 Box const& bounding_box) const noexcept
376 {
377 const auto& problo = geom.ProbLoArray();
378 const auto& dx = geom.CellSizeArray();
379 const Box& bx = levelset.box();
380 const auto& blo = amrex::lbound(bounding_box);
381 const auto& bhi = amrex::ubound(bounding_box);
382
383 const auto& a = levelset.array();
384 amrex::LoopOnCpu(bx, [&] (int i, int j, int k) noexcept
385 {
386 a(i,j,k) = m_f(RealArray{AMREX_D_DECL(problo[0]+amrex::Clamp(i,blo.x,bhi.x)*dx[0],
387 problo[1]+amrex::Clamp(j,blo.y,bhi.y)*dx[1],
388 problo[2]+amrex::Clamp(k,blo.z,bhi.z)*dx[2])});
389 });
390 }
391
392
402 void getIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
403 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
404 Array4<Real const> const&, Geometry const& geom, RunOn run_on,
405 Box const& bounding_box) const noexcept
406 requires (IsGPUable<F>::value)
407 {
408 auto const& dx = geom.CellSizeArray();
409 auto const& problo = geom.ProbLoArray();
410 const auto& blo = amrex::lbound(bounding_box);
411 const auto& bhi = amrex::ubound(bounding_box);
412 auto const& f = m_f;
413 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
414 Array4<Real> const& inter = inter_arr[idim];
415 Array4<Type_t const> const& type = type_arr[idim];
416 const Box bx{inter};
417 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, bx, i, j, k,
418 {
419 if (type(i,j,k) == Type::irregular) {
420 IntVect ivlo(AMREX_D_DECL(i,j,k));
421 IntVect ivhi(AMREX_D_DECL(i,j,k));
422 ivhi[idim] += 1;
423 inter(i,j,k) = BrentRootFinder
424 ({AMREX_D_DECL(problo[0]+amrex::Clamp(ivlo[0],blo.x,bhi.x)*dx[0],
425 problo[1]+amrex::Clamp(ivlo[1],blo.y,bhi.y)*dx[1],
426 problo[2]+amrex::Clamp(ivlo[2],blo.z,bhi.z)*dx[2])},
427 {AMREX_D_DECL(problo[0]+amrex::Clamp(ivhi[0],blo.x,bhi.x)*dx[0],
428 problo[1]+amrex::Clamp(ivhi[1],blo.y,bhi.y)*dx[1],
429 problo[2]+amrex::Clamp(ivhi[2],blo.z,bhi.z)*dx[2])},
430 idim, f);
431 } else {
432 inter(i,j,k) = std::numeric_limits<Real>::quiet_NaN();
433 }
434 });
435 }
436 }
437
439 void getIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
440 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
441 Array4<Real const> const&, Geometry const& geom, RunOn,
442 Box const& bounding_box) const noexcept
443 requires (!IsGPUable<F>::value)
444 {
445 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
446 Array4<Real> const& inter = inter_arr[idim];
447 Array4<Type_t const> const& type = type_arr[idim];
448 // When GShop is not GPUable, the intercept is either in managed memory or
449 // in device memory and needs to be copied to host and setup there.
450#ifdef AMREX_USE_GPU
451 if (!The_Arena()->isHostAccessible()) {
452 const Box bx{inter};
453 BaseFab<Real> h_inter_fab(bx, 1, The_Pinned_Arena());
454 BaseFab<Type_t> h_type_fab(bx, 1, The_Pinned_Arena());
455 Gpu::dtoh_memcpy_async(h_type_fab.dataPtr(), type.dataPtr(),
456 h_type_fab.nBytes(bx, 1));
458 const auto& h_inter = h_inter_fab.array();
459 const auto& h_type = h_type_fab.const_array();
460 getIntercept_Cpu(h_inter, h_type, geom, bounding_box, idim);
461 Gpu::htod_memcpy_async(inter.dataPtr(), h_inter.dataPtr(),
462 h_inter_fab.nBytes(bx, 1));
464 }
465 else
466#endif
467 {
468 getIntercept_Cpu(inter, type, geom, bounding_box, idim);
469 }
470
471 }
472 }
473
474 void getIntercept_Cpu (Array4<Real> const& inter,
475 Array4<Type_t const> const& type,
476 Geometry const& geom,
477 Box const& bounding_box,
478 const int idim) const noexcept
479 {
480 auto const& dx = geom.CellSizeArray();
481 auto const& problo = geom.ProbLoArray();
482 const auto& blo = amrex::lbound(bounding_box);
483 const auto& bhi = amrex::ubound(bounding_box);
484 const Box bx{inter};
485 amrex::LoopOnCpu(bx, [&] (int i, int j, int k) noexcept
486 {
487 if (type(i,j,k) == Type::irregular) {
488 IntVect ivlo(AMREX_D_DECL(i,j,k));
489 IntVect ivhi(AMREX_D_DECL(i,j,k));
490 ivhi[idim] += 1;
491 inter(i,j,k) = BrentRootFinder
492 ({AMREX_D_DECL(problo[0]+amrex::Clamp(ivlo[0],blo.x,bhi.x)*dx[0],
493 problo[1]+amrex::Clamp(ivlo[1],blo.y,bhi.y)*dx[1],
494 problo[2]+amrex::Clamp(ivlo[2],blo.z,bhi.z)*dx[2])},
495 {AMREX_D_DECL(problo[0]+amrex::Clamp(ivhi[0],blo.x,bhi.x)*dx[0],
496 problo[1]+amrex::Clamp(ivhi[1],blo.y,bhi.y)*dx[1],
497 problo[2]+amrex::Clamp(ivhi[2],blo.z,bhi.z)*dx[2])},
498 idim, m_f);
499 } else {
500 inter(i,j,k) = std::numeric_limits<Real>::quiet_NaN();
501 }
502 });
503 }
504
513 void updateIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
514 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
515 Array4<Real const> const& lst, Geometry const& geom) const noexcept
516 {
517 auto const& dx = geom.CellSizeArray();
518 auto const& problo = geom.ProbLoArray();
519 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
520 Array4<Real> const& inter = inter_arr[idim];
521 Array4<Type_t const> const& type = type_arr[idim];
522 const Box bx{inter};
524 {
525 if (type(i,j,k) == Type::irregular) {
526 bool is_nan = amrex::isnan(inter(i,j,k));
527 if (idim == 0) {
528 if (lst(i,j,k) == Real(0.0) ||
529 (lst(i,j,k) > Real(0.0) && is_nan))
530 {
531 // interp might still be quiet_nan because lst that
532 // was set to zero has been changed by FillBoundary
533 // at periodic boundaries.
534 inter(i,j,k) = problo[0] + i*dx[0];
535 }
536 else if (lst(i+1,j,k) == Real(0.0) ||
537 (lst(i+1,j,k) > Real(0.0) && is_nan))
538 {
539 inter(i,j,k) = problo[0] + (i+1)*dx[0];
540 }
541 } else if (idim == 1) {
542 if (lst(i,j,k) == Real(0.0) ||
543 (lst(i,j,k) > Real(0.0) && is_nan))
544 {
545 inter(i,j,k) = problo[1] + j*dx[1];
546 }
547 else if (lst(i,j+1,k) == Real(0.0) ||
548 (lst(i,j+1,k) > Real(0.0) && is_nan))
549 {
550 inter(i,j,k) = problo[1] + (j+1)*dx[1];
551 }
552 } else {
553 if (lst(i,j,k) == Real(0.0) ||
554 (lst(i,j,k) > Real(0.0) && is_nan))
555 {
556 inter(i,j,k) = problo[2] + k*dx[2];
557 }
558 else if (lst(i,j,k+1) == Real(0.0) ||
559 (lst(i,j,k+1) > Real(0.0) && is_nan))
560 {
561 inter(i,j,k) = problo[2] + (k+1)*dx[2];
562 }
563 }
564 }
565 });
566 }
567 }
568
569private:
570
571 F m_f;
572 R m_resource; // We use this to hold the ownership of resource for F if needed,
573 // because F needs to be a simply type suitable for GPU.
574};
575
581template <class F>
582GeometryShop<std::decay_t<F>>
584{
585 return GeometryShop<std::decay_t<F>>(std::forward<F>(f));
586}
587
595template <class F, class R>
596GeometryShop<std::decay_t<F>, std::decay_t<R>>
597makeShop (F&& f, R&& r)
598{
600 std::decay_t<R>>
601 (std::forward<F>(f), std::forward<R>(r));
602}
603
604}
605
606#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:104
#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
virtual bool isHostAccessible() const
Definition AMReX_Arena.cpp:76
A FortranArrayBox(FAB)-like object.
Definition AMReX_BaseFab.H:194
Array4< T const > const_array() const noexcept
Definition AMReX_BaseFab.H:423
std::size_t nBytes() const noexcept
Returns how many bytes used.
Definition AMReX_BaseFab.H:274
Array4< T const > array() const noexcept
Definition AMReX_BaseFab.H:387
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:360
Samples an implicit function on boxes and fills EB data.
Definition AMReX_EB2_GeometryShop.H:185
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:264
F const & GetImpFunc() const &
Immutable access to the underlying implicit function.
Definition AMReX_EB2_GeometryShop.H:218
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:354
GeometryShop(F f, R r)
Construct a GeometryShop with an implicit function and resource handle.
Definition AMReX_EB2_GeometryShop.H:213
static constexpr int in_body
Definition AMReX_EB2_GeometryShop.H:190
static constexpr bool isGPUable() noexcept
Definition AMReX_EB2_GeometryShop.H:318
static constexpr int allcovered
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:334
GeometryShop(F f)
Construct a GeometryShop that owns an implicit function f.
Definition AMReX_EB2_GeometryShop.H:203
static constexpr int in_fluid
Definition AMReX_EB2_GeometryShop.H:188
int getBoxType(const Box &bx, const Geometry &geom, RunOn) const noexcept
CPU-only fallback for getBoxType().
Definition AMReX_EB2_GeometryShop.H:312
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:402
void fillFab_Cpu(BaseFab< Real > &levelset, const Geometry &geom, Box const &bounding_box) const noexcept
Definition AMReX_EB2_GeometryShop.H:374
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:439
F FunctionType
Definition AMReX_EB2_GeometryShop.H:196
static constexpr int mixedcells
Definition AMReX_EB2_GeometryShop.H:193
static constexpr int allregular
Definition AMReX_EB2_GeometryShop.H:192
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:474
static constexpr int on_boundary
Definition AMReX_EB2_GeometryShop.H:189
static constexpr bool isGPUable() noexcept
Definition AMReX_EB2_GeometryShop.H:322
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:513
F && GetImpFunc() &&
Definition AMReX_EB2_GeometryShop.H:219
int getBoxType_Cpu(const Box &bx, Geometry const &geom) const noexcept
Definition AMReX_EB2_GeometryShop.H:221
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:75
Definition AMReX_Reduce.H:438
Type value()
Definition AMReX_Reduce.H:473
Definition AMReX_Reduce.H:597
void eval(MF const &mf, IntVect const &nghost, D &reduce_data, F &&f)
Definition AMReX_Reduce.H:731
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:1359
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Return the inclusive lower bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1345
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:44
__host__ __device__ Real IF_f(F const &f, GpuArray< Real, 3 > const &p) noexcept
Definition AMReX_EB2_GeometryShop.H:22
GeometryShop< std::decay_t< F > > makeShop(F &&f)
Helper that constructs a GeometryShop owning the supplied implicit function.
Definition AMReX_EB2_GeometryShop.H:583
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:88
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:65
__host__ __device__ constexpr const T & Clamp(const T &v, const T &lo, const T &hi)
Definition AMReX_Algorithm.H:107
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:25
void Error(const std::string &msg)
Print out message to cerr and exit via amrex::Abort().
Definition AMReX.cpp:235
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:285
__host__ __device__ constexpr T * dataPtr() const noexcept
Get raw data pointer.
Definition AMReX_Array4.H:613
Type trait that reports whether a functor derives from GPUable.
Definition AMReX_EB2_IF_Base.H:24
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:43