28 Solver(param, profile), mat(mat), matSloppy(matSloppy), matPrecon(matPrecon),
init(false) {
47 int reliable(
double &rNorm,
double &maxrx,
double &maxrr,
const double &r2,
const double &delta) {
50 if (rNorm > maxrx) maxrx = rNorm;
51 if (rNorm > maxrr) maxrr = rNorm;
54 int updateR = (rNorm < delta*maxrr) ? 1 : 0;
146 const bool use_heavy_quark_res =
149 int heavy_quark_check = 10;
165 double rNorm =
sqrt(r2);
167 double maxrr = rNorm;
168 double maxrx = rNorm;
170 PrintStats(
"BiCGstab", k, r2, b2, heavy_quark_res);
183 printfQuda(
"BiCGstab debug: x2=%e, r2=%e, v2=%e, p2=%e, tmp2=%e r0=%e t2=%e\n",
189 matSloppy(v, p, tmp);
198 if (
abs(rho) == 0.0) alpha = 0.0;
199 else alpha = rho / r0v;
204 matSloppy(t, rSloppy, tmp);
211 double t2 = omega_t2.z;
213 double s2 =
norm2(rSloppy);
216 r2 = s2 - real(omega *
conj(tr)) ;
219 updateR =
reliable(rNorm, maxrx, maxrr, r2, delta);
223 omega =
Complex(omega_t2.x / omega_t2.z, omega_t2.y / omega_t2.z);
236 rho =
Complex(rho_r2.x, rho_r2.y);
240 if (use_heavy_quark_res && k%heavy_quark_check==0) {
241 if (&x != &xSloppy) {
272 PrintStats(
"BiCGstab", k, r2, b2, heavy_quark_res);
274 printfQuda(
"BiCGstab debug: x2=%e, r2=%e, v2=%e, p2=%e, tmp2=%e r0=%e t2=%e\n",
279 if (
abs(rho*alpha) == 0.0) beta = 0.0;
280 else beta = (rho/rho0) * (alpha/omega);
307 #if (__COMPUTE_CAPABILITY__ >= 200)
330 if (&x != &xSloppy)
delete x_sloppy;
bool convergence(const double &r2, const double &hq2, const double &r2_tol, const double &hq_tol)
void setPrecision(QudaPrecision precision)
void caxpyCuda(const Complex &a, cudaColorSpinorField &x, cudaColorSpinorField &y)
static double stopping(const double &tol, const double &b2, QudaResidualType residual_type)
QudaVerbosity getVerbosity()
__host__ __device__ ValueType sqrt(ValueType x)
std::complex< double > Complex
void mat(void *out, void **fatlink, void **longlink, void *in, double kappa, int dagger_bit, QudaPrecision sPrecision, QudaPrecision gPrecision)
void operator()(cudaColorSpinorField &out, cudaColorSpinorField &in)
double resNorm(const DiracMatrix &mat, cudaColorSpinorField &b, cudaColorSpinorField &x)
BiCGstab(DiracMatrix &mat, DiracMatrix &matSloppy, DiracMatrix &matPrecon, SolverParam ¶m, TimeProfile &profile)
QudaInverterType inv_type_precondition
void fillInnerSolveParam(SolverParam &inner, const SolverParam &outer)
unsigned long long flops() const
cudaColorSpinorField * tmp
void PrintSummary(const char *name, int k, const double &r2, const double &b2)
double3 caxpbypzYmbwcDotProductUYNormYCuda(const Complex &a, cudaColorSpinorField &x, const Complex &b, cudaColorSpinorField &y, cudaColorSpinorField &z, cudaColorSpinorField &w, cudaColorSpinorField &u)
QudaResidualType residual_type
Complex cDotProductCuda(cudaColorSpinorField &, cudaColorSpinorField &)
void copyCuda(cudaColorSpinorField &dst, const cudaColorSpinorField &src)
double normCuda(const cudaColorSpinorField &b)
unsigned long long blas_flops
double3 xpyHeavyQuarkResidualNormCuda(cudaColorSpinorField &x, cudaColorSpinorField &y, cudaColorSpinorField &r)
void xpyCuda(cudaColorSpinorField &x, cudaColorSpinorField &y)
void Stop(QudaProfileType idx)
QudaPrecision Precision() const
void PrintStats(const char *, int k, const double &r2, const double &b2, const double &hq2)
double Last(QudaProfileType idx)
void reduceDouble(double &)
void cxpaypbzCuda(cudaColorSpinorField &, const Complex &b, cudaColorSpinorField &y, const Complex &c, cudaColorSpinorField &z)
int reliable(double &rNorm, double &maxrx, double &maxrr, const double &r2, const double &delta)
void zeroCuda(cudaColorSpinorField &a)
double3 cDotProductNormACuda(cudaColorSpinorField &a, cudaColorSpinorField &b)
void Start(QudaProfileType idx)
QudaUseInitGuess use_init_guess
void init(int argc, char **argv)
__host__ __device__ ValueType abs(ValueType x)
QudaPrecision precision_sloppy
bool use_sloppy_partial_accumulator
__host__ __device__ ValueType conj(ValueType x)
double3 HeavyQuarkResidualNormCuda(cudaColorSpinorField &x, cudaColorSpinorField &r)
double norm2(const ColorSpinorField &)
double xmyNormCuda(cudaColorSpinorField &a, cudaColorSpinorField &b)
void caxpbypzYmbwCuda(const Complex &, cudaColorSpinorField &, const Complex &, cudaColorSpinorField &, cudaColorSpinorField &, cudaColorSpinorField &)