15 profile_global(profile_global),
16 profile(
"MG level " +
std::to_string(
param.
level+1), false ),
17 coarse(nullptr), fine(
param.fine), coarse_solver(nullptr),
18 param_coarse(nullptr), param_presmooth(nullptr), param_postsmooth(nullptr),
19 r(nullptr), r_coarse(nullptr), x_coarse(nullptr), tmp_coarse(nullptr),
20 diracCoarseResidual(nullptr), diracCoarseSmoother(nullptr), matCoarseResidual(nullptr), matCoarseSmoother(nullptr) {
42 errorQuda(
"Cannot use preconditioned coarse grid solution without preconditioned smoother solve");
71 printfQuda(
"start creating transfer operator\n");
77 printfQuda(
"end creating transfer operator\n");
119 printfQuda(
"Creating coarse null-space vectors\n");
120 B_coarse =
new std::vector<ColorSpinorField*>();
124 for (
int i=0;
i<nVec_coarse;
i++)
136 printfQuda(
"Creating next multigrid level\n");
148 printfQuda(
"Assigned coarse solver to coarse MG operator\n");
181 printfQuda(
"Assigned coarse solver to preconditioned GCR solver\n");
277 for (
int i=0;
i<nVec_coarse;
i++)
if ((*
B_coarse)[
i])
delete (*B_coarse)[
i];
348 printfQuda(
"Vector %d: norms v_k = %e P^\\dagger v_k = %e P P^\\dagger v_k = %e\n",
352 printfQuda(
"L2 relative deviation = %e\n", deviation);
357 printfQuda(
"Checking 1 > || (1 - D P (P^\\dagger D P) P^\\dagger v_k || / || v_k || for %d vectors\n",
373 printfQuda(
"Checking 0 = (1 - P^\\dagger P) eta_c\n");
380 printfQuda(
"L2 relative deviation = %e\n", deviation);
384 printfQuda(
"Comparing native coarse operator to emulated operator\n");
409 #if 0 // enable to print out emulated and actual coarse-grid operator vectors for debugging 425 if(
fabs(delta_factor) >
tol ) {
429 deviation =
fabs(deviation);
432 printfQuda(
"L2 relative deviation = %e\n\n", deviation);
442 printfQuda(
"Smoother normal operator test (eta^dag M^dag M eta): real=%e imag=%e, relative imaginary deviation=%e\n",
443 real(
dot), imag(
dot), deviation);
451 printfQuda(
"Residual normal operator test (eta^dag M^dag M eta): real=%e imag=%e, relative imaginary deviation=%e\n",
452 real(
dot), imag(
dot), deviation);
461 printfQuda(
"Normal operator test (eta^dag M^dag M eta): real=%e imag=%e, relative imaginary deviation=%e\n",
462 real(
dot), imag(
dot), deviation);
472 double arpack_tol = 1
e-7;
473 char *which = (
char*)
malloc(256*
sizeof(
char));
487 std::vector<ColorSpinorField*> evecsBuffer;
488 evecsBuffer.reserve(nmodes);
495 void *evalsBuffer = arpPrecision ==
QUDA_DOUBLE_PRECISION ?
static_cast<void*
>(
new std::complex<double>[nmodes+1]) : static_cast<void*>(
new std::complex<float>[nmodes+1]);
499 for (
int i=0;
i<nmodes;
i++) {
501 *
tmp1 = *evecsBuffer[
i];
506 printfQuda(
"Vector %d: norms v_k = %e P^\\dagger v_k = %e P P^\\dagger v_k = %e\n",
510 printfQuda(
"L2 relative deviation = %e\n", deviation);
513 for (
unsigned int i = 0;
i < evecsBuffer.size();
i++)
delete evecsBuffer[
i];
516 else delete static_cast<std::complex<float>*
> (evalsBuffer);
539 if (
debug)
printfQuda(
"outer_solution_type = %d, inner_solution_type = %d\n", outer_solution_type, inner_solution_type);
542 errorQuda(
"Unsupported solution type combination");
545 errorQuda(
"For this coarse grid solution type, a preconditioned smoother is required");
563 dirac.prepare(
in,
out,
x, residual, outer_solution_type);
569 (*presmoother)(*
out, *
in);
571 ColorSpinorField &solution = inner_solution_type == outer_solution_type ?
x :
x.Even();
572 dirac.reconstruct(solution,
b, inner_solution_type);
576 bool use_solver_residual =
583 if (use_solver_residual) {
606 xpy(x_coarse_2_fine, solution);
609 printfQuda(
"after coarse-grid correction x2 = %e, r2 = %e\n",
624 (*postsmoother)(*
out, *
in);
626 dirac.reconstruct(
x,
b, outer_solution_type);
634 (*presmoother)(*
out, *
in);
635 dirac.reconstruct(
x,
b, outer_solution_type);
656 const int Nvec =
B.size();
659 void **
V =
new void*[Nvec];
660 for (
int i=0;
i<Nvec;
i++) {
670 B[0]->Ncolor(),
B[0]->Nspin(), Nvec, 0, (
char**)0);
672 errorQuda(
"\nQIO library was not built.\n");
675 printfQuda(
"Using %d constant nullvectors\n", Nvec);
678 for (
int i = 0;
i < (Nvec < 2 ? Nvec : 2);
i++) {
684 for (
int s=
i;
s<4;
s+=2) {
685 for (
int c=0;
c<
B[
i]->Ncolor();
c++) {
693 printfQuda(
"Using random source for nullvector = %d\n",
i);
717 const int Nvec =
B.size();
720 void **
V =
static_cast<void**
>(
safe_malloc(Nvec*
sizeof(
void*)));
721 for (
int i=0;
i<Nvec;
i++) {
729 B[0]->Ncolor(),
B[0]->Nspin(), Nvec, 0, (
char**)0);
739 errorQuda(
"\nQIO library was not built.\n");
771 csParam.setPrecision(
B[0]->Precision());
779 std::vector<ColorSpinorField*> B_gpu;
800 for(
unsigned int i=0;
i<
B.size();
i++) {
819 for (
int i=0; B_gpu[
i] !=
x; ++
i) {
825 if (nrm2 > 1
e-16)
ax(1.0 /
sqrt(nrm2), *
x);
826 else errorQuda(
"\nCannot orthogonalize %u vector\n",
i);
833 for (
int i=0;
i<(
int)
B.size();
i++) {
bool global_reduction
whether the solver acting as a preconditioner for another solver
DiracMatrix * matSmoothSloppy
QudaBoolean global_reduction
QudaMultigridParam & mg_global
QudaVerbosity verbosity_precondition
SolverParam * param_coarse_solver
std::vector< ColorSpinorField * > & B
enum QudaPrecision_s QudaPrecision
QudaFieldLocation location
void saveVectors(std::vector< ColorSpinorField *> &B)
Save the null space vectors in from file.
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
QudaVerbosity getVerbosity()
void createSmoother()
Create the smoothers.
ColorSpinorField * b_tilde
void R(ColorSpinorField &out, const ColorSpinorField &in) const
QudaSolveType smoother_solve_type[QUDA_MAX_MG_LEVEL]
void destroySmoother()
Free the smoothers.
double setup_tol[QUDA_MAX_MG_LEVEL]
__host__ __device__ ValueType sqrt(ValueType x)
Complex cDotProduct(ColorSpinorField &, ColorSpinorField &)
std::complex< double > Complex
void setOutputPrefix(const char *prefix)
Dirac * diracCoarseSmoother
DiracMatrix * matCoarseSmootherSloppy
SolverParam * param_presmooth
cudaColorSpinorField * tmp
const ColorSpinorField & Even() const
static ColorSpinorField * Create(const ColorSpinorParam ¶m)
cudaMipmappedArray_const_t unsigned int level
void arpackSolve(std::vector< ColorSpinorField *> &B, void *evals, DiracMatrix &matEigen, QudaPrecision matPrec, QudaPrecision arpackPrec, double tol, int nev, int ncv, char *target)
void P(ColorSpinorField &out, const ColorSpinorField &in) const
double pow(double, double)
double xmyNorm(ColorSpinorField &x, ColorSpinorField &y)
void loadVectors(std::vector< ColorSpinorField *> &B)
Load the null space vectors in from file.
QudaInverterType inv_type_precondition
QudaPreserveSource preserve_source
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. If site_subset is QUDA_FULL_SITE_SUBSET, the transfer operator can still be applied to single-parity fields, however, if site_subset is QUDA_PARITY_SITE_SUBSET, then the transfer operator cannot be applied to full fields, and setSiteSubset will need to be called first to reset to QUDA_FULL_SITE_SUBSET. This method exists to reduce GPU memory overhead - if only transfering single-parity fine fields then we only store a single-parity copy of the null space components on the device.
virtual void MdagM(ColorSpinorField &out, const ColorSpinorField &in) const =0
double mu_factor[QUDA_MAX_MG_LEVEL]
QudaSolveType smoother_solve_type
QudaSiteSubset siteSubset
static void axpby(Float a, Float *x, Float b, Float *y, int len)
double norm2(const CloverField &a, bool inverse=false)
void operator()(ColorSpinorField &out, ColorSpinorField &in)
int n_vec[QUDA_MAX_MG_LEVEL]
DiracMatrix * matCoarseSmoother
int strcmp(const char *__s1, const char *__s2)
QudaComputeNullVector compute_null_vector
ColorSpinorField * tmp_coarse
void ax(const double &a, GaugeField &u)
Scale the gauge field by the scalar a.
void * malloc(size_t __size) __attribute__((__warn_unused_result__)) __attribute__((alloc_size(1)))
QudaFieldLocation location
const ColorSpinorField & Vectors(QudaFieldLocation location=QUDA_CPU_FIELD_LOCATION) const
QudaResidualType residual_type
VOLATILE spinorFloat kappa
SolverParam * param_postsmooth
QudaFieldOrder fieldOrder
QudaInverterType smoother
static Solver * create(SolverParam ¶m, DiracMatrix &mat, DiracMatrix &matSloppy, DiracMatrix &matPrecon, TimeProfile &profile)
enum QudaMatPCType_s QudaMatPCType
int geoBlockSize[QUDA_MAX_DIM]
double smoother_tol[QUDA_MAX_MG_LEVEL]
Dirac * diracCoarseResidual
enum QudaSolutionType_s QudaSolutionType
bool is_preconditioner
verbosity to use for preconditioner
QudaVerbosity verbosity[QUDA_MAX_MG_LEVEL]
QudaMatPCType matpcType
NEW: used by mobius domain wall only.
std::vector< ColorSpinorField * > * B_coarse
enum QudaParity_s QudaParity
ColorSpinorField * r_coarse
void caxpy(const Complex &a, ColorSpinorField &x, ColorSpinorField &y)
DiracMatrix * matCoarseResidual
#define safe_malloc(size)
std::vector< ColorSpinorField * > * B
QudaPrecision precision_precondition
Dirac * diracCoarseSmootherSloppy
QudaFieldLocation location[QUDA_MAX_MG_LEVEL]
enum QudaSiteSubset_s QudaSiteSubset
virtual void solve(ColorSpinorField &out, ColorSpinorField &in)
char * strncpy(char *__dst, const char *__src, size_t __n)
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
int sprintf(char *, const char *,...) __attribute__((__format__(__printf__
int nu_pre[QUDA_MAX_MG_LEVEL]
#define QUDA_MAX_MG_LEVEL
Maximum number of multi-grid levels. This number may be increased if needed.
void xpy(ColorSpinorField &x, ColorSpinorField &y)
QudaTwistFlavorType TwistFlavor() const
void write_spinor_field(const char *filename, void *V[], QudaPrecision precision, const int *X, int nColor, int nSpin, int Nvec, int argc, char *argv[])
QudaUseInitGuess use_init_guess
QudaPrecision precision_sloppy
int nu_post[QUDA_MAX_MG_LEVEL]
MG(MGParam ¶m, TimeProfile &profile)
QudaComputeNullVector compute_null_vector
void generateNullVectors(std::vector< ColorSpinorField *> B)
Generate the null-space vectors.
QudaSolutionType coarse_grid_solution_type
QudaMultigridCycleType cycle_type
__device__ __host__ void zero(vector_type< scalar, n > &v)
QudaInverterType setup_inv_type[QUDA_MAX_MG_LEVEL]
void read_spinor_field(const char *filename, void *V[], QudaPrecision precision, const int *X, int nColor, int nSpin, int Nvec, int argc, char *argv[])
QudaInverterType smoother[QUDA_MAX_MG_LEVEL]
DiracMatrix * matResidual
ColorSpinorField * x_coarse
QudaInvertParam * invert_param
static void dot(sFloat *res, gFloat *a, sFloat *b)