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