Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
AMReX_MLNodeTensorLap_2D_K.H
Go to the documentation of this file.
1#ifndef AMREX_MLNODETENSORLAP_2D_K_H_
2#define AMREX_MLNODETENSORLAP_2D_K_H_
3#include <AMReX_Config.H>
4
5namespace amrex {
6
8namespace mlndts_detail {
9
11 Real ts_interp_line_x (Array4<Real const> const& crse, int ic, int jc) noexcept
12 {
13 return (crse(ic,jc,0)+crse(ic+1,jc,0))*Real(0.5);
14 }
15
17 Real ts_interp_line_y (Array4<Real const> const& crse, int ic, int jc) noexcept
18 {
19 return (crse(ic,jc,0)+crse(ic,jc+1,0))*Real(0.5);
20 }
21
23 Real ts_interp_face_xy (Array4<Real const> const& crse, int ic, int jc) noexcept
24 {
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);
29 }
30}
32
34void mlndtslap_interpadd (int i, int j, int, Array4<Real> const& fine,
35 Array4<Real const> const& crse, Array4<int const> const& msk) noexcept
36{
37 using namespace mlndts_detail;
38
39 if (!msk(i,j,0)) {
40 int ic = amrex::coarsen(i,2);
41 int jc = amrex::coarsen(j,2);
42 bool i_is_odd = (ic*2 != i);
43 bool j_is_odd = (jc*2 != j);
44 if (i_is_odd && j_is_odd) {
45 // Node on a X-Y face
46 fine(i,j,0) += ts_interp_face_xy(crse,ic,jc);
47 } else if (i_is_odd) {
48 // Node on X line
49 fine(i,j,0) += ts_interp_line_x(crse,ic,jc);
50 } else if (j_is_odd) {
51 // Node on Y line
52 fine(i,j,0) += ts_interp_line_y(crse,ic,jc);
53 } else {
54 // Node coincident with coarse node
55 fine(i,j,0) += crse(ic,jc,0);
56 }
57 }
58}
59
61void mlndtslap_semi_interpadd (int i, int j, int, Array4<Real> const& fine,
62 Array4<Real const> const& crse, Array4<int const> const& msk,
63 int semi_dir) noexcept
64{
65 using namespace mlndts_detail;
66
67 if (!msk(i,j,0)) {
68 if (semi_dir == 0) {
69 int jc = amrex::coarsen(j,2);
70 bool j_is_odd = (jc*2 != j);
71 if (j_is_odd) {
72 // Node on Y line
73 fine(i,j,0) += ts_interp_line_y(crse,i,jc);
74 } else {
75 // Node coincident with coarse node
76 fine(i,j,0) += crse(i,jc,0);
77 }
78 } else {
79 int ic = amrex::coarsen(i,2);
80 bool i_is_odd = (ic*2 != i);
81 if (i_is_odd) {
82 // Node on X line
83 fine(i,j,0) += ts_interp_line_x(crse,ic,j);
84 } else {
85 // Node coincident with coarse node
86 fine(i,j,0) += crse(ic,j,0);
87 }
88 }
89 }
90}
91
93void mlndtslap_adotx (int i, int j, int k, Array4<Real> const& y, Array4<Real const> const& x,
94 Array4<int const> const& msk, GpuArray<Real,3> const& s) noexcept
95{
96 if (msk(i,j,k)) {
97 y(i,j,k) = Real(0.0);
98 } else {
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));
103 }
104}
105
107void mlndtslap_gauss_seidel (int i, int j, int k, Array4<Real> const& sol,
108 Array4<Real const> const& rhs, Array4<int const> const& msk,
109 GpuArray<Real,3> const& s) noexcept
110{
111 if (msk(i,j,k)) {
112 sol(i,j,k) = 0.0;
113 } else {
114 constexpr Real omega = Real(1.25);
115 Real s0 = Real(-2.)*(s[0]+s[2]);
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))
118 + s0 * sol(i,j,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);
121 }
122}
123
124#if defined(AMREX_USE_HYPRE) && (AMREX_SPACEDIM > 1)
125
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, // NOLINT(readability-non-const-parameter)
131 GpuArray<Real,3> const& s) noexcept
132{
133 constexpr auto gidmax = std::numeric_limits<AtomicInt>::max();
134 HypreInt nelems = 0;
135 amrex::LoopOnCpu(ndbx, [&] (int i, int j, int k) noexcept
136 {
137 if (lid(i,j,k) >= 0)
138 {
139 HypreInt nelems_old = nelems;
140
141 cols[nelems] = gid(i,j,k);
142 mat[nelems] = Real(-2.)*(s[0]+s[2]);
143 ++nelems;
144
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];
148 ++nelems;
149 }
150
151 if (gid(i,j-1,k) < gidmax) {
152 cols[nelems] = gid(i,j-1,k);
153 mat[nelems] = s[2];
154 ++nelems;
155 }
156
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];
160 ++nelems;
161 }
162
163 if (gid(i-1,j,k) < gidmax) {
164 cols[nelems] = gid(i-1,j,k);
165 mat[nelems] = s[0];
166 ++nelems;
167 }
168
169 if (gid(i+1,j,k) < gidmax) {
170 cols[nelems] = gid(i+1,j,k);
171 mat[nelems] = s[0];
172 ++nelems;
173 }
174
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];
178 ++nelems;
179 }
180
181 if (gid(i,j+1,k) < gidmax) {
182 cols[nelems] = gid(i,j+1,k);
183 mat[nelems] = s[2];
184 ++nelems;
185 }
186
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];
190 ++nelems;
191 }
192
193 ncols[lid(i,j,k)] = nelems - nelems_old;
194 }
195 });
196}
197
198#ifdef AMREX_USE_GPU
199template <typename HypreInt, typename AtomicInt>
201void mlndtslap_fill_ijmatrix_gpu (const int ps, const int i, const int j, const int k,
202 const int offset, Box const& ndbx,
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
207{
208 if (lid(i,j,k) >= 0)
209 {
210 constexpr auto gidmax = std::numeric_limits<AtomicInt>::max();
211
212 if (offset == 0) {
213 cols[ps] = gid(i,j,k);
214 mat[ps] = Real(-2.)*(s[0]+s[2]);
215 int nc = 1;
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;
225 }
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];
229 }
230 else if (offset == 2 && gid(i ,j-1,k) < gidmax) {
231 cols[ps] = gid(i ,j-1,k);
232 mat[ps] = s[2];
233 }
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];
237 }
238 else if (offset == 4 && gid(i-1,j ,k) < gidmax) {
239 cols[ps] = gid(i-1,j ,k);
240 mat[ps] = s[0];
241 }
242 else if (offset == 5 && gid(i+1,j ,k) < gidmax) {
243 cols[ps] = gid(i+1,j ,k);
244 mat[ps] = s[0];
245 }
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];
249 }
250 else if (offset == 7 && gid(i ,j+1,k) < gidmax) {
251 cols[ps] = gid(i ,j+1,k);
252 mat[ps] = s[2];
253 }
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];
257 }
258 }
259}
260#endif
261
262#endif
263
264}
265
266#endif
#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