QUDA  v0.7.0
A library for QCD on GPUs
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Pages
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 
17 //for twisted mass only:
18  bool twisted; // whether to create twisted mass clover
19  double mu2;
20 
24  this->precision = precision;
25  order = (precision == QUDA_DOUBLE_PRECISION) ?
27  }
28  };
29 
30  std::ostream& operator<<(std::ostream& output, const CloverFieldParam& param);
31 
32  class CloverField : public LatticeField {
33 
34  protected:
35  size_t bytes; // bytes allocated per clover full field
36  size_t norm_bytes; // sizeof each norm full field
37  int length;
39  int nColor;
40  int nSpin;
41 
42  void *clover;
43  void *norm;
44  void *cloverInv;
45  void *invNorm;
46 //new!
47  bool twisted;
48  double mu2;
49 
52 
53  double *trlog;
54 
55  public:
57  virtual ~CloverField();
58 
59  void* V(bool inverse=false) { return inverse ? cloverInv : clover; }
60  void* Norm(bool inverse=false) { return inverse ? invNorm : norm; }
61  const void* V(bool inverse=false) const { return inverse ? cloverInv : clover; }
62  const void* Norm(bool inverse=false) const { return inverse ? invNorm : norm; }
63 
64  double* TrLog() const { return trlog; }
65 
66  QudaCloverFieldOrder Order() const { return order; }
67  size_t Bytes() const { return bytes; }
68  size_t NormBytes() const { return norm_bytes; }
69 //new!
70  bool Twisted() const {return twisted; }
71  double Mu2() const {return mu2; }
72  };
73 
74  class cudaCloverField : public CloverField {
75 
76  private:
77  void *even, *odd;
78  void *evenNorm, *oddNorm;
79 
80  void *evenInv, *oddInv;
81  void *evenInvNorm, *oddInvNorm;
82 
83  // computes the clover field given the input gauge field
84  void compute(const cudaGaugeField &gauge);
85 
86 #ifdef USE_TEXTURE_OBJECTS
87  cudaTextureObject_t evenTex;
88  cudaTextureObject_t evenNormTex;
89  cudaTextureObject_t oddTex;
90  cudaTextureObject_t oddNormTex;
91  cudaTextureObject_t evenInvTex;
92  cudaTextureObject_t evenInvNormTex;
93  cudaTextureObject_t oddInvTex;
94  cudaTextureObject_t oddInvNormTex;
95  void createTexObject(cudaTextureObject_t &tex, cudaTextureObject_t &texNorm, void *field, void *norm);
96  void destroyTexObject();
97 #endif
98 
99  public:
100  // create a cudaCloverField from a CloverFieldParam
102 
103  virtual ~cudaCloverField();
104 
105 #ifdef USE_TEXTURE_OBJECTS
106  const cudaTextureObject_t& EvenTex() const { return evenTex; }
107  const cudaTextureObject_t& EvenNormTex() const { return evenNormTex; }
108  const cudaTextureObject_t& OddTex() const { return oddTex; }
109  const cudaTextureObject_t& OddNormTex() const { return oddNormTex; }
110  const cudaTextureObject_t& EvenInvTex() const { return evenInvTex; }
111  const cudaTextureObject_t& EvenInvNormTex() const { return evenInvNormTex; }
112  const cudaTextureObject_t& OddInvTex() const { return oddInvTex; }
113  const cudaTextureObject_t& OddInvNormTex() const { return oddInvNormTex; }
114 #endif
115 
120  void copy(const CloverField &src, bool inverse=true);
121 
126  void loadCPUField(const cpuCloverField &cpu);
127 
128 
133  void saveCPUField(cpuCloverField &cpu);
134 
135  friend class DiracClover;
136  friend class DiracCloverPC;
137  friend struct FullClover;
138  };
139 
140  // this is a place holder for a future host-side clover object
141  class cpuCloverField : public CloverField {
142 
143  private:
144 
145  public:
147  virtual ~cpuCloverField();
148  };
149 
150  // lightweight struct used to send pointers to cuda driver code
151  struct FullClover {
152  void *even;
153  void *odd;
154  void *evenNorm;
155  void *oddNorm;
157  size_t bytes; // sizeof each clover field (per parity)
158  size_t norm_bytes; // sizeof each norm field (per parity)
159  int stride; // stride (volume + pad)
160 
161 #ifdef USE_TEXTURE_OBJECTS
162  const cudaTextureObject_t &evenTex;
163  const cudaTextureObject_t &evenNormTex;
164  const cudaTextureObject_t &oddTex;
165  const cudaTextureObject_t &oddNormTex;
166  const cudaTextureObject_t& EvenTex() const { return evenTex; }
167  const cudaTextureObject_t& EvenNormTex() const { return evenNormTex; }
168  const cudaTextureObject_t& OddTex() const { return oddTex; }
169  const cudaTextureObject_t& OddNormTex() const { return oddNormTex; }
170 #endif
171 
172  FullClover(const cudaCloverField &clover, bool inverse=false) :
173  precision(clover.precision), bytes(clover.bytes), norm_bytes(clover.norm_bytes), stride(clover.stride)
174 #ifdef USE_TEXTURE_OBJECTS
175  , evenTex(inverse ? clover.evenInvTex : clover.evenTex)
176  , evenNormTex(inverse ? clover.evenInvNormTex : clover.evenNormTex)
177  , oddTex(inverse ? clover.oddInvTex : clover.oddTex)
178  , oddNormTex(inverse ? clover.oddInvNormTex : clover.oddNormTex)
179 #endif
180  {
181  if (inverse) {
182  even = clover.evenInv;
183  evenNorm = clover.evenInvNorm;
184  odd = clover.oddInv;
185  oddNorm = clover.oddInvNorm;
186  } else {
187  even = clover.even;
188  evenNorm = clover.evenNorm;
189  odd = clover.odd;
190  oddNorm = clover.oddNorm;
191  }
192  }
193  };
194 
195  // driver for computing the clover field from the gauge field
196  void computeClover(CloverField &clover, const GaugeField &gauge, double coeff, QudaFieldLocation location);
197 
198 
199  void computeCloverSigmaTrace(GaugeField &gauge, const CloverField &clover, int dir1, int dir2, QudaFieldLocation location);
200 
213  void copyGenericClover(CloverField &out, const CloverField &in, bool inverse, QudaFieldLocation location,
214  void *Out=0, void *In=0, void *outNorm=0, void *inNorm=0);
215 
216 
217  void cloverDerivative(cudaGaugeField &out, cudaGaugeField& gauge, cudaGaugeField& oprod, int mu, int nu, double coeff, QudaParity parity, int conjugate);
218 
219 
220 
228  void cloverInvert(CloverField &clover, bool computeTraceLog, QudaFieldLocation location);
229 
230 } // namespace quda
231 
232 #endif // _CLOVER_QUDA_H
virtual ~CloverField()
QudaCloverFieldOrder order
Definition: clover_field.h:21
cudaCloverField(const CloverFieldParam &param)
enum QudaPrecision_s QudaPrecision
bool Twisted() const
Definition: clover_field.h:70
void * V(bool inverse=false)
Definition: clover_field.h:59
size_t Bytes() const
Definition: clover_field.h:67
void saveCPUField(cpuCloverField &cpu)
QudaFieldCreate create
Definition: clover_field.h:22
FullClover(const cudaCloverField &clover, bool inverse=false)
Definition: clover_field.h:172
void loadCPUField(const cpuCloverField &cpu)
__global__ void const RealA *const const RealA *const const RealA *const const RealB *const const RealB *const int int mu
QudaPrecision precision
Definition: lattice_field.h:41
QudaCloverFieldOrder order
Definition: clover_field.h:50
CloverField(const CloverFieldParam &param)
void cloverDerivative(cudaGaugeField &out, cudaGaugeField &gauge, cudaGaugeField &oprod, int mu, int nu, double coeff, QudaParity parity, int conjugate)
QudaFieldCreate create
Definition: clover_field.h:51
std::ostream & operator<<(std::ostream &output, const CloverFieldParam &param)
QudaGaugeParam param
Definition: pack_test.cpp:17
enum QudaCloverFieldOrder_s QudaCloverFieldOrder
void cloverInvert(CloverField &clover, bool computeTraceLog, QudaFieldLocation location)
const QudaFieldLocation location
Definition: pack_test.cpp:46
QudaPrecision precision
Definition: clover_field.h:156
QudaCloverFieldOrder Order() const
Definition: clover_field.h:66
cpuColorSpinorField * in
void setPrecision(QudaPrecision precision)
Definition: clover_field.h:23
__constant__ double coeff
double Mu2() const
Definition: clover_field.h:71
enum QudaParity_s QudaParity
void computeCloverSigmaTrace(GaugeField &gauge, const CloverField &clover, int dir1, int dir2, QudaFieldLocation location)
cpuCloverField(const CloverFieldParam &param)
void copy(const CloverField &src, bool inverse=true)
const void * V(bool inverse=false) const
Definition: clover_field.h:61
enum QudaFieldLocation_s QudaFieldLocation
cpuColorSpinorField * out
void * Norm(bool inverse=false)
Definition: clover_field.h:60
enum QudaFieldCreate_s QudaFieldCreate
size_t NormBytes() const
Definition: clover_field.h:68
void computeClover(CloverField &clover, const GaugeField &gauge, double coeff, QudaFieldLocation location)
Definition: clover_quda.cu:602
double * TrLog() const
Definition: clover_field.h:64
const QudaParity parity
Definition: dslash_test.cpp:29
void * gauge[4]
Definition: su3_test.cpp:15
const void * Norm(bool inverse=false) const
Definition: clover_field.h:62
void copyGenericClover(CloverField &out, const CloverField &in, bool inverse, QudaFieldLocation location, void *Out=0, void *In=0, void *outNorm=0, void *inNorm=0)
Definition: copy_clover.cu:182