1#ifndef AMREX_MLNODELAP_K_H_
2#define AMREX_MLNODELAP_K_H_
3#include <AMReX_Config.H>
15namespace nodelap_detail {
18 constexpr float eps = 1.e-30_rt;
20 constexpr double eps = 1.e-100_rt;
22 constexpr Real almostone =
Real(1.) -
Real(100.)*std::numeric_limits<Real>::epsilon();
23 constexpr Real almostzero =
Real(1.) - almostone;
33 for (
int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
36 if (blo.
smallEnd(idim) == nddom.smallEnd(idim)) {
45 if (bhi.
bigEnd(idim) == nddom.bigEnd(idim)) {
73#if (AMREX_SPACEDIM == 1)
75#elif (AMREX_SPACEDIM == 2)
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);
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);
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);
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; }
#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