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
22template <typename T = Real>
24{
25public:
26 using MF = typename R2C<T>::MF;
27 using cMF = typename R2C<T>::cMF;
28
35 explicit OpenBCSolver (Box const& domain, Info const& info = Info{});
36
43 template <class F>
44 void setGreensFunction (F const& greens_function);
45
52 void solve (MF& phi, MF const& rho);
53
59 [[nodiscard]] Box const& Domain () const { return m_domain; }
60
61private:
62 static Box make_grown_domain (Box const& domain, Info const& info);
63
64 Box m_domain;
65 Info m_info;
66 R2C<T> m_r2c;
67 cMF m_G_fft;
68 std::unique_ptr<R2C<T>> m_r2c_green;
69};
70
71template <typename T>
72Box OpenBCSolver<T>::make_grown_domain (Box const& domain, Info const& info)
73{
74 IntVect len = domain.length();
75#if (AMREX_SPACEDIM == 3)
76 if (info.twod_mode) { len[2] = 0; }
77#else
79#endif
80 return Box(domain.smallEnd(), domain.bigEnd()+len, domain.ixType());
81}
82
83template <typename T>
84OpenBCSolver<T>::OpenBCSolver (Box const& domain, Info const& info)
85 : m_domain(domain),
86 m_info(info),
87 m_r2c(OpenBCSolver<T>::make_grown_domain(domain,info),
88 m_info.setDomainStrategy(FFT::DomainStrategy::slab))
89{
90#if (AMREX_SPACEDIM == 3)
91 if (m_info.twod_mode) {
92 auto gdom = make_grown_domain(domain,m_info);
93 gdom.enclosedCells(2);
94 gdom.setSmall(2, 0);
95 int nprocs = std::min({ParallelContext::NProcsSub(),
96 m_info.nprocs,
97 m_domain.length(2)});
98 gdom.setBig(2, nprocs-1);
99 m_r2c_green = std::make_unique<R2C<T>>(gdom,m_info);
100 auto [sd, ord] = m_r2c_green->getSpectralData();
101 m_G_fft = cMF(*sd, amrex::make_alias, 0, 1);
102 } else
103#endif
104 {
105 amrex::ignore_unused(m_r2c_green);
106 auto [sd, ord] = m_r2c.getSpectralData();
108 m_G_fft.define(sd->boxArray(), sd->DistributionMap(), 1, 0);
109 }
110}
111
112template <typename T>
113template <class F>
114void OpenBCSolver<T>::setGreensFunction (F const& greens_function)
115{
116 BL_PROFILE("OpenBCSolver::setGreensFunction");
117
118 auto* infab = m_info.twod_mode ? detail::get_fab(m_r2c_green->m_rx)
119 : detail::get_fab(m_r2c.m_rx);
120 auto const& lo = m_domain.smallEnd();
121 auto const& lo3 = lo.dim3();
122 auto const& len = m_domain.length3d();
123 if (infab) {
124 auto const& a = infab->array();
125 auto box = infab->box();
126 GpuArray<int,3> nimages{1,1,1};
127 int ndims = m_info.twod_mode ? AMREX_SPACEDIM-1 : AMREX_SPACEDIM;
128 for (int idim = 0; idim < ndims; ++idim) {
129 if (box.smallEnd(idim) == lo[idim] && box.length(idim) == 2*len[idim]) {
130 box.growHi(idim, -len[idim]+1); // +1 to include the middle plane
131 nimages[idim] = 2;
132 }
133 }
134 AMREX_ASSERT(nimages[0] == 2);
135 box.shift(-lo);
136 amrex::ParallelForOMP(box, [=] AMREX_GPU_DEVICE (int i, int j, int k)
137 {
138 T G;
139 if (i == len[0] || j == len[1] || k == len[2]) {
140 G = 0;
141 } else {
142 auto ii = i;
143 auto jj = (j > len[1]) ? 2*len[1]-j : j;
144 auto kk = (k > len[2]) ? 2*len[2]-k : k;
145 G = greens_function(ii+lo3.x,jj+lo3.y,kk+lo3.z);
146 }
147 for (int koff = 0; koff < nimages[2]; ++koff) {
148 int k2 = (koff == 0) ? k : 2*len[2]-k;
149 if ((k2 == 2*len[2]) || (koff == 1 && k == len[2])) {
150 continue;
151 }
152 for (int joff = 0; joff < nimages[1]; ++joff) {
153 int j2 = (joff == 0) ? j : 2*len[1]-j;
154 if ((j2 == 2*len[1]) || (joff == 1 && j == len[1])) {
155 continue;
156 }
157 for (int ioff = 0; ioff < nimages[0]; ++ioff) {
158 int i2 = (ioff == 0) ? i : 2*len[0]-i;
159 if ((i2 == 2*len[0]) || (ioff == 1 && i == len[0])) {
160 continue;
161 }
162 a(i2+lo3.x,j2+lo3.y,k2+lo3.z) = G;
163 }
164 }
165 }
166 });
167 }
168
169 if (m_info.twod_mode) {
170 m_r2c_green->forward(m_r2c_green->m_rx);
171 } else {
172 m_r2c.forward(m_r2c.m_rx);
173 }
174
175 if (!m_info.twod_mode) {
176 auto [sd, ord] = m_r2c.getSpectralData();
178 auto const* srcfab = detail::get_fab(*sd);
179 if (srcfab) {
180 auto* dstfab = detail::get_fab(m_G_fft);
181 if (dstfab) {
182 Gpu::dtod_memcpy_async(dstfab->dataPtr(), srcfab->dataPtr(), dstfab->nBytes());
183 } else {
184 amrex::Abort("FFT::OpenBCSolver: how did this happen");
185 }
186 }
187
188 m_r2c.prepare_openbc();
189 }
190}
191
192template <typename T>
193void OpenBCSolver<T>::solve (MF& phi, MF const& rho)
194{
195 BL_PROFILE("OpenBCSolver::solve");
196
197 auto& inmf = m_r2c.m_rx;
198 inmf.setVal(T(0));
199 inmf.ParallelCopy(rho, 0, 0, 1);
200
201 m_r2c.m_openbc_half = !m_info.twod_mode;
202 m_r2c.forward(inmf);
203 m_r2c.m_openbc_half = false;
204
205 auto scaling_factor = m_r2c.scalingFactor();
206
207 auto const* gfab = detail::get_fab(m_G_fft);
208 if (gfab) {
209 auto [sd, ord] = m_r2c.getSpectralData();
211 auto* rhofab = detail::get_fab(*sd);
212 if (rhofab) {
213 auto* pdst = rhofab->dataPtr();
214 auto const* psrc = gfab->dataPtr();
215 Box const& rhobox = rhofab->box();
216#if (AMREX_SPACEDIM == 3)
217 Long leng = gfab->box().numPts();
218 if (m_info.twod_mode) {
219 AMREX_ASSERT(gfab->box().length(2) == 1 &&
220 leng == (rhobox.length(0) * rhobox.length(1)));
221 } else {
222 AMREX_ASSERT(leng == rhobox.numPts());
223 }
224#endif
226 {
227#if (AMREX_SPACEDIM == 3)
228 Long isrc = i % leng;
229#else
230 Long isrc = i;
231#endif
232 pdst[i] *= psrc[isrc] * scaling_factor;
233 });
234 } else {
235 amrex::Abort("FFT::OpenBCSolver::solve: how did this happen?");
236 }
237 }
238
239 m_r2c.m_openbc_half = !m_info.twod_mode;
240 m_r2c.backward_doit(phi, phi.nGrowVect());
241 m_r2c.m_openbc_half = false;
242}
243
244}
245
246#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
Real * pdst
Definition AMReX_HypreMLABecLap.cpp:1090
__host__ __device__ const IntVectND< dim > & bigEnd() const &noexcept
Return the inclusive upper bound of the box.
Definition AMReX_Box.H:123
__host__ __device__ Long numPts() const noexcept
Return the number of points contained in the BoxND.
Definition AMReX_Box.H:356
__host__ __device__ IntVectND< dim > length() const noexcept
Return the length of the BoxND.
Definition AMReX_Box.H:154
__host__ __device__ IndexTypeND< dim > ixType() const noexcept
Return the indexing type.
Definition AMReX_Box.H:135
__host__ __device__ const IntVectND< dim > & smallEnd() const &noexcept
Return the inclusive lower bound of the box.
Definition AMReX_Box.H:111
Convolution-based solver for open boundary conditions using Green's functions.
Definition AMReX_FFT_OpenBCSolver.H:24
Box const & Domain() const
Access the physical domain this solver was built for.
Definition AMReX_FFT_OpenBCSolver.H:59
typename R2C< T >::MF MF
Definition AMReX_FFT_OpenBCSolver.H:26
void solve(MF &phi, MF const &rho)
Solve for phi given right-hand side rho.
Definition AMReX_FFT_OpenBCSolver.H:193
void setGreensFunction(F const &greens_function)
Populate the spectral Green's function used by subsequent solves.
Definition AMReX_FFT_OpenBCSolver.H:114
typename R2C< T >::cMF cMF
Definition AMReX_FFT_OpenBCSolver.H:27
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:84
Parallel Discrete Fourier Transform.
Definition AMReX_FFT_R2C.H:47
std::conditional_t< C, cMF, std::conditional_t< std::is_same_v< T, Real >, MultiFab, FabArray< BaseFab< T > > > > MF
Definition AMReX_FFT_R2C.H:52
Open Boundary Poisson Solver.
Definition AMReX_OpenBC.H:63
OpenBCSolver()=default
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:319
Definition AMReX_FFT_Helper.H:52
DomainStrategy
Definition AMReX_FFT_Helper.H:56
void dtod_memcpy_async(void *p_d_dst, const void *p_d_src, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:329
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 &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:139
BoxND< 3 > Box
Box is an alias for amrex::BoxND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:30
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition AMReX.cpp:230
Definition AMReX_FFT_Helper.H:64
bool twod_mode
Definition AMReX_FFT_Helper.H:75
int nprocs
Max number of processes to use.
Definition AMReX_FFT_Helper.H:84
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:43