14 static bool debug =
false;
22 postsmoother(nullptr),
23 profile_global(profile_global),
24 profile(
"MG level " + std::to_string(
param.level), false),
27 param_coarse(nullptr),
28 param_presmooth(nullptr),
29 param_postsmooth(nullptr),
30 param_coarse_solver(nullptr),
38 diracResidual(
param.matResidual->Expose()),
39 diracSmoother(
param.matSmooth->Expose()),
40 diracSmootherSloppy(
param.matSmoothSloppy->Expose()),
41 diracCoarseResidual(nullptr),
42 diracCoarseSmoother(nullptr),
43 diracCoarseSmootherSloppy(nullptr),
44 matCoarseResidual(nullptr),
45 matCoarseSmoother(nullptr),
46 matCoarseSmootherSloppy(nullptr),
50 pushLevel(
param.level);
53 errorQuda(
"Level=%d is greater than limit of multigrid recursion depth",
param.level);
56 errorQuda(
"Cannot use preconditioned coarse grid solution without preconditioned smoother solve");
60 errorQuda(
"CPU solver location for MG is disabled");
93 for (
int i = 0; i < (int)
param.B.size(); i++) {
98 if (
param.mg_global.num_setup_iter[
param.level] > 0) {
100 && strcmp(
param.mg_global.vec_infile[
param.level],
"")
103 }
else if (
param.mg_global.use_eig_solver[
param.level]) {
109 }
else if (strcmp(
param.mg_global.vec_infile[
param.level],
"")
121 if (!transfer)
reset();
123 popLevel(
param.level);
127 pushLevel(param.
level);
159 if (resetTransfer || refresh) {
161 resetTransfer =
false;
191 B_coarse =
new std::vector<ColorSpinorField*>();
193 B_coarse->resize(nVec_coarse);
197 for (
int i=0; i<nVec_coarse; i++)
203 for (
int i=0; i<param.
Nvec; i++) {
204 zero(*(*B_coarse)[i]);
205 transfer->
R(*(*B_coarse)[i], *(param.
B[i]));
228 coarse->param.
delta = 1e-20;
231 coarse->param.
matSmooth = matCoarseSmoother;
233 coarse->
reset(refresh);
236 param_coarse =
new MGParam(param, *B_coarse, matCoarseResidual, matCoarseSmoother, matCoarseSmootherSloppy,
238 param_coarse->
fine =
this;
239 param_coarse->
delta = 1e-20;
242 coarse =
new MG(*param_coarse, profile_global);
257 popLevel(param.
level);
260 void MG::pushLevel(
int level)
const
267 void MG::popLevel(
int level)
const
276 pushLevel(param.
level);
280 presmoother =
nullptr;
283 if (param_presmooth) {
284 delete param_presmooth;
285 param_presmooth =
nullptr;
290 postsmoother =
nullptr;
293 if (param_postsmooth) {
294 delete param_postsmooth;
295 param_postsmooth =
nullptr;
298 popLevel(param.
level);
302 pushLevel(param.
level);
341 param_postsmooth =
new SolverParam(*param_presmooth);
359 popLevel(param.
level);
363 pushLevel(param.
level);
372 if (diracCoarseResidual)
delete diracCoarseResidual;
373 if (diracCoarseSmoother)
delete diracCoarseSmoother;
374 if (diracCoarseSmootherSloppy)
delete diracCoarseSmootherSloppy;
387 if (fine_dirac_type != dirac_type)
388 errorQuda(
"Input dirac type %d does not match smoother type %d\n", dirac_type, fine_dirac_type);
404 diracParamKD.
mass = diracSmoother->
Mass();
405 diracParamKD.
mu = diracSmoother->
Mu();
410 diracParamKD.
xInvKD = xInvKD;
412 diracParamKD.
tmp1 = tmp_coarse;
413 diracParamKD.
tmp2 = tmp2_coarse;
432 errorQuda(
"Invalid dirac_type %d", dirac_type);
442 diracParam.
dirac = preconditioned_coarsen ?
const_cast<Dirac *
>(diracSmoother) :
const_cast<Dirac *
>(diracResidual);
443 diracParam.
kappa = (param.
B[0]->Nspin() == 1) ?
453 for (
int i = 0; i <= param.
level; i++) {
463 diracParam.
tmp1 = tmp_coarse;
464 diracParam.
tmp2 = tmp2_coarse;
477 diracParam.
tmp1 = &(tmp_coarse->
Even());
478 diracParam.
tmp2 = &(tmp2_coarse->
Even());
482 for (
int i = 0; i < 4; i++) diracParam.
commDim[i] = schwarz ? 0 : 1;
487 diracParam.
tmp1 = tmp_coarse;
488 diracParam.
tmp2 = tmp2_coarse;
492 for (
int i = 0; i < 4; i++) diracParam.
commDim[i] = schwarz ? 0 : 1;
494 diracCoarseSmootherSloppy =
new DiracCoarse(
static_cast<DiracCoarse &
>(*diracCoarseSmoother), diracParam);
498 if (matCoarseResidual)
delete matCoarseResidual;
499 if (matCoarseSmoother)
delete matCoarseSmoother;
500 if (matCoarseSmootherSloppy)
delete matCoarseSmootherSloppy;
501 matCoarseResidual =
new DiracM(*diracCoarseResidual);
502 matCoarseSmoother =
new DiracM(*diracCoarseSmoother);
503 matCoarseSmootherSloppy =
new DiracM(*diracCoarseSmootherSloppy);
507 popLevel(param.
level);
511 pushLevel(param.
level);
517 auto &coarse_solver_inner =
reinterpret_cast<PreconditionedSolver *
>(coarse_solver)->ExposeSolver();
523 coarse_solver_inner.extractDeflationSpace(
evecs);
525 delete coarse_solver;
526 coarse_solver =
nullptr;
528 if (param_coarse_solver) {
529 delete param_coarse_solver;
530 param_coarse_solver =
nullptr;
536 popLevel(param.
level);
540 pushLevel(param.
level);
546 coarse_solver = coarse;
575 errorQuda(
"Coarse grid deflation not supported with coarse solver %d", param_coarse_solver->
inv_type);
582 vec_infile +=
"_level_";
583 vec_infile += std::to_string(param.
level + 1);
584 vec_infile +=
"_defl_";
593 vec_outfile +=
"_level_";
594 vec_outfile += std::to_string(param.
level + 1);
595 vec_outfile +=
"_defl_";
604 param_coarse_solver->
delta = 1e-8;
632 *matCoarseSmoother, *matCoarseSmoother, profile);
633 sprintf(coarse_prefix,
"MG level %d (%s): ", param.
level + 1,
639 *matCoarseResidual, *matCoarseResidual, profile);
640 sprintf(coarse_prefix,
"MG level %d (%s): ", param.
level + 1,
642 coarse_solver =
new PreconditionedSolver(*solver, *matCoarseResidual->
Expose(), *param_coarse_solver, profile, coarse_prefix);
650 int defl_size =
evecs.size();
651 auto &coarse_solver_inner =
reinterpret_cast<PreconditionedSolver *
>(coarse_solver)->ExposeSolver();
656 coarse_solver_inner.setRecomputeEvals(
true);
658 printfQuda(
"Transferring deflation space size %d to coarse solver\n", defl_size);
660 coarse_solver_inner.injectDeflationSpace(
evecs);
666 param_coarse_solver->
maxiter = 1;
667 (*coarse_solver)(*x_coarse, *r_coarse);
678 popLevel(param.
level);
683 pushLevel(param.
level);
686 if (coarse)
delete coarse;
688 if (coarse_solver)
delete coarse_solver;
689 if (param_coarse_solver)
delete param_coarse_solver;
694 for (
int i = 0; i < nVec_coarse; i++)
695 if ((*B_coarse)[i])
delete (*B_coarse)[i];
698 if (transfer)
delete transfer;
699 if (matCoarseSmootherSloppy)
delete matCoarseSmootherSloppy;
700 if (diracCoarseSmootherSloppy)
delete diracCoarseSmootherSloppy;
701 if (matCoarseSmoother)
delete matCoarseSmoother;
702 if (diracCoarseSmoother)
delete diracCoarseSmoother;
703 if (matCoarseResidual)
delete matCoarseResidual;
704 if (diracCoarseResidual)
delete diracCoarseResidual;
705 if (postsmoother)
delete postsmoother;
706 if (param_postsmooth)
delete param_postsmooth;
714 if (presmoother)
delete presmoother;
715 if (param_presmooth)
delete param_presmooth;
719 if (r_coarse)
delete r_coarse;
720 if (x_coarse)
delete x_coarse;
721 if (tmp_coarse)
delete tmp_coarse;
722 if (tmp2_coarse)
delete tmp2_coarse;
724 if (xInvKD)
delete xInvKD;
726 if (param_coarse)
delete param_coarse;
730 popLevel(param.
level);
737 if (param_coarse_solver) {
739 param_coarse_solver->
gflops = 0;
744 if (param_presmooth) {
746 param_presmooth->
gflops = 0;
749 if (param_postsmooth) {
751 param_postsmooth->
gflops = 0;
766 pushLevel(param.
level);
787 printfQuda(
"Checking 0 = (1 - P P^\\dagger) v_k for %d vectors\n", param.
Nvec);
789 for (
int i = 0; i < param.
Nvec; i++) {
793 transfer->
R(*r_coarse, *tmp1);
794 transfer->
P(*tmp2, *r_coarse);
799 "Vector %d: norms v_k = %e P^\\dagger v_k = %e (1 - P P^\\dagger) v_k = %e, L2 relative deviation = %e\n",
801 if (deviation >
tol)
errorQuda(
"L2 relative deviation for k=%d failed, %e > %e", i, deviation,
tol);
806 sprintf(prefix,
"MG level %d (%s): Null vector Oblique Projections : ", param.
level + 1,
812 printfQuda(
"Checking 1 > || (1 - DP(P^dagDP)P^dag) v_k || / || v_k || for %d vectors\n", param.
Nvec);
814 for (
int i = 0; i < param.
Nvec; i++) {
815 transfer->
R(*r_coarse, *(param.
B[i]));
816 (*coarse_solver)(*x_coarse, *r_coarse);
818 transfer->
P(*tmp2, *x_coarse);
820 *tmp2 = *(param.
B[i]);
826 sprintf(prefix,
"MG level %d (%s): ", param.
level + 1,
835 printfQuda(
"Checking 1 > || (1 - D P (P^\\dagger D P) P^\\dagger v_k || / || v_k || for %d vectors\n",
838 for (
int i=0; i<param.
Nvec; i++) {
839 transfer->
R(*r_coarse, *(param.
B[i]));
840 (*coarse)(*x_coarse, *r_coarse);
842 transfer->
P(*tmp2, *x_coarse);
844 *tmp2 = *(param.
B[i]);
878 transfer->
P(*tmp2, *x_coarse);
879 transfer->
R(*r_coarse, *tmp2);
885 if (deviation >
tol )
errorQuda(
"L2 relative deviation = %e > %e failed", deviation,
tol);
903 transfer->
P(*tmp1, *tmp_coarse);
907 double mass = diracResidual->
Mass();
908 if (param.
level == 0) {
909 if (tmp1->
Nspin() == 4) {
912 }
else if (tmp1->
Nspin() == 2) {
929 transfer->
R(*x_coarse, *tmp2);
930 static_cast<DiracCoarse *
>(diracCoarseResidual)->
M(*r_coarse, *tmp_coarse);
935 for (
unsigned int i=0; i<
comm_size(); i++) {
948 double r_nrm =
norm2(*r_coarse);
951 if (diracResidual->
Mu() != 0.0) {
955 if (fabs(delta_factor) >
tol) {
958 deviation -= fabs(delta_a) *
sqrt(
norm2(*tmp_coarse) /
norm2(*x_coarse));
959 deviation = fabs(deviation);
963 printfQuda(
"L2 norms: Emulated = %e, Native = %e, relative deviation = %e\n",
norm2(*x_coarse), r_nrm, deviation);
965 if (deviation >
tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation,
tol);
973 if (coarse_was_preconditioned) {
976 printfQuda(
"Checking Deo of preconditioned operator 0 = \\hat{D}_c - A^{-1} D_c\n");
983 printfQuda(
"L2 norms: Emulated = %e, Native = %e, relative deviation = %e\n",
norm2(x_coarse->
Even()), r_nrm,
985 if (deviation >
tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation,
tol);
989 printfQuda(
"Checking Doe of preconditioned operator 0 = \\hat{D}_c - A^{-1} D_c\n");
996 printfQuda(
"L2 norms: Emulated = %e, Native = %e, relative deviation = %e\n",
norm2(x_coarse->
Odd()), r_nrm,
998 if (deviation >
tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation,
tol);
1005 if (tmp2->
Nspin() == 1) {
1006 diracSmoother->
M(tmp2->
Even(), tmp1->
Odd());
1011 double deviation = std::fabs(dot.imag()) / std::fabs(dot.real());
1013 real(dot), imag(dot), deviation);
1014 if (deviation >
tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation,
tol);
1020 diracResidual->
MdagM(*tmp2, *tmp1);
1023 diracResidual->
M(*tmp2, *tmp1);
1026 double deviation = std::fabs(dot.imag()) / std::fabs(dot.real());
1028 real(dot), imag(dot), deviation);
1029 if (deviation >
tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation,
tol);
1036 sprintf(prefix,
"MG level %d (%s): eigenvector overlap : ", param.
level + 1,
1044 for (
int i = 0; i < param.
Nvec; i++) {
1047 transfer->
R(*r_coarse, *param.
B[i]);
1049 transfer->
P(*tmp2, *r_coarse);
1051 printfQuda(
"Vector %d: norms v_k = %e P^dag v_k = %e PP^dag v_k = %e\n", i,
norm2(*param.
B[i]),
1056 printfQuda(
"L2 relative deviation = %e\n", deviation);
1060 sprintf(prefix,
"MG level %d (%s): eigenvector Oblique Projections : ", param.
level + 1,
1066 printfQuda(
"Checking 1 > || (1 - DP(P^dagDP)P^dag) v_k || / || v_k || for vector %d\n", i);
1068 transfer->
R(*r_coarse, *param.
B[i]);
1069 (*coarse_solver)(*x_coarse, *r_coarse);
1071 transfer->
P(*tmp2, *x_coarse);
1075 printfQuda(
"Vector %d: norms v_k %e DP(P^dagDP)P^dag v_k %e\n", i,
norm2(*param.
B[i]),
norm2(*tmp1));
1080 sprintf(prefix,
"MG level %d (%s): ", param.
level + 1,
1092 popLevel(param.
level);
1114 if (debug)
printfQuda(
"outer_solution_type = %d, inner_solution_type = %d\n", outer_solution_type, inner_solution_type);
1117 errorQuda(
"Unsupported solution type combination");
1120 errorQuda(
"For this coarse grid solution type, a preconditioned smoother is required");
1137 diracSmoother->
prepare(in, out, x, residual, outer_solution_type);
1143 if (presmoother) (*presmoother)(*out, *in);
else zero(*out);
1146 diracSmoother->
reconstruct(solution, b, inner_solution_type);
1150 bool use_solver_residual
1158 if (use_solver_residual) {
1159 if (debug) r2 =
norm2(*r);
1165 axpby(1.0, b, -1.0, *r);
1173 transfer->
R(*r_coarse, residual);
1174 if ( debug )
printfQuda(
"after pre-smoothing x2 = %e, r2 = %e, r_coarse2 = %e\n",
norm2(x), r2,
norm2(*r_coarse));
1177 (*coarse_solver)(*x_coarse, *r_coarse);
1178 if (debug)
printfQuda(
"after coarse solve x_coarse2 = %e r_coarse2 = %e\n",
norm2(*x_coarse),
norm2(*r_coarse));
1182 transfer->
P(x_coarse_2_fine, *x_coarse);
1183 xpy(x_coarse_2_fine, solution);
1190 if (debug)
printfQuda(
"preparing to post smooth\n");
1204 if (postsmoother) (*postsmoother)(*out, *in);
1206 if (debug)
printfQuda(
"exited postsmooth, about to reconstruct\n");
1208 diracSmoother->
reconstruct(x, b, outer_solution_type);
1210 if (debug)
printfQuda(
"finished reconstruct\n");
1215 diracSmoother->
prepare(in, out, x, b, outer_solution_type);
1216 if (presmoother) (*presmoother)(*out, *in);
1217 diracSmoother->
reconstruct(x, b, outer_solution_type);
1234 warningQuda(
"Cannot load near-null vectors for top level of staggered MG solve.");
1239 pushLevel(param.
level);
1241 vec_infile +=
"_level_";
1242 vec_infile += std::to_string(param.
level);
1243 vec_infile +=
"_nvec_";
1247 popLevel(param.
level);
1256 warningQuda(
"Cannot save near-null vectors for top level of staggered MG solve.");
1261 pushLevel(param.
level);
1263 vec_outfile +=
"_level_";
1264 vec_outfile += std::to_string(param.
level);
1265 vec_outfile +=
"_nvec_";
1269 popLevel(param.
level);
1278 warningQuda(
"Cannot dump near-null vectors for top level of staggered MG solve.");
1287 pushLevel(param.
level);
1295 solverParam.
delta = 1e-1;
1311 if (param.
level == 0) {
1333 if (schwarz_reset) {
1348 solve =
Solver::create(solverParam, *mdagm, *mdagmSloppy, *mdagmSloppy, *mdagmSloppy, profile);
1359 solverParam.
omega = 1.0;
1376 printfQuda(
"Running vectors setup on level %d iter %d of %d\n", param.
level, si + 1,
1381 for(
int i=0; i<(int)B.size(); i++) {
1382 for (
int j=0; j<i; j++) {
1384 caxpy(-alpha, *B[j], *B[i]);
1386 double nrm2 =
norm2(*B[i]);
1387 if (nrm2 > 1e-16)
ax(1.0 /
sqrt(nrm2), *B[i]);
1388 else errorQuda(
"\nCannot normalize %u vector\n", i);
1393 for (
int i=0; i<(int)B.size(); i++) {
1407 (*solve)(*out, *in);
1416 for(
int i=0; i<(int)B.size(); i++) {
1417 for (
int j=0; j<i; j++) {
1419 caxpy(-alpha, *B[j], *B[i]);
1421 double nrm2 =
norm2(*B[i]);
1422 if (
sqrt(nrm2) > 1e-16)
ax(1.0/
sqrt(nrm2), *B[i]);
1423 else errorQuda(
"\nCannot normalize %u vector (nrm=%e)\n", i,
sqrt(nrm2));
1430 resetTransfer =
true;
1437 for (
int i=0; i<param.
Nvec; i++) {
1438 zero(*(*B_coarse)[i]);
1439 transfer->
R(*(*B_coarse)[i], *(param.
B[i]));
1442 coarse->resetTransfer =
true;
1453 if (mdagm)
delete mdagm;
1454 if (mdagmSloppy)
delete mdagmSloppy;
1462 if (schwarz_reset) {
1473 popLevel(param.
level);
1480 pushLevel(param.
level);
1481 const int Nvec = B.size();
1485 const int Ncolor = B[0]->Ncolor();
1486 const int Nspin = B[0]->Nspin();
1493 if (Nvec != 6)
errorQuda(
"\nError in MG::buildFreeVectors: Wilson-type fermions require Nvec = 6");
1496 printfQuda(
"Building %d free field vectors for Wilson-type fermions\n", Nvec);
1499 for (
int i = 0; i < Nvec; i++)
zero(*B[i]);
1507 for (
int c = 0; c <
Ncolor; c++) {
1508 for (
int s = 0; s < 2; s++) {
1518 }
else if (
Nspin == 1) {
1521 if (Nvec != 24)
errorQuda(
"\nError in MG::buildFreeVectors: Staggered-type fermions require Nvec = 24\n");
1524 printfQuda(
"Building %d free field vectors for Staggered-type fermions\n", Nvec);
1527 for (
int i = 0; i < Nvec; i++)
zero(*B[i]);
1535 for (
int c = 0; c < B[0]->Ncolor(); c++) {
1540 xpy(*
tmp, *B[8 * c + 0]);
1542 xpy(*
tmp, *B[8 * c + 0]);
1546 xpy(*
tmp, *B[8 * c + 1]);
1548 xpy(*
tmp, *B[8 * c + 1]);
1552 xpy(*
tmp, *B[8 * c + 2]);
1554 xpy(*
tmp, *B[8 * c + 2]);
1558 xpy(*
tmp, *B[8 * c + 3]);
1560 xpy(*
tmp, *B[8 * c + 3]);
1564 xpy(*
tmp, *B[8 * c + 4]);
1566 xpy(*
tmp, *B[8 * c + 4]);
1570 xpy(*
tmp, *B[8 * c + 5]);
1572 xpy(*
tmp, *B[8 * c + 5]);
1576 xpy(*
tmp, *B[8 * c + 6]);
1578 xpy(*
tmp, *B[8 * c + 6]);
1582 xpy(*
tmp, *B[8 * c + 7]);
1584 xpy(*
tmp, *B[8 * c + 7]);
1589 errorQuda(
"\nError in MG::buildFreeVectors: Unsupported combo of Nc %d, Nspin %d",
Ncolor,
Nspin);
1594 if (Nvec !=
Ncolor)
errorQuda(
"\nError in MG::buildFreeVectors: Coarse fermions require Nvec = Ncolor");
1599 for (
int i = 0; i < Nvec; i++)
zero(*B[i]);
1606 for (
int c = 0; c <
Ncolor; c++) {
1614 }
else if (
Nspin == 1) {
1616 if (Nvec !=
Ncolor)
errorQuda(
"\nError in MG::buildFreeVectors: Coarse fermions require Nvec = Ncolor");
1621 for (
int i = 0; i < Nvec; i++)
zero(*B[i]);
1628 for (
int c = 0; c <
Ncolor; c++) {
1635 errorQuda(
"\nError in MG::buildFreeVectors: Unexpected Nspin = %d for coarse fermions",
Nspin);
1641 for(
int i=0; i<(int)B.size(); i++) {
1642 double nrm2 =
norm2(*B[i]);
1643 if (nrm2 > 1e-16)
ax(1.0 /
sqrt(nrm2), *B[i]);
1644 else errorQuda(
"\nCannot normalize %u vector\n", i);
1650 popLevel(param.
level);
1655 pushLevel(param.
level);
1663 std::vector<Complex>
evals(n_conv, 0.0);
1665 std::vector<ColorSpinorField *> B_evecs;
1676 for (
int i = 0; i < (int)param.
B.size(); i++)
delete param.
B[i];
1679 if (!normop && !
dagger) {
1682 (*eig_solve)(B_evecs,
evals);
1685 }
else if (!normop &&
dagger) {
1688 (*eig_solve)(B_evecs,
evals);
1691 }
else if (normop && !
dagger) {
1694 (*eig_solve)(B_evecs,
evals);
1697 }
else if (normop &&
dagger) {
1700 (*eig_solve)(B_evecs,
evals);
1706 for (
int i = 0; i < (int)param.
B.size(); i++) {
1708 *param.
B[i] = *B_evecs[i];
1712 for (
auto b : B_evecs) {
delete b; }
1717 popLevel(param.
level);
virtual void PrintVector(unsigned int x) const =0
const ColorSpinorField & Odd() const
QudaTwistFlavorType TwistFlavor() const
QudaSiteSubset SiteSubset() const
static ColorSpinorField * Create(const ColorSpinorParam ¶m)
const ColorSpinorField & Even() const
virtual void Source(const QudaSourceType sourceType, const int st=0, const int s=0, const int c=0)=0
virtual void M(ColorSpinorField &out, const ColorSpinorField &in) const =0
Apply M for the dirac op. E.g. the Schur Complement operator.
double Kappa() const
accessor for Kappa (mass parameter)
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const =0
Apply MdagM operator which may be optimized.
virtual double Mu() const
accessor for twist parameter – overrride can return better value
virtual double Mass() const
accessor for Mass (in case of a factor of 2 for staggered)
virtual void prefetch(QudaFieldLocation mem_space, qudaStream_t stream=0) const
If managed memory and prefetch is enabled, prefetch the gauge field and temporary spinors to the CPU ...
void setHaloPrecision(QudaPrecision halo_precision_) const
QudaPrecision HaloPrecision() const
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType solType) const =0
virtual QudaDiracType getDiracType() const =0
returns the Dirac type
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType solType) const =0
void setCommDim(const int commDim_[QUDA_MAX_DIM]) const
Enable / disable communications for the Dirac operator.
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const =0
Xpay version of Dslash.
const Dirac * Expose() const
int commDim[QUDA_MAX_DIM]
cudaGaugeField * longGauge
cudaGaugeField * fatGauge
QudaPrecision halo_precision
This is the generic driver for launching Dslash kernels (the base kernel of which is defined in dslas...
static EigenSolver * create(QudaEigParam *eig_param, const DiracMatrix &mat, TimeProfile &profile)
Creates the eigensolver using the parameters given and the matrix.
QudaPrecision Precision() const
double flops() const
Return the total flops done on this and all coarser levels.
MG(MGParam ¶m, TimeProfile &profile)
void operator()(ColorSpinorField &out, ColorSpinorField &in)
void destroySmoother()
Destroy the smoothers.
void createCoarseDirac()
Create the coarse dirac operator.
void createSmoother()
Create the smoothers.
void saveVectors(const std::vector< ColorSpinorField * > &B) const
Save the null space vectors in from file.
void buildFreeVectors(std::vector< ColorSpinorField * > &B)
Build free-field null-space vectors.
void createCoarseSolver()
Create the solver wrapper.
void verify(bool recursively=false)
Verify the correctness of the MG method, optionally recursively starting from the top down.
void dumpNullVectors() const
Dump the null-space vectors to disk. Will recurse dumping all levels.
void loadVectors(std::vector< ColorSpinorField * > &B)
Load the null space vectors in from file.
void generateEigenVectors()
Generate lowest eigenvectors.
void destroyCoarseSolver()
Destroy the solver wrapper.
void reset(bool refresh=false)
This method resets the solver, e.g., when a parameter has changed such as the mass.
void generateNullVectors(std::vector< ColorSpinorField * > &B, bool refresh=false)
Generate the null-space vectors.
Class declaration to initialize and hold CURAND RNG states.
std::vector< ColorSpinorField * > evecs
int deflationSpaceSize() const
Returns the size of deflation space.
void setDeflateCompute(bool flag)
Sets the deflation compute boolean.
std::vector< Complex > evals
static Solver * create(SolverParam ¶m, const DiracMatrix &mat, const DiracMatrix &matSloppy, const DiracMatrix &matPrecon, const DiracMatrix &matEig, TimeProfile &profile)
Solver factory.
bool isRunning(QudaProfileType idx)
void reset()
for resetting the Transfer when the null vectors have changed
void R(ColorSpinorField &out, const ColorSpinorField &in) const
void setSiteSubset(QudaSiteSubset site_subset, QudaParity parity)
Sets whether the transfer operator is to act on full fields or single parity fields,...
void P(ColorSpinorField &out, const ColorSpinorField &in) const
const ColorSpinorField & Vectors(QudaFieldLocation location=QUDA_INVALID_FIELD_LOCATION) const
VectorIO is a simple wrapper class for loading and saving sets of vector fields using QIO.
void load(std::vector< ColorSpinorField * > &vecs)
Load vectors from filename.
void save(const std::vector< ColorSpinorField * > &vecs)
Save vectors to filename.
quda::mgarray< QudaInverterType > coarse_solver
quda::mgarray< QudaSolveType > smoother_solve_type
cudaColorSpinorField * tmp
@ QUDA_MG_CYCLE_RECURSIVE
enum QudaPrecision_s QudaPrecision
@ QUDA_CUDA_FIELD_LOCATION
@ QUDA_CPU_FIELD_LOCATION
@ QUDA_USE_INIT_GUESS_YES
@ QUDA_PARITY_SITE_SUBSET
@ QUDA_DEGRAND_ROSSI_GAMMA_BASIS
@ QUDA_L2_RELATIVE_RESIDUAL
@ QUDA_TRANSFER_COARSE_KD
@ QUDA_TRANSFER_AGGREGATE
@ QUDA_TRANSFER_OPTIMIZED_KD
enum QudaSiteSubset_s QudaSiteSubset
enum QudaSolutionType_s QudaSolutionType
@ QUDA_MATPC_EVEN_EVEN_ASYMMETRIC
enum QudaMatPCType_s QudaMatPCType
enum QudaResidualType_s QudaResidualType
@ QUDA_PRESERVE_SOURCE_NO
enum QudaParity_s QudaParity
@ QUDA_COMPUTE_NULL_VECTOR_YES
double xmyNorm(ColorSpinorField &x, ColorSpinorField &y)
void caxpy(const Complex &a, ColorSpinorField &x, ColorSpinorField &y)
void xpy(ColorSpinorField &x, ColorSpinorField &y)
Complex cDotProduct(ColorSpinorField &, ColorSpinorField &)
void axpby(double a, ColorSpinorField &x, double b, ColorSpinorField &y)
cudaGaugeField * AllocateAndBuildStaggeredKahlerDiracInverse(const cudaGaugeField &gauge, const double mass, const QudaPrecision override_prec)
Allocate and build the Kahler-Dirac inverse block for KD operators.
double norm2(const CloverField &a, bool inverse=false)
__device__ __host__ void zero(double &a)
void ax(const double &a, GaugeField &u)
Scale the gauge field by the scalar a.
std::complex< double > Complex
__host__ __device__ ValueType sqrt(ValueType x)
void spinorNoise(ColorSpinorField &src, RNG &randstates, QudaNoiseType type)
Generate a random noise spinor. This variant allows the user to manage the RNG state.
#define QUDA_MAX_MG_LEVEL
Maximum number of multi-grid levels. This number may be increased if needed.
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5.
QudaFieldLocation location
QudaPrecision cuda_prec_sloppy
QudaPrecision cuda_prec_sloppy
QudaPrecision cuda_prec_precondition
QudaBoolean pre_orthonormalize
int smoother_schwarz_cycle[QUDA_MAX_MG_LEVEL]
double coarse_solver_ca_lambda_min[QUDA_MAX_MG_LEVEL]
double setup_ca_lambda_max[QUDA_MAX_MG_LEVEL]
QudaSolutionType coarse_grid_solution_type[QUDA_MAX_MG_LEVEL]
double coarse_solver_tol[QUDA_MAX_MG_LEVEL]
int setup_maxiter_refresh[QUDA_MAX_MG_LEVEL]
QudaPrecision precision_null[QUDA_MAX_MG_LEVEL]
char vec_infile[QUDA_MAX_MG_LEVEL][256]
QudaBoolean use_eig_solver[QUDA_MAX_MG_LEVEL]
int coarse_solver_maxiter[QUDA_MAX_MG_LEVEL]
QudaEigParam * eig_param[QUDA_MAX_MG_LEVEL]
int n_vec[QUDA_MAX_MG_LEVEL]
int geo_block_size[QUDA_MAX_MG_LEVEL][QUDA_MAX_DIM]
QudaInverterType coarse_solver[QUDA_MAX_MG_LEVEL]
int coarse_solver_ca_basis_size[QUDA_MAX_MG_LEVEL]
QudaTransferType transfer_type[QUDA_MAX_MG_LEVEL]
QudaFieldLocation setup_location[QUDA_MAX_MG_LEVEL]
QudaBoolean post_orthonormalize
int num_setup_iter[QUDA_MAX_MG_LEVEL]
double mu_factor[QUDA_MAX_MG_LEVEL]
QudaCABasis setup_ca_basis[QUDA_MAX_MG_LEVEL]
QudaInverterType setup_inv_type[QUDA_MAX_MG_LEVEL]
int setup_maxiter[QUDA_MAX_MG_LEVEL]
int setup_ca_basis_size[QUDA_MAX_MG_LEVEL]
QudaBoolean preserve_deflation
double coarse_solver_ca_lambda_max[QUDA_MAX_MG_LEVEL]
char vec_outfile[QUDA_MAX_MG_LEVEL][256]
QudaBoolean run_low_mode_check
double setup_ca_lambda_min[QUDA_MAX_MG_LEVEL]
QudaCABasis coarse_solver_ca_basis[QUDA_MAX_MG_LEVEL]
QudaBoolean vec_store[QUDA_MAX_MG_LEVEL]
QudaSolveType smoother_solve_type[QUDA_MAX_MG_LEVEL]
QudaFieldLocation location[QUDA_MAX_MG_LEVEL]
QudaVerbosity verbosity[QUDA_MAX_MG_LEVEL]
QudaBoolean setup_minimize_memory
QudaBoolean generate_all_levels
QudaSchwarzType smoother_schwarz_type[QUDA_MAX_MG_LEVEL]
QudaBoolean run_oblique_proj_check
QudaInvertParam * invert_param
double setup_tol[QUDA_MAX_MG_LEVEL]
QudaPrecision smoother_halo_precision[QUDA_MAX_MG_LEVEL]
QudaBoolean vec_load[QUDA_MAX_MG_LEVEL]
DiracMatrix * matResidual
std::vector< ColorSpinorField * > & B
QudaInverterType smoother
QudaSolutionType coarse_grid_solution_type
DiracMatrix * matSmoothSloppy
QudaFieldLocation location
QudaBoolean global_reduction
QudaMultigridCycleType cycle_type
int geoBlockSize[QUDA_MAX_DIM]
QudaMultigridParam & mg_global
QudaSolveType smoother_solve_type
QudaTransferType transfer_type
QudaFieldLocation setup_location
QudaPreserveSource preserve_source
QudaComputeNullVector compute_null_vector
bool is_preconditioner
verbosity to use for preconditioner
QudaSchwarzType schwarz_type
QudaResidualType residual_type
QudaPrecision precision_precondition
QudaPrecision precision_sloppy
bool mg_instance
whether to use a global or local (node) reduction for this solver
QudaInverterType inv_type
QudaVerbosity verbosity_precondition
QudaUseInitGuess use_init_guess
QudaInverterType inv_type_precondition
void updateInvertParam(QudaInvertParam ¶m, int offset=-1)
bool global_reduction
whether the solver acting as a preconditioner for another solver
void pushVerbosity(QudaVerbosity verbosity)
Push a new verbosity onto the stack.
void popOutputPrefix()
Pop the output prefix restoring the prior one on the stack.
void popVerbosity()
Pop the verbosity restoring the prior one on the stack.
void pushOutputPrefix(const char *prefix)
Push a new output prefix onto the stack.
QudaVerbosity getVerbosity()
void setOutputPrefix(const char *prefix)