3#include <AMReX_Config.H>
35template <
typename TC,
typename Index,
typename T>
37bool pmis_blocked (CsrView<TC const,Index>
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 Index>
70bool csr_row_has_state (CsrView<T const,Index>
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 ();
492 static void allgather_rows (
double* buf, AlgPartition
const& part,
Long stride);
496 if (!m_prepared) {
setup(); }
497 else if (m_smoothers_stale) { build_smoothers(); }
500 void vcycle (AlgVector<T>&
x, AlgVector<T>
const& b);
501 void blew_up ()
const;
502 void solve_krylov (AlgVector<T>& a_sol, AlgVector<T>
const& a_rhs, T res_target);
503 void smooth (
int lev, AlgVector<T>& xvec, AlgVector<T>
const& bvec,
int nu,
504 bool with_initial_guess,
bool backward =
false);
505 void bottom_solve ();
507 static std::pair<T,T> compute_rbnorms (SpMatrix<T>
const& mat,
508 AlgVector<T>
const& xvec,
509 AlgVector<T>
const& bvec,
512 static T compute_rnorm (SpMatrix<T>
const& mat, AlgVector<T>
const& xvec,
513 AlgVector<T>
const& bvec, AlgVector<T>& rvec);
526 bool& aggressive,
Long& ncoarse,
bool& truncated,
527 Long& nnz_untruncated);
539 bool isolated_as_coarse =
false)
const;
544 int which = PointType::coarse);
550 if (!m_singular) {
return; }
551 T
const mean =
x.sum() / T(
x.partition().numGlobalRows());
567 int max_elmts, T trunc_factor);
601 bool plus_i,
int const* row_state =
nullptr,
602 int max_elmts = 0, T trunc_factor = T(0),
603 Long* nnz_untruncated =
nullptr);
614 int max_elmts, T trunc_factor,
615 Long& nnz_untruncated);
632 m_krylov_bicgstab.reset();
633 m_krylov_gmres.reset();
634 m_krylov_pcg.reset();
638void AlgMG<T>::clear_levels ()
654 m_bottom_piv.clear();
667 if (!m_bproj) { m_bproj = std::make_unique<AlgVector<T>>(m_mat->partition()); }
668 m_bproj->copy(a_rhs);
670 prhs = m_bproj.get();
674 auto norms = compute_rbnorms(*m_mat, a_sol, rhs, m_res[0]);
675 m_bnorm = norms.first;
676 m_init_rnorm = m_rnorm = norms.second;
685 if (m_verbose >= 1) {
686 amrex::Print() << m_print_ident <<
"AlgMG: Initial rhs 2-norm = " << m_bnorm <<
"\n"
687 << m_print_ident <<
"AlgMG: Initial residual 2-norm = " << m_rnorm <<
"\n";
690 T res_target = std::max(m_bnorm * m_reltol, m_abstol);
691 if (m_fixediter <= 0 && m_rnorm <= res_target) {
692 if (m_verbose >= 1) {
693 amrex::Print() << m_print_ident <<
"AlgMG: No iterations needed\n";
699 if (m_krylov != KrylovSolver::none) {
700 solve_krylov(a_sol, rhs, res_target);
705 int const niters = (m_fixediter > 0) ? m_fixediter : m_maxiter;
706 for (
int iter = 0; iter < niters; ++iter)
708 vcycle(m_cor[0], m_res[0]);
711 m_rnorm = compute_rnorm(*m_mat, a_sol, rhs, m_res[0]);
714 if (m_verbose >= 2) {
715 amrex::Print() << m_print_ident <<
"AlgMG: Iteration " << std::setw(3)
716 << iter+1 <<
": rnorm, rnorm/bnorm = "
717 << m_rnorm <<
", " << m_rnorm/m_bnorm <<
"\n";
720 bool converged = (m_fixediter <= 0) && (m_rnorm <= res_target);
722 if (m_verbose >= 1) {
723 amrex::Print() << m_print_ident <<
"AlgMG: Final Iter. " << iter+1
724 <<
": rnorm, rnorm/bnorm = "
725 << m_rnorm <<
", " << m_rnorm/m_bnorm <<
"\n";
728 }
else if (m_rnorm > T(1.e20) * m_init_rnorm ||
amrex::isnan(m_rnorm)) {
730 amrex::Print() << m_print_ident <<
"AlgMG: Failing to converge after " << iter+1 <<
" iterations."
731 <<
" rnorm, rnorm/bnorm = "
732 << m_rnorm <<
", " << m_rnorm/m_bnorm <<
"\n";
738 if (m_verbose >= 1 && m_fixediter <= 0 && m_rnorm > res_target) {
739 amrex::Print() << m_print_ident <<
"AlgMG: Not converged after " << m_niters
740 <<
" iterations. rnorm/bnorm = " << m_rnorm/m_bnorm <<
"\n";
754void AlgMG<T>::blew_up ()
const
756 if (m_throw_exception) {
757 throw std::runtime_error(
"AlgMG blew up.");
770 int const niters = (m_fixediter > 0) ? m_fixediter : m_maxiter;
771 T
const atol = (m_fixediter > 0) ? T(0) : res_target;
772 int const verbose = std::max(0, m_verbose-1);
773 T krylov_rnorm = T(0);
777 auto run = [&] (
auto& s) {
778 s.setVerbose(verbose);
779 s.setInitialGuessNonzero(
true);
782 s.setMaxIters(niters - m_niters);
783 s.solve(a_sol, a_rhs, T(0), atol);
784 m_niters += s.getNumIters();
785 krylov_rnorm = s.getResidualNorm();
786 m_rnorm = compute_rnorm(*m_mat, a_sol, a_rhs, m_res[0]);
787 if (m_fixediter > 0 || m_rnorm <= res_target || krylov_rnorm > res_target ||
788 m_niters >= niters || s.getNumIters() == 0 ||
amrex::isnan(m_rnorm)) {
791 if (m_verbose >= 2) {
792 amrex::Print() << m_print_ident <<
"AlgMG: restarting Krylov solver at iteration "
793 << m_niters <<
", rnorm/bnorm = " << m_rnorm/m_bnorm <<
"\n";
798 char const*
name =
"";
799 if (m_krylov == KrylovSolver::bicgstab) {
801 if (!m_krylov_bicgstab) {
802 m_krylov_bicgstab = std::make_unique<BiCGStab_MV<T>>(m_mat);
804 m_krylov_bicgstab->setPrecond(pc);
805 run(m_krylov_bicgstab->getSolver());
806 }
else if (m_krylov == KrylovSolver::gmres) {
808 if (!m_krylov_gmres) {
809 m_krylov_gmres = std::make_unique<GMRES_MV<T>>(m_mat);
811 m_krylov_gmres->setPrecond(pc);
812 run(m_krylov_gmres->getSolver());
816 m_krylov_pcg = std::make_unique<PCG_MV<T>>(m_mat);
818 m_krylov_pcg->setPrecond(pc);
819 run(m_krylov_pcg->getSolver());
822 if (
amrex::isnan(m_rnorm) || m_rnorm > T(1.e20) * m_init_rnorm) { blew_up(); }
824 if (m_verbose >= 1) {
825 amrex::Print() << m_print_ident <<
"AlgMG(" <<
name <<
"): Final Iter. " << m_niters
826 <<
": rnorm, rnorm/bnorm = " << m_rnorm <<
", " << m_rnorm/m_bnorm;
827 if (m_verbose >= 2) {
828 amrex::Print() <<
" (Krylov recurrence " << krylov_rnorm <<
")";
830 amrex::Print() <<
"\n";
831 if (m_fixediter <= 0 && m_rnorm > res_target) {
832 amrex::Print() << m_print_ident <<
"AlgMG(" <<
name <<
"): Not converged after "
833 << m_niters <<
" iterations.\n";
839std::pair<T,T> AlgMG<T>::compute_rbnorms (
SpMatrix<T> const& mat,
844 auto bnorm = bvec.norm2(
true);
847 auto rnorm = rvec.norm2(
true);
854 rnorm = std::sqrt(rnorm);
857 return {
bnorm, rnorm};
873 int const nlev = numLevels();
874 int const bottom = nlev-1;
877 if (&b != &m_res[0]) { m_res[0].copyAsync(b); }
879 if (&
x != &m_cor[0]) {
x.copyAsync(m_cor[0]); }
884 auto cor = [&] (
int lev) ->
AlgVector<T>& {
return (lev == 0) ?
x : m_cor[lev]; };
885 auto res = [&] (
int lev) ->
AlgVector<T> const& {
return (lev == 0) ? b : m_res[lev]; };
887 smooth(0,
x, b, numPreSmooth(),
false);
889 for (
int lev = 0; lev < bottom; ++lev) {
891 SpMV(m_res[lev+1], m_R[lev], m_rescor[lev]);
892 if (lev+1 == bottom) {
895 smooth(lev+1, m_cor[lev+1], m_res[lev+1], numPreSmooth(),
false);
899 for (
int lev = bottom-1; lev >= 0; --lev) {
900 SpMV(m_rescor[lev], m_P[lev], m_cor[lev+1]);
901 cor(lev).plusAsync(m_rescor[lev]);
902 smooth(lev, cor(lev), res(lev), numPostSmooth(),
true,
true);
908 bool with_initial_guess,
bool backward)
910 if (m_smoother_type == Smoother::l1_gauss_seidel) {
913 smoother(xvec, bvec, with_initial_guess, backward);
914 }
else if (m_smoother_type == Smoother::chebyshev) {
917 smoother(xvec, bvec, with_initial_guess);
921 smoother(xvec, bvec, with_initial_guess);
926void AlgMG<T>::bottom_solve ()
930 int const lev = numLevels()-1;
931 project_out(m_res[lev]);
932 auto smooth_bottom = [&] () {
933 if (m_smoother_type == Smoother::l1_gauss_seidel) {
935 int const nu = (numBottomSmooth()+1)/2;
936 smooth(lev, m_cor[lev], m_res[lev], nu,
false);
937 smooth(lev, m_cor[lev], m_res[lev], nu,
true,
true);
939 smooth(lev, m_cor[lev], m_res[lev], numBottomSmooth(),
false);
945 auto krylov_bottom = [&] (
auto& krylov) {
946 krylov.setVerbose(m_bottom_verbose);
947 krylov.getSolver().setMaxIters(m_bottom_maxiter);
948 krylov.solve(m_cor[lev], m_res[lev], m_bottom_reltol, T(0));
949 auto const& solver = krylov.getSolver();
950 T
const rnorm0 = solver.getInitialResidualNorm();
951 if (rnorm0 > T(0) && !(solver.getResidualNorm() < rnorm0)) {
952 if (m_verbose >= 2) {
953 amrex::Print() << m_print_ident <<
"AlgMG: bottom solve failed, smoothing instead\n";
958 auto bottom_solver = m_bottom_solver;
959 if (bottom_solver == BottomSolver::direct) {
960 if (m_bottom_n < 0) { build_direct_bottom(); }
961 if (m_bottom_n == 0) { bottom_solver = BottomSolver::jacobi; }
963 switch (bottom_solver) {
964 case BottomSolver::direct:
965 direct_bottom_solve();
967 case BottomSolver::jacobi:
970 case BottomSolver::gmres:
972 m_gmres = std::make_unique<GMRES_MV<T>>(m_a[lev]);
973 m_gmres->setPrecond(
JacobiSmoother<T>(m_a[lev], m_smoother_type != Smoother::jacobi));
975 krylov_bottom(*m_gmres);
979 m_bicgstab = std::make_unique<BiCGStab_MV<T>>(m_a[lev]);
980 m_bicgstab->setPrecond(
JacobiSmoother<T>(m_a[lev], m_smoother_type != Smoother::jacobi));
982 krylov_bottom(*m_bicgstab);
985 project_out(m_cor[lev]);
995 m_a.push_back(m_mat);
996 m_mat->setColumnPartition(m_mat->partition());
997 m_setup_peak_bytes = 0;
998 note_memory(0,
"start");
1002 for (
int lev = 0; lev+1 < m_max_levels; ++lev)
1004 auto& A = *m_a[lev];
1005 if (A.numGlobalRows() <= m_max_coarse_size) {
break; }
1009 bool aggressive =
false;
1011 bool truncated =
false;
1013 if (!
coarsen_level(lev, A, P, cpart, aggressive, ncoarse, truncated, nnz0)) {
break; }
1014 note_memory(lev,
"strength release");
1015 if (m_p_max_elmts > 0 || m_trunc_factor > T(0)) {
1019 note_memory(lev,
"truncation");
1021 if (m_verbose >= 2) {
1024 amrex::Print() << m_print_ident <<
"AlgMG: level " << lev
1025 <<
": P truncated from " << nnz[0] <<
" to "
1026 << nnz[1] <<
" nonzeros\n";
1032 if (m_verbose >= 2 && aggressive) {
1033 amrex::Print() << m_print_ident <<
"AlgMG: level " << lev
1034 <<
": aggressive coarsening, " << ncoarse <<
" -> "
1037 if (m_verbose >= 3) {
1051 amrex::Print() << m_print_ident <<
"AlgMG: level " << lev <<
": P row sums in ["
1052 << rmin <<
", " << rmax <<
"], " << nzero <<
" empty rows\n";
1055 note_memory(lev,
"R");
1057#ifndef AMREX_USE_GPU
1064 note_memory(lev,
"A*P");
1068 note_memory(lev,
"coarse operator");
1070 m_P.push_back(std::move(P));
1071 m_R.push_back(std::move(R));
1072 m_crsemat.push_back(std::make_unique<
SpMatrix<T>>(std::move(Ac)));
1073 m_a.push_back(m_crsemat.back().get());
1077 for (
int lev = 0; lev < nlev; ++lev) {
1078 m_res.emplace_back(m_a[lev]->partition());
1079 m_cor.emplace_back(m_a[lev]->partition());
1081 m_rescor.emplace_back(m_a[lev]->partition());
1086 if (m_verbose >= 1) {
1087 Long nnz0 = 0, nnz_total = 0;
1088 for (
int lev = 0; lev < nlev; ++lev) {
1089 Long nnz = m_a[lev]->numLocalNonZeros();
1091 if (lev == 0) { nnz0 = nnz; }
1093 amrex::Print() << m_print_ident <<
"AlgMG: level " << lev <<
": "
1094 << m_a[lev]->numGlobalRows() <<
" rows, "
1095 << nnz <<
" nonzeros\n";
1097 amrex::Print() << m_print_ident <<
"AlgMG: operator complexity = "
1098 << double(nnz_total)/double(std::max(nnz0,
Long(1))) <<
"\n";
1099 if (m_setup_peak_bytes > 0) {
1100 auto peak = m_setup_peak_bytes;
1102 amrex::Print() << m_print_ident <<
"AlgMG: setup peak arena use = "
1103 << double(peak)/double(1<<20) <<
" MB ("
1104 << double(peak)/double(m_mat->numGlobalRows())
1112template <
typename T>
1113void AlgMG<T>::build_direct_bottom ()
1116 int const lev = numLevels()-1;
1117 auto const& A = *m_a[lev];
1118 Long const n = A.numGlobalRows();
1120 m_bottom_lu.clear();
1121 m_bottom_piv.clear();
1122 if (n < 1 || n > max_direct_size) {
1123 if (m_verbose >= 1 && n > max_direct_size) {
1124 amrex::Print() << m_print_ident <<
"AlgMG: coarsest level has " << n
1125 <<
" rows, using smoother sweeps instead of the direct bottom solver\n";
1132 auto const pc = A.const_parcsr();
1133 auto to_host = [] (
auto const* dptr,
Long cnt) {
1134 Vector<std::remove_const_t<std::remove_pointer_t<
decltype(dptr)>>> h(cnt);
1139 auto const& c = pc.csr0;
1140 auto ro = to_host(c.row_offset, (c.nrows > 0) ? c.nrows+1 : 0);
1141 auto ci = to_host(c.col_index, c.nnz);
1142 auto va = to_host(c.mat, c.nnz);
1143 for (
Long i = 0; i < c.nrows; ++i) {
1144 Long const ig = pc.row_begin + i;
1145 for (
int idx = ro[i]; idx < ro[i+1]; ++idx) {
1146 dense[ig*n + pc.col_begin + ci[idx]] += double(va[idx]);
1150 if (pc.csr1.nrows > 0 && pc.csr1.nnz > 0) {
1151 auto const& c = pc.csr1;
1152 auto ro = to_host(c.row_offset, c.nrows+1);
1153 auto ci = to_host(c.col_index, c.nnz);
1154 auto va = to_host(c.mat, c.nnz);
1155 auto rmap = to_host(pc.row_map, pc.csr0.nrows);
1157 for (
auto j : ci) { ncols = std::max(ncols, j+1); }
1158 auto cmap = to_host(pc.col_map, ncols);
1159 for (
Long i = 0; i < pc.csr0.nrows; ++i) {
1160 Long const r = rmap[i];
1161 if (r < 0) {
continue; }
1162 Long const ig = pc.row_begin + i;
1163 for (
int idx = ro[r]; idx < ro[r+1]; ++idx) {
1164 dense[ig*n + cmap[ci[idx]]] += double(va[idx]);
1168 allgather_rows(dense.data(), A.partition(), n);
1175 for (
Long i = 0; i < n; ++i) { dsum += std::abs(dense[i*n+i]); }
1176 double const w = (dsum/double(n))/
double(n);
1177 for (
auto& v : dense) { v += w; }
1183 for (
auto v : dense) { amat = std::max(amat, std::abs(v)); }
1184 double const pivot_tol = double(n) * std::numeric_limits<double>::epsilon() * amat;
1186 std::iota(piv.begin(), piv.end(), 0);
1187 for (
Long i = 0; i < n; ++i) {
1189 double amax = std::abs(dense[i*n+i]);
1190 for (
Long k = i+1; k < n; ++k) {
1191 double const a = std::abs(dense[k*n+i]);
1192 if (a > amax) { amax = a; imax = k; }
1194 if (amax <= pivot_tol) {
1195 if (m_verbose >= 1) {
1196 amrex::Print() << m_print_ident <<
"AlgMG: coarsest matrix is singular, "
1197 <<
"using smoother sweeps instead of the direct bottom solver\n";
1202 std::swap(piv[i], piv[imax]);
1203 for (
Long j = 0; j < n; ++j) { std::swap(dense[i*n+j], dense[imax*n+j]); }
1205 double const dinv = 1.0/dense[i*n+i];
1206 for (
Long j = i+1; j < n; ++j) {
1207 double const f = dense[j*n+i]*dinv;
1209 for (
Long k = i+1; k < n; ++k) { dense[j*n+k] -= f*dense[i*n+k]; }
1212 m_bottom_lu = std::move(dense);
1213 m_bottom_piv = std::move(piv);
1217template <
typename T>
1218void AlgMG<T>::allgather_rows (
double* buf,
AlgPartition const& part,
Long stride)
1222 if (nprocs == 1) {
return; }
1224 for (
int p = 0;
p < nprocs; ++
p) {
1225 displs[
p] =
int(part[
p]*stride);
1226 counts[
p] =
int((part[
p+1]-part[
p])*stride);
1229 buf, counts.data(), displs.data(), MPI_DOUBLE,
1236template <
typename T>
1237void AlgMG<T>::direct_bottom_solve ()
1239 int const lev = numLevels()-1;
1240 auto const& r = m_res[lev];
1241 auto&
x = m_cor[lev];
1242 Long const n = m_bottom_n;
1243 Long const nl = r.numLocalRows();
1244 Long const b0 = r.globalBegin();
1251 for (
Long i = 0; i < nl; ++i) { rhs[b0+i] = double(loc[i]); }
1253 allgather_rows(rhs.data(), r.partition(), 1);
1257 auto const& lu = m_bottom_lu;
1258 for (
Long i = 0; i < n; ++i) {
1259 double v = rhs[m_bottom_piv[i]];
1260 for (
Long k = 0; k < i; ++k) { v -= lu[i*n+k]*
y[k]; }
1263 for (
Long i = n-1; i >= 0; --i) {
1265 for (
Long k = i+1; k < n; ++k) { v -= lu[i*n+k]*
y[k]; }
1270 for (
Long i = 0; i < nl; ++i) { loc[i] = T(
y[b0+i]); }
1275template <
typename T>
1276void AlgMG<T>::build_smoothers ()
1283 int const nlev = numLevels();
1284 for (
int lev = 0; lev < nlev; ++lev) {
1285 if (m_smoother_type == Smoother::l1_gauss_seidel) {
1287 amrex::Abort(
"AlgMG: L1GaussSeidel smoother is only available in CPU builds");
1290 }
else if (m_smoother_type == Smoother::chebyshev) {
1292 m_cheby.back()->setDegree(m_cheby_degree);
1293 m_cheby.back()->setEigRatio(m_cheby_ratio);
1295 bool const l1 = (m_smoother_type == Smoother::l1_jacobi);
1297 T
const w = (m_relax_weight >= T(0)) ? m_relax_weight : (l1 ? T(4./3.) : T(2./3.));
1298 m_smoother.back()->setWeight(w);
1301 m_smoothers_stale =
false;
1304template <
typename T>
1306 bool& aggressive,
Long& ncoarse,
bool& truncated,
1307 Long& nnz_untruncated)
1314 S.setColumnPartition(S.partition());
1315 note_memory(lev,
"S");
1317 ST.setColumnPartition(ST.partition());
1319 note_memory(lev,
"strength");
1322 pmis(S, ST, state, m_singular);
1323 note_memory(lev,
"PMIS");
1328 if (ncoarse == 0 || ncoarse == A.
numGlobalRows()) {
return false; }
1329#ifndef AMREX_USE_GPU
1330 bool const host_interp = m_interp_type != InterpType::direct
1331 && detail::spmat_comm_is_local(A.
partition(), cpart);
1334 if (lev < m_aggressive_num_levels) {
1337 S2.setColumnPartition(cpart);
1339 S2T.setColumnPartition(cpart);
1340 note_memory(lev,
"second-pass strength");
1342 pmis(S2, S2T, state2,
true);
1345 Long const nc2 = cpart2.numGlobalRows();
1346 if (nc2 > 0 && nc2 < ncoarse) {
1351 auto const* p1 = state.
data();
1352 auto const* c1 = cidx.
data();
1353 auto const* p2 = state2.
data();
1354 auto const* c2 = cidx2.
data();
1355 auto* pf = statef.
data();
1356 auto* cf = cidxf.
data();
1359 bool const c = (p1[i] == PointType::coarse) &&
1360 (p2[c1[i]] == PointType::coarse);
1361 pf[i] = c ? PointType::coarse : PointType::fine;
1362 cf[i] = (p1[i] == PointType::coarse) ? c2[c1[i]] :
Long(0);
1366 bool const trunc = (m_p_max_elmts > 0 || m_trunc_factor > T(0));
1368#ifndef AMREX_USE_GPU
1369 if (!m_aggressive_direct_interp && host_interp) {
1371 P1 = build_interp_mm_host(A, S, ST, state, cidx, cpart, sscale,
1372 m_interp_type == InterpType::mm_ext_i,
1373 m_p_max_elmts, m_trunc_factor, unused);
1374 }
else if (!m_aggressive_direct_interp && m_interp_type != InterpType::direct) {
1376 m_interp_type == InterpType::mm_ext_i,
nullptr,
1377 m_p_max_elmts, m_trunc_factor);
1381 P1 = m_aggressive_direct_interp ?
build_interp(A, S, state, cidx, cpart)
1383 if (trunc) { P1 =
truncate_interp(P1, cpart, m_p_max_elmts, m_trunc_factor); }
1385 note_memory(lev,
"P1");
1387#ifndef AMREX_USE_GPU
1389 m_interp_type != InterpType::mm_ext,
1390 state.
data(), m_p_max_elmts, m_trunc_factor);
1393 m_interp_type != InterpType::mm_ext,
1395 if (trunc) { Pf =
truncate_interp(Pf, cpart2, m_p_max_elmts, m_trunc_factor); }
1398 note_memory(lev,
"P2");
1405#ifndef AMREX_USE_GPU
1407 P = build_interp_mm_host(A, S, ST, state, cidx, cpart, sscale,
1408 m_interp_type == InterpType::mm_ext_i,
1409 m_p_max_elmts, m_trunc_factor, nnz_untruncated);
1411 }
else if (m_interp_type != InterpType::direct) {
1413 m_interp_type == InterpType::mm_ext_i,
nullptr,
1414 m_p_max_elmts, m_trunc_factor, &nnz_untruncated);
1422 note_memory(lev,
"interpolation");
1426template <
typename T>
1429 auto const rs = P.
rowSum();
1435 amrex::Abort(
"AlgMG: singular mode needs P 1 = 1, but " + std::to_string(nzero)
1436 +
" rows of P are empty");
1441 T
const f = T(1) / prs[i];
1442 for (
Long idx = a.csr0.row_offset[i]; idx < a.csr0.row_offset[i+1]; ++idx) {
1443 a.csr0.mat[idx] *= f;
1445 if (detail::has_remote_row(a, i)) {
1446 Long const ii = a.row_map[i];
1447 for (
Long idx = a.csr1.row_offset[ii]; idx < a.csr1.row_offset[ii+1]; ++idx) {
1448 a.csr1.mat[idx] *= f;
1455template <
typename T>
1463 return (dmax > T(0)) ? dmax : T(1);
1466template <
typename T>
1472 Long const nrows = a.csr0.nrows;
1477 auto const threshold = m_strong_threshold;
1483 T
const tol = T(64) * std::numeric_limits<T>::epsilon();
1485#ifndef AMREX_USE_GPU
1488 auto const& c0 = a.csr0;
1489 auto const& c1 = a.csr1;
1490 auto rlen1 = [&] (
Long i) ->
Long {
1491 if (!detail::has_remote_row(a, i)) {
return 0; }
1492 Long const ii = a.row_map[i];
1493 return c1.row_offset[ii+1] - c1.row_offset[ii];
1495 auto csr = detail::csr_from_rows_cpu<TS,SpMatrix<TS>::template container_type,
int>
1496 (nrows, [&] (
Long i) {
return c0.row_offset[i+1] - c0.row_offset[i] + rlen1(i); },
1499 (
Long i,
int const*& rc,
TS const*& rv)
mutable ->
Long
1501 Long const b = c0.row_offset[i];
1502 Long const e = c0.row_offset[i+1];
1503 Long const n1 = rlen1(i);
1504 Long const b1 = (n1 > 0) ?
Long(c1.row_offset[a.row_map[i]]) : 0;
1505 T
const sgn = std::copysign(T(1), A_ii[i]);
1507 T amin = std::numeric_limits<T>::max();
1508 for (
Long idx = b; idx < e; ++idx) {
1510 if (c0.col_index[idx] != i) { amin = std::min(amin, sgn*c0.mat[idx]); }
1512 for (
Long idx = b1; idx < b1+n1; ++idx) {
1514 amin = std::min(amin, sgn*c1.mat[idx]);
1517 if (!(std::abs(rs) > T(0.9)*std::abs(A_ii[i])) && amin < T(0)) {
1518 T
const thr = amin * threshold * (T(1) - tol);
1519 if (
Long(col.size()) < e-b+n1) {
1523 for (
Long idx = b; idx < e; ++idx) {
1524 if (c0.col_index[idx] != i && sgn*c0.mat[idx] <= thr) {
1525 col[n] = c0.col_index[idx];
1526 val[n] =
TS(c0.mat[idx]/
scale);
1530 for (
Long idx = b1; idx < b1+n1; ++idx) {
1531 if (sgn*c1.mat[idx] <= thr) {
1532 col[n] = nlocal + c1.col_index[idx];
1533 val[n] =
TS(c1.mat[idx]/
scale);
1549 auto const rsA = A.
rowSum();
1552 if (std::abs(prs[i]) > T(0.9)*std::abs(A_ii[i])) {
1553 return std::numeric_limits<T>::lowest();
1555 auto const& c0 = a.csr0;
1556 T
const sgn = std::copysign(T(1), A_ii[i]);
1557 T amin = std::numeric_limits<T>::max();
1558 for (
Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
1559 if (c0.col_index[idx] != i) {
1560 amin = std::min(amin, sgn*c0.mat[idx]);
1563 if (detail::has_remote_row(a, i)) {
1564 auto const& c1 = a.csr1;
1565 Long const ii = a.row_map[i];
1566 for (
Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
1567 amin = std::min(amin, sgn*c1.mat[idx]);
1570 if (amin >= T(0)) {
return std::numeric_limits<T>::lowest(); }
1571 return amin * threshold * (T(1) - tol);
1579 csr.row_offset.resize(nrows+1);
1584 if (i >= nrows) {
return 0; }
1585 T
const thr = pthr[i];
1586 T
const sgn = std::copysign(T(1), A_ii[i]);
1588 auto const& c0 = a.csr0;
1589 for (
Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
1590 if (c0.col_index[idx] != i && sgn*c0.mat[idx] <= thr) { ++n; }
1592 if (detail::has_remote_row(a, i)) {
1593 auto const& c1 = a.csr1;
1594 Long const ii = a.row_map[i];
1595 for (
Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
1596 if (sgn*c1.mat[idx] <= thr) { ++n; }
1604 csr.resize(nrows, nnz);
1609 T
const thr = pthr[i];
1610 T
const sgn = std::copysign(T(1), A_ii[i]);
1612 auto const& c0 = a.csr0;
1613 for (
Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
1614 if (c0.col_index[idx] != i && sgn*c0.mat[idx] <= thr) {
1616 pc[
p] = c0.col_index[idx];
1620 if (detail::has_remote_row(a, i)) {
1621 auto const& c1 = a.csr1;
1622 Long const ii = a.row_map[i];
1623 for (
Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
1624 if (sgn*c1.mat[idx] <= thr) {
1626 pc[
p] = nlocal + c1.col_index[idx];
1640template <
typename T>
1642 bool isolated_as_coarse)
const
1647 int const U = PointType::undecided;
1648 int const C = PointType::coarse;
1658 Long n = st.csr0.row_offset[i+1] - st.csr0.row_offset[i];
1659 if (detail::has_remote_row(st, i)) {
1660 Long const ii = st.row_map[i];
1661 n += st.csr1.row_offset[ii+1] - st.csr1.row_offset[ii];
1663 Long nout = s.csr0.row_offset[i+1] - s.csr0.row_offset[i];
1664 if (detail::has_remote_row(s, i)) {
1665 Long const ii = s.row_map[i];
1666 nout += s.csr1.row_offset[ii+1] - s.csr1.row_offset[ii];
1668 pm[i] = (n == 0 && nout == 0) ?
Long(-1) : n;
1678 Long const row_begin = s.row_begin;
1686 ps[i] = isolated_as_coarse ? PointType::coarse : PointType::fine;
1688 pw[i] = T(pm[i]) + detail::pmis_tiebreak<T>(i + row_begin);
1689 ps[i] = PointType::undecided;
1713 Long const ig = i + row_begin;
1715 detail::pmis_blocked(s.csr0, i, wi, ig, pw, ps,
nullptr, s.col_begin, U) ||
1716 detail::pmis_blocked(st.csr0, i, wi, ig, pw, ps,
nullptr, st.col_begin, U);
1717 if (!blocked && detail::has_remote_row(s, i)) {
1718 blocked = detail::pmis_blocked(s.csr1, s.row_map[i], wi, ig,
1719 pwS, psS, s.col_map, 0, U);
1721 if (!blocked && detail::has_remote_row(st, i)) {
1722 blocked = detail::pmis_blocked(st.csr1, st.row_map[i], wi, ig,
1723 pwST, psST, st.col_map, 0, U);
1725 if (!blocked) { st_i =
C; }
1738 if (detail::csr_row_has_state(s.csr0, i, ps2,
C) ||
1739 (detail::has_remote_row(s, i) &&
1740 detail::csr_row_has_state(s.csr1, s.row_map[i], ps2S,
C))) {
1741 st_i = PointType::fine;
1748 return (ps[i] == U) ? 1 : 0;
1751 if (nund == 0) {
break; }
1754 if (m_verbose >= 2) {
1758 return (ps[i] ==
C) ? 1 : 0;
1764 if (ps[i] !=
C) {
return 0; }
1765 return (detail::csr_row_has_state(s.csr0, i, ps,
C) ||
1766 (detail::has_remote_row(s, i) &&
1767 detail::csr_row_has_state(s.csr1, s.row_map[i], psS,
C))) ? 1 : 0;
1771 if (ps[i] != PointType::fine) {
return 0; }
1772 return (detail::csr_row_has_state(s.csr0, i, ps,
C) ||
1773 (detail::has_remote_row(s, i) &&
1774 detail::csr_row_has_state(s.csr1, s.row_map[i], psS,
C))) ? 0 : 1;
1776 Long counts[3] = {nc, ncc, nf0};
1778 amrex::Print() << m_print_ident <<
"AlgMG: PMIS: " << counts[0] <<
" C-points of "
1780 <<
", F-points w/o C neighbor " << counts[2] <<
"\n";
1784template <
typename T>
1805 std::partial_sum(counts.begin(), counts.end(), rows.begin()+1);
1809template <
typename T>
1820 Long const nrows = a.csr0.nrows;
1834 detail::DirectInterpRow<T,TS>
const w{a, S.
const_parcsr(), ps, pc,
1835 state_r.data(), nlocal_c,
1836 int(PointType::coarse)};
1840 csr.row_offset.resize(nrows+1);
1846 if (i >= nrows) {
return 0; }
1847 if (ps[i] == PointType::coarse) {
return 1; }
1850 w.walk(i, diag, [&] (
int, T,
bool strong_c) {
if (strong_c) { ++n; } });
1856 csr.resize(nrows, nnz);
1863 if (ps[i] == PointType::coarse) {
1864 pcol[
p] =
int(pc[i]);
1869 T sum_n_neg = T(0), sum_n_pos = T(0), sum_p_neg = T(0), sum_p_pos = T(0);
1871 w.walk(i, diag, [&] (
int, T v,
bool strong_c) {
1874 if (strong_c) { sum_p_neg += v; }
1877 if (strong_c) { sum_p_pos += v; }
1881 T alfa = T(0), beta = T(0);
1882 if (sum_p_neg != T(0)) { alfa = sum_n_neg / sum_p_neg; }
else { diag += sum_n_neg; }
1883 if (sum_p_pos != T(0)) { beta = sum_n_pos / sum_p_pos; }
else { diag += sum_n_pos; }
1884 if (diag == T(0)) { diag = T(1); }
1887 w.walk(i, unused, [&] (
int cj, T v,
bool strong_c) {
1890 pmat[
p] = ((v < T(0)) ? -alfa*v : -beta*v) / diag;
1907template <
typename T,
int KMAX>
1916 template <
typename F>
1919 auto const& c0 =
p.csr0;
1920 for (
Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
1921 f(c0.col_index[idx], c0.mat[idx]);
1923 if (has_remote_row(
p, i)) {
1924 auto const& c1 =
p.csr1;
1925 Long const ii =
p.row_map[i];
1926 for (
Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
1927 f(nlocal_cols + c1.col_index[idx], c1.mat[idx]);
1938 walk(i, [&] (
int, T v) {
1939 amax = std::max(amax, std::abs(v));
1940 if (v > T(0)) { sum_pos += v; }
else { sum_neg += v; }
1942 return trunc_factor * amax;
1949 walk(i, [&] (
int, T v) {
if (std::abs(v) >= thresh) { ++n; } });
1957 int select (
Long i, T thresh,
int* cols, T* vals)
const
1959 int const kmax = std::min(max_elmts, KMAX);
1961 walk(i, [&] (
int col, T v) {
1962 T
const av = std::abs(v);
1963 if (av < thresh) {
return; }
1964 if (n == kmax && av <= std::abs(vals[n-1])) {
return; }
1965 int pos = (n < kmax) ? n : n-1;
1966 while (pos > 0 && std::abs(vals[pos-1]) < av) {
1967 cols[pos] = cols[pos-1];
1968 vals[pos] = vals[pos-1];
1973 if (n < kmax) { ++n; }
1979#ifndef AMREX_USE_GPU
1982template <
typename T,
int KMAX>
1983Long truncate_row_host (
int* col, T* val,
Long n,
int max_elmts, T trunc_factor,
int nc)
1986 int ro[2] = {0,
int(n)};
1987 ParCsr<T const> prow;
1988 prow.csr0 = CsrView<T const,int>{val, col, ro, n, 1};
1989 TruncRow<T,KMAX>
const tr{prow, max_elmts, trunc_factor, nc};
1991 T
const thresh = tr.threshold(0, sum_pos, sum_neg);
1995 T kept_pos = T(0), kept_neg = T(0);
1996 if (max_elmts > 0) {
1997 nk = tr.select(0, thresh, tcol, tval);
1998 for (
Long k = 0; k < nk; ++k) {
1999 if (tval[k] > T(0)) { kept_pos += tval[k]; }
else { kept_neg += tval[k]; }
2002 for (
Long k = 0; k < n; ++k) {
2003 if (std::abs(val[k]) >= thresh) {
2004 if (val[k] > T(0)) { kept_pos += val[k]; }
else { kept_neg += val[k]; }
2008 T
const fpos = (kept_pos != T(0)) ? sum_pos/kept_pos : T(1);
2009 T
const fneg = (kept_neg != T(0)) ? sum_neg/kept_neg : T(1);
2010 if (max_elmts > 0) {
2012 for (
Long k = 0; k < nk; ++k) {
2014 for (; m > 0 && col[m-1] > tcol[k]; --m) {
2019 val[m] = tval[k] * ((tval[k] > T(0)) ? fpos : fneg);
2022 for (
Long k = 0; k < n; ++k) {
2023 if (std::abs(val[k]) >= thresh) {
2025 val[nk] = val[k] * ((val[k] > T(0)) ? fpos : fneg);
2042template <
typename T,
typename F>
2055 if (i >= nrows || state[i] != c) {
return 0; }
2057 auto const& c0 =
p.csr0;
2058 for (
Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
2059 if (colmap(c0.col_index[idx],
false) >= 0) { ++n; }
2061 if (has_remote_row(
p, i)) {
2062 auto const& c1 =
p.csr1;
2063 Long const ii =
p.row_map[i];
2064 for (
Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
2065 if (colmap(c1.col_index[idx],
true) >= 0) { ++n; }
2080 if (i == nrows) { prow[nc] = nnz;
return; }
2081 if (state[i] != c) {
return; }
2084 auto const& c0 =
p.csr0;
2085 for (
Long idx = c0.row_offset[i]; idx < c0.row_offset[i+1]; ++idx) {
2086 int const g = colmap(c0.col_index[idx],
false);
2087 if (g >= 0) { pcol[q] = g; pmat[q] = c0.mat[idx]; ++q; }
2089 if (has_remote_row(
p, i)) {
2090 auto const& c1 =
p.csr1;
2091 Long const ii =
p.row_map[i];
2092 for (
Long idx = c1.row_offset[ii]; idx < c1.row_offset[ii+1]; ++idx) {
2093 int const g = colmap(c1.col_index[idx],
true);
2094 if (g >= 0) { pcol[q] = g; pmat[q] = c1.mat[idx]; ++q; }
2104template <
typename T>
2110 if (m_interp_type == InterpType::direct) {
2114 m_interp_type == InterpType::mm_ext_i);
2118template <
typename T>
2130 Long const nrows = s.csr0.nrows;
2132 int const C = PointType::coarse;
2133 int const F = PointType::fine;
2137 Long const nf = fpart.numLocalRows();
2138 int const nlocal_c =
int(nc);
2139 int const nlocal_f =
int(nf);
2145 Long const fbegin = fpart.globalRowBegin();
2152 pgc[i] = pc[i] + cbegin;
2153 pgf[i] = pf[i] + fbegin;
2161 auto csr_cf = detail::compact_rows<TS>
2165 return (ps_r[j] ==
F) ? nlocal_f + j : -1;
2167 return (ps[j] ==
F) ?
int(pf[j]) : -1;
2171 SCF.
define_split(cpart, fpart, std::move(csr_cf), nlocal_f, gfidx_r.data(), 0);
2174 auto csr_fc = detail::compact_rows<TS>
2178 return (ps_r[j] ==
C) ? nlocal_c + j : -1;
2180 return (ps[j] ==
C) ?
int(pc[j]) : -1;
2184 SFC.
define_split(fpart, cpart, std::move(csr_fc), nlocal_c, gcidx_r.data(), 0);
2189 auto const& s2 = S2.const_parcsr();
2196 if (r >= nc) {
return 0; }
2198 auto const& c0 = s2.csr0;
2199 for (
Long idx = c0.row_offset[r]; idx < c0.row_offset[r+1]; ++idx) {
2200 n += (c0.col_index[idx] !=
int(r));
2202 if (detail::has_remote_row(s2, r)) {
2203 Long const rr = s2.row_map[r];
2204 n += s2.csr1.row_offset[rr+1] - s2.csr1.row_offset[rr];
2216 auto const& c0 = s2.csr0;
2217 for (
Long idx = c0.row_offset[r]; idx < c0.row_offset[r+1]; ++idx) {
2218 if (c0.col_index[idx] !=
int(r)) {
2219 ocol[q] = c0.col_index[idx];
2220 omat[q] = c0.mat[idx];
2224 if (detail::has_remote_row(s2, r)) {
2225 auto const& c1 = s2.csr1;
2226 Long const rr = s2.row_map[r];
2227 for (
Long idx = c1.row_offset[rr]; idx < c1.row_offset[rr+1]; ++idx) {
2228 ocol[q] = nlocal_c + c1.col_index[idx];
2229 omat[q] = c1.mat[idx];
2236 S2d.
define_split(cpart, cpart, std::move(out), nlocal_c, s2.col_map, 0);
2240template <
typename T>
2248 auto csr = detail::compact_rows<T>
2251 return remote ? nlocal + j : j;
2258template <
typename T>
2260 int max_elmts, T trunc_factor)
2269 Long const nrows =
p.csr0.nrows;
2271 detail::TruncRow<T,KMAX>
const tr{
p, max_elmts, trunc_factor, nlocal};
2280 pthr[i] = tr.threshold(i, ppos[i], pneg[i]);
2285 csr.row_offset.resize(nrows+1);
2290 if (i >= nrows) {
return 0; }
2291 int const n = tr.count(i, pthr[i]);
2292 return (max_elmts > 0) ? std::min(n, max_elmts) : n;
2297 csr.resize(nrows, nnz);
2302 T
const thresh = pthr[i];
2303 T
const sum_pos = ppos[i];
2304 T
const sum_neg = pneg[i];
2306 if (max_elmts > 0) {
2309 int const n = tr.select(i, thresh, cols, vals);
2310 T kept_pos = T(0), kept_neg = T(0);
2311 for (
int k = 0; k < n; ++k) {
2312 if (vals[k] > T(0)) { kept_pos += vals[k]; }
else { kept_neg += vals[k]; }
2314 T
const fpos = (kept_pos != T(0)) ? sum_pos/kept_pos : T(1);
2315 T
const fneg = (kept_neg != T(0)) ? sum_neg/kept_neg : T(1);
2316 for (
int k = 0; k < n; ++k) {
2318 pmat[q] = vals[k] * ((vals[k] > T(0)) ? fpos : fneg);
2322 T kept_pos = T(0), kept_neg = T(0);
2323 tr.walk(i, [&] (
int, T v) {
2324 if (std::abs(v) >= thresh) {
2325 if (v > T(0)) { kept_pos += v; }
else { kept_neg += v; }
2328 T
const fpos = (kept_pos != T(0)) ? sum_pos/kept_pos : T(1);
2329 T
const fneg = (kept_neg != T(0)) ? sum_neg/kept_neg : T(1);
2330 tr.walk(i, [&] (
int col, T v) {
2331 if (std::abs(v) >= thresh) {
2333 pmat[q] = v * ((v > T(0)) ? fpos : fneg);
2348template <
typename T>
2354 bool plus_i,
int const* row_state,
2355 int max_elmts, T trunc_factor,
Long* nnz_untruncated)
2364 Long const nrows = s.csr0.nrows;
2370 int const C = PointType::coarse;
2371 int const F = PointType::fine;
2375 auto const rsA = A.
rowSum();
2376 auto const rsS = S.
rowSum();
2393 auto const& sc = s.csr0;
2394 for (
Long idx = sc.row_offset[i]; idx < sc.row_offset[i+1]; ++idx) {
2395 if (ps[sc.col_index[idx]] ==
C) { d += T(sc.mat[idx]); }
2397 if (detail::has_remote_row(s, i)) {
2398 auto const& sc1 = s.csr1;
2399 Long const ii = s.row_map[i];
2400 for (
Long idx = sc1.row_offset[ii]; idx < sc1.row_offset[ii+1]; ++idx) {
2401 if (pstate_r[sc1.col_index[idx]] ==
C) { d += T(sc1.mat[idx]); }
2409 detail::MMInterpRow<T,TS>
const w{s, ST.
const_parcsr(), ps, pdb, pc,
2410 pstate_r, dbeta_r.data(), nlocal_f, nlocal_c};
2415 local_csr_type mcsr, bcsr;
2416 mcsr.row_offset.resize(nrows+1);
2417 bcsr.row_offset.resize(nrows+1);
2422 return st_j ==
F && (db_j + (plus_i ? aji : T(0))) != T(0);
2429 return prow_state ==
nullptr || prow_state[i] ==
C;
2435 if (i >= nrows || !has_row(i)) {
return 0; }
2438 w.walk(i, [&] (
int,
int, T, T aji,
int st_j, T db_j) {
2439 if (usable_f(aji, st_j, db_j)) { ++n; }
2450 if (i >= nrows) {
return 0; }
2451 if (ps[i] ==
C) {
return 1; }
2453 w.walk(i, [&] (
int,
int, T, T,
int st_j, T) {
2454 if (st_j ==
C) { ++n; }
2461 mcsr.resize(nrows, mnnz);
2462 bcsr.resize(nrows, bnnz);
2472 bool const in_p = has_row(i);
2478 pbcol[bp] =
int(pc[i]);
2482 T
const db_i = pdb[i];
2483 T
const bscale = plus_i ? T(1) : ((db_i != T(0)) ? T(1)/db_i : T(0));
2485 w.walk(i, [&] (
int,
int jc, T aij, T,
int st_j, T) {
2488 pbmat[bp] = bscale*aij;
2497 T alpha = prsA[i]/sscale - T(prsS[i]);
2499 w.walk(i, [&] (
int,
int, T aij, T aji,
int st_j, T db_j) {
2501 T
const den = db_j + (plus_i ? aji : T(0));
2503 if (plus_i) { theta += aij * aji / den; }
2509 T
const denom = alpha + theta;
2510 T
const scale = (denom != T(0)) ? T(-1)/denom : T(0);
2515 w.walk(i, [&] (
int jf,
int jc, T aij, T aji,
int st_j, T db_j) {
2518 pbmat[bp] = bscale*aij;
2520 }
else if (usable_f(aji, st_j, db_j)) {
2522 pmmat[mp] = plus_i ?
scale*aij/(db_j + aji) :
scale*aij;
2534#ifndef AMREX_USE_GPU
2535 if (max_elmts > 0 || trunc_factor > T(0)) {
2542 return detail::truncate_row_host<T,KMAX>(col, val, n, max_elmts, trunc_factor,
2545 if (nnz_untruncated) {
2546 *nnz_untruncated = std::accumulate(n0.begin(), n0.end(),
Long(0));
2556#ifndef AMREX_USE_GPU
2557template <
typename T>
2563 int max_elmts, T trunc_factor,
2564 Long& nnz_untruncated)
2570 constexpr int KMAX = max_p_elmts;
2574 Long const nrows = s.csr0.nrows;
2578 int const C = PointType::coarse;
2579 int const F = PointType::fine;
2581 auto const rsA = A.
rowSum();
2582 auto const rsS = S.
rowSum();
2590 auto const& sc = s.csr0;
2595 for (
Long idx = sc.row_offset[i]; idx < sc.row_offset[i+1]; ++idx) {
2596 if (ps[sc.col_index[idx]] ==
C) {
2597 d += T(sc.mat[idx]);
2602 boff[i+1] = (ps[i] ==
F) ? n : 0;
2604 std::partial_sum(boff.begin(), boff.end(), boff.begin());
2609 if (ps[i] !=
F) {
return; }
2610 T
const bscale = plus_i ? T(1) : ((pdb[i] != T(0)) ? T(1)/pdb[i] : T(0));
2612 for (
Long idx = sc.row_offset[i]; idx < sc.row_offset[i+1]; ++idx) {
2613 int const j = sc.col_index[idx];
2615 bcol[q] = int(pc[j]);
2616 bval[q] = bscale*T(sc.mat[idx]);
2622 detail::MMInterpRow<T,TS>
const w{s, ST.
const_parcsr(), ps, pdb, pc,
2623 nullptr,
nullptr,
int(nrows), nc};
2624 bool const trunc = (max_elmts > 0 || trunc_factor > T(0));
2625 struct Nbr {
int j; T aij; T den; };
2628 auto csr = detail::csr_from_rows_cpu<T,SpMatrix<T>::template container_type,
int>
2629 (nrows, [&] (
Long i) {
return sc.row_offset[i+1] - sc.row_offset[i] + 1; },
2633 (
Long i,
int const*& rc, T
const*& rv)
mutable ->
Long
2638 col[0] =
int(pc[i]);
2646 T alpha = prsA[i]/sscale - T(prsS[i]);
2648 Long maxlen = boff[i+1] - boff[i];
2650 w.walk(i, [&] (
int jf,
int, T aij, T aji,
int st_j, T db_j) {
2652 T
const den = db_j + (plus_i ? aji : T(0));
2654 if (plus_i) { theta += aij * aji / den; }
2655 maxlen += boff[jf+1] - boff[jf];
2656 nbr.push_back({jf, aij, den});
2662 T
const denom = alpha + theta;
2663 T
const scale = (denom != T(0)) ? T(-1)/denom : T(0);
2664 maxlen = std::min(maxlen,
Long(nc));
2665 if (
Long(col.size()) < maxlen) {
2674 auto add = [&] (
Long k, T m) {
2675 for (
Long bp = boff[k]; bp < boff[k+1]; ++bp) {
2676 int const j = bcol[bp];
2677 Long const pos = marker[j];
2678 if (pos >= 0 && pos < n && col[pos] == j) {
2679 val[pos] += m * bval[bp];
2683 val[n] = m * bval[bp];
2688 T
const mdiag = plus_i ?
scale :
scale*pdb[i];
2689 bool diag_done =
false;
2690 for (
auto const& nb : nbr) {
2691 if (!diag_done && nb.j > i) {
2695 add(nb.j, plus_i ?
scale*nb.aij/nb.den :
scale*nb.aij);
2697 if (!diag_done) { add(i, mdiag); }
2698 detail::sort_row_cpu(col.data(), val.data(), n, tmp);
2700 if (!trunc) {
return n; }
2702 return detail::truncate_row_host<T,KMAX>(col.data(), val.data(), n, max_elmts,
2707 nnz_untruncated = std::accumulate(nnz0.begin(), nnz0.end(),
Long(0));
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:593
#define AMREX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition AMReX_BLassert.H:49
#define AMREX_ALWAYS_ASSERT(EX)
Definition AMReX_BLassert.H:50
Coalescing first-fit dynamic memory arena.
Enum reflection utilities and the AMREX_ENUM macro.
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
#define AMREX_RESTRICT
Definition AMReX_Extension.H:37
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
while parser_exe_t p
Definition AMReX_Parser_Exe_Body.H:33
Algebraic multigrid solver for SpMatrix / AlgVector systems.
Definition AMReX_AlgMG.H:264
SpMatrix< TS > create_soc(SpMatrix< T > const &A, T scale) const
Definition AMReX_AlgMG.H:1467
void setAggressiveNumLevels(int n)
Number of levels with aggressive coarsening, starting at the finest. Next setup.
Definition AMReX_AlgMG.H:322
void setFixedIter(int n)
If positive, run exactly this many V-cycles regardless of tolerance.
Definition AMReX_AlgMG.H:289
void setMaxIter(int n)
Definition AMReX_AlgMG.H:287
void setChebyshevDegree(int d)
Chebyshev polynomial degree. Next solve.
Definition AMReX_AlgMG.H:298
SpMatrix< T > const & getMatrix(int lev) const
Definition AMReX_AlgMG.H:382
AlgMGInterpType InterpType
Definition AMReX_AlgMG.H:272
AlgMGKrylovSolver KrylovSolver
Definition AMReX_AlgMG.H:271
T getResidualNorm() const
Definition AMReX_AlgMG.H:375
void setMaxCoarseSize(Long n)
Stop coarsening when a level has at most this many rows. Next setup.
Definition AMReX_AlgMG.H:320
void solve(AlgVector< T > &a_sol, AlgVector< T > const &a_rhs)
Solve A x = b. a_sol holds the initial guess on entry.
Definition AMReX_AlgMG.H:658
static SpMatrix< TS > second_pass_strength(SpMatrix< TS > &S, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart)
Definition AMReX_AlgMG.H:2119
void setRelTol(T t)
Definition AMReX_AlgMG.H:290
std::conditional_t< std::is_same_v< T, double >, float, T > TS
Definition AMReX_AlgMG.H:268
void setBottomMaxIter(int n)
Definition AMReX_AlgMG.H:368
void setPreSmooth(int nu)
Number of pre-smoothing sweeps; negative restores the smoother's default.
Definition AMReX_AlgMG.H:305
static constexpr int default_p_max_elmts
Definition AMReX_AlgMG.H:343
void setBottomSolver(BottomSolver bs)
Definition AMReX_AlgMG.H:360
void define(SpMatrix< T > &a_mat)
Define the solver with a square matrix. Must be called before solve.
Definition AMReX_AlgMG.H:626
void setPMaxElmts(int n)
Definition AMReX_AlgMG.H:337
void setAbsTol(T t)
Converged when the residual 2-norm is below max(reltol*bnorm, abstol).
Definition AMReX_AlgMG.H:292
bool coarsen_level(int lev, SpMatrix< T > &A, SpMatrix< T > &P, AlgPartition &cpart, bool &aggressive, Long &ncoarse, bool &truncated, Long &nnz_untruncated)
Definition AMReX_AlgMG.H:1305
void setSingular(bool b)
Definition AMReX_AlgMG.H:330
void setPrintIndentation(std::string s)
Prefix of the lines the solver prints, e.g. to nest them in MLMG output.
Definition AMReX_AlgMG.H:286
void setPostSmooth(int nu)
Number of post-smoothing sweeps; negative restores the smoother's default.
Definition AMReX_AlgMG.H:307
void setInterpType(InterpType it)
Takes effect at the next setup.
Definition AMReX_AlgMG.H:334
SpMatrix< T > const & getInterp(int lev) const
Definition AMReX_AlgMG.H:383
AlgMGBottomSolver BottomSolver
Definition AMReX_AlgMG.H:270
void setBottomSmooth(int nu)
Definition AMReX_AlgMG.H:310
void setTruncFactor(T f)
Drop interpolation weights below this fraction of the row maximum. Next setup.
Definition AMReX_AlgMG.H:345
static AlgPartition coarse_numbering(Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > &cidx, int which=PointType::coarse)
Local index among the rows of type which and their partition.
Definition AMReX_AlgMG.H:1785
void setBottomVerbose(int v)
Definition AMReX_AlgMG.H:369
static SpMatrix< T > select_coarse_rows(SpMatrix< T > const &P, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart)
Rows of P at the C-points, in C-point numbering.
Definition AMReX_AlgMG.H:2241
static T strength_scale(SpMatrix< T > const &A)
Largest |a_ii|, the unit of the strength matrix entries.
Definition AMReX_AlgMG.H:1456
void setMaxLevels(int n)
Definition AMReX_AlgMG.H:318
static SpMatrix< T > truncate_interp(SpMatrix< T > &P, AlgPartition const &cpart, int max_elmts, T trunc_factor)
Definition AMReX_AlgMG.H:2259
int numLevels() const
Number of levels; 0 before setup.
Definition AMReX_AlgMG.H:377
Long numGlobalRows(int lev) const
Global number of rows on a level. Valid after setup.
Definition AMReX_AlgMG.H:381
static constexpr int max_p_elmts
Definition AMReX_AlgMG.H:342
void precond(AlgVector< T > &x, AlgVector< T > const &b)
Definition AMReX_AlgMG.H:746
void setStrongThreshold(T theta)
Definition AMReX_AlgMG.H:315
static SpMatrix< T > build_interp_mm(SpMatrix< T > &A, SpMatrix< TS > &S, SpMatrix< TS > const &ST, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart, T scale, bool plus_i, int const *row_state=nullptr, int max_elmts=0, T trunc_factor=T(0), Long *nnz_untruncated=nullptr)
Definition AMReX_AlgMG.H:2349
void setRelaxWeight(T w)
Relaxation weight of the smoother; negative restores the default. Next solve.
Definition AMReX_AlgMG.H:296
void setup()
Definition AMReX_AlgMG.H:989
void setAggressiveDirectInterp(bool b)
Definition AMReX_AlgMG.H:325
static SpMatrix< T > build_interp(SpMatrix< T > &A, SpMatrix< TS > const &S, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart)
Classical direct interpolation.
Definition AMReX_AlgMG.H:1810
SpMatrix< T > interpolation(SpMatrix< T > &A, SpMatrix< TS > &S, SpMatrix< TS > const &ST, Gpu::DeviceVector< int > const &state, Gpu::DeviceVector< Long > const &cidx, AlgPartition const &cpart, T scale) const
Interpolation of the selected type from the C-points in state.
Definition AMReX_AlgMG.H:2105
void pmis(SpMatrix< TS > &S, SpMatrix< TS > &ST, Gpu::DeviceVector< int > &state, bool isolated_as_coarse=false) const
Definition AMReX_AlgMG.H:1641
std::size_t getSetupPeakBytes() const
Peak bytes in use in The_Arena during setup (0 if not available).
Definition AMReX_AlgMG.H:379
void setSmoother(Smoother sm)
Takes effect at the next solve; the hierarchy is kept.
Definition AMReX_AlgMG.H:294
void setThrowException(bool b)
Definition AMReX_AlgMG.H:370
void setVerbose(int v)
Definition AMReX_AlgMG.H:284
static void normalize_rows(SpMatrix< T > &P)
Scales the rows of P to sum to one; aborts on an empty row.
Definition AMReX_AlgMG.H:1427
int getNumSetups() const
Number of setups run so far (for checks that the setup is reused).
Definition AMReX_AlgMG.H:374
void project_out(AlgVector< T > &x)
Definition AMReX_AlgMG.H:548
void setBottomTol(T t)
Definition AMReX_AlgMG.H:367
AlgMGSmoother Smoother
Definition AMReX_AlgMG.H:273
void setKrylovSolver(KrylovSolver a)
Krylov solver with the V-cycle as preconditioner (see KrylovSolver). Default None.
Definition AMReX_AlgMG.H:347
void setChebyshevRatio(T r)
Chebyshev eigenvalue ratio lambda_max / lambda_min. Next solve.
Definition AMReX_AlgMG.H:300
int getNumIters() const
Definition AMReX_AlgMG.H:372
Definition AMReX_AlgPartition.H:26
Long numGlobalRows() const
Total number of rows covered by the partition.
Definition AMReX_AlgPartition.H:55
Long globalRowBegin() const
Inclusive global index begin on this process.
Definition AMReX_AlgPartition.H:62
Long numLocalRows() const
Number of local rows.
Definition AMReX_AlgPartition.H:50
Distributed dense vector that mirrors the layout of an AlgPartition.
Definition AMReX_AlgVector.H:29
void setVal(T val)
Definition AMReX_AlgVector.H:273
Long numLocalRows() const
Number of entries stored on this rank.
Definition AMReX_AlgVector.H:74
T const * data() const
Definition AMReX_AlgVector.H:85
void plusAsync(AlgVector< T, Allocator > const &rhs)
Definition AMReX_AlgVector.H:312
Chebyshev polynomial smoother for AlgVector/SpMatrix systems.
Definition AMReX_Smoother_MV.H:163
static std::size_t freeMemAvailable()
Definition AMReX_GpuDevice.cpp:1481
Jacobi smoother for AlgVector/SpMatrix linear systems.
Definition AMReX_Smoother_MV.H:32
l1 hybrid Gauss-Seidel smoother for CPUs.
Definition AMReX_Smoother_MV.H:318
Dynamically allocated vector for trivially copyable data.
Definition AMReX_PODVector.H:308
size_type size() const noexcept
Definition AMReX_PODVector.H:654
void resize(size_type a_new_size, GrowthStrategy strategy=GrowthStrategy::Poisson)
Definition AMReX_PODVector.H:734
T * data() noexcept
Definition AMReX_PODVector.H:672
This class provides the user with a few print options.
Definition AMReX_Print.H:35
Distributed CSR matrix that manages storage and GPU-friendly partitions.
Definition AMReX_SpMatrix.H:65
void setColumnPartition(AlgPartition const &col_partition)
Set the column partition and split the matrix into local and remote blocks with 32-bit column indices...
Definition AMReX_SpMatrix.H:1019
Long numGlobalRows() const
Global row count.
Definition AMReX_SpMatrix.H:196
Gpu::DeviceVector< U > gatherRemote(U const *x)
Gather values at the columns of the off-diagonal block.
Definition AMReX_SpMatrix.H:1342
Long numLocalNonZeros() const
Number of nonzeros stored locally.
Definition AMReX_SpMatrix.H:198
AlgVector< T, AllocT > rowSum() const
Sum the values in each local row and return the result as an AlgVector.
Definition AMReX_SpMatrix.H:1164
void define_split(AlgPartition partition, AlgPartition const &col_partition, local_csr_type &&compact, Long nlocal, Long const *remote_cols, Long nremote)
Define directly in split form from a CSR with compact 32-bit columns: c < nlocal is the local column ...
Definition AMReX_SpMatrix.H:1695
AlgPartition const & columnPartition() const
Return the column partition used for matrix-vector and matrix-matrix multiplications.
Definition AMReX_SpMatrix.H:191
ParCsr< T const > const_parcsr() const
Const-qualified alias of parcsr() for convenience.
Definition AMReX_SpMatrix.H:1226
ParCsr< T > parcsr()
Build GPU-friendly CSR views split into diagonal/off-diagonal blocks.
Definition AMReX_SpMatrix.H:1201
Long numLocalRows() const
Number of rows owned by this rank.
Definition AMReX_SpMatrix.H:194
AlgPartition const & partition() const
Row partition describing how matrix rows are distributed across ranks.
Definition AMReX_SpMatrix.H:182
AlgVector< T, AllocT > const & diagonalVector() const
Return diagonal elements in a square matrix.
Definition AMReX_SpMatrix.H:1145
CSR< T, container_type, int > local_csr_type
Definition AMReX_SpMatrix.H:71
This class is a thin wrapper around std::vector. Unlike vector, Vector::operator[] provides bound che...
Definition AMReX_Vector.H:29
amrex_long Long
Definition AMReX_INT.H:30
T Min(N n, T const *v, T init_val=std::numeric_limits< T >::max())
Compute the minimum of an array of values.
Definition AMReX_Reduce.H:1102
T Max(N n, T const *v, T init_val=std::numeric_limits< T >::lowest())
Compute the maximum of an array of values.
Definition AMReX_Reduce.H:1125
T Sum(N n, T const *v, T init_val=0)
Compute the sum of an array of values.
Definition AMReX_Reduce.H:1079
void ParallelForOMP(T n, L const &f) noexcept
Performance-portable kernel launch function with optional OpenMP threading.
Definition AMReX_GpuLaunch.H:328
PODVector< T, PinnedArenaAllocator< T > > PinnedVector
A PODVector that uses the pinned memory arena.
Definition AMReX_GpuContainers.H:48
PODVector< T, ArenaAllocator< T > > DeviceVector
A PODVector that uses the standard memory Arena. Note that the memory might or might not be managed d...
Definition AMReX_GpuContainers.H:24
Arena * The_Arena()
Definition AMReX_Arena.cpp:839
void Min(KeyValuePair< K, V > &vi, MPI_Comm comm)
Definition AMReX_ParallelReduce.H:161
void Sum(Gpu::DeviceVector< T > &v, MPI_Comm comm)
Definition AMReX_GpuParallelReduce.H:37
void AllGather(const T *v, int cnt, T *vs, MPI_Comm comm)
Definition AMReX_ParallelReduce.H:106
void Max(KeyValuePair< K, V > &vi, MPI_Comm comm)
Definition AMReX_ParallelReduce.H:133
#define AMREX_ENUM(CLASS,...)
Declare a scoped enum with reflection support.
Definition AMReX_Enum.H:270
void copy(HostToDevice, InIter begin, InIter end, OutIter result) noexcept
A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous st...
Definition AMReX_GpuContainers.H:128
static constexpr DeviceToHost deviceToHost
Definition AMReX_GpuContainers.H:106
static constexpr HostToDevice hostToDevice
Definition AMReX_GpuContainers.H:105
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:329
std::string const & name()
Definition AMReX_Machine.cpp:46
constexpr T powi(T x) noexcept
Return pow(x, Power), where Power is an integer known at compile time.
Definition AMReX_Math.H:377
constexpr int get_max_threads()
Definition AMReX_OpenMP.H:36
MPI_Comm CommunicatorSub() noexcept
sub-communicator for current frame
Definition AMReX_ParallelContext.H:70
int NProcsSub() noexcept
number of ranks in current frame
Definition AMReX_ParallelContext.H:74
static constexpr struct amrex::Scan::Type::Exclusive exclusive
static constexpr RetSum retSum
Definition AMReX_Scan.H:34
T PrefixSum(N n, FIN const &fin, FOUT const &fout, TYPE, RetSum a_ret_sum=retSum)
Definition AMReX_Scan.H:815
static constexpr int MPI_DATATYPE_NULL
Definition AMReX_ccse-mpi.H:56
int verbose
Definition AMReX.cpp:113
Definition AMReX_Amr.cpp:50
__host__ __device__ void ignore_unused(const Ts &...)
No-op helper that marks variables as intentionally unused.
Definition AMReX.H:273
AlgMGInterpType
Definition AMReX_AlgMG.H:220
@ mm_ext
Definition AMReX_AlgMG.H:220
@ mm_ext_i
Definition AMReX_AlgMG.H:220
@ C
Definition AMReX_Order.H:7
@ F
Definition AMReX_Order.H:7
@ none
Definition AMReX.H:558
void SpMV(Long nrows, Long ncols, T *__restrict__ py, CsrView< T const, Index > const &A, T const *__restrict__ px)
Perform y = A * x using CSR data (GPU/CPU aware).
Definition AMReX_SpMV.H:30
__host__ __device__ bool isnan(T x) noexcept
Return true if x is NaN. Works under fast math.
Definition AMReX_Math.H:117
AlgMGBottomSolver
Definition AMReX_AlgMG.H:210
@ jacobi
Definition AMReX_AlgMG.H:210
@ direct
Definition AMReX_AlgMG.H:210
@ gmres
Definition AMReX_AlgMG.H:210
@ bicgstab
Definition AMReX_AlgMG.H:210
AlgMGKrylovSolver
Definition AMReX_AlgMG.H:216
@ pcg
Definition AMReX_AlgMG.H:216
void computeResidual(AlgVector< T, AllocV > &res, SpMatrix< T, AllocM > const &A, AlgVector< T, AllocV > const &x, AlgVector< T, AllocV > const &b)
Compute the residual res = b - A * x.
Definition AMReX_SpMV.H:345
@ y
Definition AMReX_Orientation.H:14
@ x
Definition AMReX_Orientation.H:14
@ z
Definition AMReX_Orientation.H:14
@ smoother
Definition AMReX_MLLinOp.H:43
@ bicgstab
Definition AMReX_MLLinOp.H:43
__host__ __device__ constexpr IntVectND< dim > scale(const IntVectND< dim > &p, int s) noexcept
Returns a IntVectND obtained by multiplying each of the components of this IntVectND by s.
Definition AMReX_IntVect.H:1128
SpMatrix< T, Allocator > RAP(SpMatrix< T, Allocator > const &R, SpMatrix< T, Allocator > const &A, SpMatrix< T, Allocator > const &P, AlgPartition const &col_partition)
Galerkin product R (A P), with the result of SpGEMM(R, SpGEMM(A, P, col_partition),...
Definition AMReX_SpGEMM.H:1368
SpMatrix< T, Allocator > SpGEMM(SpMatrix< T, Allocator > const &A, SpMatrix< T, Allocator > const &B, AlgPartition const &col_partition, F const &row_post)
SpGEMM with a callback on each row of the product.
Definition AMReX_SpGEMM.H:1216
void Abort(const std::string &msg)
Print a fatal-error message to stderr and abort execution.
Definition AMReX.cpp:244
const int[]
Definition AMReX_BLProfiler.cpp:1665
@ bnorm
Definition AMReX_MLMG.H:23
CSR< T, V, IO > transpose(CSR< T, V, Index > const &csr, Long ncols)
Build the transpose CSR of csr.
Definition AMReX_SpMatUtil.H:25
AlgMGSmoother
Definition AMReX_AlgMG.H:229
@ chebyshev
Definition AMReX_AlgMG.H:229
@ l1_gauss_seidel
Definition AMReX_AlgMG.H:229
@ l1_jacobi
Definition AMReX_AlgMG.H:229
V< Index > col_index
Definition AMReX_CSR.H:62
void resize(Long num_rows, Long num_non_zeros)
Resize the storage to accommodate num_rows and num_non_zeros entries.
Definition AMReX_CSR.H:80
V< T > mat
Definition AMReX_CSR.H:61
V< Index > row_offset
Definition AMReX_CSR.H:63
Long nrows() const
Number of logical rows represented by the CSR offset array.
Definition AMReX_CSR.H:67
Definition AMReX_SpMatrix.H:39