QUDA  1.0.0
All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros
gauge_field.cpp
Go to the documentation of this file.
1 #include <gauge_field.h>
2 #include <typeinfo>
3 #include <blas_quda.h>
4 
5 namespace quda {
6 
9  location(u.Location()),
10  nColor(u.Ncolor()),
11  nFace(u.Nface()),
12  reconstruct(u.Reconstruct()),
13  order(u.Order()),
14  fixed(u.GaugeFixed()),
15  link_type(u.LinkType()),
16  t_boundary(u.TBoundary()),
17  anisotropy(u.Anisotropy()),
18  tadpole(u.Tadpole()),
19  gauge(NULL),
20  create(QUDA_NULL_FIELD_CREATE),
21  geometry(u.Geometry()),
22  compute_fat_link_max(false),
23  staggeredPhaseType(u.StaggeredPhase()),
24  staggeredPhaseApplied(u.StaggeredPhaseApplied()),
25  i_mu(u.iMu()),
26  site_offset(u.SiteOffset()),
27  site_size(u.SiteSize())
28  { }
29 
31  LatticeField(param),
32  bytes(0),
33  phase_offset(0),
34  phase_bytes(0),
35  nColor(param.nColor),
36  nFace(param.nFace),
37  geometry(param.geometry),
38  reconstruct(param.reconstruct),
40  order(param.order),
41  fixed(param.fixed),
42  link_type(param.link_type),
43  t_boundary(param.t_boundary),
44  anisotropy(param.anisotropy),
45  tadpole(param.tadpole),
47  create(param.create),
50  i_mu(param.i_mu),
51  site_offset(param.site_offset),
52  site_size(param.site_size)
53  {
54  if (ghost_precision != precision) ghost_precision = precision; // gauge fields require matching precision
55 
56  if (link_type != QUDA_COARSE_LINKS && nColor != 3)
57  errorQuda("nColor must be 3, not %d for this link type", nColor);
58  if (nDim != 4)
59  errorQuda("Number of dimensions must be 4 not %d", nDim);
60  if (link_type != QUDA_WILSON_LINKS && anisotropy != 1.0)
61  errorQuda("Anisotropy only supported for Wilson links");
63  errorQuda("Temporal gauge fixing only supported for Wilson links");
66  length = 2*stride*nInternal; // two comes from being full lattice
67  } else if (geometry == QUDA_VECTOR_GEOMETRY) {
69  length = 2*nDim*stride*nInternal; // two comes from being full lattice
70  } else if (geometry == QUDA_TENSOR_GEOMETRY) {
72  length = 2*(nDim*(nDim-1)/2)*stride*nInternal; // two comes from being full lattice
73  } else if (geometry == QUDA_COARSE_GEOMETRY) {
75  length = 2*2*nDim*stride*nInternal; //two comes from being full lattice
76  }
77 
79  // Need to adjust the phase alignment as well.
80  int half_phase_bytes
81  = (length / (2 * reconstruct)) * precision; // number of bytes needed to store phases for a single parity
82  int half_gauge_bytes = (length / 2) * precision
83  - half_phase_bytes; // number of bytes needed to store the gauge field for a single parity excluding the phases
84  // Adjust the alignments for the gauge and phase separately
85  half_phase_bytes = ((half_phase_bytes + (512-1))/512)*512;
86  half_gauge_bytes = ((half_gauge_bytes + (512-1))/512)*512;
87 
88  phase_offset = half_gauge_bytes;
89  phase_bytes = half_phase_bytes*2;
90  bytes = (half_gauge_bytes + half_phase_bytes)*2;
91  } else {
93  if (isNative()) bytes = 2*ALIGNMENT_ADJUST(bytes/2);
94  }
96 
98  }
99 
101 
102  }
103 
106  int aux_string_n = TuneKey::aux_n / 2;
107  int check = snprintf(aux_string, aux_string_n, "vol=%d,stride=%d,precision=%d,geometry=%d,Nc=%d",
109  if (check < 0 || check >= aux_string_n) errorQuda("Error writing aux string");
110  }
111 
112  void GaugeField::createGhostZone(const int *R, bool no_comms_fill, bool bidir) const
113  {
114  if (typeid(*this) == typeid(cpuGaugeField)) return;
115 
116  // if this is not a bidirectional exchange then we are doing a
117  // scalar exchange, e.g., only the link matrix in the direcion we
118  // are exchanging is exchanged, and none of the orthogonal links
120 
121  // calculate size of ghost zone required
122  ghost_bytes_old = ghost_bytes; // save for subsequent resize checking
123  ghost_bytes = 0;
124  for (int i=0; i<nDim; i++) {
125  ghost_face_bytes[i] = 0;
126  if ( !(comm_dim_partitioned(i) || (no_comms_fill && R[i])) ) ghostFace[i] = 0;
127  else ghostFace[i] = surface[i] * R[i]; // includes the radius (unlike ColorSpinorField)
128 
129  ghostOffset[i][0] = (i == 0) ? 0 : ghostOffset[i-1][1] + ghostFace[i-1]*geometry_comms*nInternal;
130  ghostOffset[i][1] = (bidir ? ghostOffset[i][0] + ghostFace[i]*geometry_comms*nInternal : ghostOffset[i][0]);
131 
132  ghost_face_bytes[i] = ghostFace[i] * geometry_comms * nInternal * ghost_precision;
133  ghost_bytes += (bidir ? 2 : 1 ) * ghost_face_bytes[i]; // factor of two from direction
134  }
135 
137  } // createGhostZone
138 
140  if (staggeredPhaseApplied) errorQuda("Staggered phases already applied");
141 
142  if (phase != QUDA_STAGGERED_PHASE_INVALID) staggeredPhaseType = phase;
143  applyGaugePhase(*this);
145  if (typeid(*this)==typeid(cudaGaugeField)) {
146  static_cast<cudaGaugeField&>(*this).exchangeGhost();
147  } else {
148  static_cast<cpuGaugeField&>(*this).exchangeGhost();
149  }
150  }
151  staggeredPhaseApplied = true;
152  }
153 
155  if (!staggeredPhaseApplied) errorQuda("No staggered phases to remove");
156  applyGaugePhase(*this);
158  if (typeid(*this)==typeid(cudaGaugeField)) {
159  static_cast<cudaGaugeField&>(*this).exchangeGhost();
160  } else {
161  static_cast<cpuGaugeField&>(*this).exchangeGhost();
162  }
163  }
164  staggeredPhaseApplied = false;
165  }
166 
167  bool GaugeField::isNative() const {
169  if (order == QUDA_FLOAT2_GAUGE_ORDER) return true;
173  if (order == QUDA_FLOAT2_GAUGE_ORDER) return true;
175  if (order == QUDA_FLOAT4_GAUGE_ORDER) return true;
177  if (order == QUDA_FLOAT4_GAUGE_ORDER) return true;
178  } else if (reconstruct == QUDA_RECONSTRUCT_10) {
179  if (order == QUDA_FLOAT2_GAUGE_ORDER) return true;
180  }
181  }
182  return false;
183  }
184 
185  void GaugeField::exchange(void **ghost_link, void **link_sendbuf, QudaDirection dir) const {
186  MsgHandle *mh_send[4];
187  MsgHandle *mh_recv[4];
188  size_t bytes[4];
189 
190  for (int i=0; i<nDimComms; i++) bytes[i] = 2*nFace*surfaceCB[i]*nInternal*precision;
191 
192  // in general (standard ghost exchange) we always do the exchange
193  // even if a dimension isn't partitioned. However, this breaks
194  // GaugeField::injectGhost(), so when transferring backwards we
195  // only exchange if a dimension is partitioned. FIXME: this
196  // should probably be cleaned up.
197  bool no_comms_fill = (dir == QUDA_BACKWARDS) ? false : true;
198 
199  void *send[4];
200  void *receive[4];
202  for (int i=0; i<nDimComms; i++) {
203  if (comm_dim_partitioned(i)) {
204  send[i] = link_sendbuf[i];
205  receive[i] = ghost_link[i];
206  } else {
207  if (no_comms_fill) memcpy(ghost_link[i], link_sendbuf[i], bytes[i]);
208  }
209  }
210  } else { // FIXME for CUDA field copy back to the CPU
211  for (int i=0; i<nDimComms; i++) {
212  if (comm_dim_partitioned(i)) {
213  send[i] = pool_pinned_malloc(bytes[i]);
214  receive[i] = pool_pinned_malloc(bytes[i]);
215  qudaMemcpy(send[i], link_sendbuf[i], bytes[i], cudaMemcpyDeviceToHost);
216  } else {
217  if (no_comms_fill) qudaMemcpy(ghost_link[i], link_sendbuf[i], bytes[i], cudaMemcpyDeviceToDevice);
218  }
219  }
220  }
221 
222  for (int i=0; i<nDimComms; i++) {
223  if (!comm_dim_partitioned(i)) continue;
224  if (dir == QUDA_FORWARDS) {
225  mh_send[i] = comm_declare_send_relative(send[i], i, +1, bytes[i]);
226  mh_recv[i] = comm_declare_receive_relative(receive[i], i, -1, bytes[i]);
227  } else if (dir == QUDA_BACKWARDS) {
228  mh_send[i] = comm_declare_send_relative(send[i], i, -1, bytes[i]);
229  mh_recv[i] = comm_declare_receive_relative(receive[i], i, +1, bytes[i]);
230  } else {
231  errorQuda("Unsuported dir=%d", dir);
232  }
233 
234  }
235 
236  for (int i=0; i<nDimComms; i++) {
237  if (!comm_dim_partitioned(i)) continue;
238  comm_start(mh_send[i]);
239  comm_start(mh_recv[i]);
240  }
241 
242  for (int i=0; i<nDimComms; i++) {
243  if (!comm_dim_partitioned(i)) continue;
244  comm_wait(mh_send[i]);
245  comm_wait(mh_recv[i]);
246  }
247 
249  for (int i=0; i<nDimComms; i++) {
250  if (!comm_dim_partitioned(i)) continue;
251  qudaMemcpy(ghost_link[i], receive[i], bytes[i], cudaMemcpyHostToDevice);
252  pool_pinned_free(send[i]);
253  pool_pinned_free(receive[i]);
254  }
255  }
256 
257  for (int i=0; i<nDimComms; i++) {
258  if (!comm_dim_partitioned(i)) continue;
259  comm_free(mh_send[i]);
260  comm_free(mh_recv[i]);
261  }
262 
263  }
264 
265  void GaugeField::checkField(const LatticeField &l) const {
267  try {
268  const GaugeField &g = dynamic_cast<const GaugeField&>(l);
269  if (g.link_type != link_type) errorQuda("link_type does not match %d %d", link_type, g.link_type);
270  if (g.nColor != nColor) errorQuda("nColor does not match %d %d", nColor, g.nColor);
271  if (g.nFace != nFace) errorQuda("nFace does not match %d %d", nFace, g.nFace);
272  if (g.fixed != fixed) errorQuda("fixed does not match %d %d", fixed, g.fixed);
273  if (g.t_boundary != t_boundary) errorQuda("t_boundary does not match %d %d", t_boundary, g.t_boundary);
274  if (g.anisotropy != anisotropy) errorQuda("anisotropy does not match %e %e", anisotropy, g.anisotropy);
275  if (g.tadpole != tadpole) errorQuda("tadpole does not match %e %e", tadpole, g.tadpole);
276  }
277  catch(std::bad_cast &e) {
278  errorQuda("Failed to cast reference to GaugeField");
279  }
280  }
281 
282  std::ostream& operator<<(std::ostream& output, const GaugeFieldParam& param) {
283  output << static_cast<const LatticeFieldParam &>(param);
284  output << "nColor = " << param.nColor << std::endl;
285  output << "nFace = " << param.nFace << std::endl;
286  output << "reconstruct = " << param.reconstruct << std::endl;
287  int nInternal = (param.reconstruct != QUDA_RECONSTRUCT_NO ?
288  param.reconstruct : param.nColor * param.nColor * 2);
289  output << "nInternal = " << nInternal << std::endl;
290  output << "order = " << param.order << std::endl;
291  output << "fixed = " << param.fixed << std::endl;
292  output << "link_type = " << param.link_type << std::endl;
293  output << "t_boundary = " << param.t_boundary << std::endl;
294  output << "anisotropy = " << param.anisotropy << std::endl;
295  output << "tadpole = " << param.tadpole << std::endl;
296  output << "create = " << param.create << std::endl;
297  output << "geometry = " << param.geometry << std::endl;
298  output << "staggeredPhaseType = " << param.staggeredPhaseType << std::endl;
299  output << "staggeredPhaseApplied = " << param.staggeredPhaseApplied << std::endl;
300 
301  return output; // for multiple << operators.
302  }
303 
306  errorQuda("Not implemented for this order %d", a.FieldOrder());
307 
308  if (a.LinkType() == QUDA_COARSE_LINKS) errorQuda("Not implemented for coarse-link type");
309  if (a.Ncolor() != 3) errorQuda("Not implemented for Ncolor = %d", a.Ncolor());
310 
312  errorQuda("Casting a GaugeField into ColorSpinorField not possible in half or quarter precision");
313 
314  ColorSpinorParam spinor_param;
315  spinor_param.nColor = (a.Geometry()*a.Reconstruct())/2;
316  spinor_param.nSpin = 1;
317  spinor_param.nDim = a.Ndim();
318  for (int d=0; d<a.Ndim(); d++) spinor_param.x[d] = a.X()[d];
319  spinor_param.setPrecision(a.Precision());
320  spinor_param.pad = a.Pad();
321  spinor_param.siteSubset = QUDA_FULL_SITE_SUBSET;
322  spinor_param.siteOrder = QUDA_EVEN_ODD_SITE_ORDER;
323  spinor_param.fieldOrder = (a.Precision() == QUDA_DOUBLE_PRECISION || spinor_param.nSpin == 1) ?
325  spinor_param.gammaBasis = QUDA_UKQCD_GAMMA_BASIS;
326  spinor_param.create = QUDA_REFERENCE_FIELD_CREATE;
327  spinor_param.v = (void*)a.Gauge_p();
328  spinor_param.location = a.Location();
329  return spinor_param;
330  }
331 
332  // Return the L2 norm squared of the gauge field
333  double norm2(const GaugeField &a) {
335  double nrm2 = blas::norm2(*b);
336  delete b;
337  return nrm2;
338  }
339 
340  // Return the L1 norm of the gauge field
341  double norm1(const GaugeField &a) {
343  double nrm1 = blas::norm1(*b);
344  delete b;
345  return nrm1;
346  }
347 
348  // Scale the gauge field by the constant a
349  void ax(const double &a, GaugeField &u) {
351  blas::ax(a, *b);
352  delete b;
353  }
354 
355  uint64_t GaugeField::checksum(bool mini) const {
356  return Checksum(*this, mini);
357  }
358 
360 
361  GaugeField *field = nullptr;
362  if (param.location == QUDA_CPU_FIELD_LOCATION) {
363  field = new cpuGaugeField(param);
364  } else if (param.location== QUDA_CUDA_FIELD_LOCATION) {
365  field = new cudaGaugeField(param);
366  } else {
367  errorQuda("Invalid field location %d", param.location);
368  }
369 
370  return field;
371  }
372 
373 } // namespace quda
QudaTboundary t_boundary
Definition: gauge_field.h:20
void ax(double a, ColorSpinorField &x)
Definition: blas_quda.cu:508
#define qudaMemcpy(dst, src, count, kind)
Definition: quda_cuda_api.h:33
void setPrecision(QudaPrecision precision, QudaPrecision ghost_precision=QUDA_INVALID_PRECISION, bool force_native=false)
#define pool_pinned_free(ptr)
Definition: malloc_quda.h:128
QudaGaugeFieldOrder FieldOrder() const
Definition: gauge_field.h:257
virtual ~GaugeField()
#define errorQuda(...)
Definition: util_quda.h:121
double norm2(const ColorSpinorField &a)
Definition: reduce_quda.cu:721
QudaReconstructType reconstruct
Definition: gauge_field.h:176
bool staggeredPhaseApplied
Definition: gauge_field.h:201
uint64_t checksum(bool mini=false) const
QudaLinkType LinkType() const
Definition: gauge_field.h:255
static ColorSpinorField * Create(const ColorSpinorParam &param)
uint64_t Checksum(const GaugeField &u, bool mini=false)
Definition: checksum.cu:34
void setTuningString()
Set the vol_string and aux_string for use in tuning.
double anisotropy
Definition: test_util.cpp:1650
QudaFieldGeometry Geometry() const
Definition: gauge_field.h:258
virtual void setTuningString()
void applyGaugePhase(GaugeField &u)
Definition: gauge_phase.cu:223
QudaStaggeredPhase staggeredPhaseType
Definition: gauge_field.h:196
GaugeField(const GaugeFieldParam &param)
Definition: gauge_field.cpp:30
QudaSiteSubset siteSubset
Definition: lattice_field.h:71
std::ostream & operator<<(std::ostream &output, const CloverFieldParam &param)
QudaFieldGeometry geometry
Definition: gauge_field.h:174
QudaGaugeParam param
Definition: pack_test.cpp:17
int Ncolor() const
Definition: gauge_field.h:249
int x[QUDA_MAX_DIM]
Definition: lattice_field.h:67
const int * R() const
enum QudaDirection_s QudaDirection
void checkField(const LatticeField &) const
void ax(const double &a, GaugeField &u)
Scale the gauge field by the scalar a.
QudaFieldLocation location
#define comm_declare_send_relative(buffer, dim, dir, nbytes)
Definition: comm_quda.h:59
void checkField(const LatticeField &a) const
#define comm_declare_receive_relative(buffer, dim, dir, nbytes)
Definition: comm_quda.h:74
QudaGaugeFixed fixed
Definition: gauge_field.h:18
enum QudaStaggeredPhase_s QudaStaggeredPhase
static GaugeField * Create(const GaugeFieldParam &param)
Create the gauge field, with meta data specified in the parameter struct.
const int nColor
Definition: covdev_test.cpp:75
char aux_string[TuneKey::aux_n]
#define ALIGNMENT_ADJUST(n)
Definition: quda_internal.h:22
Generic reconstruction helper with no reconstruction.
QudaFieldLocation location
Definition: gauge_field.h:12
void comm_start(MsgHandle *mh)
Definition: comm_mpi.cpp:216
QudaGaugeFieldOrder order
Definition: gauge_field.h:17
QudaGhostExchange ghostExchange
void exchange(void **recv, void **send, QudaDirection dir) const
Exchange the buffers across all dimensions in a given direction.
void exchangeGhost(QudaLinkDirection link_direction=QUDA_LINK_BACKWARDS)
Exchange the ghost and store store in the padded region.
double norm2(int dim=-1) const
Compute the L2 norm squared of the field.
Definition: max_gauge.cu:71
void comm_free(MsgHandle *&mh)
Definition: comm_mpi.cpp:207
QudaStaggeredPhase staggeredPhaseType
Definition: gauge_field.h:36
size_t ghost_face_bytes[QUDA_MAX_DIM]
void createGhostZone(const int *R, bool no_comms_fill, bool bidir=true) const
double norm1(int dim=-1) const
Compute the L1 norm of the field.
Definition: max_gauge.cu:58
QudaFieldLocation Location() const
int surface[QUDA_MAX_DIM]
#define pool_pinned_malloc(size)
Definition: malloc_quda.h:127
void exchangeGhost(QudaLinkDirection link_direction=QUDA_LINK_BACKWARDS)
Exchange the ghost and store store in the padded region.
int ghostOffset[QUDA_MAX_DIM][2]
double norm1(const ColorSpinorField &b)
Definition: reduce_quda.cu:714
int Ncolor
Definition: blas_test.cu:46
QudaLinkType link_type
Definition: gauge_field.h:180
QudaPrecision ghost_precision
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
static const int aux_n
Definition: tune_key.h:12
ColorSpinorParam colorSpinorParam(const CloverField &a, bool inverse)
QudaReconstructType reconstruct
Definition: gauge_field.h:16
QudaFieldCreate create
Definition: gauge_field.h:26
int surfaceCB[QUDA_MAX_DIM]
QudaReconstructType Reconstruct() const
Definition: gauge_field.h:250
enum QudaFieldGeometry_s QudaFieldGeometry
virtual void * Gauge_p()
Definition: gauge_field.h:315
__device__ __host__ auto StaggeredPhase(const int coords[], int dim, int dir, const Arg &arg) -> typename Arg::real
Compute the staggered phase factor at unit shift from the current lattice coordinates. The routine below optimizes out the shift where possible, hence is only visible where we need to consider the boundary condition.
QudaFieldGeometry geometry
Definition: gauge_field.h:28
void comm_wait(MsgHandle *mh)
Definition: comm_mpi.cpp:222
QudaGaugeFieldOrder order
Definition: gauge_field.h:178
QudaPrecision Precision() const
bool isNative() const
QudaPrecision precision
QudaGaugeFixed fixed
Definition: gauge_field.h:179
float fat_link_max
void applyStaggeredPhase(QudaStaggeredPhase phase=QUDA_STAGGERED_PHASE_INVALID)
unsigned long long bytes
Definition: blas_quda.cu:23
int comm_dim_partitioned(int dim)
const int * X() const
int ghostFace[QUDA_MAX_DIM]
Definition: gauge_field.h:191
void removeStaggeredPhase()