20 postsmoother(nullptr),
21 profile_global(profile_global),
22 profile(
"MG level " +
std::to_string(param.level), false),
25 param_coarse(nullptr),
26 param_presmooth(nullptr),
27 param_postsmooth(nullptr),
28 param_coarse_solver(nullptr),
34 diracResidual(param.matResidual->Expose()),
35 diracSmoother(param.matSmooth->Expose()),
36 diracSmootherSloppy(param.matSmoothSloppy->Expose()),
37 diracCoarseResidual(nullptr),
38 diracCoarseSmoother(nullptr),
39 diracCoarseSmootherSloppy(nullptr),
40 matCoarseResidual(nullptr),
41 matCoarseSmoother(nullptr),
42 matCoarseSmootherSloppy(nullptr),
49 errorQuda(
"Level=%d is greater than limit of multigrid recursion depth", param.
level);
52 errorQuda(
"Cannot use preconditioned coarse grid solution without preconditioned smoother solve");
65 if (param.
B[0]->Nspin() == 1) csParam.
gammaBasis = param.
B[0]->GammaBasis();
76 rng =
new RNG(*param.
B[0], 1234);
84 for(
int i=0; i<(int)param.
B.size(); i++) {
87 param.
evals.push_back(0.0);
165 B_coarse =
new std::vector<ColorSpinorField*>();
171 for (
int i=0; i<nVec_coarse; i++)
388 for (
int i=0; i<4; i++) diracParam.
commDim[i] = schwarz ? 0 : 1;
398 for (
int i=0; i<4; i++) diracParam.
commDim[i] = schwarz ? 0 : 1;
479 vec_infile +=
"_level_";
480 vec_infile += std::to_string(
param.
level + 1);
481 vec_infile +=
"_defl_";
490 vec_outfile +=
"_level_";
491 vec_outfile += std::to_string(
param.
level + 1);
492 vec_outfile +=
"_defl_";
559 for (
int i = 0; i < nVec_coarse; i++)
560 if ((*
B_coarse)[i])
delete (*B_coarse)[i];
656 printfQuda(
"Vector %d: norms v_k = %e P^\\dagger v_k = %e P P^\\dagger v_k = %e, L2 relative deviation = %e\n",
658 if (deviation > tol)
errorQuda(
"L2 relative deviation for k=%d failed, %e > %e", i, deviation, tol);
663 sprintf(
prefix,
"MG level %d (%s): Null vector Oblique Projections : ",
param.
level + 1,
669 printfQuda(
"Checking 1 > || (1 - DP(P^dagDP)P^dag) v_k || / || v_k || for %d vectors\n",
param.
Nvec);
689 printfQuda(
"Checking 1 > || (1 - D P (P^\\dagger D P) P^\\dagger v_k || / || v_k || for %d vectors\n",
717 if (deviation > tol )
errorQuda(
"L2 relative deviation = %e > %e failed", deviation, tol);
743 #if 0 // enable to print out emulated and actual coarse-grid operator vectors for debugging 766 if(fabs(delta_factor) >
tol ) {
770 deviation = fabs(deviation);
774 printfQuda(
"L2 norms: Emulated = %e, Native = %e, relative deviation = %e\n",
norm2(*
x_coarse), r_nrm, deviation);
775 if (deviation > tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation, tol);
780 if (coarse_was_preconditioned) {
783 printfQuda(
"Checking Deo of preconditioned operator 0 = \\hat{D}_c - A^{-1} D_c\n");
792 if (deviation > tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation, tol);
796 printfQuda(
"Checking Doe of preconditioned operator 0 = \\hat{D}_c - A^{-1} D_c\n");
805 if (deviation > tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation, tol);
814 double deviation = std::fabs(dot.imag()) / std::fabs(dot.real());
816 real(dot), imag(dot), deviation);
817 if (deviation > tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation, tol);
824 double deviation = std::fabs(dot.imag()) / std::fabs(dot.real());
826 real(dot), imag(dot), deviation);
827 if (deviation > tol)
errorQuda(
"failed, deviation = %e (tol=%e)", deviation, tol);
832 sprintf(
prefix,
"MG level %d (%s): eigenvector overlap : ",
param.
level + 1,
852 printfQuda(
"L2 relative deviation = %e\n", deviation);
856 sprintf(
prefix,
"MG level %d (%s): eigenvector Oblique Projections : ",
param.
level + 1,
862 printfQuda(
"Checking 1 > || (1 - DP(P^dagDP)P^dag) v_k || / || v_k || for vector %d\n", i);
910 if (debug)
printfQuda(
"outer_solution_type = %d, inner_solution_type = %d\n", outer_solution_type, inner_solution_type);
913 errorQuda(
"Unsupported solution type combination");
916 errorQuda(
"For this coarse grid solution type, a preconditioned smoother is required");
944 bool use_solver_residual =
950 if (use_solver_residual) {
951 if (debug) r2 =
norm2(*
r);
955 else axpby(1.0, b, -1.0, *
r);
972 xpy(x_coarse_2_fine, solution);
1019 vec_infile +=
"_level_";
1021 vec_infile +=
"_nvec_";
1035 vec_outfile +=
"_level_";
1037 vec_outfile +=
"_nvec_";
1061 solverParam.
delta = 1e-1;
1098 if (schwarz_reset) {
1124 solverParam.
omega = 1.0;
1144 for(
int i=0; i<(int)B.size(); i++) {
1145 for (
int j=0; j<i; j++) {
1147 caxpy(-alpha, *B[j], *B[i]);
1149 double nrm2 =
norm2(*B[i]);
1150 if (nrm2 > 1e-16)
ax(1.0 /
sqrt(nrm2), *B[i]);
1151 else errorQuda(
"\nCannot normalize %u vector\n", i);
1156 for (
int i=0; i<(int)B.size(); i++) {
1170 (*solve)(*
out, *
in);
1179 for(
int i=0; i<(int)B.size(); i++) {
1180 for (
int j=0; j<i; j++) {
1182 caxpy(-alpha, *B[j], *B[i]);
1184 double nrm2 =
norm2(*B[i]);
1185 if (
sqrt(nrm2) > 1e-16)
ax(1.0/
sqrt(nrm2), *B[i]);
1186 else errorQuda(
"\nCannot normalize %u vector (nrm=%e)\n", i,
sqrt(nrm2));
1216 if (mdagm)
delete mdagm;
1217 if (mdagmSloppy)
delete mdagmSloppy;
1225 if (schwarz_reset) {
1244 const int Nvec = B.size();
1248 const int Ncolor = B[0]->Ncolor();
1249 const int Nspin = B[0]->Nspin();
1256 if (Nvec != 6)
errorQuda(
"\nError in MG::buildFreeVectors: Wilson-type fermions require Nvec = 6");
1259 printfQuda(
"Building %d free field vectors for Wilson-type fermions\n", Nvec);
1262 for (
int i = 0; i < Nvec; i++)
zero(*B[i]);
1270 for (
int c = 0; c <
Ncolor; c++) {
1271 for (
int s = 0;
s < 2;
s++) {
1273 xpy(*tmp, *B[counter]);
1275 xpy(*tmp, *B[counter]);
1281 }
else if (Nspin == 1)
1284 if (Nvec != 24)
errorQuda(
"\nError in MG::buildFreeVectors: Staggered-type fermions require Nvec = 24\n");
1287 printfQuda(
"Building %d free field vectors for Staggered-type fermions\n", Nvec);
1290 for (
int i = 0; i < Nvec; i++)
zero(*B[i]);
1298 for (
int c = 0; c < B[0]->Ncolor(); c++) {
1304 xpy(*tmp, *B[8 * c + 0]);
1306 xpy(*tmp, *B[8 * c + 0]);
1310 xpy(*tmp, *B[8 * c + 1]);
1312 xpy(*tmp, *B[8 * c + 1]);
1316 xpy(*tmp, *B[8 * c + 2]);
1318 xpy(*tmp, *B[8 * c + 2]);
1322 xpy(*tmp, *B[8 * c + 3]);
1324 xpy(*tmp, *B[8 * c + 3]);
1328 xpy(*tmp, *B[8 * c + 4]);
1330 xpy(*tmp, *B[8 * c + 4]);
1334 xpy(*tmp, *B[8 * c + 5]);
1336 xpy(*tmp, *B[8 * c + 5]);
1340 xpy(*tmp, *B[8 * c + 6]);
1342 xpy(*tmp, *B[8 * c + 6]);
1346 xpy(*tmp, *B[8 * c + 7]);
1348 xpy(*tmp, *B[8 * c + 7]);
1353 errorQuda(
"\nError in MG::buildFreeVectors: Unsupported combo of Nc %d, Nspin %d", Ncolor, Nspin);
1359 if (Nvec != Ncolor)
errorQuda(
"\nError in MG::buildFreeVectors: Coarse fermions require Nvec = Ncolor");
1364 for (
int i = 0; i < Nvec; i++)
zero(*B[i]);
1371 for (
int c = 0; c <
Ncolor; c++) {
1379 }
else if (Nspin == 1) {
1381 if (Nvec != Ncolor)
errorQuda(
"\nError in MG::buildFreeVectors: Coarse fermions require Nvec = Ncolor");
1386 for (
int i = 0; i < Nvec; i++)
zero(*B[i]);
1393 for (
int c = 0; c <
Ncolor; c++) {
1400 errorQuda(
"\nError in MG::buildFreeVectors: Unexpected Nspin = %d for coarse fermions", Nspin);
1406 for(
int i=0; i<(int)B.size(); i++) {
1407 double nrm2 =
norm2(*B[i]);
1408 if (nrm2 > 1e-16)
ax(1.0 /
sqrt(nrm2), *B[i]);
1409 else errorQuda(
"\nCannot normalize %u vector\n", i);
1428 std::vector<Complex> evals(nConv, 0.0);
1430 std::vector<ColorSpinorField *> B_evecs;
1441 for (
int i = 0; i < (int)
param.
B.size(); i++)
delete param.
B[i];
1444 if (!normop && !dagger) {
1447 (*eig_solve)(B_evecs, evals);
1450 }
else if (!normop && dagger) {
1453 (*eig_solve)(B_evecs, evals);
1456 }
else if (normop && !dagger) {
1459 (*eig_solve)(B_evecs, evals);
1462 }
else if (normop && dagger) {
1465 (*eig_solve)(B_evecs, evals);
1471 for (
int i = 0; i < (int)
param.
B.size(); i++) {
1473 *
param.
B[i] = *B_evecs[i];
1477 for (
auto b : B_evecs) {
delete b; }
cudaColorSpinorField * tmp2
bool mg_instance
whether to use a global or local (node) reduction for this solver
bool global_reduction
whether the solver acting as a preconditioner for another solver
QudaSchwarzType schwarz_type
void Init()
Initialize CURAND RNG states.
DiracMatrix * matSmoothSloppy
QudaBoolean global_reduction
void dumpNullVectors() const
Dump the null-space vectors to disk. Will recurse dumping all levels.
QudaMultigridParam & mg_global
void setPrecision(QudaPrecision precision, QudaPrecision ghost_precision=QUDA_INVALID_PRECISION, bool force_native=false)
QudaVerbosity verbosity_precondition
SolverParam * param_coarse_solver
cudaColorSpinorField * tmp1
std::vector< ColorSpinorField * > & B
QudaEigParam * eig_param[QUDA_MAX_MG_LEVEL]
enum QudaPrecision_s QudaPrecision
double setup_ca_lambda_max[QUDA_MAX_MG_LEVEL]
QudaFieldLocation location
void generateNullVectors(std::vector< ColorSpinorField *> &B, bool refresh=false)
Generate the null-space vectors.
virtual void Source(const QudaSourceType sourceType, const int st=0, const int s=0, const int c=0)=0
TimeProfile & profile_global
enum QudaResidualType_s QudaResidualType
QudaInverterType inv_type
const Dirac * diracResidual
QudaVerbosity getVerbosity()
QudaFieldLocation setup_location
QudaBoolean use_eig_solver[QUDA_MAX_MG_LEVEL]
void pushLevel(int level) const
Helper function called on entry to each MG function.
void createSmoother()
Create the smoothers.
ColorSpinorField * b_tilde
void destroyCoarseSolver()
Destroy the solver wrapper.
void generateEigenVectors()
Generate lowest eigenvectors.
void R(ColorSpinorField &out, const ColorSpinorField &in) const
QudaSolveType smoother_solve_type[QUDA_MAX_MG_LEVEL]
void destroySmoother()
Destroy the smoothers.
double setup_tol[QUDA_MAX_MG_LEVEL]
__host__ __device__ ValueType sqrt(ValueType x)
Complex cDotProduct(ColorSpinorField &, ColorSpinorField &)
int num_setup_iter[QUDA_MAX_MG_LEVEL]
void setOutputPrefix(const char *prefix)
Dirac * diracCoarseSmoother
DiracMatrix * matCoarseSmootherSloppy
SolverParam * param_presmooth
cudaColorSpinorField * tmp
const ColorSpinorField & Even() const
double coarse_solver_tol[QUDA_MAX_MG_LEVEL]
const ColorSpinorField & Odd() const
static ColorSpinorField * Create(const ColorSpinorParam ¶m)
static void loadVectors(std::vector< ColorSpinorField *> &eig_vecs, std::string file)
Load vectors from file.
QudaBoolean pre_orthonormalize
QudaSchwarzType smoother_schwarz_type[QUDA_MAX_MG_LEVEL]
QudaBoolean vec_load[QUDA_MAX_MG_LEVEL]
void P(ColorSpinorField &out, const ColorSpinorField &in) const
double setup_ca_lambda_min[QUDA_MAX_MG_LEVEL]
QudaPrecision HaloPrecision() const
double xmyNorm(ColorSpinorField &x, ColorSpinorField &y)
QudaPrecision smoother_halo_precision[QUDA_MAX_MG_LEVEL]
void loadVectors(std::vector< ColorSpinorField *> &B)
Load the null space vectors in from file.
QudaInverterType inv_type_precondition
void setHaloPrecision(QudaPrecision halo_precision_) const
QudaPreserveSource preserve_source
void reset(bool refresh=false)
This method resets the solver, e.g., when a parameter has changed such as the mass.
virtual double Mu() const
void setSiteSubset(QudaSiteSubset site_subset, QudaParity parity)
Sets whether the transfer operator is to act on full fields or single parity fields, and if single-parity which parity.
QudaBoolean setup_minimize_memory
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const =0
double mu_factor[QUDA_MAX_MG_LEVEL]
QudaSolveType smoother_solve_type
QudaSiteSubset siteSubset
QudaCABasis setup_ca_basis[QUDA_MAX_MG_LEVEL]
static void axpby(Float a, Float *x, Float b, Float *y, int len)
double norm2(const CloverField &a, bool inverse=false)
const Dirac * diracSmootherSloppy
void operator()(ColorSpinorField &out, ColorSpinorField &in)
int n_vec[QUDA_MAX_MG_LEVEL]
void popLevel(int level) const
Helper function called on exit to each MG member function.
QudaFieldLocation setup_location[QUDA_MAX_MG_LEVEL]
static void saveVectors(const std::vector< ColorSpinorField *> &eig_vecs, std::string file)
Save vectors to file.
DiracMatrix * matCoarseSmoother
void saveVectors(const std::vector< ColorSpinorField *> &B) const
Save the null space vectors in from file.
QudaComputeNullVector compute_null_vector
int coarse_solver_maxiter[QUDA_MAX_MG_LEVEL]
ColorSpinorField * tmp_coarse
void ax(const double &a, GaugeField &u)
Scale the gauge field by the scalar a.
__device__ __host__ void caxpy(const complex< Float > &a, const complex< Float > &x, complex< Float > &y)
QudaFieldLocation location
void buildFreeVectors(std::vector< ColorSpinorField *> &B)
Build free-field null-space vectors.
QudaBoolean post_orthonormalize
static EigenSolver * create(QudaEigParam *eig_param, const DiracMatrix &mat, TimeProfile &profile)
Creates the eigensolver using the parameters given and the matrix.
QudaResidualType residual_type
SolverParam * param_postsmooth
void pushOutputPrefix(const char *prefix)
Push a new output prefix onto the stack.
void updateInvertParam(QudaInvertParam ¶m, int offset=-1)
QudaPrecision cuda_prec_sloppy
void Release()
Release Device memory for CURAND RNG states.
int setup_maxiter[QUDA_MAX_MG_LEVEL]
const Dirac * diracSmoother
virtual void prepare(ColorSpinorField *&src, ColorSpinorField *&sol, ColorSpinorField &x, ColorSpinorField &b, const QudaSolutionType) const =0
int commDim[QUDA_MAX_DIM]
std::vector< Complex > evals
QudaInverterType smoother
virtual void DslashXpay(ColorSpinorField &out, const ColorSpinorField &in, const QudaParity parity, const ColorSpinorField &x, const double &k) const =0
static Solver * create(SolverParam ¶m, DiracMatrix &mat, DiracMatrix &matSloppy, DiracMatrix &matPrecon, TimeProfile &profile)
QudaSiteSubset SiteSubset() const
Class declaration to initialize and hold CURAND RNG states.
enum QudaMatPCType_s QudaMatPCType
QudaGammaBasis gammaBasis
int geoBlockSize[QUDA_MAX_DIM]
Dirac * diracCoarseResidual
enum QudaSolutionType_s QudaSolutionType
bool is_preconditioner
verbosity to use for preconditioner
QudaPrecision precision_null[QUDA_MAX_MG_LEVEL]
void popOutputPrefix()
Pop the output prefix restoring the prior one on the stack.
QudaVerbosity verbosity[QUDA_MAX_MG_LEVEL]
QudaInverterType coarse_solver[QUDA_MAX_MG_LEVEL]
std::complex< double > Complex
std::vector< ColorSpinorField * > * B_coarse
enum QudaParity_s QudaParity
ColorSpinorField * tmp2_coarse
ColorSpinorField * r_coarse
QudaInverterType coarse_solver[QUDA_MAX_MG_LEVEL]
#define MAX_BLOCK_FLOAT_NC
DiracMatrix * matCoarseResidual
QudaBoolean vec_store[QUDA_MAX_MG_LEVEL]
QudaPrecision halo_precision
void pushVerbosity(QudaVerbosity verbosity)
Push a new verbosity onto the stack.
QudaPrecision precision_precondition
QudaBoolean run_oblique_proj_check
Dirac * diracCoarseSmootherSloppy
void setCommDim(const int commDim_[QUDA_MAX_DIM]) const
Enable / disable communications for the Dirac operator.
QudaFieldLocation location[QUDA_MAX_MG_LEVEL]
enum QudaSiteSubset_s QudaSiteSubset
int smoother_schwarz_cycle[QUDA_MAX_MG_LEVEL]
char vec_outfile[QUDA_MAX_MG_LEVEL][256]
QudaBoolean run_low_mode_check
double coarse_solver_ca_lambda_max[QUDA_MAX_MG_LEVEL]
static int commDim[QUDA_MAX_DIM]
cpuColorSpinorField * out
QudaPrecision cuda_prec_precondition
QudaSolutionType coarse_grid_solution_type[QUDA_MAX_MG_LEVEL]
double flops() const
Return the total flops done on this and all coarser levels.
int geo_block_size[QUDA_MAX_MG_LEVEL][QUDA_MAX_DIM]
QudaBoolean generate_all_levels
#define QUDA_MAX_MG_LEVEL
Maximum number of multi-grid levels. This number may be increased if needed.
QudaPrecision Precision() const
int coarse_solver_ca_basis_size[QUDA_MAX_MG_LEVEL]
QudaCABasis coarse_solver_ca_basis[QUDA_MAX_MG_LEVEL]
void createCoarseDirac()
Create the coarse dirac operator.
char vec_infile[QUDA_MAX_MG_LEVEL][256]
void xpy(ColorSpinorField &x, ColorSpinorField &y)
QudaTwistFlavorType TwistFlavor() const
int setup_maxiter_refresh[QUDA_MAX_MG_LEVEL]
virtual void PrintVector(unsigned int x) const =0
QudaUseInitGuess use_init_guess
void popVerbosity()
Pop the verbosity restoring the prior one on the stack.
QudaPrecision precision_sloppy
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5...
MG(MGParam ¶m, TimeProfile &profile)
QudaComputeNullVector compute_null_vector
void createCoarseSolver()
Create the solver wrapper.
void mat(void *out, void **link, void *in, int dagger_bit, int mu, QudaPrecision sPrecision, QudaPrecision gPrecision)
virtual void reconstruct(ColorSpinorField &x, const ColorSpinorField &b, const QudaSolutionType) const =0
QudaPrecision Precision() const
QudaSolutionType coarse_grid_solution_type
int setup_ca_basis_size[QUDA_MAX_MG_LEVEL]
double coarse_solver_ca_lambda_min[QUDA_MAX_MG_LEVEL]
const ColorSpinorField & Vectors(QudaFieldLocation location=QUDA_INVALID_FIELD_LOCATION) const
QudaMultigridCycleType cycle_type
__device__ __host__ void zero(vector_type< scalar, n > &v)
QudaInverterType setup_inv_type[QUDA_MAX_MG_LEVEL]
const Dirac * Expose() const
DiracMatrix * matResidual
void spinorNoise(ColorSpinorField &src, RNG &randstates, QudaNoiseType type)
Generate a random noise spinor. This variant allows the user to manage the RNG state.
ColorSpinorField * x_coarse
QudaInvertParam * invert_param
static void dot(sFloat *res, gFloat *a, sFloat *b)
void reset()
for resetting the Transfer when the null vectors have changed