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 template <class U=F>
265 requires (IsGPUable<U>::value)
266 [[nodiscard]] int getBoxType (const Box& bx, const Geometry& geom, RunOn run_on) const noexcept
267 {
268 if (run_on == RunOn::Gpu && Gpu::inLaunchRegion())
269 {
270 const auto& problo = geom.ProbLoArray();
271 const auto& dx = geom.CellSizeArray();
272 auto const& f = m_f;
274 ReduceData<int,int,int> reduce_data(reduce_op);
275 using ReduceTuple = typename decltype(reduce_data)::Type;
276 reduce_op.eval(bx, reduce_data,
277 [=] AMREX_GPU_DEVICE (int i, int j, int k) -> ReduceTuple
278 {
280 Real v = f(AMREX_D_DECL(problo[0]+i*dx[0],
281 problo[1]+j*dx[1],
282 problo[2]+k*dx[2]));
283 int nbody = 0, nzero = 0, nfluid = 0;
284 if (v == 0.0_rt) {
285 ++nzero;
286 } else if (v > 0.0_rt) {
287 ++nbody;
288 } else {
289 ++nfluid;
290 }
291 return {nbody,nzero,nfluid};
292 });
293 ReduceTuple hv = reduce_data.value(reduce_op);
294 int nbody = amrex::get<0>(hv);
295 // int nzero = amrex::get<1>(hv);
296 int nfluid = amrex::get<2>(hv);
297
298 if (nbody == 0) {
299 return allregular;
300 } else if (nfluid == 0) {
301 return allcovered;
302 } else {
303 return mixedcells;
304 }
305 }
306 else
307 {
308 return getBoxType_Cpu(bx, geom);
309 }
310 }
311
313 template <class U=F>
314 requires (!IsGPUable<U>::value)
315 [[nodiscard]] int getBoxType (const Box& bx, const Geometry& geom, RunOn) const noexcept
316 {
317 return getBoxType_Cpu(bx, geom);
318 }
319
320 template <class U=F>
321 requires (IsGPUable<U>::value)
322 static constexpr bool isGPUable () noexcept { return true; }
323
324 template <class U=F>
325 requires (!IsGPUable<U>::value)
326 static constexpr bool isGPUable () noexcept { return false; }
327
336 template <class U=F>
337 requires (IsGPUable<U>::value)
338 void fillFab (BaseFab<Real>& levelset, const Geometry& geom, RunOn run_on,
339 Box const& bounding_box) const noexcept
340 {
341 const auto& problo = geom.ProbLoArray();
342 const auto& dx = geom.CellSizeArray();
343 const Box& bx = levelset.box();
344 const auto& a = levelset.array();
345 const auto& blo = amrex::lbound(bounding_box);
346 const auto& bhi = amrex::ubound(bounding_box);
347 const auto& f = m_f;
348 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, bx, i, j, k,
349 {
350 a(i,j,k) = f(AMREX_D_DECL(problo[0]+amrex::Clamp(i,blo.x,bhi.x)*dx[0],
351 problo[1]+amrex::Clamp(j,blo.y,bhi.y)*dx[1],
352 problo[2]+amrex::Clamp(k,blo.z,bhi.z)*dx[2]));
353 });
354 }
355
357 template <class U=F>
358 requires (!IsGPUable<U>::value)
360 Box const& bounding_box) const noexcept
361 {
362#ifdef AMREX_USE_GPU
363 if (!levelset.arena()->isHostAccessible()) {
364 const Box& bx = levelset.box();
365 BaseFab<Real> h_levelset(bx, levelset.nComp(), The_Pinned_Arena());
366 fillFab_Cpu(h_levelset, geom, bounding_box);
367 Gpu::htod_memcpy_async(levelset.dataPtr(), h_levelset.dataPtr(),
368 levelset.nBytes(bx, levelset.nComp()));
370 }
371 else
372#endif
373 {
374 fillFab_Cpu(levelset, geom, bounding_box);
375 }
376 }
377
379 Box const& bounding_box) const noexcept
380 {
381 const auto& problo = geom.ProbLoArray();
382 const auto& dx = geom.CellSizeArray();
383 const Box& bx = levelset.box();
384 const auto& blo = amrex::lbound(bounding_box);
385 const auto& bhi = amrex::ubound(bounding_box);
386
387 const auto& a = levelset.array();
388 amrex::LoopOnCpu(bx, [&] (int i, int j, int k) noexcept
389 {
390 a(i,j,k) = m_f(RealArray{AMREX_D_DECL(problo[0]+amrex::Clamp(i,blo.x,bhi.x)*dx[0],
391 problo[1]+amrex::Clamp(j,blo.y,bhi.y)*dx[1],
392 problo[2]+amrex::Clamp(k,blo.z,bhi.z)*dx[2])});
393 });
394 }
395
396
406 template <class U=F>
407 requires (IsGPUable<U>::value)
408 void getIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
409 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
410 Array4<Real const> const&, Geometry const& geom, RunOn run_on,
411 Box const& bounding_box) const noexcept
412 {
413 auto const& dx = geom.CellSizeArray();
414 auto const& problo = geom.ProbLoArray();
415 const auto& blo = amrex::lbound(bounding_box);
416 const auto& bhi = amrex::ubound(bounding_box);
417 auto const& f = m_f;
418 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
419 Array4<Real> const& inter = inter_arr[idim];
420 Array4<Type_t const> const& type = type_arr[idim];
421 const Box bx{inter};
422 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, bx, i, j, k,
423 {
424 if (type(i,j,k) == Type::irregular) {
425 IntVect ivlo(AMREX_D_DECL(i,j,k));
426 IntVect ivhi(AMREX_D_DECL(i,j,k));
427 ivhi[idim] += 1;
428 inter(i,j,k) = BrentRootFinder
429 ({AMREX_D_DECL(problo[0]+amrex::Clamp(ivlo[0],blo.x,bhi.x)*dx[0],
430 problo[1]+amrex::Clamp(ivlo[1],blo.y,bhi.y)*dx[1],
431 problo[2]+amrex::Clamp(ivlo[2],blo.z,bhi.z)*dx[2])},
432 {AMREX_D_DECL(problo[0]+amrex::Clamp(ivhi[0],blo.x,bhi.x)*dx[0],
433 problo[1]+amrex::Clamp(ivhi[1],blo.y,bhi.y)*dx[1],
434 problo[2]+amrex::Clamp(ivhi[2],blo.z,bhi.z)*dx[2])},
435 idim, f);
436 } else {
437 inter(i,j,k) = std::numeric_limits<Real>::quiet_NaN();
438 }
439 });
440 }
441 }
442
444 template <class U=F>
445 requires (!IsGPUable<U>::value)
446 void getIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
447 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
448 Array4<Real const> const&, Geometry const& geom, RunOn,
449 Box const& bounding_box) const noexcept
450 {
451 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
452 Array4<Real> const& inter = inter_arr[idim];
453 Array4<Type_t const> const& type = type_arr[idim];
454 // When GShop is not GPUable, the intercept is either in managed memory or
455 // in device memory and needs to be copied to host and setup there.
456#ifdef AMREX_USE_GPU
457 if (!The_Arena()->isHostAccessible()) {
458 const Box bx{inter};
459 BaseFab<Real> h_inter_fab(bx, 1, The_Pinned_Arena());
460 BaseFab<Type_t> h_type_fab(bx, 1, The_Pinned_Arena());
461 Gpu::dtoh_memcpy_async(h_type_fab.dataPtr(), type.dataPtr(),
462 h_type_fab.nBytes(bx, 1));
464 const auto& h_inter = h_inter_fab.array();
465 const auto& h_type = h_type_fab.const_array();
466 getIntercept_Cpu(h_inter, h_type, geom, bounding_box, idim);
467 Gpu::htod_memcpy_async(inter.dataPtr(), h_inter.dataPtr(),
468 h_inter_fab.nBytes(bx, 1));
470 }
471 else
472#endif
473 {
474 getIntercept_Cpu(inter, type, geom, bounding_box, idim);
475 }
476
477 }
478 }
479
480 void getIntercept_Cpu (Array4<Real> const& inter,
481 Array4<Type_t const> const& type,
482 Geometry const& geom,
483 Box const& bounding_box,
484 const int idim) const noexcept
485 {
486 auto const& dx = geom.CellSizeArray();
487 auto const& problo = geom.ProbLoArray();
488 const auto& blo = amrex::lbound(bounding_box);
489 const auto& bhi = amrex::ubound(bounding_box);
490 const Box bx{inter};
491 amrex::LoopOnCpu(bx, [&] (int i, int j, int k) noexcept
492 {
493 if (type(i,j,k) == Type::irregular) {
494 IntVect ivlo(AMREX_D_DECL(i,j,k));
495 IntVect ivhi(AMREX_D_DECL(i,j,k));
496 ivhi[idim] += 1;
497 inter(i,j,k) = BrentRootFinder
498 ({AMREX_D_DECL(problo[0]+amrex::Clamp(ivlo[0],blo.x,bhi.x)*dx[0],
499 problo[1]+amrex::Clamp(ivlo[1],blo.y,bhi.y)*dx[1],
500 problo[2]+amrex::Clamp(ivlo[2],blo.z,bhi.z)*dx[2])},
501 {AMREX_D_DECL(problo[0]+amrex::Clamp(ivhi[0],blo.x,bhi.x)*dx[0],
502 problo[1]+amrex::Clamp(ivhi[1],blo.y,bhi.y)*dx[1],
503 problo[2]+amrex::Clamp(ivhi[2],blo.z,bhi.z)*dx[2])},
504 idim, m_f);
505 } else {
506 inter(i,j,k) = std::numeric_limits<Real>::quiet_NaN();
507 }
508 });
509 }
510
519 void updateIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
520 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
521 Array4<Real const> const& lst, Geometry const& geom) const noexcept
522 {
523 auto const& dx = geom.CellSizeArray();
524 auto const& problo = geom.ProbLoArray();
525 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
526 Array4<Real> const& inter = inter_arr[idim];
527 Array4<Type_t const> const& type = type_arr[idim];
528 const Box bx{inter};
530 {
531 if (type(i,j,k) == Type::irregular) {
532 bool is_nan = amrex::isnan(inter(i,j,k));
533 if (idim == 0) {
534 if (lst(i,j,k) == Real(0.0) ||
535 (lst(i,j,k) > Real(0.0) && is_nan))
536 {
537 // interp might still be quiet_nan because lst that
538 // was set to zero has been changed by FillBoundary
539 // at periodic boundaries.
540 inter(i,j,k) = problo[0] + i*dx[0];
541 }
542 else if (lst(i+1,j,k) == Real(0.0) ||
543 (lst(i+1,j,k) > Real(0.0) && is_nan))
544 {
545 inter(i,j,k) = problo[0] + (i+1)*dx[0];
546 }
547 } else if (idim == 1) {
548 if (lst(i,j,k) == Real(0.0) ||
549 (lst(i,j,k) > Real(0.0) && is_nan))
550 {
551 inter(i,j,k) = problo[1] + j*dx[1];
552 }
553 else if (lst(i,j+1,k) == Real(0.0) ||
554 (lst(i,j+1,k) > Real(0.0) && is_nan))
555 {
556 inter(i,j,k) = problo[1] + (j+1)*dx[1];
557 }
558 } else {
559 if (lst(i,j,k) == Real(0.0) ||
560 (lst(i,j,k) > Real(0.0) && is_nan))
561 {
562 inter(i,j,k) = problo[2] + k*dx[2];
563 }
564 else if (lst(i,j,k+1) == Real(0.0) ||
565 (lst(i,j,k+1) > Real(0.0) && is_nan))
566 {
567 inter(i,j,k) = problo[2] + (k+1)*dx[2];
568 }
569 }
570 }
571 });
572 }
573 }
574
575private:
576
577 F m_f;
578 R m_resource; // We use this to hold the ownership of resource for F if needed,
579 // because F needs to be a simply type suitable for GPU.
580};
581
587template <class F>
588GeometryShop<std::decay_t<F>>
590{
591 return GeometryShop<std::decay_t<F>>(std::forward<F>(f));
592}
593
601template <class F, class R>
602GeometryShop<std::decay_t<F>, std::decay_t<R>>
603makeShop (F&& f, R&& r)
604{
606 std::decay_t<R>>
607 (std::forward<F>(f), std::forward<R>(r));
608}
609
610}
611
612#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
GpuArray< Real, 3 > CellSizeArray() const noexcept
Definition AMReX_CoordSys.H:76
Samples an implicit function on boxes and fills EB data.
Definition AMReX_EB2_GeometryShop.H:185
F const & GetImpFunc() const &
Immutable access to the underlying implicit function.
Definition AMReX_EB2_GeometryShop.H:218
GeometryShop(F f, R r)
Construct a GeometryShop with an implicit function and resource handle.
Definition AMReX_EB2_GeometryShop.H:213
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:408
static constexpr int in_body
Definition AMReX_EB2_GeometryShop.H:190
static constexpr int allcovered
Definition AMReX_EB2_GeometryShop.H:194
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:266
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
static constexpr bool isGPUable() noexcept
Definition AMReX_EB2_GeometryShop.H:322
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:338
void fillFab_Cpu(BaseFab< Real > &levelset, const Geometry &geom, Box const &bounding_box) const noexcept
Definition AMReX_EB2_GeometryShop.H:378
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:480
static constexpr int on_boundary
Definition AMReX_EB2_GeometryShop.H:189
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:519
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
GpuArray< Real, 3 > ProbLoArray() const noexcept
Definition AMReX_Geometry.H:192
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:589
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