18 ghost( ), ghostNorm( ), ghostFace( ),
19 bytes(0), norm_bytes(0), even(0), odd(0),
20 composite_descr(
param.is_composite,
param.composite_dim,
param.is_component,
param.component_id),
31 ghost( ), ghostNorm( ), ghostFace( ),
32 bytes(0), norm_bytes(0), even(0), odd(0),
33 composite_descr(field.composite_descr), components(0)
49 int nSpinGhost = (
nSpin == 4 && spin_project) ? 2 :
nSpin;
55 int x5 =
nDim == 5 ?
x[4] : 1;
59 for (
int i=0; i<dims; i++) {
63 for (
int j=0; j<dims; j++) {
97 for (
int j=0; j<4; j++) face[j] =
X[j];
145 void ColorSpinorField::create(
int Ndim,
const int *X,
int Nc,
int Ns,
int Nvec,
QudaTwistFlavorType Twistflavor,
171 for (
int d=0; d<
nDim; d++) {
177 errorQuda(
"Must be two flavors for non-degenerate twisted mass spinor (while provided with %d number of components)\n",
x[4]);
244 for (
int d=1; d<
nDim; d++) {
253 char aux_tmp[aux_string_n];
254 int check = snprintf(
aux_string, aux_string_n,
"vol=%lu,stride=%lu,precision=%d,order=%d,Ns=%d,Nc=%d",
volume,
256 if (check < 0 || check >= aux_string_n)
errorQuda(
"Error writing aux string");
260 if (check < 0 || check >= aux_string_n)
errorQuda(
"Error writing aux string");
265 void ColorSpinorField::destroy() {
311 for (
int d=0; d<
nDim; d++) {
318 errorQuda(
"Must be two flavors for non-degenerate twisted mass spinor (provided with %d)\n",
x[4]);
375 if (!
init)
errorQuda(
"Shouldn't be resetting a non-inited field\n");
393 param.is_component =
false;
394 param.component_id = 0;
422 void *total_send =
nullptr;
423 void *total_recv =
nullptr;
430 bool no_comms_fill =
false;
436 bool fine_grained_memcpy =
false;
441 send_back[i] = sendbuf[2*i + 0];
442 send_fwd[i] = sendbuf[2*i + 1];
443 recv_fwd[i] =
ghost[2*i + 1];
444 recv_back[i] =
ghost[2*i + 0];
445 }
else if (no_comms_fill) {
446 memcpy(
ghost[2*i+1], sendbuf[2*i+0],
bytes[i]);
447 memcpy(
ghost[2*i+0], sendbuf[2*i+1],
bytes[i]);
458 send_back[i] =
static_cast<char*
>(total_send) + offset;
459 recv_back[i] =
static_cast<char*
>(total_recv) + offset;
461 send_fwd[i] =
static_cast<char*
>(total_send) + offset;
462 recv_fwd[i] =
static_cast<char*
>(total_recv) + offset;
464 if (fine_grained_memcpy) {
465 qudaMemcpy(send_back[i], sendbuf[2*i + 0],
bytes[i], cudaMemcpyDeviceToHost);
466 qudaMemcpy(send_fwd[i], sendbuf[2*i + 1],
bytes[i], cudaMemcpyDeviceToHost);
468 }
else if (no_comms_fill) {
475 void *send_ptr =
nullptr;
478 send_ptr = sendbuf[2*i];
514 if (fine_grained_memcpy) {
522 void *ghost_ptr =
nullptr;
525 ghost_ptr =
ghost[2*i];
574 errorQuda(
"Cannot return even subset of QDPJIT field");
582 errorQuda(
"Cannot return even subset of QDPJIT field");
590 errorQuda(
"Cannot return even subset of QDPJIT field");
598 errorQuda(
"Cannot return even subset of QDPJIT field");
608 errorQuda(
"Incorrect component index...");
611 errorQuda(
"Cannot get requested component");
621 errorQuda(
"Incorrect component index...");
624 errorQuda(
"Cannot get requested component");
668 for (
int d=0; d<
nDim; d++) {
678 for (
int d=1; d<
nDim; d++) oddBit += y[d];
681 y[0] = 2*y[0] + oddBit;
704 for (
int d=
nDim-1; d>=0; d--) {
745 errorQuda(
"Cannot create an alias to source with lower precision than the alias");
758 param.norm =
Norm() ?
Norm() :
static_cast<char *
>(
V()) + norm_offset;
763 if (alias->Bytes() >
Bytes())
errorQuda(
"Alias footprint %lu greater than source %lu", alias->Bytes(),
Bytes());
766 if (
static_cast<char *
>(alias->V()) + alias->Bytes() > alias->Norm())
767 errorQuda(
"Overlap between alias body and norm");
769 if (
static_cast<char *
>(alias->Norm()) + alias->NormBytes() >
static_cast<char *
>(
V()) +
Bytes())
770 errorQuda(
"Norm is not contained in the srouce field");
780 for (
int d=0; d<
nDim; d++) coarseParam.
x[d] =
x[d]/geoBlockSize[d];
782 int geoBlockVolume = 1;
783 for (
int d = 0; d <
nDim; d++) { geoBlockVolume *= geoBlockSize[d]; }
820 errorQuda(
"Invalid field location %d", new_location);
830 for (
int d=0; d<
nDim; d++) fineParam.
x[d] =
x[d] * geoBlockSize[d];
844 fineParam.
setPrecision(new_precision, new_precision,
true);
860 errorQuda(
"Invalid field location %d", new_location);
866 out <<
"typedid = " <<
typeid(a).name() << std::endl;
867 out <<
"nColor = " << a.
nColor << std::endl;
868 out <<
"nSpin = " << a.
nSpin << std::endl;
869 out <<
"twistFlavor = " << a.
twistFlavor << std::endl;
870 out <<
"nDim = " << a.
nDim << std::endl;
871 for (
int d=0; d<a.
nDim; d++) out <<
"x[" << d <<
"] = " << a.
x[d] << std::endl;
872 out <<
"volume = " << a.
volume << std::endl;
873 out <<
"pc_type = " << a.
pc_type << std::endl;
875 out <<
"precision = " << a.
precision << std::endl;
877 out <<
"pad = " << a.
pad << std::endl;
878 out <<
"stride = " << a.
stride << std::endl;
879 out <<
"real_length = " << a.
real_length << std::endl;
880 out <<
"length = " << a.
length << std::endl;
881 out <<
"bytes = " << a.
bytes << std::endl;
882 out <<
"norm_bytes = " << a.
norm_bytes << std::endl;
883 out <<
"siteSubset = " << a.
siteSubset << std::endl;
884 out <<
"siteOrder = " << a.
siteOrder << std::endl;
885 out <<
"fieldOrder = " << a.
fieldOrder << std::endl;
886 out <<
"gammaBasis = " << a.
gammaBasis << std::endl;
897 out <<
"pc_type = " << a.
pc_type << std::endl;
void exchange(void **ghost, void **sendbuf, int nFace=1) const
void fill(ColorSpinorParam &) const
void LatticeIndex(int *y, int i) const
int ghostFace[QUDA_MAX_DIM]
void *const * Ghost() const
ColorSpinorField * CreateAlias(const ColorSpinorParam ¶m)
Create a field that aliases this field's storage. The alias field can use a different precision than ...
virtual ColorSpinorField & operator=(const ColorSpinorField &)
void * ghost_buf[2 *QUDA_MAX_DIM]
QudaParity suggested_parity
const ColorSpinorField & Odd() const
CompositeColorSpinorFieldDescriptor composite_descr
used for deflation eigenvector sets etc.:
ColorSpinorField(const ColorSpinorField &)
QudaTwistFlavorType TwistFlavor() const
QudaTwistFlavorType twistFlavor
ColorSpinorField * CreateFine(const int *geoblockSize, int spinBlockSize, int Nvec, QudaPrecision precision=QUDA_INVALID_PRECISION, QudaFieldLocation location=QUDA_INVALID_FIELD_LOCATION, QudaMemoryType mem_type=QUDA_MEMORY_INVALID)
Create a fine color-spinor field, using this field to set the meta data.
CompositeColorSpinorField components
QudaGammaBasis gammaBasis
void * ghost[2][QUDA_MAX_DIM]
static ColorSpinorField * Create(const ColorSpinorParam ¶m)
void reset(const ColorSpinorParam &)
void createGhostZone(int nFace, bool spin_project=true) const
static void checkField(const ColorSpinorField &, const ColorSpinorField &)
QudaPrecision ghost_precision_allocated
ColorSpinorField * CreateCoarse(const int *geoBlockSize, int spinBlockSize, int Nvec, QudaPrecision precision=QUDA_INVALID_PRECISION, QudaFieldLocation location=QUDA_INVALID_FIELD_LOCATION, QudaMemoryType mem_Type=QUDA_MEMORY_INVALID)
Create a coarse color-spinor field, using this field to set the meta data.
QudaFieldOrder fieldOrder
void OffsetIndex(int &i, int *y) const
void * GhostNorm(const int i)
const ColorSpinorField & Even() const
QudaSiteSubset siteSubset
ColorSpinorField & Component(const int idx) const
int ghostFaceCB[QUDA_MAX_DIM]
void * ghostNorm[2][QUDA_MAX_DIM]
void setTuningString()
Set the vol_string and aux_string for use in tuning.
virtual ~ColorSpinorField()
DslashConstant dslash_constant
void setPrecision(QudaPrecision precision, QudaPrecision ghost_precision=QUDA_INVALID_PRECISION, bool force_native=false)
QudaFieldOrder fieldOrder
MsgHandle * mh_send_fwd[2][QUDA_MAX_DIM]
size_t ghost_offset[QUDA_MAX_DIM][2]
QudaPrecision ghost_precision
QudaPrecision Precision() const
QudaFieldLocation Location() const
size_t ghost_face_bytes[QUDA_MAX_DIM]
char aux_string[TuneKey::aux_n]
char vol_string[TuneKey::volume_n]
int surfaceCB[QUDA_MAX_DIM]
MsgHandle * mh_send_back[2][QUDA_MAX_DIM]
size_t ghost_face_bytes_aligned[QUDA_MAX_DIM]
void comm_start(MsgHandle *mh)
int comm_dim_partitioned(int dim)
#define comm_declare_receive_relative(buffer, dim, dir, nbytes)
void comm_wait(MsgHandle *mh)
void comm_free(MsgHandle *&mh)
#define comm_declare_send_relative(buffer, dim, dir, nbytes)
enum QudaSiteOrder_s QudaSiteOrder
enum QudaPrecision_s QudaPrecision
@ QUDA_CUDA_FIELD_LOCATION
@ QUDA_CPU_FIELD_LOCATION
@ QUDA_INVALID_FIELD_LOCATION
enum QudaTwistFlavorType_s QudaTwistFlavorType
@ QUDA_INVALID_SITE_SUBSET
@ QUDA_PARITY_SITE_SUBSET
@ QUDA_INVALID_GAMMA_BASIS
enum QudaPCType_s QudaPCType
enum QudaFieldOrder_s QudaFieldOrder
enum QudaSiteSubset_s QudaSiteSubset
enum QudaFieldLocation_s QudaFieldLocation
@ QUDA_INVALID_SITE_ORDER
enum QudaMemoryType_s QudaMemoryType
@ QUDA_INVALID_FIELD_ORDER
@ QUDA_FLOAT2_FIELD_ORDER
@ QUDA_QDPJIT_FIELD_ORDER
@ QUDA_SPACE_SPIN_COLOR_FIELD_ORDER
enum QudaGammaBasis_s QudaGammaBasis
@ QUDA_INVALID_FIELD_CREATE
@ QUDA_REFERENCE_FIELD_CREATE
@ QUDA_TWIST_NONDEG_DOUBLET
enum QudaParity_s QudaParity
#define pool_pinned_malloc(size)
#define pool_pinned_free(ptr)
void init()
Create the BLAS context.
__host__ __device__ ValueType norm(const complex< ValueType > &z)
Returns the magnitude of z squared.
std::ostream & operator<<(std::ostream &output, const CloverFieldParam ¶m)
#define qudaMemcpy(dst, src, count, kind)
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5.
#define ALIGNMENT_ADJUST(n)
QudaFieldLocation location
int_fastdiv Xh[QUDA_MAX_DIM]
int ghostFace[QUDA_MAX_DIM+1]
int ghostFaceCB[QUDA_MAX_DIM+1]
int_fastdiv X[QUDA_MAX_DIM]
QudaSiteSubset siteSubset
QudaPrecision Precision() const
static const int volume_n