Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_AlgMG.H
Go to the documentation of this file.
1#ifndef AMREX_ALGMG_H_
2#define AMREX_ALGMG_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_Algebra.H>
6#include <AMReX_Enum.H>
7#include <AMReX_BiCGStab_MV.H>
8#include <AMReX_BLProfiler.H>
9#include <AMReX_CArena.H>
10#include <AMReX_GMRES_MV.H>
11#include <AMReX_PCG_MV.H>
13#include <AMReX_Scan.H>
14#include <AMReX_Smoother_MV.H>
15#include <AMReX_Vector.H>
16
17#include <cmath>
18#include <cstdint>
19#include <iomanip>
20#include <limits>
21#include <numeric>
22#include <memory>
23#include <stdexcept>
24#include <string>
25#include <type_traits>
26#include <utility>
27
28namespace amrex {
29
30namespace detail {
31
35template <typename TC, typename Index, typename T>
37bool pmis_blocked (CsrView<TC const,Index> const& csr, Long row, T wi, Long ig,
38 T const* AMREX_RESTRICT w, int const* AMREX_RESTRICT state,
39 Long const* AMREX_RESTRICT gmap, Long goff, int undecided)
40{
41 for (Long idx = csr.row_offset[row]; idx < csr.row_offset[row+1]; ++idx) {
42 Long j = csr.col_index[idx];
43 Long jg = gmap ? gmap[j] : j + goff;
44 if (jg != ig && state[j] == undecided &&
45 (w[j] > wi || (w[j] == wi && jg > ig))) {
46 return true;
47 }
48 }
49 return false;
50}
51
54template <typename T>
56T pmis_tiebreak (Long ig)
57{
58 auto z = static_cast<std::uint64_t>(ig) + 0x9E3779B97F4A7C15ULL;
59 z = (z ^ (z >> 30)) * 0xBF58476D1CE4E5B9ULL;
60 z = (z ^ (z >> 27)) * 0x94D049BB133111EBULL;
61 z ^= (z >> 31);
62 constexpr int digits = std::numeric_limits<T>::digits;
63 static_assert(digits < 64, "pmis_tiebreak: T must be float or double");
64 return T(z >> (64 - digits)) / T(std::uint64_t(1) << digits);
65}
66
68template <typename T, typename Index>
70bool csr_row_has_state (CsrView<T const,Index> const& csr, Long row,
71 int const* AMREX_RESTRICT state, int c)
72{
73 for (Long idx = csr.row_offset[row]; idx < csr.row_offset[row+1]; ++idx) {
74 if (state[csr.col_index[idx]] == c) { return true; }
75 }
76 return false;
77}
78
80inline std::pair<std::size_t,std::size_t> arena_bytes ()
81{
82 if (auto* carena = dynamic_cast<CArena*>(The_Arena())) {
83 return {carena->heap_space_actually_used(), carena->heap_space_used()};
84 }
85 return {0, 0};
86}
87
90template <typename T, typename TS>
91struct DirectInterpRow
92{
93 ParCsr<T const> a;
94 ParCsr<TS const> s;
95 int const* AMREX_RESTRICT state; // local rows
96 Long const* AMREX_RESTRICT cidx; // local coarse index, local rows
97 int const* AMREX_RESTRICT state_r; // A's remote columns
98 int nlocal_c; // number of local coarse columns
99 int coarse;
100
104 template <typename F>
105 AMREX_GPU_HOST_DEVICE void walk (Long i, T& diag, F const& f) const
106 {
107 diag = T(0);
108 {
109 auto const& ac = a.csr0;
110 auto const& sc = s.csr0;
111 Long sp = sc.row_offset[i];
112 Long const se = sc.row_offset[i+1];
113 for (Long idx = ac.row_offset[i]; idx < ac.row_offset[i+1]; ++idx) {
114 Long const j = ac.col_index[idx];
115 T const v = ac.mat[idx];
116 if (j == i) { diag = v; continue; }
117 while (sp < se && sc.col_index[sp] < j) { ++sp; }
118 bool const strong_c = (sp < se && sc.col_index[sp] == j
119 && state[j] == coarse);
120 f(int(cidx[j]), v, strong_c);
121 }
122 }
123 if (has_remote_row(a, i)) {
124 auto const& ac = a.csr1;
125 auto const& sc = s.csr1;
126 Long const ii = a.row_map[i];
127 Long sp = 0, se = 0;
128 if (has_remote_row(s, i)) {
129 Long const jj = s.row_map[i];
130 sp = sc.row_offset[jj];
131 se = sc.row_offset[jj+1];
132 }
133 for (Long idx = ac.row_offset[ii]; idx < ac.row_offset[ii+1]; ++idx) {
134 Long const j = ac.col_index[idx];
135 Long const jg = a.col_map[j];
136 T const v = ac.mat[idx];
137 while (sp < se && s.col_map[sc.col_index[sp]] < jg) { ++sp; }
138 bool const strong_c = (sp < se && s.col_map[sc.col_index[sp]] == jg
139 && state_r[j] == coarse);
140 f(int(nlocal_c + j), v, strong_c);
141 }
142 }
143 }
144};
145
150template <typename T, typename TS>
151struct MMInterpRow
152{
153 ParCsr<TS const> s;
154 ParCsr<TS const> st;
155 int const* AMREX_RESTRICT state;
156 T const* AMREX_RESTRICT dbeta;
157 Long const* AMREX_RESTRICT cidx; // local coarse index, local rows
158 int const* AMREX_RESTRICT state_r; // S's remote columns
159 T const* AMREX_RESTRICT dbeta_r;
160 int nlocal_f; // local fine columns
161 int nlocal_c; // local coarse columns
162
166 template <typename F>
167 AMREX_GPU_HOST_DEVICE void walk (Long i, F const& f) const
168 {
169 {
170 auto const& sc = s.csr0;
171 auto const& tc = st.csr0;
172 Long tp = tc.row_offset[i];
173 Long const te = tc.row_offset[i+1];
174 for (Long idx = sc.row_offset[i]; idx < sc.row_offset[i+1]; ++idx) {
175 Long const j = sc.col_index[idx];
176 while (tp < te && tc.col_index[tp] < j) { ++tp; }
177 T const aji = (tp < te && tc.col_index[tp] == j) ? T(tc.mat[tp]) : T(0);
178 f(int(j), int(cidx[j]), T(sc.mat[idx]), aji, state[j], dbeta[j]);
179 }
180 }
181 if (has_remote_row(s, i)) {
182 auto const& sc = s.csr1;
183 auto const& tc = st.csr1;
184 Long const ii = s.row_map[i];
185 Long tp = 0, te = 0;
186 if (has_remote_row(st, i)) {
187 Long const jj = st.row_map[i];
188 tp = tc.row_offset[jj];
189 te = tc.row_offset[jj+1];
190 }
191 for (Long idx = sc.row_offset[ii]; idx < sc.row_offset[ii+1]; ++idx) {
192 Long const j = sc.col_index[idx];
193 Long const jg = s.col_map[j];
194 while (tp < te && st.col_map[tc.col_index[tp]] < jg) { ++tp; }
195 T const aji = (tp < te && st.col_map[tc.col_index[tp]] == jg)
196 ? T(tc.mat[tp]) : T(0);
197 f(int(nlocal_f + j), int(nlocal_c + j), T(sc.mat[idx]), aji,
198 state_r[j], dbeta_r[j]);
199 }
200 }
201 }
202};
203
204}
205
206
211
217
221
262template <typename T>
263class AlgMG
264{
265public:
268 using TS = std::conditional_t<std::is_same_v<T,double>, float, T>;
269
274
275 AlgMG () = default;
276 explicit AlgMG (SpMatrix<T>& a_mat);
277
279 void define (SpMatrix<T>& a_mat);
280
296
298 void solve (AlgVector<T>& a_sol, AlgVector<T> const& a_rhs);
299
300 void setVerbose (int v) { m_verbose = v; }
302 void setPrintIndentation (std::string s) { m_print_ident = std::move(s); }
303 void setMaxIter (int n) { m_maxiter = n; }
305 void setFixedIter (int n) { m_fixediter = n; }
306 void setRelTol (T t) { m_reltol = t; }
308 void setAbsTol (T t) { m_abstol = t; }
310 void setSmoother (Smoother sm) { if (m_smoother_type != sm) { m_smoother_type = sm; m_smoothers_stale = true; } }
312 void setRelaxWeight (T w) { if (m_relax_weight != w) { m_relax_weight = w; m_smoothers_stale = true; } }
314 void setChebyshevDegree (int d) { if (m_cheby_degree != d) { m_cheby_degree = d; m_smoothers_stale = true; } }
316 void setChebyshevRatio (T r) {
317 AMREX_ALWAYS_ASSERT_WITH_MESSAGE(r > T(1), "AlgMG::setChebyshevRatio: the ratio must exceed 1");
318 if (m_cheby_ratio != r) { m_cheby_ratio = r; m_smoothers_stale = true; }
319 }
320
321 void setPreSmooth (int nu) { m_nu1 = nu; }
323 void setPostSmooth (int nu) { m_nu2 = nu; }
326 void setBottomSmooth (int nu) { m_nu_bottom = nu; }
331 void setStrongThreshold (T theta) { if (m_strong_threshold != theta) { m_strong_threshold = theta; m_prepared = false; } }
334 void setMaxLevels (int n) { if (m_max_levels != n) { m_max_levels = n; m_prepared = false; } }
336 void setMaxCoarseSize (Long n) { if (m_max_coarse_size != n) { m_max_coarse_size = n; m_prepared = false; } }
338 void setAggressiveNumLevels (int n) { if (m_aggressive_num_levels != n) { m_aggressive_num_levels = n; m_prepared = false; } }
341 void setAggressiveDirectInterp (bool b) { if (m_aggressive_direct_interp != b) { m_aggressive_direct_interp = b; m_prepared = false; } }
346 void setSingular (bool b) {
347 if (m_singular != b) { m_singular = b; m_prepared = false; }
348 }
349
350 void setInterpType (InterpType it) { if (m_interp_type != it) { m_interp_type = it; m_prepared = false; } }
353 void setPMaxElmts (int n) {
354 AMREX_ALWAYS_ASSERT_WITH_MESSAGE(n <= max_p_elmts, "AlgMG::setPMaxElmts: at most 32");
355 int const p = (n < 0) ? default_p_max_elmts : n;
356 if (m_p_max_elmts != p) { m_p_max_elmts = p; m_prepared = false; }
357 }
358 static constexpr int max_p_elmts = 32;
359 static constexpr int default_p_max_elmts = 4;
361 void setTruncFactor (T f) { if (m_trunc_factor != f) { m_trunc_factor = f; m_prepared = false; } }
364 if (m_krylov == a) { return; }
365 m_krylov = a;
366 m_krylov_bicgstab.reset();
367 m_krylov_gmres.reset();
368 m_krylov_pcg.reset();
369 }
370
377 if (m_bottom_solver == bs) { return; }
378 m_bottom_solver = bs;
379 m_gmres.reset();
380 m_bicgstab.reset();
381 m_bottom_n = -1;
382 }
383 void setBottomTol (T t) { m_bottom_reltol = t; }
384 void setBottomMaxIter (int n) { m_bottom_maxiter = n; }
385 void setBottomVerbose (int v) { m_bottom_verbose = v; }
386 void setThrowException (bool b) { m_throw_exception = b; }
387
388 [[nodiscard]] int getNumIters () const { return m_niters; }
390 [[nodiscard]] int getNumSetups () const { return m_num_setups; }
391 [[nodiscard]] T getResidualNorm () const { return m_rnorm; }
393 [[nodiscard]] int numLevels () const { return int(m_a.size()); }
395 [[nodiscard]] std::size_t getSetupPeakBytes () const { return m_setup_peak_bytes; }
397 [[nodiscard]] Long numGlobalRows (int lev) const { return m_a[lev]->numGlobalRows(); }
398 [[nodiscard]] SpMatrix<T> const& getMatrix (int lev) const { return *m_a[lev]; }
399 [[nodiscard]] SpMatrix<T> const& getInterp (int lev) const { return m_P[lev]; }
400
401private:
402
403 enum PointType : int { undecided = 0, fine = 1, coarse = 2 };
404
405 bool m_prepared = false;
406 int m_num_setups = 0;
407 bool m_smoothers_stale = false;
408 bool m_throw_exception = false;
409 int m_verbose = 0;
410 int m_maxiter = 200;
411 int m_fixediter = 0;
412 int m_niters = 0;
413
414 Smoother m_smoother_type = Smoother::chebyshev;
415 T m_relax_weight = T(-1);
416 int m_cheby_degree = 2;
417 T m_cheby_ratio = T(6);
418 int m_nu1 = -1;
419 int m_nu2 = -1;
420 int m_nu_bottom = -1;
421
422 [[nodiscard]] int numPreSmooth () const {
423 return (m_nu1 >= 0) ? m_nu1 : ((m_smoother_type == Smoother::chebyshev) ? 1 : 2);
424 }
425 [[nodiscard]] int numPostSmooth () const {
426 return (m_nu2 >= 0) ? m_nu2 : ((m_smoother_type == Smoother::chebyshev) ? 1 : 2);
427 }
428 [[nodiscard]] int numBottomSmooth () const {
429 return (m_nu_bottom >= 0) ? m_nu_bottom : ((m_smoother_type == Smoother::chebyshev) ? 4 : 8);
430 }
431
432 int m_max_levels = 25;
433 Long m_max_coarse_size = 9;
434 int m_aggressive_num_levels = 0;
435 bool m_aggressive_direct_interp = false;
436 bool m_singular = false;
437 std::unique_ptr<AlgVector<T>> m_bproj; // rhs with the mean removed
438 T m_strong_threshold = T(0.25);
439
440 InterpType m_interp_type = InterpType::mm_ext_i;
441 int m_p_max_elmts = default_p_max_elmts;
442 T m_trunc_factor = T(0);
443 BottomSolver m_bottom_solver = BottomSolver::direct;
444 // Safety limit for the dense LU on every process; the coarsest level
445 // normally has a few rows (max_coarse_size). Beyond it smoother sweeps
446 // are used, which keeps the cycle linear for PCG and GMRES.
447 static constexpr Long max_direct_size = 1024;
448 Long m_bottom_n = -1; // rows of the factored coarsest matrix; -1: not built, 0: not applicable
449 Vector<double> m_bottom_lu; // dense LU, row-major, pivoted
450 Vector<int> m_bottom_piv;
451 T m_bottom_reltol = T(1.e-4);
452 int m_bottom_maxiter = 200;
453 int m_bottom_verbose = 0;
454
455 T m_reltol = std::is_same_v<T,float> ? T(1.e-4) : T(1.e-8);
456 T m_abstol = T(0);
457 T m_init_rnorm = std::numeric_limits<T>::max();
458 T m_rnorm = std::numeric_limits<T>::max();
459 T m_bnorm = std::numeric_limits<T>::max();
460
461 std::string m_print_ident;
462 std::size_t m_setup_peak_bytes = 0;
463
465 void note_memory (int lev, char const* stage)
466 {
467 auto const [bytes, reserved] = detail::arena_bytes();
468 m_setup_peak_bytes = std::max(m_setup_peak_bytes, bytes);
469 if (m_verbose >= 2 && bytes > 0) {
470 amrex::Print() << m_print_ident << "AlgMG: level " << lev << ": after " << stage
471 << ": arena in use " << double(bytes)/double(1<<20)
472 << " MB, reserved " << double(reserved)/double(1<<20) << " MB"
473#ifdef AMREX_USE_GPU
474 << ", device free " << double(Gpu::Device::freeMemAvailable())/double(1<<20) << " MB"
475#endif
476 << "\n";
477 }
478 }
479
480 SpMatrix<T>* m_mat = nullptr;
481 Vector<SpMatrix<T>*> m_a; // matrix on each level
482 Vector<std::unique_ptr<SpMatrix<T>>> m_crsemat; // owns levels 1..n-1
483 Vector<SpMatrix<T>> m_P; // m_P[lev] interpolates lev+1 -> lev
484 Vector<SpMatrix<T>> m_R; // m_R[lev] restricts lev -> lev+1
485
486 Vector<AlgVector<T>> m_res; // rhs - A(sol)
487 Vector<AlgVector<T>> m_cor; // cor = A^-1 (res)
488 Vector<AlgVector<T>> m_rescor; // res - A(cor)
489
490 Vector<std::unique_ptr<JacobiSmoother<T>>> m_smoother;
491 Vector<std::unique_ptr<ChebyshevSmoother<T>>> m_cheby;
492 Vector<std::unique_ptr<L1GaussSeidelSmoother<T>>> m_l1gs;
493 std::unique_ptr<GMRES_MV<T>> m_gmres;
494 std::unique_ptr<BiCGStab_MV<T>> m_bicgstab;
495
496 KrylovSolver m_krylov = KrylovSolver::none;
497 std::unique_ptr<BiCGStab_MV<T>> m_krylov_bicgstab;
498 std::unique_ptr<GMRES_MV<T>> m_krylov_gmres;
499 std::unique_ptr<PCG_MV<T>> m_krylov_pcg;
500
501 void clear_levels ();
502 void build_smoothers ();
504 void build_direct_bottom ();
505 void direct_bottom_solve ();
506 // In-place all-gather of a row-distributed dense array (stride doubles
507 // per row): each rank's rows already sit at their global offset in buf.
508 static void allgather_rows (double* buf, AlgPartition const& part, Long stride);
510 void prepare ()
511 {
512 if (!m_prepared) { setup(); }
513 else if (m_smoothers_stale) { build_smoothers(); }
514 }
516 void vcycle (AlgVector<T>& x, AlgVector<T> const& b);
517 void blew_up () const;
518 void solve_krylov (AlgVector<T>& a_sol, AlgVector<T> const& a_rhs, T res_target);
519 void smooth (int lev, AlgVector<T>& xvec, AlgVector<T> const& bvec, int nu,
520 bool with_initial_guess, bool backward = false);
521 void bottom_solve ();
522
523 static std::pair<T,T> compute_rbnorms (SpMatrix<T> const& mat,
524 AlgVector<T> const& xvec,
525 AlgVector<T> const& bvec,
526 AlgVector<T>& rvec);
527
528 static T compute_rnorm (SpMatrix<T> const& mat, AlgVector<T> const& xvec,
529 AlgVector<T> const& bvec, AlgVector<T>& rvec);
530
531//Public for CUDA
532public: // NOLINT
533
534 void setup ();
535
541 bool coarsen_level (int lev, SpMatrix<T>& A, SpMatrix<T>& P, AlgPartition& cpart,
542 bool& aggressive, Long& ncoarse, bool& truncated,
543 Long& nnz_untruncated);
544
547 [[nodiscard]] SpMatrix<TS> create_soc (SpMatrix<T> const& A, T scale) const;
548
550 static T strength_scale (SpMatrix<T> const& A);
551
555 bool isolated_as_coarse = false) const;
556
560 int which = PointType::coarse);
561
565 {
566 if (!m_singular) { return; }
567 T const mean = x.sum() / T(x.partition().numGlobalRows());
568 auto* AMREX_RESTRICT px = x.data();
569 ParallelForOMP(x.numLocalRows(), [=] AMREX_GPU_DEVICE (Long i) noexcept { px[i] -= mean; });
571 }
572
575 Gpu::DeviceVector<int> const& state,
576 Gpu::DeviceVector<Long> const& cidx,
577 AlgPartition const& cpart);
578
583 int max_elmts, T trunc_factor);
585 static void normalize_rows (SpMatrix<T>& P);
586
590 Gpu::DeviceVector<int> const& state,
591 Gpu::DeviceVector<Long> const& cidx,
592 AlgPartition const& cpart);
593
596 Gpu::DeviceVector<int> const& state,
597 Gpu::DeviceVector<Long> const& cidx,
598 AlgPartition const& cpart);
599
602 Gpu::DeviceVector<int> const& state,
603 Gpu::DeviceVector<Long> const& cidx,
604 AlgPartition const& cpart, T scale) const;
605
613 SpMatrix<TS> const& ST,
614 Gpu::DeviceVector<int> const& state,
615 Gpu::DeviceVector<Long> const& cidx,
616 AlgPartition const& cpart, T scale,
617 bool plus_i, int const* row_state = nullptr,
618 int max_elmts = 0, T trunc_factor = T(0),
619 Long* nnz_untruncated = nullptr);
620
621#ifndef AMREX_USE_GPU
625 static SpMatrix<T> build_interp_mm_host (SpMatrix<T>& A, SpMatrix<TS>& S,
626 SpMatrix<TS> const& ST,
627 Gpu::DeviceVector<int> const& state,
628 Gpu::DeviceVector<Long> const& cidx,
629 AlgPartition const& cpart, T scale, bool plus_i,
630 int max_elmts, T trunc_factor,
631 Long& nnz_untruncated);
632#endif
633};
634
635template <typename T>
637{
638 define(a_mat);
639}
640
641template <typename T>
643{
644 m_mat = &a_mat;
645 m_prepared = false;
646 clear_levels();
647 m_bproj.reset();
648 m_krylov_bicgstab.reset();
649 m_krylov_gmres.reset();
650 m_krylov_pcg.reset();
651}
652
653template <typename T>
655{
656 if (!m_prepared || m_a.empty()) {
657 define(a_mat);
658 return false;
659 }
660 AMREX_ALWAYS_ASSERT_WITH_MESSAGE(a_mat.partition() == m_mat->partition(),
661 "AlgMG::updateMatrix: the matrix must have the same row partition");
662 m_mat = &a_mat;
663 m_mat->setColumnPartition(m_mat->partition()); // as in setup
664 m_a[0] = m_mat;
665 // The Krylov solvers hold the matrix
666 m_krylov_bicgstab.reset();
667 m_krylov_gmres.reset();
668 m_krylov_pcg.reset();
669 if (numLevels() == 1) { // the finest level is the bottom
670 m_bottom_n = -1;
671 m_bottom_lu.clear();
672 m_bottom_piv.clear();
673 }
674 m_smoothers_stale = true;
675 return true;
676}
677
678template <typename T>
679void AlgMG<T>::clear_levels ()
680{
681 m_a.clear();
682 m_crsemat.clear();
683 m_P.clear();
684 m_R.clear();
685 m_res.clear();
686 m_cor.clear();
687 m_rescor.clear();
688 m_smoother.clear();
689 m_cheby.clear();
690 m_l1gs.clear();
691 m_gmres.reset();
692 m_bicgstab.reset();
693 m_bottom_n = -1;
694 m_bottom_lu.clear();
695 m_bottom_piv.clear();
696}
697
698template <typename T>
699void AlgMG<T>::solve (AlgVector<T>& a_sol, AlgVector<T> const& a_rhs)
700{
701 BL_PROFILE("AlgMG::solve");
702
703 prepare();
704
705 // Singular: iterate on the rhs with its mean removed.
706 AlgVector<T> const* prhs = &a_rhs;
707 if (m_singular) {
708 if (!m_bproj) { m_bproj = std::make_unique<AlgVector<T>>(m_mat->partition()); }
709 m_bproj->copy(a_rhs);
710 project_out(*m_bproj);
711 prhs = m_bproj.get();
712 }
713 AlgVector<T> const& rhs = *prhs;
714
715 auto norms = compute_rbnorms(*m_mat, a_sol, rhs, m_res[0]);
716 m_bnorm = norms.first;
717 m_init_rnorm = m_rnorm = norms.second;
718 m_niters = 0;
719
720 if (m_bnorm == 0) {
721 a_sol.setVal(0); // also the zero-mean solution of a singular system
722 m_rnorm = 0;
723 return;
724 }
725
726 if (m_verbose >= 1) {
727 amrex::Print() << m_print_ident << "AlgMG: Initial rhs 2-norm = " << m_bnorm << "\n"
728 << m_print_ident << "AlgMG: Initial residual 2-norm = " << m_rnorm << "\n";
729 }
730
731 T res_target = std::max(m_bnorm * m_reltol, m_abstol);
732 if (m_fixediter <= 0 && m_rnorm <= res_target) {
733 if (m_verbose >= 1) {
734 amrex::Print() << m_print_ident << "AlgMG: No iterations needed\n";
735 }
736 project_out(a_sol);
737 return;
738 }
739
740 if (m_krylov != KrylovSolver::none) {
741 solve_krylov(a_sol, rhs, res_target);
742 project_out(a_sol);
743 return;
744 }
745
746 int const niters = (m_fixediter > 0) ? m_fixediter : m_maxiter;
747 for (int iter = 0; iter < niters; ++iter)
748 {
749 vcycle(m_cor[0], m_res[0]);
750 a_sol.plusAsync(m_cor[0]);
751
752 m_rnorm = compute_rnorm(*m_mat, a_sol, rhs, m_res[0]);
753 m_niters = iter+1;
754
755 if (m_verbose >= 2) {
756 amrex::Print() << m_print_ident << "AlgMG: Iteration " << std::setw(3)
757 << iter+1 << ": rnorm, rnorm/bnorm = "
758 << m_rnorm << ", " << m_rnorm/m_bnorm << "\n";
759 }
760
761 bool converged = (m_fixediter <= 0) && (m_rnorm <= res_target);
762 if (converged) {
763 if (m_verbose >= 1) {
764 amrex::Print() << m_print_ident << "AlgMG: Final Iter. " << iter+1
765 << ": rnorm, rnorm/bnorm = "
766 << m_rnorm << ", " << m_rnorm/m_bnorm << "\n";
767 }
768 break;
769 } else if (m_rnorm > T(1.e20) * m_init_rnorm || amrex::isnan(m_rnorm)) {
770 if (m_verbose > 0) {
771 amrex::Print() << m_print_ident << "AlgMG: Failing to converge after " << iter+1 << " iterations."
772 << " rnorm, rnorm/bnorm = "
773 << m_rnorm << ", " << m_rnorm/m_bnorm << "\n";
774 }
775 blew_up();
776 }
777 }
778
779 if (m_verbose >= 1 && m_fixediter <= 0 && m_rnorm > res_target) {
780 amrex::Print() << m_print_ident << "AlgMG: Not converged after " << m_niters
781 << " iterations. rnorm/bnorm = " << m_rnorm/m_bnorm << "\n";
782 }
783 project_out(a_sol);
784}
785
786template <typename T>
788{
789 prepare();
790 vcycle(x, b);
791 project_out(x);
792}
793
794template <typename T>
795void AlgMG<T>::blew_up () const
796{
797 if (m_throw_exception) {
798 throw std::runtime_error("AlgMG blew up.");
799 } else {
800 amrex::Abort("AlgMG failing so lets stop here");
801 }
802}
803
804template <typename T>
805void AlgMG<T>::solve_krylov (AlgVector<T>& a_sol, AlgVector<T> const& a_rhs, T res_target)
806{
807 BL_PROFILE("AlgMG::solve_krylov");
808
809 // Bound on every solve so that a moved AlgMG object stays usable.
810 auto pc = [this] (AlgVector<T>& x, AlgVector<T> const& b) { this->precond(x, b); };
811 int const niters = (m_fixediter > 0) ? m_fixediter : m_maxiter;
812 T const atol = (m_fixediter > 0) ? T(0) : res_target;
813 int const verbose = std::max(0, m_verbose-1);
814 T krylov_rnorm = T(0);
815
816 // Converge on the absolute target reltol*|b|, as the V-cycle iteration.
817 // Restart if the recurrence residual converged before the true one.
818 auto run = [&] (auto& s) {
819 s.setVerbose(verbose);
820 s.setInitialGuessNonzero(true);
821 m_niters = 0;
822 while (true) {
823 s.setMaxIters(niters - m_niters);
824 s.solve(a_sol, a_rhs, T(0), atol);
825 m_niters += s.getNumIters();
826 krylov_rnorm = s.getResidualNorm();
827 m_rnorm = compute_rnorm(*m_mat, a_sol, a_rhs, m_res[0]);
828 if (m_fixediter > 0 || m_rnorm <= res_target || krylov_rnorm > res_target ||
829 m_niters >= niters || s.getNumIters() == 0 || amrex::isnan(m_rnorm)) {
830 break;
831 }
832 if (m_verbose >= 2) {
833 amrex::Print() << m_print_ident << "AlgMG: restarting Krylov solver at iteration "
834 << m_niters << ", rnorm/bnorm = " << m_rnorm/m_bnorm << "\n";
835 }
836 }
837 };
838
839 char const* name = "";
840 if (m_krylov == KrylovSolver::bicgstab) {
841 name = "BiCGStab";
842 if (!m_krylov_bicgstab) {
843 m_krylov_bicgstab = std::make_unique<BiCGStab_MV<T>>(m_mat);
844 }
845 m_krylov_bicgstab->setPrecond(pc);
846 run(m_krylov_bicgstab->getSolver());
847 } else if (m_krylov == KrylovSolver::gmres) {
848 name = "GMRES";
849 if (!m_krylov_gmres) {
850 m_krylov_gmres = std::make_unique<GMRES_MV<T>>(m_mat);
851 }
852 m_krylov_gmres->setPrecond(pc);
853 run(m_krylov_gmres->getSolver());
854 } else {
855 name = "PCG";
856 if (!m_krylov_pcg) {
857 m_krylov_pcg = std::make_unique<PCG_MV<T>>(m_mat);
858 }
859 m_krylov_pcg->setPrecond(pc);
860 run(m_krylov_pcg->getSolver());
861 }
862
863 if (amrex::isnan(m_rnorm) || m_rnorm > T(1.e20) * m_init_rnorm) { blew_up(); }
864
865 if (m_verbose >= 1) {
866 amrex::Print() << m_print_ident << "AlgMG(" << name << "): Final Iter. " << m_niters
867 << ": rnorm, rnorm/bnorm = " << m_rnorm << ", " << m_rnorm/m_bnorm;
868 if (m_verbose >= 2) {
869 amrex::Print() << " (Krylov recurrence " << krylov_rnorm << ")";
870 }
871 amrex::Print() << "\n";
872 if (m_fixediter <= 0 && m_rnorm > res_target) {
873 amrex::Print() << m_print_ident << "AlgMG(" << name << "): Not converged after "
874 << m_niters << " iterations.\n";
875 }
876 }
877}
878
879template <typename T>
880std::pair<T,T> AlgMG<T>::compute_rbnorms (SpMatrix<T> const& mat,
881 AlgVector<T> const& xvec,
882 AlgVector<T> const& bvec,
883 AlgVector<T>& rvec)
884{
885 auto bnorm = bvec.norm2(true);
886
887 amrex::computeResidual(rvec, mat, xvec, bvec);
888 auto rnorm = rvec.norm2(true);
889
890 if (ParallelContext::NProcsSub() > 1) {
892 rnorm = Math::powi<2>(rnorm);
894 bnorm = std::sqrt(bnorm);
895 rnorm = std::sqrt(rnorm);
896 }
897
898 return {bnorm, rnorm};
899}
900
901template <typename T>
902T AlgMG<T>::compute_rnorm (SpMatrix<T> const& mat, AlgVector<T> const& xvec,
903 AlgVector<T> const& bvec, AlgVector<T>& rvec)
904{
905 amrex::computeResidual(rvec, mat, xvec, bvec);
906 return rvec.norm2();
907}
908
909template <typename T>
910void AlgMG<T>::vcycle (AlgVector<T>& x, AlgVector<T> const& b)
911{
912 BL_PROFILE("AlgMG::vcycle");
913
914 int const nlev = numLevels();
915 int const bottom = nlev-1;
916
917 if (nlev == 1) { // bottom_solve works on the level vectors
918 if (&b != &m_res[0]) { m_res[0].copyAsync(b); }
919 bottom_solve();
920 if (&x != &m_cor[0]) { x.copyAsync(m_cor[0]); }
921 return;
922 }
923
924 // Level 0 uses the caller's vectors, the coarse levels their own.
925 auto cor = [&] (int lev) -> AlgVector<T>& { return (lev == 0) ? x : m_cor[lev]; };
926 auto res = [&] (int lev) -> AlgVector<T> const& { return (lev == 0) ? b : m_res[lev]; };
927
928 smooth(0, x, b, numPreSmooth(), false);
929
930 for (int lev = 0; lev < bottom; ++lev) {
931 computeResidual(m_rescor[lev], *m_a[lev], cor(lev), res(lev));
932 SpMV(m_res[lev+1], m_R[lev], m_rescor[lev]);
933 if (lev+1 == bottom) {
934 bottom_solve();
935 } else {
936 smooth(lev+1, m_cor[lev+1], m_res[lev+1], numPreSmooth(), false);
937 }
938 }
939
940 for (int lev = bottom-1; lev >= 0; --lev) {
941 SpMV(m_rescor[lev], m_P[lev], m_cor[lev+1]);
942 cor(lev).plusAsync(m_rescor[lev]);
943 smooth(lev, cor(lev), res(lev), numPostSmooth(), true, true);
944 }
945}
946
947template <typename T>
948void AlgMG<T>::smooth (int lev, AlgVector<T>& xvec, AlgVector<T> const& bvec, int nu,
949 bool with_initial_guess, bool backward)
950{
951 if (m_smoother_type == Smoother::l1_gauss_seidel) {
952 auto& smoother = *m_l1gs[lev];
953 smoother.setNumIters(nu);
954 smoother(xvec, bvec, with_initial_guess, backward);
955 } else if (m_smoother_type == Smoother::chebyshev) {
956 auto& smoother = *m_cheby[lev];
957 smoother.setNumIters(nu);
958 smoother(xvec, bvec, with_initial_guess);
959 } else {
960 auto& smoother = *m_smoother[lev];
961 smoother.setNumIters(nu);
962 smoother(xvec, bvec, with_initial_guess);
963 }
964}
965
966template <typename T>
967void AlgMG<T>::bottom_solve ()
968{
969 BL_PROFILE("AlgMG::bottom_solve");
970
971 int const lev = numLevels()-1;
972 project_out(m_res[lev]); // keep the coarsest solve consistent
973 auto smooth_bottom = [&] () {
974 if (m_smoother_type == Smoother::l1_gauss_seidel) {
975 // Forward then backward sweeps keep the cycle symmetric.
976 int const nu = (numBottomSmooth()+1)/2;
977 smooth(lev, m_cor[lev], m_res[lev], nu, false);
978 smooth(lev, m_cor[lev], m_res[lev], nu, true, true);
979 } else {
980 smooth(lev, m_cor[lev], m_res[lev], numBottomSmooth(), false);
981 }
982 };
983 // As in MLMG, a Krylov solve that does not reduce the residual is
984 // discarded; its iterate can be far worse than no correction. The
985 // solver starts from zero unless told otherwise.
986 auto krylov_bottom = [&] (auto& krylov) {
987 krylov.setVerbose(m_bottom_verbose);
988 krylov.getSolver().setMaxIters(m_bottom_maxiter);
989 krylov.solve(m_cor[lev], m_res[lev], m_bottom_reltol, T(0));
990 auto const& solver = krylov.getSolver();
991 T const rnorm0 = solver.getInitialResidualNorm();
992 if (rnorm0 > T(0) && !(solver.getResidualNorm() < rnorm0)) {
993 if (m_verbose >= 2) {
994 amrex::Print() << m_print_ident << "AlgMG: bottom solve failed, smoothing instead\n";
995 }
996 smooth_bottom();
997 }
998 };
999 auto bottom_solver = m_bottom_solver;
1000 if (bottom_solver == BottomSolver::direct) {
1001 if (m_bottom_n < 0) { build_direct_bottom(); }
1002 if (m_bottom_n == 0) { bottom_solver = BottomSolver::jacobi; }
1003 }
1004 switch (bottom_solver) {
1005 case BottomSolver::direct:
1006 direct_bottom_solve();
1007 break;
1008 case BottomSolver::jacobi:
1009 smooth_bottom();
1010 break;
1011 case BottomSolver::gmres:
1012 if (!m_gmres) {
1013 m_gmres = std::make_unique<GMRES_MV<T>>(m_a[lev]);
1014 m_gmres->setPrecond(JacobiSmoother<T>(m_a[lev], m_smoother_type != Smoother::jacobi));
1015 }
1016 krylov_bottom(*m_gmres);
1017 break;
1019 if (!m_bicgstab) {
1020 m_bicgstab = std::make_unique<BiCGStab_MV<T>>(m_a[lev]);
1021 m_bicgstab->setPrecond(JacobiSmoother<T>(m_a[lev], m_smoother_type != Smoother::jacobi));
1022 }
1023 krylov_bottom(*m_bicgstab);
1024 break;
1025 }
1026 project_out(m_cor[lev]);
1027}
1028
1029template <typename T>
1031{
1032 BL_PROFILE("AlgMG::setup");
1033
1034 ++m_num_setups;
1035 clear_levels();
1036 m_a.push_back(m_mat);
1037 m_mat->setColumnPartition(m_mat->partition()); // square: local column indexing
1038 m_setup_peak_bytes = 0;
1039 note_memory(0, "start");
1040
1041 // The coarsening follows De Sterck, Yang & Heys [2].
1042
1043 for (int lev = 0; lev+1 < m_max_levels; ++lev)
1044 {
1045 auto& A = *m_a[lev];
1046 if (A.numGlobalRows() <= m_max_coarse_size) { break; }
1047
1048 SpMatrix<T> P;
1049 AlgPartition cpart;
1050 bool aggressive = false;
1051 Long ncoarse = 0;
1052 bool truncated = false;
1053 Long nnz0 = 0;
1054 if (!coarsen_level(lev, A, P, cpart, aggressive, ncoarse, truncated, nnz0)) { break; }
1055 note_memory(lev, "strength release");
1056 if (m_p_max_elmts > 0 || m_trunc_factor > T(0)) {
1057 if (!truncated) {
1058 nnz0 = P.numLocalNonZeros();
1059 P = truncate_interp(P, cpart, m_p_max_elmts, m_trunc_factor);
1060 note_memory(lev, "truncation");
1061 }
1062 if (m_verbose >= 2) {
1063 Long nnz[2] = {nnz0, P.numLocalNonZeros()};
1065 amrex::Print() << m_print_ident << "AlgMG: level " << lev
1066 << ": P truncated from " << nnz[0] << " to "
1067 << nnz[1] << " nonzeros\n";
1068 }
1069 }
1070 // P 1 = 1 only to single precision, since the weights come from the
1071 // strength matrix; singular problems need it exactly.
1072 if (m_singular) { normalize_rows(P); }
1073 if (m_verbose >= 2 && aggressive) {
1074 amrex::Print() << m_print_ident << "AlgMG: level " << lev
1075 << ": aggressive coarsening, " << ncoarse << " -> "
1076 << cpart.numGlobalRows() << " C-points\n";
1077 }
1078 if (m_verbose >= 3) {
1079 // P should reproduce constants: row sums close to one.
1080 AlgVector<T> ones(cpart), rs(A.partition());
1081 ones.setVal(T(1));
1082 SpMV(rs, P, ones);
1083 Long nzero = Reduce::Sum<Long>(rs.numLocalRows(),
1084 [q = rs.data()] AMREX_GPU_DEVICE (Long i) -> Long { return q[i] == T(0); });
1085 T rmin = Reduce::Min<T>(rs.numLocalRows(),
1086 [q = rs.data()] AMREX_GPU_DEVICE (Long i) -> T { return q[i]; });
1087 T rmax = Reduce::Max<T>(rs.numLocalRows(),
1088 [q = rs.data()] AMREX_GPU_DEVICE (Long i) -> T { return q[i]; });
1092 amrex::Print() << m_print_ident << "AlgMG: level " << lev << ": P row sums in ["
1093 << rmin << ", " << rmax << "], " << nzero << " empty rows\n";
1094 }
1095 auto R = amrex::transpose(P, cpart);
1096 note_memory(lev, "R");
1097 SpMatrix<T> Ac;
1098#ifndef AMREX_USE_GPU
1099 Ac = amrex::RAP(R, A, P, cpart);
1100#else
1101 {
1102 // R (A P): measured smaller and faster than (R A) P, whose
1103 // intermediate gathers several fine rows per coarse row.
1104 auto AP = amrex::SpGEMM(A, P, cpart);
1105 note_memory(lev, "A*P");
1106 Ac = amrex::SpGEMM(R, AP, cpart);
1107 }
1108#endif
1109 note_memory(lev, "coarse operator");
1110
1111 m_P.push_back(std::move(P));
1112 m_R.push_back(std::move(R));
1113 m_crsemat.push_back(std::make_unique<SpMatrix<T>>(std::move(Ac)));
1114 m_a.push_back(m_crsemat.back().get());
1115 }
1116
1117 int const nlev = numLevels();
1118 for (int lev = 0; lev < nlev; ++lev) {
1119 m_res.emplace_back(m_a[lev]->partition());
1120 m_cor.emplace_back(m_a[lev]->partition());
1121 if (lev < nlev-1) {
1122 m_rescor.emplace_back(m_a[lev]->partition());
1123 }
1124 }
1125 build_smoothers();
1126
1127 if (m_verbose >= 1) {
1128 Long nnz0 = 0, nnz_total = 0;
1129 for (int lev = 0; lev < nlev; ++lev) {
1130 Long nnz = m_a[lev]->numLocalNonZeros();
1132 if (lev == 0) { nnz0 = nnz; }
1133 nnz_total += nnz;
1134 amrex::Print() << m_print_ident << "AlgMG: level " << lev << ": "
1135 << m_a[lev]->numGlobalRows() << " rows, "
1136 << nnz << " nonzeros\n";
1137 }
1138 amrex::Print() << m_print_ident << "AlgMG: operator complexity = "
1139 << double(nnz_total)/double(std::max(nnz0,Long(1))) << "\n";
1140 if (m_setup_peak_bytes > 0) {
1141 auto peak = m_setup_peak_bytes;
1143 amrex::Print() << m_print_ident << "AlgMG: setup peak arena use = "
1144 << double(peak)/double(1<<20) << " MB ("
1145 << double(peak)/double(m_mat->numGlobalRows())
1146 << " B per row)\n";
1147 }
1148 }
1149
1150 m_prepared = true;
1151}
1152
1153template <typename T>
1154void AlgMG<T>::build_direct_bottom ()
1155{
1156 BL_PROFILE("AlgMG::build_direct_bottom");
1157 int const lev = numLevels()-1;
1158 auto const& A = *m_a[lev];
1159 Long const n = A.numGlobalRows();
1160 m_bottom_n = 0;
1161 m_bottom_lu.clear();
1162 m_bottom_piv.clear();
1163 if (n < 1 || n > max_direct_size) {
1164 if (m_verbose >= 1 && n > max_direct_size) {
1165 amrex::Print() << m_print_ident << "AlgMG: coarsest level has " << n
1166 << " rows, using smoother sweeps instead of the direct bottom solver\n";
1167 }
1168 return;
1169 }
1170
1171 // Dense copy of the local rows at their global offsets; gathered below.
1172 Vector<double> dense(n*n, 0.0);
1173 auto const pc = A.const_parcsr();
1174 auto to_host = [] (auto const* dptr, Long cnt) {
1175 Vector<std::remove_const_t<std::remove_pointer_t<decltype(dptr)>>> h(cnt);
1176 if (cnt > 0) { Gpu::copy(Gpu::deviceToHost, dptr, dptr+cnt, h.begin()); }
1177 return h;
1178 };
1179 {
1180 auto const& c = pc.csr0;
1181 auto ro = to_host(c.row_offset, (c.nrows > 0) ? c.nrows+1 : 0);
1182 auto ci = to_host(c.col_index, c.nnz);
1183 auto va = to_host(c.mat, c.nnz);
1184 for (Long i = 0; i < c.nrows; ++i) {
1185 Long const ig = pc.row_begin + i;
1186 for (int idx = ro[i]; idx < ro[i+1]; ++idx) {
1187 dense[ig*n + pc.col_begin + ci[idx]] += double(va[idx]);
1188 }
1189 }
1190 }
1191 if (pc.csr1.nrows > 0 && pc.csr1.nnz > 0) {
1192 auto const& c = pc.csr1;
1193 auto ro = to_host(c.row_offset, c.nrows+1);
1194 auto ci = to_host(c.col_index, c.nnz);
1195 auto va = to_host(c.mat, c.nnz);
1196 auto rmap = to_host(pc.row_map, pc.csr0.nrows); // local row -> csr1 row or -1
1197 int ncols = 0;
1198 for (auto j : ci) { ncols = std::max(ncols, j+1); }
1199 auto cmap = to_host(pc.col_map, ncols);
1200 for (Long i = 0; i < pc.csr0.nrows; ++i) {
1201 Long const r = rmap[i];
1202 if (r < 0) { continue; }
1203 Long const ig = pc.row_begin + i;
1204 for (int idx = ro[r]; idx < ro[r+1]; ++idx) {
1205 dense[ig*n + cmap[ci[idx]]] += double(va[idx]);
1206 }
1207 }
1208 }
1209 allgather_rows(dense.data(), A.partition(), n);
1210
1211 // Singular: the constant vector is in the null space. Adding (1/n) 1 1^T
1212 // makes the matrix regular; the coarsest residual and correction have
1213 // their mean removed, so the added term never acts on them.
1214 if (m_singular) {
1215 double dsum = 0.0;
1216 for (Long i = 0; i < n; ++i) { dsum += std::abs(dense[i*n+i]); }
1217 double const w = (dsum/double(n))/double(n);
1218 for (auto& v : dense) { v += w; }
1219 }
1220
1221 // LU with partial pivoting, in place (Doolittle). A pivot at roundoff
1222 // level relative to the matrix means a singular matrix.
1223 double amat = 0.0;
1224 for (auto v : dense) { amat = std::max(amat, std::abs(v)); }
1225 double const pivot_tol = double(n) * std::numeric_limits<double>::epsilon() * amat;
1226 Vector<int> piv(n);
1227 std::iota(piv.begin(), piv.end(), 0);
1228 for (Long i = 0; i < n; ++i) {
1229 Long imax = i;
1230 double amax = std::abs(dense[i*n+i]);
1231 for (Long k = i+1; k < n; ++k) {
1232 double const a = std::abs(dense[k*n+i]);
1233 if (a > amax) { amax = a; imax = k; }
1234 }
1235 if (amax <= pivot_tol) {
1236 if (m_verbose >= 1) {
1237 amrex::Print() << m_print_ident << "AlgMG: coarsest matrix is singular, "
1238 << "using smoother sweeps instead of the direct bottom solver\n";
1239 }
1240 return;
1241 }
1242 if (imax != i) {
1243 std::swap(piv[i], piv[imax]);
1244 for (Long j = 0; j < n; ++j) { std::swap(dense[i*n+j], dense[imax*n+j]); }
1245 }
1246 double const dinv = 1.0/dense[i*n+i];
1247 for (Long j = i+1; j < n; ++j) {
1248 double const f = dense[j*n+i]*dinv;
1249 dense[j*n+i] = f;
1250 for (Long k = i+1; k < n; ++k) { dense[j*n+k] -= f*dense[i*n+k]; }
1251 }
1252 }
1253 m_bottom_lu = std::move(dense);
1254 m_bottom_piv = std::move(piv);
1255 m_bottom_n = n;
1256}
1257
1258template <typename T>
1259void AlgMG<T>::allgather_rows (double* buf, AlgPartition const& part, Long stride)
1260{
1261#ifdef AMREX_USE_MPI
1262 int const nprocs = ParallelContext::NProcsSub();
1263 if (nprocs == 1) { return; }
1264 Vector<int> counts(nprocs), displs(nprocs);
1265 for (int p = 0; p < nprocs; ++p) {
1266 displs[p] = int(part[p]*stride);
1267 counts[p] = int((part[p+1]-part[p])*stride);
1268 }
1269 BL_MPI_REQUIRE(MPI_Allgatherv(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL,
1270 buf, counts.data(), displs.data(), MPI_DOUBLE,
1272#else
1273 amrex::ignore_unused(buf, part, stride);
1274#endif
1275}
1276
1277template <typename T>
1278void AlgMG<T>::direct_bottom_solve ()
1279{
1280 int const lev = numLevels()-1;
1281 auto const& r = m_res[lev];
1282 auto& x = m_cor[lev];
1283 Long const n = m_bottom_n;
1284 Long const nl = r.numLocalRows();
1285 Long const b0 = r.globalBegin();
1286
1287 // Gather the residual.
1288 PODVector<double> rhs(n); // uninitialized; every entry is written below
1289 Gpu::PinnedVector<T> loc(nl);
1290 if (nl > 0) {
1291 Gpu::copy(Gpu::deviceToHost, r.data(), r.data()+nl, loc.begin());
1292 for (Long i = 0; i < nl; ++i) { rhs[b0+i] = double(loc[i]); }
1293 }
1294 allgather_rows(rhs.data(), r.partition(), 1);
1295
1296 // P A = L U: forward substitution with the unit lower part, then back.
1297 Vector<double> y(n);
1298 auto const& lu = m_bottom_lu;
1299 for (Long i = 0; i < n; ++i) {
1300 double v = rhs[m_bottom_piv[i]];
1301 for (Long k = 0; k < i; ++k) { v -= lu[i*n+k]*y[k]; }
1302 y[i] = v;
1303 }
1304 for (Long i = n-1; i >= 0; --i) {
1305 double v = y[i];
1306 for (Long k = i+1; k < n; ++k) { v -= lu[i*n+k]*y[k]; }
1307 y[i] = v/lu[i*n+i];
1308 }
1309
1310 if (nl > 0) {
1311 for (Long i = 0; i < nl; ++i) { loc[i] = T(y[b0+i]); }
1312 Gpu::copy(Gpu::hostToDevice, loc.begin(), loc.end(), x.data());
1313 }
1314}
1315
1316template <typename T>
1317void AlgMG<T>::build_smoothers ()
1318{
1319 m_smoother.clear();
1320 m_cheby.clear();
1321 m_l1gs.clear();
1322 m_gmres.reset(); // their preconditioner depends on the smoother type
1323 m_bicgstab.reset();
1324 int const nlev = numLevels();
1325 for (int lev = 0; lev < nlev; ++lev) {
1326 if (m_smoother_type == Smoother::l1_gauss_seidel) {
1327#ifdef AMREX_USE_GPU
1328 amrex::Abort("AlgMG: L1GaussSeidel smoother is only available in CPU builds");
1329#endif
1330 m_l1gs.push_back(std::make_unique<L1GaussSeidelSmoother<T>>(m_a[lev]));
1331 } else if (m_smoother_type == Smoother::chebyshev) {
1332 m_cheby.push_back(std::make_unique<ChebyshevSmoother<T>>(m_a[lev], true));
1333 m_cheby.back()->setDegree(m_cheby_degree);
1334 m_cheby.back()->setEigRatio(m_cheby_ratio);
1335 } else {
1336 bool const l1 = (m_smoother_type == Smoother::l1_jacobi);
1337 m_smoother.push_back(std::make_unique<JacobiSmoother<T>>(m_a[lev], l1));
1338 T const w = (m_relax_weight >= T(0)) ? m_relax_weight : (l1 ? T(4./3.) : T(2./3.));
1339 m_smoother.back()->setWeight(w);
1340 }
1341 }
1342 m_smoothers_stale = false;
1343}
1344
1345template <typename T>
1347 bool& aggressive, Long& ncoarse, bool& truncated,
1348 Long& nnz_untruncated)
1349{
1350#ifdef AMREX_USE_GPU
1351 amrex::ignore_unused(truncated, nnz_untruncated); // set by the host paths only
1352#endif
1353 T const sscale = strength_scale(A);
1354 auto S = create_soc(A, sscale);
1355 S.setColumnPartition(S.partition());
1356 note_memory(lev, "S");
1357 auto ST = amrex::transpose(S, S.partition());
1358 ST.setColumnPartition(ST.partition());
1359
1360 note_memory(lev, "strength");
1362 // Singular: an isolated point becomes C, so every row of P sums to one.
1363 pmis(S, ST, state, m_singular);
1364 note_memory(lev, "PMIS");
1365
1367 cpart = coarse_numbering(state, cidx);
1368 ncoarse = cpart.numGlobalRows();
1369 if (ncoarse == 0 || ncoarse == A.numGlobalRows()) { return false; }
1370#ifndef AMREX_USE_GPU
1371 bool const host_interp = m_interp_type != InterpType::direct
1372 && detail::spmat_comm_is_local(A.partition(), cpart);
1373#endif
1374
1375 if (lev < m_aggressive_num_levels) {
1376 // Second PMIS pass on the C-points connected at distance two.
1377 auto S2 = second_pass_strength(S, state, cidx, cpart);
1378 S2.setColumnPartition(cpart);
1379 auto S2T = amrex::transpose(S2, cpart);
1380 S2T.setColumnPartition(cpart);
1381 note_memory(lev, "second-pass strength");
1383 pmis(S2, S2T, state2, true); // isolated C-points stay coarse
1385 auto cpart2 = coarse_numbering(state2, cidx2);
1386 Long const nc2 = cpart2.numGlobalRows();
1387 if (nc2 > 0 && nc2 < ncoarse) {
1388 // C/F split of the fine points with respect to C2.
1389 Long const nrows = A.numLocalRows();
1390 Gpu::DeviceVector<int> statef(nrows);
1391 Gpu::DeviceVector<Long> cidxf(nrows);
1392 auto const* p1 = state.data();
1393 auto const* c1 = cidx.data();
1394 auto const* p2 = state2.data();
1395 auto const* c2 = cidx2.data();
1396 auto* pf = statef.data();
1397 auto* cf = cidxf.data();
1398 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
1399 {
1400 bool const c = (p1[i] == PointType::coarse) &&
1401 (p2[c1[i]] == PointType::coarse);
1402 pf[i] = c ? PointType::coarse : PointType::fine;
1403 cf[i] = (p1[i] == PointType::coarse) ? c2[c1[i]] : Long(0);
1404 });
1406
1407 bool const trunc = (m_p_max_elmts > 0 || m_trunc_factor > T(0));
1408 SpMatrix<T> P1;
1409#ifndef AMREX_USE_GPU
1410 if (!m_aggressive_direct_interp && host_interp) {
1411 Long unused;
1412 P1 = build_interp_mm_host(A, S, ST, state, cidx, cpart, sscale,
1413 m_interp_type == InterpType::mm_ext_i,
1414 m_p_max_elmts, m_trunc_factor, unused);
1415 } else if (!m_aggressive_direct_interp && m_interp_type != InterpType::direct) {
1416 P1 = build_interp_mm(A, S, ST, state, cidx, cpart, sscale,
1417 m_interp_type == InterpType::mm_ext_i, nullptr,
1418 m_p_max_elmts, m_trunc_factor);
1419 } else
1420#endif
1421 {
1422 P1 = m_aggressive_direct_interp ? build_interp(A, S, state, cidx, cpart)
1423 : interpolation(A, S, ST, state, cidx, cpart, sscale);
1424 if (trunc) { P1 = truncate_interp(P1, cpart, m_p_max_elmts, m_trunc_factor); }
1425 }
1426 note_memory(lev, "P1");
1427 // P2 needs only the rows of the first-pass C-points.
1428#ifndef AMREX_USE_GPU
1429 auto Pf = build_interp_mm(A, S, ST, statef, cidxf, cpart2, sscale,
1430 m_interp_type != InterpType::mm_ext,
1431 state.data(), m_p_max_elmts, m_trunc_factor);
1432#else
1433 auto Pf = build_interp_mm(A, S, ST, statef, cidxf, cpart2, sscale,
1434 m_interp_type != InterpType::mm_ext,
1435 state.data());
1436 if (trunc) { Pf = truncate_interp(Pf, cpart2, m_p_max_elmts, m_trunc_factor); }
1437#endif
1438 auto P2 = select_coarse_rows(Pf, state, cidx, cpart);
1439 note_memory(lev, "P2");
1440 P = amrex::SpGEMM(P1, P2, cpart2);
1441 cpart = cpart2;
1442 aggressive = true;
1443 }
1444 }
1445 if (!aggressive) {
1446#ifndef AMREX_USE_GPU
1447 if (host_interp) {
1448 P = build_interp_mm_host(A, S, ST, state, cidx, cpart, sscale,
1449 m_interp_type == InterpType::mm_ext_i,
1450 m_p_max_elmts, m_trunc_factor, nnz_untruncated);
1451 truncated = true;
1452 } else if (m_interp_type != InterpType::direct) {
1453 P = build_interp_mm(A, S, ST, state, cidx, cpart, sscale,
1454 m_interp_type == InterpType::mm_ext_i, nullptr,
1455 m_p_max_elmts, m_trunc_factor, &nnz_untruncated);
1456 truncated = true;
1457 } else
1458#endif
1459 {
1460 P = interpolation(A, S, ST, state, cidx, cpart, sscale);
1461 }
1462 }
1463 note_memory(lev, "interpolation");
1464 return true;
1465}
1466
1467template <typename T>
1469{
1470 auto const rs = P.rowSum();
1471 auto const* AMREX_RESTRICT prs = rs.data();
1472 Long nzero = Reduce::Sum<Long>(rs.numLocalRows(),
1473 [=] AMREX_GPU_DEVICE (Long i) -> Long { return prs[i] == T(0); });
1475 if (nzero > 0) {
1476 amrex::Abort("AlgMG: singular mode needs P 1 = 1, but " + std::to_string(nzero)
1477 + " rows of P are empty");
1478 }
1479 auto a = P.parcsr();
1480 ParallelForOMP(rs.numLocalRows(), [=] AMREX_GPU_DEVICE (Long i) noexcept
1481 {
1482 T const f = T(1) / prs[i];
1483 for (Long idx = a.csr0.row_offset[i]; idx < a.csr0.row_offset[i+1]; ++idx) {
1484 a.csr0.mat[idx] *= f;
1485 }
1486 if (detail::has_remote_row(a, i)) {
1487 Long const ii = a.row_map[i];
1488 for (Long idx = a.csr1.row_offset[ii]; idx < a.csr1.row_offset[ii+1]; ++idx) {
1489 a.csr1.mat[idx] *= f;
1490 }
1491 }
1492 });
1494}
1495
1496template <typename T>
1498{
1499 auto const& Ad = A.diagonalVector();
1500 auto const* AMREX_RESTRICT p = Ad.data();
1501 T dmax = Reduce::Max<T>(Ad.numLocalRows(),
1502 [=] AMREX_GPU_DEVICE (Long i) -> T { return std::abs(p[i]); });
1504 return (dmax > T(0)) ? dmax : T(1);
1505}
1506
1507template <typename T>
1509{
1510 BL_PROFILE("AlgMG::create_soc");
1511
1512 auto const& a = A.const_parcsr();
1513 Long const nrows = a.csr0.nrows;
1514 int const nlocal = int(A.numLocalRows());
1515
1516 auto const& Ad = A.diagonalVector();
1517 auto const* AMREX_RESTRICT A_ii = Ad.data();
1518 auto const threshold = m_strong_threshold;
1519
1520 // Threshold of row i: theta times the strongest connection, relaxed
1521 // by a few ulps so that entries nominally on the threshold are strong
1522 // in every row [1]. Rows dominated by their diagonal (row sum above
1523 // 0.9 |a_ii|) have no strong connection.
1524 T const tol = T(64) * std::numeric_limits<T>::epsilon();
1525
1526#ifndef AMREX_USE_GPU
1527 {
1528 // One pass per row on the host.
1529 auto const& c0 = a.csr0;
1530 auto const& c1 = a.csr1;
1531 auto rlen1 = [&] (Long i) -> Long {
1532 if (!detail::has_remote_row(a, i)) { return 0; }
1533 Long const ii = a.row_map[i];
1534 return c1.row_offset[ii+1] - c1.row_offset[ii];
1535 };
1536 auto csr = detail::csr_from_rows_cpu<TS,SpMatrix<TS>::template container_type,int>
1537 (nrows, [&] (Long i) { return c0.row_offset[i+1] - c0.row_offset[i] + rlen1(i); },
1538 [&] (int) {
1539 return [&, col = Vector<int>(), val = Vector<TS>()]
1540 (Long i, int const*& rc, TS const*& rv) mutable -> Long
1541 {
1542 Long const b = c0.row_offset[i];
1543 Long const e = c0.row_offset[i+1];
1544 Long const n1 = rlen1(i);
1545 Long const b1 = (n1 > 0) ? Long(c1.row_offset[a.row_map[i]]) : 0;
1546 T const sgn = std::copysign(T(1), A_ii[i]);
1547 T rs = T(0);
1548 T amin = std::numeric_limits<T>::max();
1549 for (Long idx = b; idx < e; ++idx) {
1550 rs += c0.mat[idx];
1551 if (c0.col_index[idx] != i) { amin = std::min(amin, sgn*c0.mat[idx]); }
1552 }
1553 for (Long idx = b1; idx < b1+n1; ++idx) {
1554 rs += c1.mat[idx];
1555 amin = std::min(amin, sgn*c1.mat[idx]);
1556 }
1557 Long n = 0;
1558 if (!(std::abs(rs) > T(0.9)*std::abs(A_ii[i])) && amin < T(0)) {
1559 T const thr = amin * threshold * (T(1) - tol);
1560 if (Long(col.size()) < e-b+n1) {
1561 col.resize(e-b+n1);
1562 val.resize(e-b+n1);
1563 }
1564 for (Long idx = b; idx < e; ++idx) {
1565 if (c0.col_index[idx] != i && sgn*c0.mat[idx] <= thr) {
1566 col[n] = c0.col_index[idx];
1567 val[n] = TS(c0.mat[idx]/scale);
1568 ++n;
1569 }
1570 }
1571 for (Long idx = b1; idx < b1+n1; ++idx) {
1572 if (sgn*c1.mat[idx] <= thr) {
1573 col[n] = nlocal + c1.col_index[idx];
1574 val[n] = TS(c1.mat[idx]/scale);
1575 ++n;
1576 }
1577 }
1578 }
1579 rc = col.data();
1580 rv = val.data();
1581 return n;
1582 };
1583 });
1584 SpMatrix<TS> S;
1585 S.define_split(A.partition(), A.partition(), std::move(csr), nlocal, a.col_map, 0);
1586 return S;
1587 }
1588#else
1589
1590 auto const rsA = A.rowSum();
1591 auto const* AMREX_RESTRICT prs = rsA.data();
1592 auto row_thresh = [=] AMREX_GPU_DEVICE (Long i) -> T {
1593 if (std::abs(prs[i]) > T(0.9)*std::abs(A_ii[i])) {
1594 return std::numeric_limits<T>::lowest();
1595 }
1596 auto const& c0 = a.csr0;
1597 T const sgn = std::copysign(T(1), A_ii[i]);
1598 T amin = std::numeric_limits<T>::max();
1599 for (Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
1600 if (c0.col_index[idx] != i) {
1601 amin = std::min(amin, sgn*c0.mat[idx]);
1602 }
1603 }
1604 if (detail::has_remote_row(a, i)) {
1605 auto const& c1 = a.csr1;
1606 Long const ii = a.row_map[i];
1607 for (Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
1608 amin = std::min(amin, sgn*c1.mat[idx]);
1609 }
1610 }
1611 if (amin >= T(0)) { return std::numeric_limits<T>::lowest(); }
1612 return amin * threshold * (T(1) - tol);
1613 };
1614 Gpu::DeviceVector<T> rthr(nrows);
1615 auto* AMREX_RESTRICT pthr = rthr.data();
1616 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept { pthr[i] = row_thresh(i); });
1617
1618 using local_csr_type = typename SpMatrix<TS>::local_csr_type;
1619 local_csr_type csr;
1620 csr.row_offset.resize(nrows+1);
1621 auto* AMREX_RESTRICT psrow = csr.row_offset.data();
1622 Long const nnz = Scan::PrefixSum<Long>
1623 (nrows+1,
1624 [=] AMREX_GPU_DEVICE (Long i) -> Long {
1625 if (i >= nrows) { return 0; }
1626 T const thr = pthr[i];
1627 T const sgn = std::copysign(T(1), A_ii[i]);
1628 Long n = 0;
1629 auto const& c0 = a.csr0;
1630 for (Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
1631 if (c0.col_index[idx] != i && sgn*c0.mat[idx] <= thr) { ++n; }
1632 }
1633 if (detail::has_remote_row(a, i)) {
1634 auto const& c1 = a.csr1;
1635 Long const ii = a.row_map[i];
1636 for (Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
1637 if (sgn*c1.mat[idx] <= thr) { ++n; }
1638 }
1639 }
1640 return n;
1641 },
1642 [=] AMREX_GPU_DEVICE (Long i, Long const& x) { psrow[i] = x; },
1644
1645 csr.resize(nrows, nnz);
1646 auto* AMREX_RESTRICT ps = csr.mat.data();
1647 auto* AMREX_RESTRICT pc = csr.col_index.data();
1648 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
1649 {
1650 T const thr = pthr[i];
1651 T const sgn = std::copysign(T(1), A_ii[i]);
1652 Long p = psrow[i];
1653 auto const& c0 = a.csr0;
1654 for (Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
1655 if (c0.col_index[idx] != i && sgn*c0.mat[idx] <= thr) {
1656 ps[p] = TS(c0.mat[idx]/scale);
1657 pc[p] = c0.col_index[idx];
1658 ++p;
1659 }
1660 }
1661 if (detail::has_remote_row(a, i)) {
1662 auto const& c1 = a.csr1;
1663 Long const ii = a.row_map[i];
1664 for (Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
1665 if (sgn*c1.mat[idx] <= thr) {
1666 ps[p] = TS(c1.mat[idx]/scale);
1667 pc[p] = nlocal + c1.col_index[idx];
1668 ++p;
1669 }
1670 }
1671 }
1672 });
1674
1675 SpMatrix<TS> S;
1676 S.define_split(A.partition(), A.partition(), std::move(csr), nlocal, a.col_map, 0);
1677 return S;
1678#endif
1679}
1680
1681template <typename T>
1683 bool isolated_as_coarse) const
1684{
1685 BL_PROFILE("AlgMG::pmis");
1686
1687 Long const nrows = S.numLocalRows();
1688 int const U = PointType::undecided;
1689 int const C = PointType::coarse;
1690
1691 // Measure: number of points that strongly depend on i. A negative
1692 // measure marks a point with no strong connection at all.
1693 auto const& s = S.const_parcsr();
1694 auto const& st = ST.const_parcsr();
1695 Gpu::DeviceVector<Long> measure(nrows);
1696 auto* AMREX_RESTRICT pm = measure.data();
1697 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
1698 {
1699 Long n = st.csr0.row_offset[i+1] - st.csr0.row_offset[i];
1700 if (detail::has_remote_row(st, i)) {
1701 Long const ii = st.row_map[i];
1702 n += st.csr1.row_offset[ii+1] - st.csr1.row_offset[ii];
1703 }
1704 Long nout = s.csr0.row_offset[i+1] - s.csr0.row_offset[i];
1705 if (detail::has_remote_row(s, i)) {
1706 Long const ii = s.row_map[i];
1707 nout += s.csr1.row_offset[ii+1] - s.csr1.row_offset[ii];
1708 }
1709 pm[i] = (n == 0 && nout == 0) ? Long(-1) : n;
1710 });
1711
1712 Gpu::DeviceVector<T> weight(nrows);
1713 state.resize(nrows);
1714 Gpu::DeviceVector<int> state2(nrows);
1715 auto* AMREX_RESTRICT pw = weight.data();
1716 auto* AMREX_RESTRICT ps = state.data();
1717 auto* AMREX_RESTRICT ps2 = state2.data();
1718
1719 Long const row_begin = s.row_begin;
1720
1721 // A point nobody depends on gets a weight below one, so it becomes C
1722 // only if all its dependencies are F.
1723 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
1724 {
1725 if (pm[i] < 0) {
1726 pw[i] = T(0);
1727 ps[i] = isolated_as_coarse ? PointType::coarse : PointType::fine;
1728 } else {
1729 pw[i] = T(pm[i]) + detail::pmis_tiebreak<T>(i + row_begin);
1730 ps[i] = PointType::undecided;
1731 }
1732 });
1733
1734 // Ghost weights are fixed; ghost states are gathered every sweep.
1735 auto const wS = S.gatherRemote(pw);
1736 auto const wST = ST.gatherRemote(pw);
1737 auto const* AMREX_RESTRICT pwS = wS.data();
1738 auto const* AMREX_RESTRICT pwST = wST.data();
1739
1740 Gpu::DeviceVector<int> sS, sST, s2S;
1741
1742 while (true) {
1743 S.gatherRemote(ps, sS);
1744 ST.gatherRemote(ps, sST);
1745 auto const* AMREX_RESTRICT psS = sS.data();
1746 auto const* AMREX_RESTRICT psST = sST.data();
1747
1748 // Independent set in S union S^T: keep only own index writes.
1749 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
1750 {
1751 int st_i = ps[i];
1752 if (st_i == U) {
1753 T const wi = pw[i];
1754 Long const ig = i + row_begin;
1755 bool blocked =
1756 detail::pmis_blocked(s.csr0, i, wi, ig, pw, ps, nullptr, s.col_begin, U) ||
1757 detail::pmis_blocked(st.csr0, i, wi, ig, pw, ps, nullptr, st.col_begin, U);
1758 if (!blocked && detail::has_remote_row(s, i)) {
1759 blocked = detail::pmis_blocked(s.csr1, s.row_map[i], wi, ig,
1760 pwS, psS, s.col_map, 0, U);
1761 }
1762 if (!blocked && detail::has_remote_row(st, i)) {
1763 blocked = detail::pmis_blocked(st.csr1, st.row_map[i], wi, ig,
1764 pwST, psST, st.col_map, 0, U);
1765 }
1766 if (!blocked) { st_i = C; }
1767 }
1768 ps2[i] = st_i;
1769 });
1770
1771 S.gatherRemote(ps2, s2S);
1772 auto const* AMREX_RESTRICT ps2S = s2S.data();
1773
1774 // Undecided points that strongly depend on a C-point become F.
1775 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
1776 {
1777 int st_i = ps2[i];
1778 if (st_i == U) {
1779 if (detail::csr_row_has_state(s.csr0, i, ps2, C) ||
1780 (detail::has_remote_row(s, i) &&
1781 detail::csr_row_has_state(s.csr1, s.row_map[i], ps2S, C))) {
1782 st_i = PointType::fine;
1783 }
1784 }
1785 ps[i] = st_i;
1786 });
1787
1788 Long nund = Reduce::Sum<Long>(nrows, [=] AMREX_GPU_DEVICE (Long i) -> Long {
1789 return (ps[i] == U) ? 1 : 0;
1790 });
1792 if (nund == 0) { break; }
1793 }
1794
1795 if (m_verbose >= 2) {
1796 S.gatherRemote(ps, sS);
1797 auto const* AMREX_RESTRICT psS = sS.data();
1798 Long nc = Reduce::Sum<Long>(nrows, [=] AMREX_GPU_DEVICE (Long i) -> Long {
1799 return (ps[i] == C) ? 1 : 0;
1800 });
1801 // C-points that strongly depend on another C-point. A point becomes
1802 // F only through its own dependencies, so a few one-directional
1803 // C-C connections remain where the strength is asymmetric.
1804 Long ncc = Reduce::Sum<Long>(nrows, [=] AMREX_GPU_DEVICE (Long i) -> Long {
1805 if (ps[i] != C) { return 0; }
1806 return (detail::csr_row_has_state(s.csr0, i, ps, C) ||
1807 (detail::has_remote_row(s, i) &&
1808 detail::csr_row_has_state(s.csr1, s.row_map[i], psS, C))) ? 1 : 0;
1809 });
1810 // F-points without a strong C neighbor get no interpolation.
1811 Long nf0 = Reduce::Sum<Long>(nrows, [=] AMREX_GPU_DEVICE (Long i) -> Long {
1812 if (ps[i] != PointType::fine) { return 0; }
1813 return (detail::csr_row_has_state(s.csr0, i, ps, C) ||
1814 (detail::has_remote_row(s, i) &&
1815 detail::csr_row_has_state(s.csr1, s.row_map[i], psS, C))) ? 0 : 1;
1816 });
1817 Long counts[3] = {nc, ncc, nf0};
1819 amrex::Print() << m_print_ident << "AlgMG: PMIS: " << counts[0] << " C-points of "
1820 << S.numGlobalRows() << ", C-C strong pairs " << counts[1]
1821 << ", F-points w/o C neighbor " << counts[2] << "\n";
1822 }
1823}
1824
1825template <typename T>
1827 Gpu::DeviceVector<Long>& cidx, int which)
1828{
1829 Long const nrows = Long(state.size());
1830 auto const* AMREX_RESTRICT ps = state.data();
1831 cidx.resize(nrows);
1832 auto* AMREX_RESTRICT pc = cidx.data();
1833 Long const nc = Scan::PrefixSum<Long>
1834 (nrows,
1835 [=] AMREX_GPU_DEVICE (Long i) -> Long { return (ps[i] == which) ? 1 : 0; },
1836 [=] AMREX_GPU_DEVICE (Long i, Long const& x) { pc[i] = x; },
1838
1839 int const nprocs = ParallelContext::NProcsSub();
1840 Vector<Long> counts(nprocs, nc);
1841#ifdef AMREX_USE_MPI
1842 // One value per rank: an all-gather, not a reduction of a P-vector.
1844#endif
1845 Vector<Long> rows(nprocs+1, 0); // exclusive prefix sum of counts
1846 std::partial_sum(counts.begin(), counts.end(), rows.begin()+1);
1847 return AlgPartition(std::move(rows));
1848}
1849
1850template <typename T>
1852 Gpu::DeviceVector<int> const& state,
1853 Gpu::DeviceVector<Long> const& cidx,
1854 AlgPartition const& cpart)
1855{
1856 BL_PROFILE("AlgMG::build_interp");
1857
1858 // Classical direct interpolation [1].
1859
1860 auto const& a = A.const_parcsr();
1861 Long const nrows = a.csr0.nrows;
1862 auto const* AMREX_RESTRICT ps = state.data();
1863 auto const* AMREX_RESTRICT pc = cidx.data();
1864 Long const cbegin = cpart.globalRowBegin();
1865 int const nlocal_c = int(cpart.numLocalRows());
1866
1867 // Global coarse index of local rows, and ghost copies for A's remote
1868 // columns (the remote column list of P).
1869 Gpu::DeviceVector<Long> gcidx(nrows);
1870 auto* AMREX_RESTRICT pgc = gcidx.data();
1871 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept { pgc[i] = pc[i] + cbegin; });
1872 auto const state_r = A.gatherRemote(ps);
1873 auto const gcidx_r = A.gatherRemote(pgc);
1874
1875 detail::DirectInterpRow<T,TS> const w{a, S.const_parcsr(), ps, pc,
1876 state_r.data(), nlocal_c,
1877 int(PointType::coarse)};
1878
1879 using local_csr_type = typename SpMatrix<T>::local_csr_type;
1880 local_csr_type csr;
1881 csr.row_offset.resize(nrows+1);
1882 auto* AMREX_RESTRICT prow = csr.row_offset.data();
1883
1884 Long const nnz = Scan::PrefixSum<Long>
1885 (nrows+1,
1886 [=] AMREX_GPU_DEVICE (Long i) -> Long {
1887 if (i >= nrows) { return 0; }
1888 if (ps[i] == PointType::coarse) { return 1; }
1889 Long n = 0;
1890 T diag;
1891 w.walk(i, diag, [&] (int, T, bool strong_c) { if (strong_c) { ++n; } });
1892 return n;
1893 },
1894 [=] AMREX_GPU_DEVICE (Long i, Long const& x) { prow[i] = x; },
1896
1897 csr.resize(nrows, nnz);
1898 auto* AMREX_RESTRICT pmat = csr.mat.data();
1899 auto* AMREX_RESTRICT pcol = csr.col_index.data();
1900
1901 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
1902 {
1903 Long p = prow[i];
1904 if (ps[i] == PointType::coarse) {
1905 pcol[p] = int(pc[i]);
1906 pmat[p] = T(1);
1907 return;
1908 }
1909
1910 T sum_n_neg = T(0), sum_n_pos = T(0), sum_p_neg = T(0), sum_p_pos = T(0);
1911 T diag;
1912 w.walk(i, diag, [&] (int, T v, bool strong_c) {
1913 if (v < T(0)) {
1914 sum_n_neg += v;
1915 if (strong_c) { sum_p_neg += v; }
1916 } else {
1917 sum_n_pos += v;
1918 if (strong_c) { sum_p_pos += v; }
1919 }
1920 });
1921
1922 T alfa = T(0), beta = T(0);
1923 if (sum_p_neg != T(0)) { alfa = sum_n_neg / sum_p_neg; } else { diag += sum_n_neg; }
1924 if (sum_p_pos != T(0)) { beta = sum_n_pos / sum_p_pos; } else { diag += sum_n_pos; }
1925 if (diag == T(0)) { diag = T(1); } // no interpolation weights anyway
1926
1927 T unused;
1928 w.walk(i, unused, [&] (int cj, T v, bool strong_c) {
1929 if (strong_c) {
1930 pcol[p] = cj;
1931 pmat[p] = ((v < T(0)) ? -alfa*v : -beta*v) / diag;
1932 ++p;
1933 }
1934 });
1935 });
1937
1938 SpMatrix<T> P;
1939 P.define_split(A.partition(), cpart, std::move(csr), nlocal_c, gcidx_r.data(), 0);
1940 return P;
1941}
1942
1943namespace detail {
1944
1948template <typename T, int KMAX>
1949struct TruncRow
1950{
1952 int max_elmts;
1953 T trunc_factor;
1954 int nlocal_cols;
1955
1957 template <typename F>
1958 AMREX_GPU_HOST_DEVICE void walk (Long i, F const& f) const
1959 {
1960 auto const& c0 = p.csr0;
1961 for (Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
1962 f(c0.col_index[idx], c0.mat[idx]);
1963 }
1964 if (has_remote_row(p, i)) {
1965 auto const& c1 = p.csr1;
1966 Long const ii = p.row_map[i];
1967 for (Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
1968 f(nlocal_cols + c1.col_index[idx], c1.mat[idx]);
1969 }
1970 }
1971 }
1972
1974 AMREX_GPU_HOST_DEVICE T threshold (Long i, T& sum_pos, T& sum_neg) const
1975 {
1976 T amax = T(0);
1977 sum_pos = T(0);
1978 sum_neg = T(0);
1979 walk(i, [&] (int, T v) {
1980 amax = std::max(amax, std::abs(v));
1981 if (v > T(0)) { sum_pos += v; } else { sum_neg += v; }
1982 });
1983 return trunc_factor * amax;
1984 }
1985
1987 [[nodiscard]] AMREX_GPU_HOST_DEVICE int count (Long i, T thresh) const
1988 {
1989 int n = 0;
1990 walk(i, [&] (int, T v) { if (std::abs(v) >= thresh) { ++n; } });
1991 return n;
1992 }
1993
1998 int select (Long i, T thresh, int* cols, T* vals) const
1999 {
2000 int const kmax = std::min(max_elmts, KMAX);
2001 int n = 0;
2002 walk(i, [&] (int col, T v) {
2003 T const av = std::abs(v);
2004 if (av < thresh) { return; }
2005 if (n == kmax && av <= std::abs(vals[n-1])) { return; }
2006 int pos = (n < kmax) ? n : n-1;
2007 while (pos > 0 && std::abs(vals[pos-1]) < av) {
2008 cols[pos] = cols[pos-1];
2009 vals[pos] = vals[pos-1];
2010 --pos;
2011 }
2012 cols[pos] = col;
2013 vals[pos] = v;
2014 if (n < kmax) { ++n; }
2015 });
2016 return n;
2017 }
2018};
2019
2020#ifndef AMREX_USE_GPU
2023template <typename T, int KMAX>
2024Long truncate_row_host (int* col, T* val, Long n, int max_elmts, T trunc_factor, int nc)
2025{
2026 // The row as a one-row CSR for TruncRow.
2027 int ro[2] = {0, int(n)};
2028 ParCsr<T const> prow;
2029 prow.csr0 = CsrView<T const,int>{val, col, ro, n, 1};
2030 TruncRow<T,KMAX> const tr{prow, max_elmts, trunc_factor, nc};
2031 T sum_pos, sum_neg;
2032 T const thresh = tr.threshold(0, sum_pos, sum_neg);
2033 int tcol[KMAX];
2034 T tval[KMAX];
2035 Long nk = 0;
2036 T kept_pos = T(0), kept_neg = T(0);
2037 if (max_elmts > 0) {
2038 nk = tr.select(0, thresh, tcol, tval);
2039 for (Long k = 0; k < nk; ++k) {
2040 if (tval[k] > T(0)) { kept_pos += tval[k]; } else { kept_neg += tval[k]; }
2041 }
2042 } else {
2043 for (Long k = 0; k < n; ++k) {
2044 if (std::abs(val[k]) >= thresh) {
2045 if (val[k] > T(0)) { kept_pos += val[k]; } else { kept_neg += val[k]; }
2046 }
2047 }
2048 }
2049 T const fpos = (kept_pos != T(0)) ? sum_pos/kept_pos : T(1);
2050 T const fneg = (kept_neg != T(0)) ? sum_neg/kept_neg : T(1);
2051 if (max_elmts > 0) {
2052 // Kept by decreasing magnitude; back to column order.
2053 for (Long k = 0; k < nk; ++k) {
2054 Long m = k;
2055 for (; m > 0 && col[m-1] > tcol[k]; --m) {
2056 col[m] = col[m-1];
2057 val[m] = val[m-1];
2058 }
2059 col[m] = tcol[k];
2060 val[m] = tval[k] * ((tval[k] > T(0)) ? fpos : fneg);
2061 }
2062 } else {
2063 for (Long k = 0; k < n; ++k) {
2064 if (std::abs(val[k]) >= thresh) {
2065 col[nk] = col[k];
2066 val[nk] = val[k] * ((val[k] > T(0)) ? fpos : fneg);
2067 ++nk;
2068 }
2069 }
2070 }
2071 return nk;
2072}
2073#endif
2074
2075}
2076
2077namespace detail {
2078
2083template <typename T, typename F>
2085compact_rows (ParCsr<T const> const& p, Long nc, int const* AMREX_RESTRICT state,
2086 Long const* AMREX_RESTRICT cidx, int c, F const& colmap)
2087{
2088 Long const nrows = p.csr0.nrows;
2089
2090 // Offsets by fine row, then gathered at the selected rows.
2091 Gpu::DeviceVector<Long> off(nrows+1);
2092 auto* AMREX_RESTRICT poff = off.data();
2093 Long const nnz = Scan::PrefixSum<Long>
2094 (nrows+1,
2095 [=] AMREX_GPU_DEVICE (Long i) -> Long {
2096 if (i >= nrows || state[i] != c) { return 0; }
2097 Long n = 0;
2098 auto const& c0 = p.csr0;
2099 for (Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
2100 if (colmap(c0.col_index[idx], false) >= 0) { ++n; }
2101 }
2102 if (has_remote_row(p, i)) {
2103 auto const& c1 = p.csr1;
2104 Long const ii = p.row_map[i];
2105 for (Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
2106 if (colmap(c1.col_index[idx], true) >= 0) { ++n; }
2107 }
2108 }
2109 return n;
2110 },
2111 [=] AMREX_GPU_DEVICE (Long i, Long const& x) { poff[i] = x; },
2113
2114 typename SpMatrix<T>::local_csr_type csr;
2115 csr.resize(nc, nnz);
2116 auto* AMREX_RESTRICT prow = csr.row_offset.data();
2117 auto* AMREX_RESTRICT pcol = csr.col_index.data();
2118 auto* AMREX_RESTRICT pmat = csr.mat.data();
2119 ParallelForOMP(nrows+1, [=] AMREX_GPU_DEVICE (Long i) noexcept
2120 {
2121 if (i == nrows) { prow[nc] = nnz; return; }
2122 if (state[i] != c) { return; }
2123 Long q = poff[i];
2124 prow[cidx[i]] = q;
2125 auto const& c0 = p.csr0;
2126 for (Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
2127 int const g = colmap(c0.col_index[idx], false);
2128 if (g >= 0) { pcol[q] = g; pmat[q] = c0.mat[idx]; ++q; }
2129 }
2130 if (has_remote_row(p, i)) {
2131 auto const& c1 = p.csr1;
2132 Long const ii = p.row_map[i];
2133 for (Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
2134 int const g = colmap(c1.col_index[idx], true);
2135 if (g >= 0) { pcol[q] = g; pmat[q] = c1.mat[idx]; ++q; }
2136 }
2137 }
2138 });
2140 return csr;
2141}
2142
2143}
2144
2145template <typename T>
2147 Gpu::DeviceVector<int> const& state,
2148 Gpu::DeviceVector<Long> const& cidx,
2149 AlgPartition const& cpart, T scale) const
2150{
2151 if (m_interp_type == InterpType::direct) {
2152 return build_interp(A, S, state, cidx, cpart);
2153 } else {
2154 return build_interp_mm(A, S, ST, state, cidx, cpart, scale,
2155 m_interp_type == InterpType::mm_ext_i);
2156 }
2157}
2158
2159template <typename T>
2161 Gpu::DeviceVector<int> const& state,
2162 Gpu::DeviceVector<Long> const& cidx,
2163 AlgPartition const& cpart)
2164{
2165 BL_PROFILE("AlgMG::second_pass_strength");
2166
2167 // Distance-two C-C connections: S_CF S_FC. The few direct C-C
2168 // connections PMIS leaves are ignored. Both factors are built in
2169 // compact C and F numberings, so the product never forms the full S^2.
2170 auto const& s = S.const_parcsr();
2171 Long const nrows = s.csr0.nrows;
2172 Long const nc = cpart.numLocalRows();
2173 int const C = PointType::coarse;
2174 int const F = PointType::fine;
2175
2177 auto const fpart = coarse_numbering(state, fidx, PointType::fine);
2178 Long const nf = fpart.numLocalRows();
2179 int const nlocal_c = int(nc);
2180 int const nlocal_f = int(nf);
2181
2182 auto const* AMREX_RESTRICT ps = state.data();
2183 auto const* AMREX_RESTRICT pc = cidx.data();
2184 auto const* AMREX_RESTRICT pf = fidx.data();
2185 Long const cbegin = cpart.globalRowBegin();
2186 Long const fbegin = fpart.globalRowBegin();
2187 Gpu::DeviceVector<Long> gcidx(nrows);
2188 Gpu::DeviceVector<Long> gfidx(nrows);
2189 auto* AMREX_RESTRICT pgc = gcidx.data();
2190 auto* AMREX_RESTRICT pgf = gfidx.data();
2191 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
2192 {
2193 pgc[i] = pc[i] + cbegin;
2194 pgf[i] = pf[i] + fbegin;
2195 });
2196 auto const state_r = S.gatherRemote(ps);
2197 auto const gcidx_r = S.gatherRemote(pgc);
2198 auto const gfidx_r = S.gatherRemote(pgf);
2199 auto const* AMREX_RESTRICT ps_r = state_r.data();
2200
2201 // S_CF: C rows, F columns in F numbering (remote: nlocal_f + remote index).
2202 auto csr_cf = detail::compact_rows<TS>
2203 (s, nc, ps, pc, C,
2204 [=] AMREX_GPU_DEVICE (int j, bool remote) -> int {
2205 if (remote) {
2206 return (ps_r[j] == F) ? nlocal_f + j : -1;
2207 } else {
2208 return (ps[j] == F) ? int(pf[j]) : -1;
2209 }
2210 });
2211 SpMatrix<TS> SCF;
2212 SCF.define_split(cpart, fpart, std::move(csr_cf), nlocal_f, gfidx_r.data(), 0);
2213
2214 // S_FC: F rows, C columns in C numbering.
2215 auto csr_fc = detail::compact_rows<TS>
2216 (s, nf, ps, pf, F,
2217 [=] AMREX_GPU_DEVICE (int j, bool remote) -> int {
2218 if (remote) {
2219 return (ps_r[j] == C) ? nlocal_c + j : -1;
2220 } else {
2221 return (ps[j] == C) ? int(pc[j]) : -1;
2222 }
2223 });
2224 SpMatrix<TS> SFC;
2225 SFC.define_split(fpart, cpart, std::move(csr_fc), nlocal_c, gcidx_r.data(), 0);
2226
2227 auto S2 = amrex::SpGEMM(SCF, SFC, cpart);
2228
2229 // Drop the diagonal: a C-point is not its own distance-two neighbor.
2230 auto const& s2 = S2.const_parcsr();
2231 typename SpMatrix<TS>::local_csr_type out;
2232 out.row_offset.resize(nc+1);
2233 auto* AMREX_RESTRICT orow = out.row_offset.data();
2234 Long const nnz = Scan::PrefixSum<Long>
2235 (nc+1,
2236 [=] AMREX_GPU_DEVICE (Long r) -> Long {
2237 if (r >= nc) { return 0; }
2238 Long n = 0;
2239 auto const& c0 = s2.csr0;
2240 for (Long idx = c0.row_offset[r]; idx < c0.row_offset[r+1]; ++idx) {
2241 n += (c0.col_index[idx] != int(r));
2242 }
2243 if (detail::has_remote_row(s2, r)) {
2244 Long const rr = s2.row_map[r];
2245 n += s2.csr1.row_offset[rr+1] - s2.csr1.row_offset[rr];
2246 }
2247 return n;
2248 },
2249 [=] AMREX_GPU_DEVICE (Long r, Long const& x) { orow[r] = x; },
2251 out.resize(nc, nnz);
2252 auto* AMREX_RESTRICT ocol = out.col_index.data();
2253 auto* AMREX_RESTRICT omat = out.mat.data();
2254 ParallelForOMP(nc, [=] AMREX_GPU_DEVICE (Long r) noexcept
2255 {
2256 Long q = orow[r];
2257 auto const& c0 = s2.csr0;
2258 for (Long idx = c0.row_offset[r]; idx < c0.row_offset[r+1]; ++idx) {
2259 if (c0.col_index[idx] != int(r)) {
2260 ocol[q] = c0.col_index[idx];
2261 omat[q] = c0.mat[idx];
2262 ++q;
2263 }
2264 }
2265 if (detail::has_remote_row(s2, r)) {
2266 auto const& c1 = s2.csr1;
2267 Long const rr = s2.row_map[r];
2268 for (Long idx = c1.row_offset[rr]; idx < c1.row_offset[rr+1]; ++idx) {
2269 ocol[q] = nlocal_c + c1.col_index[idx];
2270 omat[q] = c1.mat[idx];
2271 ++q;
2272 }
2273 }
2274 });
2276 SpMatrix<TS> S2d;
2277 S2d.define_split(cpart, cpart, std::move(out), nlocal_c, s2.col_map, 0);
2278 return S2d;
2279}
2280
2281template <typename T>
2283 Gpu::DeviceVector<int> const& state,
2284 Gpu::DeviceVector<Long> const& cidx,
2285 AlgPartition const& cpart)
2286{
2287 auto const& p = P.const_parcsr();
2288 int const nlocal = int(P.columnPartition().numLocalRows());
2289 auto csr = detail::compact_rows<T>
2290 (p, cpart.numLocalRows(), state.data(), cidx.data(), int(PointType::coarse),
2291 [=] AMREX_GPU_DEVICE (int j, bool remote) -> int {
2292 return remote ? nlocal + j : j;
2293 });
2294 SpMatrix<T> Q;
2295 Q.define_split(cpart, P.columnPartition(), std::move(csr), nlocal, p.col_map, 0);
2296 return Q;
2297}
2298
2299template <typename T>
2301 int max_elmts, T trunc_factor)
2302{
2303 BL_PROFILE("AlgMG::truncate_interp");
2304
2305 constexpr int KMAX = max_p_elmts; // rows of P are short
2306 AMREX_ALWAYS_ASSERT(max_elmts <= KMAX);
2307
2308 P.setColumnPartition(cpart);
2309 auto const& p = P.const_parcsr();
2310 Long const nrows = p.csr0.nrows;
2311 int const nlocal = int(cpart.numLocalRows());
2312 detail::TruncRow<T,KMAX> const tr{p, max_elmts, trunc_factor, nlocal};
2313
2314 // Drop threshold and signed sums of each row, computed once.
2315 Gpu::DeviceVector<T> rthr(nrows), rpos(nrows), rneg(nrows);
2316 auto* AMREX_RESTRICT pthr = rthr.data();
2317 auto* AMREX_RESTRICT ppos = rpos.data();
2318 auto* AMREX_RESTRICT pneg = rneg.data();
2319 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
2320 {
2321 pthr[i] = tr.threshold(i, ppos[i], pneg[i]);
2322 });
2323
2324 using local_csr_type = typename SpMatrix<T>::local_csr_type;
2325 local_csr_type csr;
2326 csr.row_offset.resize(nrows+1);
2327 auto* AMREX_RESTRICT prow = csr.row_offset.data();
2328 Long const nnz = Scan::PrefixSum<Long>
2329 (nrows+1,
2330 [=] AMREX_GPU_DEVICE (Long i) -> Long {
2331 if (i >= nrows) { return 0; }
2332 int const n = tr.count(i, pthr[i]);
2333 return (max_elmts > 0) ? std::min(n, max_elmts) : n;
2334 },
2335 [=] AMREX_GPU_DEVICE (Long i, Long const& x) { prow[i] = x; },
2337
2338 csr.resize(nrows, nnz);
2339 auto* AMREX_RESTRICT pmat = csr.mat.data();
2340 auto* AMREX_RESTRICT pcol = csr.col_index.data();
2341 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
2342 {
2343 T const thresh = pthr[i];
2344 T const sum_pos = ppos[i];
2345 T const sum_neg = pneg[i];
2346 Long q = prow[i];
2347 if (max_elmts > 0) {
2348 int cols[KMAX];
2349 T vals[KMAX];
2350 int const n = tr.select(i, thresh, cols, vals);
2351 T kept_pos = T(0), kept_neg = T(0);
2352 for (int k = 0; k < n; ++k) {
2353 if (vals[k] > T(0)) { kept_pos += vals[k]; } else { kept_neg += vals[k]; }
2354 }
2355 T const fpos = (kept_pos != T(0)) ? sum_pos/kept_pos : T(1);
2356 T const fneg = (kept_neg != T(0)) ? sum_neg/kept_neg : T(1);
2357 for (int k = 0; k < n; ++k) {
2358 pcol[q] = cols[k];
2359 pmat[q] = vals[k] * ((vals[k] > T(0)) ? fpos : fneg);
2360 ++q;
2361 }
2362 } else { // threshold only, no limit on the count
2363 T kept_pos = T(0), kept_neg = T(0);
2364 tr.walk(i, [&] (int, T v) {
2365 if (std::abs(v) >= thresh) {
2366 if (v > T(0)) { kept_pos += v; } else { kept_neg += v; }
2367 }
2368 });
2369 T const fpos = (kept_pos != T(0)) ? sum_pos/kept_pos : T(1);
2370 T const fneg = (kept_neg != T(0)) ? sum_neg/kept_neg : T(1);
2371 tr.walk(i, [&] (int col, T v) {
2372 if (std::abs(v) >= thresh) {
2373 pcol[q] = col;
2374 pmat[q] = v * ((v > T(0)) ? fpos : fneg);
2375 ++q;
2376 }
2377 });
2378 }
2379 });
2381
2382 // Compact columns sort like the global ones within each block.
2383 csr.sort();
2384 SpMatrix<T> Pt;
2385 Pt.define_split(P.partition(), cpart, std::move(csr), nlocal, p.col_map, 0);
2386 return Pt;
2387}
2388
2389template <typename T>
2391 SpMatrix<TS> const& ST,
2392 Gpu::DeviceVector<int> const& state,
2393 Gpu::DeviceVector<Long> const& cidx,
2394 AlgPartition const& cpart, T sscale,
2395 bool plus_i, int const* row_state,
2396 int max_elmts, T trunc_factor, Long* nnz_untruncated)
2397{
2398 BL_PROFILE("AlgMG::build_interp_mm");
2399
2400 // W = -M_F B_F of Li, Sjögreen & Yang (2021), Eqs. (3.4)-(3.12). Both
2401 // factors keep the natural row numbering with unit C rows, so P is one
2402 // product SpGEMM(M, B).
2403
2404 auto const& s = S.const_parcsr();
2405 Long const nrows = s.csr0.nrows;
2406 auto const* AMREX_RESTRICT ps = state.data();
2407 auto const* AMREX_RESTRICT pc = cidx.data();
2408 Long const cbegin = cpart.globalRowBegin();
2409 int const nlocal_f = int(A.numLocalRows());
2410 int const nlocal_c = int(cpart.numLocalRows());
2411 int const C = PointType::coarse;
2412 int const F = PointType::fine;
2413
2414 // alpha_i = a_ii + weak off-diagonal row sum = rowsum(A) - rowsum(S),
2415 // in the units of S. The weights do not depend on the scaling.
2416 auto const rsA = A.rowSum();
2417 auto const rsS = S.rowSum();
2418 auto const* AMREX_RESTRICT prsA = rsA.data();
2419 auto const* AMREX_RESTRICT prsS = rsS.data();
2420
2421 Gpu::DeviceVector<Long> gcidx(nrows);
2422 auto* AMREX_RESTRICT pgc = gcidx.data();
2423 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept { pgc[i] = pc[i] + cbegin; });
2424 auto const state_r = S.gatherRemote(ps);
2425 auto const gcidx_r = S.gatherRemote(pgc);
2426 auto const* AMREX_RESTRICT pstate_r = state_r.data();
2427
2428 // D_beta: row sum of strong connections to C-points.
2429 Gpu::DeviceVector<T> dbeta(nrows);
2430 auto* AMREX_RESTRICT pdb = dbeta.data();
2431 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
2432 {
2433 T d = T(0);
2434 auto const& sc = s.csr0;
2435 for (Long idx = sc.row_offset[i]; idx < sc.row_offset[i+1]; ++idx) {
2436 if (ps[sc.col_index[idx]] == C) { d += T(sc.mat[idx]); }
2437 }
2438 if (detail::has_remote_row(s, i)) {
2439 auto const& sc1 = s.csr1;
2440 Long const ii = s.row_map[i];
2441 for (Long idx = sc1.row_offset[ii]; idx < sc1.row_offset[ii+1]; ++idx) {
2442 if (pstate_r[sc1.col_index[idx]] == C) { d += T(sc1.mat[idx]); }
2443 }
2444 }
2445 pdb[i] = d;
2446 });
2448 auto const dbeta_r = S.gatherRemote(pdb);
2449
2450 detail::MMInterpRow<T,TS> const w{s, ST.const_parcsr(), ps, pdb, pc,
2451 pstate_r, dbeta_r.data(), nlocal_f, nlocal_c};
2452
2453 // Counts: M gets a diagonal plus the usable strong F neighbors; B gets
2454 // the strong C neighbors. C rows get one unit entry each.
2455 using local_csr_type = typename SpMatrix<T>::local_csr_type;
2456 local_csr_type mcsr, bcsr;
2457 mcsr.row_offset.resize(nrows+1);
2458 bcsr.row_offset.resize(nrows+1);
2459 auto* AMREX_RESTRICT pmrow = mcsr.row_offset.data();
2460 auto* AMREX_RESTRICT pbrow = bcsr.row_offset.data();
2461
2462 auto usable_f = [=] AMREX_GPU_DEVICE (T aji, int st_j, T db_j) -> bool {
2463 return st_j == F && (db_j + (plus_i ? aji : T(0))) != T(0);
2464 };
2465
2466 // Rows of P: all, or the C-points of `row_state` (the others stay
2467 // empty in M and so in P, but B needs every row).
2468 auto const* AMREX_RESTRICT prow_state = row_state;
2469 auto has_row = [=] AMREX_GPU_DEVICE (Long i) -> bool {
2470 return prow_state == nullptr || prow_state[i] == C;
2471 };
2472
2473 Long const mnnz = Scan::PrefixSum<Long>
2474 (nrows+1,
2475 [=] AMREX_GPU_DEVICE (Long i) -> Long {
2476 if (i >= nrows || !has_row(i)) { return 0; }
2477 Long n = 1;
2478 if (ps[i] == F) {
2479 w.walk(i, [&] (int, int, T, T aji, int st_j, T db_j) {
2480 if (usable_f(aji, st_j, db_j)) { ++n; }
2481 });
2482 }
2483 return n;
2484 },
2485 [=] AMREX_GPU_DEVICE (Long i, Long const& x) { pmrow[i] = x; },
2487
2488 Long const bnnz = Scan::PrefixSum<Long>
2489 (nrows+1,
2490 [=] AMREX_GPU_DEVICE (Long i) -> Long {
2491 if (i >= nrows) { return 0; }
2492 if (ps[i] == C) { return 1; }
2493 Long n = 0;
2494 w.walk(i, [&] (int, int, T, T, int st_j, T) {
2495 if (st_j == C) { ++n; }
2496 });
2497 return n;
2498 },
2499 [=] AMREX_GPU_DEVICE (Long i, Long const& x) { pbrow[i] = x; },
2501
2502 mcsr.resize(nrows, mnnz);
2503 bcsr.resize(nrows, bnnz);
2504 auto* AMREX_RESTRICT pmmat = mcsr.mat.data();
2505 auto* AMREX_RESTRICT pmcol = mcsr.col_index.data();
2506 auto* AMREX_RESTRICT pbmat = bcsr.mat.data();
2507 auto* AMREX_RESTRICT pbcol = bcsr.col_index.data();
2508
2509 ParallelForOMP(nrows, [=] AMREX_GPU_DEVICE (Long i) noexcept
2510 {
2511 Long mp = pmrow[i];
2512 Long bp = pbrow[i];
2513 bool const in_p = has_row(i);
2514 if (ps[i] != F) {
2515 if (in_p) {
2516 pmcol[mp] = int(i);
2517 pmmat[mp] = T(1);
2518 }
2519 pbcol[bp] = int(pc[i]);
2520 pbmat[bp] = T(1);
2521 return;
2522 }
2523 T const db_i = pdb[i];
2524 T const bscale = plus_i ? T(1) : ((db_i != T(0)) ? T(1)/db_i : T(0));
2525 if (!in_p) {
2526 w.walk(i, [&] (int, int jc, T aij, T, int st_j, T) {
2527 if (st_j == C) {
2528 pbcol[bp] = jc;
2529 pbmat[bp] = bscale*aij;
2530 ++bp;
2531 }
2532 });
2533 return;
2534 }
2535
2536 // Strong F neighbors without a usable denominator are lumped into
2537 // the diagonal as weak connections (remark after Eq. (3.8)).
2538 T alpha = prsA[i]/sscale - T(prsS[i]);
2539 T theta = T(0);
2540 w.walk(i, [&] (int, int, T aij, T aji, int st_j, T db_j) {
2541 if (st_j == F) {
2542 T const den = db_j + (plus_i ? aji : T(0));
2543 if (den != T(0)) {
2544 if (plus_i) { theta += aij * aji / den; }
2545 } else {
2546 alpha += aij;
2547 }
2548 }
2549 });
2550 T const denom = alpha + theta;
2551 T const scale = (denom != T(0)) ? T(-1)/denom : T(0);
2552
2553 pmcol[mp] = int(i);
2554 pmmat[mp] = plus_i ? scale : scale*db_i;
2555 ++mp;
2556 w.walk(i, [&] (int jf, int jc, T aij, T aji, int st_j, T db_j) {
2557 if (st_j == C) {
2558 pbcol[bp] = jc;
2559 pbmat[bp] = bscale*aij;
2560 ++bp;
2561 } else if (usable_f(aji, st_j, db_j)) {
2562 pmcol[mp] = jf;
2563 pmmat[mp] = plus_i ? scale*aij/(db_j + aji) : scale*aij;
2564 ++mp;
2565 }
2566 });
2567 });
2569
2570 // The diagonal of M is emitted first and must be sorted into place.
2571 mcsr.sort();
2572 SpMatrix<T> M, B;
2573 M.define_split(A.partition(), A.partition(), std::move(mcsr), nlocal_f, s.col_map, 0);
2574 B.define_split(A.partition(), cpart, std::move(bcsr), nlocal_c, gcidx_r.data(), 0);
2575#ifndef AMREX_USE_GPU
2576 if (max_elmts > 0 || trunc_factor > T(0)) {
2577 // Each row is truncated as the product makes it.
2578 constexpr int KMAX = max_p_elmts;
2579 AMREX_ALWAYS_ASSERT(max_elmts <= KMAX);
2580 Vector<Long> n0(nrows, 0);
2581 auto P = amrex::SpGEMM(M, B, cpart, [&] (Long i, int* col, T* val, Long n) {
2582 n0[i] = n;
2583 return detail::truncate_row_host<T,KMAX>(col, val, n, max_elmts, trunc_factor,
2584 nlocal_c);
2585 });
2586 if (nnz_untruncated) {
2587 *nnz_untruncated = std::accumulate(n0.begin(), n0.end(), Long(0));
2588 }
2589 return P;
2590 }
2591#else
2592 amrex::ignore_unused(max_elmts, trunc_factor, nnz_untruncated);
2593#endif
2594 return amrex::SpGEMM(M, B, cpart);
2595}
2596
2597#ifndef AMREX_USE_GPU
2598template <typename T>
2600 SpMatrix<TS> const& ST,
2601 Gpu::DeviceVector<int> const& state,
2602 Gpu::DeviceVector<Long> const& cidx,
2603 AlgPartition const& cpart, T sscale, bool plus_i,
2604 int max_elmts, T trunc_factor,
2605 Long& nnz_untruncated)
2606{
2607 BL_PROFILE("AlgMG::build_interp_mm_host");
2608
2609 // Row i of P = M B (see build_interp_mm) is summed from the rows of B
2610 // in the order of the SpGEMM, then truncated as in truncate_interp.
2611 constexpr int KMAX = max_p_elmts;
2612 AMREX_ALWAYS_ASSERT(max_elmts <= KMAX);
2613
2614 auto const& s = S.const_parcsr();
2615 Long const nrows = s.csr0.nrows;
2616 auto const* AMREX_RESTRICT ps = state.data();
2617 auto const* AMREX_RESTRICT pc = cidx.data();
2618 int const nc = int(cpart.numLocalRows());
2619 int const C = PointType::coarse;
2620 int const F = PointType::fine;
2621
2622 auto const rsA = A.rowSum();
2623 auto const rsS = S.rowSum();
2624 auto const* AMREX_RESTRICT prsA = rsA.data();
2625 auto const* AMREX_RESTRICT prsS = rsS.data();
2626
2627 // D_beta and the rows of B at the F-points: the strong C neighbors.
2628 Gpu::DeviceVector<T> dbeta(nrows);
2629 auto* AMREX_RESTRICT pdb = dbeta.data();
2630 Vector<Long> boff(nrows+1, 0);
2631 auto const& sc = s.csr0;
2632 ParallelForOMP(nrows, [&] (Long i) noexcept
2633 {
2634 T d = T(0);
2635 Long n = 0;
2636 for (Long idx = sc.row_offset[i]; idx < sc.row_offset[i+1]; ++idx) {
2637 if (ps[sc.col_index[idx]] == C) {
2638 d += T(sc.mat[idx]);
2639 ++n;
2640 }
2641 }
2642 pdb[i] = d;
2643 boff[i+1] = (ps[i] == F) ? n : 0;
2644 });
2645 std::partial_sum(boff.begin(), boff.end(), boff.begin());
2646 Vector<int> bcol(boff[nrows]);
2647 Vector<T> bval(boff[nrows]);
2648 ParallelForOMP(nrows, [&] (Long i) noexcept
2649 {
2650 if (ps[i] != F) { return; }
2651 T const bscale = plus_i ? T(1) : ((pdb[i] != T(0)) ? T(1)/pdb[i] : T(0));
2652 Long q = boff[i];
2653 for (Long idx = sc.row_offset[i]; idx < sc.row_offset[i+1]; ++idx) {
2654 int const j = sc.col_index[idx];
2655 if (ps[j] == C) {
2656 bcol[q] = int(pc[j]);
2657 bval[q] = bscale*T(sc.mat[idx]);
2658 ++q;
2659 }
2660 }
2661 });
2662
2663 detail::MMInterpRow<T,TS> const w{s, ST.const_parcsr(), ps, pdb, pc,
2664 nullptr, nullptr, int(nrows), nc};
2665 bool const trunc = (max_elmts > 0 || trunc_factor > T(0));
2666 struct Nbr { int j; T aij; T den; };
2667
2669 auto csr = detail::csr_from_rows_cpu<T,SpMatrix<T>::template container_type,int>
2670 (nrows, [&] (Long i) { return sc.row_offset[i+1] - sc.row_offset[i] + 1; },
2671 [&] (int t) {
2672 return [&, t, marker = Vector<int>(nc, -1), col = Vector<int>(1), val = Vector<T>(1),
2673 tmp = Vector<std::pair<int,T>>(), nbr = Vector<Nbr>()]
2674 (Long i, int const*& rc, T const*& rv) mutable -> Long
2675 {
2676 rc = col.data();
2677 rv = val.data();
2678 if (ps[i] != F) {
2679 col[0] = int(pc[i]);
2680 val[0] = T(1);
2681 ++nnz0[t];
2682 return 1;
2683 }
2684
2685 // Scale of row i, as in build_interp_mm, and the usable strong
2686 // F neighbors (column order) for M.
2687 T alpha = prsA[i]/sscale - T(prsS[i]);
2688 T theta = T(0);
2689 Long maxlen = boff[i+1] - boff[i];
2690 nbr.clear();
2691 w.walk(i, [&] (int jf, int, T aij, T aji, int st_j, T db_j) {
2692 if (st_j == F) {
2693 T const den = db_j + (plus_i ? aji : T(0));
2694 if (den != T(0)) {
2695 if (plus_i) { theta += aij * aji / den; }
2696 maxlen += boff[jf+1] - boff[jf];
2697 nbr.push_back({jf, aij, den});
2698 } else {
2699 alpha += aij;
2700 }
2701 }
2702 });
2703 T const denom = alpha + theta;
2704 T const scale = (denom != T(0)) ? T(-1)/denom : T(0);
2705 maxlen = std::min(maxlen, Long(nc));
2706 if (Long(col.size()) < maxlen) {
2707 col.resize(maxlen);
2708 val.resize(maxlen);
2709 rc = col.data();
2710 rv = val.data();
2711 }
2712
2713 // M_ik times row k of B, k in increasing order (M's row sorted).
2714 Long n = 0;
2715 auto add = [&] (Long k, T m) {
2716 for (Long bp = boff[k]; bp < boff[k+1]; ++bp) {
2717 int const j = bcol[bp];
2718 Long const pos = marker[j];
2719 if (pos >= 0 && pos < n && col[pos] == j) {
2720 val[pos] += m * bval[bp];
2721 } else {
2722 marker[j] = int(n);
2723 col[n] = j;
2724 val[n] = m * bval[bp];
2725 ++n;
2726 }
2727 }
2728 };
2729 T const mdiag = plus_i ? scale : scale*pdb[i];
2730 bool diag_done = false;
2731 for (auto const& nb : nbr) {
2732 if (!diag_done && nb.j > i) {
2733 add(i, mdiag);
2734 diag_done = true;
2735 }
2736 add(nb.j, plus_i ? scale*nb.aij/nb.den : scale*nb.aij);
2737 }
2738 if (!diag_done) { add(i, mdiag); }
2739 detail::sort_row_cpu(col.data(), val.data(), n, tmp);
2740 nnz0[t] += n;
2741 if (!trunc) { return n; }
2742
2743 return detail::truncate_row_host<T,KMAX>(col.data(), val.data(), n, max_elmts,
2744 trunc_factor, nc);
2745 };
2746 });
2747
2748 nnz_untruncated = std::accumulate(nnz0.begin(), nnz0.end(), Long(0));
2749 SpMatrix<T> P;
2750 P.define_split(A.partition(), cpart, std::move(csr), nc, nullptr, 0);
2751 return P;
2752}
2753#endif
2754
2755}
2756
2757#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
Coalescing first-fit dynamic memory arena.
Enum reflection utilities and the AMREX_ENUM macro.
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
#define AMREX_RESTRICT
Definition AMReX_Extension.H:37
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
while parser_exe_t p
Definition AMReX_Parser_Exe_Body.H:33
Algebraic multigrid solver for SpMatrix / AlgVector systems.
Definition AMReX_AlgMG.H:264
SpMatrix< TS > create_soc(SpMatrix< T > const &A, T scale) const
Definition AMReX_AlgMG.H:1508
void setAggressiveNumLevels(int n)
Number of levels with aggressive coarsening, starting at the finest. Next setup.
Definition AMReX_AlgMG.H:338
void setFixedIter(int n)
If positive, run exactly this many V-cycles regardless of tolerance.
Definition AMReX_AlgMG.H:305
void setMaxIter(int n)
Definition AMReX_AlgMG.H:303
void setChebyshevDegree(int d)
Chebyshev polynomial degree. Next solve.
Definition AMReX_AlgMG.H:314
SpMatrix< T > const & getMatrix(int lev) const
Definition AMReX_AlgMG.H:398
AlgMGInterpType InterpType
Definition AMReX_AlgMG.H:272
AlgMGKrylovSolver KrylovSolver
Definition AMReX_AlgMG.H:271
T getResidualNorm() const
Definition AMReX_AlgMG.H:391
void setMaxCoarseSize(Long n)
Stop coarsening when a level has at most this many rows. Next setup.
Definition AMReX_AlgMG.H:336
void solve(AlgVector< T > &a_sol, AlgVector< T > const &a_rhs)
Solve A x = b. a_sol holds the initial guess on entry.
Definition AMReX_AlgMG.H:699
static SpMatrix< TS > second_pass_strength(SpMatrix< TS > &S, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart)
Definition AMReX_AlgMG.H:2160
void setRelTol(T t)
Definition AMReX_AlgMG.H:306
std::conditional_t< std::is_same_v< T, double >, float, T > TS
Definition AMReX_AlgMG.H:268
void setBottomMaxIter(int n)
Definition AMReX_AlgMG.H:384
void setPreSmooth(int nu)
Number of pre-smoothing sweeps; negative restores the smoother's default.
Definition AMReX_AlgMG.H:321
static constexpr int default_p_max_elmts
Definition AMReX_AlgMG.H:359
void setBottomSolver(BottomSolver bs)
Definition AMReX_AlgMG.H:376
void define(SpMatrix< T > &a_mat)
Define the solver with a square matrix. Must be called before solve.
Definition AMReX_AlgMG.H:642
void setPMaxElmts(int n)
Definition AMReX_AlgMG.H:353
void setAbsTol(T t)
Converged when the residual 2-norm is below max(reltol*bnorm, abstol).
Definition AMReX_AlgMG.H:308
bool coarsen_level(int lev, SpMatrix< T > &A, SpMatrix< T > &P, AlgPartition &cpart, bool &aggressive, Long &ncoarse, bool &truncated, Long &nnz_untruncated)
Definition AMReX_AlgMG.H:1346
void setSingular(bool b)
Definition AMReX_AlgMG.H:346
void setPrintIndentation(std::string s)
Prefix of the lines the solver prints, e.g. to nest them in MLMG output.
Definition AMReX_AlgMG.H:302
AlgMG()=default
void setPostSmooth(int nu)
Number of post-smoothing sweeps; negative restores the smoother's default.
Definition AMReX_AlgMG.H:323
void setInterpType(InterpType it)
Takes effect at the next setup.
Definition AMReX_AlgMG.H:350
SpMatrix< T > const & getInterp(int lev) const
Definition AMReX_AlgMG.H:399
AlgMGBottomSolver BottomSolver
Definition AMReX_AlgMG.H:270
void setBottomSmooth(int nu)
Definition AMReX_AlgMG.H:326
bool updateMatrix(SpMatrix< T > &a_mat)
Replace the matrix, keeping the hierarchy.
Definition AMReX_AlgMG.H:654
void setTruncFactor(T f)
Drop interpolation weights below this fraction of the row maximum. Next setup.
Definition AMReX_AlgMG.H:361
static AlgPartition coarse_numbering(Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > &cidx, int which=PointType::coarse)
Local index among the rows of type which and their partition.
Definition AMReX_AlgMG.H:1826
void setBottomVerbose(int v)
Definition AMReX_AlgMG.H:385
static SpMatrix< T > select_coarse_rows(SpMatrix< T > const &P, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart)
Rows of P at the C-points, in C-point numbering.
Definition AMReX_AlgMG.H:2282
static T strength_scale(SpMatrix< T > const &A)
Largest |a_ii|, the unit of the strength matrix entries.
Definition AMReX_AlgMG.H:1497
void setMaxLevels(int n)
Definition AMReX_AlgMG.H:334
static SpMatrix< T > truncate_interp(SpMatrix< T > &P, AlgPartition const &cpart, int max_elmts, T trunc_factor)
Definition AMReX_AlgMG.H:2300
int numLevels() const
Number of levels; 0 before setup.
Definition AMReX_AlgMG.H:393
Long numGlobalRows(int lev) const
Global number of rows on a level. Valid after setup.
Definition AMReX_AlgMG.H:397
static constexpr int max_p_elmts
Definition AMReX_AlgMG.H:358
void precond(AlgVector< T > &x, AlgVector< T > const &b)
Definition AMReX_AlgMG.H:787
void setStrongThreshold(T theta)
Definition AMReX_AlgMG.H:331
static SpMatrix< T > build_interp_mm(SpMatrix< T > &A, SpMatrix< TS > &S, SpMatrix< TS > const &ST, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart, T scale, bool plus_i, int const *row_state=nullptr, int max_elmts=0, T trunc_factor=T(0), Long *nnz_untruncated=nullptr)
Definition AMReX_AlgMG.H:2390
void setRelaxWeight(T w)
Relaxation weight of the smoother; negative restores the default. Next solve.
Definition AMReX_AlgMG.H:312
void setup()
Definition AMReX_AlgMG.H:1030
void setAggressiveDirectInterp(bool b)
Definition AMReX_AlgMG.H:341
static SpMatrix< T > build_interp(SpMatrix< T > &A, SpMatrix< TS > const &S, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart)
Classical direct interpolation.
Definition AMReX_AlgMG.H:1851
SpMatrix< T > interpolation(SpMatrix< T > &A, SpMatrix< TS > &S, SpMatrix< TS > const &ST, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart, T scale) const
Interpolation of the selected type from the C-points in state.
Definition AMReX_AlgMG.H:2146
void pmis(SpMatrix< TS > &S, SpMatrix< TS > &ST, Gpu::DeviceVector< int > &state, bool isolated_as_coarse=false) const
Definition AMReX_AlgMG.H:1682
std::size_t getSetupPeakBytes() const
Peak bytes in use in The_Arena during setup (0 if not available).
Definition AMReX_AlgMG.H:395
void setSmoother(Smoother sm)
Takes effect at the next solve; the hierarchy is kept.
Definition AMReX_AlgMG.H:310
void setThrowException(bool b)
Definition AMReX_AlgMG.H:386
void setVerbose(int v)
Definition AMReX_AlgMG.H:300
static void normalize_rows(SpMatrix< T > &P)
Scales the rows of P to sum to one; aborts on an empty row.
Definition AMReX_AlgMG.H:1468
int getNumSetups() const
Number of setups run so far (for checks that the setup is reused).
Definition AMReX_AlgMG.H:390
void project_out(AlgVector< T > &x)
Definition AMReX_AlgMG.H:564
void setBottomTol(T t)
Definition AMReX_AlgMG.H:383
AlgMGSmoother Smoother
Definition AMReX_AlgMG.H:273
void setKrylovSolver(KrylovSolver a)
Krylov solver with the V-cycle as preconditioner (see KrylovSolver). Default None.
Definition AMReX_AlgMG.H:363
void setChebyshevRatio(T r)
Chebyshev eigenvalue ratio lambda_max / lambda_min. Next solve.
Definition AMReX_AlgMG.H:316
int getNumIters() const
Definition AMReX_AlgMG.H:388
Definition AMReX_AlgPartition.H:26
Long numGlobalRows() const
Total number of rows covered by the partition.
Definition AMReX_AlgPartition.H:55
Long globalRowBegin() const
Inclusive global index begin on this process.
Definition AMReX_AlgPartition.H:62
Long numLocalRows() const
Number of local rows.
Definition AMReX_AlgPartition.H:50
Distributed dense vector that mirrors the layout of an AlgPartition.
Definition AMReX_AlgVector.H:29
void setVal(T val)
Definition AMReX_AlgVector.H:273
Long numLocalRows() const
Number of entries stored on this rank.
Definition AMReX_AlgVector.H:74
T const * data() const
Definition AMReX_AlgVector.H:85
void plusAsync(AlgVector< T, Allocator > const &rhs)
Definition AMReX_AlgVector.H:312
Chebyshev polynomial smoother for AlgVector/SpMatrix systems.
Definition AMReX_Smoother_MV.H:163
static std::size_t freeMemAvailable()
Definition AMReX_GpuDevice.cpp:1481
Jacobi smoother for AlgVector/SpMatrix linear systems.
Definition AMReX_Smoother_MV.H:32
l1 hybrid Gauss-Seidel smoother for CPUs.
Definition AMReX_Smoother_MV.H:318
Dynamically allocated vector for trivially copyable data.
Definition AMReX_PODVector.H:308
size_type size() const noexcept
Definition AMReX_PODVector.H:654
void resize(size_type a_new_size, GrowthStrategy strategy=GrowthStrategy::Poisson)
Definition AMReX_PODVector.H:734
T * data() noexcept
Definition AMReX_PODVector.H:672
This class provides the user with a few print options.
Definition AMReX_Print.H:35
Distributed CSR matrix that manages storage and GPU-friendly partitions.
Definition AMReX_SpMatrix.H:65
void setColumnPartition(AlgPartition const &col_partition)
Set the column partition and split the matrix into local and remote blocks with 32-bit column indices...
Definition AMReX_SpMatrix.H:1019
Long numGlobalRows() const
Global row count.
Definition AMReX_SpMatrix.H:196
Gpu::DeviceVector< U > gatherRemote(U const *x)
Gather values at the columns of the off-diagonal block.
Definition AMReX_SpMatrix.H:1342
Long numLocalNonZeros() const
Number of nonzeros stored locally.
Definition AMReX_SpMatrix.H:198
AlgVector< T, AllocT > rowSum() const
Sum the values in each local row and return the result as an AlgVector.
Definition AMReX_SpMatrix.H:1164
void define_split(AlgPartition partition, AlgPartition const &col_partition, local_csr_type &&compact, Long nlocal, Long const *remote_cols, Long nremote)
Define directly in split form from a CSR with compact 32-bit columns: c < nlocal is the local column ...
Definition AMReX_SpMatrix.H:1695
AlgPartition const & columnPartition() const
Return the column partition used for matrix-vector and matrix-matrix multiplications.
Definition AMReX_SpMatrix.H:191
ParCsr< T const > const_parcsr() const
Const-qualified alias of parcsr() for convenience.
Definition AMReX_SpMatrix.H:1226
ParCsr< T > parcsr()
Build GPU-friendly CSR views split into diagonal/off-diagonal blocks.
Definition AMReX_SpMatrix.H:1201
Long numLocalRows() const
Number of rows owned by this rank.
Definition AMReX_SpMatrix.H:194
AlgPartition const & partition() const
Row partition describing how matrix rows are distributed across ranks.
Definition AMReX_SpMatrix.H:182
AlgVector< T, AllocT > const & diagonalVector() const
Return diagonal elements in a square matrix.
Definition AMReX_SpMatrix.H:1145
CSR< T, container_type, int > local_csr_type
Definition AMReX_SpMatrix.H:71
This class is a thin wrapper around std::vector. Unlike vector, Vector::operator[] provides bound che...
Definition AMReX_Vector.H:29
amrex_long Long
Definition AMReX_INT.H:30
T Min(N n, T const *v, T init_val=std::numeric_limits< T >::max())
Compute the minimum of an array of values.
Definition AMReX_Reduce.H:1102
T Max(N n, T const *v, T init_val=std::numeric_limits< T >::lowest())
Compute the maximum of an array of values.
Definition AMReX_Reduce.H:1125
T Sum(N n, T const *v, T init_val=0)
Compute the sum of an array of values.
Definition AMReX_Reduce.H:1079
void ParallelForOMP(T n, L const &f) noexcept
Performance-portable kernel launch function with optional OpenMP threading.
Definition AMReX_GpuLaunch.H:328
PODVector< T, PinnedArenaAllocator< T > > PinnedVector
A PODVector that uses the pinned memory arena.
Definition AMReX_GpuContainers.H:48
PODVector< T, ArenaAllocator< T > > DeviceVector
A PODVector that uses the standard memory Arena. Note that the memory might or might not be managed d...
Definition AMReX_GpuContainers.H:24
Arena * The_Arena()
Definition AMReX_Arena.cpp:839
void Min(KeyValuePair< K, V > &vi, MPI_Comm comm)
Definition AMReX_ParallelReduce.H:161
void Sum(Gpu::DeviceVector< T > &v, MPI_Comm comm)
Definition AMReX_GpuParallelReduce.H:37
void AllGather(const T *v, int cnt, T *vs, MPI_Comm comm)
Definition AMReX_ParallelReduce.H:106
void Max(KeyValuePair< K, V > &vi, MPI_Comm comm)
Definition AMReX_ParallelReduce.H:133
#define AMREX_ENUM(CLASS,...)
Declare a scoped enum with reflection support.
Definition AMReX_Enum.H:270
void copy(HostToDevice, InIter begin, InIter end, OutIter result) noexcept
A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous st...
Definition AMReX_GpuContainers.H:128
static constexpr DeviceToHost deviceToHost
Definition AMReX_GpuContainers.H:106
static constexpr HostToDevice hostToDevice
Definition AMReX_GpuContainers.H:105
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:329
std::string const & name()
Definition AMReX_Machine.cpp:46
constexpr T powi(T x) noexcept
Return pow(x, Power), where Power is an integer known at compile time.
Definition AMReX_Math.H:377
constexpr int get_max_threads()
Definition AMReX_OpenMP.H:36
MPI_Comm CommunicatorSub() noexcept
sub-communicator for current frame
Definition AMReX_ParallelContext.H:70
int NProcsSub() noexcept
number of ranks in current frame
Definition AMReX_ParallelContext.H:74
static constexpr struct amrex::Scan::Type::Exclusive exclusive
static constexpr RetSum retSum
Definition AMReX_Scan.H:34
T PrefixSum(N n, FIN const &fin, FOUT const &fout, TYPE, RetSum a_ret_sum=retSum)
Definition AMReX_Scan.H:815
static constexpr int MPI_DATATYPE_NULL
Definition AMReX_ccse-mpi.H:56
int verbose
Definition AMReX.cpp:113
Definition AMReX_Amr.cpp:50
__host__ __device__ void ignore_unused(const Ts &...)
No-op helper that marks variables as intentionally unused.
Definition AMReX.H:273
AlgMGInterpType
Definition AMReX_AlgMG.H:220
@ mm_ext
Definition AMReX_AlgMG.H:220
@ mm_ext_i
Definition AMReX_AlgMG.H:220
@ C
Definition AMReX_Order.H:7
@ F
Definition AMReX_Order.H:7
@ none
Definition AMReX.H:558
void SpMV(Long nrows, Long ncols, T *__restrict__ py, CsrView< T const, Index > const &A, T const *__restrict__ px)
Perform y = A * x using CSR data (GPU/CPU aware).
Definition AMReX_SpMV.H:30
__host__ __device__ bool isnan(T x) noexcept
Return true if x is NaN. Works under fast math.
Definition AMReX_Math.H:117
AlgMGBottomSolver
Definition AMReX_AlgMG.H:210
@ jacobi
Definition AMReX_AlgMG.H:210
@ direct
Definition AMReX_AlgMG.H:210
@ gmres
Definition AMReX_AlgMG.H:210
@ bicgstab
Definition AMReX_AlgMG.H:210
AlgMGKrylovSolver
Definition AMReX_AlgMG.H:216
@ pcg
Definition AMReX_AlgMG.H:216
void computeResidual(AlgVector< T, AllocV > &res, SpMatrix< T, AllocM > const &A, AlgVector< T, AllocV > const &x, AlgVector< T, AllocV > const &b)
Compute the residual res = b - A * x.
Definition AMReX_SpMV.H:345
@ y
Definition AMReX_Orientation.H:14
@ x
Definition AMReX_Orientation.H:14
@ z
Definition AMReX_Orientation.H:14
@ smoother
Definition AMReX_MLLinOp.H:43
@ bicgstab
Definition AMReX_MLLinOp.H:43
__host__ __device__ constexpr IntVectND< dim > scale(const IntVectND< dim > &p, int s) noexcept
Returns a IntVectND obtained by multiplying each of the components of this IntVectND by s.
Definition AMReX_IntVect.H:1128
SpMatrix< T, Allocator > RAP(SpMatrix< T, Allocator > const &R, SpMatrix< T, Allocator > const &A, SpMatrix< T, Allocator > const &P, AlgPartition const &col_partition)
Galerkin product R (A P), with the result of SpGEMM(R, SpGEMM(A, P, col_partition),...
Definition AMReX_SpGEMM.H:1368
SpMatrix< T, Allocator > SpGEMM(SpMatrix< T, Allocator > const &A, SpMatrix< T, Allocator > const &B, AlgPartition const &col_partition, F const &row_post)
SpGEMM with a callback on each row of the product.
Definition AMReX_SpGEMM.H:1216
void Abort(const std::string &msg)
Print a fatal-error message to stderr and abort execution.
Definition AMReX.cpp:244
const int[]
Definition AMReX_BLProfiler.cpp:1665
@ bnorm
Definition AMReX_MLMG.H:23
CSR< T, V, IO > transpose(CSR< T, V, Index > const &csr, Long ncols)
Build the transpose CSR of csr.
Definition AMReX_SpMatUtil.H:25
AlgMGSmoother
Definition AMReX_AlgMG.H:229
@ chebyshev
Definition AMReX_AlgMG.H:229
@ l1_gauss_seidel
Definition AMReX_AlgMG.H:229
@ l1_jacobi
Definition AMReX_AlgMG.H:229
V< Index > col_index
Definition AMReX_CSR.H:62
void resize(Long num_rows, Long num_non_zeros)
Resize the storage to accommodate num_rows and num_non_zeros entries.
Definition AMReX_CSR.H:80
V< T > mat
Definition AMReX_CSR.H:61
V< Index > row_offset
Definition AMReX_CSR.H:63
Long nrows() const
Number of logical rows represented by the CSR offset array.
Definition AMReX_CSR.H:67
Definition AMReX_SpMatrix.H:39