Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
AMReX_algoim_K.H
Go to the documentation of this file.
1/*
2 * This file contains a modified version of Algoim developed by
3 * R. Saye, SIAM J. Sci. Comput., Vol. 37, No. 2, pp. A993-A1019,
4 * http://dx.doi.org/10.1137/140966290, https://algoim.github.io/.
5 *
6 * Algoim Copyright (c) 2018, The Regents of the University of
7 * California, through Lawrence Berkeley National Laboratory (subject
8 * to receipt of any required approvals from the U.S. Dept. of
9 * Energy). All rights reserved.
10 */
11
12#ifndef AMREX_ALGOIM_K_H_
13#define AMREX_ALGOIM_K_H_
14#include <AMReX_Config.H>
15
16#include <AMReX_algoim.H>
17#include <AMReX_Array.H>
18#include <limits>
19#include <cfloat>
20
21namespace amrex::algoim {
22
23struct EBPlane
24{
27
28 EBPlane () = default;
29
31 constexpr
32 EBPlane (GpuArray<Real,3> const& c, GpuArray<Real,3> const n) noexcept
33 : cent(c), norm(n)
34 {}
35
37 constexpr
38 EBPlane (Real cx, Real cy, Real cz, Real nx, Real ny, Real nz) noexcept
39 : cent{cx,cy,cz}, norm{nx,ny,nz}
40 {}
41
43 Real operator() (Real x, Real y, Real z) const noexcept
44 {
45 return (x-cent[0])*norm[0] + (y-cent[1])*norm[1] + (z-cent[2])*norm[2];
46 }
47
49 Real operator() (GpuArray<Real,3> const& p) const noexcept
50 {
51 return (p[0]-cent[0])*norm[0] + (p[1]-cent[1])*norm[1]
52 + (p[2]-cent[2])*norm[2];
53 }
54
56 Real grad (int d) const noexcept
57 {
58 return norm[d];
59 }
60};
61
63{
64 struct Node
65 {
66 Real x, y, z, w;
67 };
68 int nnodes = 0;
70
71 template<typename F>
73 Real
74 operator() (const F& f) const noexcept
75 {
76 Real r = 0;
77 for (int i = 0; i < nnodes; ++i) {
78 r += f(nodes[i].x,nodes[i].y,nodes[i].z) * nodes[i].w;
79 }
80 return r;
81 }
82
83 template<typename F>
85 Real
86 eval (const F& f) const noexcept
87 {
88 Real r = 0;
89 for (int i = 0; i < nnodes; ++i) {
90 r += f(nodes[i].x,nodes[i].y,nodes[i].z) * nodes[i].w;
91 }
92 return r;
93 }
94
95 // evalIntegrand records quadrature nodes when it is invoked by
96 // ImplicitIntegral
98 void
99 evalIntegrand (const GpuArray<Real,3>& x_, Real w_) noexcept
100 {
101 nodes[nnodes].x = x_[0];
102 nodes[nnodes].y = x_[1];
103 nodes[nnodes].z = x_[2];
104 nodes[nnodes].w = w_;
105 ++nnodes;
106 }
107};
108
110namespace detail
111{
112 template <class T>
114 void swap (T& a, T& b) noexcept {
115 T c(a); a=b; b=c;
116 }
117
118 // Determines the sign conditions for restricting a (possibly
119 // already restricted) level set function, i.e., sgn_L and sgn_U
120 // in [R. Saye, High-Order Quadrature Methods for Implicitly
121 // Defined Surfaces and Volumes in Hyperrectangles, SIAM
122 // J. Sci. Comput., Vol. 37, No. 2, pp. A993-A1019,
123 // http://dx.doi.org/10.1137/140966290].
124 template<bool S>
126 void
127 determineSigns (bool positiveAbove, int sign, int& bottomSign, int& topSign) noexcept
128 {
129 // When evaluating a volume integral:
130 // If 0 integrating over the positive part:
131 // if positive above the height function: bottom = +/-, top = +
132 // if positive below the height function: bottom = +, top = +/-
133 // If integrating over the negative part:
134 // if positive above the height function: bottom = -, top = +/-
135 // if positive below the height function: bottom = +/-, top = -
136 // If integrating over both the positive and negative part (i.e. unrestricted in sign), keep it alive
137 if (S)
138 {
139 // When evaluating a surface integral,
140 // if the function is positive above the height function, then
141 // the bottom side must be negative and the top side must be positive;
142 // if the function is positive below the height function,
143 // then the bottom side must be positive and the top side must be negative.
144 bottomSign = positiveAbove? -1 : 1;
145 topSign = positiveAbove? 1 : -1;
146 } else {
147 if (sign == 1)
148 {
149 bottomSign = positiveAbove? 0 : 1;
150 topSign = positiveAbove? 1 : 0;
151 }
152 else if (sign == -1)
153 {
154 bottomSign = positiveAbove? -1 : 0;
155 topSign = positiveAbove? 0 : -1;
156 }
157 else {
158 bottomSign = topSign = 0;
159 }
160 }
161 }
162} // namespace detail
164
165// PsiCode encodes sign information of restricted level set functions
166// on particular sides of a hyperrectangle in a packed array of
167// bits. The first N bits encodes side information, the N+1st bit is
168// true iff the sign == 0, while the N+2nd bit stores the sign if sign
169// != 0.
170template<int N>
172{
173 unsigned char bits = 0;
174
176 PsiCode () noexcept = default;
177
179 PsiCode (const GpuArray<int,N>& sides, int sign) noexcept
180 {
181 static_assert(N <= 3, "algoim::PsiCode: N must be <= 3");
182 for (int dim = 0; dim < N; ++dim) {
183 if (sides[dim] == 1) {
184 bits |= (1 << dim);
185 }
186 }
187 if (sign == 0) {
188 bits |= (1 << N);
189 } else {
190 bits &= ~(1 << N);
191 if (sign == 1) {
192 bits |= (1 << (N+1));
193 }
194 }
195 }
196
197 // Modify an existing code by restriction in a particular dimension.
199 PsiCode (const PsiCode& i, int dim, int side, int sign) noexcept
200 : bits(i.bits)
201 {
202 if (side == 1) {
203 bits |= (1 << dim);
204 }
205 if (sign == 0) {
206 bits |= (1 << N);
207 } else {
208 bits &= ~(1 << N);
209 if (sign == 1) {
210 bits |= (1 << (N+1));
211 }
212 }
213 }
214
216 int side (int dim) const noexcept
217 {
218 return bits & (1 << dim);
219 }
220
222 int sign () const noexcept
223 {
224 return (bits & (1 << N)) ? 0 : ((bits & (1 << (N+1))) ? 1 : -1);
225 }
226};
227
228template <int N>
230{
231 static constexpr Real min (int /*dim*/) noexcept { return -0.5_rt; }
232
233 static constexpr Real max (int /*dim*/) noexcept { return 0.5_rt; }
234
235 static constexpr Real extent (int /*dim*/) noexcept { return 1.0_rt; }
236
237 static constexpr GpuArray<Real,N> midpoint () noexcept {
238 return GpuArray<Real,N>{};
239 }
240
241 static constexpr Real midpoint (int /*dim*/) noexcept { return 0.0_rt; }
242
243 constexpr Real operator() (int side, int /*dim*/) const noexcept {
244 return (side == 0) ? -0.5_rt : 0.5_rt;
245 }
246};
247
249{
250 Real alpha = std::numeric_limits<Real>::lowest();
251};
252
253// M-dimensional integral of an N-dimensional function restricted to
254// given implicitly defined domains
255template<int M, int N, typename Phi, typename F, bool S>
257{
258 const Phi* phi;
259 F* f;
264 int e0 = -1;
266 static constexpr int p = 4;
267
268 // Prune the given set of functions by checking for the existence
269 // of the interface. If a function is uniformly positive or
270 // negative and is consistent with specified sign, it can be
271 // removed. If a function is uniformly positive or negative but
272 // inconsistent with specified sign, the domain of integration is
273 // empty.
275 bool prune () noexcept
276 {
277 static constexpr Real eps = std::numeric_limits<Real>::epsilon();
278 static constexpr Real almostone = 1.0-10.*eps;
279
280 for (int i = 0; i < psiCount; )
281 {
282 GpuArray<Real,N> mid = xrange.midpoint();
283 Real dphi_max = 0.0_rt;
284 for (int dim = 0; dim < N; ++dim) {
285 if (free[dim]) {
286 dphi_max += std::abs(phi->grad(dim));
287 } else {
288 mid[dim] = xrange(psi[i].side(dim),dim);
289 }
290 }
291 dphi_max *= 0.5*almostone;
292 const Real phi_0 = (*phi)(mid);
293 bool uniform_sign = (phi_0 > dphi_max) || (phi_0 < -dphi_max);
294 if (uniform_sign)
295 {
296 if ((phi_0 >= 0.0 && psi[i].sign() >= 0) ||
297 (phi_0 <= 0.0 && psi[i].sign() <= 0) )
298 {
299 --psiCount;
300 detail::swap(psi[i], psi[psiCount]);
301 }
302 else {
303 return false;
304 }
305 }
306 else {
307 ++i;
308 }
309 }
310 return true;
311 }
312
313 // Gaussian quadrature for when the domain of integration is
314 // determined to be the entire M-dimensional cube.
316 void tensorProductIntegral () noexcept
317 {
318 constexpr Real gauss_x[4]={0.069431844202973712388026755553595247452_rt,
319 0.33000947820757186759866712044837765640_rt,
320 0.66999052179242813240133287955162234360_rt,
321 0.93056815579702628761197324444640475255_rt};
322 constexpr Real gauss_w[4]={0.173927422568726928686531974610999703618_rt,
323 0.326072577431273071313468025389000296382_rt,
324 0.326072577431273071313468025389000296382_rt,
325 0.173927422568726928686531974610999703618_rt};
326
327 int nloops = 1;
328 for (int i = 0; i < M; ++i) {
329 nloops *= p;
330 }
331 GpuArray<int,M> i{}; // i is initialized to zero.
332 for (int iloop = 0; iloop < nloops; ++iloop)
333 {
335 Real w = 1.0_rt;
336 for (int dim = 0, k = 0; dim < N; ++dim) {
337 if (free[dim]) {
338 x[dim] = xrange.min(dim) + xrange.extent(dim) * gauss_x[i[k]];
339 w *= xrange.extent(dim) * gauss_w[i[k]];
340 ++k;
341 }
342 }
343 f->evalIntegrand(x, w);
344
345 for (int m = M-1; m >= 0; --m) {
346 ++(i[m]);
347 if (i[m] < p) {
348 break;
349 } else {
350 i[m] = 0;
351 }
352 }
353 }
354 }
355
356 // Given x, valid in all free variables barring e0, root find in
357 // the e0 direction on each of the implicit functions, and apply
358 // Gaussian quadrature to each segment. Weights are multiplied
359 // upon going back up the tree of recursive calls.
361 void evalIntegrand (GpuArray<Real,N> x, Real w) const noexcept
362 {
363 constexpr Real gauss_x[4]={0.069431844202973712388026755553595247452_rt,
364 0.33000947820757186759866712044837765640_rt,
365 0.66999052179242813240133287955162234360_rt,
366 0.93056815579702628761197324444640475255_rt};
367 constexpr Real gauss_w[4]={0.173927422568726928686531974610999703618_rt,
368 0.326072577431273071313468025389000296382_rt,
369 0.326072577431273071313468025389000296382_rt,
370 0.173927422568726928686531974610999703618_rt};
371
372 // Each thread has its own storage for computing roots. These
373 // are not corrupted by the recursive chain of evalIntegrand()
374 // calls since each call is in a different templated
375 // namespace. Moreover, if using OpenMP tasking, each task is
376 // assumed tied and so one thread can only execute
377 // "evalIntegrand" from start to end uninterrupted.
378 GpuArray<Real,4> roots;
379
380 // Surface integral
381 if (S)
382 {
383 AMREX_ASSERT(M == N && N >= 2); // x is thus valid in all variables except e0
384
385 Real x_min = xrange.min(e0);
386 Real x_max = xrange.max(e0);
387 int nroots = 0;
388
389 x[e0] = x_min;
390 Real phi_lo = (*phi)(x);
391 Real xroot = x[e0] - phi_lo / phi->grad(e0);
392 if (xroot > x_min && xroot < x_max) {
393 roots[nroots++] = xroot;
394 }
395
396 AMREX_ASSERT(nroots <= 1);
397
398 for (int i = 0; i < nroots; ++i)
399 {
400 x[e0] = roots[i];
401 Real mag = 0._rt;
402 for (int dim = 0; dim < N; ++dim) {
403 mag += phi->grad(dim)*phi->grad(dim);
404 }
405 mag = std::sqrt(mag);
406 f->evalIntegrand(x, (mag/std::abs(phi->grad(e0)))*w);
407 }
408
409 return;
410 }
411
412 // Partition [xmin(e0), xmax(e0)] by roots found
413 Real x_min = xrange.min(e0);
414 Real x_max = xrange.max(e0);
415 roots[0] = x_min;
416 int nroots = 1;
417 if (phi->grad(e0) != Real(0.0)) {
418 for (int i = 0; i < psiCount; ++i)
419 {
420 for (int dim = 0; dim < N; ++dim) {
421 if (!free[dim]) {
422 x[dim] = xrange(psi[i].side(dim),dim);
423 }
424 }
425 // x is now valid in all variables except e0
426 x[e0] = x_min;
427 Real phi_lo = (*phi)(x);
428 Real xroot = x[e0] - phi_lo / phi->grad(e0);
429 if (xroot > x_min && xroot < x_max) {
430 roots[nroots++] = xroot;
431 }
432 }
433 AMREX_ASSERT(nroots <= 3);
434 }
435 if (nroots == 3 && roots[1] > roots[2]) {
436 detail::swap(roots[1],roots[2]);
437 }
438 roots[nroots++] = x_max;
439
440 // In rare cases, degenerate segments can be found, filter out with a tolerance
441 static constexpr Real eps = std::numeric_limits<Real>::epsilon();
442 constexpr Real tol = 50.0*eps;
443
444 // Loop over segments of divided interval
445 for (int i = 0; i < nroots - 1; ++i)
446 {
447 if (roots[i+1] - roots[i] < tol) { continue; }
448
449 // Evaluate sign of phi within segment and check for consistency with psi
450 bool okay = true;
451 x[e0] = (roots[i] + roots[i+1])*0.5;
452 for (int j = 0; j < psiCount && okay; ++j)
453 {
454 for (int dim = 0; dim < N; ++dim) {
455 if (!free[dim]) {
456 x[dim] = xrange(psi[j].side(dim), dim);
457 }
458 }
459 bool new_ok = ((*phi)(x) > 0.0) ? (psi[j].sign() >= 0) : (psi[j].sign() <= 0);
460 okay = okay && new_ok;
461 }
462 if (!okay) { continue; }
463
464 for (int j = 0; j < p; ++j)
465 {
466 x[e0] = roots[i] + (roots[i+1] - roots[i]) * gauss_x[j];
467 Real gw = (roots[i+1] - roots[i]) * gauss_w[j];
468 f->evalIntegrand(x, w * gw);
469 }
470 }
471 }
472
473 // Main calling engine; parameters with underscores are copied
474 // upon entry but modified internally in the ctor
475 template <int K = M, std::enable_if_t<K==1,int> = 0>
477 ImplicitIntegral (const Phi& phi_, F& f_, const GpuArray<bool,N>& free_,
478 const GpuArray<PsiCode<N>,1 << (N-1)>& psi_,
479 int psiCount_) noexcept
480 : phi(&phi_), f(&f_), free(free_), psi(psi_), psiCount(psiCount_),
481 xrange()
482 {
483 // For the one-dimensional base case, evaluate the
484 // bottom-level integral.
485 for (int dim = 0; dim < N; ++dim) {
486 if (free[dim]) {
487 e0 = dim;
488 }
489 }
491 }
492
493 // Main calling engine; parameters with underscores are copied
494 // upon entry but modified internally in the ctor
495 template <int K = M, std::enable_if_t<(K>1),int> = 0>
497 ImplicitIntegral (const Phi& phi_, F& f_, const GpuArray<bool,N>& free_,
498 const GpuArray<PsiCode<N>,1 << (N-1)>& psi_,
499 int psiCount_) noexcept
500 : phi(&phi_), f(&f_), free(free_), psi(psi_), psiCount(psiCount_),
501 xrange()
502 {
503 // Establish interval bounds for prune() and remaining part of ctor.
504 for (int dim = 0; dim < N; ++dim)
505 {
506 if (free[dim])
507 {
508 xint[dim].alpha = xrange.midpoint(dim);
509 }
510 else
511 {
512 xint[dim].alpha = 0.0_rt;
513 }
514 }
515
516 // Prune list of psi functions: if prune procedure returns
517 // false, then the domain of integration is empty.
518 if (!prune()) {
519 return;
520 }
521
522 // If all psi functions were pruned, then the volumetric
523 // integral domain is the entire hyperrectangle.
524 if (psiCount == 0)
525 {
526 if (!S) {
528 }
529 return;
530 }
531
532 // Among all monotone height function directions, choose the
533 // one that makes the associated height function look as flat
534 // as possible. This is a modification to the criterion
535 // presented in [R. Saye, High-Order Quadrature Methods for
536 // Implicitly Defined Surfaces and Volumes in Hyperrectangles,
537 // SIAM J. Sci. Comput., Vol. 37, No. 2, pp. A993-A1019,
538 // http://dx.doi.org/10.1137/140966290].
539 e0 = -1;
540 {
541 Real gmax = -1.;
542 for (int dim = 0; dim < N; ++dim) {
543 if (free[dim] && std::abs(phi->grad(dim)) > gmax) {
544 gmax = std::abs(phi->grad(dim));
545 e0 = dim;
546 }
547 }
548 }
549
550 // Check compatibility with all implicit functions whilst
551 // simultaneously constructing new implicit functions.
552 GpuArray<PsiCode<N>,1 << (N-1)> newPsi;
553 int newPsiCount = 0;
554 for (int i = 0; i < psiCount; ++i)
555 {
556 // Evaluate gradient in an interval
557 for (int dim = 0; dim < N; ++dim) {
558 if (!free[dim]) {
559 xint[dim].alpha = xrange(psi[i].side(dim), dim);
560 }
561 }
562
563 // w.z. We assume grad does not change.
564 int bottomSign, topSign;
565 detail::determineSigns<S>(phi->grad(e0) > 0.0, psi[i].sign(),
566 bottomSign, topSign);
567 // w.z. There are sides 0 and 1.
568 newPsi[newPsiCount++] = PsiCode<N>(psi[i], e0, 0, bottomSign);
569 newPsi[newPsiCount++] = PsiCode<N>(psi[i], e0, 1, topSign);
570 }
571
572 // Dimension reduction call
573 GpuArray<bool,N> new_free = free;
574 new_free[e0] = false;
576 (*phi, *this, new_free, newPsi, newPsiCount);
577 }
578};
579
580// Partial specialisation on M=0 as a dummy base case for the compiler
581template<int N, typename Phi, typename F, bool S>
582struct ImplicitIntegral<0,N,Phi,F,S>
583{
585 ImplicitIntegral (const Phi&, F&, const GpuArray<bool,N>&,
586 const GpuArray<PsiCode<N>,1 << (N-1)>&,
587 int) noexcept {}
588};
589
591QuadratureRule
592quadGen (EBPlane const& phi) noexcept
593{
595 GpuArray<bool,3> free{true,true,true};
596 GpuArray<PsiCode<3>,4> psi;
597 psi[0] = PsiCode<3>({0,0,0}, -1);
599 return q;
600}
601
603QuadratureRule
604quadGenSurf (EBPlane const& phi) noexcept
605{
607 GpuArray<bool,3> free{true,true,true};
608 GpuArray<PsiCode<3>,4> psi;
609 psi[0] = PsiCode<3>({0,0,0}, -1);
611 return q;
612}
613
615void set_regular (int i, int j, int k, Array4<Real> const& intg) noexcept
616{
617 constexpr Real twelfth = 1._rt/12._rt;
618 constexpr Real offth = 1._rt/144._rt;
619 intg(i,j,k,i_S_x ) = 0.0_rt;
620 intg(i,j,k,i_S_y ) = 0.0_rt;
621 intg(i,j,k,i_S_z ) = 0.0_rt;
622 intg(i,j,k,i_S_x2 ) = twelfth;
623 intg(i,j,k,i_S_y2 ) = twelfth;
624 intg(i,j,k,i_S_z2 ) = twelfth;
625 intg(i,j,k,i_S_x_y ) = 0.0_rt;
626 intg(i,j,k,i_S_x_z ) = 0.0_rt;
627 intg(i,j,k,i_S_y_z ) = 0.0_rt;
628 intg(i,j,k,i_S_x2_y ) = 0.0_rt;
629 intg(i,j,k,i_S_x2_z ) = 0.0_rt;
630 intg(i,j,k,i_S_x_y2 ) = 0.0_rt;
631 intg(i,j,k,i_S_y2_z ) = 0.0_rt;
632 intg(i,j,k,i_S_x_z2 ) = 0.0_rt;
633 intg(i,j,k,i_S_y_z2 ) = 0.0_rt;
634 intg(i,j,k,i_S_x2_y2) = offth;
635 intg(i,j,k,i_S_x2_z2) = offth;
636 intg(i,j,k,i_S_y2_z2) = offth;
637 intg(i,j,k,i_S_xyz ) = 0.0_rt;
638}
639
641void set_regular_surface (int i, int j, int k, Array4<Real> const& sintg) noexcept
642{
643 sintg(i,j,k,i_B_x ) = 0.0_rt;
644 sintg(i,j,k,i_B_y ) = 0.0_rt;
645 sintg(i,j,k,i_B_z ) = 0.0_rt;
646 sintg(i,j,k,i_B_x_y ) = 0.0_rt;
647 sintg(i,j,k,i_B_x_z ) = 0.0_rt;
648 sintg(i,j,k,i_B_y_z ) = 0.0_rt;
649 sintg(i,j,k,i_B_xyz ) = 0.0_rt;
650}
651
652}
653
654#endif
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
#define AMREX_GPU_HOST
Definition AMReX_GpuQualifiers.H:17
amrex_real Real
Floating Point Type for Fields.
Definition AMReX_REAL.H:79
Definition AMReX_algoim.cpp:6
static constexpr int i_B_xyz
Definition AMReX_algoim.H:42
static constexpr int i_S_y_z2
Definition AMReX_algoim.H:26
static constexpr int i_S_y_z
Definition AMReX_algoim.H:20
static constexpr int i_S_x2
Definition AMReX_algoim.H:15
static constexpr int i_S_x_y
Definition AMReX_algoim.H:18
static constexpr int i_S_y
Definition AMReX_algoim.H:13
static constexpr int i_S_x_z
Definition AMReX_algoim.H:19
static constexpr int i_B_z
Definition AMReX_algoim.H:38
__host__ __device__ QuadratureRule quadGenSurf(EBPlane const &phi) noexcept
Definition AMReX_algoim_K.H:604
__host__ __device__ void set_regular(int i, int j, int k, Array4< Real > const &intg) noexcept
Definition AMReX_algoim_K.H:615
static constexpr int i_S_x_y2
Definition AMReX_algoim.H:23
static constexpr int i_S_xyz
Definition AMReX_algoim.H:30
__host__ __device__ QuadratureRule quadGen(EBPlane const &phi) noexcept
Definition AMReX_algoim_K.H:592
static constexpr int i_S_y2_z2
Definition AMReX_algoim.H:29
static constexpr int i_B_x_y
Definition AMReX_algoim.H:39
static constexpr int i_B_x_z
Definition AMReX_algoim.H:40
static constexpr int i_S_x2_y2
Definition AMReX_algoim.H:27
static constexpr int i_S_x2_z2
Definition AMReX_algoim.H:28
static constexpr int i_S_x
Definition AMReX_algoim.H:12
static constexpr int i_S_y2_z
Definition AMReX_algoim.H:24
static constexpr int i_S_x2_y
Definition AMReX_algoim.H:21
static constexpr int i_B_y_z
Definition AMReX_algoim.H:41
static constexpr int i_B_x
Definition AMReX_algoim.H:36
static constexpr int i_S_z
Definition AMReX_algoim.H:14
static constexpr int i_S_z2
Definition AMReX_algoim.H:17
static constexpr int i_S_x_z2
Definition AMReX_algoim.H:25
static constexpr int i_S_x2_z
Definition AMReX_algoim.H:22
static constexpr int i_B_y
Definition AMReX_algoim.H:37
static constexpr int i_S_y2
Definition AMReX_algoim.H:16
__host__ __device__ void set_regular_surface(int i, int j, int k, Array4< Real > const &sintg) noexcept
Definition AMReX_algoim_K.H:641
Definition AMReX_Array4.H:61
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:40
Definition AMReX_algoim_K.H:230
static constexpr Real min(int) noexcept
Definition AMReX_algoim_K.H:231
static constexpr GpuArray< Real, N > midpoint() noexcept
Definition AMReX_algoim_K.H:237
static constexpr Real max(int) noexcept
Definition AMReX_algoim_K.H:233
constexpr Real operator()(int side, int) const noexcept
Definition AMReX_algoim_K.H:243
static constexpr Real midpoint(int) noexcept
Definition AMReX_algoim_K.H:241
static constexpr Real extent(int) noexcept
Definition AMReX_algoim_K.H:235
Definition AMReX_algoim_K.H:24
__host__ __device__ Real grad(int d) const noexcept
Definition AMReX_algoim_K.H:56
__host__ __device__ constexpr EBPlane(GpuArray< Real, 3 > const &c, GpuArray< Real, 3 > const n) noexcept
Definition AMReX_algoim_K.H:32
GpuArray< Real, 3 > cent
Definition AMReX_algoim_K.H:25
__host__ __device__ Real operator()(Real x, Real y, Real z) const noexcept
Definition AMReX_algoim_K.H:43
__host__ __device__ constexpr EBPlane(Real cx, Real cy, Real cz, Real nx, Real ny, Real nz) noexcept
Definition AMReX_algoim_K.H:38
GpuArray< Real, 3 > norm
Definition AMReX_algoim_K.H:26
__host__ __device__ ImplicitIntegral(const Phi &, F &, const GpuArray< bool, N > &, const GpuArray< PsiCode< N >, 1<<(N-1)> &, int) noexcept
Definition AMReX_algoim_K.H:585
Definition AMReX_algoim_K.H:257
static constexpr int p
Definition AMReX_algoim_K.H:266
BoundingBox< N > xrange
Definition AMReX_algoim_K.H:263
F * f
Definition AMReX_algoim_K.H:259
const Phi * phi
Definition AMReX_algoim_K.H:258
__host__ __device__ ImplicitIntegral(const Phi &phi_, F &f_, const GpuArray< bool, N > &free_, const GpuArray< PsiCode< N >, 1<<(N-1)> &psi_, int psiCount_) noexcept
Definition AMReX_algoim_K.H:477
__host__ __device__ bool prune() noexcept
Definition AMReX_algoim_K.H:275
int e0
Definition AMReX_algoim_K.H:264
int psiCount
Definition AMReX_algoim_K.H:262
__host__ __device__ void tensorProductIntegral() noexcept
Definition AMReX_algoim_K.H:316
__host__ __device__ void evalIntegrand(GpuArray< Real, N > x, Real w) const noexcept
Definition AMReX_algoim_K.H:361
GpuArray< bool, N > free
Definition AMReX_algoim_K.H:260
GpuArray< Interval, N > xint
Definition AMReX_algoim_K.H:265
GpuArray< PsiCode< N >, 1<<(N - 1)> psi
Definition AMReX_algoim_K.H:261
Definition AMReX_algoim_K.H:249
Real alpha
Definition AMReX_algoim_K.H:250
Definition AMReX_algoim_K.H:172
__host__ __device__ PsiCode(const PsiCode &i, int dim, int side, int sign) noexcept
Definition AMReX_algoim_K.H:199
__host__ __device__ int side(int dim) const noexcept
Definition AMReX_algoim_K.H:216
__host__ __device__ constexpr PsiCode() noexcept=default
unsigned char bits
Definition AMReX_algoim_K.H:173
__host__ __device__ int sign() const noexcept
Definition AMReX_algoim_K.H:222
Definition AMReX_algoim_K.H:65
Real w
Definition AMReX_algoim_K.H:66
Real z
Definition AMReX_algoim_K.H:66
Real y
Definition AMReX_algoim_K.H:66
Real x
Definition AMReX_algoim_K.H:66
Definition AMReX_algoim_K.H:63
__host__ __device__ Real operator()(const F &f) const noexcept
Definition AMReX_algoim_K.H:74
__host__ Real eval(const F &f) const noexcept
Definition AMReX_algoim_K.H:86
__host__ __device__ void evalIntegrand(const GpuArray< Real, 3 > &x_, Real w_) noexcept
Definition AMReX_algoim_K.H:99
int nnodes
Definition AMReX_algoim_K.H:68
GpuArray< Node, 256 > nodes
Definition AMReX_algoim_K.H:69