1#ifndef AMREX_FFT_STOKES_H_
2#define AMREX_FFT_STOKES_H_
14template <
typename MF = MultiFab>
19 static_assert(AMREX_SPACEDIM >= 2,
"FFT::Stokes requires 2D or 3D");
34 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM>
const& bc)
35 requires (IsFabArray_v<MF>)
46 requires (IsFabArray_v<MF>)
60 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM>
const& bc)
61 requires (IsFabArray_v<MF>)
63 m_domain_lo = geom.Domain().smallEnd();
64 m_geom = detail::shift_geom(geom);
67 bool all_periodic =
true;
68 for (
int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
69 all_periodic = all_periodic
77 m_r2c = std::make_unique<R2C<typename MF::value_type>>(m_geom.
Domain());
86 requires (IsFabArray_v<MF>)
115 AMREX_D_DECL(MF
const& rhsx, MF
const& rhsy, MF
const& rhsz),
116 typename MF::value_type alpha,
typename MF::value_type eta);
122 std::unique_ptr<R2C<typename MF::value_type>> m_r2c;
126template <
typename MF>
138 typename MF::value_type alpha,
139 typename MF::value_type eta)
143 using T =
typename MF::value_type;
146 amrex::Abort(
"FFT::Stokes::solve called before define()");
151 rhsx.
ixType() == U.ixType(),
153 rhsy.
ixType() == V.ixType(),
155 rhsz.
ixType() == W.ixType()));
159#if (BL_SPACEDIM == 3)
163 MF
const* rhsxmf = &rhsx;
164 MF
const* rhsymf = &rhsy;
165#if (BL_SPACEDIM == 3)
166 MF
const* rhszmf = &rhsz;
170#if (BL_SPACEDIM == 3)
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);
180 detail::shift_mf(m_domain_lo, p, ptmp);
183#if (BL_SPACEDIM == 3)
189#if (BL_SPACEDIM == 3)
195 auto const& dxinv = m_geom.InvCellSizeArray();
196 auto const scaling = r2c.scalingFactor();
197 auto const& [cba, cdm] = r2c.getSpectralDataLayout();
199 cMF phat(cba, cdm, 1, 0);
201 cMF rxhat(cba,cdm,1,0);
202 r2c.forward(*rhsxmf, rxhat);
204 cMF ryhat(cba,cdm,1,0);
205 r2c.forward(*rhsymf, ryhat);
207#if (BL_SPACEDIM == 3)
208 cMF rzhat(cba,cdm,1,0);
209 r2c.forward(*rhszmf, rzhat);
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);
218#if (AMREX_SPACEDIM == 3)
219 int const nz = m_geom.Domain().length(2);
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();
230 T kwy = T(2)*Math::pi<T>()/T(ny);,
231 T kwz = T(2)*Math::pi<T>()/T(nz);)
235 int jk = (j <= ny/2) ? j : j - ny;,
236 int kk = (k <= nz/2) ? k : k - nz);
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]));
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]))};
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]))};
261 Complex
const ryk = ry(i,j,k);,
262 Complex
const rzk = rz(i,j,k);)
264 Complex rhsdotdp = scaling * (
AMREX_D_TERM(rxk*delkp[0],
268 if (std::abs(delsqk) > tol) {
269 parr(i,j,k)= rhsdotdp / delsqk;
274 T diffop = alpha - eta*delsqk;
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;
284#if (BL_SPACEDIM == 3)
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);
#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