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