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