Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_FFT_Stokes.H
Go to the documentation of this file.
1#ifndef AMREX_FFT_STOKES_H_
2#define AMREX_FFT_STOKES_H_
3
4#include <AMReX_FFT.H>
5#include <AMReX_Geometry.H>
6
7namespace amrex::FFT
8{
9
14template <typename MF = MultiFab>
15class Stokes
16{
17public:
18
19 static_assert(AMREX_SPACEDIM >= 2, "FFT::Stokes requires 2D or 3D");
20
22
23 Stokes () = default;
24
33 Stokes (Geometry const& geom,
34 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> const& bc)
35 requires (IsFabArray_v<MF>)
36 {
37 define(geom, bc);
38 }
39
45 explicit Stokes (Geometry const& geom)
46 requires (IsFabArray_v<MF>)
47 {
48 define(geom);
49 }
50
59 void define (Geometry const& geom,
60 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> const& bc)
61 requires (IsFabArray_v<MF>)
62 {
63 m_domain_lo = geom.Domain().smallEnd();
64 m_geom = detail::shift_geom(geom);
65 m_bc = bc;
66
67 bool all_periodic = true;
68 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
69 all_periodic = all_periodic
70 && (bc[idim].first == Boundary::periodic)
71 && (bc[idim].second == Boundary::periodic);
72 }
73 if (!all_periodic || !m_geom.isAllPeriodic()) {
74 amrex::Abort("FFT::Stokes: only supports periodic BC");
75 }
76
77 m_r2c = std::make_unique<R2C<typename MF::value_type>>(m_geom.Domain());
78 }
79
85 void define (Geometry const& geom)
86 requires (IsFabArray_v<MF>)
87 {
90 std::make_pair(Boundary::periodic,Boundary::periodic))});
91 }
92
114 void solve (AMREX_D_DECL(MF& U, MF& V, MF& W), MF& p,
115 AMREX_D_DECL(MF const& rhsx, MF const& rhsy, MF const& rhsz),
116 typename MF::value_type alpha, typename MF::value_type eta);
117
118private:
119 IntVect m_domain_lo;
120 Geometry m_geom;
121 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> m_bc;
122 std::unique_ptr<R2C<typename MF::value_type>> m_r2c;
123};
124
125
126template <typename MF>
128 MF& V,
129#if (BL_SPACEDIM == 3)
130 MF& W,
131#endif
132 MF& p,
133 MF const& rhsx,
134 MF const& rhsy,
135#if (BL_SPACEDIM == 3)
136 MF const& rhsz,
137#endif
138 typename MF::value_type alpha,
139 typename MF::value_type eta)
140{
141 BL_PROFILE("FFT::Stokes::solve");
142
143 using T = typename MF::value_type;
144
145 if (!m_r2c) {
146 amrex::Abort("FFT::Stokes::solve called before define()");
147 }
148
149 AMREX_ASSERT(p.ixType() == IndexType::TheCellType() &&
151 rhsx.ixType() == U.ixType(),
152 && V.ixType() == IndexType(IntVect::TheDimensionVector(1)) &&
153 rhsy.ixType() == V.ixType(),
154 && W.ixType() == IndexType(IntVect::TheDimensionVector(2)) &&
155 rhsz.ixType() == W.ixType()));
156
157 MF* Umf = &U;
158 MF* Vmf = &V;
159#if (BL_SPACEDIM == 3)
160 MF* Wmf = &W;
161#endif
162 MF* pmf = &p;
163 MF const* rhsxmf = &rhsx;
164 MF const* rhsymf = &rhsy;
165#if (BL_SPACEDIM == 3)
166 MF const* rhszmf = &rhsz;
167#endif
168 MF Utmp, rhsxtmp;
169 MF Vtmp, rhsytmp;
170#if (BL_SPACEDIM == 3)
171 MF Wtmp, rhsztmp;
172#endif
173 MF ptmp;
174 if (m_domain_lo != 0) {
175 detail::shift_mfs(m_domain_lo, U, rhsx, Utmp, rhsxtmp);
176 detail::shift_mfs(m_domain_lo, V, rhsy, Vtmp, rhsytmp);
177#if (BL_SPACEDIM == 3)
178 detail::shift_mfs(m_domain_lo, W, rhsz, Wtmp, rhsztmp);
179#endif
180 detail::shift_mf(m_domain_lo, p, ptmp);
181 Umf = &Utmp;
182 Vmf = &Vtmp;
183#if (BL_SPACEDIM == 3)
184 Wmf = &Wtmp;
185#endif
186 pmf = &ptmp;
187 rhsxmf = &rhsxtmp;
188 rhsymf = &rhsytmp;
189#if (BL_SPACEDIM == 3)
190 rhszmf = &rhsztmp;
191#endif
192 }
193
194 auto& r2c = *m_r2c;
195 auto const& dxinv = m_geom.InvCellSizeArray();
196 auto const scaling = r2c.scalingFactor();
197 auto const& [cba, cdm] = r2c.getSpectralDataLayout();
198
199 cMF phat(cba, cdm, 1, 0);
200
201 cMF rxhat(cba,cdm,1,0);
202 r2c.forward(*rhsxmf, rxhat);
203
204 cMF ryhat(cba,cdm,1,0);
205 r2c.forward(*rhsymf, ryhat);
206
207#if (BL_SPACEDIM == 3)
208 cMF rzhat(cba,cdm,1,0);
209 r2c.forward(*rhszmf, rzhat);
210#endif
211
212 using Complex = GpuComplex<T>;
213 T constexpr tol = std::numeric_limits<T>::epsilon() * T(10);
214 int const nx = m_geom.Domain().length(0);
215#if (AMREX_SPACEDIM >= 2)
216 int const ny = m_geom.Domain().length(1);
217#endif
218#if (AMREX_SPACEDIM == 3)
219 int const nz = m_geom.Domain().length(2);
220#endif
221 for (MFIter mfi(phat); mfi.isValid(); ++mfi) {
222 auto const& pb = phat[mfi].box();
223 auto const& parr = phat[mfi].array();
224 auto const& rx = rxhat[mfi].array();
225 auto const& ry = ryhat[mfi].array();
226#if (BL_SPACEDIM == 3)
227 auto const& rz = rzhat[mfi].array();
228#endif
229 AMREX_D_TERM(T kwx = T(2)*Math::pi<T>()/T(nx);,
230 T kwy = T(2)*Math::pi<T>()/T(ny);,
231 T kwz = T(2)*Math::pi<T>()/T(nz);)
232 ParallelFor(pb, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
233 {
234 AMREX_D_TERM(int ik = (i <= nx/2) ? i : i - nx;,
235 int jk = (j <= ny/2) ? j : j - ny;,
236 int kk = (k <= nz/2) ? k : k - nz);
237
238 GpuArray<T,AMREX_SPACEDIM> kwave{AMREX_D_DECL(ik*kwx,jk*kwy,kk*kwz)};
239
240 T delsqk = AMREX_D_TERM(T(2)*(std::cos(kwave[0])-T(1))*(dxinv[0]*dxinv[0]),
241 + T(2)*(std::cos(kwave[1])-T(1))*(dxinv[1]*dxinv[1]),
242 + T(2)*(std::cos(kwave[2])-T(1))*(dxinv[2]*dxinv[2]));
243
245 {AMREX_D_DECL(Complex((std::cos(kwave[0])-T(1))*dxinv[0],
246 std::sin(kwave[0]) *dxinv[0]),
247 Complex((std::cos(kwave[1])-T(1))*dxinv[1],
248 std::sin(kwave[1]) *dxinv[1]),
249 Complex((std::cos(kwave[2])-T(1))*dxinv[2],
250 std::sin(kwave[2]) *dxinv[2]))};
251
253 {AMREX_D_DECL(Complex((T(1)-std::cos(kwave[0]))*dxinv[0],
254 std::sin(kwave[0]) *dxinv[0]),
255 Complex((T(1)-std::cos(kwave[1]))*dxinv[1],
256 std::sin(kwave[1]) *dxinv[1]),
257 Complex((T(1)-std::cos(kwave[2]))*dxinv[2],
258 std::sin(kwave[2]) *dxinv[2]))};
259
260 AMREX_D_TERM(Complex const rxk = rx(i,j,k);,
261 Complex const ryk = ry(i,j,k);,
262 Complex const rzk = rz(i,j,k);)
263
264 Complex rhsdotdp = scaling * (AMREX_D_TERM(rxk*delkp[0],
265 +ryk*delkp[1],
266 +rzk*delkp[2]));
267
268 if (std::abs(delsqk) > tol) {
269 parr(i,j,k)= rhsdotdp / delsqk;
270 } else {
271 parr(i,j,k)= T(0.);
272 }
273
274 T diffop = alpha - eta*delsqk;
275 if (diffop > tol) {
276 rx(i,j,k) = (scaling*rxk - parr(i,j,k)*delkm[0])/diffop;
277 ry(i,j,k) = (scaling*ryk - parr(i,j,k)*delkm[1])/diffop;
278#if (BL_SPACEDIM == 3)
279 rz(i,j,k) = (scaling*rzk - parr(i,j,k)*delkm[2])/diffop;
280#endif
281 } else {
282 rx(i,j,k) = T(0.);
283 ry(i,j,k) = T(0.);
284#if (BL_SPACEDIM == 3)
285 rz(i,j,k) = T(0.);
286#endif
287 }
288
289 });
290 }
291
292 r2c.backward(phat, *pmf);
293 r2c.backward(rxhat, *Umf);
294 r2c.backward(ryhat, *Vmf);
295#if (BL_SPACEDIM == 3)
296 r2c.backward(rzhat, *Wmf);
297#endif
298}
299
300} // namespace amrex::FFT
301
302#endif
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:551
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:172
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
#define BL_SPACEDIM
Definition AMReX_SPACE.H:15
Stokes solver for periodic domains using FFT.
Definition AMReX_FFT_Stokes.H:16
Stokes(Geometry const &geom, Array< std::pair< Boundary, Boundary >, 3 > const &bc)
Construct a Stokes solver with explicit boundary types.
Definition AMReX_FFT_Stokes.H:33
void define(Geometry const &geom)
Define a purely periodic Stokes solver.
Definition AMReX_FFT_Stokes.H:85
Stokes(Geometry const &geom)
Construct a purely periodic Stokes solver.
Definition AMReX_FFT_Stokes.H:45
void define(Geometry const &geom, Array< std::pair< Boundary, Boundary >, 3 > const &bc)
Define a Stokes solver with explicit boundary types.
Definition AMReX_FFT_Stokes.H:59
void solve(MF &U, MF &V, MF &W, MF &p, MF const &rhsx, MF const &rhsy, MF const &rhsz, typename MF::value_type alpha, typename MF::value_type eta)
Solve the generalized Stokes problem in spectral space.
Definition AMReX_FFT_Stokes.H:127
Box box(int K) const noexcept
Return the Kth Box in the BoxArray. That is, the valid region of the Kth grid.
Definition AMReX_FabArrayBase.H:101
Array4< typename FabArray< FAB >::value_type const > array(const MFIter &mfi) const noexcept
Definition AMReX_FabArray.H:561
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:75
const Box & Domain() const noexcept
Returns our rectangular domain.
Definition AMReX_Geometry.H:216
bool isAllPeriodic() const noexcept
Is domain periodic in all directions?
Definition AMReX_Geometry.H:344
__host__ __device__ constexpr CellIndex ixType(int dir) const noexcept
Returns the CellIndex in direction dir.
Definition AMReX_IndexType.H:117
__host__ static __device__ constexpr IndexTypeND< dim > TheCellType() noexcept
This static member function returns an IndexTypeND object of value IndexTypeND::CELL....
Definition AMReX_IndexType.H:150
__host__ static __device__ constexpr IntVectND< dim > TheDimensionVector(int d) noexcept
This static member function returns a reference to a constant IntVectND object, all of whose dim argu...
Definition AMReX_IntVect.H:790
Iterator for looping ever tiles and boxes of amrex::FabArray based containers.
Definition AMReX_MFIter.H:88
bool isValid() const noexcept
Is the iterator valid i.e. is it associated with a FAB?
Definition AMReX_MFIter.H:172
std::array< T, N > Array
Definition AMReX_Array.H:31
Definition AMReX_FFT_Helper.H:53
void ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition AMReX_CTOParallelForImpl.H:202
double second() noexcept
Definition AMReX_Utility.cpp:919
void Abort(const std::string &msg)
Print a fatal-error message to stderr and abort execution.
Definition AMReX.cpp:241
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:52
A host / device complex number type, because std::complex doesn't work in device code with Cuda yet.
Definition AMReX_GpuComplex.H:30