QUDA  1.0.0
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
clover_field.h
Go to the documentation of this file.
1 #ifndef _CLOVER_QUDA_H
2 #define _CLOVER_QUDA_H
3 
4 #include <quda_internal.h>
5 #include <lattice_field.h>
6 
7 namespace quda {
8 
10  bool direct; // whether to create the direct clover
11  bool inverse; // whether to create the inverse clover
12  void *clover;
13  void *norm;
14  void *cloverInv;
15  void *invNorm;
16  double csw;
17  bool twisted; // whether to create twisted mass clover
18  double mu2;
19  double rho;
20 
24  this->precision = precision;
25  this->ghost_precision = precision;
26  order = (precision == QUDA_DOUBLE_PRECISION) ?
28  }
29 
31  direct(true), inverse(true), clover(nullptr), norm(nullptr),
32  cloverInv(nullptr), invNorm(nullptr), twisted(false), mu2(0.0), rho(0.0) { }
33 
35  direct(param.direct), inverse(param.inverse),
36  clover(param.clover), norm(param.norm),
37  cloverInv(param.cloverInv), invNorm(param.invNorm),
38  twisted(param.twisted), mu2(param.mu2), rho(param.rho) { }
39 
40  CloverFieldParam(const CloverField &field);
41  };
42 
43  std::ostream& operator<<(std::ostream& output, const CloverFieldParam& param);
44 
45  class CloverField : public LatticeField {
46 
47  protected:
48  size_t bytes; // bytes allocated per clover full field
49  size_t norm_bytes; // sizeof each norm full field
50  size_t length;
51  size_t real_length;
52  int nColor;
53  int nSpin;
54 
55  void *clover;
56  void *norm;
57  void *cloverInv;
58  void *invNorm;
59 
60  double csw;
61  bool twisted;
62  double mu2;
63  double rho;
64 
67 
68  mutable double trlog[2];
69 
70  public:
71  CloverField(const CloverFieldParam &param);
72  virtual ~CloverField();
73 
74  void* V(bool inverse=false) { return inverse ? cloverInv : clover; }
75  void* Norm(bool inverse=false) { return inverse ? invNorm : norm; }
76  const void* V(bool inverse=false) const { return inverse ? cloverInv : clover; }
77  const void* Norm(bool inverse=false) const { return inverse ? invNorm : norm; }
78 
83  bool isNative() const;
84 
88  double* TrLog() const { return trlog; }
89 
93  QudaCloverFieldOrder Order() const { return order; }
94 
98  size_t Bytes() const { return bytes; }
99 
103  size_t NormBytes() const { return norm_bytes; }
104 
108  int Ncolor() const { return nColor; }
109 
113  int Nspin() const { return nSpin; }
114 
118  double Csw() const { return csw; }
119 
123  bool Twisted() const { return twisted; }
124 
128  double Mu2() const { return mu2; }
129 
134  double Rho() const { return rho; }
135 
140  void setRho(double rho);
141 
146  double norm1() const;
147 
152  double norm2() const;
153 
158  double abs_max() const;
159 
164  double abs_min() const;
165 
166  };
167 
168  class cudaCloverField : public CloverField {
169 
170  private:
171  void *even, *odd;
172  void *evenNorm, *oddNorm;
173 
174  void *evenInv, *oddInv;
175  void *evenInvNorm, *oddInvNorm;
176 
177  // computes the clover field given the input gauge field
178  void compute(const cudaGaugeField &gauge);
179 
180 #ifdef USE_TEXTURE_OBJECTS
181  cudaTextureObject_t tex;
182  cudaTextureObject_t normTex;
183  cudaTextureObject_t invTex;
184  cudaTextureObject_t invNormTex;
185  cudaTextureObject_t evenTex;
186  cudaTextureObject_t evenNormTex;
187  cudaTextureObject_t oddTex;
188  cudaTextureObject_t oddNormTex;
189  cudaTextureObject_t evenInvTex;
190  cudaTextureObject_t evenInvNormTex;
191  cudaTextureObject_t oddInvTex;
192  cudaTextureObject_t oddInvNormTex;
193  void createTexObject(cudaTextureObject_t &tex, cudaTextureObject_t &texNorm, void *field, void *norm, bool full);
194  void destroyTexObject();
195 #endif
196 
197  public:
198  // create a cudaCloverField from a CloverFieldParam
199  cudaCloverField(const CloverFieldParam &param);
200 
201  virtual ~cudaCloverField();
202 
203 #ifdef USE_TEXTURE_OBJECTS
204  const cudaTextureObject_t& Tex() const { return tex; }
205  const cudaTextureObject_t& NormTex() const { return normTex; }
206  const cudaTextureObject_t& InvTex() const { return invTex; }
207  const cudaTextureObject_t& InvNormTex() const { return invNormTex; }
208  const cudaTextureObject_t& EvenTex() const { return evenTex; }
209  const cudaTextureObject_t& EvenNormTex() const { return evenNormTex; }
210  const cudaTextureObject_t& OddTex() const { return oddTex; }
211  const cudaTextureObject_t& OddNormTex() const { return oddNormTex; }
212  const cudaTextureObject_t& EvenInvTex() const { return evenInvTex; }
213  const cudaTextureObject_t& EvenInvNormTex() const { return evenInvNormTex; }
214  const cudaTextureObject_t& OddInvTex() const { return oddInvTex; }
215  const cudaTextureObject_t& OddInvNormTex() const { return oddInvNormTex; }
216 #endif
217 
223  void copy(const CloverField &src, bool inverse=true);
224 
229  void loadCPUField(const cpuCloverField &cpu);
230 
231 
236  void saveCPUField(cpuCloverField &cpu) const;
237 
238  friend class DiracClover;
239  friend class DiracCloverPC;
240  friend struct FullClover;
241  };
242 
243  // this is a place holder for a future host-side clover object
244  class cpuCloverField : public CloverField {
245 
246  private:
247 
248  public:
249  cpuCloverField(const CloverFieldParam &param);
250  virtual ~cpuCloverField();
251  };
252 
259  double norm1(const CloverField &u, bool inverse=false);
260 
267  double norm2(const CloverField &a, bool inverse=false);
268 
269 
270  // lightweight struct used to send pointers to cuda driver code
271  struct FullClover {
272  void *even;
273  void *odd;
274  void *evenNorm;
275  void *oddNorm;
277  size_t bytes; // sizeof each clover field (per parity)
278  size_t norm_bytes; // sizeof each norm field (per parity)
279  int stride; // stride (volume + pad)
280  double rho; // rho additive factor
281 
282 #ifdef USE_TEXTURE_OBJECTS
283  const cudaTextureObject_t &evenTex;
284  const cudaTextureObject_t &evenNormTex;
285  const cudaTextureObject_t &oddTex;
286  const cudaTextureObject_t &oddNormTex;
287  const cudaTextureObject_t& EvenTex() const { return evenTex; }
288  const cudaTextureObject_t& EvenNormTex() const { return evenNormTex; }
289  const cudaTextureObject_t& OddTex() const { return oddTex; }
290  const cudaTextureObject_t& OddNormTex() const { return oddNormTex; }
291 #endif
292 
293  FullClover(const cudaCloverField &clover, bool inverse=false) :
294  precision(clover.precision), bytes(clover.bytes), norm_bytes(clover.norm_bytes),
295  stride(clover.stride), rho(clover.rho)
296 #ifdef USE_TEXTURE_OBJECTS
297  , evenTex(inverse ? clover.evenInvTex : clover.evenTex)
298  , evenNormTex(inverse ? clover.evenInvNormTex : clover.evenNormTex)
299  , oddTex(inverse ? clover.oddInvTex : clover.oddTex)
300  , oddNormTex(inverse ? clover.oddInvNormTex : clover.oddNormTex)
301 #endif
302  {
303  if (inverse) {
304  even = clover.evenInv;
305  evenNorm = clover.evenInvNorm;
306  odd = clover.oddInv;
307  oddNorm = clover.oddInvNorm;
308  } else {
309  even = clover.even;
310  evenNorm = clover.evenNorm;
311  odd = clover.odd;
312  oddNorm = clover.oddNorm;
313  }
314  }
315  };
316 
317 
318  // driver for computing the clover field from the gauge field
319  void computeClover(CloverField &clover, const GaugeField &gauge, double coeff, QudaFieldLocation location);
320 
321 
336  QudaFieldLocation location, void *Out=0, void *In=0, void *outNorm=0, void *inNorm=0);
337 
338 
339 
347  void cloverInvert(CloverField &clover, bool computeTraceLog);
348 
355  void cloverRho(CloverField &clover, double rho);
356 
374  void computeCloverForce(GaugeField& force, const GaugeField& U,
375  std::vector<ColorSpinorField*> &x, std::vector<ColorSpinorField*> &p,
376  std::vector<double> &coeff);
388  std::vector<ColorSpinorField*> &x,
389  std::vector<ColorSpinorField*> &p,
390  std::vector< std::vector<double> > &coeff);
399  void computeCloverSigmaTrace(GaugeField &output, const CloverField &clover, double coeff);
400 
412  void cloverDerivative(cudaGaugeField &force, cudaGaugeField& gauge, cudaGaugeField& oprod, double coeff, QudaParity parity);
413 
414 } // namespace quda
415 
416 #endif // _CLOVER_QUDA_H
QudaCloverFieldOrder order
Definition: clover_field.h:21
void cloverDerivative(cudaGaugeField &force, cudaGaugeField &gauge, cudaGaugeField &oprod, double coeff, QudaParity parity)
Compute the derivative of the clover matrix in the direction mu,nu and compute the resulting force gi...
double * TrLog() const
Definition: clover_field.h:88
enum QudaPrecision_s QudaPrecision
void * V(bool inverse=false)
Definition: clover_field.h:74
void computeCloverForce(GaugeField &force, const GaugeField &U, std::vector< ColorSpinorField *> &x, std::vector< ColorSpinorField *> &p, std::vector< double > &coeff)
Compute the force contribution from the solver solution fields.
QudaFieldCreate create
Definition: clover_field.h:22
size_t Bytes() const
Definition: clover_field.h:98
FullClover(const cudaCloverField &clover, bool inverse=false)
Definition: clover_field.h:293
__host__ __device__ void copy(T1 &a, const T2 &b)
QudaPrecision precision
Definition: lattice_field.h:51
QudaCloverFieldOrder order
Definition: clover_field.h:65
QudaCloverFieldOrder Order() const
Definition: clover_field.h:93
QudaFieldCreate create
Definition: clover_field.h:66
std::ostream & operator<<(std::ostream &output, const CloverFieldParam &param)
void cloverInvert(CloverField &clover, bool computeTraceLog)
This function compute the Cholesky decomposition of each clover matrix and stores the clover inverse ...
void computeCloverSigmaTrace(GaugeField &output, const CloverField &clover, double coeff)
Compute the matrix tensor field necessary for the force calculation from the clover trace action...
double norm2(const CloverField &a, bool inverse=false)
double Csw() const
Definition: clover_field.h:118
QudaGaugeParam param
Definition: pack_test.cpp:17
int x[QUDA_MAX_DIM]
Definition: lattice_field.h:67
double norm1(const CloverField &u, bool inverse=false)
enum QudaCloverFieldOrder_s QudaCloverFieldOrder
int Ncolor() const
Definition: clover_field.h:108
QudaPrecision precision
Definition: clover_field.h:276
const int nColor
Definition: covdev_test.cpp:75
cpuColorSpinorField * in
void setPrecision(QudaPrecision precision)
Definition: clover_field.h:23
int Nspin() const
Definition: clover_field.h:113
enum QudaParity_s QudaParity
enum QudaFieldLocation_s QudaFieldLocation
cpuColorSpinorField * out
bool twisted
Clover coefficient.
Definition: clover_field.h:17
void cloverRho(CloverField &clover, double rho)
This function adds a real scalar onto the clover diagonal (only to the direct field not the inverse) ...
size_t NormBytes() const
Definition: clover_field.h:103
void * Norm(bool inverse=false)
Definition: clover_field.h:75
double Mu2() const
Definition: clover_field.h:128
const void * V(bool inverse=false) const
Definition: clover_field.h:76
double Rho() const
Definition: clover_field.h:134
QudaPrecision ghost_precision
Definition: lattice_field.h:54
enum QudaFieldCreate_s QudaFieldCreate
const void * Norm(bool inverse=false) const
Definition: clover_field.h:77
void computeCloverSigmaOprod(GaugeField &oprod, std::vector< ColorSpinorField *> &x, std::vector< ColorSpinorField *> &p, std::vector< std::vector< double > > &coeff)
Compute the outer product from the solver solution fields arising from the diagonal term of the fermi...
void computeClover(CloverField &clover, const GaugeField &gauge, double coeff, QudaFieldLocation location)
Definition: clover_quda.cu:204
CloverFieldParam(const CloverFieldParam &param)
Definition: clover_field.h:34
QudaParity parity
Definition: covdev_test.cpp:54
bool Twisted() const
Definition: clover_field.h:123
unsigned long long bytes
Definition: blas_quda.cu:23
void copyGenericClover(CloverField &out, const CloverField &in, bool inverse, QudaFieldLocation location, void *Out=0, void *In=0, void *outNorm=0, void *inNorm=0)
This generic function is used for copying the clover field where in the input and output can be in an...
Definition: copy_clover.cu:175