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 <limits>
15#include <utility>
16
17namespace amrex::EB2 {
18
19template <class F>
20requires (IsGPUable<F>::value)
22Real
23IF_f (F const& f, GpuArray<Real,AMREX_SPACEDIM> const& p) noexcept
24{
25 return f(AMREX_D_DECL(p[0],p[1],p[2]));
26}
27
28template <class F>
29requires (!IsGPUable<F>::value)
31Real
33{
36 amrex::Error("EB2::GeometryShop: how did this happen?");
37 return 0.0;
38 ))
39 AMREX_IF_ON_HOST((return f({AMREX_D_DECL(p[0],p[1],p[2])});))
40}
41
42template <class F>
44Real
47 int rangedir, F const& f)
48{
49 // The tolerance must not exceed a fraction of the edge length.
50 const Real h = std::abs(hi[rangedir] - lo[rangedir]);
51#ifdef AMREX_USE_FLOAT
52 const Real tol = amrex::min(1.e-4_rt, 1.e-2_rt*h);
53 const Real EPS = 1.0e-6_rt;
54#else
55 const Real tol = amrex::min(1.e-12, 1.e-10*h);
56 const Real EPS = 3.0e-15;
57#endif
58 const int MAXITER = 100;
59
60 Real p, q, r, s;
61
64
65 Real fa = IF_f(f, aPt);
66 Real fb = IF_f(f, bPt);
67 Real c = bPt[rangedir];
68 Real fc = fb;
69
70 if (fb*fa > 0.0_rt) {
71// amrex::AllPrint() << "fa " << fa << " fb " << fb << "\n";
72 amrex::Error("BrentRootFinder. Root must be bracketed, but instead the supplied end points have the same sign.");
73 return 0.0_rt;
74 } else if (fa == 0.0_rt) {
75 return aPt[rangedir];
76 } else if (fb == 0.0_rt) {
77 return bPt[rangedir];
78 }
79
80 Real d = 0.0_rt, e = 0.0_rt;
81 int i;
82 for (i = 0; i < MAXITER; ++i)
83 {
84 if (fb*fc > 0)
85 {
86 // Rename a, b, c and adjust bounding interval d
87 c = aPt[rangedir];
88 fc = fa;
89 d = bPt[rangedir] - aPt[rangedir];
90 e = d;
91 }
92
93 if (std::abs(fc) < std::abs(fb))
94 {
95 aPt[rangedir] = bPt[rangedir];
96 bPt[rangedir] = c;
97 c = aPt[rangedir];
98 fa = fb;
99 fb = fc;
100 fc = fa;
101 }
102
103 // Convergence check
104 Real tol1 = 2.0_rt * EPS * std::abs(bPt[rangedir]) + 0.5_rt * tol;
105 Real xm = 0.5_rt * (c - bPt[rangedir]);
106
107 if (std::abs(xm) <= tol1 || fb == 0.0_rt)
108 {
109 break;
110 }
111
112 if (std::abs(e) >= tol1 && std::abs(fa) > std::abs(fb))
113 {
114 // Attempt inverse quadratic interpolation
115 s = fb / fa;
116 if (aPt[rangedir] == c)
117 {
118 p = 2.0_rt * xm * s;
119 q = 1.0_rt - s;
120 }
121 else
122 {
123 q = fa / fc;
124 r = fb / fc;
125 p = s * (2.0_rt * xm * q * (q-r) - (bPt[rangedir]-aPt[rangedir]) * (r-1.0_rt));
126 q = (q-1.0_rt) * (r-1.0_rt) * (s-1.0_rt);
127 }
128
129 // Check whether in bounds
130 if (p > 0) { q = -q; }
131
132 p = std::abs(p);
133
134 if (2.0_rt * p < amrex::min(3.0_rt*xm*q-std::abs(tol1*q), 1.0_rt*std::abs(e*q)))
135 {
136 // Accept interpolation
137 e = d;
138 d = p / q;
139 }
140 else
141 {
142 // Interpolation failed, use bisection
143 d = xm;
144 e = d;
145 }
146 }
147 else
148 {
149 // Bounds decreasing too slowly, use bisection
150 d = xm;
151 e = d;
152 }
153
154 // Move last best guess to a
155 aPt[rangedir] = bPt[rangedir];
156 fa = fb;
157
158 // Evaluate new trial root
159 if (std::abs(d) > tol1)
160 {
161 bPt[rangedir] = bPt[rangedir] + d;
162 }
163 else
164 {
165 if (xm < 0) { bPt[rangedir] = bPt[rangedir] - tol1; }
166 else { bPt[rangedir] = bPt[rangedir] + tol1; }
167 }
168
169 fb = IF_f(f, bPt);
170 }
171
172 if (i >= MAXITER)
173 {
174 amrex::Error("BrentRootFinder: exceeding maximum iterations.");
175 }
176
177 return bPt[rangedir];
178}
179
180// Inverse squared distance to a neighboring node, by the bits of its
181// nonzero index offsets, for IF_roundoff.
184{
185 Real const d2x = dx[0]*dx[0];
186 Real const d2y = dx[1]*dx[1];
187 Real const d2z = (AMREX_SPACEDIM == 3) ? dx[AMREX_SPACEDIM-1]*dx[AMREX_SPACEDIM-1] : 0.0_rt;
188 GpuArray<Real,8> invd2;
189 invd2[0] = 0.0_rt;
190 for (int m = 1; m < 8; ++m) {
191 Real const den = Real(m&1)*d2x + Real((m>>1)&1)*d2y + Real((m>>2)&1)*d2z;
192 invd2[m] = (den > 0.0_rt) ? Real(1.0)/den : 0.0_rt; // z entries unused in 2D
193 }
194 return invd2;
195}
196
197// Size of the coordinate operands anywhere in the domain, for IF_roundoff.
200 GpuArray<Real,AMREX_SPACEDIM> const& probhi)
201{
202 Real size = 0.0_rt;
203 for (int d = 0; d < AMREX_SPACEDIM; ++d) {
204 size += std::abs(problo[d]) + (probhi[d]-problo[d]);
205 }
206 return size;
207}
208
209// Upper bound of IF_roundoff over a fab whose values lie in [vmin,vmax]
211Real IF_tolmax (Real vmin, Real vmax, GpuArray<Real,8> const& invd2, Real size)
212{
213 Real const dv = vmax - vmin;
214 Real invd2max = invd2[1];
215 for (int m = 2; m < 8; ++m) { invd2max = amrex::max(invd2max, invd2[m]); }
216 return Real(16.0) * std::numeric_limits<Real>::epsilon() * std::sqrt(dv*dv*invd2max) * size;
217}
218
219// Tolerance below which f at node iv is roundoff: 16 epsilon times the
220// largest slope of f to a neighboring node, diagonals included so that a
221// kink of the surface at the node does not hide the slope, times size (see
222// IF_size). size is uniform over the domain so that a surface at the same
223// distance from a whole row of nodes is classified the same along the row.
224// Neighbor values are taken from fab where it has them. Nodes outside
225// [blo,bhi] take the neighbor values of the clamped node.
226template <class F>
228Real
229IF_roundoff (F const& f, IntVect const& iv, Real fv, Array4<Real const> const& fab,
230 GpuArray<Real,AMREX_SPACEDIM> const& problo,
231 GpuArray<Real,AMREX_SPACEDIM> const& dx, GpuArray<Real,8> const& invd2,
232 Real size, IntVect const& blo, IntVect const& bhi)
233{
234 IntVect ivc = iv;
235 for (int d = 0; d < AMREX_SPACEDIM; ++d) { ivc[d] = amrex::Clamp(iv[d], blo[d], bhi[d]); }
236 constexpr int kmax = (AMREX_SPACEDIM == 3) ? 1 : 0;
237 int const ic = ivc[0];
238 int const jc = ivc[1];
239 int const kc = (AMREX_SPACEDIM == 3) ? ivc[AMREX_SPACEDIM-1] : 0;
240 Real slope2 = 0.0_rt;
241 IntVect const ivlo = ivc - 1;
242 IntVect const ivhi = ivc + 1;
243 if (ivlo.allGE(blo) && ivhi.allLE(bhi) && fab.contains(ivlo) && fab.contains(ivhi)) {
244 // Same result as the loop below, grouped by invd2.
245 Real const* q = fab.ptr(ic,jc,kc);
246 Long const js = fab.template get_stride<1>();
247 auto sq = [=] (Real a) { return (a-fv)*(a-fv); };
248 slope2 = amrex::max(sq(q[-1]), sq(q[1])) * invd2[1];
249 slope2 = amrex::max(slope2, amrex::max(sq(q[-js]), sq(q[js])) * invd2[2]);
250 slope2 = amrex::max(slope2, amrex::max(sq(q[-1-js]), sq(q[1-js]),
251 sq(q[-1+js]), sq(q[1+js])) * invd2[3]);
252#if (AMREX_SPACEDIM == 3)
253 Long const ks = fab.template get_stride<2>();
254 slope2 = amrex::max(slope2, amrex::max(sq(q[-ks]), sq(q[ks])) * invd2[4]);
255 slope2 = amrex::max(slope2, amrex::max(sq(q[-1-ks]), sq(q[1-ks]),
256 sq(q[-1+ks]), sq(q[1+ks])) * invd2[5]);
257 slope2 = amrex::max(slope2, amrex::max(sq(q[-js-ks]), sq(q[js-ks]),
258 sq(q[-js+ks]), sq(q[js+ks])) * invd2[6]);
259 slope2 = amrex::max(slope2, amrex::max(sq(q[-1-js-ks]), sq(q[1-js-ks]),
260 sq(q[-1+js-ks]), sq(q[1+js-ks]),
261 sq(q[-1-js+ks]), sq(q[1-js+ks]),
262 sq(q[-1+js+ks]), sq(q[1+js+ks])) * invd2[7]);
263#endif
264 } else {
265 for (int kk = -kmax; kk <= kmax; ++kk) {
266 for (int jj = -1; jj <= 1; ++jj) {
267 for (int ii = -1; ii <= 1; ++ii) {
268 int const m = int(ii != 0) | (int(jj != 0) << 1) | (int(kk != 0) << 2);
269 if (m != 0) {
270 IntVect ivn = ivc + IntVect(AMREX_D_DECL(ii,jj,kk));
271 for (int d = 0; d < AMREX_SPACEDIM; ++d) {
272 ivn[d] = amrex::Clamp(ivn[d], blo[d], bhi[d]);
273 }
274 Real fn;
275 if (fab.contains(ivn)) {
276 fn = fab(ivn);
277 } else {
279 for (int d = 0; d < AMREX_SPACEDIM; ++d) {
280 p[d] = node_coord(problo[d], ivn[d], dx[d]);
281 }
282 fn = IF_f(f, p);
283 }
284 Real const df = fn - fv;
285 slope2 = amrex::max(slope2, df*df*invd2[m]);
286 }
287 }}}
288 }
289 return Real(16.0) * std::numeric_limits<Real>::epsilon() * std::sqrt(slope2) * size;
290}
291
292// Intercept on the edge from node ivz, which is on the surface (zero level
293// set), to the fluid node one cell away in direction rangedir. This is
294// where f leaves the surface, which is not ivz if the edge runs along it.
295template <class F>
297Real
298ZeroEndIntercept (F const& f, IntVect const& ivz, int rangedir, int side,
299 Array4<Real const> const& lst,
300 GpuArray<Real,AMREX_SPACEDIM> const& problo,
301 GpuArray<Real,AMREX_SPACEDIM> const& dx, Real size,
302 IntVect const& blo, IntVect const& bhi)
303{
305 for (int d = 0; d < AMREX_SPACEDIM; ++d) {
306 p[d] = node_coord(problo[d], amrex::Clamp(ivz[d],blo[d],bhi[d]), dx[d]);
307 }
308 Real const tol = IF_roundoff(f, ivz, IF_f(f, p), lst, problo, dx, IF_invdist2(dx), size,
309 blo, bhi);
310 IntVect ivf = ivz;
311 ivf[rangedir] += side;
312 Real a = p[rangedir];
313 Real b = node_coord(problo[rangedir], amrex::Clamp(ivf[rangedir],blo[rangedir],bhi[rangedir]),
314 dx[rangedir]);
315 Real const xtol = Real(4.0) * std::numeric_limits<Real>::epsilon()
316 * amrex::max(std::abs(a), std::abs(b), std::abs(problo[rangedir]));
317 Real const x0 = a;
318 p[rangedir] = b;
319 Real const fb = IF_f(f, p);
320 Real const h = std::abs(b - a);
321 // Within roundoff of the node, the surface just passes through it.
322 // Roundoff here is where f drops below -tol at the slope along the edge.
323 Real const node_tol = (fb != 0.0_rt) ? Real(2.0)*tol*h/std::abs(fb) : 0.0_rt;
324 // The usual case: f is already below -tol at node_tol, where the
325 // bisection would end.
326 if (node_tol >= h) { return x0; }
327 p[rangedir] = x0 + Real(side)*node_tol;
328 if (IF_f(f, p) < -tol) { return x0; }
329 for (int iter = 0; iter < 200 && std::abs(b-a) > xtol; ++iter) {
330 p[rangedir] = Real(0.5)*(a+b);
331 if (IF_f(f, p) >= -tol) {
332 a = p[rangedir];
333 } else {
334 b = p[rangedir];
335 }
336 }
337 return (std::abs(a - x0) <= node_tol) ? x0 : a;
338}
339
346template <class F, class R = int>
348{
349public:
350
351 static constexpr int in_fluid = -1;
352 static constexpr int on_boundary = 0;
353 static constexpr int in_body = 1;
354 //
355 static constexpr int allregular = -1;
356 static constexpr int mixedcells = 0;
357 static constexpr int allcovered = 1;
358
360
366 explicit GeometryShop (F f)
367 : m_f(std::move(f)), m_resource()
368 {}
369
377 : m_f(std::move(f)), m_resource(std::move(r))
378 {}
379
381 [[nodiscard]] F const& GetImpFunc () const& { return m_f; }
382 F&& GetImpFunc () && { return std::move(m_f); }
383
384 // Nodes are classified as in fillFab, with the same neighbors and
385 // clamping for the roundoff tolerance, so that a box is regular or
386 // covered only if fillFab would make no cut cells in it. Both signs
387 // make the box mixed. If all nodes are fluid, only those within
388 // IF_tolmax can be body, which needs the roundoff stencil.
389 [[nodiscard]] int getBoxType_Cpu (const Box& bx, Geometry const& geom,
390 Box const& bounding_box) const noexcept
391 {
392 const auto& problo = geom.ProbLoArray();
393 const auto& dx = geom.CellSizeArray();
394 const Box gbx = amrex::grow(bx, 1);
395 const auto& blo = amrex::lbound(bounding_box);
396 const auto& bhi = amrex::ubound(bounding_box);
397 BaseFab<Real> ffab(gbx, 1, The_Cpu_Arena());
398 const auto& fv = ffab.array();
399 const auto& fc = ffab.const_array();
400 const auto& lo = amrex::lbound(bx);
401 const auto& hi = amrex::ubound(bx);
402 int npos = 0, nneg = 0;
403 for (int k = lo.z; k <= hi.z; ++k) {
404 for (int j = lo.y; j <= hi.y; ++j) {
405 for (int i = lo.x; i <= hi.x; ++i) {
406 Real const v = m_f(RealArray{AMREX_D_DECL(node_coord(problo[0], amrex::Clamp(i,blo.x,bhi.x), dx[0]),
407 node_coord(problo[1], amrex::Clamp(j,blo.y,bhi.y), dx[1]),
408 node_coord(problo[2], amrex::Clamp(k,blo.z,bhi.z), dx[2]))});
409 fv(i,j,k) = v;
410 if (v >= 0.0_rt) { ++npos; } else { ++nneg; }
411 if (npos > 0 && nneg > 0) { return mixedcells; }
412 }}}
413 if (nneg == 0) { return allcovered; }
414
415 Real vmin = std::numeric_limits<Real>::max();
416 Real vmax = std::numeric_limits<Real>::lowest();
417 amrex::LoopOnCpu(gbx, [&] (int i, int j, int k) noexcept
418 {
420 if (!bx.contains(IntVect(AMREX_D_DECL(i,j,k)))) {
421 fv(i,j,k) = m_f(RealArray{AMREX_D_DECL(node_coord(problo[0], amrex::Clamp(i,blo.x,bhi.x), dx[0]),
422 node_coord(problo[1], amrex::Clamp(j,blo.y,bhi.y), dx[1]),
423 node_coord(problo[2], amrex::Clamp(k,blo.z,bhi.z), dx[2]))});
424 }
425 vmin = amrex::min(vmin, fc(i,j,k));
426 vmax = amrex::max(vmax, fc(i,j,k));
427 });
428 IntVect const glo = bounding_box.smallEnd();
429 IntVect const ghi = bounding_box.bigEnd();
430 auto const invd2 = IF_invdist2(dx);
431 auto const size = IF_size(problo, geom.ProbHiArray());
432 Real const tolmax = IF_tolmax(vmin, vmax, invd2, size);
433 if (vmax < -tolmax) { return allregular; }
434 for (int k = lo.z; k <= hi.z; ++k) {
435 for (int j = lo.y; j <= hi.y; ++j) {
436 for (int i = lo.x; i <= hi.x; ++i) {
437 IntVect const iv(AMREX_D_DECL(i,j,k));
438 Real const v = fc(i,j,k);
439 if (std::abs(v) <= tolmax &&
440 std::abs(v) <= IF_roundoff(m_f, iv, v, fc, problo, dx, invd2, size, glo, ghi)) {
441 return mixedcells;
442 }
443 }}}
444 return allregular;
445 }
446
456 template <class U=F>
457 requires (IsGPUable<U>::value)
458 [[nodiscard]] int getBoxType (const Box& bx, const Geometry& geom, RunOn run_on,
459 Box const& bounding_box) const noexcept
460 {
461 if (run_on == RunOn::Gpu && Gpu::inLaunchRegion())
462 {
463 const auto& problo = geom.ProbLoArray();
464 const auto& dx = geom.CellSizeArray();
465 auto const& f = m_f;
466 // Node classification as in getBoxType_Cpu.
467 const Box gbx = amrex::grow(bx, 1);
468 const auto& blo = amrex::lbound(bounding_box);
469 const auto& bhi = amrex::ubound(bounding_box);
470 BaseFab<Real> ffab(gbx, 1, The_Async_Arena());
471 auto const& fv = ffab.array();
472 auto const& fc = ffab.const_array();
474 ReduceData<int,int,Real,Real> reduce_data(reduce_op);
475 using ReduceTuple = typename decltype(reduce_data)::Type;
476 reduce_op.eval(gbx, reduce_data,
477 [=] AMREX_GPU_DEVICE (int i, int j, int k) -> ReduceTuple
478 {
480 Real const v = f(AMREX_D_DECL(node_coord(problo[0], amrex::Clamp(i,blo.x,bhi.x), dx[0]),
481 node_coord(problo[1], amrex::Clamp(j,blo.y,bhi.y), dx[1]),
482 node_coord(problo[2], amrex::Clamp(k,blo.z,bhi.z), dx[2])));
483 fv(i,j,k) = v;
484 bool const inbx = bx.contains(IntVect(AMREX_D_DECL(i,j,k)));
485 return {int(inbx && v >= 0.0_rt), int(inbx && v < 0.0_rt), v, v};
486 });
487 ReduceTuple hv = reduce_data.value(reduce_op);
488 int const npos = amrex::get<0>(hv);
489 int const nneg = amrex::get<1>(hv);
490 if (npos > 0 && nneg > 0) { return mixedcells; }
491 if (nneg == 0) { return allcovered; }
492
493 IntVect const glo = bounding_box.smallEnd();
494 IntVect const ghi = bounding_box.bigEnd();
495 auto const invd2 = IF_invdist2(dx);
496 auto const size = IF_size(problo, geom.ProbHiArray());
497 Real const tolmax = IF_tolmax(amrex::get<2>(hv), amrex::get<3>(hv), invd2, size);
498 if (amrex::get<3>(hv) < -tolmax) { return allregular; }
499 ReduceOps<ReduceOpSum> reduce_op2;
500 ReduceData<int> reduce_data2(reduce_op2);
501 using ReduceTuple2 = typename decltype(reduce_data2)::Type;
502 reduce_op2.eval(bx, reduce_data2,
503 [=] AMREX_GPU_DEVICE (int i, int j, int k) -> ReduceTuple2
504 {
506 IntVect const iv(AMREX_D_DECL(i,j,k));
507 Real const v = fc(iv);
508 bool const body = std::abs(v) <= tolmax &&
509 std::abs(v) <= IF_roundoff(f, iv, v, fc, problo, dx, invd2, size, glo, ghi);
510 return {int(body)};
511 });
512 int const nbody = amrex::get<0>(reduce_data2.value(reduce_op2));
513 return (nbody > 0) ? mixedcells : allregular;
514 }
515 else
516 {
517 return getBoxType_Cpu(bx, geom, bounding_box);
518 }
519 }
520
522 template <class U=F>
523 requires (!IsGPUable<U>::value)
524 [[nodiscard]] int getBoxType (const Box& bx, const Geometry& geom, RunOn,
525 Box const& bounding_box) const noexcept
526 {
527 return getBoxType_Cpu(bx, geom, bounding_box);
528 }
529
530 template <class U=F>
531 requires (IsGPUable<U>::value)
532 static constexpr bool isGPUable () noexcept { return true; }
533
534 template <class U=F>
535 requires (!IsGPUable<U>::value)
536 static constexpr bool isGPUable () noexcept { return false; }
537
546 template <class U=F>
547 requires (IsGPUable<U>::value)
548 void fillFab (BaseFab<Real>& levelset, const Geometry& geom, RunOn run_on,
549 Box const& bounding_box) const noexcept
550 {
551 const auto& problo = geom.ProbLoArray();
552 const auto& dx = geom.CellSizeArray();
553 const Box& bx = levelset.box();
554 // One more node layer for the neighbors in the roundoff tolerance.
555 const Box gbx = amrex::grow(bx, 1);
556 // Nodes outside the bounding box take the values of the nearest node in it.
557 const Box ibx = bx & bounding_box;
558 const Box ibx2 = gbx & bounding_box;
559 IntVect const bloiv = bounding_box.smallEnd();
560 IntVect const bhiiv = bounding_box.bigEnd();
561 const auto& f = m_f;
562 bool const on_gpu = (run_on == RunOn::Gpu) && Gpu::inLaunchRegion();
563 Arena* const arena = on_gpu ? The_Async_Arena() : The_Cpu_Arena();
564 BaseFab<Real> ffab(gbx, 1, arena);
565 auto const& fv = ffab.array();
566 auto const& fc = ffab.const_array();
567 // Range of the sampled values, for IF_tolmax, without a host sync.
568 auto* const pmm = static_cast<Real*>(arena->alloc(2*sizeof(Real)));
569#ifdef AMREX_USE_GPU
570 if (on_gpu) {
572 {
573 pmm[0] = std::numeric_limits<Real>::max();
574 pmm[1] = std::numeric_limits<Real>::lowest();
575 });
576 amrex::ParallelFor(Gpu::KernelInfo().setReduction(true), ibx2,
577 [=] AMREX_GPU_DEVICE (int i, int j, int k, Gpu::Handler const& h) noexcept
578 {
580 Real const v = f(AMREX_D_DECL(node_coord(problo[0], i, dx[0]),
581 node_coord(problo[1], j, dx[1]),
582 node_coord(problo[2], k, dx[2])));
583 fv(i,j,k) = v;
584 Gpu::deviceReduceMin(pmm, v, h);
585 Gpu::deviceReduceMax(pmm+1, v, h);
586 });
587 } else
588#endif
589 {
590 Real vmin = std::numeric_limits<Real>::max();
591 Real vmax = std::numeric_limits<Real>::lowest();
592 amrex::LoopOnCpu(ibx2, [&] (int i, int j, int k) noexcept
593 {
595 Real const v = f(AMREX_D_DECL(node_coord(problo[0], i, dx[0]),
596 node_coord(problo[1], j, dx[1]),
597 node_coord(problo[2], k, dx[2])));
598 fv(i,j,k) = v;
599 vmin = amrex::min(vmin, v);
600 vmax = amrex::max(vmax, v);
601 });
602 pmm[0] = vmin;
603 pmm[1] = vmax;
604 }
605 if (ibx2 != gbx) {
606 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, gbx, i, j, k,
607 {
609 IntVect const iv(AMREX_D_DECL(i,j,k));
610 if (!ibx2.contains(iv)) {
611 IntVect ivc = iv;
612 for (int d = 0; d < AMREX_SPACEDIM; ++d) {
613 ivc[d] = amrex::Clamp(iv[d], bloiv[d], bhiiv[d]);
614 }
615 fv(iv) = fv(ivc);
616 }
617 });
618 }
619
620 // Nodes on the surface within roundoff get a zero level set.
621 auto const& s = levelset.array();
622 auto const invd2 = IF_invdist2(dx);
623 auto const size = IF_size(problo, geom.ProbHiArray());
624 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, ibx, i, j, k,
625 {
626 IntVect const iv(AMREX_D_DECL(i,j,k));
627 Real const v = fc(i,j,k);
628 Real const tolmax = IF_tolmax(pmm[0], pmm[1], invd2, size);
629 s(i,j,k) = (v != 0.0_rt && std::abs(v) <= tolmax &&
630 std::abs(v) <= IF_roundoff(f, iv, v, fc, problo, dx, invd2, size,
631 bloiv, bhiiv))
632 ? 0.0_rt : v;
633 });
634 arena->free(pmm);
635 if (ibx != bx) {
636 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, bx, i, j, k,
637 {
639 IntVect const iv(AMREX_D_DECL(i,j,k));
640 if (!ibx.contains(iv)) {
641 IntVect ivc = iv;
642 for (int d = 0; d < AMREX_SPACEDIM; ++d) {
643 ivc[d] = amrex::Clamp(iv[d], bloiv[d], bhiiv[d]);
644 }
645 s(iv) = s(ivc);
646 }
647 });
648 }
649 }
650
652 template <class U=F>
653 requires (!IsGPUable<U>::value)
655 Box const& bounding_box) const noexcept
656 {
657#ifdef AMREX_USE_GPU
658 if (!levelset.arena()->isHostAccessible()) {
659 const Box& bx = levelset.box();
660 BaseFab<Real> h_levelset(bx, levelset.nComp(), The_Pinned_Arena());
661 fillFab_Cpu(h_levelset, geom, bounding_box);
662 Gpu::htod_memcpy_async(levelset.dataPtr(), h_levelset.dataPtr(),
663 levelset.nBytes(bx, levelset.nComp()));
665 }
666 else
667#endif
668 {
669 fillFab_Cpu(levelset, geom, bounding_box);
670 }
671 }
672
674 Box const& bounding_box) const noexcept
675 {
676 const auto& problo = geom.ProbLoArray();
677 const auto& dx = geom.CellSizeArray();
678 const Box& bx = levelset.box();
679 // See the GPU version.
680 const Box gbx = amrex::grow(bx, 1);
681 const Box ibx = bx & bounding_box;
682 const Box ibx2 = gbx & bounding_box;
683 IntVect const bloiv = bounding_box.smallEnd();
684 IntVect const bhiiv = bounding_box.bigEnd();
685
686 BaseFab<Real> ffab(gbx, 1, The_Cpu_Arena());
687 const auto& fv = ffab.array();
688 const auto& fc = ffab.const_array();
689 Real vmin = std::numeric_limits<Real>::max();
690 Real vmax = std::numeric_limits<Real>::lowest();
691 amrex::LoopOnCpu(ibx2, [&] (int i, int j, int k) noexcept
692 {
694 Real const v = m_f(RealArray{AMREX_D_DECL(node_coord(problo[0], i, dx[0]),
695 node_coord(problo[1], j, dx[1]),
696 node_coord(problo[2], k, dx[2]))});
697 fv(i,j,k) = v;
698 vmin = amrex::min(vmin, v);
699 vmax = amrex::max(vmax, v);
700 });
701 if (ibx2 != gbx) {
702 amrex::LoopOnCpu(gbx, [&] (int i, int j, int k) noexcept
703 {
705 IntVect const iv(AMREX_D_DECL(i,j,k));
706 if (!ibx2.contains(iv)) {
707 IntVect ivc = iv;
708 for (int d = 0; d < AMREX_SPACEDIM; ++d) {
709 ivc[d] = amrex::Clamp(iv[d], bloiv[d], bhiiv[d]);
710 }
711 fv(iv) = fv(ivc);
712 }
713 });
714 }
715
716 // Nodes on the surface within roundoff get a zero level set.
717 auto const& s = levelset.array();
718 auto const invd2 = IF_invdist2(dx);
719 auto const size = IF_size(problo, geom.ProbHiArray());
720 Real const tolmax = IF_tolmax(vmin, vmax, invd2, size);
721 amrex::LoopOnCpu(ibx, [&] (int i, int j, int k) noexcept
722 {
723 IntVect const iv(AMREX_D_DECL(i,j,k));
724 Real const v = fc(i,j,k);
725 s(i,j,k) = (v != 0.0_rt && std::abs(v) <= tolmax &&
726 std::abs(v) <= IF_roundoff(m_f, iv, v, fc, problo, dx, invd2, size,
727 bloiv, bhiiv))
728 ? 0.0_rt : v;
729 });
730 if (ibx != bx) {
731 amrex::LoopOnCpu(bx, [&] (int i, int j, int k) noexcept
732 {
734 IntVect const iv(AMREX_D_DECL(i,j,k));
735 if (!ibx.contains(iv)) {
736 IntVect ivc = iv;
737 for (int d = 0; d < AMREX_SPACEDIM; ++d) {
738 ivc[d] = amrex::Clamp(iv[d], bloiv[d], bhiiv[d]);
739 }
740 s(iv) = s(ivc);
741 }
742 });
743 }
744 }
745
746
758 template <class U=F>
759 requires (IsGPUable<U>::value)
760 void getIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
761 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
762 Array4<Real const> const& lst, Geometry const& geom, RunOn run_on,
763 Box const& bounding_box) const noexcept
764 {
765 auto const& dx = geom.CellSizeArray();
766 auto const& problo = geom.ProbLoArray();
767 const auto& blo = amrex::lbound(bounding_box);
768 const auto& bhi = amrex::ubound(bounding_box);
769 IntVect const bloiv = bounding_box.smallEnd();
770 IntVect const bhiiv = bounding_box.bigEnd();
771 Real const size = IF_size(problo, geom.ProbHiArray());
772 auto const& f = m_f;
773 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
774 Array4<Real> const& inter = inter_arr[idim];
775 Array4<Type_t const> const& type = type_arr[idim];
776 const Box bx{inter};
777 AMREX_HOST_DEVICE_FOR_3D_FLAG(run_on, bx, i, j, k,
778 {
779 if (type(i,j,k) == Type::irregular) {
780 IntVect ivlo(AMREX_D_DECL(i,j,k));
781 IntVect ivhi(AMREX_D_DECL(i,j,k));
782 ivhi[idim] += 1;
783 if (lst(ivlo) == 0.0_rt) {
784 inter(i,j,k) = ZeroEndIntercept(f, ivlo, idim, 1, lst, problo, dx, size,
785 bloiv, bhiiv);
786 } else if (lst(ivhi) == 0.0_rt) {
787 inter(i,j,k) = ZeroEndIntercept(f, ivhi, idim, -1, lst, problo, dx, size,
788 bloiv, bhiiv);
789 } else {
790 inter(i,j,k) = BrentRootFinder
791 ({AMREX_D_DECL(node_coord(problo[0], amrex::Clamp(ivlo[0],blo.x,bhi.x), dx[0]),
792 node_coord(problo[1], amrex::Clamp(ivlo[1],blo.y,bhi.y), dx[1]),
793 node_coord(problo[2], amrex::Clamp(ivlo[2],blo.z,bhi.z), dx[2]))},
794 {AMREX_D_DECL(node_coord(problo[0], amrex::Clamp(ivhi[0],blo.x,bhi.x), dx[0]),
795 node_coord(problo[1], amrex::Clamp(ivhi[1],blo.y,bhi.y), dx[1]),
796 node_coord(problo[2], amrex::Clamp(ivhi[2],blo.z,bhi.z), dx[2]))},
797 idim, f);
798 }
799 } else {
800 inter(i,j,k) = no_intercept;
801 }
802 });
803 }
804 }
805
807 template <class U=F>
808 requires (!IsGPUable<U>::value)
809 void getIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
810 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
811 Array4<Real const> const& lst, Geometry const& geom, RunOn,
812 Box const& bounding_box) const noexcept
813 {
814#ifdef AMREX_USE_GPU
815 BaseFab<Real> h_lst_fab;
816 Array4<Real const> h_lst = lst;
817 if (!The_Arena()->isHostAccessible()) {
818 const Box lbx{lst};
819 h_lst_fab.resize(lbx, 1, The_Pinned_Arena());
820 Gpu::dtoh_memcpy_async(h_lst_fab.dataPtr(), lst.dataPtr(), h_lst_fab.nBytes(lbx, 1));
822 h_lst = h_lst_fab.const_array();
823 }
824#endif
825 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
826 Array4<Real> const& inter = inter_arr[idim];
827 Array4<Type_t const> const& type = type_arr[idim];
828 // When GShop is not GPUable, the intercept is either in managed memory or
829 // in device memory and needs to be copied to host and setup there.
830#ifdef AMREX_USE_GPU
831 if (!The_Arena()->isHostAccessible()) {
832 const Box bx{inter};
833 BaseFab<Real> h_inter_fab(bx, 1, The_Pinned_Arena());
834 BaseFab<Type_t> h_type_fab(bx, 1, The_Pinned_Arena());
835 Gpu::dtoh_memcpy_async(h_type_fab.dataPtr(), type.dataPtr(),
836 h_type_fab.nBytes(bx, 1));
838 const auto& h_inter = h_inter_fab.array();
839 const auto& h_type = h_type_fab.const_array();
840 getIntercept_Cpu(h_inter, h_type, h_lst, geom, bounding_box, idim);
841 Gpu::htod_memcpy_async(inter.dataPtr(), h_inter.dataPtr(),
842 h_inter_fab.nBytes(bx, 1));
844 }
845 else
846#endif
847 {
848 getIntercept_Cpu(inter, type, lst, geom, bounding_box, idim);
849 }
850
851 }
852 }
853
854 void getIntercept_Cpu (Array4<Real> const& inter,
855 Array4<Type_t const> const& type,
856 Array4<Real const> const& lst,
857 Geometry const& geom,
858 Box const& bounding_box,
859 const int idim) const noexcept
860 {
861 auto const& dx = geom.CellSizeArray();
862 auto const& problo = geom.ProbLoArray();
863 const auto& blo = amrex::lbound(bounding_box);
864 const auto& bhi = amrex::ubound(bounding_box);
865 IntVect const bloiv = bounding_box.smallEnd();
866 IntVect const bhiiv = bounding_box.bigEnd();
867 Real const size = IF_size(problo, geom.ProbHiArray());
868 const Box bx{inter};
869 amrex::LoopOnCpu(bx, [&] (int i, int j, int k) noexcept
870 {
871 if (type(i,j,k) == Type::irregular) {
872 IntVect ivlo(AMREX_D_DECL(i,j,k));
873 IntVect ivhi(AMREX_D_DECL(i,j,k));
874 ivhi[idim] += 1;
875 if (lst(ivlo) == 0.0_rt) {
876 inter(i,j,k) = ZeroEndIntercept(m_f, ivlo, idim, 1, lst, problo, dx, size,
877 bloiv, bhiiv);
878 } else if (lst(ivhi) == 0.0_rt) {
879 inter(i,j,k) = ZeroEndIntercept(m_f, ivhi, idim, -1, lst, problo, dx, size,
880 bloiv, bhiiv);
881 } else {
882 inter(i,j,k) = BrentRootFinder
883 ({AMREX_D_DECL(node_coord(problo[0], amrex::Clamp(ivlo[0],blo.x,bhi.x), dx[0]),
884 node_coord(problo[1], amrex::Clamp(ivlo[1],blo.y,bhi.y), dx[1]),
885 node_coord(problo[2], amrex::Clamp(ivlo[2],blo.z,bhi.z), dx[2]))},
886 {AMREX_D_DECL(node_coord(problo[0], amrex::Clamp(ivhi[0],blo.x,bhi.x), dx[0]),
887 node_coord(problo[1], amrex::Clamp(ivhi[1],blo.y,bhi.y), dx[1]),
888 node_coord(problo[2], amrex::Clamp(ivhi[2],blo.z,bhi.z), dx[2]))},
889 idim, m_f);
890 }
891 } else {
892 inter(i,j,k) = no_intercept;
893 }
894 });
895 }
896
905 void updateIntercept (Array<Array4<Real>,AMREX_SPACEDIM> const& inter_arr,
906 Array<Array4<Type_t const>,AMREX_SPACEDIM> const& type_arr,
907 Array4<Real const> const& lst, Geometry const& geom) const noexcept
908 {
909 auto const& dx = geom.CellSizeArray();
910 auto const& problo = geom.ProbLoArray();
911 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
912 Array4<Real> const& inter = inter_arr[idim];
913 Array4<Type_t const> const& type = type_arr[idim];
914 const Box bx{inter};
916 {
917 if (type(i,j,k) == Type::irregular) {
918 if (inter(i,j,k) == no_intercept) {
919 // Only edges without an intercept are updated: they
920 // became irregular because small cell fixing set lst
921 // to zero, or lst changed in FillBoundary at periodic
922 // boundaries. A node that was zero from the start
923 // keeps the intercept from getIntercept. Other
924 // edges are reset, so an edge that is irregular
925 // again gets a new intercept.
926 IntVect ivlo(AMREX_D_DECL(i,j,k));
927 IntVect ivhi = ivlo;
928 ivhi[idim] += 1;
929 // An irregular edge has one end below zero and one at or
930 // above. The intercept is at the body end.
931 bool const at_lo = lst(ivlo) >= Real(0.0);
932 int const inode = at_lo ? ivlo[idim] : ivhi[idim];
933 inter(i,j,k) = node_coord(problo[idim], inode, dx[idim]);
934 }
935 } else {
936 inter(i,j,k) = no_intercept;
937 }
938 });
939 }
940 }
941
942private:
943
944 F m_f;
945 R m_resource; // We use this to hold the ownership of resource for F if needed,
946 // because F needs to be a simply type suitable for GPU.
947};
948
950
954template <class F>
955GeometryShop<std::decay_t<F>>
956makeShop (F&& f)
957{
958 return GeometryShop<std::decay_t<F>>(std::forward<F>(f));
959}
960
965 * \param f Implicit function callable.
966 * \param r Extra resource (e.g., Parser data) kept alive for \p f.
967 */
968template <class F, class R>
969GeometryShop<std::decay_t<F>, std::decay_t<R>>
970makeShop (F&& f, R&& r)
971{
973 std::decay_t<R>>
974 (std::forward<F>(f), std::forward<R>(r));
975}
976
977}
978
979#endif
Fixed-size array types for use on GPU and CPU.
BaseFab container template providing box-based field storage.
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
Problem-domain geometry: maps between index space and physical space.
#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:102
#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
while parser_exe_t p
Definition AMReX_Parser_Exe_Body.H:33
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
Abstract base class for memory arenas.
Definition AMReX_Arena.H:154
virtual void free(void *pt)=0
Free a previously allocated block pointed to by pt.
virtual void * alloc(std::size_t sz)=0
Allocate sz bytes from this arena.
A FortranArrayBox(FAB)-like object.
Definition AMReX_BaseFab.H:222
Array4< T const > const_array() const noexcept
Return a const-qualified Array4 view over all components.
Definition AMReX_BaseFab.H:547
void resize(const Box &b, int N=1, Arena *ar=nullptr)
This function resizes a BaseFab so it covers the Box b with N components.
Definition AMReX_BaseFab.H:1776
std::size_t nBytes() const noexcept
Returns how many bytes used.
Definition AMReX_BaseFab.H:332
Array4< T const > array() const noexcept
Create an Array4 view over all components.
Definition AMReX_BaseFab.H:475
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:418
__host__ __device__ const IntVectND< dim > & bigEnd() const &noexcept
Return the inclusive upper bound of the box.
Definition AMReX_Box.H:136
__host__ __device__ bool contains(const IntVectND< dim > &p) const noexcept
Return true if argument is contained within BoxND.
Definition AMReX_Box.H:233
__host__ __device__ const IntVectND< dim > & smallEnd() const &noexcept
Return the inclusive lower bound of the box.
Definition AMReX_Box.H:124
GpuArray< Real, 3 > CellSizeArray() const noexcept
Returns the cell sizes as a GpuArray for use on host or device.
Definition AMReX_CoordSys.H:85
Samples an implicit function on boxes and fills EB data.
Definition AMReX_EB2_GeometryShop.H:348
F const & GetImpFunc() const &
Immutable access to the underlying implicit function.
Definition AMReX_EB2_GeometryShop.H:381
GeometryShop(F f, R r)
Construct a GeometryShop with an implicit function and resource handle.
Definition AMReX_EB2_GeometryShop.H:376
static constexpr int in_body
Definition AMReX_EB2_GeometryShop.H:353
int getBoxType(const Box &bx, const Geometry &geom, RunOn run_on, Box const &bounding_box) const noexcept
Classify bx as regular, covered, or mixed (GPU-capable case). Nodes on the surface within roundoff co...
Definition AMReX_EB2_GeometryShop.H:458
static constexpr int allcovered
Definition AMReX_EB2_GeometryShop.H:357
GeometryShop(F f)
Construct a GeometryShop that owns an implicit function f.
Definition AMReX_EB2_GeometryShop.H:366
static constexpr int in_fluid
Definition AMReX_EB2_GeometryShop.H:351
static constexpr bool isGPUable() noexcept
Definition AMReX_EB2_GeometryShop.H:532
int getBoxType_Cpu(const Box &bx, Geometry const &geom, Box const &bounding_box) const noexcept
Definition AMReX_EB2_GeometryShop.H:389
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:548
void fillFab_Cpu(BaseFab< Real > &levelset, const Geometry &geom, Box const &bounding_box) const noexcept
Definition AMReX_EB2_GeometryShop.H:673
void getIntercept_Cpu(Array4< Real > const &inter, Array4< Type_t const > const &type, Array4< Real const > const &lst, Geometry const &geom, Box const &bounding_box, const int idim) const noexcept
Definition AMReX_EB2_GeometryShop.H:854
F FunctionType
Definition AMReX_EB2_GeometryShop.H:359
static constexpr int mixedcells
Definition AMReX_EB2_GeometryShop.H:356
static constexpr int allregular
Definition AMReX_EB2_GeometryShop.H:355
void getIntercept(Array< Array4< Real >, 3 > const &inter_arr, Array< Array4< Type_t const >, 3 > const &type_arr, Array4< Real const > const &lst, 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:760
static constexpr int on_boundary
Definition AMReX_EB2_GeometryShop.H:352
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:905
F && GetImpFunc() &&
Definition AMReX_EB2_GeometryShop.H:382
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:87
GpuArray< Real, 3 > ProbHiArray() const noexcept
Return the hi end of the problem domain in a GpuArray for device code.
Definition AMReX_Geometry.H:224
GpuArray< Real, 3 > ProbLoArray() const noexcept
Return the lo end of the problem domain in a GpuArray for device code.
Definition AMReX_Geometry.H:219
Definition AMReX_GpuKernelInfo.H:8
__host__ __device__ constexpr bool allGE(const IntVectND< dim > &rhs) const noexcept
Returns true if this is greater than or equal to argument for all components. NOTE: This is NOT a str...
Definition AMReX_IntVect.H:542
__host__ __device__ constexpr bool allLE(const IntVectND< dim > &rhs) const noexcept
Returns true if this is less than or equal to argument for all components. NOTE: This is NOT a strict...
Definition AMReX_IntVect.H:492
Definition AMReX_Reduce.H:556
Type value()
Definition AMReX_Reduce.H:591
Definition AMReX_Reduce.H:715
void eval(MF const &mf, IntVect const &nghost, D &reduce_data, F &&f)
Definition AMReX_Reduce.H:852
amrex_real Real
Floating Point Type for Fields.
Definition AMReX_REAL.H:80
amrex_long Long
Definition AMReX_INT.H:30
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Return the inclusive upper bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1365
ArrayND< T, 4, true > Array4
Convenient alias for a 4D ArrayND with three spatial dimensions plus a component dimension.
Definition AMReX_Array4.H:1340
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Return the inclusive lower bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1351
__host__ __device__ BoxND< dim > grow(const BoxND< dim > &b, int i) noexcept
Return a copy of b grown uniformly by i cells in every direction.
Definition AMReX_Box.H:1328
std::array< T, N > Array
Definition AMReX_Array.H:31
Array< Real, 3 > RealArray
Definition AMReX_Array.H:35
Arena * The_Pinned_Arena()
Definition AMReX_Arena.cpp:879
Arena * The_Async_Arena()
Definition AMReX_Arena.cpp:849
Arena * The_Cpu_Arena()
Definition AMReX_Arena.cpp:889
Arena * The_Arena()
Definition AMReX_Arena.cpp:839
__host__ __device__ T node_coord(T problo, int i, T dx) noexcept
Coordinate of node i: problo + i*dx.
Definition AMReX_Geometry.H:625
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:31
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:54
__host__ __device__ constexpr T Clamp(const T &v, const T &lo, const T &hi)
Definition AMReX_Algorithm.H:123
Type flags representing regular/covered/etc.
Definition AMReX_EB2_Graph.H:49
static constexpr Type_t irregular
Definition AMReX_EB2_Graph.H:52
Definition AMReX_FabArrayBase.H:38
__host__ __device__ Real IF_tolmax(Real vmin, Real vmax, GpuArray< Real, 8 > const &invd2, Real size)
Definition AMReX_EB2_GeometryShop.H:211
__host__ __device__ GpuArray< Real, 8 > IF_invdist2(GpuArray< Real, 3 > const &dx)
Definition AMReX_EB2_GeometryShop.H:183
__host__ __device__ Real IF_roundoff(F const &f, IntVect const &iv, Real fv, Array4< Real const > const &fab, GpuArray< Real, 3 > const &problo, GpuArray< Real, 3 > const &dx, GpuArray< Real, 8 > const &invd2, Real size, IntVect const &blo, IntVect const &bhi)
Definition AMReX_EB2_GeometryShop.H:229
__host__ __device__ Real ZeroEndIntercept(F const &f, IntVect const &ivz, int rangedir, int side, Array4< Real const > const &lst, GpuArray< Real, 3 > const &problo, GpuArray< Real, 3 > const &dx, Real size, IntVect const &blo, IntVect const &bhi)
Definition AMReX_EB2_GeometryShop.H:298
__host__ __device__ Real IF_size(GpuArray< Real, 3 > const &problo, GpuArray< Real, 3 > const &probhi)
Definition AMReX_EB2_GeometryShop.H:199
__host__ __device__ Real BrentRootFinder(GpuArray< Real, 3 > const &lo, GpuArray< Real, 3 > const &hi, int rangedir, F const &f)
Definition AMReX_EB2_GeometryShop.H:45
__host__ __device__ Real IF_f(F const &f, GpuArray< Real, 3 > const &p) noexcept
Definition AMReX_EB2_GeometryShop.H:23
static constexpr Real no_intercept
Marker stored in an edge intercept array when the edge has no intercept.
Definition AMReX_EB2_Graph.H:57
GeometryShop< std::decay_t< F > > makeShop(F &&f)
Helper that constructs a GeometryShop owning the supplied implicit function.
Definition AMReX_EB2_GeometryShop.H:965
__device__ void deviceReduceMax(T *dest, T source, Gpu::Handler const &h) noexcept
Definition AMReX_GpuReduce.H:645
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:329
void dtoh_memcpy_async(void *p_h, const void *p_d, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:454
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:440
__device__ void deviceReduceMin(T *dest, T source, Gpu::Handler const &h) noexcept
Definition AMReX_GpuReduce.H:626
__host__ __device__ void ignore_unused(const Ts &...)
No-op helper that marks variables as intentionally unused.
Definition AMReX.H:273
@ F
Definition AMReX_Order.H:7
RunOn
Definition AMReX_GpuControl.H:65
@ Gpu
Definition AMReX_GpuControl.H:65
void ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition AMReX_CTOParallelForImpl.H:202
BoxND< 3 > Box
Box is an alias for amrex::BoxND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:35
void single_task(L &&f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1245
void Error(const std::string &msg)
Print a message to stderr and abort the program.
Definition AMReX.cpp:238
IntVectND< 3 > IntVect
IntVect is an alias for amrex::IntVectND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:38
const int[]
Definition AMReX_BLProfiler.cpp:1665
void LoopOnCpu(Dim3 lo, Dim3 hi, F const &f) noexcept
Definition AMReX_Loop.H:365
@ levelset
Definition AMReX_EBData.H:20
__host__ __device__ constexpr int get(IntVectND< dim > const &iv) noexcept
Get the Idx-th element of IntVectND<dim>.
Definition AMReX_IntVect.H:1360
__host__ __device__ constexpr bool contains(idx... i) const noexcept
Test whether an index tuple lies inside the ArrayND bounds.
Definition AMReX_Array4.H:691
__host__ __device__ constexpr T * dataPtr() const noexcept
Get raw data pointer.
Definition AMReX_Array4.H:617
__host__ __device__ T * ptr(idx... i) const noexcept
Multi-index ptr() for accessing pointer to element.
Definition AMReX_Array4.H:561
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:52
Definition AMReX_GpuTypes.H:86