3#include <AMReX_Config.H>
35template <
typename TC,
typename I,
typename T>
37bool pmis_blocked (CsrView<TC const,I>
const& csr,
Long row, T wi,
Long ig,
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))) {
56T pmis_tiebreak (
Long ig)
58 auto z =
static_cast<std::uint64_t
>(ig) + 0x9E3779B97F4A7C15ULL;
59 z = (
z ^ (
z >> 30)) * 0xBF58476D1CE4E5B9ULL;
60 z = (
z ^ (
z >> 27)) * 0x94D049BB133111EBULL;
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);
68template <
typename T,
typename I>
70bool csr_row_has_state (CsrView<T const,I>
const& csr,
Long row,
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; }
80inline std::pair<std::size_t,std::size_t> arena_bytes ()
82 if (
auto* carena =
dynamic_cast<CArena*
>(
The_Arena())) {
83 return {carena->heap_space_actually_used(), carena->heap_space_used()};
90template <
typename T,
typename TS>
104 template <
typename F>
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);
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];
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];
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);
150template <
typename T,
typename TS>
166 template <
typename F>
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]);
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];
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];
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]);
268 using TS = std::conditional_t<std::is_same_v<T,double>, float, 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; } }
302 if (m_cheby_ratio != r) { m_cheby_ratio = r; m_smoothers_stale =
true; }
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; } }
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; } }
331 if (m_singular != b) { m_singular = b; m_prepared =
false; }
340 if (m_p_max_elmts != p) { m_p_max_elmts = p; m_prepared =
false; }
345 void setTruncFactor (T f) {
if (m_trunc_factor != f) { m_trunc_factor = f; m_prepared =
false; } }
348 if (m_krylov == a) {
return; }
350 m_krylov_bicgstab.reset();
351 m_krylov_gmres.reset();
352 m_krylov_pcg.reset();
361 if (m_bottom_solver == bs) {
return; }
362 m_bottom_solver = bs;
387 enum PointType :
int { undecided = 0, fine = 1, coarse = 2 };
389 bool m_prepared =
false;
390 int m_num_setups = 0;
391 bool m_smoothers_stale =
false;
392 bool m_throw_exception =
false;
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);
404 int m_nu_bottom = -1;
406 [[nodiscard]]
int numPreSmooth ()
const {
407 return (m_nu1 >= 0) ? m_nu1 : ((m_smoother_type == Smoother::chebyshev) ? 1 : 2);
409 [[nodiscard]]
int numPostSmooth ()
const {
410 return (m_nu2 >= 0) ? m_nu2 : ((m_smoother_type == Smoother::chebyshev) ? 1 : 2);
412 [[nodiscard]]
int numBottomSmooth ()
const {
413 return (m_nu_bottom >= 0) ? m_nu_bottom : ((m_smoother_type == Smoother::chebyshev) ? 4 : 8);
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;
422 T m_strong_threshold = T(0.25);
424 InterpType m_interp_type = InterpType::mm_ext_i;
426 T m_trunc_factor = T(0);
431 static constexpr Long max_direct_size = 1024;
432 Long m_bottom_n = -1;
433 Vector<double> m_bottom_lu;
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;
439 T m_reltol = std::is_same_v<T,float> ? T(1.e-4) : T(1.e-8);
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();
445 std::string m_print_ident;
446 std::size_t m_setup_peak_bytes = 0;
449 void note_memory (
int lev,
char const* stage)
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"
464 SpMatrix<T>* m_mat =
nullptr;
465 Vector<SpMatrix<T>*> m_a;
466 Vector<std::unique_ptr<SpMatrix<T>>> m_crsemat;
467 Vector<SpMatrix<T>> m_P;
468 Vector<SpMatrix<T>> m_R;
470 Vector<AlgVector<T>> m_res;
471 Vector<AlgVector<T>> m_cor;
472 Vector<AlgVector<T>> m_rescor;
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;
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;
485 void clear_levels ();
486 void build_smoothers ();
488 void build_direct_bottom ();
489 void direct_bottom_solve ();
493 if (!m_prepared) {
setup(); }
494 else if (m_smoothers_stale) { build_smoothers(); }
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 ();
504 static std::pair<T,T> compute_rbnorms (SpMatrix<T>
const& mat,
505 AlgVector<T>
const& xvec,
506 AlgVector<T>
const& bvec,
509 static T compute_rnorm (SpMatrix<T>
const& mat, AlgVector<T>
const& xvec,
510 AlgVector<T>
const& bvec, AlgVector<T>& rvec);
523 bool& aggressive,
Long& ncoarse,
bool& truncated,
524 Long& nnz_untruncated);
536 bool isolated_as_coarse =
false)
const;
541 int which = PointType::coarse);
547 if (!m_singular) {
return; }
548 T
const mean =
x.sum() / T(
x.partition().numGlobalRows());
564 int max_elmts, T trunc_factor);
598 bool plus_i,
int const* row_state =
nullptr,
599 int max_elmts = 0, T trunc_factor = T(0),
600 Long* nnz_untruncated =
nullptr);
611 int max_elmts, T trunc_factor,
612 Long& nnz_untruncated);
629 m_krylov_bicgstab.reset();
630 m_krylov_gmres.reset();
631 m_krylov_pcg.reset();
651 m_bottom_piv.clear();
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();
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;
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";
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";
696 if (m_krylov != KrylovSolver::none) {
697 solve_krylov(a_sol, rhs, res_target);
702 int const niters = (m_fixediter > 0) ? m_fixediter : m_maxiter;
703 for (
int iter = 0; iter < niters; ++iter)
705 vcycle(m_cor[0], m_res[0]);
708 m_rnorm = compute_rnorm(*m_mat, a_sol, rhs, m_res[0]);
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";
717 bool converged = (m_fixediter <= 0) && (m_rnorm <= res_target);
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";
725 }
else if (m_rnorm > T(1.e20) * m_init_rnorm || amrex::isnan(m_rnorm)) {
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";
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";
753 if (m_throw_exception) {
754 throw std::runtime_error(
"AlgMG blew up.");
761void AlgMG<T>::solve_krylov (AlgVector<T>& a_sol, AlgVector<T>
const& a_rhs, T res_target)
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);
774 auto run = [&] (
auto&
s) {
775 s.setVerbose(verbose);
776 s.setInitialGuessNonzero(
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)) {
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";
795 char const*
name =
"";
796 if (m_krylov == KrylovSolver::bicgstab) {
798 if (!m_krylov_bicgstab) {
799 m_krylov_bicgstab = std::make_unique<BiCGStab_MV<T>>(m_mat);
801 m_krylov_bicgstab->setPrecond(pc);
802 run(m_krylov_bicgstab->getSolver());
803 }
else if (m_krylov == KrylovSolver::gmres) {
805 if (!m_krylov_gmres) {
806 m_krylov_gmres = std::make_unique<GMRES_MV<T>>(m_mat);
808 m_krylov_gmres->setPrecond(pc);
809 run(m_krylov_gmres->getSolver());
813 m_krylov_pcg = std::make_unique<PCG_MV<T>>(m_mat);
815 m_krylov_pcg->setPrecond(pc);
816 run(m_krylov_pcg->getSolver());
819 if (amrex::isnan(m_rnorm) || m_rnorm > T(1.e20) * m_init_rnorm) { blew_up(); }
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 <<
")";
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";
836std::pair<T,T> AlgMG<T>::compute_rbnorms (SpMatrix<T>
const& mat,
837 AlgVector<T>
const& xvec,
838 AlgVector<T>
const& bvec,
841 auto bnorm = bvec.norm2(
true);
844 auto rnorm = rvec.norm2(
true);
848 rnorm = Math::powi<2>(rnorm);
851 rnorm = std::sqrt(rnorm);
854 return {
bnorm, rnorm};
858T AlgMG<T>::compute_rnorm (SpMatrix<T>
const& mat, AlgVector<T>
const& xvec,
859 AlgVector<T>
const& bvec, AlgVector<T>& rvec)
866void AlgMG<T>::vcycle (AlgVector<T>&
x, AlgVector<T>
const& b)
870 int const nlev = numLevels();
871 int const bottom = nlev-1;
874 if (&b != &m_res[0]) { m_res[0].copyAsync(b); }
876 if (&
x != &m_cor[0]) {
x.copyAsync(m_cor[0]); }
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]; };
884 smooth(0,
x, b, numPreSmooth(),
false);
886 for (
int lev = 0; lev < bottom; ++lev) {
888 SpMV(m_res[lev+1], m_R[lev], m_rescor[lev]);
889 if (lev+1 == bottom) {
892 smooth(lev+1, m_cor[lev+1], m_res[lev+1], numPreSmooth(),
false);
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);
904void AlgMG<T>::smooth (
int lev, AlgVector<T>& xvec, AlgVector<T>
const& bvec,
int nu,
905 bool with_initial_guess,
bool backward)
907 if (m_smoother_type == Smoother::l1_gauss_seidel) {
910 smoother(xvec, bvec, with_initial_guess, backward);
911 }
else if (m_smoother_type == Smoother::chebyshev) {
914 smoother(xvec, bvec, with_initial_guess);
918 smoother(xvec, bvec, with_initial_guess);
923void AlgMG<T>::bottom_solve ()
927 int const lev = numLevels()-1;
928 project_out(m_res[lev]);
929 auto smooth_bottom = [&] () {
930 if (m_smoother_type == Smoother::l1_gauss_seidel) {
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);
936 smooth(lev, m_cor[lev], m_res[lev], numBottomSmooth(),
false);
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";
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; }
960 switch (bottom_solver) {
961 case BottomSolver::direct:
962 direct_bottom_solve();
964 case BottomSolver::jacobi:
967 case BottomSolver::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));
972 krylov_bottom(*m_gmres);
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));
979 krylov_bottom(*m_bicgstab);
982 project_out(m_cor[lev]);
992 m_a.push_back(m_mat);
993 m_mat->setColumnPartition(m_mat->partition());
994 m_setup_peak_bytes = 0;
995 note_memory(0,
"start");
999 for (
int lev = 0; lev+1 < m_max_levels; ++lev)
1001 auto& A = *m_a[lev];
1002 if (A.numGlobalRows() <= m_max_coarse_size) {
break; }
1006 bool aggressive =
false;
1008 bool truncated =
false;
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)) {
1015 P = truncate_interp(P, cpart, m_p_max_elmts, m_trunc_factor);
1016 note_memory(lev,
"truncation");
1018 if (m_verbose >= 2) {
1021 amrex::Print() << m_print_ident <<
"AlgMG: level " << lev
1022 <<
": P truncated from " << nnz[0] <<
" to "
1023 << nnz[1] <<
" nonzeros\n";
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 <<
" -> "
1034 if (m_verbose >= 3) {
1048 amrex::Print() << m_print_ident <<
"AlgMG: level " << lev <<
": P row sums in ["
1049 << rmin <<
", " << rmax <<
"], " << nzero <<
" empty rows\n";
1052 note_memory(lev,
"R");
1054#ifndef AMREX_USE_GPU
1061 note_memory(lev,
"A*P");
1065 note_memory(lev,
"coarse operator");
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());
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());
1078 m_rescor.emplace_back(m_a[lev]->partition());
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; }
1090 amrex::Print() << m_print_ident <<
"AlgMG: level " << lev <<
": "
1091 << m_a[lev]->numGlobalRows() <<
" rows, "
1092 << nnz <<
" nonzeros\n";
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())
1109template <
typename T>
1113 int const lev = numLevels()-1;
1114 auto const& A = *m_a[lev];
1115 Long const n = A.numGlobalRows();
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";
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);
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]);
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);
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]);
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; }
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;
1183 std::iota(piv.begin(), piv.end(), 0);
1184 for (
Long i = 0; i < n; ++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; }
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";
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]); }
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;
1206 for (
Long k = i+1; k < n; ++k) { dense[j*n+k] -= f*dense[i*n+k]; }
1209 m_bottom_lu = std::move(dense);
1210 m_bottom_piv = std::move(piv);
1214template <
typename T>
1215void AlgMG<T>::direct_bottom_solve ()
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();
1225 Vector<double> rhs(n, 0.0);
1229 for (
Long i = 0; i < nl; ++i) { rhs[b0+i] = double(loc[i]); }
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]; }
1241 for (
Long i = n-1; i >= 0; --i) {
1243 for (
Long k = i+1; k < n; ++k) { v -= lu[i*n+k]*
y[k]; }
1248 for (
Long i = 0; i < nl; ++i) { loc[i] = T(
y[b0+i]); }
1253template <
typename T>
1254void AlgMG<T>::build_smoothers ()
1261 int const nlev = numLevels();
1262 for (
int lev = 0; lev < nlev; ++lev) {
1263 if (m_smoother_type == Smoother::l1_gauss_seidel) {
1265 amrex::Abort(
"AlgMG: L1GaussSeidel smoother is only available in CPU builds");
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);
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);
1279 m_smoothers_stale =
false;
1282template <
typename T>
1284 bool& aggressive,
Long& ncoarse,
bool& truncated,
1285 Long& nnz_untruncated)
1290 T
const sscale = strength_scale(A);
1291 auto S = create_soc(A, sscale);
1292 S.setColumnPartition(S.partition());
1293 note_memory(lev,
"S");
1295 ST.setColumnPartition(ST.partition());
1297 note_memory(lev,
"strength");
1300 pmis(S, ST, state, m_singular);
1301 note_memory(lev,
"PMIS");
1304 cpart = coarse_numbering(state, cidx);
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);
1312 if (lev < m_aggressive_num_levels) {
1314 auto S2 = second_pass_strength(S, state, cidx, cpart);
1315 S2.setColumnPartition(cpart);
1317 S2T.setColumnPartition(cpart);
1318 note_memory(lev,
"second-pass strength");
1320 pmis(S2, S2T, state2,
true);
1322 auto cpart2 = coarse_numbering(state2, cidx2);
1323 Long const nc2 = cpart2.numGlobalRows();
1324 if (nc2 > 0 && nc2 < ncoarse) {
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();
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);
1344 bool const trunc = (m_p_max_elmts > 0 || m_trunc_factor > T(0));
1346#ifndef AMREX_USE_GPU
1347 if (!m_aggressive_direct_interp && host_interp) {
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);
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); }
1363 note_memory(lev,
"P1");
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);
1370 auto Pf = build_interp_mm(A, S, ST, statef, cidxf, cpart2, sscale,
1371 m_interp_type != InterpType::mm_ext,
1373 if (trunc) { Pf = truncate_interp(Pf, cpart2, m_p_max_elmts, m_trunc_factor); }
1375 auto P2 = select_coarse_rows(Pf, state, cidx, cpart);
1376 note_memory(lev,
"P2");
1383#ifndef AMREX_USE_GPU
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);
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);
1397 P = interpolation(A, S, ST, state, cidx, cpart, sscale);
1400 note_memory(lev,
"interpolation");
1404template <
typename T>
1407 auto const rs = P.
rowSum();
1409 Long nzero = Reduce::Sum<Long>(rs.numLocalRows(),
1413 amrex::Abort(
"AlgMG: singular mode needs P 1 = 1, but " + std::to_string(nzero)
1414 +
" rows of P are empty");
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;
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;
1433template <
typename T>
1438 T dmax = Reduce::Max<T>(Ad.numLocalRows(),
1441 return (dmax > T(0)) ? dmax : T(1);
1444template <
typename T>
1450 Long const nrows = a.csr0.nrows;
1455 auto const threshold = m_strong_threshold;
1461 T
const tol = T(64) * std::numeric_limits<T>::epsilon();
1463#ifndef AMREX_USE_GPU
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];
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); },
1477 (
Long i,
int const*& rc,
TS const*& rv)
mutable ->
Long
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]);
1485 T amin = std::numeric_limits<T>::max();
1486 for (
Long idx = b; idx < e; ++idx) {
1488 if (c0.col_index[idx] != i) { amin = std::min(amin, sgn*c0.mat[idx]); }
1490 for (
Long idx = b1; idx < b1+n1; ++idx) {
1492 amin = std::min(amin, sgn*c1.mat[idx]);
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) {
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);
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);
1527 auto const rsA = A.
rowSum();
1530 if (std::abs(prs[i]) > T(0.9)*std::abs(A_ii[i])) {
1531 return std::numeric_limits<T>::lowest();
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]);
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]);
1548 if (amin >= T(0)) {
return std::numeric_limits<T>::lowest(); }
1549 return amin * threshold * (T(1) - tol);
1559 Long const nnz = Scan::PrefixSum<Long>
1562 if (i >= nrows) {
return 0; }
1563 T
const thr = pthr[i];
1564 T
const sgn = std::copysign(T(1), A_ii[i]);
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; }
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; }
1582 csr.resize(nrows, nnz);
1587 T
const thr = pthr[i];
1588 T
const sgn = std::copysign(T(1), A_ii[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];
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];
1618template <
typename T>
1620 bool isolated_as_coarse)
const
1625 int const U = PointType::undecided;
1626 int const C = PointType::coarse;
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];
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];
1646 pm[i] = (n == 0 && nout == 0) ?
Long(-1) : n;
1656 Long const row_begin =
s.row_begin;
1664 ps[i] = isolated_as_coarse ? PointType::coarse : PointType::fine;
1666 pw[i] = T(pm[i]) + detail::pmis_tiebreak<T>(i + row_begin);
1667 ps[i] = PointType::undecided;
1691 Long const ig = i + row_begin;
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);
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);
1703 if (!blocked) { st_i =
C; }
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;
1726 return (ps[i] == U) ? 1 : 0;
1729 if (nund == 0) {
break; }
1732 if (m_verbose >= 2) {
1736 return (ps[i] ==
C) ? 1 : 0;
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;
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;
1754 Long counts[3] = {nc, ncc, nf0};
1756 amrex::Print() << m_print_ident <<
"AlgMG: PMIS: " << counts[0] <<
" C-points of "
1758 <<
", F-points w/o C neighbor " << counts[2] <<
"\n";
1762template <
typename T>
1770 Long const nc = Scan::PrefixSum<Long>
1783 std::partial_sum(counts.begin(), counts.end(), rows.begin()+1);
1787template <
typename T>
1798 Long const nrows = a.csr0.nrows;
1812 detail::DirectInterpRow<T,TS>
const w{a, S.
const_parcsr(), ps, pc,
1813 state_r.data(), nlocal_c,
1814 int(PointType::coarse)};
1821 Long const nnz = Scan::PrefixSum<Long>
1824 if (i >= nrows) {
return 0; }
1825 if (ps[i] == PointType::coarse) {
return 1; }
1828 w.walk(i, diag, [&] (
int, T,
bool strong_c) {
if (strong_c) { ++n; } });
1834 csr.resize(nrows, nnz);
1841 if (ps[i] == PointType::coarse) {
1842 pcol[p] =
int(pc[i]);
1847 T sum_n_neg = T(0), sum_n_pos = T(0), sum_p_neg = T(0), sum_p_pos = T(0);
1849 w.walk(i, diag, [&] (
int, T v,
bool strong_c) {
1852 if (strong_c) { sum_p_neg += v; }
1855 if (strong_c) { sum_p_pos += v; }
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); }
1865 w.walk(i, unused, [&] (
int cj, T v,
bool strong_c) {
1868 pmat[p] = ((v < T(0)) ? -alfa*v : -
beta*v) / diag;
1885template <
typename T,
int KMAX>
1894 template <
typename F>
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]);
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]);
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; }
1920 return trunc_factor * amax;
1927 walk(i, [&] (
int, T v) {
if (std::abs(v) >= thresh) { ++n; } });
1935 int select (
Long i, T thresh,
int* cols, T* vals)
const
1937 int const kmax = std::min(max_elmts, KMAX);
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];
1951 if (n < kmax) { ++n; }
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)
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};
1969 T
const thresh = tr.threshold(0, sum_pos, sum_neg);
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]; }
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]; }
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) {
1990 for (
Long k = 0; k < nk; ++k) {
1992 for (; m > 0 && col[m-1] > tcol[k]; --m) {
1997 val[m] = tval[k] * ((tval[k] > T(0)) ? fpos : fneg);
2000 for (
Long k = 0; k < n; ++k) {
2001 if (std::abs(val[k]) >= thresh) {
2003 val[nk] = val[k] * ((val[k] > T(0)) ? fpos : fneg);
2020template <
typename T,
typename F>
2025 Long const nrows = p.csr0.nrows;
2028 Gpu::DeviceVector<Long> off(nrows+1);
2030 Long const nnz = Scan::PrefixSum<Long>
2033 if (i >= nrows || state[i] != c) {
return 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; }
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; }
2058 if (i == nrows) { prow[nc] = nnz;
return; }
2059 if (state[i] != c) {
return; }
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; }
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; }
2082template <
typename T>
2088 if (m_interp_type == InterpType::direct) {
2089 return build_interp(A, S, state, cidx, cpart);
2091 return build_interp_mm(A, S, ST, state, cidx, cpart,
scale,
2092 m_interp_type == InterpType::mm_ext_i);
2096template <
typename T>
2108 Long const nrows =
s.csr0.nrows;
2110 int const C = PointType::coarse;
2111 int const F = PointType::fine;
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);
2123 Long const fbegin = fpart.globalRowBegin();
2130 pgc[i] = pc[i] + cbegin;
2131 pgf[i] = pf[i] + fbegin;
2139 auto csr_cf = detail::compact_rows<TS>
2143 return (ps_r[j] ==
F) ? nlocal_f + j : -1;
2145 return (ps[j] ==
F) ?
int(pf[j]) : -1;
2149 SCF.
define_split(cpart, fpart, std::move(csr_cf), nlocal_f, gfidx_r.data(), 0);
2152 auto csr_fc = detail::compact_rows<TS>
2156 return (ps_r[j] ==
C) ? nlocal_c + j : -1;
2158 return (ps[j] ==
C) ?
int(pc[j]) : -1;
2162 SFC.
define_split(fpart, cpart, std::move(csr_fc), nlocal_c, gcidx_r.data(), 0);
2167 auto const& s2 = S2.const_parcsr();
2171 Long const nnz = Scan::PrefixSum<Long>
2174 if (r >= nc) {
return 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));
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];
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];
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];
2214 S2d.
define_split(cpart, cpart, std::move(out), nlocal_c, s2.col_map, 0);
2218template <
typename T>
2226 auto csr = detail::compact_rows<T>
2229 return remote ? nlocal + j : j;
2236template <
typename T>
2238 int max_elmts, T trunc_factor)
2242 constexpr int KMAX = max_p_elmts;
2247 Long const nrows = p.csr0.nrows;
2249 detail::TruncRow<T,KMAX>
const tr{p, max_elmts, trunc_factor, nlocal};
2258 pthr[i] = tr.threshold(i, ppos[i], pneg[i]);
2265 Long const nnz = Scan::PrefixSum<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;
2275 csr.resize(nrows, nnz);
2280 T
const thresh = pthr[i];
2281 T
const sum_pos = ppos[i];
2282 T
const sum_neg = pneg[i];
2284 if (max_elmts > 0) {
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]; }
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) {
2296 pmat[q] = vals[k] * ((vals[k] > T(0)) ? fpos : fneg);
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; }
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) {
2311 pmat[q] = v * ((v > T(0)) ? fpos : fneg);
2326template <
typename T>
2332 bool plus_i,
int const* row_state,
2333 int max_elmts, T trunc_factor,
Long* nnz_untruncated)
2342 Long const nrows =
s.csr0.nrows;
2348 int const C = PointType::coarse;
2349 int const F = PointType::fine;
2353 auto const rsA = A.
rowSum();
2354 auto const rsS = S.
rowSum();
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]); }
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]); }
2387 detail::MMInterpRow<T,TS>
const w{
s, ST.
const_parcsr(), ps, pdb, pc,
2388 pstate_r, dbeta_r.data(), nlocal_f, nlocal_c};
2393 local_csr_type mcsr, bcsr;
2395 bcsr.row_offset.resize(nrows+1);
2400 return st_j ==
F && (db_j + (plus_i ? aji : T(0))) != T(0);
2407 return prow_state ==
nullptr || prow_state[i] ==
C;
2410 Long const mnnz = Scan::PrefixSum<Long>
2413 if (i >= nrows || !has_row(i)) {
return 0; }
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; }
2425 Long const bnnz = Scan::PrefixSum<Long>
2428 if (i >= nrows) {
return 0; }
2429 if (ps[i] ==
C) {
return 1; }
2431 w.walk(i, [&] (
int,
int, T, T,
int st_j, T) {
2432 if (st_j ==
C) { ++n; }
2439 mcsr.resize(nrows, mnnz);
2440 bcsr.resize(nrows, bnnz);
2450 bool const in_p = has_row(i);
2456 pbcol[bp] =
int(pc[i]);
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));
2463 w.walk(i, [&] (
int,
int jc, T aij, T,
int st_j, T) {
2466 pbmat[bp] = bscale*aij;
2475 T alpha = prsA[i]/sscale - T(prsS[i]);
2477 w.walk(i, [&] (
int,
int, T aij, T aji,
int st_j, T db_j) {
2479 T
const den = db_j + (plus_i ? aji : T(0));
2481 if (plus_i) { theta += aij * aji / den; }
2487 T
const denom = alpha + theta;
2488 T
const scale = (denom != T(0)) ? T(-1)/denom : T(0);
2493 w.walk(i, [&] (
int jf,
int jc, T aij, T aji,
int st_j, T db_j) {
2496 pbmat[bp] = bscale*aij;
2498 }
else if (usable_f(aji, st_j, db_j)) {
2500 pmmat[mp] = plus_i ?
scale*aij/(db_j + aji) :
scale*aij;
2512#ifndef AMREX_USE_GPU
2513 if (max_elmts > 0 || trunc_factor > T(0)) {
2515 constexpr int KMAX = max_p_elmts;
2520 return detail::truncate_row_host<T,KMAX>(col, val, n, max_elmts, trunc_factor,
2523 if (nnz_untruncated) {
2524 *nnz_untruncated = std::accumulate(n0.begin(), n0.end(),
Long(0));
2534#ifndef AMREX_USE_GPU
2535template <
typename T>
2541 int max_elmts, T trunc_factor,
2542 Long& nnz_untruncated)
2548 constexpr int KMAX = max_p_elmts;
2552 Long const nrows =
s.csr0.nrows;
2556 int const C = PointType::coarse;
2557 int const F = PointType::fine;
2559 auto const rsA = A.
rowSum();
2560 auto const rsS = S.
rowSum();
2568 auto const& sc =
s.csr0;
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]);
2580 boff[i+1] = (ps[i] ==
F) ? n : 0;
2582 std::partial_sum(boff.begin(), boff.end(), boff.begin());
2583 Vector<int> bcol(boff[nrows]);
2584 Vector<T> bval(boff[nrows]);
2587 if (ps[i] !=
F) {
return; }
2588 T
const bscale = plus_i ? T(1) : ((pdb[i] != T(0)) ? T(1)/pdb[i] : T(0));
2590 for (
Long idx = sc.row_offset[i]; idx < sc.row_offset[i+1]; ++idx) {
2591 int const j = sc.col_index[idx];
2593 bcol[q] =
int(pc[j]);
2594 bval[q] = bscale*T(sc.mat[idx]);
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; };
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; },
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
2616 col[0] =
int(pc[i]);
2624 T alpha = prsA[i]/sscale - T(prsS[i]);
2626 Long maxlen = boff[i+1] - boff[i];
2628 w.walk(i, [&] (
int jf,
int, T aij, T aji,
int st_j, T db_j) {
2630 T
const den = db_j + (plus_i ? aji : T(0));
2632 if (plus_i) { theta += aij * aji / den; }
2633 maxlen += boff[jf+1] - boff[jf];
2634 nbr.push_back({jf, aij, den});
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) {
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];
2661 val[n] = m * bval[bp];
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) {
2673 add(nb.j, plus_i ?
scale*nb.aij/nb.den :
scale*nb.aij);
2675 if (!diag_done) { add(i, mdiag); }
2676 detail::sort_row_cpu(col.data(), val.data(), n, tmp);
2678 if (!trunc) {
return n; }
2680 return detail::truncate_row_host<T,KMAX>(col.data(), val.data(), n, max_elmts,
2685 nnz_untruncated = std::accumulate(nnz0.begin(), nnz0.end(),
Long(0));
2687 P.define_split(A.
partition(), cpart, std::move(csr), nc,
nullptr, 0);
#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
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