QUDA  v1.1.0
A library for QCD on GPUs
inv_multi_cg_quda.cpp
Go to the documentation of this file.
1 #include <cstdio>
2 #include <cstdlib>
3 #include <cmath>
4 #include <limits>
5 
6 #include <quda_internal.h>
7 #include <color_spinor_field.h>
8 #include <blas_quda.h>
9 #include <dslash_quda.h>
10 #include <invert_quda.h>
11 #include <util_quda.h>
12 
23 #include <worker.h>
24 
25 namespace quda {
26 
49  class ShiftUpdate : public Worker {
50 
52  std::vector<ColorSpinorField*> p;
53  std::vector<ColorSpinorField*> x;
54 
55  double *alpha;
56  double *beta;
57  double *zeta;
58  double *zeta_old;
59 
60  const int j_low;
61  int n_shift;
62 
69  int n_update;
70 
71  public:
72  ShiftUpdate(ColorSpinorField *r, std::vector<ColorSpinorField*> p, std::vector<ColorSpinorField*> x,
73  double *alpha, double *beta, double *zeta, double *zeta_old, int j_low, int n_shift) :
74  r(r), p(p), x(x), alpha(alpha), beta(beta), zeta(zeta), zeta_old(zeta_old), j_low(j_low),
75  n_shift(n_shift), n_update( (r->Nspin()==4) ? 4 : 2 ) {
76 
77  }
78  virtual ~ShiftUpdate() { }
79 
80  void updateNshift(int new_n_shift) { n_shift = new_n_shift; }
81  void updateNupdate(int new_n_update) { n_update = new_n_update; }
82 
83  // note that we can't set the stream parameter here so it is
84  // ignored. This is more of a future design direction to consider
85  void apply(const qudaStream_t &stream)
86  {
87  static int count = 0;
88 
89 #if 0
90  // on the first call do the first half of the update
91  for (int j= (count*n_shift)/n_update+1; j<=((count+1)*n_shift)/n_update && j<n_shift; j++) {
92  beta[j] = beta[j_low] * zeta[j] * alpha[j] / ( zeta_old[j] * alpha[j_low] );
93  // update p[i] and x[i]
94  blas::axpyBzpcx(alpha[j], *(p[j]), *(x[j]), zeta[j], *r, beta[j]);
95  }
96 #else
97  int zero = (count*n_shift)/n_update+1;
98  std::vector<ColorSpinorField*> P, X;
99  for (int j= (count*n_shift)/n_update+1; j<=((count+1)*n_shift)/n_update && j<n_shift; j++) {
100  beta[j] = beta[j_low] * zeta[j] * alpha[j] / ( zeta_old[j] * alpha[j_low] );
101  P.push_back(p[j]);
102  X.push_back(x[j]);
103  }
104  if (P.size()) blas::axpyBzpcx(&alpha[zero], P, X, &zeta[zero], *r, &beta[zero]);
105 #endif
106  if (++count == n_update) count = 0;
107  }
108  };
109 
110  // this is the Worker pointer that the dslash uses to launch the shifted updates
111  namespace dslash {
113  }
114 
116  TimeProfile &profile) :
117  MultiShiftSolver(mat, matSloppy, param, profile)
118  {
119  }
120 
122 
126  void updateAlphaZeta(double *alpha, double *zeta, double *zeta_old,
127  const double *r2, const double *beta, const double pAp,
128  const double *offset, const int nShift, const int j_low) {
129  double alpha_old[QUDA_MAX_MULTI_SHIFT];
130  for (int j=0; j<nShift; j++) alpha_old[j] = alpha[j];
131 
132  alpha[0] = r2[0] / pAp;
133  zeta[0] = 1.0;
134  for (int j=1; j<nShift; j++) {
135  double c0 = zeta[j] * zeta_old[j] * alpha_old[j_low];
136  double c1 = alpha[j_low] * beta[j_low] * (zeta_old[j]-zeta[j]);
137  double c2 = zeta_old[j] * alpha_old[j_low] * (1.0+(offset[j]-offset[0])*alpha[j_low]);
138 
139  zeta_old[j] = zeta[j];
140  if (c1+c2 != 0.0){
141  zeta[j] = c0 / (c1 + c2);
142  }
143  else {
144  zeta[j] = 0.0;
145  }
146  if (zeta[j] != 0.0){
147  alpha[j] = alpha[j_low] * zeta[j] / zeta_old[j];
148  }
149  else {
150  alpha[j] = 0.0;
151  }
152  }
153  }
154 
155  void MultiShiftCG::operator()(std::vector<ColorSpinorField*>x, ColorSpinorField &b, std::vector<ColorSpinorField*> &p, double* r2_old_array )
156  {
157  if (checkLocation(*(x[0]), b) != QUDA_CUDA_FIELD_LOCATION)
158  errorQuda("Not supported");
159 
160  profile.TPSTART(QUDA_PROFILE_INIT);
161 
162  int num_offset = param.num_offset;
163  double *offset = param.offset;
164 
165  if (num_offset == 0) return;
166 
167  const double b2 = blas::norm2(b);
168  // Check to see that we're not trying to invert on a zero-field source
169  if(b2 == 0){
170  profile.TPSTOP(QUDA_PROFILE_INIT);
171  printfQuda("Warning: inverting on zero-field source\n");
172  for(int i=0; i<num_offset; ++i){
173  *(x[i]) = b;
174  param.true_res_offset[i] = 0.0;
175  param.true_res_hq_offset[i] = 0.0;
176  }
177  return;
178  }
179 
180  bool exit_early = false;
181  bool mixed = param.precision_sloppy != param.precision;
182  // whether we will switch to refinement on unshifted system after other shifts have converged
183  bool zero_refinement = param.precision_refinement_sloppy != param.precision;
184 
185  // this is the limit of precision possible
186  const double sloppy_tol= param.precision_sloppy == 8 ? std::numeric_limits<double>::epsilon() :
188  const double fine_tol = pow(10.,(-2*(int)b.Precision()+1));
189  std::unique_ptr<double[]> prec_tol(new double[num_offset]);
190 
191  prec_tol[0] = mixed ? sloppy_tol : fine_tol;
192  for (int i=1; i<num_offset; i++) {
193  prec_tol[i] = std::min(sloppy_tol,std::max(fine_tol,sqrt(param.tol_offset[i]*sloppy_tol)));
194  }
195 
196  double zeta[QUDA_MAX_MULTI_SHIFT];
197  double zeta_old[QUDA_MAX_MULTI_SHIFT];
198  double alpha[QUDA_MAX_MULTI_SHIFT];
199  double beta[QUDA_MAX_MULTI_SHIFT];
200 
201  int j_low = 0;
202  int num_offset_now = num_offset;
203  for (int i=0; i<num_offset; i++) {
204  zeta[i] = zeta_old[i] = 1.0;
205  beta[i] = 0.0;
206  alpha[i] = 1.0;
207  }
208 
209  // flag whether we will be using reliable updates or not
210  bool reliable = false;
211  for (int j=0; j<num_offset; j++)
212  if (param.tol_offset[j] < param.delta) reliable = true;
213 
214 
215  auto *r = new cudaColorSpinorField(b);
216  std::vector<ColorSpinorField*> x_sloppy;
217  x_sloppy.resize(num_offset);
218  std::vector<ColorSpinorField*> y;
219 
222 
223  if (reliable) {
224  y.resize(num_offset);
225  for (int i=0; i<num_offset; i++) y[i] = new cudaColorSpinorField(*r, csParam);
226  }
227 
228  csParam.setPrecision(param.precision_sloppy);
229 
230  cudaColorSpinorField *r_sloppy;
231  if (param.precision_sloppy == x[0]->Precision()) {
232  r_sloppy = r;
233  } else {
235  r_sloppy = new cudaColorSpinorField(*r, csParam);
236  }
237 
238  if (param.precision_sloppy == x[0]->Precision() ||
240  for (int i=0; i<num_offset; i++){
241  x_sloppy[i] = x[i];
242  blas::zero(*x_sloppy[i]);
243  }
244  } else {
246  for (int i=0; i<num_offset; i++)
247  x_sloppy[i] = new cudaColorSpinorField(*x[i], csParam);
248  }
249 
250  p.resize(num_offset);
251  for (int i=0; i<num_offset; i++) p[i] = new cudaColorSpinorField(*r_sloppy);
252 
254  auto* Ap = new cudaColorSpinorField(*r_sloppy, csParam);
255 
256  cudaColorSpinorField tmp1(*Ap, csParam);
257 
258  // tmp2 only needed for multi-gpu Wilson-like kernels
259  cudaColorSpinorField *tmp2_p = !mat.isStaggered() ?
260  new cudaColorSpinorField(*Ap, csParam) : &tmp1;
261  cudaColorSpinorField &tmp2 = *tmp2_p;
262 
263  // additional high-precision temporary if Wilson and mixed-precision
264  csParam.setPrecision(param.precision);
265  cudaColorSpinorField *tmp3_p =
267  new cudaColorSpinorField(*r, csParam) : &tmp1;
268  cudaColorSpinorField &tmp3 = *tmp3_p;
269 
270  profile.TPSTOP(QUDA_PROFILE_INIT);
272 
273  // stopping condition of each shift
274  double stop[QUDA_MAX_MULTI_SHIFT];
275  double r2[QUDA_MAX_MULTI_SHIFT];
276  int iter[QUDA_MAX_MULTI_SHIFT+1]; // record how many iterations for each shift
277  for (int i=0; i<num_offset; i++) {
278  r2[i] = b2;
280  iter[i] = 0;
281  }
282  // this initial condition ensures that the heaviest shift can be removed
283  iter[num_offset] = 1;
284 
285  double r2_old;
286  double pAp;
287 
288  double rNorm[QUDA_MAX_MULTI_SHIFT];
289  double r0Norm[QUDA_MAX_MULTI_SHIFT];
290  double maxrx[QUDA_MAX_MULTI_SHIFT];
291  double maxrr[QUDA_MAX_MULTI_SHIFT];
292  for (int i=0; i<num_offset; i++) {
293  rNorm[i] = sqrt(r2[i]);
294  r0Norm[i] = rNorm[i];
295  maxrx[i] = rNorm[i];
296  maxrr[i] = rNorm[i];
297  }
298  double delta = param.delta;
299 
300  // this parameter determines how many consective reliable update
301  // reisudal increases we tolerate before terminating the solver,
302  // i.e., how long do we want to keep trying to converge
303  const int maxResIncrease = param.max_res_increase; // check if we reached the limit of our tolerance
304  const int maxResIncreaseTotal = param.max_res_increase_total;
305 
306  int resIncrease = 0;
307  int resIncreaseTotal[QUDA_MAX_MULTI_SHIFT];
308  for (int i=0; i<num_offset; i++) {
309  resIncreaseTotal[i]=0;
310  }
311 
312  int k = 0;
313  int rUpdate = 0;
314  blas::flops = 0;
315 
316  bool aux_update = false;
317 
318  // now create the worker class for updating the shifted solutions and gradient vectors
319  ShiftUpdate shift_update(r_sloppy, p, x_sloppy, alpha, beta, zeta, zeta_old, j_low, num_offset_now);
320 
322  profile.TPSTART(QUDA_PROFILE_COMPUTE);
323 
324  if (getVerbosity() >= QUDA_VERBOSE)
325  printfQuda("MultiShift CG: %d iterations, <r,r> = %e, |r|/|b| = %e\n", k, r2[0], sqrt(r2[0]/b2));
326 
327  while ( !convergence(r2, stop, num_offset_now) && !exit_early && k < param.maxiter) {
328 
329  if (aux_update) dslash::aux_worker = &shift_update;
330  matSloppy(*Ap, *p[0], tmp1, tmp2);
331  dslash::aux_worker = nullptr;
332  aux_update = false;
333 
334  // update number of shifts now instead of end of previous
335  // iteration so that all shifts are updated during the dslash
336  shift_update.updateNshift(num_offset_now);
337 
338  // at some point we should curry these into the Dirac operator
339  if (r->Nspin()==4) pAp = blas::axpyReDot(offset[0], *p[0], *Ap);
340  else pAp = blas::reDotProduct(*p[0], *Ap);
341 
342  // compute zeta and alpha
343  for (int j=1; j<num_offset_now; j++) r2_old_array[j] = zeta[j] * zeta[j] * r2[0];
344  updateAlphaZeta(alpha, zeta, zeta_old, r2, beta, pAp, offset, num_offset_now, j_low);
345 
346  r2_old = r2[0];
347  r2_old_array[0] = r2_old;
348 
349  Complex cg_norm = blas::axpyCGNorm(-alpha[j_low], *Ap, *r_sloppy);
350  r2[0] = real(cg_norm);
351  double zn = imag(cg_norm);
352 
353  // reliable update conditions
354  rNorm[0] = sqrt(r2[0]);
355  for (int j=1; j<num_offset_now; j++) rNorm[j] = rNorm[0] * zeta[j];
356 
357  int updateX=0, updateR=0;
358  //fixme: with the current implementation of the reliable update it is sufficient to trigger it only for shift 0
359  //fixme: The loop below is unnecessary but I don't want to delete it as we still might find a better reliable update
360  int reliable_shift = -1; // this is the shift that sets the reliable_shift
361  for (int j=0; j>=0; j--) {
362  if (rNorm[j] > maxrx[j]) maxrx[j] = rNorm[j];
363  if (rNorm[j] > maxrr[j]) maxrr[j] = rNorm[j];
364  updateX = (rNorm[j] < delta*r0Norm[j] && r0Norm[j] <= maxrx[j]) ? 1 : updateX;
365  updateR = ((rNorm[j] < delta*maxrr[j] && r0Norm[j] <= maxrr[j]) || updateX) ? 1 : updateR;
366  if ((updateX || updateR) && reliable_shift == -1) reliable_shift = j;
367  }
368 
369  if ( !(updateR || updateX) || !reliable) {
370  //beta[0] = r2[0] / r2_old;
371  beta[0] = zn / r2_old;
372  // update p[0] and x[0]
373  blas::axpyZpbx(alpha[0], *p[0], *x_sloppy[0], *r_sloppy, beta[0]);
374 
375  // this should trigger the shift update in the subsequent sloppy dslash
376  aux_update = true;
377  /*
378  for (int j=1; j<num_offset_now; j++) {
379  beta[j] = beta[j_low] * zeta[j] * alpha[j] / (zeta_old[j] * alpha[j_low]);
380  // update p[i] and x[i]
381  blas::axpyBzpcx(alpha[j], *p[j], *x_sloppy[j], zeta[j], *r_sloppy, beta[j]);
382  }
383  */
384  } else {
385  for (int j=0; j<num_offset_now; j++) {
386  blas::axpy(alpha[j], *p[j], *x_sloppy[j]);
387  blas::xpy(*x_sloppy[j], *y[j]);
388  }
389 
390  mat(*r, *y[0], *x[0], tmp3); // here we can use x as tmp
391  if (r->Nspin()==4) blas::axpy(offset[0], *y[0], *r);
392 
393  r2[0] = blas::xmyNorm(b, *r);
394  for (int j=1; j<num_offset_now; j++) r2[j] = zeta[j] * zeta[j] * r2[0];
395  for (int j=0; j<num_offset_now; j++) blas::zero(*x_sloppy[j]);
396 
397  blas::copy(*r_sloppy, *r);
398 
399  // break-out check if we have reached the limit of the precision
400  if (sqrt(r2[reliable_shift]) > r0Norm[reliable_shift]) { // reuse r0Norm for this
401  resIncrease++;
402  resIncreaseTotal[reliable_shift]++;
403  warningQuda("MultiShiftCG: Shift %d, updated residual %e is greater than previous residual %e (total #inc %i)",
404  reliable_shift, sqrt(r2[reliable_shift]), r0Norm[reliable_shift], resIncreaseTotal[reliable_shift]);
405 
406  if (resIncrease > maxResIncrease or resIncreaseTotal[reliable_shift] > maxResIncreaseTotal) {
407  warningQuda("MultiShiftCG: solver exiting due to too many true residual norm increases");
408  break;
409  }
410  } else {
411  resIncrease = 0;
412  }
413 
414  // explicitly restore the orthogonality of the gradient vector
415  for (int j=0; j<num_offset_now; j++) {
416  Complex rp = blas::cDotProduct(*r_sloppy, *p[j]) / (r2[0]);
417  blas::caxpy(-rp, *r_sloppy, *p[j]);
418  }
419 
420  // update beta and p
421  beta[0] = r2[0] / r2_old;
422  blas::xpay(*r_sloppy, beta[0], *p[0]);
423  for (int j=1; j<num_offset_now; j++) {
424  beta[j] = beta[j_low] * zeta[j] * alpha[j] / (zeta_old[j] * alpha[j_low]);
425  blas::axpby(zeta[j], *r_sloppy, beta[j], *p[j]);
426  }
427 
428  // update reliable update parameters for the system that triggered the update
429  int m = reliable_shift;
430  rNorm[m] = sqrt(r2[0]) * zeta[m];
431  maxrr[m] = rNorm[m];
432  maxrx[m] = rNorm[m];
433  r0Norm[m] = rNorm[m];
434  rUpdate++;
435  }
436 
437  // now we can check if any of the shifts have converged and remove them
438  int converged = 0;
439  for (int j=num_offset_now-1; j>=1; j--) {
440  if (zeta[j] == 0.0 && r2[j+1] < stop[j+1]) {
441  converged++;
442  if (getVerbosity() >= QUDA_VERBOSE)
443  printfQuda("MultiShift CG: Shift %d converged after %d iterations\n", j, k+1);
444  } else {
445  r2[j] = zeta[j] * zeta[j] * r2[0];
446  // only remove if shift above has converged
447  if ((r2[j] < stop[j] || sqrt(r2[j] / b2) < prec_tol[j]) && iter[j+1] ) {
448  converged++;
449  iter[j] = k+1;
450  if (getVerbosity() >= QUDA_VERBOSE)
451  printfQuda("MultiShift CG: Shift %d converged after %d iterations\n", j, k+1);
452  }
453  }
454  }
455  num_offset_now -= converged;
456 
457  // exit early so that we can finish of shift 0 using CG and allowing for mixed precison refinement
458  if ( (mixed || zero_refinement) and param.compute_true_res and num_offset_now==1) {
459  exit_early=true;
460  num_offset_now--;
461  }
462 
463  // this ensure we do the update on any shifted systems that
464  // happen to converge when the un-shifted system converges
465  if ( (convergence(r2, stop, num_offset_now) || exit_early || k == param.maxiter) && aux_update == true) {
466  if (getVerbosity() >= QUDA_VERBOSE)
467  printfQuda("Convergence of unshifted system so trigger shiftUpdate\n");
468 
469  // set worker to do all updates at once
470  shift_update.updateNupdate(1);
471  shift_update.apply(0);
472 
473  for (int j=0; j<num_offset_now; j++) iter[j] = k+1;
474  }
475 
476  k++;
477 
478  if (getVerbosity() >= QUDA_VERBOSE)
479  printfQuda("MultiShift CG: %d iterations, <r,r> = %e, |r|/|b| = %e\n", k, r2[0], sqrt(r2[0]/b2));
480  }
481 
482  for (int i=0; i<num_offset; i++) {
483  if (iter[i] == 0) iter[i] = k;
484  blas::copy(*x[i], *x_sloppy[i]);
485  if (reliable) blas::xpy(*y[i], *x[i]);
486  }
487 
490 
491  if (getVerbosity() >= QUDA_VERBOSE)
492  printfQuda("MultiShift CG: Reliable updates = %d\n", rUpdate);
493 
494  if (k==param.maxiter) warningQuda("Exceeded maximum iterations %d\n", param.maxiter);
495 
497  double gflops = (blas::flops + mat.flops() + matSloppy.flops())*1e-9;
498  param.gflops = gflops;
499  param.iter += k;
500 
501  if (param.compute_true_res){
502  // only allocate temporaries if necessary
503  csParam.setPrecision(param.precision);
504  ColorSpinorField *tmp4_p = reliable ? y[0] : tmp1.Precision() == x[0]->Precision() ? &tmp1 : ColorSpinorField::Create(csParam);
505  ColorSpinorField *tmp5_p = mat.isStaggered() ? tmp4_p :
506  reliable ? y[1] : (tmp2.Precision() == x[0]->Precision() && &tmp1 != tmp2_p) ? tmp2_p : ColorSpinorField::Create(csParam);
507 
508  for (int i = 0; i < num_offset; i++) {
509  // only calculate true residual if we need to:
510  // 1.) For higher shifts if we did not use mixed precision
511  // 2.) For shift 0 if we did not exit early (we went to the full solution)
512  if ( (i > 0 and not mixed) or (i == 0 and not exit_early) ) {
513  mat(*r, *x[i], *tmp4_p, *tmp5_p);
514  if (r->Nspin() == 4) {
515  blas::axpy(offset[i], *x[i], *r); // Offset it.
516  } else if (i != 0) {
517  blas::axpy(offset[i] - offset[0], *x[i], *r); // Offset it.
518  }
519  double true_res = blas::xmyNorm(b, *r);
520  param.true_res_offset[i] = sqrt(true_res / b2);
521  param.iter_res_offset[i] = sqrt(r2[i] / b2);
523  } else {
524  param.iter_res_offset[i] = sqrt(r2[i] / b2);
525  param.true_res_offset[i] = std::numeric_limits<double>::infinity();
526  param.true_res_hq_offset[i] = std::numeric_limits<double>::infinity();
527  }
528  }
529 
530  if (getVerbosity() >= QUDA_SUMMARIZE) {
531  printfQuda("MultiShift CG: Converged after %d iterations\n", k);
532  for (int i = 0; i < num_offset; i++) {
533  if(std::isinf(param.true_res_offset[i])){
534  printfQuda(" shift=%d, %d iterations, relative residual: iterated = %e\n",
535  i, iter[i], param.iter_res_offset[i]);
536  } else {
537  printfQuda(" shift=%d, %d iterations, relative residual: iterated = %e, true = %e\n",
538  i, iter[i], param.iter_res_offset[i], param.true_res_offset[i]);
539  }
540  }
541  }
542 
543  if (tmp5_p != tmp4_p && tmp5_p != tmp2_p && (reliable ? tmp5_p != y[1] : 1)) delete tmp5_p;
544  if (tmp4_p != &tmp1 && (reliable ? tmp4_p != y[0] : 1)) delete tmp4_p;
545  } else {
546  if (getVerbosity() >= QUDA_SUMMARIZE)
547  {
548  printfQuda("MultiShift CG: Converged after %d iterations\n", k);
549  for (int i = 0; i < num_offset; i++) {
550  param.iter_res_offset[i] = sqrt(r2[i] / b2);
551  printfQuda(" shift=%d, %d iterations, relative residual: iterated = %e\n",
552  i, iter[i], param.iter_res_offset[i]);
553  }
554  }
555  }
556 
557  // reset the flops counters
558  blas::flops = 0;
559  mat.flops();
560  matSloppy.flops();
561 
563  profile.TPSTART(QUDA_PROFILE_FREE);
564 
565  if (&tmp3 != &tmp1) delete tmp3_p;
566  if (&tmp2 != &tmp1) delete tmp2_p;
567 
568  if (r_sloppy->Precision() != r->Precision()) delete r_sloppy;
569  for (int i=0; i<num_offset; i++)
570  if (x_sloppy[i]->Precision() != x[i]->Precision()) delete x_sloppy[i];
571 
572  delete r;
573 
574  if (reliable) for (int i=0; i<num_offset; i++) delete y[i];
575 
576  delete Ap;
577 
578  profile.TPSTOP(QUDA_PROFILE_FREE);
579 
580  return;
581  }
582 
583 } // namespace quda
int Nspin
Definition: blas_test.cpp:50
static ColorSpinorField * Create(const ColorSpinorParam &param)
bool isStaggered() const
return if the operator is a staggered operator
Definition: dirac_quda.h:1935
unsigned long long flops() const
Definition: dirac_quda.h:1909
QudaPrecision Precision() const
MultiShiftCG(const DiracMatrix &mat, const DiracMatrix &matSloppy, SolverParam &param, TimeProfile &profile)
void operator()(std::vector< ColorSpinorField * >x, ColorSpinorField &b, std::vector< ColorSpinorField * > &p, double *r2_old_array)
Run multi-shift and return Krylov-space at the end of the solve in p and r2_old_arry.
const DiracMatrix & matSloppy
Definition: invert_quda.h:1236
TimeProfile & profile
Definition: invert_quda.h:1238
const DiracMatrix & mat
Definition: invert_quda.h:1235
bool convergence(const double *r2, const double *r2_tol, int n) const
Definition: solver.cpp:427
void updateNshift(int new_n_shift)
void updateNupdate(int new_n_update)
void apply(const qudaStream_t &stream)
ShiftUpdate(ColorSpinorField *r, std::vector< ColorSpinorField * > p, std::vector< ColorSpinorField * > x, double *alpha, double *beta, double *zeta, double *zeta_old, int j_low, int n_shift)
static double stopping(double tol, double b2, QudaResidualType residual_type)
Set the solver L2 stopping condition.
Definition: solver.cpp:311
double Last(QudaProfileType idx)
Definition: timer.h:254
double epsilon
void mat(void *out, void **link, void *in, int dagger_bit, int mu, QudaPrecision sPrecision, QudaPrecision gPrecision)
@ QUDA_CUDA_FIELD_LOCATION
Definition: enum_quda.h:326
@ QUDA_SUMMARIZE
Definition: enum_quda.h:266
@ QUDA_VERBOSE
Definition: enum_quda.h:267
@ QUDA_ZERO_FIELD_CREATE
Definition: enum_quda.h:361
@ QUDA_COPY_FIELD_CREATE
Definition: enum_quda.h:362
#define checkLocation(...)
double axpyReDot(double a, ColorSpinorField &x, ColorSpinorField &y)
Complex axpyCGNorm(double a, ColorSpinorField &x, ColorSpinorField &y)
void axpyZpbx(double a, ColorSpinorField &x, ColorSpinorField &y, ColorSpinorField &z, double b)
double3 HeavyQuarkResidualNorm(ColorSpinorField &x, ColorSpinorField &r)
double xmyNorm(ColorSpinorField &x, ColorSpinorField &y)
Definition: blas_quda.h:79
unsigned long long flops
void axpyBzpcx(double a, ColorSpinorField &x, ColorSpinorField &y, double b, ColorSpinorField &z, double c)
void xpay(ColorSpinorField &x, double a, ColorSpinorField &y)
Definition: blas_quda.h:45
void zero(ColorSpinorField &a)
double norm2(const ColorSpinorField &a)
double reDotProduct(ColorSpinorField &x, ColorSpinorField &y)
void axpy(double a, ColorSpinorField &x, ColorSpinorField &y)
Definition: blas_quda.h:43
void caxpy(const Complex &a, ColorSpinorField &x, ColorSpinorField &y)
void xpy(ColorSpinorField &x, ColorSpinorField &y)
Definition: blas_quda.h:41
void copy(ColorSpinorField &dst, const ColorSpinorField &src)
Definition: blas_quda.h:24
Complex cDotProduct(ColorSpinorField &, ColorSpinorField &)
void axpby(double a, ColorSpinorField &x, double b, ColorSpinorField &y)
Definition: blas_quda.h:44
void stop()
Stop profiling.
Definition: device.cpp:228
int reliable(double &rNorm, double &maxrx, double &maxrr, const double &r2, const double &delta)
__device__ __host__ void zero(double &a)
Definition: float_vector.h:318
void updateAlphaZeta(double *alpha, double *zeta, double *zeta_old, const double *r2, const double *beta, const double pAp, const double *offset, const int nShift, const int j_low)
std::complex< double > Complex
Definition: quda_internal.h:86
__host__ __device__ ValueType sqrt(ValueType x)
Definition: complex_quda.h:120
qudaStream_t * stream
@ QUDA_PROFILE_INIT
Definition: timer.h:106
@ QUDA_PROFILE_EPILOGUE
Definition: timer.h:110
@ QUDA_PROFILE_COMPUTE
Definition: timer.h:108
@ QUDA_PROFILE_FREE
Definition: timer.h:111
@ QUDA_PROFILE_PREAMBLE
Definition: timer.h:107
__host__ __device__ ValueType pow(ValueType x, ExponentType e)
Definition: complex_quda.h:111
ColorSpinorParam csParam
Definition: pack_test.cpp:25
QudaGaugeParam param
Definition: pack_test.cpp:18
void updateR()
update the radius for halos.
cudaStream_t qudaStream_t
Definition: quda_api.h:9
#define QUDA_MAX_MULTI_SHIFT
Maximum number of shifts supported by the multi-shift solver. This number may be changed if need be.
QudaPrecision precision
Definition: invert_quda.h:136
double iter_res_offset[QUDA_MAX_MULTI_SHIFT]
Definition: invert_quda.h:181
bool use_sloppy_partial_accumulator
Definition: invert_quda.h:70
int max_res_increase_total
Definition: invert_quda.h:90
QudaResidualType residual_type
Definition: invert_quda.h:49
QudaPrecision precision_refinement_sloppy
Definition: invert_quda.h:142
double offset[QUDA_MAX_MULTI_SHIFT]
Definition: invert_quda.h:169
QudaPrecision precision_sloppy
Definition: invert_quda.h:139
double tol_offset[QUDA_MAX_MULTI_SHIFT]
Definition: invert_quda.h:172
double true_res_offset[QUDA_MAX_MULTI_SHIFT]
Definition: invert_quda.h:178
double true_res_hq_offset[QUDA_MAX_MULTI_SHIFT]
Definition: invert_quda.h:184
#define printfQuda(...)
Definition: util_quda.h:114
QudaVerbosity getVerbosity()
Definition: util_quda.cpp:21
#define warningQuda(...)
Definition: util_quda.h:132
#define errorQuda(...)
Definition: util_quda.h:120