1 #ifndef _COLOR_SPINOR_FIELD_H 2 #define _COLOR_SPINOR_FIELD_H 47 : is_composite(false), is_component(false), dim(0), id(0), volume(0), volumeCB(0),
48 stride(0), real_length(0), length(0), bytes(0), norm_bytes(0) {};
51 : is_composite(is_composite), is_component(is_component), dim(dim), id(id), volume(0), volumeCB(0),
52 stride(0), real_length(0), length(0), bytes(0), norm_bytes(0)
54 if(is_composite && is_component)
errorQuda(
"\nComposite type is not implemented.\n");
55 else if(is_composite && dim == 0) is_composite =
false;
63 if(is_composite && is_component)
errorQuda(
"\nComposite type is not implemented.\n");
144 gammaBasis(inv_param.gamma_basis),
155 for (
int d = 0; d < nDim; d++) x[d] = X[d];
205 nColor(cpuParam.nColor),
206 nSpin(cpuParam.nSpin),
208 twistFlavor(cpuParam.twistFlavor),
213 pc_type(cpuParam.pc_type),
215 is_composite(cpuParam.is_composite),
216 composite_dim(cpuParam.composite_dim),
233 bool native = force_native ? true :
false;
239 this->precision = precision;
250 printfQuda(
"twistFlavor = %d\n", twistFlavor);
252 for (
int d=0; d<nDim; d++)
printfQuda(
"x[%d] = %d\n", d, x[d]);
254 printfQuda(
"ghost_precision = %d\n", ghost_precision);
262 printfQuda(
"norm = %lx\n", (
unsigned long)norm);
264 if(is_composite)
printfQuda(
"Number of elements = %d\n", composite_dim);
314 void create(
int nDim,
const int *x,
int Nc,
int Ns,
int Nvec,
QudaTwistFlavorType Twistflavor,
384 void createGhostZone(
int nFace,
bool spin_project=
true)
const;
394 void setTuningString();
407 int Nvec()
const {
return nVec; }
409 int Ndim()
const {
return nDim; }
410 const int*
X()
const {
return x; }
411 int X(
int d)
const {
return x[d]; }
417 int Pad()
const {
return pad; }
424 void*
V() {
return v;}
425 const void*
V()
const {
return v;}
428 virtual const void*
Ghost2()
const {
return nullptr; }
438 void exchange(
void **ghost,
void **sendbuf,
int nFace=1)
const;
455 const MemoryLocation *halo_location=
nullptr,
bool gdr_send=
false,
bool gdr_recv=
false,
463 bool isNative()
const;
489 int GhostOffset(
const int i,
const int j)
const {
return ghostOffset[i][j]; }
493 void* Ghost(
const int i);
494 const void* Ghost(
const int i)
const;
495 void* GhostNorm(
const int i);
496 const void* GhostNorm(
const int i)
const;
501 void*
const* Ghost()
const;
521 virtual void Source(
const QudaSourceType sourceType,
const int st=0,
const int s=0,
const int c=0) = 0;
523 virtual void PrintVector(
unsigned int x)
const = 0;
530 void LatticeIndex(
int *y,
int i)
const;
537 void OffsetIndex(
int &i,
int *y)
const;
551 ColorSpinorField* CreateCoarse(
const int *geoBlockSize,
int spinBlockSize,
int Nvec,
565 ColorSpinorField* CreateFine(
const int *geoblockSize,
int spinBlockSize,
int Nvec,
587 #ifdef USE_TEXTURE_OBJECTS 588 cudaTextureObject_t tex;
589 cudaTextureObject_t texNorm;
590 void createTexObject();
591 void destroyTexObject();
592 mutable cudaTextureObject_t ghostTex[4];
593 mutable cudaTextureObject_t ghostTexNorm[4];
594 void createGhostTexObject()
const;
595 void destroyGhostTexObject()
const;
600 mutable void *ghost_field_tex[4];
635 void switchBufferPinned();
642 void createComms(
int nFace,
bool spin_project=
true);
649 void allocateGhostBuffer(
int nFace,
bool spin_project=
true)
const;
671 bool spin_project,
double a = 0,
double b = 0,
double c = 0);
674 const int dagger,cudaStream_t*
stream,
bool zero_copy=
false);
685 void sendGhost(
void *ghost_spinor,
const int nFace,
const int dim,
const QudaDirection dir,
697 void unpackGhost(
const void* ghost_spinor,
const int nFace,
const int dim,
711 void unpackGhostExtended(
const void* ghost_spinor,
const int nFace,
const QudaParity parity,
715 void streamInit(cudaStream_t *stream_p);
734 MemoryLocation location_label,
bool spin_project =
true,
double a = 0,
double b = 0,
double c = 0);
736 void packExtended(
const int nFace,
const int R[],
const int parity,
const int dagger,
737 const int dim, cudaStream_t *stream_p,
const bool zeroCopyPack=
false);
739 void gather(
int nFace,
int dagger,
int dir, cudaStream_t *stream_p=NULL);
750 void recvStart(
int nFace,
int dir,
int dagger=0, cudaStream_t *stream_p=
nullptr,
bool gdr=
false);
763 void sendStart(
int nFace,
int d,
int dagger=0, cudaStream_t *stream_p=
nullptr,
bool gdr=
false,
bool remote_write=
false);
775 void commsStart(
int nFace,
int d,
int dagger=0, cudaStream_t *stream_p=
nullptr,
bool gdr_send=
false,
bool gdr_recv=
false);
787 int commsQuery(
int nFace,
int d,
int dagger=0, cudaStream_t *stream_p=
nullptr,
bool gdr_send=
false,
bool gdr_recv=
false);
799 void commsWait(
int nFace,
int d,
int dagger=0, cudaStream_t *stream_p=
nullptr,
bool gdr_send=
false,
bool gdr_recv=
false);
801 void scatter(
int nFace,
int dagger,
int dir, cudaStream_t *stream_p);
802 void scatter(
int nFace,
int dagger,
int dir);
804 void scatterExtended(
int nFace,
int parity,
int dagger,
int dir);
807 if (bufferIndex < 2) {
808 return ghost_recv_buffer_d[bufferIndex];
810 return ghost_pinned_recv_buffer_hd[bufferIndex%2];
829 const MemoryLocation *halo_location=
nullptr,
bool gdr_send=
false,
bool gdr_recv=
false,
832 #ifdef USE_TEXTURE_OBJECTS 833 inline const cudaTextureObject_t& Tex()
const {
return tex; }
834 inline const cudaTextureObject_t& TexNorm()
const {
return texNorm; }
835 inline const cudaTextureObject_t& GhostTex()
const {
return ghostTex[bufferIndex]; }
836 inline const cudaTextureObject_t& GhostTexNorm()
const {
return ghostTexNorm[bufferIndex]; }
840 CompositeColorSpinorField& Components()
const;
847 void getTexObjectInfo()
const;
849 void Source(
const QudaSourceType sourceType,
const int st=0,
const int s=0,
const int c=0);
851 void PrintVector(
unsigned int x)
const;
861 void restore()
const;
898 void Source(
const QudaSourceType sourceType,
const int st=0,
const int s=0,
const int c=0);
913 void PrintVector(
unsigned int x)
const;
919 void allocateGhostBuffer(
int nFace)
const;
920 static void freeGhostBuffer(
void);
923 void unpackGhost(
void* ghost_spinor,
const int dim,
944 const MemoryLocation *halo_location=
nullptr,
bool gdr_send=
false,
bool gdr_recv=
false,
955 void restore()
const;
960 void *dstNorm=0,
void*srcNorm=0);
1018 errorQuda(
"PCTypes %d %d do not match (%s:%d in %s())\n", a.
PCType(), b.
PCType(), file, line, func);
1029 template <
typename... Args>
1036 #define checkPCType(...) PCType_(__func__, __FILE__, __LINE__, __VA_ARGS__) 1040 #endif // _COLOR_SPINOR_FIELD_H QudaDslashType dslash_type
QudaDiracFieldOrder dirac_order
static int initGhostFaceBuffer
int genericCompare(const cpuColorSpinorField &a, const cpuColorSpinorField &b, int tol)
CompositeColorSpinorField components
void genericPrintVector(const cpuColorSpinorField &a, unsigned int x)
int ComponentVolumeCB() const
void setPrecision(QudaPrecision precision, QudaPrecision ghost_precision=QUDA_INVALID_PRECISION, bool force_native=false)
QudaFieldOrder fieldOrder
enum QudaPrecision_s QudaPrecision
void destroy()
Destroy the CUBLAS context.
const void * Ghost2() const
const void * Norm() const
Constants used by dslash and packing kernels.
__host__ __device__ ValueType norm(const complex< ValueType > &z)
Returns the magnitude of z squared.
enum QudaNoiseType_s QudaNoiseType
int GhostOffset(const int i, const int j) const
size_t ComponentBytes() const
QudaDslashType dslash_type
enum QudaPCType_s QudaPCType
enum QudaFieldOrder_s QudaFieldOrder
CompositeColorSpinorFieldDescriptor composite_descr
used for deflation eigenvector sets etc.:
QudaPrecision & cuda_prec
CompositeColorSpinorFieldDescriptor()
enum QudaSiteOrder_s QudaSiteOrder
CompositeColorSpinorField & Components()
QudaGammaBasis gammaBasis
QudaGammaBasis GammaBasis() const
__host__ __device__ void copy(T1 &a, const T2 &b)
size_t ComponentLength() const
DslashConstant dslash_constant
void copyGenericColorSpinor(ColorSpinorField &dst, const ColorSpinorField &src, QudaFieldLocation location, void *Dst=0, void *Src=0, void *dstNorm=0, void *srcNorm=0)
ColorSpinorParam(ColorSpinorParam &cpuParam, QudaInvertParam &inv_param, QudaFieldLocation location=QUDA_CUDA_FIELD_LOCATION)
QudaSiteSubset siteSubset
std::ostream & operator<<(std::ostream &output, const CloverFieldParam ¶m)
virtual ~ColorSpinorParam()
enum QudaSourceType_s QudaSourceType
size_t RealLength() const
QudaPCType PCType_(const char *func, const char *file, int line, const ColorSpinorField &a, const ColorSpinorField &b)
Helper function for determining if the preconditioning type of the fields is the same.
bool is_composite
for deflation solvers:
enum QudaDirection_s QudaDirection
void genericPackGhost(void **ghost, const ColorSpinorField &a, QudaParity parity, int nFace, int dagger, MemoryLocation *destination=nullptr)
Generic ghost packing routine.
std::vector< ColorSpinorField * > CompositeColorSpinorField
QudaFieldLocation location
QudaInvertParam inv_param
__device__ __host__ __forceinline__ void packGhost(Arg &arg, int x_cb, int parity, int spinor_parity, int spin_block, int color_block)
QudaFieldOrder fieldOrder
int GhostNormOffset(const int i, const int j) const
QudaSiteSubset SiteSubset() const
Class declaration to initialize and hold CURAND RNG states.
QudaGammaBasis gammaBasis
const int * GhostFaceCB() const
QudaPCType PCType() const
QudaPrecision ghost_precision_tex
void exchangeExtendedGhost(cudaColorSpinorField *spinor, int R[], int parity, cudaStream_t *stream_p)
enum QudaParity_s QudaParity
void genericSource(cpuColorSpinorField &a, QudaSourceType sourceType, int x, int s, int c)
QudaTwistFlavorType twistFlavor
int GhostOffset(const int i) const
int ComponentVolume() const
enum QudaSiteSubset_s QudaSiteSubset
const DslashConstant & getDslashConstant() const
Get the dslash_constant structure from this field.
QudaPrecision ghost_precision_allocated
CompositeColorSpinorFieldDescriptor(bool is_composite, int dim, bool is_component=false, int id=0)
void wuppertalStep(ColorSpinorField &out, const ColorSpinorField &in, int parity, const GaugeField &U, double A, double B)
enum QudaFieldLocation_s QudaFieldLocation
cpuColorSpinorField * out
__device__ __host__ void pack(Arg &arg, int ghost_idx, int s, int parity)
void copyExtendedColorSpinor(ColorSpinorField &dst, const ColorSpinorField &src, QudaFieldLocation location, const int parity, void *Dst, void *Src, void *dstNorm, void *srcNorm)
enum QudaGammaBasis_s QudaGammaBasis
Main header file for the QUDA library.
int ComponentStride() const
size_t ComponentNormBytes() const
QudaTwistFlavorType twistFlavor
QudaSiteOrder SiteOrder() const
QudaTwistFlavorType TwistFlavor() const
QudaTwistFlavorType twist_flavor
const int * GhostFace() const
size_t ComponentRealLength() const
enum QudaFieldCreate_s QudaFieldCreate
virtual const void * Ghost2() const
#define QUDA_MAX_DIM
Maximum number of dimensions supported by QUDA. In practice, no routines make use of more than 5...
ColorSpinorParam(void *V, QudaInvertParam &inv_param, const int *X, const bool pc_solution, QudaFieldLocation location=QUDA_CPU_FIELD_LOCATION)
int GhostNormOffset(const int i) const
size_t GhostBytes() const
QudaSiteSubset siteSubset
QudaFieldOrder FieldOrder() const
__device__ __host__ void zero(vector_type< scalar, n > &v)
CompositeColorSpinorFieldDescriptor(const CompositeColorSpinorFieldDescriptor &descr)
enum QudaMemoryType_s QudaMemoryType
cpuColorSpinorField * spinor
size_t GhostNormBytes() const
void spinorNoise(ColorSpinorField &src, RNG &randstates, QudaNoiseType type)
Generate a random noise spinor. This variant allows the user to manage the RNG state.
void genericCudaPrintVector(const cudaColorSpinorField &a, unsigned x)
enum QudaTwistFlavorType_s QudaTwistFlavorType