Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
AMReX_FFT_R2X.H
Go to the documentation of this file.
1#ifndef AMREX_FFT_R2X_H_
2#define AMREX_FFT_R2X_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_MultiFab.H>
6#include <AMReX_FFT_Helper.H>
7#include <algorithm>
8#include <numeric>
9#include <tuple>
10
11namespace amrex::FFT
12{
13
14template <typename T> class Poisson;
15template <typename T> class PoissonHybrid;
16
24template <typename T = Real>
25class R2X
26{
27public:
28 using MF = std::conditional_t<std::is_same_v<T,Real>,
31
32 template <typename U> friend class Poisson;
33 template <typename U> friend class PoissonHybrid;
34
35 R2X (Box const& domain,
36 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> const& bc,
37 Info const& info = Info{});
38
39 ~R2X ();
40
41 R2X (R2X const&) = delete;
42 R2X (R2X &&) = delete;
43 R2X& operator= (R2X const&) = delete;
44 R2X& operator= (R2X &&) = delete;
45
46 [[nodiscard]] T scalingFactor () const;
47
48 template <typename F>
49 void forwardThenBackward (MF const& inmf, MF& outmf, F const& post_forward);
50
51 // public for cuda
52 template <int dim, typename FAB, typename F>
53 void post_forward_doit (FAB* fab, F const& f);
54
55 // private function made public for cuda
56 template <typename F>
57 void forwardThenBackward_doit_0 (MF const& inmf, MF& outmf, F const& post_forward,
58 IntVect const& ngout = IntVect(0),
59 Periodicity const& period = Periodicity::NonPeriodic());
60 template <typename F>
61 void forwardThenBackward_doit_1 (MF const& inmf, MF& outmf, F const& post_forward,
62 IntVect const& ngout = IntVect(0),
63 Periodicity const& period = Periodicity::NonPeriodic());
64
65private:
66
67 void forward (MF const& inmf, MF& outmf);
68 void forward (MF const& inmf, cMF& outmf);
69 void forward (MF const& inmf);
70 void backward (MF const& inmf, MF& outmf, IntVect const& ngout,
71 Periodicity const& period);
72 void backward (cMF const& inmf, MF& outmf, IntVect const& ngout,
73 Periodicity const& period);
74 void backward ();
75
76 Box m_dom_0;
77 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> m_bc;
78
79 Plan<T> m_fft_fwd_x{};
80 Plan<T> m_fft_bwd_x{};
81 Plan<T> m_fft_fwd_y{};
82 Plan<T> m_fft_bwd_y{};
83 Plan<T> m_fft_fwd_z{};
84 Plan<T> m_fft_bwd_z{};
85
86 std::unique_ptr<MultiBlockCommMetaData> m_cmd_cx2cy;
87 std::unique_ptr<MultiBlockCommMetaData> m_cmd_rx2ry;
88 std::unique_ptr<MultiBlockCommMetaData> m_cmd_cy2cz;
89 std::unique_ptr<MultiBlockCommMetaData> m_cmd_ry2rz;
90
91 std::unique_ptr<MultiBlockCommMetaData> m_cmd_cy2cx;
92 std::unique_ptr<MultiBlockCommMetaData> m_cmd_ry2rx;
93 std::unique_ptr<MultiBlockCommMetaData> m_cmd_cz2cy;
94 std::unique_ptr<MultiBlockCommMetaData> m_cmd_rz2ry;
95
96 Swap01 m_dtos_x2y{};
97 Swap01 m_dtos_y2x{};
98 Swap02 m_dtos_y2z{};
99 Swap02 m_dtos_z2y{};
100
101 MF m_rx;
102 MF m_ry;
103 MF m_rz;
104 cMF m_cx;
105 cMF m_cy;
106 cMF m_cz;
107
108 std::unique_ptr<char,DataDeleter> m_data_1;
109 std::unique_ptr<char,DataDeleter> m_data_2;
110
111 Box m_dom_rx;
112 Box m_dom_ry;
113 Box m_dom_rz;
114 Box m_dom_cx;
115 Box m_dom_cy;
116 Box m_dom_cz;
117
118 std::unique_ptr<R2X<T>> m_r2x_sub;
119 detail::SubHelper m_sub_helper;
120
121 Info m_info;
122};
123
124template <typename T>
125R2X<T>::R2X (Box const& domain,
126 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> const& bc,
127 Info const& info)
128 : m_dom_0(domain),
129 m_bc(bc),
130 m_sub_helper(domain),
131 m_info(info)
132{
133 BL_PROFILE("FFT::R2X");
134
135 static_assert(std::is_same_v<float,T> || std::is_same_v<double,T>);
136
137 AMREX_ALWAYS_ASSERT((m_dom_0.numPts() > 1) && (m_info.batch_size == 1));
138#if (AMREX_SPACEDIM == 2)
140#else
141 if (m_info.twod_mode) {
142 AMREX_ALWAYS_ASSERT((int(domain.length(0) > 1) +
143 int(domain.length(1) > 1) +
144 int(domain.length(2) > 1)) >= 2);
145 }
146#endif
147
148 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
149 if (bc[idim].first == Boundary::periodic ||
150 bc[idim].second == Boundary::periodic) {
151 AMREX_ALWAYS_ASSERT(bc[idim].first == bc[idim].second);
152 }
153 }
154
155 {
156 Box subbox = m_sub_helper.make_box(m_dom_0);
157 if (subbox.size() != m_dom_0.size()) {
158 m_r2x_sub = std::make_unique<R2X<T>>
159 (subbox, m_sub_helper.make_array(bc), info);
160 return;
161 }
162 }
163
164 int myproc = ParallelContext::MyProcSub();
165 int nprocs = std::min(ParallelContext::NProcsSub(), m_info.nprocs);
166
167 //
168 // make data containers
169 //
170
171 m_dom_rx = m_dom_0;
172 auto bax = amrex::decompose(m_dom_rx, nprocs, {AMREX_D_DECL(false,true,true)});
173 DistributionMapping dmx = detail::make_iota_distromap(bax.size());
174 m_rx.define(bax, dmx, 1, 0, MFInfo().SetAlloc(false));
175
176 // x-direction
177 if (bc[0].first == Boundary::periodic) {
178 // x-fft: r2c(m_rx->m_cx)
179 m_dom_cx = Box(IntVect(0), IntVect(AMREX_D_DECL(domain.length(0)/2,
180 domain.bigEnd(1),
181 domain.bigEnd(2))));
182 BoxList bl = bax.boxList();
183 for (auto & b : bl) {
184 b.setBig(0, m_dom_cx.bigEnd(0));
185 }
186 BoxArray cbax(std::move(bl));
187 m_cx.define(cbax, dmx, 1, 0, MFInfo().SetAlloc(false));
188 } // else: x-fft: r2r(m_rx)
189
190#if (AMREX_SPACEDIM >= 2)
191 if (domain.length(1) > 1 && !m_info.oned_mode) {
192 if (! m_cx.empty()) {
193 // copy(m_cx->m_cy)
194 m_dom_cy = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_cx.bigEnd(1),
195 m_dom_cx.bigEnd(0),
196 m_dom_cx.bigEnd(2))));
197 auto ba = amrex::decompose(m_dom_cy, nprocs, {AMREX_D_DECL(false,true,true)});
199 if (ba.size() == m_cx.size()) {
200 dm = m_cx.DistributionMap();
201 } else {
202 dm = detail::make_iota_distromap(ba.size());
203 }
204 m_cy.define(ba, dm, 1, 0, MFInfo().SetAlloc(false));
205 // if bc[1] is periodic:
206 // c2c(m_cy->m_cy)
207 // else:
208 // r2r(m_cy.re) & r2r(m_cy.im)
209 } else {
210 // copy(m_rx->m_ry)
211 m_dom_ry = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_rx.bigEnd(1),
212 m_dom_rx.bigEnd(0),
213 m_dom_rx.bigEnd(2))));
214 auto ba = amrex::decompose(m_dom_ry, nprocs, {AMREX_D_DECL(false,true,true)});
216 if (ba.size() == m_rx.size()) {
217 dm = m_rx.DistributionMap();
218 } else {
219 dm = detail::make_iota_distromap(ba.size());
220 }
221 m_ry.define(ba, dm, 1, 0, MFInfo().SetAlloc(false));
222 // if bc[1] is periodic:
223 // r2c(m_ry->m_cy)
224 // else:
225 // r2r(m_ry)
226 if (bc[1].first == Boundary::periodic) {
227 m_dom_cy = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_ry.length(0)/2,
228 m_dom_ry.bigEnd(1),
229 m_dom_ry.bigEnd(2))));
230 BoxList bl = ba.boxList();
231 for (auto & b : bl) {
232 b.setBig(0, m_dom_cy.bigEnd(0));
233 }
234 BoxArray cba(std::move(bl));
235 m_cy.define(cba, dm, 1, 0, MFInfo().SetAlloc(false));
236 }
237 }
238 }
239#endif
240
241#if (AMREX_SPACEDIM == 3)
242 if (domain.length(2) > 1 && !m_info.twod_mode) {
243 if (! m_cy.empty()) {
244 // copy(m_cy, m_cz)
245 m_dom_cz = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_cy.bigEnd(2),
246 m_dom_cy.bigEnd(1),
247 m_dom_cy.bigEnd(0))));
248 auto ba = amrex::decompose(m_dom_cz, nprocs, {AMREX_D_DECL(false,true,true)});
250 if (ba.size() == m_cy.size()) {
251 dm = m_cy.DistributionMap();
252 } else {
253 dm = detail::make_iota_distromap(ba.size());
254 }
255 m_cz.define(ba, dm, 1, 0, MFInfo().SetAlloc(false));
256 // if bc[2] is periodic:
257 // c2c(m_cz->m_cz)
258 // else:
259 // r2r(m_cz.re) & r2r(m_cz.im)
260 } else {
261 // copy(m_ry, m_rz)
262 m_dom_rz = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_ry.bigEnd(2),
263 m_dom_ry.bigEnd(1),
264 m_dom_ry.bigEnd(0))));
265 auto ba = amrex::decompose(m_dom_rz, nprocs, {AMREX_D_DECL(false,true,true)});
267 if (ba.size() == m_ry.size()) {
268 dm = m_ry.DistributionMap();
269 } else {
270 dm = detail::make_iota_distromap(ba.size());
271 }
272 m_rz.define(ba, dm, 1, 0, MFInfo().SetAlloc(false));
273 // if bc[2] is periodic:
274 // r2c(m_rz->m_cz)
275 // else:
276 // r2r(m_rz)
277 if (bc[2].first == Boundary::periodic) {
278 m_dom_cz = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_rz.length(0)/2,
279 m_dom_rz.bigEnd(1),
280 m_dom_rz.bigEnd(2))));
281 BoxList bl = ba.boxList();
282 for (auto & b : bl) {
283 b.setBig(0, m_dom_cz.bigEnd(0));
284 }
285 BoxArray cba(std::move(bl));
286 m_cz.define(cba, dm, 1, 0, MFInfo().SetAlloc(false));
287 }
288 }
289 }
290#endif
291
292 // There are several different execution paths.
293 //
294 // (1) x-r2c(m_rx->m_cx), copy(m_cx->m_cy), y-fft(m_cy),
295 // copy(m_cy->m_cz), z-fft(m_cz)
296 // In this case, we have m_rx, m_cx, m_cy, & m_cz.
297 // we can alias(m_rx,m_cy) and alias(m_cx,m_cz).
298 //
299 // (2) x_r2r(m_rx), copy(m_rx->m_ry), y-r2c(m_ry->m_cy),
300 // copy(m_cy->m_cz), z-fft(m_cz)
301 // In this case, we have m_rx, m_ry, m_cy, & m_cz.
302 // We can alias(m_rx,m_cy) and alias(m_ry,m_cz).
303 //
304 // (3) x_r2r(m_rx), copy(m_rx->m_ry), y-r2r(m_ry),
305 // copy(m_ry->m_rz), z-r2c(m_rz->m_rz)
306 // In this case, we have m_rx, m_ry, m_rz, & m_cz
307 // We can alias(m_rx,m_rz) and alias(m_ry,m_cz)
308 //
309 // (4) x_r2r(m_rx), copy(m_rx->m_ry), y-r2r(m_ry),
310 // copy(m_ry->m_rz), z-r2r(m_rz)
311 // In this case, we have m_rx, m_ry, & m_rz.
312 // We can alias(m_rx,m_rz).
313
314 if (! m_cx.empty()) {
315 m_data_1 = detail::make_mfs_share(m_rx, m_cy);
316 m_data_2 = detail::make_mfs_share(m_cx, m_cz);
317 } else if (! m_cy.empty()) {
318 m_data_1 = detail::make_mfs_share(m_rx, m_cy);
319 m_data_2 = detail::make_mfs_share(m_ry, m_cz);
320 } else if (! m_cz.empty()) {
321 m_data_1 = detail::make_mfs_share(m_rx, m_rz);
322 m_data_2 = detail::make_mfs_share(m_ry, m_cz);
323 } else {
324 m_data_1 = detail::make_mfs_share(m_rx, m_rz);
325 m_data_2 = detail::make_mfs_share(m_ry, m_cz); // It's okay m_cz is empty.
326 }
327
328 //
329 // make copiers
330 //
331
332#if (AMREX_SPACEDIM >= 2)
333 if (!m_cy.empty() || !m_ry.empty()) {
334 if (! m_cx.empty()) {
335 // copy(m_cx->m_cy)
336 m_cmd_cx2cy = std::make_unique<MultiBlockCommMetaData>
337 (m_cy, m_dom_cy, m_cx, IntVect(0), m_dtos_x2y);
338 m_cmd_cy2cx = std::make_unique<MultiBlockCommMetaData>
339 (m_cx, m_dom_cx, m_cy, IntVect(0), m_dtos_y2x);
340 } else {
341 // copy(m_rx->m_ry)
342 m_cmd_rx2ry = std::make_unique<MultiBlockCommMetaData>
343 (m_ry, m_dom_ry, m_rx, IntVect(0), m_dtos_x2y);
344 m_cmd_ry2rx = std::make_unique<MultiBlockCommMetaData>
345 (m_rx, m_dom_rx, m_ry, IntVect(0), m_dtos_y2x);
346 }
347 }
348#endif
349
350#if (AMREX_SPACEDIM == 3)
351 if (!m_cz.empty() || !m_rz.empty()) {
352 if (! m_cy.empty()) {
353 // copy(m_cy, m_cz)
354 m_cmd_cy2cz = std::make_unique<MultiBlockCommMetaData>
355 (m_cz, m_dom_cz, m_cy, IntVect(0), m_dtos_y2z);
356 m_cmd_cz2cy = std::make_unique<MultiBlockCommMetaData>
357 (m_cy, m_dom_cy, m_cz, IntVect(0), m_dtos_z2y);
358 } else {
359 // copy(m_ry, m_rz)
360 m_cmd_ry2rz = std::make_unique<MultiBlockCommMetaData>
361 (m_rz, m_dom_rz, m_ry, IntVect(0), m_dtos_y2z);
362 m_cmd_rz2ry = std::make_unique<MultiBlockCommMetaData>
363 (m_ry, m_dom_ry, m_rz, IntVect(0), m_dtos_z2y);
364 }
365 }
366#endif
367
368 //
369 // make plans
370 //
371
372 using VendorComplex = typename Plan<T>::VendorComplex;
373
374 if (myproc < m_rx.size())
375 {
376 Box const& box = m_rx.box(myproc);
377 auto* pf = m_rx[myproc].dataPtr();
378 if (bc[0].first == Boundary::periodic) {
379 auto* pb = (VendorComplex*) m_cx[myproc].dataPtr();
380 m_fft_fwd_x.template init_r2c<Direction::forward>(box, pf, pb);
381#if defined(AMREX_USE_SYCL)
382 m_fft_bwd_x = m_fft_fwd_x;
383#else
384 m_fft_bwd_x.template init_r2c<Direction::backward>(box, pf, pb);
385#endif
386 } else {
387 m_fft_fwd_x.template init_r2r<Direction::forward>(box, pf, bc[0]);
388#if defined(AMREX_USE_GPU)
389 if ((bc[0].first == Boundary::even && bc[0].second == Boundary::odd) ||
390 (bc[0].first == Boundary::odd && bc[0].second == Boundary::even)) {
391 m_fft_bwd_x = m_fft_fwd_x;
392 } else
393#endif
394 {
395 m_fft_bwd_x.template init_r2r<Direction::backward>(box, pf, bc[0]);
396 }
397 }
398 }
399
400#if (AMREX_SPACEDIM >= 2)
401 if (m_ry.empty() && m_bc[1].first == Boundary::periodic) {
402 if (myproc < m_cy.size()) {
403 Box const& box = m_cy.box(myproc);
404 auto* p = (VendorComplex *)m_cy[myproc].dataPtr();
405 m_fft_fwd_y.template init_c2c<Direction::forward>(box, p);
406#if defined(AMREX_USE_SYCL)
407 m_fft_bwd_y = m_fft_fwd_y;
408#else
409 m_fft_bwd_y.template init_c2c<Direction::backward>(box, p);
410#endif
411 }
412 } else if (!m_ry.empty() && m_bc[1].first == Boundary::periodic) {
413 if (myproc < m_ry.size()) {
414 Box const& box = m_ry.box(myproc);
415 auto* pr = m_ry[myproc].dataPtr();
416 auto* pc = (VendorComplex*)m_cy[myproc].dataPtr();
417 m_fft_fwd_y.template init_r2c<Direction::forward>(box, pr, pc);
418#if defined(AMREX_USE_SYCL)
419 m_fft_bwd_y = m_fft_fwd_y;
420#else
421 m_fft_bwd_y.template init_r2c<Direction::backward>(box, pr, pc);
422#endif
423 }
424 } else if (!m_cy.empty()) {
425 if (myproc < m_cy.size()) {
426 Box const& box = m_cy.box(myproc);
427 auto* p = (VendorComplex*) m_cy[myproc].dataPtr();
428 m_fft_fwd_y.template init_r2r<Direction::forward>(box, p, bc[1]);
429#if defined(AMREX_USE_GPU)
430 if ((bc[1].first == Boundary::even && bc[1].second == Boundary::odd) ||
431 (bc[1].first == Boundary::odd && bc[1].second == Boundary::even)) {
432 m_fft_bwd_y = m_fft_fwd_y;
433 } else
434#endif
435 {
436 m_fft_bwd_y.template init_r2r<Direction::backward>(box, p, bc[1]);
437 }
438 }
439 } else {
440 if (myproc < m_ry.size()) {
441 Box const& box = m_ry.box(myproc);
442 auto* p = m_ry[myproc].dataPtr();
443 m_fft_fwd_y.template init_r2r<Direction::forward>(box, p, bc[1]);
444#if defined(AMREX_USE_GPU)
445 if ((bc[1].first == Boundary::even && bc[1].second == Boundary::odd) ||
446 (bc[1].first == Boundary::odd && bc[1].second == Boundary::even)) {
447 m_fft_bwd_y = m_fft_fwd_y;
448 } else
449#endif
450 {
451 m_fft_bwd_y.template init_r2r<Direction::backward>(box, p, bc[1]);
452 }
453 }
454 }
455#endif
456
457#if (AMREX_SPACEDIM == 3)
458 if (m_rz.empty() && m_bc[2].first == Boundary::periodic) {
459 if (myproc < m_cz.size()) {
460 Box const& box = m_cz.box(myproc);
461 auto* p = (VendorComplex*)m_cz[myproc].dataPtr();
462 m_fft_fwd_z.template init_c2c<Direction::forward>(box, p);
463#if defined(AMREX_USE_SYCL)
464 m_fft_bwd_z = m_fft_fwd_z;
465#else
466 m_fft_bwd_z.template init_c2c<Direction::backward>(box, p);
467#endif
468 }
469 } else if (!m_rz.empty() && m_bc[2].first == Boundary::periodic) {
470 if (myproc < m_rz.size()) {
471 Box const& box = m_rz.box(myproc);
472 auto* pr = m_rz[myproc].dataPtr();
473 auto* pc = (VendorComplex*)m_cz[myproc].dataPtr();
474 m_fft_fwd_z.template init_r2c<Direction::forward>(box, pr, pc);
475#if defined(AMREX_USE_SYCL)
476 m_fft_bwd_z = m_fft_fwd_z;
477#else
478 m_fft_bwd_z.template init_r2c<Direction::backward>(box, pr, pc);
479#endif
480 }
481 } else if (!m_cz.empty()) {
482 if (myproc < m_cz.size()) {
483 Box const& box = m_cz.box(myproc);
484 auto* p = (VendorComplex*) m_cz[myproc].dataPtr();
485 m_fft_fwd_z.template init_r2r<Direction::forward>(box, p, bc[2]);
486#if defined(AMREX_USE_GPU)
487 if ((bc[2].first == Boundary::even && bc[2].second == Boundary::odd) ||
488 (bc[2].first == Boundary::odd && bc[2].second == Boundary::even)) {
489 m_fft_bwd_z = m_fft_fwd_z;
490 } else
491#endif
492 {
493 m_fft_bwd_z.template init_r2r<Direction::backward>(box, p, bc[2]);
494 }
495 }
496 } else {
497 if (myproc < m_rz.size()) {
498 Box const& box = m_rz.box(myproc);
499 auto* p = m_rz[myproc].dataPtr();
500 m_fft_fwd_z.template init_r2r<Direction::forward>(box, p, bc[2]);
501#if defined(AMREX_USE_GPU)
502 if ((bc[2].first == Boundary::even && bc[2].second == Boundary::odd) ||
503 (bc[2].first == Boundary::odd && bc[2].second == Boundary::even)) {
504 m_fft_bwd_z = m_fft_fwd_z;
505 } else
506#endif
507 {
508 m_fft_bwd_z.template init_r2r<Direction::backward>(box, p, bc[2]);
509 }
510 }
511 }
512#endif
513}
514
515template <typename T>
517{
518 if (m_fft_bwd_x.plan != m_fft_fwd_x.plan) {
519 m_fft_bwd_x.destroy();
520 }
521 if (m_fft_bwd_y.plan != m_fft_fwd_y.plan) {
522 m_fft_bwd_y.destroy();
523 }
524 if (m_fft_bwd_z.plan != m_fft_fwd_z.plan) {
525 m_fft_bwd_z.destroy();
526 }
527 m_fft_fwd_x.destroy();
528 m_fft_fwd_y.destroy();
529 m_fft_fwd_z.destroy();
530}
531
532template <typename T>
534{
535 Long r = 1;
536 int ndims = m_info.twod_mode ? AMREX_SPACEDIM-1 : AMREX_SPACEDIM;
537#if (AMREX_SPACEDIM == 3)
538 if (m_info.twod_mode && m_dom_0.length(2) == 1) { ndims = 1; };
539#endif
540 for (int idim = 0; idim < ndims; ++idim) {
541 r *= m_dom_0.length(idim);
542 if (m_bc[idim].first != Boundary::periodic && (m_dom_0.length(idim) > 1)) {
543 r *= 2;
544 }
545 }
546 return T(1)/T(r);
547}
548
549template <typename T>
550template <typename F>
551void R2X<T>::forwardThenBackward (MF const& inmf, MF& outmf, F const& post_forward)
552{
553 forwardThenBackward_doit_0(inmf, outmf, post_forward);
554}
555
556template <typename T>
557template <typename F>
558void R2X<T>::forwardThenBackward_doit_0 (MF const& inmf, MF& outmf,
559 F const& post_forward,
560 IntVect const& ngout,
561 Periodicity const& period)
562{
563 BL_PROFILE("FFT::R2X::forwardbackward_0");
564
565 if (m_r2x_sub) {
566 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
567 MF inmf_sub, inmf_tmp;
568 if (inmf_safe) {
569 inmf_sub = m_sub_helper.make_alias_mf(inmf);
570 } else {
571 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
572 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
573 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
574 }
575
576 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
577 MF outmf_sub, outmf_tmp;
578 if (outmf_safe) {
579 outmf_sub = m_sub_helper.make_alias_mf(outmf);
580 } else {
581 IntVect const& ngtmp = m_sub_helper.make_safe_ghost(outmf.nGrowVect());
582 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, ngtmp);
583 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
584 }
585
586 IntVect const& subngout = m_sub_helper.make_iv(ngout);
587 Periodicity const& subperiod = m_sub_helper.make_periodicity(period);
588 GpuArray<int,3> const& order = m_sub_helper.xyz_order();
589 m_r2x_sub->forwardThenBackward_doit_1
590 (inmf_sub, outmf_sub,
591 [=] AMREX_GPU_DEVICE (int i, int j, int k, auto& sp)
592 {
593 GpuArray<int,3> idx{i,j,k};
594 post_forward(idx[order[0]], idx[order[1]], idx[order[2]], sp);
595 },
596 subngout, subperiod);
597
598 if (!outmf_safe) {
599 outmf.LocalCopy(outmf_tmp, 0, 0, 1, outmf_tmp.nGrowVect());
600 }
601 }
602 else
603 {
604 this->forwardThenBackward_doit_1(inmf, outmf, post_forward, ngout, period);
605 }
606}
607
608template <typename T>
609template <typename F>
610void R2X<T>::forwardThenBackward_doit_1 (MF const& inmf, MF& outmf,
611 F const& post_forward,
612 IntVect const& ngout,
613 Periodicity const& period)
614{
615 BL_PROFILE("FFT::R2X::forwardbackward_1");
616
617 if (m_r2x_sub) {
618 amrex::Abort("R2X::forwardThenBackward_doit_1: How did this happen?");
619 }
620 else
621 {
622 this->forward(inmf);
623
624 // post-forward
625
626 int actual_dim = AMREX_SPACEDIM;
627#if (AMREX_SPACEDIM >= 2)
628 if (m_dom_0.length(1) == 1) { actual_dim = 1; }
629#endif
630#if (AMREX_SPACEDIM == 3)
631 if ((m_dom_0.length(2) == 1) && (m_dom_0.length(1) > 1)) { actual_dim = 2; }
632#endif
633
634 if (actual_dim == 1) {
635 if (m_cx.empty()) {
636 post_forward_doit<0>(detail::get_fab(m_rx), post_forward);
637 } else {
638 post_forward_doit<0>(detail::get_fab(m_cx), post_forward);
639 }
640 }
641#if (AMREX_SPACEDIM >= 2)
642 else if (actual_dim == 2) {
643 if (m_cy.empty()) {
644 post_forward_doit<1>(detail::get_fab(m_ry), post_forward);
645 } else {
646 post_forward_doit<1>(detail::get_fab(m_cy), post_forward);
647 }
648 }
649#endif
650#if (AMREX_SPACEDIM == 3)
651 else if (actual_dim == 3) {
652 if (m_cz.empty()) {
653 post_forward_doit<2>(detail::get_fab(m_rz), post_forward);
654 } else {
655 post_forward_doit<2>(detail::get_fab(m_cz), post_forward);
656 }
657 }
658#endif
659
660 this->backward();
661
662 outmf.ParallelCopy(m_rx, 0, 0, 1, IntVect(0),
663 amrex::elemwiseMin(ngout,outmf.nGrowVect()), period);
664 }
665}
666
667template <typename T>
668void R2X<T>::forward (MF const& inmf)
669{
670 BL_PROFILE("FFT::R2X::forward");
671
672 if (m_r2x_sub) {
673 if (m_sub_helper.ghost_safe(inmf.nGrowVect())) {
674 m_r2x_sub->forward(m_sub_helper.make_alias_mf(inmf));
675 } else {
676 MF tmp(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
677 tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
678 m_r2x_sub->forward(m_sub_helper.make_alias_mf(tmp));
679 }
680 return;
681 }
682
683 m_rx.ParallelCopy(inmf, 0, 0, 1);
684 if (m_bc[0].first == Boundary::periodic) {
685 m_fft_fwd_x.template compute_r2c<Direction::forward>();
686 } else {
687 m_fft_fwd_x.template compute_r2r<Direction::forward>();
688 }
689
690#if (AMREX_SPACEDIM >= 2)
691 if ( m_cmd_cx2cy) {
692 ParallelCopy(m_cy, m_cx, *m_cmd_cx2cy, 0, 0, 1, m_dtos_x2y);
693 } else if ( m_cmd_rx2ry) {
694 ParallelCopy(m_ry, m_rx, *m_cmd_rx2ry, 0, 0, 1, m_dtos_x2y);
695 }
696 if (m_bc[1].first != Boundary::periodic)
697 {
698 m_fft_fwd_y.template compute_r2r<Direction::forward>();
699 }
700 else if (m_bc[0].first == Boundary::periodic)
701 {
702 m_fft_fwd_y.template compute_c2c<Direction::forward>();
703 }
704 else
705 {
706 m_fft_fwd_y.template compute_r2c<Direction::forward>();
707 }
708#endif
709
710#if (AMREX_SPACEDIM == 3)
711 if ( m_cmd_cy2cz) {
712 ParallelCopy(m_cz, m_cy, *m_cmd_cy2cz, 0, 0, 1, m_dtos_y2z);
713 } else if ( m_cmd_ry2rz) {
714 ParallelCopy(m_rz, m_ry, *m_cmd_ry2rz, 0, 0, 1, m_dtos_y2z);
715 }
716 if (m_bc[2].first != Boundary::periodic)
717 {
718 m_fft_fwd_z.template compute_r2r<Direction::forward>();
719 }
720 else if (m_bc[0].first == Boundary::periodic ||
721 m_bc[1].first == Boundary::periodic)
722 {
723 m_fft_fwd_z.template compute_c2c<Direction::forward>();
724 }
725 else
726 {
727 m_fft_fwd_z.template compute_r2c<Direction::forward>();
728 }
729#endif
730}
731
732template <typename T>
733void R2X<T>::forward (MF const& inmf, MF& outmf)
734{
735 if (m_r2x_sub)
736 {
737 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
738 MF inmf_sub, inmf_tmp;
739 if (inmf_safe) {
740 inmf_sub = m_sub_helper.make_alias_mf(inmf);
741 } else {
742 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
743 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
744 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
745 }
746
747 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
748 MF outmf_sub, outmf_tmp;
749 if (outmf_safe) {
750 outmf_sub = m_sub_helper.make_alias_mf(outmf);
751 } else {
752 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, 0);
753 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
754 }
755
756 m_r2x_sub->forward(inmf_sub, outmf_sub);
757
758 if (!outmf_safe) {
759 outmf.LocalCopy(outmf_tmp, 0, 0, 1, IntVect(0));
760 }
761 }
762 else
763 {
764 this->forward(inmf);
765
766#if (AMREX_SPACEDIM == 3)
767 if (m_info.twod_mode) {
768 if (m_cy.empty() && !m_ry.empty()) {
769 ParallelCopy(outmf, m_dom_rx, m_ry, 0, 0, 1, IntVect(0), Swap01{});
770 } else if (m_ry.empty() && m_cy.empty() && m_cx.empty()) {
771 outmf.ParallelCopy(m_rx, 0, 0, 1);
772 } else {
773 amrex::Abort("R2X::forward(MF,MF): How did this happen?");
774 }
775 } else
776#endif
777 {
779 amrex::Abort("R2X::forward(MF,MF): TODO");
780 }
781 }
782}
783
784template <typename T>
785void R2X<T>::forward (MF const& inmf, cMF& outmf)
786{
787 if (m_r2x_sub)
788 {
789 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
790 MF inmf_sub, inmf_tmp;
791 if (inmf_safe) {
792 inmf_sub = m_sub_helper.make_alias_mf(inmf);
793 } else {
794 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
795 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
796 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
797 }
798
799 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
800 cMF outmf_sub, outmf_tmp;
801 if (outmf_safe) {
802 outmf_sub = m_sub_helper.make_alias_mf(outmf);
803 } else {
804 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, 0);
805 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
806 }
807
808 m_r2x_sub->forward(inmf_sub, outmf_sub);
809
810 if (!outmf_safe) {
811 outmf.LocalCopy(outmf_tmp, 0, 0, 1, IntVect(0));
812 }
813 }
814 else
815 {
816 this->forward(inmf);
817
818#if (AMREX_SPACEDIM == 3)
819 if (m_info.twod_mode) {
820 if (!m_cy.empty()) {
821 auto lo = m_dom_cy.smallEnd();
822 auto hi = m_dom_cy.bigEnd();
823 std::swap(lo[0],lo[1]);
824 std::swap(hi[0],hi[1]);
825 Box dom(lo,hi);
826 ParallelCopy(outmf, dom, m_cy, 0, 0, 1, IntVect(0), Swap01{});
827 } else if (m_ry.empty() && m_cy.empty() && !m_cx.empty()) {
828 outmf.ParallelCopy(m_cx, 0, 0, 1);
829 } else {
830 amrex::Abort("R2X::forward(MF,cMF): How did this happen?");
831 }
832 } else
833#endif
834 {
836 amrex::Abort("R2X::forward(MF,cMF): TODO");
837 }
838 }
839}
840
841template <typename T>
842void R2X<T>::backward ()
843{
844 BL_PROFILE("FFT::R2X::backward");
845
846 AMREX_ALWAYS_ASSERT(m_r2x_sub == nullptr);
847
848#if (AMREX_SPACEDIM == 3)
849 if (m_bc[2].first != Boundary::periodic)
850 {
851 m_fft_bwd_z.template compute_r2r<Direction::backward>();
852 }
853 else if (m_bc[0].first == Boundary::periodic ||
854 m_bc[1].first == Boundary::periodic)
855 {
856 m_fft_bwd_z.template compute_c2c<Direction::backward>();
857 }
858 else
859 {
860 m_fft_bwd_z.template compute_r2c<Direction::backward>();
861 }
862 if ( m_cmd_cz2cy) {
863 ParallelCopy(m_cy, m_cz, *m_cmd_cz2cy, 0, 0, 1, m_dtos_z2y);
864 } else if ( m_cmd_rz2ry) {
865 ParallelCopy(m_ry, m_rz, *m_cmd_rz2ry, 0, 0, 1, m_dtos_z2y);
866 }
867#endif
868
869#if (AMREX_SPACEDIM >= 2)
870 if (m_bc[1].first != Boundary::periodic)
871 {
872 m_fft_bwd_y.template compute_r2r<Direction::backward>();
873 }
874 else if (m_bc[0].first == Boundary::periodic)
875 {
876 m_fft_bwd_y.template compute_c2c<Direction::backward>();
877 }
878 else
879 {
880 m_fft_bwd_y.template compute_r2c<Direction::backward>();
881 }
882 if ( m_cmd_cy2cx) {
883 ParallelCopy(m_cx, m_cy, *m_cmd_cy2cx, 0, 0, 1, m_dtos_y2x);
884 } else if ( m_cmd_ry2rx) {
885 ParallelCopy(m_rx, m_ry, *m_cmd_ry2rx, 0, 0, 1, m_dtos_y2x);
886 }
887#endif
888
889 if (m_bc[0].first == Boundary::periodic) {
890 m_fft_bwd_x.template compute_r2c<Direction::backward>();
891 } else {
892 m_fft_bwd_x.template compute_r2r<Direction::backward>();
893 }
894}
895
896template <typename T>
897void R2X<T>::backward (MF const& inmf, MF& outmf, IntVect const& ngout,
898 Periodicity const& period)
899{
900 if (m_r2x_sub)
901 {
902 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
903 MF inmf_sub, inmf_tmp;
904 if (inmf_safe) {
905 inmf_sub = m_sub_helper.make_alias_mf(inmf);
906 } else {
907 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
908 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
909 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
910 }
911
912 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
913 MF outmf_sub, outmf_tmp;
914 if (outmf_safe) {
915 outmf_sub = m_sub_helper.make_alias_mf(outmf);
916 } else {
917 IntVect const& ngtmp = m_sub_helper.make_safe_ghost(outmf.nGrowVect());
918 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, ngtmp);
919 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
920 }
921
922 IntVect const& subngout = m_sub_helper.make_iv(ngout);
923 Periodicity const& subperiod = m_sub_helper.make_periodicity(period);
924 m_r2x_sub->backward(inmf_sub, outmf_sub, subngout, subperiod);
925
926 if (!outmf_safe) {
927 outmf.LocalCopy(outmf_tmp, 0, 0, 1, outmf_tmp.nGrowVect());
928 }
929 }
930 else
931 {
932#if (AMREX_SPACEDIM == 3)
933 if (m_info.twod_mode) {
934 if (m_cy.empty() && !m_ry.empty()) {
935 ParallelCopy(m_ry, m_dom_ry, inmf, 0, 0, 1, IntVect(0), Swap01{});
936 } else if (m_ry.empty() && m_cy.empty() && m_cx.empty()) {
937 m_rx.ParallelCopy(inmf, 0, 0, 1);
938 } else {
939 amrex::Abort("R2X::backward(MF,MF): How did this happen?");
940 }
941 } else
942#endif
943 {
944 amrex::ignore_unused(inmf,outmf,ngout,period);
945 amrex::Abort("R2X::backward(MF,MF): TODO");
946 }
947
948 this->backward();
949
950 outmf.ParallelCopy(m_rx, 0, 0, 1, IntVect(0),
951 amrex::elemwiseMin(ngout,outmf.nGrowVect()), period);
952 }
953}
954
955template <typename T>
956void R2X<T>::backward (cMF const& inmf, MF& outmf, IntVect const& ngout,
957 Periodicity const& period)
958{
959 if (m_r2x_sub)
960 {
961 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
962 cMF inmf_sub, inmf_tmp;
963 if (inmf_safe) {
964 inmf_sub = m_sub_helper.make_alias_mf(inmf);
965 } else {
966 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
967 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
968 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
969 }
970
971 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
972 MF outmf_sub, outmf_tmp;
973 if (outmf_safe) {
974 outmf_sub = m_sub_helper.make_alias_mf(outmf);
975 } else {
976 IntVect const& ngtmp = m_sub_helper.make_safe_ghost(outmf.nGrowVect());
977 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, ngtmp);
978 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
979 }
980
981 IntVect const& subngout = m_sub_helper.make_iv(ngout);
982 Periodicity const& subperiod = m_sub_helper.make_periodicity(period);
983 m_r2x_sub->backward(inmf_sub, outmf_sub, subngout, subperiod);
984
985 if (!outmf_safe) {
986 outmf.LocalCopy(outmf_tmp, 0, 0, 1, outmf_tmp.nGrowVect());
987 }
988 }
989 else
990 {
991#if (AMREX_SPACEDIM == 3)
992 if (m_info.twod_mode) {
993 if (!m_cy.empty()) {
994 ParallelCopy(m_cy, m_dom_cy, inmf, 0, 0, 1, IntVect(0), Swap01{});
995 } else if (m_ry.empty() && m_cy.empty() && !m_cx.empty()) {
996 m_cx.ParallelCopy(inmf, 0, 0, 1);
997 } else {
998 amrex::Abort("R2X::backward(cMF,MF): How did this happen?");
999 }
1000 } else
1001#endif
1002 {
1003 amrex::ignore_unused(inmf,outmf,ngout,period);
1004 amrex::Abort("R2X::backward(cMF,MF): TODO");
1005 }
1006
1007 this->backward();
1008
1009 outmf.ParallelCopy(m_rx, 0, 0, 1, IntVect(0),
1010 amrex::elemwiseMin(ngout,outmf.nGrowVect()), period);
1011 }
1012}
1013
1014template <typename T>
1015template <int dim, typename FAB, typename F>
1016void R2X<T>::post_forward_doit (FAB* fab, F const& f)
1017{
1018 if (m_info.twod_mode) {
1019 amrex::Abort("xxxxx post_forward_doit: todo");
1020 }
1021 if (fab) {
1022 auto const& a = fab->array();
1023 ParallelForOMP(fab->box(),
1024 [f=f,a=a] AMREX_GPU_DEVICE (int i, int j, int k)
1025 {
1026 if constexpr (dim == 0) {
1027 f(i,j,k,a(i,j,k));
1028 } else if constexpr (dim == 1) {
1029 f(j,i,k,a(i,j,k));
1030 } else {
1031 f(j,k,i,a(i,j,k));
1032 }
1033 });
1034 }
1035}
1036
1037}
1038
1039#endif
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:551
#define AMREX_ALWAYS_ASSERT(EX)
Definition AMReX_BLassert.H:50
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
A collection of Boxes stored in an Array.
Definition AMReX_BoxArray.H:567
A class for managing a List of Boxes that share a common IndexType. This class implements operations ...
Definition AMReX_BoxList.H:52
__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__ IntVectND< dim > size() const noexcept
Return the length of the BoxND.
Definition AMReX_Box.H:147
Calculates the distribution of FABs to MPI processes.
Definition AMReX_DistributionMapping.H:43
3D Poisson solver for periodic, Dirichlet & Neumann boundaries in the first two dimensions,...
Definition AMReX_FFT_Poisson.H:151
Poisson solver for periodic, Dirichlet & Neumann boundaries using FFT.
Definition AMReX_FFT_Poisson.H:61
Discrete Fourier Transform.
Definition AMReX_FFT_R2X.H:26
void post_forward_doit(FAB *fab, F const &f)
Definition AMReX_FFT_R2X.H:1016
std::conditional_t< std::is_same_v< T, Real >, MultiFab, FabArray< BaseFab< T > > > MF
Definition AMReX_FFT_R2X.H:29
~R2X()
Definition AMReX_FFT_R2X.H:516
void forwardThenBackward(MF const &inmf, MF &outmf, F const &post_forward)
Definition AMReX_FFT_R2X.H:551
R2X(Box const &domain, Array< std::pair< Boundary, Boundary >, 3 > const &bc, Info const &info=Info{})
Definition AMReX_FFT_R2X.H:125
FabArray< BaseFab< GpuComplex< T > > > cMF
Definition AMReX_FFT_R2X.H:30
R2X(R2X &&)=delete
T scalingFactor() const
Definition AMReX_FFT_R2X.H:533
R2X(R2X const &)=delete
R2X & operator=(R2X const &)=delete
void forwardThenBackward_doit_1(MF const &inmf, MF &outmf, F const &post_forward, IntVect const &ngout=IntVect(0), Periodicity const &period=Periodicity::NonPeriodic())
Definition AMReX_FFT_R2X.H:610
void forwardThenBackward_doit_0(MF const &inmf, MF &outmf, F const &post_forward, IntVect const &ngout=IntVect(0), Periodicity const &period=Periodicity::NonPeriodic())
Definition AMReX_FFT_R2X.H:558
int size() const noexcept
Return the number of FABs in the FabArray.
Definition AMReX_FabArrayBase.H:110
const DistributionMapping & DistributionMap() const noexcept
Return constant reference to associated DistributionMapping.
Definition AMReX_FabArrayBase.H:131
bool empty() const noexcept
Definition AMReX_FabArrayBase.H:89
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
An Array of FortranArrayBox(FAB)-like Objects.
Definition AMReX_FabArray.H:347
void define(const BoxArray &bxs, const DistributionMapping &dm, int nvar, int ngrow, const MFInfo &info=MFInfo(), const FabFactory< FAB > &factory=DefaultFabFactory< FAB >())
Define this FabArray identically to that performed by the constructor having an analogous function si...
Definition AMReX_FabArray.H:2171
A collection (stored as an array) of FArrayBox objects.
Definition AMReX_MultiFab.H:40
This provides length of period for periodic domains. 0 means it is not periodic in that direction....
Definition AMReX_Periodicity.H:17
static const Periodicity & NonPeriodic() noexcept
Definition AMReX_Periodicity.cpp:52
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:243
std::array< T, N > Array
Definition AMReX_Array.H:25
Definition AMReX_FFT_Helper.H:46
int MyProcSub() noexcept
my sub-rank in current frame
Definition AMReX_ParallelContext.H:76
int NProcsSub() noexcept
number of ranks in current frame
Definition AMReX_ParallelContext.H:74
__host__ __device__ void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:138
BoxND< 3 > Box
Box is an alias for amrex::BoxND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:27
double second() noexcept
Definition AMReX_Utility.cpp:940
BoxArray decompose(Box const &domain, int nboxes, Array< bool, 3 > const &decomp, bool no_overlap)
Decompose domain box into BoxArray.
Definition AMReX_BoxArray.cpp:1940
IntVectND< 3 > IntVect
IntVect is an alias for amrex::IntVectND 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
void ParallelCopy(MF &dst, MF const &src, int scomp, int dcomp, int ncomp, IntVect const &ng_src=IntVect(0), IntVect const &ng_dst=IntVect(0), Periodicity const &period=Periodicity::NonPeriodic())
dst = src w/ MPI communication
Definition AMReX_FabArrayUtility.H:2019
__host__ __device__ constexpr T elemwiseMin(T const &a, T const &b) noexcept
Definition AMReX_Algorithm.H:49
Definition AMReX_FFT_Helper.H:58
bool twod_mode
Definition AMReX_FFT_Helper.H:69
bool oned_mode
We might have a special twod_mode: nx or ny == 1 && nz > 1.
Definition AMReX_FFT_Helper.H:72
int batch_size
Batched FFT size. Only support in R2C, not R2X.
Definition AMReX_FFT_Helper.H:75
int nprocs
Max number of processes to use.
Definition AMReX_FFT_Helper.H:78
Definition AMReX_FFT_Helper.H:134
std::conditional_t< std::is_same_v< float, T >, cuComplex, cuDoubleComplex > VendorComplex
Definition AMReX_FFT_Helper.H:138
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:40
FabArray memory allocation information.
Definition AMReX_FabArray.H:66