Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
AMReX_MLNodeLap_K.H
Go to the documentation of this file.
1#ifndef AMREX_MLNODELAP_K_H_
2#define AMREX_MLNODELAP_K_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_FArrayBox.H>
6#include <AMReX_LO_BCTYPES.H>
7#ifdef AMREX_USE_EB
8#include <AMReX_EBCellFlag.H>
9#endif
10#include <AMReX_MLNodeLinOp_K.H>
11
12namespace amrex {
13
15namespace nodelap_detail {
16
17#if (BL_USE_FLOAT)
18 constexpr float eps = 1.e-30_rt;
19#else
20 constexpr double eps = 1.e-100_rt;
21#endif
22 constexpr Real almostone = Real(1.) - Real(100.)*std::numeric_limits<Real>::epsilon();
23 constexpr Real almostzero = Real(1.) - almostone;
24
25}
27
28inline void
29mlndlap_scale_neumann_bc (Real s, Box const& bx, Array4<Real> const& rhs, Box const& nddom,
31 GpuArray<LinOpBCType,AMREX_SPACEDIM> const& hibc) noexcept
32{
33 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
34 if (lobc[idim] == LinOpBCType::Neumann || lobc[idim] == LinOpBCType::inflow) {
35 Box const& blo = amrex::bdryLo(bx, idim);
36 if (blo.smallEnd(idim) == nddom.smallEnd(idim)) {
38 {
39 rhs(i,j,k) *= s;
40 });
41 }
42 }
43 if (hibc[idim] == LinOpBCType::Neumann || hibc[idim] == LinOpBCType::inflow) {
44 Box const& bhi = amrex::bdryHi(bx, idim);
45 if (bhi.bigEnd(idim) == nddom.bigEnd(idim)) {
47 {
48 rhs(i,j,k) *= s;
49 });
50 }
51 }
52 }
53}
54
55inline void
56mlndlap_impose_neumann_bc (Box const& bx, Array4<Real> const& rhs, Box const& nddom,
58 GpuArray<LinOpBCType,AMREX_SPACEDIM> const& hibc) noexcept
59{
60 mlndlap_scale_neumann_bc(2.0, bx, rhs, nddom, lobc, hibc);
61}
62
63inline void
64mlndlap_unimpose_neumann_bc (Box const& bx, Array4<Real> const& rhs, Box const& nddom,
66 GpuArray<LinOpBCType,AMREX_SPACEDIM> const& hibc) noexcept
67{
68 mlndlap_scale_neumann_bc(0.5, bx, rhs, nddom, lobc, hibc);
69}
70
71}
72
73#if (AMREX_SPACEDIM == 1)
75#elif (AMREX_SPACEDIM == 2)
77#else
79#endif
80
81namespace amrex {
82
84void mlndlap_normalize_sten (int i, int j, int k, Array4<Real> const& x,
85 Array4<Real const> const& sten,
86 Array4<int const> const& msk, Real s0_norm0) noexcept
87{
88 if (!msk(i,j,k) && std::abs(sten(i,j,k,0)) > s0_norm0) {
89 x(i,j,k) /= sten(i,j,k,0);
90 }
91}
92
94void mlndlap_jacobi_sten (int i, int j, int k, Array4<Real> const& sol,
95 Real Ax, Array4<Real const> const& rhs,
96 Array4<Real const> const& sten,
97 Array4<int const> const& msk) noexcept
98{
99 if (msk(i,j,k)) {
100 sol(i,j,k) = Real(0.0);
101 } else if (sten(i,j,k,0) != Real(0.0)) {
102 sol(i,j,k) += Real(2./3.) * (rhs(i,j,k) - Ax) / sten(i,j,k,0);
103 }
104}
105
107void mlndlap_jacobi_sten (Box const& bx, Array4<Real> const& sol,
108 Array4<Real const> const& Ax,
109 Array4<Real const> const& rhs,
110 Array4<Real const> const& sten,
111 Array4<int const> const& msk) noexcept
112{
113 amrex::LoopConcurrent(bx, [=] (int i, int j, int k) noexcept
114 {
115 if (msk(i,j,k)) {
116 sol(i,j,k) = Real(0.0);
117 } else if (sten(i,j,k,0) != Real(0.0)) {
118 sol(i,j,k) += Real(2./3.) * (rhs(i,j,k) - Ax(i,j,k)) / sten(i,j,k,0);
119 }
120 });
121}
122
124bool mlndlap_any_fine_sync_cells (Box const& bx, Array4<int const> const& msk, int fine_flag) noexcept
125{
126 const auto lo = amrex::lbound(bx);
127 const auto hi = amrex::ubound(bx);
128 for (int k = lo.z; k <= hi.z; ++k) {
129 for (int j = lo.y; j <= hi.y; ++j) {
130 for (int i = lo.x; i <= hi.x; ++i) {
131 if (msk(i,j,k) == fine_flag) { return true; }
132 }}}
133 return false;
134}
135
136}
137
138#endif
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_HOST_DEVICE_PARALLEL_FOR_3D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:110
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
__host__ __device__ const IntVectND< dim > & bigEnd() const &noexcept
Return the inclusive upper bound of the box.
Definition AMReX_Box.H:123
__host__ __device__ const IntVectND< dim > & smallEnd() const &noexcept
Return the inclusive lower bound of the box.
Definition AMReX_Box.H:111
amrex_real Real
Floating Point Type for Fields.
Definition AMReX_REAL.H:79
Definition AMReX_Amr.cpp:49
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:319
__host__ __device__ void mlndlap_jacobi_sten(int i, int j, int k, Array4< Real > const &sol, Real Ax, Array4< Real const > const &rhs, Array4< Real const > const &sten, Array4< int const > const &msk) noexcept
Definition AMReX_MLNodeLap_K.H:94
__host__ __device__ void mlndlap_normalize_sten(int i, int j, int k, Array4< Real > const &x, Array4< Real const > const &sten, Array4< int const > const &msk, Real s0_norm0) noexcept
Definition AMReX_MLNodeLap_K.H:84
void mlndlap_unimpose_neumann_bc(Box const &bx, Array4< Real > const &rhs, Box const &nddom, GpuArray< LinOpBCType, 3 > const &lobc, GpuArray< LinOpBCType, 3 > const &hibc) noexcept
Definition AMReX_MLNodeLap_K.H:64
__host__ __device__ BoxND< dim > bdryLo(const BoxND< dim > &b, int dir, int len=1) noexcept
Return the edge-centered BoxND (in direction dir) defining the low side of BoxND b.
Definition AMReX_Box.H:1625
void mlndlap_impose_neumann_bc(Box const &bx, Array4< Real > const &rhs, Box const &nddom, GpuArray< LinOpBCType, 3 > const &lobc, GpuArray< LinOpBCType, 3 > const &hibc) noexcept
Definition AMReX_MLNodeLap_K.H:56
void mlndlap_scale_neumann_bc(Real s, Box const &bx, Array4< Real > const &rhs, Box const &nddom, GpuArray< LinOpBCType, 3 > const &lobc, GpuArray< LinOpBCType, 3 > const &hibc) noexcept
Definition AMReX_MLNodeLap_K.H:29
bool mlndlap_any_fine_sync_cells(Box const &bx, Array4< int const > const &msk, int fine_flag) noexcept
Definition AMReX_MLNodeLap_K.H:124
__host__ __device__ BoxND< dim > bdryHi(const BoxND< dim > &b, int dir, int len=1) noexcept
Return the edge-centered BoxND (in direction dir) defining the high side of BoxND b.
Definition AMReX_Box.H:1648
__host__ __device__ void LoopConcurrent(Dim3 lo, Dim3 hi, F const &f) noexcept
Definition AMReX_Loop.H:152
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:312
Definition AMReX_Array4.H:61
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:40