9 location(u.Location()),
14 fixed(u.GaugeFixed()),
15 link_type(u.LinkType()),
16 t_boundary(u.TBoundary()),
21 geometry(u.Geometry()),
22 compute_fat_link_max(false),
24 staggeredPhaseApplied(u.StaggeredPhaseApplied()),
26 site_offset(u.SiteOffset()),
27 site_size(u.SiteSize())
61 errorQuda(
"Anisotropy only supported for Wilson links");
63 errorQuda(
"Temporal gauge fixing only supported for Wilson links");
85 half_phase_bytes = ((half_phase_bytes + (512-1))/512)*512;
86 half_gauge_bytes = ((half_gauge_bytes + (512-1))/512)*512;
90 bytes = (half_gauge_bytes + half_phase_bytes)*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");
124 for (
int i=0; i<
nDim; i++) {
204 send[i] = link_sendbuf[i];
205 receive[i] = ghost_link[i];
207 if (no_comms_fill) memcpy(ghost_link[i], link_sendbuf[i], bytes[i]);
215 qudaMemcpy(send[i], link_sendbuf[i], bytes[i], cudaMemcpyDeviceToHost);
217 if (no_comms_fill)
qudaMemcpy(ghost_link[i], link_sendbuf[i], bytes[i], cudaMemcpyDeviceToDevice);
251 qudaMemcpy(ghost_link[i], receive[i], bytes[i], cudaMemcpyHostToDevice);
277 catch(std::bad_cast &e) {
278 errorQuda(
"Failed to cast reference to GaugeField");
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;
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;
312 errorQuda(
"Casting a GaugeField into ColorSpinorField not possible in half or quarter precision");
316 spinor_param.
nSpin = 1;
318 for (
int d=0; d<a.
Ndim(); d++) spinor_param.
x[d] = a.
X()[d];
320 spinor_param.
pad = a.
Pad();
327 spinor_param.
v = (
void*)a.
Gauge_p();
void ax(double a, ColorSpinorField &x)
#define qudaMemcpy(dst, src, count, kind)
void setPrecision(QudaPrecision precision, QudaPrecision ghost_precision=QUDA_INVALID_PRECISION, bool force_native=false)
#define pool_pinned_free(ptr)
QudaGaugeFieldOrder FieldOrder() const
double norm2(const ColorSpinorField &a)
QudaReconstructType reconstruct
bool staggeredPhaseApplied
uint64_t checksum(bool mini=false) const
QudaLinkType LinkType() const
static ColorSpinorField * Create(const ColorSpinorParam ¶m)
uint64_t Checksum(const GaugeField &u, bool mini=false)
void setTuningString()
Set the vol_string and aux_string for use in tuning.
QudaFieldGeometry Geometry() const
virtual void setTuningString()
void applyGaugePhase(GaugeField &u)
QudaStaggeredPhase staggeredPhaseType
GaugeField(const GaugeFieldParam ¶m)
QudaSiteSubset siteSubset
std::ostream & operator<<(std::ostream &output, const CloverFieldParam ¶m)
QudaFieldGeometry geometry
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)
void checkField(const LatticeField &a) const
#define comm_declare_receive_relative(buffer, dim, dir, nbytes)
enum QudaStaggeredPhase_s QudaStaggeredPhase
static GaugeField * Create(const GaugeFieldParam ¶m)
Create the gauge field, with meta data specified in the parameter struct.
QudaFieldOrder fieldOrder
bool staggeredPhaseApplied
char aux_string[TuneKey::aux_n]
#define ALIGNMENT_ADJUST(n)
Generic reconstruction helper with no reconstruction.
QudaFieldLocation location
QudaGammaBasis gammaBasis
void comm_start(MsgHandle *mh)
QudaGaugeFieldOrder order
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.
void comm_free(MsgHandle *&mh)
QudaStaggeredPhase staggeredPhaseType
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.
QudaFieldLocation Location() const
int surface[QUDA_MAX_DIM]
#define pool_pinned_malloc(size)
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)
QudaPrecision ghost_precision
GaugeFieldParam(void *const h_gauge=NULL)
ColorSpinorParam colorSpinorParam(const CloverField &a, bool inverse)
QudaReconstructType reconstruct
int surfaceCB[QUDA_MAX_DIM]
QudaReconstructType Reconstruct() const
enum QudaFieldGeometry_s QudaFieldGeometry
__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
void comm_wait(MsgHandle *mh)
QudaGaugeFieldOrder order
QudaPrecision Precision() const
void applyStaggeredPhase(QudaStaggeredPhase phase=QUDA_STAGGERED_PHASE_INVALID)
int comm_dim_partitioned(int dim)
int ghostFace[QUDA_MAX_DIM]
void removeStaggeredPhase()