QUDA  v0.7.0
A library for QCD on GPUs
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
inv_cg3_quda.cpp
Go to the documentation of this file.
1 #include <stdio.h>
2 #include <stdlib.h>
3 #include <math.h>
4 
5 #include <quda_internal.h>
6 #include <color_spinor_field.h>
7 #include <blas_quda.h>
8 #include <dslash_quda.h>
9 #include <invert_quda.h>
10 #include <util_quda.h>
11 #include <sys/time.h>
12 
13 #include <face_quda.h>
14 
15 #include <iostream>
16 
17 namespace quda {
18 
19  CG3::CG3(DiracMatrix &mat, SolverParam &param, TimeProfile &profile) :
20  Solver(param, profile), mat(mat)
21  {
22  }
23 
24  CG3::~CG3() {
25  }
26 
27  void CG3::operator()(cudaColorSpinorField &x, cudaColorSpinorField &b)
28  {
29 
30  // Check to see that we're not trying to invert on a zero-field source
31  const double b2 = norm2(b);
32  if(b2 == 0){
33  profile.Stop(QUDA_PROFILE_INIT);
34  printfQuda("Warning: inverting on zero-field source\n");
35  x=b;
36  param.true_res = 0.0;
37  param.true_res_hq = 0.0;
38  return;
39  }
40 
41  ColorSpinorParam csParam(x);
43 
44 
45  cudaColorSpinorField x_prev(b, csParam);
46  cudaColorSpinorField r_prev(b, csParam);
47  cudaColorSpinorField temp(b, csParam);
48 
49  cudaColorSpinorField r(b);
50  cudaColorSpinorField w(b);
51 
52 
53  mat(r, x, temp); // r = Mx
54  double r2 = xmyNormCuda(b,r); // r = b - Mx
55  PrintStats("CG3", 0, r2, b2, 0.0);
56 
57 
58  double stop = stopping(param.tol, b2, param.residual_type);
59  if(convergence(r2, 0.0, stop, 0.0)) return;
60  // First iteration
61  mat(w, r, temp);
62  double rAr = reDotProductCuda(r,w);
63  double rho = 1.0;
64  double gamma_prev = 0.0;
65  double gamma = r2/rAr;
66 
67 
68  cudaColorSpinorField x_new(x);
69  cudaColorSpinorField r_new(r);
70  axpyCuda(gamma, r, x_new); // x_new += gamma*r
71  axpyCuda(-gamma, w, r_new); // r_new -= gamma*w
72  // end of first iteration
73 
74  // axpbyCuda(a,b,x,y) => y = a*x + b*y
75 
76  int k = 1; // First iteration performed above
77 
78  double r2_prev;
79  while(!convergence(r2, 0.0, stop, 0.0) && k<param.maxiter){
80  x_prev = x; x = x_new;
81  r_prev = r; r = r_new;
82  mat(w, r, temp);
83  rAr = reDotProductCuda(r,w);
84  r2_prev = r2;
85  r2 = norm2(r);
86 
87  // Need to rearrange this!
88  PrintStats("CG3", k, r2, b2, 0.0);
89 
90  gamma_prev = gamma;
91  gamma = r2/rAr;
92  rho = 1.0/(1. - (gamma/gamma_prev)*(r2/r2_prev)*(1.0/rho));
93 
94  x_new = x;
95  axCuda(rho,x_new);
96  axpyCuda(rho*gamma,r,x_new);
97  axpyCuda((1. - rho),x_prev,x_new);
98 
99  r_new = r;
100  axCuda(rho,r_new);
101  axpyCuda(-rho*gamma,w,r_new);
102  axpyCuda((1.-rho),r_prev,r_new);
103 
104 
105  double rr_old = reDotProductCuda(r_new,r);
106  printfQuda("rr_old = %1.14lf\n", rr_old);
107 
108 
109 
110  k++;
111  }
112 
113 
114  if(k == param.maxiter)
115  warningQuda("Exceeded maximum iterations %d", param.maxiter);
116 
117  // compute the true residual
118  mat(r, x, temp);
119  param.true_res = sqrt(xmyNormCuda(b, r)/b2);
120 
121  PrintSummary("CG3", k, r2, b2);
122 
123  return;
124  }
125 
126 } // namespace quda
__host__ __device__ ValueType sqrt(ValueType x)
Definition: complex_quda.h:105
void mat(void *out, void **fatlink, void **longlink, void *in, double kappa, int dagger_bit, QudaPrecision sPrecision, QudaPrecision gPrecision)
QudaGaugeParam param
Definition: pack_test.cpp:17
ColorSpinorParam csParam
Definition: pack_test.cpp:24
#define warningQuda(...)
Definition: util_quda.h:84
void axpyCuda(const double &a, cudaColorSpinorField &x, cudaColorSpinorField &y)
Definition: blas_quda.cu:115
int x[4]
double reDotProductCuda(cudaColorSpinorField &a, cudaColorSpinorField &b)
Definition: reduce_quda.cu:170
#define printfQuda(...)
Definition: util_quda.h:67
void axCuda(const double &a, cudaColorSpinorField &x)
Definition: blas_quda.cu:171
double norm2(const ColorSpinorField &)
double xmyNormCuda(cudaColorSpinorField &a, cudaColorSpinorField &b)
Definition: reduce_quda.cu:343