1#ifndef AMREX_MLNODETENSORLAP_2D_K_H_
2#define AMREX_MLNODETENSORLAP_2D_K_H_
3#include <AMReX_Config.H>
8namespace mlndts_detail {
11 Real ts_interp_line_x (Array4<Real const>
const&
crse,
int ic,
int jc)
noexcept
17 Real ts_interp_line_y (Array4<Real const>
const&
crse,
int ic,
int jc)
noexcept
23 Real ts_interp_face_xy (Array4<Real const>
const&
crse,
int ic,
int jc)
noexcept
25 return (ts_interp_line_y(
crse,ic ,jc ) +
26 ts_interp_line_y(
crse,ic+1,jc ) +
27 ts_interp_line_x(
crse,ic ,jc ) +
28 ts_interp_line_x(
crse,ic ,jc+1)) *
Real(0.25);
35 Array4<Real const>
const&
crse, Array4<int const>
const& msk)
noexcept
37 using namespace mlndts_detail;
42 bool i_is_odd = (ic*2 != i);
43 bool j_is_odd = (jc*2 != j);
44 if (i_is_odd && j_is_odd) {
46 fine(i,j,0) += ts_interp_face_xy(
crse,ic,jc);
47 }
else if (i_is_odd) {
49 fine(i,j,0) += ts_interp_line_x(
crse,ic,jc);
50 }
else if (j_is_odd) {
52 fine(i,j,0) += ts_interp_line_y(
crse,ic,jc);
62 Array4<Real const>
const&
crse, Array4<int const>
const& msk,
63 int semi_dir)
noexcept
65 using namespace mlndts_detail;
70 bool j_is_odd = (jc*2 != j);
73 fine(i,j,0) += ts_interp_line_y(
crse,i,jc);
80 bool i_is_odd = (ic*2 != i);
83 fine(i,j,0) += ts_interp_line_x(
crse,ic,j);
99 y(i,j,k) = s[0] * (
x(i-1,j,0) +
x(i+1,j,0))
100 + s[2] * (
x(i,j-1,0) +
x(i,j+1,0))
101 -
Real(2.)*(s[0]+s[2]) *
x(i,j,0)
102 +
Real(0.5)*s[1] * (
x(i-1,j-1,0) +
x(i+1,j+1,0) -
x(i-1,j+1,0) -
x(i+1,j-1,0));
116 Real Ax = s[0] * (sol(i-1,j,0) + sol(i+1,j,0))
117 + s[2] * (sol(i,j-1,0) + sol(i,j+1,0))
119 +
Real(0.5)*s[1] * (sol(i-1,j-1,0) + sol(i+1,j+1,0) - sol(i-1,j+1,0) - sol(i+1,j-1,0));
120 sol(i,j,k) += (rhs(i,j,k) - Ax) * (omega/s0);
124#if defined(AMREX_USE_HYPRE) && (AMREX_SPACEDIM > 1)
126template <
typename HypreInt,
typename AtomicInt>
127void mlndtslap_fill_ijmatrix_cpu (
Box const& ndbx,
128 Array4<AtomicInt const>
const& gid,
129 Array4<int const>
const& lid,
130 HypreInt*
const ncols, HypreInt*
const cols,
Real*
const mat,
131 GpuArray<Real,3>
const& s)
noexcept
133 constexpr auto gidmax = std::numeric_limits<AtomicInt>::max();
139 HypreInt nelems_old = nelems;
141 cols[nelems] = gid(i,j,k);
142 mat[nelems] =
Real(-2.)*(s[0]+s[2]);
145 if (gid(i-1,j-1,k) < gidmax) {
146 cols[nelems] = gid(i-1,j-1,k);
147 mat[nelems] =
Real(0.5)*s[1];
151 if (gid(i,j-1,k) < gidmax) {
152 cols[nelems] = gid(i,j-1,k);
157 if (gid(i+1,j-1,k) < gidmax) {
158 cols[nelems] = gid(i+1,j-1,k);
159 mat[nelems] =
Real(-0.5)*s[1];
163 if (gid(i-1,j,k) < gidmax) {
164 cols[nelems] = gid(i-1,j,k);
169 if (gid(i+1,j,k) < gidmax) {
170 cols[nelems] = gid(i+1,j,k);
175 if (gid(i-1,j+1,k) < gidmax) {
176 cols[nelems] = gid(i-1,j+1,k);
177 mat[nelems] =
Real(-0.5)*s[1];
181 if (gid(i,j+1,k) < gidmax) {
182 cols[nelems] = gid(i,j+1,k);
187 if (gid(i+1,j+1,k) < gidmax) {
188 cols[nelems] = gid(i+1,j+1,k);
189 mat[nelems] =
Real(0.5)*s[1];
193 ncols[lid(i,j,k)] = nelems - nelems_old;
199template <
typename HypreInt,
typename AtomicInt>
201void mlndtslap_fill_ijmatrix_gpu (
const int ps,
const int i,
const int j,
const int k,
203 Array4<AtomicInt const>
const& gid,
204 Array4<int const>
const& lid,
205 HypreInt*
const ncols, HypreInt*
const cols,
Real*
const mat,
206 GpuArray<Real,3>
const& s)
noexcept
210 constexpr auto gidmax = std::numeric_limits<AtomicInt>::max();
213 cols[ps] = gid(i,j,k);
214 mat[ps] =
Real(-2.)*(s[0]+s[2]);
216 if (gid(i-1,j-1,k) < gidmax) { ++nc; }
217 if (gid(i ,j-1,k) < gidmax) { ++nc; }
218 if (gid(i+1,j-1,k) < gidmax) { ++nc; }
219 if (gid(i-1,j ,k) < gidmax) { ++nc; }
220 if (gid(i+1,j ,k) < gidmax) { ++nc; }
221 if (gid(i-1,j+1,k) < gidmax) { ++nc; }
222 if (gid(i ,j+1,k) < gidmax) { ++nc; }
223 if (gid(i+1,j+1,k) < gidmax) { ++nc; }
224 ncols[lid(i,j,k)] = nc;
226 else if (
offset == 1 && gid(i-1,j-1,k) < gidmax) {
227 cols[ps] = gid(i-1,j-1,k);
228 mat[ps] =
Real(0.5)*s[1];
230 else if (
offset == 2 && gid(i ,j-1,k) < gidmax) {
231 cols[ps] = gid(i ,j-1,k);
234 else if (
offset == 3 && gid(i+1,j-1,k) < gidmax) {
235 cols[ps] = gid(i+1,j-1,k);
236 mat[ps] =
Real(-0.5)*s[1];
238 else if (
offset == 4 && gid(i-1,j ,k) < gidmax) {
239 cols[ps] = gid(i-1,j ,k);
242 else if (
offset == 5 && gid(i+1,j ,k) < gidmax) {
243 cols[ps] = gid(i+1,j ,k);
246 else if (
offset == 6 && gid(i-1,j+1,k) < gidmax) {
247 cols[ps] = gid(i-1,j+1,k);
248 mat[ps] =
Real(-0.5)*s[1];
250 else if (
offset == 7 && gid(i ,j+1,k) < gidmax) {
251 cols[ps] = gid(i ,j+1,k);
254 else if (
offset == 8 && gid(i+1,j+1,k) < gidmax) {
255 cols[ps] = gid(i+1,j+1,k);
256 mat[ps] =
Real(0.5)*s[1];
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
Array4< int const > offset
Definition AMReX_HypreMLABecLap.cpp:1089
Array4< Real > fine
Definition AMReX_InterpFaceRegister.cpp:90
Array4< Real const > crse
Definition AMReX_InterpFaceRegister.cpp:92
amrex_real Real
Floating Point Type for Fields.
Definition AMReX_REAL.H:79
__host__ __device__ BoxND< dim > coarsen(const BoxND< dim > &b, int ref_ratio) noexcept
Coarsen BoxND by given (positive) coarsening ratio.
Definition AMReX_Box.H:1409
Definition AMReX_Amr.cpp:49
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlndtslap_adotx(int i, int j, int k, Array4< Real > const &y, Array4< Real const > const &x, Array4< int const > const &msk, GpuArray< Real, 3 > const &s) noexcept
Definition AMReX_MLNodeTensorLap_2D_K.H:93
BoxND< 3 > Box
Box is an alias for amrex::BoxND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:27
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlndtslap_gauss_seidel(int i, int j, int k, Array4< Real > const &sol, Array4< Real const > const &rhs, Array4< int const > const &msk, GpuArray< Real, 3 > const &s) noexcept
Definition AMReX_MLNodeTensorLap_2D_K.H:107
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlndtslap_semi_interpadd(int, int, int, Array4< Real > const &, Array4< Real const > const &, Array4< int const > const &, int) noexcept
Definition AMReX_MLNodeTensorLap_1D_K.H:13
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlndtslap_interpadd(int, int, int, Array4< Real > const &, Array4< Real const > const &, Array4< int const > const &) noexcept
Definition AMReX_MLNodeTensorLap_1D_K.H:8
void LoopOnCpu(Dim3 lo, Dim3 hi, F const &f) noexcept
Definition AMReX_Loop.H:365
Definition AMReX_Array4.H:61
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:40