Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_FFT_OpenBCSolver.H
Go to the documentation of this file.
1#ifndef AMREX_FFT_OPENBC_SOLVER_H_
2#define AMREX_FFT_OPENBC_SOLVER_H_
3
4#include <AMReX_FFT_R2C.H>
5
6namespace amrex::FFT
7{
8
24template <typename T = Real>
26{
27public:
28 using MF = typename R2C<T>::MF;
29 using cMF = typename R2C<T>::cMF;
30
37 explicit OpenBCSolver (Box const& domain, Info const& info = Info{});
38
45 template <class F>
46 void setGreensFunction (F const& greens_function);
47
54 void solve (MF& phi, MF const& rho);
55
61 [[nodiscard]] Box const& Domain () const { return m_domain; }
62
68 [[nodiscard]] IntVect const& PaddedLength () const { return m_padded_length; }
69
84 [[nodiscard]] cMF& greensFunctionFFT () noexcept { return m_G_fft; }
85
91 [[nodiscard]] cMF const& greensFunctionFFT () const noexcept { return m_G_fft; }
92
93private:
94 static IntVect make_padded_length (Box const& domain, Info const& info);
95 static Box make_grown_domain (Box const& domain, IntVect const& padded_len,
96 Info const& info);
97
98 Box m_domain;
99 Info m_info;
100 IntVect m_padded_length;
101 R2C<T> m_r2c;
102 cMF m_G_fft;
103 std::unique_ptr<R2C<T>> m_r2c_green;
104};
105
106template <typename T>
107IntVect OpenBCSolver<T>::make_padded_length (Box const& domain, Info const& info)
108{
109 IntVect len = domain.length();
110 int ndims = AMREX_SPACEDIM;
111#if (AMREX_SPACEDIM == 3)
112 if (info.twod_mode) { ndims = 2; }
113#else
115#endif
116 if (info.openbc_padding) {
117 for (int idim = 0; idim < ndims; ++idim) {
118 len[idim] = FFT::nextFastLen(len[idim], info.openbc_padding_nfactors);
119 }
120 }
121 return len;
122}
123
124template <typename T>
125Box OpenBCSolver<T>::make_grown_domain (Box const& domain, IntVect const& padded_len,
126 Info const& info)
127{
128 IntVect len = padded_len;
129 int ndims = AMREX_SPACEDIM;
130#if (AMREX_SPACEDIM == 3)
131 if (info.twod_mode) { ndims = 2; }
132#else
134#endif
135 for (int idim = 0; idim < ndims; ++idim) {
137 len[idim] <= std::numeric_limits<int>::max()/2,
138 "FFT::OpenBCSolver: padded domain length exceeds int range");
139 len[idim] *= 2;
140 }
141 return Box(domain.smallEnd(), domain.smallEnd()+len-IntVect(1), domain.ixType());
142}
143
144template <typename T>
145OpenBCSolver<T>::OpenBCSolver (Box const& domain, Info const& info)
146 : m_domain(domain),
147 m_info(info),
148 m_padded_length(OpenBCSolver<T>::make_padded_length(domain, info)),
149 m_r2c(OpenBCSolver<T>::make_grown_domain(domain, m_padded_length, info),
150 m_info.setDomainStrategy(FFT::DomainStrategy::slab))
151{
153 "FFT::OpenBCSolver does not support FFT::Info::batch_size > 1");
154
155#if (AMREX_SPACEDIM == 3)
156 if (m_info.twod_mode) {
157 auto gdom = make_grown_domain(domain, m_padded_length, m_info);
158 gdom.enclosedCells(2);
159 gdom.setSmall(2, 0);
160 int nprocs = std::min({ParallelContext::NProcsSub(),
161 m_info.nprocs,
162 m_domain.length(2)});
163 gdom.setBig(2, nprocs-1);
164 m_r2c_green = std::make_unique<R2C<T>>(gdom,m_info);
165 auto [sd, ord] = m_r2c_green->getSpectralData();
166 m_G_fft = cMF(*sd, amrex::make_alias, 0, 1);
167 } else
168#endif
169 {
170 amrex::ignore_unused(m_r2c_green);
171 auto [sd, ord] = m_r2c.getSpectralData();
173 m_G_fft.define(sd->boxArray(), sd->DistributionMap(), 1, 0);
174 }
175}
176
177template <typename T>
178template <class F>
179void OpenBCSolver<T>::setGreensFunction (F const& greens_function)
180{
181 BL_PROFILE("OpenBCSolver::setGreensFunction");
182
183 auto* infab = m_info.twod_mode ? detail::get_fab(m_r2c_green->m_rx)
184 : detail::get_fab(m_r2c.m_rx);
185 auto lo = m_domain.smallEnd();
186#if (AMREX_SPACEDIM == 3)
187 // m_r2c_green's domain starts at z = 0.
188 if (m_info.twod_mode) { lo[2] = 0; }
189#endif
190 auto const& lo3 = lo.dim3();
191 // Hidden dimensions must have length 1, not 0, so that the mirror plane
192 // checks below cannot fire on them.
193 auto const len3d = m_padded_length.dim3(1);
194 GpuArray<int,3> len{len3d.x, len3d.y, len3d.z};
195 if (infab) {
196 auto const& a = infab->array();
197 auto box = infab->box();
198 GpuArray<int,3> nimages{1,1,1};
199 int ndims = m_info.twod_mode ? AMREX_SPACEDIM-1 : AMREX_SPACEDIM;
200 for (int idim = 0; idim < ndims; ++idim) {
201 if (box.smallEnd(idim) == lo[idim] && box.length(idim) == 2*len[idim]) {
202 box.growHi(idim, -len[idim]+1); // +1 to include the middle plane
203 nimages[idim] = 2;
204 }
205 }
206 AMREX_ASSERT(nimages[0] == 2);
207 box.shift(-lo);
208 amrex::ParallelForOMP(box, [=] AMREX_GPU_DEVICE (int i, int j, int k)
209 {
210 T G;
211 if (i == len[0] || j == len[1] || k == len[2]) {
212 G = 0;
213 } else {
214 auto ii = i;
215 auto jj = (j > len[1]) ? 2*len[1]-j : j;
216 auto kk = (k > len[2]) ? 2*len[2]-k : k;
217 G = greens_function(ii+lo3.x,jj+lo3.y,kk+lo3.z);
218 }
219 for (int koff = 0; koff < nimages[2]; ++koff) {
220 int k2 = (koff == 0) ? k : 2*len[2]-k;
221 if ((k2 == 2*len[2]) || (koff == 1 && k == len[2])) {
222 continue;
223 }
224 for (int joff = 0; joff < nimages[1]; ++joff) {
225 int j2 = (joff == 0) ? j : 2*len[1]-j;
226 if ((j2 == 2*len[1]) || (joff == 1 && j == len[1])) {
227 continue;
228 }
229 for (int ioff = 0; ioff < nimages[0]; ++ioff) {
230 int i2 = (ioff == 0) ? i : 2*len[0]-i;
231 if ((i2 == 2*len[0]) || (ioff == 1 && i == len[0])) {
232 continue;
233 }
234 a(i2+lo3.x,j2+lo3.y,k2+lo3.z) = G;
235 }
236 }
237 }
238 });
239 }
240
241 if (m_info.twod_mode) {
242 m_r2c_green->forward(m_r2c_green->m_rx);
243 } else {
244 m_r2c.forward(m_r2c.m_rx);
245 }
246
247 if (!m_info.twod_mode) {
248 auto [sd, ord] = m_r2c.getSpectralData();
250 auto const* srcfab = detail::get_fab(*sd);
251 if (srcfab) {
252 auto* dstfab = detail::get_fab(m_G_fft);
253 if (dstfab) {
254 Gpu::dtod_memcpy_async(dstfab->dataPtr(), srcfab->dataPtr(), dstfab->nBytes());
255 } else {
256 amrex::Abort("FFT::OpenBCSolver: how did this happen");
257 }
258 }
259
260 m_r2c.prepare_openbc();
261 }
262}
263
264template <typename T>
265void OpenBCSolver<T>::solve (MF& phi, MF const& rho)
266{
267 BL_PROFILE("OpenBCSolver::solve");
268
269 auto& inmf = m_r2c.m_rx;
270 inmf.setVal(T(0));
271 inmf.ParallelCopy(rho, 0, 0, 1);
272
273 m_r2c.m_openbc_half = (AMREX_SPACEDIM == 3) && !m_info.twod_mode;
274 m_r2c.forward(inmf);
275 m_r2c.m_openbc_half = false;
276
277 auto scaling_factor = m_r2c.scalingFactor();
278
279 auto const* gfab = detail::get_fab(m_G_fft);
280 if (gfab) {
281 auto [sd, ord] = m_r2c.getSpectralData();
283 auto* rhofab = detail::get_fab(*sd);
284 if (rhofab) {
285 auto* pdst = rhofab->dataPtr();
286 auto const* psrc = gfab->dataPtr();
287 Box const& rhobox = rhofab->box();
288#if (AMREX_SPACEDIM == 3)
289 Long leng = gfab->box().numPts();
290 if (m_info.twod_mode) {
291 AMREX_ASSERT(gfab->box().length(2) == 1 &&
292 leng == (rhobox.length(0) * rhobox.length(1)));
293 } else {
294 AMREX_ASSERT(leng == rhobox.numPts());
295 }
296#endif
298 {
299#if (AMREX_SPACEDIM == 3)
300 Long isrc = i % leng;
301#else
302 Long isrc = i;
303#endif
304 pdst[i] *= psrc[isrc] * scaling_factor;
305 });
306 } else {
307 amrex::Abort("FFT::OpenBCSolver::solve: how did this happen?");
308 }
309 }
310
311 m_r2c.m_openbc_half = (AMREX_SPACEDIM == 3) && !m_info.twod_mode;
312 m_r2c.backward_doit(phi, phi.nGrowVect());
313 m_r2c.m_openbc_half = false;
314}
315
316}
317
318#endif
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:562
#define AMREX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition AMReX_BLassert.H:49
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
Real * pdst
Definition AMReX_HypreMLABecLap.cpp:1132
__host__ __device__ Long numPts() const noexcept
Return the number of points contained in the BoxND.
Definition AMReX_Box.H:385
__host__ __device__ IntVectND< dim > length() const noexcept
Return the length of the BoxND.
Definition AMReX_Box.H:167
Convolution-based solver for open boundary conditions using Green's functions.
Definition AMReX_FFT_OpenBCSolver.H:26
cMF const & greensFunctionFFT() const noexcept
Access the spectral Green's function held by this solver.
Definition AMReX_FFT_OpenBCSolver.H:91
Box const & Domain() const
Access the physical domain this solver was built for.
Definition AMReX_FFT_OpenBCSolver.H:61
typename R2C< T >::MF MF
Definition AMReX_FFT_OpenBCSolver.H:28
void solve(MF &phi, MF const &rho)
Solve for phi given right-hand side rho.
Definition AMReX_FFT_OpenBCSolver.H:265
void setGreensFunction(F const &greens_function)
Populate the spectral Green's function used by subsequent solves.
Definition AMReX_FFT_OpenBCSolver.H:179
IntVect const & PaddedLength() const
Access the one-sided padded length used to build the internal FFT domain.
Definition AMReX_FFT_OpenBCSolver.H:68
typename R2C< T >::cMF cMF
Definition AMReX_FFT_OpenBCSolver.H:29
OpenBCSolver(Box const &domain, Info const &info=Info{})
Build a solver over domain using the FFT Info settings in info.
Definition AMReX_FFT_OpenBCSolver.H:145
cMF & greensFunctionFFT() noexcept
Access the spectral Green's function held by this solver.
Definition AMReX_FFT_OpenBCSolver.H:84
Parallel Discrete Fourier Transform.
Definition AMReX_FFT_R2C.H:48
std::conditional_t< C, cMF, std::conditional_t< std::is_same_v< T, Real >, MultiFab, FabArray< BaseFab< T > > > > MF
Definition AMReX_FFT_R2C.H:53
Open Boundary Poisson Solver.
Definition AMReX_OpenBC.H:70
OpenBCSolver()=default
Construct an empty solver; call define() before solving.
amrex_long Long
Definition AMReX_INT.H:30
void ParallelForOMP(T n, L const &f) noexcept
Performance-portable kernel launch function with optional OpenMP threading.
Definition AMReX_GpuLaunch.H:328
Definition AMReX_FFT_Helper.H:53
int nextFastLen(int target, int nfactors=FastNumPrimeFactors())
Return the smallest fast FFT length greater than or equal to target.
Definition AMReX_FFT_Helper.H:286
DomainStrategy
Definition AMReX_FFT_Helper.H:57
void dtod_memcpy_async(void *p_d_dst, const void *p_d_src, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:449
int NProcsSub() noexcept
number of ranks in current frame
Definition AMReX_ParallelContext.H:74
@ make_alias
Definition AMReX_MakeType.H:7
__host__ __device__ void ignore_unused(const Ts &...)
No-op helper that marks variables as intentionally unused.
Definition AMReX.H:259
BoxND< 3 > Box
Box is an alias for amrex::BoxND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:35
IntVectND< 3 > IntVect
IntVect is an alias for amrex::IntVectND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:38
void Abort(const std::string &msg)
Print a fatal-error message to stderr and abort execution.
Definition AMReX.cpp:242
Definition AMReX_FFT_Helper.H:83
bool twod_mode
Definition AMReX_FFT_Helper.H:94
int batch_size
Batched FFT size. Only support in R2C, not R2X.
Definition AMReX_FFT_Helper.H:106
int nprocs
Max number of processes to use.
Definition AMReX_FFT_Helper.H:109
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:52