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
19
20template <typename T> class Poisson;
21template <typename T> class PoissonHybrid;
22
30template <typename T = Real>
31class R2X
32{
33public:
34 using MF = std::conditional_t<std::is_same_v<T,Real>,
37
38 template <typename U> friend class Poisson;
39 template <typename U> friend class PoissonHybrid;
40
48 R2X (Box const& domain,
49 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> const& bc,
50 Info const& info = Info{});
51
55 ~R2X ();
56
57 R2X (R2X const&) = delete;
58 R2X (R2X &&) = delete;
59 R2X& operator= (R2X const&) = delete;
60 R2X& operator= (R2X &&) = delete;
61
67 [[nodiscard]] T scalingFactor () const;
68
69 template <typename F>
77 void forwardThenBackward (MF const& inmf, MF& outmf, F const& post_forward);
78
79 // public for cuda
80 template <int dim, typename FAB, typename F>
87 void post_forward_doit (FAB* fab, F const& f);
88
89 // private function made public for cuda
90 template <typename F>
100 void forwardThenBackward_doit_0 (MF const& inmf, MF& outmf, F const& post_forward,
101 IntVect const& ngout = IntVect(0),
102 Periodicity const& period = Periodicity::NonPeriodic());
103 template <typename F>
113 void forwardThenBackward_doit_1 (MF const& inmf, MF& outmf, F const& post_forward,
114 IntVect const& ngout = IntVect(0),
115 Periodicity const& period = Periodicity::NonPeriodic());
116
117private:
118
125 void forward (MF const& inmf, MF& outmf);
132 void forward (MF const& inmf, cMF& outmf);
138 void forward (MF const& inmf);
147 void backward (MF const& inmf, MF& outmf, IntVect const& ngout,
148 Periodicity const& period);
157 void backward (cMF const& inmf, MF& outmf, IntVect const& ngout,
158 Periodicity const& period);
162 void backward ();
163
164 Box m_dom_0;
165 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> m_bc;
166
167 Plan<T> m_fft_fwd_x{};
168 Plan<T> m_fft_bwd_x{};
169 Plan<T> m_fft_fwd_y{};
170 Plan<T> m_fft_bwd_y{};
171 Plan<T> m_fft_fwd_z{};
172 Plan<T> m_fft_bwd_z{};
173
174 std::unique_ptr<detail::TransposeMetaData> m_cmd_cx2cy;
175 std::unique_ptr<detail::TransposeMetaData> m_cmd_rx2ry;
176 std::unique_ptr<detail::TransposeMetaData> m_cmd_cy2cz;
177 std::unique_ptr<detail::TransposeMetaData> m_cmd_ry2rz;
178
179 std::unique_ptr<detail::TransposeMetaData> m_cmd_cy2cx;
180 std::unique_ptr<detail::TransposeMetaData> m_cmd_ry2rx;
181 std::unique_ptr<detail::TransposeMetaData> m_cmd_cz2cy;
182 std::unique_ptr<detail::TransposeMetaData> m_cmd_rz2ry;
183
184 Swap01 m_dtos_x2y{};
185 Swap01 m_dtos_y2x{};
186 Swap02 m_dtos_y2z{};
187 Swap02 m_dtos_z2y{};
188
189 MF m_rx;
190 MF m_ry;
191 MF m_rz;
192 cMF m_cx;
193 cMF m_cy;
194 cMF m_cz;
195
196 std::unique_ptr<char,DataDeleter> m_data_1;
197 std::unique_ptr<char,DataDeleter> m_data_2;
198
199 Box m_dom_rx;
200 Box m_dom_ry;
201 Box m_dom_rz;
202 Box m_dom_cx;
203 Box m_dom_cy;
204 Box m_dom_cz;
205
206 std::unique_ptr<R2X<T>> m_r2x_sub;
207 detail::SubHelper m_sub_helper;
208
211 MPI_Comm m_comm_active = MPI_COMM_NULL;
212 bool m_owns_comm_active = false; // true if it must be freed
213
214 Info m_info;
215};
216
217template <typename T>
218R2X<T>::R2X (Box const& domain,
219 Array<std::pair<Boundary,Boundary>,AMREX_SPACEDIM> const& bc,
220 Info const& info)
221 : m_dom_0(domain),
222 m_bc(bc),
223 m_sub_helper(domain),
224 m_info(info)
225{
226 BL_PROFILE("FFT::R2X");
227
228 static_assert(std::is_same_v<float,T> || std::is_same_v<double,T>);
229
230 AMREX_ALWAYS_ASSERT((m_dom_0.numPts() > 1) && (m_info.batch_size == 1));
232 m_dom_0.smallEnd() == IntVect(AMREX_D_DECL(0,0,0)),
233 "FFT::R2X currently requires domain.smallEnd() == 0");
234#if (AMREX_SPACEDIM == 2)
235 AMREX_ALWAYS_ASSERT(!m_info.twod_mode);
236#else
237 if (m_info.twod_mode) {
238 AMREX_ALWAYS_ASSERT((int(domain.length(0) > 1) +
239 int(domain.length(1) > 1) +
240 int(domain.length(2) > 1)) >= 2);
241 }
242#endif
243
244 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
245 if (bc[idim].first == Boundary::periodic ||
246 bc[idim].second == Boundary::periodic) {
247 AMREX_ALWAYS_ASSERT(bc[idim].first == bc[idim].second);
248 }
249 }
250
251 {
252 Box subbox = m_sub_helper.make_box(m_dom_0);
253 if (subbox.size() != m_dom_0.size()) {
254 m_r2x_sub = std::make_unique<R2X<T>>
255 (subbox, m_sub_helper.make_array(bc), info);
256 return;
257 }
258 }
259
260 int myproc = ParallelContext::MyProcSub();
261 int nprocs = std::min(ParallelContext::NProcsSub(), m_info.nprocs);
262
263 //
264 // make data containers
265 //
266
267 m_dom_rx = m_dom_0;
268 auto bax = amrex::decompose(m_dom_rx, nprocs, {AMREX_D_DECL(false,true,true)});
269 DistributionMapping dmx = detail::make_iota_distromap(bax.size());
270 m_rx.define(bax, dmx, 1, 0, MFInfo().SetAlloc(false));
271
272 // x-direction
273 if (bc[0].first == Boundary::periodic) {
274 // x-fft: r2c(m_rx->m_cx)
275 m_dom_cx = Box(IntVect(0), IntVect(AMREX_D_DECL(domain.length(0)/2,
276 domain.bigEnd(1),
277 domain.bigEnd(2))));
278 BoxList bl = bax.boxList();
279 for (auto & b : bl) {
280 b.setBig(0, m_dom_cx.bigEnd(0));
281 }
282 BoxArray cbax(std::move(bl));
283 m_cx.define(cbax, dmx, 1, 0, MFInfo().SetAlloc(false));
284 } // else: x-fft: r2r(m_rx)
285
286#if (AMREX_SPACEDIM >= 2)
287 if (domain.length(1) > 1 && !m_info.oned_mode) {
288 if (! m_cx.empty()) {
289 // copy(m_cx->m_cy)
290 m_dom_cy = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_cx.bigEnd(1),
291 m_dom_cx.bigEnd(0),
292 m_dom_cx.bigEnd(2))));
293 auto ba = amrex::decompose(m_dom_cy, nprocs, {AMREX_D_DECL(false,true,true)});
295 if (ba.size() == m_cx.size()) {
296 dm = m_cx.DistributionMap();
297 } else {
298 dm = detail::make_iota_distromap(ba.size());
299 }
300 m_cy.define(ba, dm, 1, 0, MFInfo().SetAlloc(false));
301 // if bc[1] is periodic:
302 // c2c(m_cy->m_cy)
303 // else:
304 // r2r(m_cy.re) & r2r(m_cy.im)
305 } else {
306 // copy(m_rx->m_ry)
307 m_dom_ry = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_rx.bigEnd(1),
308 m_dom_rx.bigEnd(0),
309 m_dom_rx.bigEnd(2))));
310 auto ba = amrex::decompose(m_dom_ry, nprocs, {AMREX_D_DECL(false,true,true)});
312 if (ba.size() == m_rx.size()) {
313 dm = m_rx.DistributionMap();
314 } else {
315 dm = detail::make_iota_distromap(ba.size());
316 }
317 m_ry.define(ba, dm, 1, 0, MFInfo().SetAlloc(false));
318 // if bc[1] is periodic:
319 // r2c(m_ry->m_cy)
320 // else:
321 // r2r(m_ry)
322 if (bc[1].first == Boundary::periodic) {
323 m_dom_cy = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_ry.length(0)/2,
324 m_dom_ry.bigEnd(1),
325 m_dom_ry.bigEnd(2))));
326 BoxList bl = ba.boxList();
327 for (auto & b : bl) {
328 b.setBig(0, m_dom_cy.bigEnd(0));
329 }
330 BoxArray cba(std::move(bl));
331 m_cy.define(cba, dm, 1, 0, MFInfo().SetAlloc(false));
332 }
333 }
334 }
335#endif
336
337#if (AMREX_SPACEDIM == 3)
338 if (domain.length(2) > 1 && !m_info.twod_mode) {
339 if (! m_cy.empty()) {
340 // copy(m_cy, m_cz)
341 m_dom_cz = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_cy.bigEnd(2),
342 m_dom_cy.bigEnd(1),
343 m_dom_cy.bigEnd(0))));
344 auto ba = amrex::decompose(m_dom_cz, nprocs, {AMREX_D_DECL(false,true,true)});
346 if (ba.size() == m_cy.size()) {
347 dm = m_cy.DistributionMap();
348 } else {
349 dm = detail::make_iota_distromap(ba.size());
350 }
351 m_cz.define(ba, dm, 1, 0, MFInfo().SetAlloc(false));
352 // if bc[2] is periodic:
353 // c2c(m_cz->m_cz)
354 // else:
355 // r2r(m_cz.re) & r2r(m_cz.im)
356 } else {
357 // copy(m_ry, m_rz)
358 m_dom_rz = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_ry.bigEnd(2),
359 m_dom_ry.bigEnd(1),
360 m_dom_ry.bigEnd(0))));
361 auto ba = amrex::decompose(m_dom_rz, nprocs, {AMREX_D_DECL(false,true,true)});
363 if (ba.size() == m_ry.size()) {
364 dm = m_ry.DistributionMap();
365 } else {
366 dm = detail::make_iota_distromap(ba.size());
367 }
368 m_rz.define(ba, dm, 1, 0, MFInfo().SetAlloc(false));
369 // if bc[2] is periodic:
370 // r2c(m_rz->m_cz)
371 // else:
372 // r2r(m_rz)
373 if (bc[2].first == Boundary::periodic) {
374 m_dom_cz = Box(IntVect(0), IntVect(AMREX_D_DECL(m_dom_rz.length(0)/2,
375 m_dom_rz.bigEnd(1),
376 m_dom_rz.bigEnd(2))));
377 BoxList bl = ba.boxList();
378 for (auto & b : bl) {
379 b.setBig(0, m_dom_cz.bigEnd(0));
380 }
381 BoxArray cba(std::move(bl));
382 m_cz.define(cba, dm, 1, 0, MFInfo().SetAlloc(false));
383 }
384 }
385 }
386#endif
387
388 // There are several different execution paths.
389 //
390 // (1) x-r2c(m_rx->m_cx), copy(m_cx->m_cy), y-fft(m_cy),
391 // copy(m_cy->m_cz), z-fft(m_cz)
392 // In this case, we have m_rx, m_cx, m_cy, & m_cz.
393 // we can alias(m_rx,m_cy) and alias(m_cx,m_cz).
394 //
395 // (2) x_r2r(m_rx), copy(m_rx->m_ry), y-r2c(m_ry->m_cy),
396 // copy(m_cy->m_cz), z-fft(m_cz)
397 // In this case, we have m_rx, m_ry, m_cy, & m_cz.
398 // We can alias(m_rx,m_cy) and alias(m_ry,m_cz).
399 //
400 // (3) x_r2r(m_rx), copy(m_rx->m_ry), y-r2r(m_ry),
401 // copy(m_ry->m_rz), z-r2c(m_rz->m_rz)
402 // In this case, we have m_rx, m_ry, m_rz, & m_cz
403 // We can alias(m_rx,m_rz) and alias(m_ry,m_cz)
404 //
405 // (4) x_r2r(m_rx), copy(m_rx->m_ry), y-r2r(m_ry),
406 // copy(m_ry->m_rz), z-r2r(m_rz)
407 // In this case, we have m_rx, m_ry, & m_rz.
408 // We can alias(m_rx,m_rz).
409
410 if (! m_cx.empty()) {
411 m_data_1 = detail::make_mfs_share(m_rx, m_cy);
412 m_data_2 = detail::make_mfs_share(m_cx, m_cz);
413 } else if (! m_cy.empty()) {
414 m_data_1 = detail::make_mfs_share(m_rx, m_cy);
415 m_data_2 = detail::make_mfs_share(m_ry, m_cz);
416 } else if (! m_cz.empty()) {
417 m_data_1 = detail::make_mfs_share(m_rx, m_rz);
418 m_data_2 = detail::make_mfs_share(m_ry, m_cz);
419 } else {
420 m_data_1 = detail::make_mfs_share(m_rx, m_rz);
421 m_data_2 = detail::make_mfs_share(m_ry, m_cz); // It's okay m_cz is empty.
422 }
423
424 //
425 // make copiers
426 //
427
428 int const nprocs_active = detail::count_active_procs({&m_rx, &m_ry, &m_rz, &m_cx, &m_cy, &m_cz});
429
430#if (AMREX_SPACEDIM >= 2)
431 if (!m_cy.empty() || !m_ry.empty()) {
432 if (! m_cx.empty()) {
433 // copy(m_cx->m_cy)
434 m_cmd_cx2cy = detail::make_transpose
435 (nprocs_active, m_cy, m_dom_cy, m_cx, IntVect(0), m_dtos_x2y);
436 m_cmd_cy2cx = detail::make_transpose
437 (nprocs_active, m_cx, m_dom_cx, m_cy, IntVect(0), m_dtos_y2x);
438 } else {
439 // copy(m_rx->m_ry)
440 m_cmd_rx2ry = detail::make_transpose
441 (nprocs_active, m_ry, m_dom_ry, m_rx, IntVect(0), m_dtos_x2y);
442 m_cmd_ry2rx = detail::make_transpose
443 (nprocs_active, m_rx, m_dom_rx, m_ry, IntVect(0), m_dtos_y2x);
444 }
445 }
446#endif
447
448#if (AMREX_SPACEDIM == 3)
449 if (!m_cz.empty() || !m_rz.empty()) {
450 if (! m_cy.empty()) {
451 // copy(m_cy, m_cz)
452 m_cmd_cy2cz = detail::make_transpose
453 (nprocs_active, m_cz, m_dom_cz, m_cy, IntVect(0), m_dtos_y2z);
454 m_cmd_cz2cy = detail::make_transpose
455 (nprocs_active, m_cy, m_dom_cy, m_cz, IntVect(0), m_dtos_z2y);
456 } else {
457 // copy(m_ry, m_rz)
458 m_cmd_ry2rz = detail::make_transpose
459 (nprocs_active, m_rz, m_dom_rz, m_ry, IntVect(0), m_dtos_y2z);
460 m_cmd_rz2ry = detail::make_transpose
461 (nprocs_active, m_ry, m_dom_ry, m_rz, IntVect(0), m_dtos_z2y);
462 }
463 }
464#endif
465
466 m_info.comm_strategy = detail::resolve_comm_strategy
467 (m_info.comm_strategy, std::max({detail::max_fab_bytes(m_rx),
468 detail::max_fab_bytes(m_ry),
469 detail::max_fab_bytes(m_rz),
470 detail::max_fab_bytes(m_cx),
471 detail::max_fab_bytes(m_cy),
472 detail::max_fab_bytes(m_cz)}),
473 detail::max_block_bytes({&m_cmd_cx2cy, &m_cmd_rx2ry, &m_cmd_cy2cz,
474 &m_cmd_ry2rz, &m_cmd_cy2cx, &m_cmd_ry2rx,
475 &m_cmd_cz2cy, &m_cmd_rz2ry}),
476 detail::transposes_dense({&m_cmd_cx2cy, &m_cmd_rx2ry, &m_cmd_cy2cz,
477 &m_cmd_ry2rz, &m_cmd_cy2cx, &m_cmd_ry2rx,
478 &m_cmd_cz2cy, &m_cmd_rz2ry}));
479 if (m_info.comm_strategy != CommStrategy::pointtopoint) {
480 m_comm_active = detail::make_active_comm(nprocs_active, m_owns_comm_active);
481 }
482
483 //
484 // make plans
485 //
486
487 using VendorComplex = typename Plan<T>::VendorComplex;
488
489 if (myproc < m_rx.size())
490 {
491 Box const& box = m_rx.box(myproc);
492 auto* pf = m_rx[myproc].dataPtr();
493 if (bc[0].first == Boundary::periodic) {
494 auto* pb = (VendorComplex*) m_cx[myproc].dataPtr();
495 m_fft_fwd_x.template init_r2c<Direction::forward>(box, pf, pb);
496#if defined(AMREX_USE_SYCL)
497 m_fft_bwd_x = m_fft_fwd_x;
498#else
499 m_fft_bwd_x.template init_r2c<Direction::backward>(box, pf, pb);
500#endif
501 } else {
502 m_fft_fwd_x.template init_r2r<Direction::forward>(box, pf, bc[0]);
503#if defined(AMREX_USE_GPU)
504 if ((bc[0].first == Boundary::even && bc[0].second == Boundary::odd) ||
505 (bc[0].first == Boundary::odd && bc[0].second == Boundary::even)) {
506 m_fft_bwd_x = m_fft_fwd_x;
507 } else
508#endif
509 {
510 m_fft_bwd_x.template init_r2r<Direction::backward>(box, pf, bc[0]);
511 }
512 }
513 }
514
515#if (AMREX_SPACEDIM >= 2)
516 if (m_ry.empty() && m_bc[1].first == Boundary::periodic) {
517 if (myproc < m_cy.size()) {
518 Box const& box = m_cy.box(myproc);
519 auto* p = (VendorComplex *)m_cy[myproc].dataPtr();
520 m_fft_fwd_y.template init_c2c<Direction::forward>(box, p);
521#if defined(AMREX_USE_SYCL)
522 m_fft_bwd_y = m_fft_fwd_y;
523#else
524 m_fft_bwd_y.template init_c2c<Direction::backward>(box, p);
525#endif
526 }
527 } else if (!m_ry.empty() && m_bc[1].first == Boundary::periodic) {
528 if (myproc < m_ry.size()) {
529 Box const& box = m_ry.box(myproc);
530 auto* pr = m_ry[myproc].dataPtr();
531 auto* pc = (VendorComplex*)m_cy[myproc].dataPtr();
532 m_fft_fwd_y.template init_r2c<Direction::forward>(box, pr, pc);
533#if defined(AMREX_USE_SYCL)
534 m_fft_bwd_y = m_fft_fwd_y;
535#else
536 m_fft_bwd_y.template init_r2c<Direction::backward>(box, pr, pc);
537#endif
538 }
539 } else if (!m_cy.empty()) {
540 if (myproc < m_cy.size()) {
541 Box const& box = m_cy.box(myproc);
542 auto* p = (VendorComplex*) m_cy[myproc].dataPtr();
543 m_fft_fwd_y.template init_r2r<Direction::forward>(box, p, bc[1]);
544#if defined(AMREX_USE_GPU)
545 if ((bc[1].first == Boundary::even && bc[1].second == Boundary::odd) ||
546 (bc[1].first == Boundary::odd && bc[1].second == Boundary::even)) {
547 m_fft_bwd_y = m_fft_fwd_y;
548 } else
549#endif
550 {
551 m_fft_bwd_y.template init_r2r<Direction::backward>(box, p, bc[1]);
552 }
553 }
554 } else {
555 if (myproc < m_ry.size()) {
556 Box const& box = m_ry.box(myproc);
557 auto* p = m_ry[myproc].dataPtr();
558 m_fft_fwd_y.template init_r2r<Direction::forward>(box, p, bc[1]);
559#if defined(AMREX_USE_GPU)
560 if ((bc[1].first == Boundary::even && bc[1].second == Boundary::odd) ||
561 (bc[1].first == Boundary::odd && bc[1].second == Boundary::even)) {
562 m_fft_bwd_y = m_fft_fwd_y;
563 } else
564#endif
565 {
566 m_fft_bwd_y.template init_r2r<Direction::backward>(box, p, bc[1]);
567 }
568 }
569 }
570#endif
571
572#if (AMREX_SPACEDIM == 3)
573 if (m_rz.empty() && m_bc[2].first == Boundary::periodic) {
574 if (myproc < m_cz.size()) {
575 Box const& box = m_cz.box(myproc);
576 auto* p = (VendorComplex*)m_cz[myproc].dataPtr();
577 m_fft_fwd_z.template init_c2c<Direction::forward>(box, p);
578#if defined(AMREX_USE_SYCL)
579 m_fft_bwd_z = m_fft_fwd_z;
580#else
581 m_fft_bwd_z.template init_c2c<Direction::backward>(box, p);
582#endif
583 }
584 } else if (!m_rz.empty() && m_bc[2].first == Boundary::periodic) {
585 if (myproc < m_rz.size()) {
586 Box const& box = m_rz.box(myproc);
587 auto* pr = m_rz[myproc].dataPtr();
588 auto* pc = (VendorComplex*)m_cz[myproc].dataPtr();
589 m_fft_fwd_z.template init_r2c<Direction::forward>(box, pr, pc);
590#if defined(AMREX_USE_SYCL)
591 m_fft_bwd_z = m_fft_fwd_z;
592#else
593 m_fft_bwd_z.template init_r2c<Direction::backward>(box, pr, pc);
594#endif
595 }
596 } else if (!m_cz.empty()) {
597 if (myproc < m_cz.size()) {
598 Box const& box = m_cz.box(myproc);
599 auto* p = (VendorComplex*) m_cz[myproc].dataPtr();
600 m_fft_fwd_z.template init_r2r<Direction::forward>(box, p, bc[2]);
601#if defined(AMREX_USE_GPU)
602 if ((bc[2].first == Boundary::even && bc[2].second == Boundary::odd) ||
603 (bc[2].first == Boundary::odd && bc[2].second == Boundary::even)) {
604 m_fft_bwd_z = m_fft_fwd_z;
605 } else
606#endif
607 {
608 m_fft_bwd_z.template init_r2r<Direction::backward>(box, p, bc[2]);
609 }
610 }
611 } else {
612 if (myproc < m_rz.size()) {
613 Box const& box = m_rz.box(myproc);
614 auto* p = m_rz[myproc].dataPtr();
615 m_fft_fwd_z.template init_r2r<Direction::forward>(box, p, bc[2]);
616#if defined(AMREX_USE_GPU)
617 if ((bc[2].first == Boundary::even && bc[2].second == Boundary::odd) ||
618 (bc[2].first == Boundary::odd && bc[2].second == Boundary::even)) {
619 m_fft_bwd_z = m_fft_fwd_z;
620 } else
621#endif
622 {
623 m_fft_bwd_z.template init_r2r<Direction::backward>(box, p, bc[2]);
624 }
625 }
626 }
627#endif
628}
629
630template <typename T>
632{
633 if (m_fft_bwd_x.plan != m_fft_fwd_x.plan) {
634 m_fft_bwd_x.destroy();
635 }
636 if (m_fft_bwd_y.plan != m_fft_fwd_y.plan) {
637 m_fft_bwd_y.destroy();
638 }
639 if (m_fft_bwd_z.plan != m_fft_fwd_z.plan) {
640 m_fft_bwd_z.destroy();
641 }
642 m_fft_fwd_x.destroy();
643 m_fft_fwd_y.destroy();
644 m_fft_fwd_z.destroy();
645 detail::free_active_comm(m_comm_active, m_owns_comm_active);
646}
647
648template <typename T>
650{
651 // The sub-object knows which axes are actually transformed.
652 if (m_r2x_sub) { return m_r2x_sub->scalingFactor(); }
653
654 // Same rules as the constructor: y is skipped in oned_mode, z in twod_mode.
655 Long r = 1;
656 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
657 if ((idim == 1 && m_info.oned_mode) ||
658 (idim == 2 && m_info.twod_mode) ||
659 (m_dom_0.length(idim) == 1)) { continue; }
660 r *= m_dom_0.length(idim);
661 if (m_bc[idim].first != Boundary::periodic) { r *= 2; }
662 }
663 return T(1)/T(r);
664}
665
666template <typename T>
667template <typename F>
668void R2X<T>::forwardThenBackward (MF const& inmf, MF& outmf, F const& post_forward)
669{
671 !m_info.twod_mode,
672 "FFT::R2X::forwardThenBackward(post_forward) does not support twod_mode yet");
673 forwardThenBackward_doit_0(inmf, outmf, post_forward);
674}
675
676template <typename T>
677template <typename F>
678void R2X<T>::forwardThenBackward_doit_0 (MF const& inmf, MF& outmf,
679 F const& post_forward,
680 IntVect const& ngout,
681 Periodicity const& period)
682{
683 BL_PROFILE("FFT::R2X::forwardbackward_0");
684
685 if (m_r2x_sub) {
686 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
687 MF inmf_sub, inmf_tmp;
688 if (inmf_safe) {
689 inmf_sub = m_sub_helper.make_alias_mf(inmf);
690 } else {
691 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
692 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
693 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
694 }
695
696 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
697 MF outmf_sub, outmf_tmp;
698 if (outmf_safe) {
699 outmf_sub = m_sub_helper.make_alias_mf(outmf);
700 } else {
701 IntVect const& ngtmp = m_sub_helper.make_safe_ghost(outmf.nGrowVect());
702 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, ngtmp);
703 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
704 }
705
706 IntVect const& subngout = m_sub_helper.make_iv(ngout);
707 Periodicity const& subperiod = m_sub_helper.make_periodicity(period);
708 GpuArray<int,3> const& order = m_sub_helper.xyz_order();
709 m_r2x_sub->forwardThenBackward_doit_1
710 (inmf_sub, outmf_sub,
711 [=] AMREX_GPU_DEVICE (int i, int j, int k, auto& sp)
712 {
713 GpuArray<int,3> idx{i,j,k};
714 post_forward(idx[order[0]], idx[order[1]], idx[order[2]], sp);
715 },
716 subngout, subperiod);
717
718 if (!outmf_safe) {
719 outmf.LocalCopy(outmf_tmp, 0, 0, 1, outmf_tmp.nGrowVect());
720 }
721 }
722 else
723 {
724 this->forwardThenBackward_doit_1(inmf, outmf, post_forward, ngout, period);
725 }
726}
727
728template <typename T>
729template <typename F>
730void R2X<T>::forwardThenBackward_doit_1 (MF const& inmf, MF& outmf,
731 F const& post_forward,
732 IntVect const& ngout,
733 Periodicity const& period)
734{
735 BL_PROFILE("FFT::R2X::forwardbackward_1");
736
737 if (m_r2x_sub) {
738 amrex::Abort("R2X::forwardThenBackward_doit_1: How did this happen?");
739 }
740 else
741 {
742 this->forward(inmf);
743
744 // post-forward
745
746 // The post-forward data live in the last transformed direction. In
747 // oned_mode, y (and z) are batched, so their MultiFabs are not defined.
748 int actual_dim = 1;
749#if (AMREX_SPACEDIM >= 2)
750 if (!m_ry.empty() || !m_cy.empty()) { actual_dim = 2; }
751#endif
752#if (AMREX_SPACEDIM == 3)
753 if (!m_rz.empty() || !m_cz.empty()) { actual_dim = 3; }
754#endif
755
756 if (actual_dim == 1) {
757 if (m_cx.empty()) {
758 post_forward_doit<0>(detail::get_fab(m_rx), post_forward);
759 } else {
760 post_forward_doit<0>(detail::get_fab(m_cx), post_forward);
761 }
762 }
763#if (AMREX_SPACEDIM >= 2)
764 else if (actual_dim == 2) {
765 if (m_cy.empty()) {
766 post_forward_doit<1>(detail::get_fab(m_ry), post_forward);
767 } else {
768 post_forward_doit<1>(detail::get_fab(m_cy), post_forward);
769 }
770 }
771#endif
772#if (AMREX_SPACEDIM == 3)
773 else if (actual_dim == 3) {
774 if (m_cz.empty()) {
775 post_forward_doit<2>(detail::get_fab(m_rz), post_forward);
776 } else {
777 post_forward_doit<2>(detail::get_fab(m_cz), post_forward);
778 }
779 }
780#endif
781
782 this->backward();
783
784 outmf.ParallelCopy(m_rx, 0, 0, 1, IntVect(0),
785 amrex::elemwiseMin(ngout,outmf.nGrowVect()), period);
786 }
787}
788
789template <typename T>
790void R2X<T>::forward (MF const& inmf)
791{
792 BL_PROFILE("FFT::R2X::forward");
793
794 if (m_r2x_sub) {
795 if (m_sub_helper.ghost_safe(inmf.nGrowVect())) {
796 m_r2x_sub->forward(m_sub_helper.make_alias_mf(inmf));
797 } else {
798 MF tmp(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
799 tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
800 m_r2x_sub->forward(m_sub_helper.make_alias_mf(tmp));
801 }
802 return;
803 }
804
805 m_rx.ParallelCopy(inmf, 0, 0, 1);
806 if (m_bc[0].first == Boundary::periodic) {
807 m_fft_fwd_x.template compute_r2c<Direction::forward>();
808 } else {
809 m_fft_fwd_x.template compute_r2r<Direction::forward>();
810 }
811
812#if (AMREX_SPACEDIM >= 2)
813 if ( m_cmd_cx2cy) {
814 detail::transpose(m_cy, m_cx, *m_cmd_cx2cy, 1, m_dtos_x2y, m_info.comm_strategy, m_comm_active);
815 } else if ( m_cmd_rx2ry) {
816 detail::transpose(m_ry, m_rx, *m_cmd_rx2ry, 1, m_dtos_x2y, m_info.comm_strategy, m_comm_active);
817 }
818 if (m_bc[1].first != Boundary::periodic)
819 {
820 m_fft_fwd_y.template compute_r2r<Direction::forward>();
821 }
822 else if (m_bc[0].first == Boundary::periodic)
823 {
824 m_fft_fwd_y.template compute_c2c<Direction::forward>();
825 }
826 else
827 {
828 m_fft_fwd_y.template compute_r2c<Direction::forward>();
829 }
830#endif
831
832#if (AMREX_SPACEDIM == 3)
833 if ( m_cmd_cy2cz) {
834 detail::transpose(m_cz, m_cy, *m_cmd_cy2cz, 1, m_dtos_y2z, m_info.comm_strategy, m_comm_active);
835 } else if ( m_cmd_ry2rz) {
836 detail::transpose(m_rz, m_ry, *m_cmd_ry2rz, 1, m_dtos_y2z, m_info.comm_strategy, m_comm_active);
837 }
838 if (m_bc[2].first != Boundary::periodic)
839 {
840 m_fft_fwd_z.template compute_r2r<Direction::forward>();
841 }
842 else if (m_bc[0].first == Boundary::periodic ||
843 m_bc[1].first == Boundary::periodic)
844 {
845 m_fft_fwd_z.template compute_c2c<Direction::forward>();
846 }
847 else
848 {
849 m_fft_fwd_z.template compute_r2c<Direction::forward>();
850 }
851#endif
852}
853
854template <typename T>
855void R2X<T>::forward (MF const& inmf, MF& outmf)
856{
857 if (m_r2x_sub)
858 {
859 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
860 MF inmf_sub, inmf_tmp;
861 if (inmf_safe) {
862 inmf_sub = m_sub_helper.make_alias_mf(inmf);
863 } else {
864 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
865 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
866 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
867 }
868
869 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
870 MF outmf_sub, outmf_tmp;
871 if (outmf_safe) {
872 outmf_sub = m_sub_helper.make_alias_mf(outmf);
873 } else {
874 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, 0);
875 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
876 }
877
878 m_r2x_sub->forward(inmf_sub, outmf_sub);
879
880 if (!outmf_safe) {
881 outmf.LocalCopy(outmf_tmp, 0, 0, 1, IntVect(0));
882 }
883 }
884 else
885 {
886 this->forward(inmf);
887
888#if (AMREX_SPACEDIM == 3)
889 if (m_info.twod_mode) {
890 if (m_cy.empty() && !m_ry.empty()) {
891 ParallelCopy(outmf, m_dom_rx, m_ry, 0, 0, 1, IntVect(0), Swap01{});
892 } else if (m_ry.empty() && m_cy.empty() && m_cx.empty()) {
893 outmf.ParallelCopy(m_rx, 0, 0, 1);
894 } else {
895 amrex::Abort("R2X::forward(MF,MF): How did this happen?");
896 }
897 } else
898#endif
899 {
901 amrex::Abort("R2X::forward(MF,MF): TODO");
902 }
903 }
904}
905
906template <typename T>
907void R2X<T>::forward (MF const& inmf, cMF& outmf)
908{
909 if (m_r2x_sub)
910 {
911 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
912 MF inmf_sub, inmf_tmp;
913 if (inmf_safe) {
914 inmf_sub = m_sub_helper.make_alias_mf(inmf);
915 } else {
916 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
917 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
918 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
919 }
920
921 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
922 cMF outmf_sub, outmf_tmp;
923 if (outmf_safe) {
924 outmf_sub = m_sub_helper.make_alias_mf(outmf);
925 } else {
926 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, 0);
927 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
928 }
929
930 m_r2x_sub->forward(inmf_sub, outmf_sub);
931
932 if (!outmf_safe) {
933 outmf.LocalCopy(outmf_tmp, 0, 0, 1, IntVect(0));
934 }
935 }
936 else
937 {
938 this->forward(inmf);
939
940#if (AMREX_SPACEDIM == 3)
941 if (m_info.twod_mode) {
942 if (!m_cy.empty()) {
943 auto lo = m_dom_cy.smallEnd();
944 auto hi = m_dom_cy.bigEnd();
945 std::swap(lo[0],lo[1]);
946 std::swap(hi[0],hi[1]);
947 Box dom(lo,hi);
948 ParallelCopy(outmf, dom, m_cy, 0, 0, 1, IntVect(0), Swap01{});
949 } else if (m_ry.empty() && m_cy.empty() && !m_cx.empty()) {
950 outmf.ParallelCopy(m_cx, 0, 0, 1);
951 } else {
952 amrex::Abort("R2X::forward(MF,cMF): How did this happen?");
953 }
954 } else
955#endif
956 {
958 amrex::Abort("R2X::forward(MF,cMF): TODO");
959 }
960 }
961}
962
963template <typename T>
964void R2X<T>::backward ()
965{
966 BL_PROFILE("FFT::R2X::backward");
967
968 AMREX_ALWAYS_ASSERT(m_r2x_sub == nullptr);
969
970#if (AMREX_SPACEDIM == 3)
971 if (m_bc[2].first != Boundary::periodic)
972 {
973 m_fft_bwd_z.template compute_r2r<Direction::backward>();
974 }
975 else if (m_bc[0].first == Boundary::periodic ||
976 m_bc[1].first == Boundary::periodic)
977 {
978 m_fft_bwd_z.template compute_c2c<Direction::backward>();
979 }
980 else
981 {
982 m_fft_bwd_z.template compute_r2c<Direction::backward>();
983 }
984 if ( m_cmd_cz2cy) {
985 detail::transpose(m_cy, m_cz, *m_cmd_cz2cy, 1, m_dtos_z2y, m_info.comm_strategy, m_comm_active);
986 } else if ( m_cmd_rz2ry) {
987 detail::transpose(m_ry, m_rz, *m_cmd_rz2ry, 1, m_dtos_z2y, m_info.comm_strategy, m_comm_active);
988 }
989#endif
990
991#if (AMREX_SPACEDIM >= 2)
992 if (m_bc[1].first != Boundary::periodic)
993 {
994 m_fft_bwd_y.template compute_r2r<Direction::backward>();
995 }
996 else if (m_bc[0].first == Boundary::periodic)
997 {
998 m_fft_bwd_y.template compute_c2c<Direction::backward>();
999 }
1000 else
1001 {
1002 m_fft_bwd_y.template compute_r2c<Direction::backward>();
1003 }
1004 if ( m_cmd_cy2cx) {
1005 detail::transpose(m_cx, m_cy, *m_cmd_cy2cx, 1, m_dtos_y2x, m_info.comm_strategy, m_comm_active);
1006 } else if ( m_cmd_ry2rx) {
1007 detail::transpose(m_rx, m_ry, *m_cmd_ry2rx, 1, m_dtos_y2x, m_info.comm_strategy, m_comm_active);
1008 }
1009#endif
1010
1011 if (m_bc[0].first == Boundary::periodic) {
1012 m_fft_bwd_x.template compute_r2c<Direction::backward>();
1013 } else {
1014 m_fft_bwd_x.template compute_r2r<Direction::backward>();
1015 }
1016}
1017
1018template <typename T>
1019void R2X<T>::backward (MF const& inmf, MF& outmf, IntVect const& ngout,
1020 Periodicity const& period)
1021{
1022 if (m_r2x_sub)
1023 {
1024 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
1025 MF inmf_sub, inmf_tmp;
1026 if (inmf_safe) {
1027 inmf_sub = m_sub_helper.make_alias_mf(inmf);
1028 } else {
1029 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
1030 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
1031 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
1032 }
1033
1034 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
1035 MF outmf_sub, outmf_tmp;
1036 if (outmf_safe) {
1037 outmf_sub = m_sub_helper.make_alias_mf(outmf);
1038 } else {
1039 IntVect const& ngtmp = m_sub_helper.make_safe_ghost(outmf.nGrowVect());
1040 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, ngtmp);
1041 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
1042 }
1043
1044 IntVect const& subngout = m_sub_helper.make_iv(ngout);
1045 Periodicity const& subperiod = m_sub_helper.make_periodicity(period);
1046 m_r2x_sub->backward(inmf_sub, outmf_sub, subngout, subperiod);
1047
1048 if (!outmf_safe) {
1049 outmf.LocalCopy(outmf_tmp, 0, 0, 1, outmf_tmp.nGrowVect());
1050 }
1051 }
1052 else
1053 {
1054#if (AMREX_SPACEDIM == 3)
1055 if (m_info.twod_mode) {
1056 if (m_cy.empty() && !m_ry.empty()) {
1057 ParallelCopy(m_ry, m_dom_ry, inmf, 0, 0, 1, IntVect(0), Swap01{});
1058 } else if (m_ry.empty() && m_cy.empty() && m_cx.empty()) {
1059 m_rx.ParallelCopy(inmf, 0, 0, 1);
1060 } else {
1061 amrex::Abort("R2X::backward(MF,MF): How did this happen?");
1062 }
1063 } else
1064#endif
1065 {
1066 amrex::ignore_unused(inmf,outmf,ngout,period);
1067 amrex::Abort("R2X::backward(MF,MF): TODO");
1068 }
1069
1070 this->backward();
1071
1072 outmf.ParallelCopy(m_rx, 0, 0, 1, IntVect(0),
1073 amrex::elemwiseMin(ngout,outmf.nGrowVect()), period);
1074 }
1075}
1076
1077template <typename T>
1078void R2X<T>::backward (cMF const& inmf, MF& outmf, IntVect const& ngout,
1079 Periodicity const& period)
1080{
1081 if (m_r2x_sub)
1082 {
1083 bool inmf_safe = m_sub_helper.ghost_safe(inmf.nGrowVect());
1084 cMF inmf_sub, inmf_tmp;
1085 if (inmf_safe) {
1086 inmf_sub = m_sub_helper.make_alias_mf(inmf);
1087 } else {
1088 inmf_tmp.define(inmf.boxArray(), inmf.DistributionMap(), 1, 0);
1089 inmf_tmp.LocalCopy(inmf, 0, 0, 1, IntVect(0));
1090 inmf_sub = m_sub_helper.make_alias_mf(inmf_tmp);
1091 }
1092
1093 bool outmf_safe = m_sub_helper.ghost_safe(outmf.nGrowVect());
1094 MF outmf_sub, outmf_tmp;
1095 if (outmf_safe) {
1096 outmf_sub = m_sub_helper.make_alias_mf(outmf);
1097 } else {
1098 IntVect const& ngtmp = m_sub_helper.make_safe_ghost(outmf.nGrowVect());
1099 outmf_tmp.define(outmf.boxArray(), outmf.DistributionMap(), 1, ngtmp);
1100 outmf_sub = m_sub_helper.make_alias_mf(outmf_tmp);
1101 }
1102
1103 IntVect const& subngout = m_sub_helper.make_iv(ngout);
1104 Periodicity const& subperiod = m_sub_helper.make_periodicity(period);
1105 m_r2x_sub->backward(inmf_sub, outmf_sub, subngout, subperiod);
1106
1107 if (!outmf_safe) {
1108 outmf.LocalCopy(outmf_tmp, 0, 0, 1, outmf_tmp.nGrowVect());
1109 }
1110 }
1111 else
1112 {
1113#if (AMREX_SPACEDIM == 3)
1114 if (m_info.twod_mode) {
1115 if (!m_cy.empty()) {
1116 ParallelCopy(m_cy, m_dom_cy, inmf, 0, 0, 1, IntVect(0), Swap01{});
1117 } else if (m_ry.empty() && m_cy.empty() && !m_cx.empty()) {
1118 m_cx.ParallelCopy(inmf, 0, 0, 1);
1119 } else {
1120 amrex::Abort("R2X::backward(cMF,MF): How did this happen?");
1121 }
1122 } else
1123#endif
1124 {
1125 amrex::ignore_unused(inmf,outmf,ngout,period);
1126 amrex::Abort("R2X::backward(cMF,MF): TODO");
1127 }
1128
1129 this->backward();
1130
1131 outmf.ParallelCopy(m_rx, 0, 0, 1, IntVect(0),
1132 amrex::elemwiseMin(ngout,outmf.nGrowVect()), period);
1133 }
1134}
1135
1136template <typename T>
1137template <int dim, typename FAB, typename F>
1138void R2X<T>::post_forward_doit (FAB* fab, F const& f)
1139{
1140 if (m_info.twod_mode) {
1141 amrex::Abort("xxxxx post_forward_doit: todo");
1142 }
1143 if (fab) {
1144 auto const& a = fab->array();
1145 ParallelForOMP(fab->box(),
1146 [f=f,a=a] AMREX_GPU_DEVICE (int i, int j, int k)
1147 {
1148 if constexpr (dim == 0) {
1149 f(i,j,k,a(i,j,k));
1150 } else if constexpr (dim == 1) {
1151 f(j,i,k,a(i,j,k));
1152 } else {
1153 f(j,k,i,a(i,j,k));
1154 }
1155 });
1156 }
1157}
1158
1159}
1160
1161#endif
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:593
#define AMREX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition AMReX_BLassert.H:49
#define AMREX_ALWAYS_ASSERT(EX)
Definition AMReX_BLassert.H:50
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
while parser_exe_t p
Definition AMReX_Parser_Exe_Body.H:33
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
Reference-counted collection of Boxes.
Definition AMReX_BoxArray.H:681
A list of Boxes sharing a common IndexType.
Definition AMReX_BoxList.H:109
__host__ __device__ const IntVectND< dim > & bigEnd() const &noexcept
Return the inclusive upper bound of the box.
Definition AMReX_Box.H:136
__host__ __device__ IntVectND< dim > length() const noexcept
Return the length of the BoxND.
Definition AMReX_Box.H:167
__host__ __device__ IntVectND< dim > size() const noexcept
Return the length of the BoxND.
Definition AMReX_Box.H:160
Calculates the distribution of FABs to MPI processes.
Definition AMReX_DistributionMapping.H:51
void define(const BoxArray &boxes, int nprocs=ParallelDescriptor::NProcs())
Build a mapping from a BoxArray using the current strategy.
Definition AMReX_DistributionMapping.cpp:347
3D Poisson solver for periodic, Dirichlet & Neumann boundaries in the first two dimensions,...
Definition AMReX_FFT_Poisson.H:174
Poisson solver for periodic, Dirichlet & Neumann boundaries using FFT.
Definition AMReX_FFT_Poisson.H:32
void post_forward_doit(FAB *fab, F const &f)
Apply a user functor to a FAB after the forward transform along dimension dim.
Definition AMReX_FFT_R2X.H:1138
friend class Poisson
Definition AMReX_FFT_R2X.H:38
~R2X()
Destroy any FFT resources held by this object.
Definition AMReX_FFT_R2X.H:631
void forwardThenBackward(MF const &inmf, MF &outmf, F const &post_forward)
Execute forward transform, apply post_forward, then backward transform.
Definition AMReX_FFT_R2X.H:668
R2X(Box const &domain, Array< std::pair< Boundary, Boundary >, 3 > const &bc, Info const &info=Info{})
Build an FFT plan for the given domain and boundary types.
Definition AMReX_FFT_R2X.H:218
FabArray< BaseFab< GpuComplex< T > > > cMF
Definition AMReX_FFT_R2X.H:36
R2X(R2X &&)=delete
T scalingFactor() const
Normalization applied after a forward/backward sequence.
Definition AMReX_FFT_R2X.H:649
R2X(R2X const &)=delete
std::conditional_t< std::is_same_v< T, Real >, MultiFab, FabArray< BaseFab< T > > > MF
Definition AMReX_FFT_R2X.H:34
R2X & operator=(R2X const &)=delete
friend class PoissonHybrid
Definition AMReX_FFT_R2X.H:39
void forwardThenBackward_doit_1(MF const &inmf, MF &outmf, F const &post_forward, IntVect const &ngout=IntVect(0), Periodicity const &period=Periodicity::NonPeriodic())
CUDA-visible helper that operates on the complex FAB path.
Definition AMReX_FFT_R2X.H:730
void forwardThenBackward_doit_0(MF const &inmf, MF &outmf, F const &post_forward, IntVect const &ngout=IntVect(0), Periodicity const &period=Periodicity::NonPeriodic())
CUDA-visible helper that performs the forward/modify/backward cycle.
Definition AMReX_FFT_R2X.H:678
An Array of FortranArrayBox(FAB)-like Objects.
Definition AMReX_FabArray.H:356
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:328
std::array< T, N > Array
Definition AMReX_Array.H:31
__host__ __device__ constexpr T elemwiseMin(T const &a, T const &b) noexcept
Definition AMReX_Algorithm.H:75
Definition AMReX_FFT_CrossProj.H:8
@ pointtopoint
Definition AMReX_FFT_Helper.H:62
@ backward
Definition AMReX_FFT_Helper.H:57
@ forward
Definition AMReX_FFT_Helper.H:57
@ odd
Definition AMReX_FFT_Helper.H:64
@ even
Definition AMReX_FFT_Helper.H:64
@ periodic
Definition AMReX_FFT_Helper.H:64
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
int MPI_Comm
Definition AMReX_ccse-mpi.H:51
static constexpr int MPI_COMM_NULL
Definition AMReX_ccse-mpi.H:59
__host__ __device__ void ignore_unused(const Ts &...)
No-op helper that marks variables as intentionally unused.
Definition AMReX.H:273
@ F
Definition AMReX_Order.H:7
BoxND< 3 > Box
Box is an alias for amrex::BoxND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:35
double second() noexcept
Definition AMReX_Utility.cpp:919
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:2256
BoxArray decompose(Box const &domain, int nboxes, Array< bool, 3 > const &decomp, bool no_overlap)
Decompose domain box into BoxArray.
Definition AMReX_BoxArray.cpp:1961
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:244
Definition AMReX_FFT_Helper.H:88
Definition AMReX_FFT_Helper.H:377
std::conditional_t< std::is_same_v< float, T >, cuComplex, cuDoubleComplex > VendorComplex
Definition AMReX_FFT_Helper.H:380
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:52
FabArray memory allocation information.
Definition AMReX_FabArray.H:73