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