QUDA  1.0.0
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
gauge_field.h
Go to the documentation of this file.
1 #ifndef _GAUGE_QUDA_H
2 #define _GAUGE_QUDA_H
3 
4 #include <quda_internal.h>
5 #include <quda.h>
6 #include <lattice_field.h>
7 
8 namespace quda {
9 
11 
12  QudaFieldLocation location; // where are we storing the field (CUDA or GPU)?
13  int nColor;
14  int nFace;
15 
21 
22  double anisotropy;
23  double tadpole;
24  void *gauge; // used when we use a reference to an external field
25 
26  QudaFieldCreate create; // used to determine the type of field created
27 
28  QudaFieldGeometry geometry; // whether the field is a scale, vector or tensor
29 
30  // whether we need to compute the fat link maxima
31  // FIXME temporary flag until we have a kernel that can do this, then we just do this in copy()
32  // always set to false, requires external override
34 
37 
40 
42  double i_mu;
43 
45  size_t site_offset;
46 
48  size_t site_size;
49 
50  // Default constructor
51  GaugeFieldParam(void *const h_gauge = NULL) :
54  nColor(3),
55  nFace(0),
56  reconstruct(QUDA_RECONSTRUCT_NO),
58  fixed(QUDA_GAUGE_FIXED_NO),
59  link_type(QUDA_WILSON_LINKS),
60  t_boundary(QUDA_INVALID_T_BOUNDARY),
61  anisotropy(1.0),
62  tadpole(1.0),
63  gauge(h_gauge),
65  geometry(QUDA_VECTOR_GEOMETRY),
66  compute_fat_link_max(false),
67  staggeredPhaseType(QUDA_STAGGERED_PHASE_NO),
68  staggeredPhaseApplied(false),
69  i_mu(0.0),
70  site_offset(0),
71  site_size(0)
72  {
73  }
74 
75  GaugeFieldParam(const GaugeField &u);
76 
77  GaugeFieldParam(const int *x, const QudaPrecision precision, const QudaReconstructType reconstruct, const int pad,
79  LatticeFieldParam(4, x, pad, precision, ghostExchange),
81  nColor(3),
82  nFace(0),
83  reconstruct(reconstruct),
85  fixed(QUDA_GAUGE_FIXED_NO),
86  link_type(QUDA_WILSON_LINKS),
87  t_boundary(QUDA_INVALID_T_BOUNDARY),
88  anisotropy(1.0),
89  tadpole(1.0),
90  gauge(0),
91  create(QUDA_NULL_FIELD_CREATE),
92  geometry(geometry),
93  compute_fat_link_max(false),
94  staggeredPhaseType(QUDA_STAGGERED_PHASE_NO),
95  staggeredPhaseApplied(false),
96  i_mu(0.0),
97  site_offset(0),
98  site_size(0)
99  {
100  }
101 
103  : LatticeFieldParam(param), location(QUDA_CPU_FIELD_LOCATION),
104  nColor(3), nFace(0), reconstruct(QUDA_RECONSTRUCT_NO),
105  order(param.gauge_order), fixed(param.gauge_fix),
106  link_type(link_type_ != QUDA_INVALID_LINKS ? link_type_ : param.type), t_boundary(param.t_boundary),
107  anisotropy(param.anisotropy), tadpole(param.tadpole_coeff), gauge(h_gauge),
109  compute_fat_link_max(false), staggeredPhaseType(param.staggered_phase_type),
110  staggeredPhaseApplied(param.staggered_phase_applied), i_mu(param.i_mu),
111  site_offset(param.gauge_offset), site_size(param.site_size)
112  {
113  switch(link_type) {
114  case QUDA_SU3_LINKS:
115  case QUDA_GENERAL_LINKS:
116  case QUDA_SMEARED_LINKS:
117  case QUDA_MOMENTUM_LINKS:
118  nFace = 1; break;
119  case QUDA_THREE_LINKS:
120  nFace = 3; break;
121  default:
122  errorQuda("Error: invalid link type(%d)\n", link_type);
123  }
124  }
125 
131  void setPrecision(QudaPrecision precision, bool force_native = false)
132  {
133  // is the current status in native field order?
134  bool native = force_native ? true : false;
135  if (precision == QUDA_DOUBLE_PRECISION) {
136  if (order == QUDA_FLOAT2_GAUGE_ORDER) native = true;
137  } else if (precision == QUDA_SINGLE_PRECISION || precision == QUDA_HALF_PRECISION
138  || precision == QUDA_QUARTER_PRECISION) {
139  if (reconstruct == QUDA_RECONSTRUCT_NO) {
140  if (order == QUDA_FLOAT2_GAUGE_ORDER) native = true;
141  } else if (reconstruct == QUDA_RECONSTRUCT_12 || reconstruct == QUDA_RECONSTRUCT_13) {
142  if (order == QUDA_FLOAT4_GAUGE_ORDER) native = true;
143  } else if (reconstruct == QUDA_RECONSTRUCT_8 || reconstruct == QUDA_RECONSTRUCT_9) {
144  if (order == QUDA_FLOAT4_GAUGE_ORDER) native = true;
145  } else if (reconstruct == QUDA_RECONSTRUCT_10) {
146  if (order == QUDA_FLOAT2_GAUGE_ORDER) native = true;
147  }
148  }
149 
150  this->precision = precision;
151  this->ghost_precision = precision;
152 
153  if (native) {
154  order = (precision == QUDA_DOUBLE_PRECISION || reconstruct == QUDA_RECONSTRUCT_NO
155  || reconstruct == QUDA_RECONSTRUCT_10) ?
158  }
159  }
160  };
161 
162  std::ostream& operator<<(std::ostream& output, const GaugeFieldParam& param);
163 
164  class GaugeField : public LatticeField {
165 
166  protected:
167  size_t bytes; // bytes allocated per full field
168  size_t phase_offset; // offset in bytes to gauge phases - useful to keep track of texture alignment
169  size_t phase_bytes; // bytes needed to store the phases
170  size_t length;
171  size_t real_length;
172  int nColor;
173  int nFace;
174  QudaFieldGeometry geometry; // whether the field is a scale, vector or tensor
175 
177  int nInternal; // number of degrees of freedom per link matrix
182 
183  double anisotropy;
184  double tadpole;
185  double fat_link_max;
186 
187  QudaFieldCreate create; // used to determine the type of field created
188 
189  mutable void *ghost[2 * QUDA_MAX_DIM]; // stores the ghost zone of the gauge field (non-native fields only)
190 
191  mutable int ghostFace[QUDA_MAX_DIM]; // the size of each face
192 
197 
202 
209  void exchange(void **recv, void **send, QudaDirection dir) const;
210 
214  double i_mu;
215 
219  size_t site_offset;
220 
224  size_t site_size;
225 
234  void createGhostZone(const int *R, bool no_comms_fill, bool bidir = true) const;
235 
239  void setTuningString();
240 
241  public:
242  GaugeField(const GaugeFieldParam &param);
243  virtual ~GaugeField();
244 
245  virtual void exchangeGhost(QudaLinkDirection = QUDA_LINK_BACKWARDS) = 0;
246  virtual void injectGhost(QudaLinkDirection = QUDA_LINK_BACKWARDS) = 0;
247 
248  size_t Length() const { return length; }
249  int Ncolor() const { return nColor; }
251  QudaGaugeFieldOrder Order() const { return order; }
252  double Anisotropy() const { return anisotropy; }
253  double Tadpole() const { return tadpole; }
254  QudaTboundary TBoundary() const { return t_boundary; }
255  QudaLinkType LinkType() const { return link_type; }
256  QudaGaugeFixed GaugeFixed() const { return fixed; }
258  QudaFieldGeometry Geometry() const { return geometry; }
261 
268  void applyStaggeredPhase(QudaStaggeredPhase phase=QUDA_STAGGERED_PHASE_INVALID);
269 
273  void removeStaggeredPhase();
274 
278  double iMu() const { return i_mu; }
279 
280  const double& LinkMax() const { return fat_link_max; }
281  int Nface() const { return nFace; }
282 
290  virtual void exchangeExtendedGhost(const int *R, bool no_comms_fill = false) = 0;
291 
301  virtual void exchangeExtendedGhost(const int *R, TimeProfile &profile, bool no_comms_fill = false) = 0;
302 
303  void checkField(const LatticeField &) const;
304 
309  bool isNative() const;
310 
311  size_t Bytes() const { return bytes; }
312  size_t PhaseBytes() const { return phase_bytes; }
313  size_t PhaseOffset() const { return phase_offset; }
314 
315  virtual void* Gauge_p() { errorQuda("Not implemented"); return (void*)0;}
316  virtual void* Even_p() { errorQuda("Not implemented"); return (void*)0;}
317  virtual void* Odd_p() { errorQuda("Not implemented"); return (void*)0;}
318 
319  virtual const void* Gauge_p() const { errorQuda("Not implemented"); return (void*)0;}
320  virtual const void* Even_p() const { errorQuda("Not implemented"); return (void*)0;}
321  virtual const void* Odd_p() const { errorQuda("Not implemented"); return (void*)0;}
322 
323  const void** Ghost() const {
324  if ( isNative() ) errorQuda("No ghost zone pointer for quda-native gauge fields");
325  return (const void**)ghost;
326  }
327 
328  void** Ghost() {
329  if ( isNative() ) errorQuda("No ghost zone pointer for quda-native gauge fields");
330  return ghost;
331  }
332 
337  size_t SiteOffset() const { return site_offset; }
338 
343  size_t SiteSize() const { return site_size; }
344 
348  virtual void zero() = 0;
349 
354  virtual void copy(const GaugeField &src) = 0;
355 
361  double norm1(int dim=-1) const;
362 
368  double norm2(int dim=-1) const;
369 
375  double abs_max(int dim=-1) const;
376 
382  double abs_min(int dim=-1) const;
383 
392  uint64_t checksum(bool mini=false) const;
393 
400  static GaugeField* Create(const GaugeFieldParam &param);
401 
402  };
403 
404  class cudaGaugeField : public GaugeField {
405 
406  private:
407  void *gauge;
408  void *gauge_h; // mapped-memory pointer when allocating on the host
409  void *even;
410  void *odd;
411 
415  void zeroPad();
416 
417 #ifdef USE_TEXTURE_OBJECTS
418  cudaTextureObject_t tex;
419  cudaTextureObject_t evenTex;
420  cudaTextureObject_t oddTex;
421  cudaTextureObject_t phaseTex;
422  cudaTextureObject_t evenPhaseTex;
423  cudaTextureObject_t oddPhaseTex;
424  void createTexObject(cudaTextureObject_t &tex, void *gauge, bool full, bool isPhase=false);
425  void destroyTexObject();
426 #endif
427 
428  public:
430  virtual ~cudaGaugeField();
431 
437  void exchangeGhost(QudaLinkDirection link_direction = QUDA_LINK_BACKWARDS);
438 
445  void injectGhost(QudaLinkDirection link_direction = QUDA_LINK_BACKWARDS);
446 
456  void createComms(const int *R, bool no_comms_fill, bool bidir=true);
457 
466  void allocateGhostBuffer(const int *R, bool no_comms_fill, bool bidir=true) const;
467 
473  void recvStart(int dim, int dir);
474 
482  void sendStart(int dim, int dir, cudaStream_t *stream_p=nullptr);
483 
489  void commsComplete(int dim, int dir);
490 
498  void exchangeExtendedGhost(const int *R, bool no_comms_fill=false);
499 
509  void exchangeExtendedGhost(const int *R, TimeProfile &profile, bool no_comms_fill=false);
510 
515  void copy(const GaugeField &src);
516 
521  void loadCPUField(const cpuGaugeField &cpu);
522 
529  void loadCPUField(const cpuGaugeField &cpu, TimeProfile &profile);
530 
535  void saveCPUField(cpuGaugeField &cpu) const;
536 
543  void saveCPUField(cpuGaugeField &cpu, TimeProfile &profile) const;
544 
545  // (ab)use with care
546  void* Gauge_p() { return gauge; }
547  void* Even_p() { return even; }
548  void* Odd_p() { return odd; }
549 
550  const void* Gauge_p() const { return gauge; }
551  const void* Even_p() const { return even; }
552  const void *Odd_p() const { return odd; }
553 
554 #ifdef USE_TEXTURE_OBJECTS
555  const cudaTextureObject_t& Tex() const { return tex; }
556  const cudaTextureObject_t& EvenTex() const { return evenTex; }
557  const cudaTextureObject_t& OddTex() const { return oddTex; }
558  const cudaTextureObject_t& EvenPhaseTex() const { return evenPhaseTex; }
559  const cudaTextureObject_t& OddPhaseTex() const { return oddPhaseTex; }
560 #endif
561 
562  void setGauge(void* _gauge); //only allowed when create== QUDA_REFERENCE_FIELD_CREATE
563 
567  void zero();
568 
572  void backup() const;
573 
577  void restore() const;
578  };
579 
580  class cpuGaugeField : public GaugeField {
581 
582  friend void cudaGaugeField::copy(const GaugeField &cpu);
583  friend void cudaGaugeField::loadCPUField(const cpuGaugeField &cpu);
584  friend void cudaGaugeField::saveCPUField(cpuGaugeField &cpu) const;
585 
586  private:
587  void **gauge; // the actual gauge field
588 
589  public:
597  cpuGaugeField(const GaugeFieldParam &param);
598  virtual ~cpuGaugeField();
599 
605  void exchangeGhost(QudaLinkDirection link_direction = QUDA_LINK_BACKWARDS);
606 
613  void injectGhost(QudaLinkDirection link_direction = QUDA_LINK_BACKWARDS);
614 
623  void exchangeExtendedGhost(const int *R, bool no_comms_fill=false);
624 
634  void exchangeExtendedGhost(const int *R, TimeProfile &profile, bool no_comms_fill=false);
635 
640  void copy(const GaugeField &src);
641 
642  void* Gauge_p() { return gauge; }
643  const void* Gauge_p() const { return gauge; }
644 
645  void setGauge(void** _gauge); //only allowed when create== QUDA_REFERENCE_FIELD_CREATE
646 
650  void zero();
651 
655  void backup() const;
656 
660  void restore() const;
661  };
662 
669  double norm1(const GaugeField &u);
670 
677  double norm2(const GaugeField &u);
678 
684  void ax(const double &a, GaugeField &u);
685 
698  void copyGenericGauge(GaugeField &out, const GaugeField &in, QudaFieldLocation location, void *Out = 0, void *In = 0,
699  void **ghostOut = 0, void **ghostIn = 0, int type = 0);
709  void copyExtendedGauge(GaugeField &out, const GaugeField &in,
710  QudaFieldLocation location, void *Out=0, void *In=0);
711 
724  void extractGaugeGhost(const GaugeField &u, void **ghost, bool extract=true, int offset=0);
725 
734  void extractExtendedGaugeGhost(const GaugeField &u, int dim, const int *R, void **ghost, bool extract);
735 
740  void applyGaugePhase(GaugeField &u);
741 
751  uint64_t Checksum(const GaugeField &u, bool mini=false);
752 
753 } // namespace quda
754 
755 #endif // _GAUGE_QUDA_H
virtual void * Even_p()
Definition: gauge_field.h:316
QudaTboundary t_boundary
Definition: gauge_field.h:20
void extractGaugeGhost(const GaugeField &u, void **ghost, bool extract=true, int offset=0)
QudaGhostExchange ghostExchange
Definition: lattice_field.h:76
enum QudaPrecision_s QudaPrecision
void saveCPUField(cpuGaugeField &cpu) const
Upload from this field into a CPU field.
QudaGaugeFieldOrder FieldOrder() const
Definition: gauge_field.h:257
void copyGenericGauge(GaugeField &out, const GaugeField &in, QudaFieldLocation location, void *Out=0, void *In=0, void **ghostOut=0, void **ghostIn=0, int type=0)
Definition: copy_gauge.cu:41
const void * Gauge_p() const
Definition: gauge_field.h:643
const void * Gauge_p() const
Definition: gauge_field.h:550
#define errorQuda(...)
Definition: util_quda.h:121
enum QudaLinkDirection_s QudaLinkDirection
virtual void * Odd_p()
Definition: gauge_field.h:317
QudaReconstructType reconstruct
Definition: gauge_field.h:176
bool staggeredPhaseApplied
Definition: gauge_field.h:201
QudaLinkType LinkType() const
Definition: gauge_field.h:255
uint64_t Checksum(const GaugeField &u, bool mini=false)
Definition: checksum.cu:34
__host__ __device__ void copy(T1 &a, const T2 &b)
QudaPrecision precision
Definition: lattice_field.h:51
static int R[4]
QudaFieldGeometry Geometry() const
Definition: gauge_field.h:258
void applyGaugePhase(GaugeField &u)
Definition: gauge_phase.cu:223
int length[]
enum QudaTboundary_s QudaTboundary
QudaStaggeredPhase staggeredPhaseType
Definition: gauge_field.h:196
int Nface() const
Definition: gauge_field.h:281
std::ostream & operator<<(std::ostream &output, const CloverFieldParam &param)
double norm2(const CloverField &a, bool inverse=false)
QudaFieldGeometry geometry
Definition: gauge_field.h:174
QudaGaugeParam param
Definition: pack_test.cpp:17
virtual const void * Odd_p() const
Definition: gauge_field.h:321
int Ncolor() const
Definition: gauge_field.h:249
int x[QUDA_MAX_DIM]
Definition: lattice_field.h:67
double norm1(const CloverField &u, bool inverse=false)
enum QudaDirection_s QudaDirection
void loadCPUField(const cpuGaugeField &cpu)
Download into this field from a CPU field.
void ax(const double &a, GaugeField &u)
Scale the gauge field by the scalar a.
QudaGaugeFixed fixed
Definition: gauge_field.h:18
enum QudaStaggeredPhase_s QudaStaggeredPhase
size_t Bytes() const
Definition: gauge_field.h:311
bool StaggeredPhaseApplied() const
Definition: gauge_field.h:260
double iMu() const
Definition: gauge_field.h:278
cpuColorSpinorField * in
void extractExtendedGaugeGhost(const GaugeField &u, int dim, const int *R, void **ghost, bool extract)
enum QudaGhostExchange_s QudaGhostExchange
QudaFieldLocation location
Definition: gauge_field.h:12
double Anisotropy() const
Definition: gauge_field.h:252
QudaGaugeFieldOrder order
Definition: gauge_field.h:17
void exchangeExtendedGhost(cudaColorSpinorField *spinor, int R[], int parity, cudaStream_t *stream_p)
enum QudaGaugeFixed_s QudaGaugeFixed
const void ** Ghost() const
Definition: gauge_field.h:323
enum QudaGaugeFieldOrder_s QudaGaugeFieldOrder
QudaStaggeredPhase staggeredPhaseType
Definition: gauge_field.h:36
enum QudaLinkType_s QudaLinkType
size_t SiteSize() const
Definition: gauge_field.h:343
size_t PhaseOffset() const
Definition: gauge_field.h:313
enum QudaFieldLocation_s QudaFieldLocation
void setPrecision(QudaPrecision precision, bool force_native=false)
Helper function for setting the precision and corresponding field order for QUDA internal fields...
Definition: gauge_field.h:131
size_t PhaseBytes() const
Definition: gauge_field.h:312
cpuColorSpinorField * out
QudaLinkType link_type
Definition: gauge_field.h:180
const double & LinkMax() const
Definition: gauge_field.h:280
QudaGaugeFieldOrder gauge_order
size_t Length() const
Definition: gauge_field.h:248
enum QudaReconstructType_s QudaReconstructType
Main header file for the QUDA library.
QudaTboundary t_boundary
Definition: gauge_field.h:181
GaugeFieldParam(void *const h_gauge=NULL)
Definition: gauge_field.h:51
QudaLinkType link_type
Definition: gauge_field.h:19
GaugeFieldParam(const int *x, const QudaPrecision precision, const QudaReconstructType reconstruct, const int pad, const QudaFieldGeometry geometry, const QudaGhostExchange ghostExchange=QUDA_GHOST_EXCHANGE_PAD)
Definition: gauge_field.h:77
bool commsComplete(cudaColorSpinorField &in, const Dslash &dslash, int dim, int dir, bool gdr_send, bool gdr_recv, bool zero_copy_recv, bool async, int scatterIndex=-1)
Wrapper for querying if communication is finished in the dslash, and if it is take the appropriate ac...
const void * Odd_p() const
Definition: gauge_field.h:552
double Tadpole() const
Definition: gauge_field.h:253
QudaFieldCreate create
Definition: gauge_field.h:187
virtual const void * Even_p() const
Definition: gauge_field.h:320
QudaReconstructType reconstruct
Definition: gauge_field.h:16
virtual const void * Gauge_p() const
Definition: gauge_field.h:319
QudaFieldCreate create
Definition: gauge_field.h:26
QudaPrecision ghost_precision
Definition: lattice_field.h:54
QudaReconstructType Reconstruct() const
Definition: gauge_field.h:250
enum QudaFieldCreate_s QudaFieldCreate
enum QudaFieldGeometry_s QudaFieldGeometry
virtual void * Gauge_p()
Definition: gauge_field.h:315
QudaGaugeFieldOrder Order() const
Definition: gauge_field.h:251
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5...
QudaFieldGeometry geometry
Definition: gauge_field.h:28
QudaTboundary TBoundary() const
Definition: gauge_field.h:254
QudaStaggeredPhase StaggeredPhase() const
Definition: gauge_field.h:259
QudaGaugeFieldOrder order
Definition: gauge_field.h:178
QudaGaugeFixed GaugeFixed() const
Definition: gauge_field.h:256
void copy(const GaugeField &src)
QudaGaugeFixed fixed
Definition: gauge_field.h:179
__device__ __host__ void zero(vector_type< scalar, n > &v)
Definition: cub_helper.cuh:54
GaugeFieldParam(void *h_gauge, const QudaGaugeParam &param, QudaLinkType link_type_=QUDA_INVALID_LINKS)
Definition: gauge_field.h:102
void ** Ghost()
Definition: gauge_field.h:328
void copyExtendedGauge(GaugeField &out, const GaugeField &in, QudaFieldLocation location, void *Out=0, void *In=0)
float fat_link_max
const void * Even_p() const
Definition: gauge_field.h:551
unsigned long long bytes
Definition: blas_quda.cu:23
size_t SiteOffset() const
Definition: gauge_field.h:337