QUDA  0.9.0
copy_gauge_inc.cu
Go to the documentation of this file.
1 #include <gauge_field_order.h>
2 #include <copy_gauge_helper.cuh>
3 
4 namespace quda {
5 
6  template <typename FloatOut, typename FloatIn, int length, typename InOrder>
7  void copyGauge(const InOrder &inOrder, const GaugeField &out, const GaugeField &in,
8  QudaFieldLocation location, FloatOut *Out, FloatOut **outGhost, int type) {
9  int faceVolumeCB[QUDA_MAX_DIM];
10  for (int i=0; i<4; i++) faceVolumeCB[i] = out.SurfaceCB(i) * out.Nface();
11  if (out.isNative()) {
12  // this overrides the check that the texture maps to the gauge
13  // pointer - this is safe here since it only occurs when running
14  // the copier on the host when we will not be using texture
15  // reads
16  const bool override = true;
17  if (out.Reconstruct() == QUDA_RECONSTRUCT_NO) {
18  if (typeid(FloatOut)==typeid(short) && out.LinkType() == QUDA_ASQTAD_FAT_LINKS) {
19  copyGauge<short,FloatIn,length>
20  (FloatNOrder<short,length,2,19>(out, (short*)Out, (short**)outGhost, override), inOrder,
21  out.Volume(), faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
22  } else {
23  typedef typename gauge_mapper<FloatOut,QUDA_RECONSTRUCT_NO>::type G;
24  copyGauge<FloatOut,FloatIn,length>
25  (G(out, Out, outGhost, override), inOrder, out.Volume(), faceVolumeCB,
26  out.Ndim(), out.Geometry(), out, in, location, type);
27  }
28  } else if (out.Reconstruct() == QUDA_RECONSTRUCT_12) {
29  typedef typename gauge_mapper<FloatOut,QUDA_RECONSTRUCT_12>::type G;
30  copyGauge<FloatOut,FloatIn,length>
31  (G(out, Out, outGhost, override), inOrder, out.Volume(), faceVolumeCB,
32  out.Ndim(), out.Geometry(), out, in, location, type);
33  } else if (out.Reconstruct() == QUDA_RECONSTRUCT_8) {
34  typedef typename gauge_mapper<FloatOut,QUDA_RECONSTRUCT_8>::type G;
35  copyGauge<FloatOut,FloatIn,length>
36  (G(out, Out, outGhost, override), inOrder, out.Volume(), faceVolumeCB,
37  out.Ndim(), out.Geometry(), out, in, location, type);
38 #ifdef GPU_STAGGERED_DIRAC
39  } else if (out.Reconstruct() == QUDA_RECONSTRUCT_13) {
40  typedef typename gauge_mapper<FloatOut,QUDA_RECONSTRUCT_13>::type G;
41  copyGauge<FloatOut,FloatIn,length>
42  (G(out, Out, outGhost, override), inOrder, out.Volume(), faceVolumeCB,
43  out.Ndim(), out.Geometry(), out, in, location, type);
44  } else if (out.Reconstruct() == QUDA_RECONSTRUCT_9) {
45  typedef typename gauge_mapper<FloatOut,QUDA_RECONSTRUCT_9>::type G;
46  copyGauge<FloatOut,FloatIn,length>
47  (G(out, Out, outGhost, override), inOrder, out.Volume(), faceVolumeCB,
48  out.Ndim(), out.Geometry(), out, in, location, type);
49 #endif
50  } else {
51  errorQuda("Reconstruction %d and order %d not supported", out.Reconstruct(), out.Order());
52  }
53  } else if (out.Order() == QUDA_QDP_GAUGE_ORDER) {
54 
55 #ifdef BUILD_QDP_INTERFACE
56  copyGauge<FloatOut,FloatIn,length>
57  (QDPOrder<FloatOut,length>(out, Out, outGhost), inOrder, out.Volume(),
58  faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
59 #else
60  errorQuda("QDP interface has not been built\n");
61 #endif
62 
63  } else if (out.Order() == QUDA_QDPJIT_GAUGE_ORDER) {
64 
65 #ifdef BUILD_QDPJIT_INTERFACE
66  copyGauge<FloatOut,FloatIn,length>
67  (QDPJITOrder<FloatOut,length>(out, Out, outGhost), inOrder, out.Volume(),
68  faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
69 #else
70  errorQuda("QDPJIT interface has not been built\n");
71 #endif
72 
73  } else if (out.Order() == QUDA_CPS_WILSON_GAUGE_ORDER) {
74 
75 #ifdef BUILD_CPS_INTERFACE
76  copyGauge<FloatOut,FloatIn,length>
77  (CPSOrder<FloatOut,length>(out, Out, outGhost), inOrder, out.Volume(),
78  faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
79 #else
80  errorQuda("CPS interface has not been built\n");
81 #endif
82 
83  } else if (out.Order() == QUDA_MILC_GAUGE_ORDER) {
84 
85 #ifdef BUILD_MILC_INTERFACE
86  copyGauge<FloatOut,FloatIn,length>
87  (MILCOrder<FloatOut,length>(out, Out, outGhost), inOrder, out.Volume(),
88  faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
89 #else
90  errorQuda("MILC interface has not been built\n");
91 #endif
92 
93  } else if (out.Order() == QUDA_MILC_SITE_GAUGE_ORDER) {
94 
95 #ifdef BUILD_MILC_INTERFACE
96  copyGauge<FloatOut,FloatIn,length>
97  (MILCSiteOrder<FloatOut,length>(out, Out, outGhost), inOrder, out.Volume(),
98  faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
99 #else
100  errorQuda("MILC interface has not been built\n");
101 #endif
102 
103  } else if (out.Order() == QUDA_BQCD_GAUGE_ORDER) {
104 
105 #ifdef BUILD_BQCD_INTERFACE
106  copyGauge<FloatOut,FloatIn,length>
107  (BQCDOrder<FloatOut,length>(out, Out, outGhost), inOrder, out.Volume(),
108  faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
109 #else
110  errorQuda("BQCD interface has not been built\n");
111 #endif
112 
113  } else if (out.Order() == QUDA_TIFR_GAUGE_ORDER) {
114 
115 #ifdef BUILD_TIFR_INTERFACE
116  copyGauge<FloatOut,FloatIn,length>
117  (TIFROrder<FloatOut,length>(out, Out, outGhost), inOrder, out.Volume(),
118  faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
119 #else
120  errorQuda("TIFR interface has not been built\n");
121 #endif
122 
123  } else if (out.Order() == QUDA_TIFR_PADDED_GAUGE_ORDER) {
124 
125 #ifdef BUILD_TIFR_INTERFACE
126  copyGauge<FloatOut,FloatIn,length>
127  (TIFRPaddedOrder<FloatOut,length>(out, Out, outGhost), inOrder, out.Volume(),
128  faceVolumeCB, out.Ndim(), out.Geometry(), out, in, location, type);
129 #else
130  errorQuda("TIFR interface has not been built\n");
131 #endif
132 
133  } else {
134  errorQuda("Gauge field %d order not supported", out.Order());
135  }
136 
137  }
138 
139  template <typename FloatOut, typename FloatIn, int length>
140  void copyGauge(GaugeField &out, const GaugeField &in, QudaFieldLocation location,
141  FloatOut *Out, FloatIn *In, FloatOut **outGhost, FloatIn **inGhost, int type) {
142 
143  // reconstruction only supported on FloatN fields currently
144  if (in.isNative()) {
145  // this overrides the check that the texture maps to the gauge
146  // pointer - this is safe here since it only occurs when running
147  // the copier on the host when we will not be using texture
148  // reads
149  const bool override = true;
150  if (in.Reconstruct() == QUDA_RECONSTRUCT_NO) {
151  if (typeid(FloatIn)==typeid(short) && in.LinkType() == QUDA_ASQTAD_FAT_LINKS) {
152  copyGauge<FloatOut,short,length> (FloatNOrder<short,length,2,19>
153  (in,(short*)In,(short**)inGhost,override),
154  out, in, location, Out, outGhost, type);
155  } else {
156  typedef typename gauge_mapper<FloatIn,QUDA_RECONSTRUCT_NO>::type G;
157  copyGauge<FloatOut,FloatIn,length> (G(in,In,inGhost,override), out, in, location, Out, outGhost, type);
158  }
159  } else if (in.Reconstruct() == QUDA_RECONSTRUCT_12) {
160  typedef typename gauge_mapper<FloatIn,QUDA_RECONSTRUCT_12>::type G;
161  copyGauge<FloatOut,FloatIn,length> (G(in,In,inGhost,override), out, in, location, Out, outGhost, type);
162  } else if (in.Reconstruct() == QUDA_RECONSTRUCT_8) {
163  typedef typename gauge_mapper<FloatIn,QUDA_RECONSTRUCT_8>::type G;
164  copyGauge<FloatOut,FloatIn,length> (G(in,In,inGhost,override), out, in, location, Out, outGhost, type);
165 #ifdef GPU_STAGGERED_DIRAC
166  } else if (in.Reconstruct() == QUDA_RECONSTRUCT_13) {
167  typedef typename gauge_mapper<FloatIn,QUDA_RECONSTRUCT_13>::type G;
168  copyGauge<FloatOut,FloatIn,length> (G(in,In,inGhost,override), out, in, location, Out, outGhost, type);
169  } else if (in.Reconstruct() == QUDA_RECONSTRUCT_9) {
170  typedef typename gauge_mapper<FloatIn,QUDA_RECONSTRUCT_9>::type G;
171  copyGauge<FloatOut,FloatIn,length> (G(in,In,inGhost,override), out, in, location, Out, outGhost, type);
172 #endif
173  } else {
174  errorQuda("Reconstruction %d and order %d not supported", in.Reconstruct(), in.Order());
175  }
176  } else if (in.Order() == QUDA_QDP_GAUGE_ORDER) {
177 
178 #ifdef BUILD_QDP_INTERFACE
179  copyGauge<FloatOut,FloatIn,length>(QDPOrder<FloatIn,length>(in, In, inGhost),
180  out, in, location, Out, outGhost, type);
181 #else
182  errorQuda("QDP interface has not been built\n");
183 #endif
184 
185  } else if (in.Order() == QUDA_QDPJIT_GAUGE_ORDER) {
186 
187 #ifdef BUILD_QDPJIT_INTERFACE
188  copyGauge<FloatOut,FloatIn,length>(QDPJITOrder<FloatIn,length>(in, In, inGhost),
189  out, in, location, Out, outGhost, type);
190 #else
191  errorQuda("QDPJIT interface has not been built\n");
192 #endif
193 
194  } else if (in.Order() == QUDA_CPS_WILSON_GAUGE_ORDER) {
195 
196 #ifdef BUILD_CPS_INTERFACE
197  copyGauge<FloatOut,FloatIn,length>(CPSOrder<FloatIn,length>(in, In, inGhost),
198  out, in, location, Out, outGhost, type);
199 #else
200  errorQuda("CPS interface has not been built\n");
201 #endif
202 
203  } else if (in.Order() == QUDA_MILC_GAUGE_ORDER) {
204 
205 #ifdef BUILD_MILC_INTERFACE
206  copyGauge<FloatOut,FloatIn,length>(MILCOrder<FloatIn,length>(in, In, inGhost),
207  out, in, location, Out, outGhost, type);
208 #else
209  errorQuda("MILC interface has not been built\n");
210 #endif
211 
212  } else if (in.Order() == QUDA_MILC_SITE_GAUGE_ORDER) {
213 
214 #ifdef BUILD_MILC_INTERFACE
215  copyGauge<FloatOut,FloatIn,length>(MILCSiteOrder<FloatIn,length>(in, In, inGhost),
216  out, in, location, Out, outGhost, type);
217 #else
218  errorQuda("MILC interface has not been built\n");
219 #endif
220 
221  } else if (in.Order() == QUDA_BQCD_GAUGE_ORDER) {
222 
223 #ifdef BUILD_BQCD_INTERFACE
224  copyGauge<FloatOut,FloatIn,length>(BQCDOrder<FloatIn,length>(in, In, inGhost),
225  out, in, location, Out, outGhost, type);
226 #else
227  errorQuda("BQCD interface has not been built\n");
228 #endif
229 
230  } else if (in.Order() == QUDA_TIFR_GAUGE_ORDER) {
231 
232 #ifdef BUILD_TIFR_INTERFACE
233  copyGauge<FloatOut,FloatIn,length>(TIFROrder<FloatIn,length>(in, In, inGhost),
234  out, in, location, Out, outGhost, type);
235 #else
236  errorQuda("TIFR interface has not been built\n");
237 #endif
238 
239  } else if (in.Order() == QUDA_TIFR_PADDED_GAUGE_ORDER) {
240 
241 #ifdef BUILD_TIFR_INTERFACE
242  copyGauge<FloatOut,FloatIn,length>(TIFRPaddedOrder<FloatIn,length>(in, In, inGhost),
243  out, in, location, Out, outGhost, type);
244 #else
245  errorQuda("TIFR interface has not been built\n");
246 #endif
247 
248  } else {
249  errorQuda("Gauge field order %d not supported", in.Order());
250  }
251 
252  }
253 
254  void checkMomOrder(const GaugeField &u);
255 
256  template <typename FloatOut, typename FloatIn, int length, typename Out, typename In, typename Arg>
257  void copyMom(Arg &arg, const GaugeField &out, const GaugeField &in, QudaFieldLocation location) {
258 
259  if (location == QUDA_CPU_FIELD_LOCATION) {
260  copyGauge<FloatOut,FloatIn,length>(arg);
261  } else if (location == QUDA_CUDA_FIELD_LOCATION) {
262  CopyGauge<FloatOut,FloatIn,length, Out, In, 0> momCopier(arg, out, in);
263  momCopier.apply(0);
264  } else {
265  errorQuda("Undefined field location %d for copyMom", location);
266  }
267 
268  }
269 
270  template <typename FloatOut, typename FloatIn>
271  void copyGauge(GaugeField &out, const GaugeField &in, QudaFieldLocation location, FloatOut *Out,
272  FloatIn *In, FloatOut **outGhost, FloatIn **inGhost, int type) {
273 
274  if (in.Ncolor() != 3 && out.Ncolor() != 3) {
275  errorQuda("Unsupported number of colors; out.Nc=%d, in.Nc=%d", out.Ncolor(), in.Ncolor());
276  }
277 
278  if (out.Geometry() != in.Geometry()) {
279  errorQuda("Field geometries %d %d do not match", out.Geometry(), in.Geometry());
280  }
281 
282  if (in.LinkType() != QUDA_ASQTAD_MOM_LINKS && out.LinkType() != QUDA_ASQTAD_MOM_LINKS) {
283  // we are doing gauge field packing
284  copyGauge<FloatOut,FloatIn,18>(out, in, location, Out, In, outGhost, inGhost, type);
285  } else {
286  if (out.Geometry() != QUDA_VECTOR_GEOMETRY) errorQuda("Unsupported geometry %d", out.Geometry());
287 
288  checkMomOrder(in);
289  checkMomOrder(out);
290 
291  int faceVolumeCB[QUDA_MAX_DIM];
292  for (int d=0; d<in.Ndim(); d++) faceVolumeCB[d] = in.SurfaceCB(d) * in.Nface();
293 
294  // this overrides the check that the texture maps to the gauge
295  // pointer - this is safe here since it only occurs when running
296  // the copier on the host when we will not be using texture
297  // reads
298  const bool override = true;
299 
300  // momentum only currently supported on MILC (10), TIFR (18) and Float2 (10) fields currently
301  if (out.Order() == QUDA_FLOAT2_GAUGE_ORDER) {
302  if (in.Order() == QUDA_FLOAT2_GAUGE_ORDER) {
303  typedef FloatNOrder<FloatOut,10,2,10> momOut;
304  typedef FloatNOrder<FloatIn,10,2,10> momIn;
305  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out, 0, override), momIn(in, In, 0, override), in.Volume(),
306  faceVolumeCB, in.Ndim(), in.Geometry());
307  copyMom<FloatOut,FloatIn,10,momOut,momIn>(arg,out,in,location);
308  } else if (in.Order() == QUDA_MILC_GAUGE_ORDER) {
309 #ifdef BUILD_MILC_INTERFACE
310  typedef FloatNOrder<FloatOut,10,2,10> momOut;
311  typedef MILCOrder<FloatIn,10> momIn;
312  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out, 0, override), momIn(in, In), in.Volume(),
313  faceVolumeCB, in.Ndim(), in.Geometry());
314  copyMom<FloatOut,FloatIn,10,momOut,momIn>(arg,out,in,location);
315 #else
316  errorQuda("MILC interface has not been built\n");
317 #endif
318  } else if (in.Order() == QUDA_MILC_SITE_GAUGE_ORDER) {
319 #ifdef BUILD_MILC_INTERFACE
320  typedef FloatNOrder<FloatOut,10,2,10> momOut;
321  typedef MILCSiteOrder<FloatIn,10> momIn;
322  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out, 0, override), momIn(in, In), in.Volume(),
323  faceVolumeCB, in.Ndim(), in.Geometry());
324  copyMom<FloatOut,FloatIn,10,momOut,momIn>(arg,out,in,location);
325 #else
326  errorQuda("MILC interface has not been built\n");
327 #endif
328  } else if (in.Order() == QUDA_TIFR_GAUGE_ORDER) {
329 #ifdef BUILD_TIFR_INTERFACE
330  typedef FloatNOrder<FloatOut,18,2,11> momOut;
331  typedef TIFROrder<FloatIn,18> momIn;
332  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out, 0, override), momIn(in, In), in.Volume(),
333  faceVolumeCB, in.Ndim(), in.Geometry());
334  copyMom<FloatOut,FloatIn,18,momOut,momIn>(arg,out,in,location);
335 #else
336  errorQuda("TIFR interface has not been built\n");
337 #endif
338  } else if (in.Order() == QUDA_TIFR_PADDED_GAUGE_ORDER) {
339 #ifdef BUILD_TIFR_INTERFACE
340  typedef FloatNOrder<FloatOut,18,2,11> momOut;
341  typedef TIFRPaddedOrder<FloatIn,18> momIn;
342  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out, 0, override), momIn(in, In), in.Volume(),
343  faceVolumeCB, in.Ndim(), in.Geometry());
344  copyMom<FloatOut,FloatIn,18,momOut,momIn>(arg,out,in,location);
345 #else
346  errorQuda("TIFR interface has not been built\n");
347 #endif
348  } else {
349  errorQuda("Gauge field orders %d not supported", in.Order());
350  }
351  } else if (out.Order() == QUDA_MILC_GAUGE_ORDER) {
352  typedef MILCOrder<FloatOut,10> momOut;
353 #ifdef BUILD_MILC_INTERFACE
354  if (in.Order() == QUDA_FLOAT2_GAUGE_ORDER) {
355  typedef FloatNOrder<FloatIn,10,2,10> momIn;
356  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out), momIn(in, In, 0, override), in.Volume(),
357  faceVolumeCB, in.Ndim(), in.Geometry());
358  copyMom<FloatOut,FloatIn,10,momOut,momIn>(arg,out,in,location);
359  } else if (in.Order() == QUDA_MILC_GAUGE_ORDER) {
360  typedef MILCOrder<FloatIn,10> momIn;
361  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out), momIn(in, In), in.Volume(),
362  faceVolumeCB, in.Ndim(), in.Geometry());
363  copyMom<FloatOut,FloatIn,10,momOut,momIn>(arg,out,in,location);
364  } else {
365  errorQuda("Gauge field orders %d not supported", in.Order());
366  }
367 #else
368  errorQuda("MILC interface has not been built\n");
369 #endif
370  } else if (out.Order() == QUDA_MILC_SITE_GAUGE_ORDER) {
371  typedef MILCSiteOrder<FloatOut,10> momOut;
372 #ifdef BUILD_MILC_INTERFACE
373  if (in.Order() == QUDA_FLOAT2_GAUGE_ORDER) {
374  typedef FloatNOrder<FloatIn,10,2,10> momIn;
375  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out), momIn(in, In, 0, override), in.Volume(),
376  faceVolumeCB, in.Ndim(), in.Geometry());
377  copyMom<FloatOut,FloatIn,10,momOut,momIn>(arg,out,in,location);
378  } else if (in.Order() == QUDA_MILC_SITE_GAUGE_ORDER) {
379  typedef MILCSiteOrder<FloatIn,10> momIn;
380  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out), momIn(in, In), in.Volume(),
381  faceVolumeCB, in.Ndim(), in.Geometry());
382  copyMom<FloatOut,FloatIn,10,momOut,momIn>(arg,out,in,location);
383  } else {
384  errorQuda("Gauge field orders %d not supported", in.Order());
385  }
386 #else
387  errorQuda("MILC interface has not been built\n");
388 #endif
389  } else if (out.Order() == QUDA_TIFR_GAUGE_ORDER) {
390  typedef TIFROrder<FloatOut,18> momOut;
391 #ifdef BUILD_TIFR_INTERFACE
392  if (in.Order() == QUDA_FLOAT2_GAUGE_ORDER) {
393  // FIX ME - 11 is a misnomer to avoid confusion in template instantiation
394  typedef FloatNOrder<FloatIn,18,2,11> momIn;
395  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out), momIn(in, In, 0, override), in.Volume(),
396  faceVolumeCB, in.Ndim(), in.Geometry());
397  copyMom<FloatOut,FloatIn,18,momOut,momIn>(arg,out,in,location);
398  } else if (in.Order() == QUDA_TIFR_GAUGE_ORDER) {
399  typedef TIFROrder<FloatIn,18> momIn;
400  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out), momIn(in, In), in.Volume(),
401  faceVolumeCB, in.Ndim(), in.Geometry());
402  copyMom<FloatOut,FloatIn,18,momOut,momIn>(arg,out,in,location);
403  } else {
404  errorQuda("Gauge field orders %d not supported", in.Order());
405  }
406 #else
407  errorQuda("TIFR interface has not been built\n");
408 #endif
409  } else if (out.Order() == QUDA_TIFR_PADDED_GAUGE_ORDER) {
410  typedef TIFRPaddedOrder<FloatOut,18> momOut;
411 #ifdef BUILD_TIFR_INTERFACE
412  if (in.Order() == QUDA_FLOAT2_GAUGE_ORDER) {
413  // FIX ME - 11 is a misnomer to avoid confusion in template instantiation
414  typedef FloatNOrder<FloatIn,18,2,11> momIn;
415  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out), momIn(in, In, 0, override), in.Volume(),
416  faceVolumeCB, in.Ndim(), in.Geometry());
417  copyMom<FloatOut,FloatIn,18,momOut,momIn>(arg,out,in,location);
418  } else if (in.Order() == QUDA_TIFR_PADDED_GAUGE_ORDER) {
419  typedef TIFRPaddedOrder<FloatIn,18> momIn;
420  CopyGaugeArg<momOut,momIn> arg(momOut(out, Out), momIn(in, In), in.Volume(),
421  faceVolumeCB, in.Ndim(), in.Geometry());
422  copyMom<FloatOut,FloatIn,18,momOut,momIn>(arg,out,in,location);
423  } else {
424  errorQuda("Gauge field orders %d not supported", in.Order());
425  }
426 #else
427  errorQuda("TIFR interface has not been built\n");
428 #endif
429  } else {
430  errorQuda("Gauge field orders %d not supported", out.Order());
431  }
432  }
433  }
434 
435 
436 } // namespace quda
QUDA_RECONSTRUCT_NO
Definition: enum_quda.h:66
quda::CopyGauge::apply
void apply(const cudaStream_t &stream)
Definition: copy_gauge_helper.cuh:222
quda::copyGauge
void copyGauge(CopyGaugeArg< OutOrder, InOrder > arg)
Definition: copy_gauge_helper.cuh:32
QUDA_ASQTAD_MOM_LINKS
Definition: enum_quda.h:32
errorQuda
#define errorQuda(...)
Definition: util_quda.h:90
QUDA_BQCD_GAUGE_ORDER
Definition: enum_quda.h:46
QUDA_CUDA_FIELD_LOCATION
Definition: enum_quda.h:297
QUDA_QDP_GAUGE_ORDER
Definition: enum_quda.h:41
QUDA_FLOAT2_GAUGE_ORDER
Definition: enum_quda.h:39
quda::LatticeField::SurfaceCB
const int * SurfaceCB() const
Definition: lattice_field.h:431
QUDA_TIFR_PADDED_GAUGE_ORDER
Definition: enum_quda.h:48
QUDA_MILC_SITE_GAUGE_ORDER
Definition: enum_quda.h:45
copy_gauge_helper.cuh
quda::gauge::MILCSiteOrder
struct to define gauge fields packed into an opaque MILC site struct:
Definition: gauge_field_order.h:1597
QUDA_RECONSTRUCT_9
Definition: enum_quda.h:69
QUDA_RECONSTRUCT_12
Definition: enum_quda.h:67
QUDA_MILC_GAUGE_ORDER
Definition: enum_quda.h:44
in
cpuColorSpinorField * in
Definition: staggered_invert_test.cpp:44
fused_exterior_ndeg_tm_dslash_cuda_gen.i
int i
start here
Definition: fused_exterior_ndeg_tm_dslash_cuda_gen.py:816
QUDA_CPS_WILSON_GAUGE_ORDER
Definition: enum_quda.h:43
gauge_field_order.h
Main header file for host and device accessors to GaugeFields.
QUDA_VECTOR_GEOMETRY
Definition: enum_quda.h:427
QUDA_TIFR_GAUGE_ORDER
Definition: enum_quda.h:47
QUDA_RECONSTRUCT_8
Definition: enum_quda.h:68
QudaFieldLocation
enum QudaFieldLocation_s QudaFieldLocation
out
cpuColorSpinorField * out
Definition: staggered_invert_test.cpp:45
quda::gauge_mapper
Definition: gauge_field_order.h:2083
quda::copyMom
void copyMom(Arg &arg, const GaugeField &out, const GaugeField &in, QudaFieldLocation location)
Definition: copy_gauge_inc.cu:257
QUDA_QDPJIT_GAUGE_ORDER
Definition: enum_quda.h:42
quda::arg
__host__ __device__ ValueType arg(const complex< ValueType > &z)
Returns the phase angle of z.
Definition: complex_quda.h:880
QUDA_RECONSTRUCT_13
Definition: enum_quda.h:70
quda::clover::FloatNOrder
Accessor routine for CloverFields in native field order.
Definition: clover_field_order.h:367
QUDA_MAX_DIM
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5...
Definition: quda_constants.h:17
quda::checkMomOrder
void checkMomOrder(const GaugeField &u)
Definition: copy_gauge.cu:19
d
static __inline__ size_t size_t d
Definition: CMakeCUDACompilerId.cpp1.ii:3019
QUDA_CPU_FIELD_LOCATION
Definition: enum_quda.h:296
QUDA_ASQTAD_FAT_LINKS
Definition: enum_quda.h:30
Generated by   doxygen 1.8.14