mirror of
https://github.com/paboyle/Grid.git
synced 2025-06-13 04:37:05 +01:00
Compare commits
5 Commits
b461184797
...
feature/sy
| Author | SHA1 | Date | |
|---|---|---|---|
| 229709a980 | |||
| 9295eeadfe | |||
| 36f471e333 | |||
| ca4eadd4ab | |||
| d954595922 |
@ -54,7 +54,7 @@ Version.h: version-cache
|
||||
include Make.inc
|
||||
include Eigen.inc
|
||||
|
||||
extra_sources+=$(ZWILS_FERMION_FILES)
|
||||
#extra_sources+=$(ZWILS_FERMION_FILES)
|
||||
extra_sources+=$(WILS_FERMION_FILES)
|
||||
extra_sources+=$(STAG_FERMION_FILES)
|
||||
if BUILD_GPARITY
|
||||
|
||||
@ -36,7 +36,7 @@ static const int CbBlack=1;
|
||||
static const int Even =CbRed;
|
||||
static const int Odd =CbBlack;
|
||||
|
||||
accelerator_inline int RedBlackCheckerBoardFromOindex (int oindex, Coordinate &rdim, Coordinate &chk_dim_msk)
|
||||
accelerator_inline int RedBlackCheckerBoardFromOindex (int oindex,const Coordinate &rdim,const Coordinate &chk_dim_msk)
|
||||
{
|
||||
int nd=rdim.size();
|
||||
Coordinate coor(nd);
|
||||
|
||||
@ -67,8 +67,13 @@ public:
|
||||
accelerator_inline const vobj & operator()(size_t i) const { return this->_odata[i]; }
|
||||
#endif
|
||||
|
||||
accelerator_inline const vobj & operator[](size_t i) const { return this->_odata[i]; };
|
||||
accelerator_inline vobj & operator[](size_t i) { return this->_odata[i]; };
|
||||
#if 1
|
||||
// accelerator_inline const vobj & operator[](size_t i) const { return this->_odata[i]; };
|
||||
accelerator_inline vobj & operator[](size_t i) const { return this->_odata[i]; };
|
||||
#else
|
||||
// accelerator_inline const vobj & operator[](size_t i) const { return this->_odata[i]; };
|
||||
// accelerator_inline vobj & operator[](size_t i) { return this->_odata[i]; };
|
||||
#endif
|
||||
|
||||
accelerator_inline uint64_t begin(void) const { return 0;};
|
||||
accelerator_inline uint64_t end(void) const { return this->_odata_size; };
|
||||
|
||||
@ -88,7 +88,7 @@ public:
|
||||
const _Spinor &chi,
|
||||
int mu,
|
||||
StencilEntry *SE,
|
||||
StencilView &St)
|
||||
const StencilView &St)
|
||||
{
|
||||
int direction = St._directions[mu];
|
||||
int distance = St._distances[mu];
|
||||
|
||||
@ -56,8 +56,12 @@ template<class Impl> class StaggeredKernels : public FermionOperator<Impl> , pub
|
||||
DoubledGaugeField &U,
|
||||
const FermionField &in, FermionField &out, int dag, int interior,int exterior);
|
||||
|
||||
void DhopDirKernel(StencilImpl &st, DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, SiteSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out, int dir,int disp);
|
||||
void DhopDirKernel(StencilImpl &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU, SiteSpinor * buf,
|
||||
int sF, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out, int dir,int disp);
|
||||
protected:
|
||||
|
||||
///////////////////////////////////////////////////////////////////////////////////////
|
||||
@ -65,53 +69,67 @@ template<class Impl> class StaggeredKernels : public FermionOperator<Impl> , pub
|
||||
///////////////////////////////////////////////////////////////////////////////////////
|
||||
template<int Naik>
|
||||
static accelerator_inline
|
||||
void DhopSiteGeneric(StencilView &st,
|
||||
DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,
|
||||
void DhopSiteGeneric(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag);
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag);
|
||||
|
||||
template<int Naik> static accelerator_inline
|
||||
void DhopSiteGenericInt(StencilView &st,
|
||||
DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,
|
||||
void DhopSiteGenericInt(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag);
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag);
|
||||
|
||||
template<int Naik> static accelerator_inline
|
||||
void DhopSiteGenericExt(StencilView &st,
|
||||
DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag);
|
||||
void DhopSiteGenericExt(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag);
|
||||
|
||||
///////////////////////////////////////////////////////////////////////////////////////
|
||||
// Nc=3 specific kernels
|
||||
///////////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
template<int Naik> static accelerator_inline
|
||||
void DhopSiteHand(StencilView &st,
|
||||
DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag);
|
||||
void DhopSiteHand(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag);
|
||||
|
||||
template<int Naik> static accelerator_inline
|
||||
void DhopSiteHandInt(StencilView &st,
|
||||
DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag);
|
||||
void DhopSiteHandInt(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag);
|
||||
|
||||
template<int Naik> static accelerator_inline
|
||||
void DhopSiteHandExt(StencilView &st,
|
||||
DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag);
|
||||
void DhopSiteHandExt(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag);
|
||||
|
||||
///////////////////////////////////////////////////////////////////////////////////////
|
||||
// Asm Nc=3 specific kernels
|
||||
///////////////////////////////////////////////////////////////////////////////////////
|
||||
|
||||
void DhopSiteAsm(StencilView &st,
|
||||
DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag);
|
||||
void DhopSiteAsm(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf, int LLs, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag);
|
||||
|
||||
public:
|
||||
|
||||
|
||||
@ -61,7 +61,7 @@ public:
|
||||
typedef typename SiteHalfSpinor::vector_type vComplexHigh;
|
||||
constexpr static int Nw=sizeof(SiteHalfSpinor)/sizeof(vComplexHigh);
|
||||
|
||||
accelerator_inline int CommDatumSize(void) {
|
||||
accelerator_inline int CommDatumSize(void) const {
|
||||
return sizeof(SiteHalfCommSpinor);
|
||||
}
|
||||
|
||||
@ -69,7 +69,7 @@ public:
|
||||
/* Compress includes precision change if mpi data is not same */
|
||||
/*****************************************************/
|
||||
template<class _SiteHalfSpinor, class _SiteSpinor>
|
||||
accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) {
|
||||
accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) const {
|
||||
_SiteHalfSpinor tmp;
|
||||
projector::Proj(tmp,in,mu,dag);
|
||||
vstream(buf[o],tmp);
|
||||
@ -81,7 +81,7 @@ public:
|
||||
accelerator_inline void Exchange(SiteHalfSpinor *mp,
|
||||
const SiteHalfSpinor * __restrict__ vp0,
|
||||
const SiteHalfSpinor * __restrict__ vp1,
|
||||
Integer type,Integer o){
|
||||
Integer type,Integer o) const {
|
||||
SiteHalfSpinor tmp1;
|
||||
SiteHalfSpinor tmp2;
|
||||
exchange(tmp1,tmp2,vp0[o],vp1[o],type);
|
||||
@ -93,7 +93,7 @@ public:
|
||||
/* Have a decompression step if mpi data is not same */
|
||||
/*****************************************************/
|
||||
accelerator_inline void Decompress(SiteHalfSpinor * __restrict__ out,
|
||||
SiteHalfSpinor * __restrict__ in, Integer o) {
|
||||
SiteHalfSpinor * __restrict__ in, Integer o) const {
|
||||
assert(0);
|
||||
}
|
||||
|
||||
@ -103,7 +103,7 @@ public:
|
||||
accelerator_inline void CompressExchange(SiteHalfSpinor * __restrict__ out0,
|
||||
SiteHalfSpinor * __restrict__ out1,
|
||||
const SiteSpinor * __restrict__ in,
|
||||
Integer j,Integer k, Integer m,Integer type)
|
||||
Integer j,Integer k, Integer m,Integer type) const
|
||||
{
|
||||
SiteHalfSpinor temp1, temp2;
|
||||
SiteHalfSpinor temp3, temp4;
|
||||
@ -117,7 +117,7 @@ public:
|
||||
/*****************************************************/
|
||||
/* Pass the info to the stencil */
|
||||
/*****************************************************/
|
||||
accelerator_inline bool DecompressionStep(void) { return false; }
|
||||
accelerator_inline bool DecompressionStep(void) const { return false; }
|
||||
|
||||
};
|
||||
|
||||
@ -142,7 +142,7 @@ public:
|
||||
typedef typename SiteHalfSpinor::vector_type vComplexHigh;
|
||||
constexpr static int Nw=sizeof(SiteHalfSpinor)/sizeof(vComplexHigh);
|
||||
|
||||
accelerator_inline int CommDatumSize(void) {
|
||||
accelerator_inline int CommDatumSize(void) const {
|
||||
return sizeof(SiteHalfCommSpinor);
|
||||
}
|
||||
|
||||
@ -150,7 +150,7 @@ public:
|
||||
/* Compress includes precision change if mpi data is not same */
|
||||
/*****************************************************/
|
||||
template<class _SiteHalfSpinor, class _SiteSpinor>
|
||||
accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) {
|
||||
accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) const {
|
||||
_SiteHalfSpinor hsp;
|
||||
SiteHalfCommSpinor *hbuf = (SiteHalfCommSpinor *)buf;
|
||||
projector::Proj(hsp,in,mu,dag);
|
||||
@ -163,7 +163,7 @@ public:
|
||||
accelerator_inline void Exchange(SiteHalfSpinor *mp,
|
||||
SiteHalfSpinor *vp0,
|
||||
SiteHalfSpinor *vp1,
|
||||
Integer type,Integer o){
|
||||
Integer type,Integer o) const {
|
||||
SiteHalfSpinor vt0,vt1;
|
||||
SiteHalfCommSpinor *vpp0 = (SiteHalfCommSpinor *)vp0;
|
||||
SiteHalfCommSpinor *vpp1 = (SiteHalfCommSpinor *)vp1;
|
||||
@ -175,7 +175,7 @@ public:
|
||||
/*****************************************************/
|
||||
/* Have a decompression step if mpi data is not same */
|
||||
/*****************************************************/
|
||||
accelerator_inline void Decompress(SiteHalfSpinor *out, SiteHalfSpinor *in, Integer o){
|
||||
accelerator_inline void Decompress(SiteHalfSpinor *out, SiteHalfSpinor *in, Integer o) const {
|
||||
SiteHalfCommSpinor *hin=(SiteHalfCommSpinor *)in;
|
||||
precisionChange((vComplexHigh *)&out[o],(vComplexLow *)&hin[o],Nw);
|
||||
}
|
||||
@ -186,7 +186,7 @@ public:
|
||||
accelerator_inline void CompressExchange(SiteHalfSpinor *out0,
|
||||
SiteHalfSpinor *out1,
|
||||
const SiteSpinor *in,
|
||||
Integer j,Integer k, Integer m,Integer type){
|
||||
Integer j,Integer k, Integer m,Integer type) const {
|
||||
SiteHalfSpinor temp1, temp2,temp3,temp4;
|
||||
SiteHalfCommSpinor *hout0 = (SiteHalfCommSpinor *)out0;
|
||||
SiteHalfCommSpinor *hout1 = (SiteHalfCommSpinor *)out1;
|
||||
@ -200,7 +200,7 @@ public:
|
||||
/*****************************************************/
|
||||
/* Pass the info to the stencil */
|
||||
/*****************************************************/
|
||||
accelerator_inline bool DecompressionStep(void) { return true; }
|
||||
accelerator_inline bool DecompressionStep(void) const { return true; }
|
||||
|
||||
};
|
||||
|
||||
|
||||
@ -95,7 +95,7 @@ public:
|
||||
const _Spinor &chi,
|
||||
int mu,
|
||||
StencilEntry *SE,
|
||||
StencilView &St)
|
||||
const StencilView &St)
|
||||
{
|
||||
multLink(phi,U,chi,mu);
|
||||
}
|
||||
|
||||
@ -49,6 +49,7 @@ public:
|
||||
|
||||
INHERIT_IMPL_TYPES(Impl);
|
||||
typedef FermionOperator<Impl> Base;
|
||||
typedef AcceleratorVector<int,STENCIL_MAX> StencilVector;
|
||||
|
||||
public:
|
||||
|
||||
@ -68,73 +69,87 @@ public:
|
||||
|
||||
private:
|
||||
|
||||
static accelerator_inline void DhopDirK(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out, int dirdisp, int gamma);
|
||||
static accelerator_inline void DhopDirK(const StencilView &st, const DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor * buf, int sF, int sU,
|
||||
const FermionFieldView &in,const FermionFieldView &out, int dirdisp, int gamma);
|
||||
|
||||
static accelerator_inline void DhopDirXp(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirYp(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirZp(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirTp(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirXm(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirYm(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirZm(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirTm(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirXp(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirYp(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirZp(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirTp(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirXm(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirYm(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirZm(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dirdisp);
|
||||
static accelerator_inline void DhopDirTm(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dirdisp);
|
||||
|
||||
// Specialised variants
|
||||
static accelerator void GenericDhopSite(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteInt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteDagInt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteExt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteDagExt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
static accelerator void GenericDhopSite(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void GenericDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static void AsmDhopSite(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,FermionFieldView &out);
|
||||
|
||||
static void AsmDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, FermionFieldView &out);
|
||||
|
||||
static void AsmDhopSiteInt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,FermionFieldView &out);
|
||||
|
||||
static void AsmDhopSiteDagInt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, FermionFieldView &out);
|
||||
|
||||
static void AsmDhopSiteExt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,FermionFieldView &out);
|
||||
|
||||
static void AsmDhopSiteDagExt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, FermionFieldView &out);
|
||||
// Keep Hand unrolled
|
||||
static accelerator void HandDhopSiteSycl(StencilVector st_perm, StencilEntry *st_p, SiteDoubledGaugeField *U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const SiteSpinor *in, SiteSpinor *out);
|
||||
|
||||
// Keep Hand unrolled temporarily
|
||||
static accelerator void HandDhopSite(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
static accelerator void HandDhopSite(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in,const FermionFieldView &out);
|
||||
|
||||
static accelerator void HandDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void HandDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void HandDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
static accelerator void HandDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void HandDhopSiteInt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
static accelerator void HandDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void HandDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, const FermionFieldView &out);
|
||||
//AVX 512 ASM
|
||||
static void AsmDhopSite(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,const FermionFieldView &out);
|
||||
|
||||
static accelerator void HandDhopSiteDagInt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
static void AsmDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static accelerator void HandDhopSiteExt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
static void AsmDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,const FermionFieldView &out);
|
||||
|
||||
static accelerator void HandDhopSiteDagExt(StencilView &st, DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out);
|
||||
static void AsmDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
static void AsmDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,const FermionFieldView &out);
|
||||
|
||||
static void AsmDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf,
|
||||
int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, const FermionFieldView &out);
|
||||
|
||||
public:
|
||||
WilsonKernels(const ImplParams &p = ImplParams()) : Base(p){};
|
||||
};
|
||||
|
||||
@ -618,11 +618,13 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
template <class Impl>
|
||||
void StaggeredKernels<Impl>::DhopSiteAsm(StencilView &st,
|
||||
DoubledGaugeFieldView &U,
|
||||
DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out,int dag)
|
||||
void StaggeredKernels<Impl>::DhopSiteAsm(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag)
|
||||
{
|
||||
assert(0);
|
||||
};
|
||||
@ -683,11 +685,13 @@ void StaggeredKernels<Impl>::DhopSiteAsm(StencilView &st,
|
||||
|
||||
// This is the single precision 5th direction vectorised kernel
|
||||
#include <Grid/simd/Intel512single.h>
|
||||
template <> void StaggeredKernels<StaggeredVec5dImplF>::DhopSiteAsm(StencilView &st,
|
||||
DoubledGaugeFieldView &U,
|
||||
DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out,int dag)
|
||||
template <> void StaggeredKernels<StaggeredVec5dImplF>::DhopSiteAsm(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag)
|
||||
{
|
||||
#ifdef AVX512
|
||||
uint64_t gauge0,gauge1,gauge2,gauge3;
|
||||
@ -738,11 +742,13 @@ template <> void StaggeredKernels<StaggeredVec5dImplF>::DhopSiteAsm(StencilView
|
||||
}
|
||||
|
||||
#include <Grid/simd/Intel512double.h>
|
||||
template <> void StaggeredKernels<StaggeredVec5dImplD>::DhopSiteAsm(StencilView &st,
|
||||
DoubledGaugeFieldView &U,
|
||||
DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out, int dag)
|
||||
template <> void StaggeredKernels<StaggeredVec5dImplD>::DhopSiteAsm(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out, int dag)
|
||||
{
|
||||
#ifdef AVX512
|
||||
uint64_t gauge0,gauge1,gauge2,gauge3;
|
||||
@ -824,11 +830,13 @@ template <> void StaggeredKernels<StaggeredVec5dImplD>::DhopSiteAsm(StencilView
|
||||
// This is the single precision 5th direction vectorised kernel
|
||||
|
||||
#include <Grid/simd/Intel512single.h>
|
||||
template <> void StaggeredKernels<StaggeredImplF>::DhopSiteAsm(StencilView &st,
|
||||
DoubledGaugeFieldView &U,
|
||||
DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out,int dag)
|
||||
template <> void StaggeredKernels<StaggeredImplF>::DhopSiteAsm(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag)
|
||||
{
|
||||
#ifdef AVX512
|
||||
uint64_t gauge0,gauge1,gauge2,gauge3;
|
||||
@ -893,11 +901,13 @@ template <> void StaggeredKernels<StaggeredImplF>::DhopSiteAsm(StencilView &st,
|
||||
}
|
||||
|
||||
#include <Grid/simd/Intel512double.h>
|
||||
template <> void StaggeredKernels<StaggeredImplD>::DhopSiteAsm(StencilView &st,
|
||||
DoubledGaugeFieldView &U,
|
||||
DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out,int dag)
|
||||
template <> void StaggeredKernels<StaggeredImplD>::DhopSiteAsm(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF,
|
||||
int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag)
|
||||
{
|
||||
#ifdef AVX512
|
||||
uint64_t gauge0,gauge1,gauge2,gauge3;
|
||||
|
||||
@ -147,10 +147,12 @@ NAMESPACE_BEGIN(Grid);
|
||||
|
||||
template <class Impl>
|
||||
template <int Naik> accelerator_inline
|
||||
void StaggeredKernels<Impl>::DhopSiteHand(StencilView &st,
|
||||
DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag)
|
||||
void StaggeredKernels<Impl>::DhopSiteHand(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag)
|
||||
{
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
@ -222,10 +224,12 @@ void StaggeredKernels<Impl>::DhopSiteHand(StencilView &st,
|
||||
|
||||
template <class Impl>
|
||||
template <int Naik> accelerator_inline
|
||||
void StaggeredKernels<Impl>::DhopSiteHandInt(StencilView &st,
|
||||
DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag)
|
||||
void StaggeredKernels<Impl>::DhopSiteHandInt(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag)
|
||||
{
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
@ -301,10 +305,12 @@ void StaggeredKernels<Impl>::DhopSiteHandInt(StencilView &st,
|
||||
|
||||
template <class Impl>
|
||||
template <int Naik> accelerator_inline
|
||||
void StaggeredKernels<Impl>::DhopSiteHandExt(StencilView &st,
|
||||
DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag)
|
||||
void StaggeredKernels<Impl>::DhopSiteHandExt(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag)
|
||||
{
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
|
||||
@ -79,10 +79,10 @@ StaggeredKernels<Impl>::StaggeredKernels(const ImplParams &p) : Base(p){};
|
||||
////////////////////////////////////////////////////////////////////////////////////
|
||||
template <class Impl>
|
||||
template <int Naik> accelerator_inline
|
||||
void StaggeredKernels<Impl>::DhopSiteGeneric(StencilView &st,
|
||||
DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out, int dag)
|
||||
void StaggeredKernels<Impl>::DhopSiteGeneric(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U, const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out, int dag)
|
||||
{
|
||||
const SiteSpinor *chi_p;
|
||||
SiteSpinor chi;
|
||||
@ -127,10 +127,11 @@ void StaggeredKernels<Impl>::DhopSiteGeneric(StencilView &st,
|
||||
///////////////////////////////////////////////////
|
||||
template <class Impl>
|
||||
template <int Naik> accelerator_inline
|
||||
void StaggeredKernels<Impl>::DhopSiteGenericInt(StencilView &st,
|
||||
DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag) {
|
||||
void StaggeredKernels<Impl>::DhopSiteGenericInt(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U, const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in, const FermionFieldView &out,int dag)
|
||||
{
|
||||
const SiteSpinor *chi_p;
|
||||
SiteSpinor chi;
|
||||
SiteSpinor Uchi;
|
||||
@ -175,10 +176,13 @@ void StaggeredKernels<Impl>::DhopSiteGenericInt(StencilView &st,
|
||||
///////////////////////////////////////////////////
|
||||
template <class Impl>
|
||||
template <int Naik> accelerator_inline
|
||||
void StaggeredKernels<Impl>::DhopSiteGenericExt(StencilView &st,
|
||||
DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in, FermionFieldView &out,int dag) {
|
||||
void StaggeredKernels<Impl>::DhopSiteGenericExt(const StencilView &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor *buf, int sF, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out,int dag)
|
||||
{
|
||||
const SiteSpinor *chi_p;
|
||||
// SiteSpinor chi;
|
||||
SiteSpinor Uchi;
|
||||
@ -225,8 +229,13 @@ void StaggeredKernels<Impl>::DhopSiteGenericExt(StencilView &st,
|
||||
// Driving / wrapping routine to select right kernel
|
||||
////////////////////////////////////////////////////////////////////////////////////
|
||||
template <class Impl>
|
||||
void StaggeredKernels<Impl>::DhopDirKernel(StencilImpl &st, DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, SiteSpinor * buf,
|
||||
int sF, int sU, const FermionFieldView &in, FermionFieldView &out, int dir,int disp)
|
||||
void StaggeredKernels<Impl>::DhopDirKernel(StencilImpl &st,
|
||||
const DoubledGaugeFieldView &U,
|
||||
const DoubledGaugeFieldView &UUU,
|
||||
SiteSpinor * buf,
|
||||
int sF, int sU,
|
||||
const FermionFieldView &in,
|
||||
const FermionFieldView &out, int dir,int disp)
|
||||
{
|
||||
// Disp should be either +1,-1,+3,-3
|
||||
// What about "dag" ?
|
||||
@ -254,7 +263,8 @@ void StaggeredKernels<Impl>::DhopDirKernel(StencilImpl &st, DoubledGaugeFieldVie
|
||||
});
|
||||
|
||||
template <class Impl>
|
||||
void StaggeredKernels<Impl>::DhopImproved(StencilImpl &st, LebesgueOrder &lo,
|
||||
void StaggeredKernels<Impl>::DhopImproved(StencilImpl &st,
|
||||
LebesgueOrder &lo,
|
||||
DoubledGaugeField &U, DoubledGaugeField &UUU,
|
||||
const FermionField &in, FermionField &out, int dag, int interior,int exterior)
|
||||
{
|
||||
|
||||
@ -38,46 +38,46 @@ NAMESPACE_BEGIN(Grid);
|
||||
|
||||
///////////////////////////////////////////////////////////
|
||||
// Default to no assembler implementation
|
||||
// Will specialise to
|
||||
// Will specialise to AVX512 if available
|
||||
///////////////////////////////////////////////////////////
|
||||
template<class Impl> void
|
||||
WilsonKernels<Impl >::AsmDhopSite(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl >::AsmDhopSite(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
assert(0);
|
||||
}
|
||||
|
||||
template<class Impl> void
|
||||
WilsonKernels<Impl >::AsmDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl >::AsmDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
assert(0);
|
||||
}
|
||||
|
||||
template<class Impl> void
|
||||
WilsonKernels<Impl >::AsmDhopSiteInt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl >::AsmDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
assert(0);
|
||||
}
|
||||
|
||||
template<class Impl> void
|
||||
WilsonKernels<Impl >::AsmDhopSiteDagInt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl >::AsmDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
assert(0);
|
||||
}
|
||||
|
||||
template<class Impl> void
|
||||
WilsonKernels<Impl >::AsmDhopSiteExt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl >::AsmDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
assert(0);
|
||||
}
|
||||
|
||||
template<class Impl> void
|
||||
WilsonKernels<Impl >::AsmDhopSiteDagExt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl >::AsmDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
assert(0);
|
||||
}
|
||||
|
||||
@ -646,9 +646,14 @@ NAMESPACE_BEGIN(Grid);
|
||||
HAND_RESULT_EXT(ss,F)
|
||||
|
||||
#define HAND_SPECIALISE_GPARITY(IMPL) \
|
||||
template<> accelerator_inline void \
|
||||
WilsonKernels<IMPL>::HandDhopSiteSycl(StencilVector st_perm, StencilEntry *st_p, \
|
||||
SiteDoubledGaugeField *U, SiteHalfSpinor * buf, \
|
||||
int sF, int sU, const SiteSpinor *in, SiteSpinor *out) {} \
|
||||
\
|
||||
template<> accelerator_inline void \
|
||||
WilsonKernels<IMPL>::HandDhopSite(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \
|
||||
WilsonKernels<IMPL>::HandDhopSite(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \
|
||||
{ \
|
||||
typedef IMPL Impl; \
|
||||
typedef typename Simd::scalar_type S; \
|
||||
@ -663,8 +668,8 @@ NAMESPACE_BEGIN(Grid);
|
||||
} \
|
||||
\
|
||||
template<> accelerator_inline void \
|
||||
WilsonKernels<IMPL>::HandDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \
|
||||
WilsonKernels<IMPL>::HandDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \
|
||||
{ \
|
||||
typedef IMPL Impl; \
|
||||
typedef typename Simd::scalar_type S; \
|
||||
@ -679,8 +684,8 @@ NAMESPACE_BEGIN(Grid);
|
||||
} \
|
||||
\
|
||||
template<> accelerator_inline void \
|
||||
WilsonKernels<IMPL>::HandDhopSiteInt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \
|
||||
WilsonKernels<IMPL>::HandDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \
|
||||
{ \
|
||||
typedef IMPL Impl; \
|
||||
typedef typename Simd::scalar_type S; \
|
||||
@ -695,8 +700,8 @@ NAMESPACE_BEGIN(Grid);
|
||||
} \
|
||||
\
|
||||
template<> accelerator_inline void \
|
||||
WilsonKernels<IMPL>::HandDhopSiteDagInt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \
|
||||
WilsonKernels<IMPL>::HandDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \
|
||||
{ \
|
||||
typedef IMPL Impl; \
|
||||
typedef typename Simd::scalar_type S; \
|
||||
@ -711,8 +716,8 @@ NAMESPACE_BEGIN(Grid);
|
||||
} \
|
||||
\
|
||||
template<> accelerator_inline void \
|
||||
WilsonKernels<IMPL>::HandDhopSiteExt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \
|
||||
WilsonKernels<IMPL>::HandDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \
|
||||
{ \
|
||||
typedef IMPL Impl; \
|
||||
typedef typename Simd::scalar_type S; \
|
||||
@ -728,8 +733,8 @@ NAMESPACE_BEGIN(Grid);
|
||||
HAND_DOP_SITE_EXT(1, LOAD_CHI_GPARITY,LOAD_CHIMU_GPARITY,MULT_2SPIN_GPARITY); \
|
||||
} \
|
||||
template<> accelerator_inline void \
|
||||
WilsonKernels<IMPL>::HandDhopSiteDagExt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \
|
||||
WilsonKernels<IMPL>::HandDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \
|
||||
{ \
|
||||
typedef IMPL Impl; \
|
||||
typedef typename Simd::scalar_type S; \
|
||||
|
||||
@ -496,8 +496,8 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
template<class Impl> accelerator_inline void
|
||||
WilsonKernels<Impl>::HandDhopSite(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl>::HandDhopSite(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
// T==0, Z==1, Y==2, Z==3 expect 1,2,2,2 simd layout etc...
|
||||
typedef typename Simd::scalar_type S;
|
||||
@ -520,8 +520,8 @@ WilsonKernels<Impl>::HandDhopSite(StencilView &st, DoubledGaugeFieldView &U,Site
|
||||
}
|
||||
|
||||
template<class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::HandDhopSiteDag(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::HandDhopSiteDag(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
@ -543,8 +543,8 @@ void WilsonKernels<Impl>::HandDhopSiteDag(StencilView &st,DoubledGaugeFieldView
|
||||
}
|
||||
|
||||
template<class Impl> accelerator_inline void
|
||||
WilsonKernels<Impl>::HandDhopSiteInt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl>::HandDhopSiteInt(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
// T==0, Z==1, Y==2, Z==3 expect 1,2,2,2 simd layout etc...
|
||||
typedef typename Simd::scalar_type S;
|
||||
@ -567,8 +567,8 @@ WilsonKernels<Impl>::HandDhopSiteInt(StencilView &st,DoubledGaugeFieldView &U,Si
|
||||
}
|
||||
|
||||
template<class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::HandDhopSiteDagInt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::HandDhopSiteDagInt(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
@ -590,8 +590,8 @@ void WilsonKernels<Impl>::HandDhopSiteDagInt(StencilView &st,DoubledGaugeFieldVi
|
||||
}
|
||||
|
||||
template<class Impl> accelerator_inline void
|
||||
WilsonKernels<Impl>::HandDhopSiteExt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
WilsonKernels<Impl>::HandDhopSiteExt(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
// T==0, Z==1, Y==2, Z==3 expect 1,2,2,2 simd layout etc...
|
||||
typedef typename Simd::scalar_type S;
|
||||
@ -615,8 +615,8 @@ WilsonKernels<Impl>::HandDhopSiteExt(StencilView &st,DoubledGaugeFieldView &U,Si
|
||||
}
|
||||
|
||||
template<class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::HandDhopSiteDagExt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::HandDhopSiteDagExt(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
@ -682,3 +682,4 @@ NAMESPACE_END(Grid);
|
||||
#undef HAND_RESULT
|
||||
#undef HAND_RESULT_INT
|
||||
#undef HAND_RESULT_EXT
|
||||
#undef HAND_DECLARATIONS
|
||||
|
||||
@ -0,0 +1,598 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernelsHand.cc
|
||||
|
||||
Copyright (C) 2015
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
|
||||
|
||||
#undef LOAD_CHIMU
|
||||
#undef LOAD_CHI
|
||||
#undef MULT_2SPIN
|
||||
#undef PERMUTE_DIR
|
||||
#undef XP_PROJ
|
||||
#undef YP_PROJ
|
||||
#undef ZP_PROJ
|
||||
#undef TP_PROJ
|
||||
#undef XM_PROJ
|
||||
#undef YM_PROJ
|
||||
#undef ZM_PROJ
|
||||
#undef TM_PROJ
|
||||
#undef XP_RECON
|
||||
#undef XP_RECON_ACCUM
|
||||
#undef XM_RECON
|
||||
#undef XM_RECON_ACCUM
|
||||
#undef YP_RECON_ACCUM
|
||||
#undef YM_RECON_ACCUM
|
||||
#undef ZP_RECON_ACCUM
|
||||
#undef ZM_RECON_ACCUM
|
||||
#undef TP_RECON_ACCUM
|
||||
#undef TM_RECON_ACCUM
|
||||
#undef ZERO_RESULT
|
||||
#undef Chimu_00
|
||||
#undef Chimu_01
|
||||
#undef Chimu_02
|
||||
#undef Chimu_10
|
||||
#undef Chimu_11
|
||||
#undef Chimu_12
|
||||
#undef Chimu_20
|
||||
#undef Chimu_21
|
||||
#undef Chimu_22
|
||||
#undef Chimu_30
|
||||
#undef Chimu_31
|
||||
#undef Chimu_32
|
||||
#undef HAND_STENCIL_LEG
|
||||
#undef HAND_STENCIL_LEG_INT
|
||||
#undef HAND_STENCIL_LEG_EXT
|
||||
#undef HAND_RESULT
|
||||
#undef HAND_RESULT_INT
|
||||
#undef HAND_RESULT_EXT
|
||||
|
||||
#define REGISTER
|
||||
|
||||
#ifdef GRID_SIMT
|
||||
#define LOAD_CHIMU(ptype) \
|
||||
{const SiteSpinor & ref (in[offset]); \
|
||||
Chimu_00=coalescedReadPermute<ptype>(ref()(0)(0),perm); \
|
||||
Chimu_01=coalescedReadPermute<ptype>(ref()(0)(1),perm); \
|
||||
Chimu_02=coalescedReadPermute<ptype>(ref()(0)(2),perm); \
|
||||
Chimu_10=coalescedReadPermute<ptype>(ref()(1)(0),perm); \
|
||||
Chimu_11=coalescedReadPermute<ptype>(ref()(1)(1),perm); \
|
||||
Chimu_12=coalescedReadPermute<ptype>(ref()(1)(2),perm); \
|
||||
Chimu_20=coalescedReadPermute<ptype>(ref()(2)(0),perm); \
|
||||
Chimu_21=coalescedReadPermute<ptype>(ref()(2)(1),perm); \
|
||||
Chimu_22=coalescedReadPermute<ptype>(ref()(2)(2),perm); \
|
||||
Chimu_30=coalescedReadPermute<ptype>(ref()(3)(0),perm); \
|
||||
Chimu_31=coalescedReadPermute<ptype>(ref()(3)(1),perm); \
|
||||
Chimu_32=coalescedReadPermute<ptype>(ref()(3)(2),perm); }
|
||||
|
||||
#define PERMUTE_DIR(dir) ;
|
||||
#else
|
||||
#define LOAD_CHIMU(ptype) \
|
||||
{const SiteSpinor & ref (in[offset]); \
|
||||
Chimu_00=coalescedRead(ref()(0)(0)); \
|
||||
Chimu_01=coalescedRead(ref()(0)(1)); \
|
||||
Chimu_02=coalescedRead(ref()(0)(2)); \
|
||||
Chimu_10=coalescedRead(ref()(1)(0)); \
|
||||
Chimu_11=coalescedRead(ref()(1)(1)); \
|
||||
Chimu_12=coalescedRead(ref()(1)(2)); \
|
||||
Chimu_20=coalescedRead(ref()(2)(0)); \
|
||||
Chimu_21=coalescedRead(ref()(2)(1)); \
|
||||
Chimu_22=coalescedRead(ref()(2)(2)); \
|
||||
Chimu_30=coalescedRead(ref()(3)(0)); \
|
||||
Chimu_31=coalescedRead(ref()(3)(1)); \
|
||||
Chimu_32=coalescedRead(ref()(3)(2)); }
|
||||
|
||||
#define PERMUTE_DIR(dir) \
|
||||
permute##dir(Chi_00,Chi_00); \
|
||||
permute##dir(Chi_01,Chi_01);\
|
||||
permute##dir(Chi_02,Chi_02);\
|
||||
permute##dir(Chi_10,Chi_10); \
|
||||
permute##dir(Chi_11,Chi_11);\
|
||||
permute##dir(Chi_12,Chi_12);
|
||||
#endif
|
||||
|
||||
#define MULT_2SPIN(A)\
|
||||
{auto & ref(U[sU](A)); \
|
||||
U_00=coalescedRead(ref()(0,0)); \
|
||||
U_10=coalescedRead(ref()(1,0)); \
|
||||
U_20=coalescedRead(ref()(2,0)); \
|
||||
U_01=coalescedRead(ref()(0,1)); \
|
||||
U_11=coalescedRead(ref()(1,1)); \
|
||||
U_21=coalescedRead(ref()(2,1)); \
|
||||
UChi_00 = U_00*Chi_00; \
|
||||
UChi_10 = U_00*Chi_10; \
|
||||
UChi_01 = U_10*Chi_00; \
|
||||
UChi_11 = U_10*Chi_10; \
|
||||
UChi_02 = U_20*Chi_00; \
|
||||
UChi_12 = U_20*Chi_10; \
|
||||
UChi_00+= U_01*Chi_01; \
|
||||
UChi_10+= U_01*Chi_11; \
|
||||
UChi_01+= U_11*Chi_01; \
|
||||
UChi_11+= U_11*Chi_11; \
|
||||
UChi_02+= U_21*Chi_01; \
|
||||
UChi_12+= U_21*Chi_11; \
|
||||
U_00=coalescedRead(ref()(0,2)); \
|
||||
U_10=coalescedRead(ref()(1,2)); \
|
||||
U_20=coalescedRead(ref()(2,2)); \
|
||||
UChi_00+= U_00*Chi_02; \
|
||||
UChi_10+= U_00*Chi_12; \
|
||||
UChi_01+= U_10*Chi_02; \
|
||||
UChi_11+= U_10*Chi_12; \
|
||||
UChi_02+= U_20*Chi_02; \
|
||||
UChi_12+= U_20*Chi_12;}
|
||||
|
||||
#define LOAD_CHI \
|
||||
{const SiteHalfSpinor &ref(buf[offset]); \
|
||||
Chi_00 = coalescedRead(ref()(0)(0)); \
|
||||
Chi_01 = coalescedRead(ref()(0)(1)); \
|
||||
Chi_02 = coalescedRead(ref()(0)(2)); \
|
||||
Chi_10 = coalescedRead(ref()(1)(0)); \
|
||||
Chi_11 = coalescedRead(ref()(1)(1)); \
|
||||
Chi_12 = coalescedRead(ref()(1)(2));}
|
||||
|
||||
// hspin(0)=fspin(0)+timesI(fspin(3));
|
||||
// hspin(1)=fspin(1)+timesI(fspin(2));
|
||||
#define XP_PROJ \
|
||||
Chi_00 = Chimu_00+timesI(Chimu_30);\
|
||||
Chi_01 = Chimu_01+timesI(Chimu_31);\
|
||||
Chi_02 = Chimu_02+timesI(Chimu_32);\
|
||||
Chi_10 = Chimu_10+timesI(Chimu_20);\
|
||||
Chi_11 = Chimu_11+timesI(Chimu_21);\
|
||||
Chi_12 = Chimu_12+timesI(Chimu_22);
|
||||
|
||||
#define YP_PROJ \
|
||||
Chi_00 = Chimu_00-Chimu_30;\
|
||||
Chi_01 = Chimu_01-Chimu_31;\
|
||||
Chi_02 = Chimu_02-Chimu_32;\
|
||||
Chi_10 = Chimu_10+Chimu_20;\
|
||||
Chi_11 = Chimu_11+Chimu_21;\
|
||||
Chi_12 = Chimu_12+Chimu_22;
|
||||
|
||||
#define ZP_PROJ \
|
||||
Chi_00 = Chimu_00+timesI(Chimu_20); \
|
||||
Chi_01 = Chimu_01+timesI(Chimu_21); \
|
||||
Chi_02 = Chimu_02+timesI(Chimu_22); \
|
||||
Chi_10 = Chimu_10-timesI(Chimu_30); \
|
||||
Chi_11 = Chimu_11-timesI(Chimu_31); \
|
||||
Chi_12 = Chimu_12-timesI(Chimu_32);
|
||||
|
||||
#define TP_PROJ \
|
||||
Chi_00 = Chimu_00+Chimu_20; \
|
||||
Chi_01 = Chimu_01+Chimu_21; \
|
||||
Chi_02 = Chimu_02+Chimu_22; \
|
||||
Chi_10 = Chimu_10+Chimu_30; \
|
||||
Chi_11 = Chimu_11+Chimu_31; \
|
||||
Chi_12 = Chimu_12+Chimu_32;
|
||||
|
||||
|
||||
// hspin(0)=fspin(0)-timesI(fspin(3));
|
||||
// hspin(1)=fspin(1)-timesI(fspin(2));
|
||||
#define XM_PROJ \
|
||||
Chi_00 = Chimu_00-timesI(Chimu_30);\
|
||||
Chi_01 = Chimu_01-timesI(Chimu_31);\
|
||||
Chi_02 = Chimu_02-timesI(Chimu_32);\
|
||||
Chi_10 = Chimu_10-timesI(Chimu_20);\
|
||||
Chi_11 = Chimu_11-timesI(Chimu_21);\
|
||||
Chi_12 = Chimu_12-timesI(Chimu_22);
|
||||
|
||||
#define YM_PROJ \
|
||||
Chi_00 = Chimu_00+Chimu_30;\
|
||||
Chi_01 = Chimu_01+Chimu_31;\
|
||||
Chi_02 = Chimu_02+Chimu_32;\
|
||||
Chi_10 = Chimu_10-Chimu_20;\
|
||||
Chi_11 = Chimu_11-Chimu_21;\
|
||||
Chi_12 = Chimu_12-Chimu_22;
|
||||
|
||||
#define ZM_PROJ \
|
||||
Chi_00 = Chimu_00-timesI(Chimu_20); \
|
||||
Chi_01 = Chimu_01-timesI(Chimu_21); \
|
||||
Chi_02 = Chimu_02-timesI(Chimu_22); \
|
||||
Chi_10 = Chimu_10+timesI(Chimu_30); \
|
||||
Chi_11 = Chimu_11+timesI(Chimu_31); \
|
||||
Chi_12 = Chimu_12+timesI(Chimu_32);
|
||||
|
||||
#define TM_PROJ \
|
||||
Chi_00 = Chimu_00-Chimu_20; \
|
||||
Chi_01 = Chimu_01-Chimu_21; \
|
||||
Chi_02 = Chimu_02-Chimu_22; \
|
||||
Chi_10 = Chimu_10-Chimu_30; \
|
||||
Chi_11 = Chimu_11-Chimu_31; \
|
||||
Chi_12 = Chimu_12-Chimu_32;
|
||||
|
||||
// fspin(0)=hspin(0);
|
||||
// fspin(1)=hspin(1);
|
||||
// fspin(2)=timesMinusI(hspin(1));
|
||||
// fspin(3)=timesMinusI(hspin(0));
|
||||
#define XP_RECON\
|
||||
result_00 = UChi_00;\
|
||||
result_01 = UChi_01;\
|
||||
result_02 = UChi_02;\
|
||||
result_10 = UChi_10;\
|
||||
result_11 = UChi_11;\
|
||||
result_12 = UChi_12;\
|
||||
result_20 = timesMinusI(UChi_10);\
|
||||
result_21 = timesMinusI(UChi_11);\
|
||||
result_22 = timesMinusI(UChi_12);\
|
||||
result_30 = timesMinusI(UChi_00);\
|
||||
result_31 = timesMinusI(UChi_01);\
|
||||
result_32 = timesMinusI(UChi_02);
|
||||
|
||||
#define XP_RECON_ACCUM\
|
||||
result_00+=UChi_00;\
|
||||
result_01+=UChi_01;\
|
||||
result_02+=UChi_02;\
|
||||
result_10+=UChi_10;\
|
||||
result_11+=UChi_11;\
|
||||
result_12+=UChi_12;\
|
||||
result_20-=timesI(UChi_10);\
|
||||
result_21-=timesI(UChi_11);\
|
||||
result_22-=timesI(UChi_12);\
|
||||
result_30-=timesI(UChi_00);\
|
||||
result_31-=timesI(UChi_01);\
|
||||
result_32-=timesI(UChi_02);
|
||||
|
||||
#define XM_RECON\
|
||||
result_00 = UChi_00;\
|
||||
result_01 = UChi_01;\
|
||||
result_02 = UChi_02;\
|
||||
result_10 = UChi_10;\
|
||||
result_11 = UChi_11;\
|
||||
result_12 = UChi_12;\
|
||||
result_20 = timesI(UChi_10);\
|
||||
result_21 = timesI(UChi_11);\
|
||||
result_22 = timesI(UChi_12);\
|
||||
result_30 = timesI(UChi_00);\
|
||||
result_31 = timesI(UChi_01);\
|
||||
result_32 = timesI(UChi_02);
|
||||
|
||||
#define XM_RECON_ACCUM\
|
||||
result_00+= UChi_00;\
|
||||
result_01+= UChi_01;\
|
||||
result_02+= UChi_02;\
|
||||
result_10+= UChi_10;\
|
||||
result_11+= UChi_11;\
|
||||
result_12+= UChi_12;\
|
||||
result_20+= timesI(UChi_10);\
|
||||
result_21+= timesI(UChi_11);\
|
||||
result_22+= timesI(UChi_12);\
|
||||
result_30+= timesI(UChi_00);\
|
||||
result_31+= timesI(UChi_01);\
|
||||
result_32+= timesI(UChi_02);
|
||||
|
||||
#define YP_RECON_ACCUM\
|
||||
result_00+= UChi_00;\
|
||||
result_01+= UChi_01;\
|
||||
result_02+= UChi_02;\
|
||||
result_10+= UChi_10;\
|
||||
result_11+= UChi_11;\
|
||||
result_12+= UChi_12;\
|
||||
result_20+= UChi_10;\
|
||||
result_21+= UChi_11;\
|
||||
result_22+= UChi_12;\
|
||||
result_30-= UChi_00;\
|
||||
result_31-= UChi_01;\
|
||||
result_32-= UChi_02;
|
||||
|
||||
#define YM_RECON_ACCUM\
|
||||
result_00+= UChi_00;\
|
||||
result_01+= UChi_01;\
|
||||
result_02+= UChi_02;\
|
||||
result_10+= UChi_10;\
|
||||
result_11+= UChi_11;\
|
||||
result_12+= UChi_12;\
|
||||
result_20-= UChi_10;\
|
||||
result_21-= UChi_11;\
|
||||
result_22-= UChi_12;\
|
||||
result_30+= UChi_00;\
|
||||
result_31+= UChi_01;\
|
||||
result_32+= UChi_02;
|
||||
|
||||
#define ZP_RECON_ACCUM\
|
||||
result_00+= UChi_00;\
|
||||
result_01+= UChi_01;\
|
||||
result_02+= UChi_02;\
|
||||
result_10+= UChi_10;\
|
||||
result_11+= UChi_11;\
|
||||
result_12+= UChi_12;\
|
||||
result_20-= timesI(UChi_00); \
|
||||
result_21-= timesI(UChi_01); \
|
||||
result_22-= timesI(UChi_02); \
|
||||
result_30+= timesI(UChi_10); \
|
||||
result_31+= timesI(UChi_11); \
|
||||
result_32+= timesI(UChi_12);
|
||||
|
||||
#define ZM_RECON_ACCUM\
|
||||
result_00+= UChi_00;\
|
||||
result_01+= UChi_01;\
|
||||
result_02+= UChi_02;\
|
||||
result_10+= UChi_10;\
|
||||
result_11+= UChi_11;\
|
||||
result_12+= UChi_12;\
|
||||
result_20+= timesI(UChi_00); \
|
||||
result_21+= timesI(UChi_01); \
|
||||
result_22+= timesI(UChi_02); \
|
||||
result_30-= timesI(UChi_10); \
|
||||
result_31-= timesI(UChi_11); \
|
||||
result_32-= timesI(UChi_12);
|
||||
|
||||
#define TP_RECON_ACCUM\
|
||||
result_00+= UChi_00;\
|
||||
result_01+= UChi_01;\
|
||||
result_02+= UChi_02;\
|
||||
result_10+= UChi_10;\
|
||||
result_11+= UChi_11;\
|
||||
result_12+= UChi_12;\
|
||||
result_20+= UChi_00; \
|
||||
result_21+= UChi_01; \
|
||||
result_22+= UChi_02; \
|
||||
result_30+= UChi_10; \
|
||||
result_31+= UChi_11; \
|
||||
result_32+= UChi_12;
|
||||
|
||||
#define TM_RECON_ACCUM\
|
||||
result_00+= UChi_00;\
|
||||
result_01+= UChi_01;\
|
||||
result_02+= UChi_02;\
|
||||
result_10+= UChi_10;\
|
||||
result_11+= UChi_11;\
|
||||
result_12+= UChi_12;\
|
||||
result_20-= UChi_00; \
|
||||
result_21-= UChi_01; \
|
||||
result_22-= UChi_02; \
|
||||
result_30-= UChi_10; \
|
||||
result_31-= UChi_11; \
|
||||
result_32-= UChi_12;
|
||||
|
||||
#define HAND_STENCIL_LEGA(PROJ,PERM,DIR,RECON) \
|
||||
SE=&st_p[DIR+8*ss]; \
|
||||
ptype=st_perm[DIR]; \
|
||||
offset = SE->_offset; \
|
||||
local = SE->_is_local; \
|
||||
perm = SE->_permute; \
|
||||
if ( local ) { \
|
||||
LOAD_CHIMU(PERM); \
|
||||
PROJ; \
|
||||
if ( perm) { \
|
||||
PERMUTE_DIR(PERM); \
|
||||
} \
|
||||
} else { \
|
||||
LOAD_CHI; \
|
||||
} \
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON;
|
||||
|
||||
#define HAND_STENCIL_LEG(PROJ,PERM,DIR,RECON) \
|
||||
SE=&st_p[DIR+8*ss]; \
|
||||
ptype=st_perm[DIR]; \
|
||||
offset = SE->_offset; \
|
||||
local = SE->_is_local; \
|
||||
perm = SE->_permute; \
|
||||
LOAD_CHIMU(PERM); \
|
||||
PROJ; \
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON;
|
||||
|
||||
|
||||
#define HAND_STENCIL_LEG_INT(PROJ,PERM,DIR,RECON) \
|
||||
SE=&st_p[DIR+8*ss]; \
|
||||
ptype=st_perm[DIR]; \
|
||||
offset = SE->_offset; \
|
||||
local = SE->_is_local; \
|
||||
perm = SE->_permute; \
|
||||
if ( local ) { \
|
||||
LOAD_CHIMU; \
|
||||
PROJ; \
|
||||
if ( perm) { \
|
||||
PERMUTE_DIR(PERM); \
|
||||
} \
|
||||
} else if ( st.same_node[DIR] ) { \
|
||||
LOAD_CHI; \
|
||||
} \
|
||||
if (local || st.same_node[DIR] ) { \
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON; \
|
||||
}
|
||||
|
||||
#define HAND_STENCIL_LEG_EXT(PROJ,PERM,DIR,RECON) \
|
||||
SE=st.GetEntry(ptype,DIR,ss); \
|
||||
offset = SE->_offset; \
|
||||
if((!SE->_is_local)&&(!st.same_node[DIR]) ) { \
|
||||
LOAD_CHI; \
|
||||
MULT_2SPIN(DIR); \
|
||||
RECON; \
|
||||
nmu++; \
|
||||
}
|
||||
|
||||
#define HAND_RESULT(ss) \
|
||||
{ \
|
||||
SiteSpinor & ref (out[ss]); \
|
||||
coalescedWrite(ref()(0)(0),result_00); \
|
||||
coalescedWrite(ref()(0)(1),result_01); \
|
||||
coalescedWrite(ref()(0)(2),result_02); \
|
||||
coalescedWrite(ref()(1)(0),result_10); \
|
||||
coalescedWrite(ref()(1)(1),result_11); \
|
||||
coalescedWrite(ref()(1)(2),result_12); \
|
||||
coalescedWrite(ref()(2)(0),result_20); \
|
||||
coalescedWrite(ref()(2)(1),result_21); \
|
||||
coalescedWrite(ref()(2)(2),result_22); \
|
||||
coalescedWrite(ref()(3)(0),result_30); \
|
||||
coalescedWrite(ref()(3)(1),result_31); \
|
||||
coalescedWrite(ref()(3)(2),result_32); \
|
||||
}
|
||||
|
||||
#define HAND_RESULT_EXT(ss) \
|
||||
if (nmu){ \
|
||||
SiteSpinor & ref (out[ss]); \
|
||||
ref()(0)(0)+=result_00; \
|
||||
ref()(0)(1)+=result_01; \
|
||||
ref()(0)(2)+=result_02; \
|
||||
ref()(1)(0)+=result_10; \
|
||||
ref()(1)(1)+=result_11; \
|
||||
ref()(1)(2)+=result_12; \
|
||||
ref()(2)(0)+=result_20; \
|
||||
ref()(2)(1)+=result_21; \
|
||||
ref()(2)(2)+=result_22; \
|
||||
ref()(3)(0)+=result_30; \
|
||||
ref()(3)(1)+=result_31; \
|
||||
ref()(3)(2)+=result_32; \
|
||||
}
|
||||
|
||||
#define HAND_DECLARATIONS(Simd) \
|
||||
Simd result_00; \
|
||||
Simd result_01; \
|
||||
Simd result_02; \
|
||||
Simd result_10; \
|
||||
Simd result_11; \
|
||||
Simd result_12; \
|
||||
Simd result_20; \
|
||||
Simd result_21; \
|
||||
Simd result_22; \
|
||||
Simd result_30; \
|
||||
Simd result_31; \
|
||||
Simd result_32; \
|
||||
Simd Chi_00; \
|
||||
Simd Chi_01; \
|
||||
Simd Chi_02; \
|
||||
Simd Chi_10; \
|
||||
Simd Chi_11; \
|
||||
Simd Chi_12; \
|
||||
Simd UChi_00; \
|
||||
Simd UChi_01; \
|
||||
Simd UChi_02; \
|
||||
Simd UChi_10; \
|
||||
Simd UChi_11; \
|
||||
Simd UChi_12; \
|
||||
Simd U_00; \
|
||||
Simd U_10; \
|
||||
Simd U_20; \
|
||||
Simd U_01; \
|
||||
Simd U_11; \
|
||||
Simd U_21;
|
||||
|
||||
#define ZERO_RESULT \
|
||||
result_00=Zero(); \
|
||||
result_01=Zero(); \
|
||||
result_02=Zero(); \
|
||||
result_10=Zero(); \
|
||||
result_11=Zero(); \
|
||||
result_12=Zero(); \
|
||||
result_20=Zero(); \
|
||||
result_21=Zero(); \
|
||||
result_22=Zero(); \
|
||||
result_30=Zero(); \
|
||||
result_31=Zero(); \
|
||||
result_32=Zero();
|
||||
|
||||
#define Chimu_00 Chi_00
|
||||
#define Chimu_01 Chi_01
|
||||
#define Chimu_02 Chi_02
|
||||
#define Chimu_10 Chi_10
|
||||
#define Chimu_11 Chi_11
|
||||
#define Chimu_12 Chi_12
|
||||
#define Chimu_20 UChi_00
|
||||
#define Chimu_21 UChi_01
|
||||
#define Chimu_22 UChi_02
|
||||
#define Chimu_30 UChi_10
|
||||
#define Chimu_31 UChi_11
|
||||
#define Chimu_32 UChi_12
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
template<class Impl> accelerator_inline void
|
||||
WilsonKernels<Impl>::HandDhopSiteSycl(StencilVector st_perm,StencilEntry *st_p, SiteDoubledGaugeField *U,SiteHalfSpinor *buf,
|
||||
int ss,int sU,const SiteSpinor *in, SiteSpinor *out)
|
||||
{
|
||||
// T==0, Z==1, Y==2, Z==3 expect 1,2,2,2 simd layout etc...
|
||||
typedef typename Simd::scalar_type S;
|
||||
typedef typename Simd::vector_type V;
|
||||
typedef iSinglet<Simd> vCplx;
|
||||
// typedef decltype( coalescedRead( vCplx()()() )) Simt;
|
||||
typedef decltype( coalescedRead( in[0]()(0)(0) )) Simt;
|
||||
|
||||
HAND_DECLARATIONS(Simt);
|
||||
|
||||
int offset,local,perm, ptype;
|
||||
StencilEntry *SE;
|
||||
HAND_STENCIL_LEG(XM_PROJ,3,Xp,XM_RECON);
|
||||
HAND_STENCIL_LEG(YM_PROJ,2,Yp,YM_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG(ZM_PROJ,1,Zp,ZM_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG(TM_PROJ,0,Tp,TM_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG(XP_PROJ,3,Xm,XP_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG(YP_PROJ,2,Ym,YP_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG(ZP_PROJ,1,Zm,ZP_RECON_ACCUM);
|
||||
HAND_STENCIL_LEG(TP_PROJ,0,Tm,TP_RECON_ACCUM);
|
||||
HAND_RESULT(ss);
|
||||
}
|
||||
|
||||
////////////// Wilson ; uses this implementation /////////////////////
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
#undef LOAD_CHIMU
|
||||
#undef LOAD_CHI
|
||||
#undef MULT_2SPIN
|
||||
#undef PERMUTE_DIR
|
||||
#undef XP_PROJ
|
||||
#undef YP_PROJ
|
||||
#undef ZP_PROJ
|
||||
#undef TP_PROJ
|
||||
#undef XM_PROJ
|
||||
#undef YM_PROJ
|
||||
#undef ZM_PROJ
|
||||
#undef TM_PROJ
|
||||
#undef XP_RECON
|
||||
#undef XP_RECON_ACCUM
|
||||
#undef XM_RECON
|
||||
#undef XM_RECON_ACCUM
|
||||
#undef YP_RECON_ACCUM
|
||||
#undef YM_RECON_ACCUM
|
||||
#undef ZP_RECON_ACCUM
|
||||
#undef ZM_RECON_ACCUM
|
||||
#undef TP_RECON_ACCUM
|
||||
#undef TM_RECON_ACCUM
|
||||
#undef ZERO_RESULT
|
||||
#undef Chimu_00
|
||||
#undef Chimu_01
|
||||
#undef Chimu_02
|
||||
#undef Chimu_10
|
||||
#undef Chimu_11
|
||||
#undef Chimu_12
|
||||
#undef Chimu_20
|
||||
#undef Chimu_21
|
||||
#undef Chimu_22
|
||||
#undef Chimu_30
|
||||
#undef Chimu_31
|
||||
#undef Chimu_32
|
||||
#undef HAND_STENCIL_LEG
|
||||
#undef HAND_STENCIL_LEG_INT
|
||||
#undef HAND_STENCIL_LEG_EXT
|
||||
#undef HAND_RESULT
|
||||
#undef HAND_RESULT_INT
|
||||
#undef HAND_RESULT_EXT
|
||||
#undef HAND_DECLARATIONS
|
||||
@ -115,9 +115,9 @@ accelerator_inline void get_stencil(StencilEntry * mem, StencilEntry &chip)
|
||||
// All legs kernels ; comms then compute
|
||||
////////////////////////////////////////////////////////////////////
|
||||
template <class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::GenericDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::GenericDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;
|
||||
typedef decltype(coalescedRead(in[0])) calcSpinor;
|
||||
@ -141,9 +141,9 @@ void WilsonKernels<Impl>::GenericDhopSiteDag(StencilView &st, DoubledGaugeFieldV
|
||||
};
|
||||
|
||||
template <class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::GenericDhopSite(StencilView &st, DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::GenericDhopSite(const StencilView &st, const DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;
|
||||
typedef decltype(coalescedRead(in[0])) calcSpinor;
|
||||
@ -170,9 +170,9 @@ void WilsonKernels<Impl>::GenericDhopSite(StencilView &st, DoubledGaugeFieldView
|
||||
// Interior kernels
|
||||
////////////////////////////////////////////////////////////////////
|
||||
template <class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::GenericDhopSiteDagInt(StencilView &st, DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::GenericDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;
|
||||
typedef decltype(coalescedRead(in[0])) calcSpinor;
|
||||
@ -198,9 +198,9 @@ void WilsonKernels<Impl>::GenericDhopSiteDagInt(StencilView &st, DoubledGaugeFi
|
||||
};
|
||||
|
||||
template <class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::GenericDhopSiteInt(StencilView &st, DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::GenericDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;
|
||||
typedef decltype(coalescedRead(in[0])) calcSpinor;
|
||||
@ -228,9 +228,9 @@ void WilsonKernels<Impl>::GenericDhopSiteInt(StencilView &st, DoubledGaugeField
|
||||
// Exterior kernels
|
||||
////////////////////////////////////////////////////////////////////
|
||||
template <class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::GenericDhopSiteDagExt(StencilView &st, DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::GenericDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;
|
||||
typedef decltype(coalescedRead(in[0])) calcSpinor;
|
||||
@ -259,9 +259,9 @@ void WilsonKernels<Impl>::GenericDhopSiteDagExt(StencilView &st, DoubledGaugeFi
|
||||
};
|
||||
|
||||
template <class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::GenericDhopSiteExt(StencilView &st, DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out)
|
||||
void WilsonKernels<Impl>::GenericDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U,
|
||||
SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, const FermionFieldView &out)
|
||||
{
|
||||
typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;
|
||||
typedef decltype(coalescedRead(in[0])) calcSpinor;
|
||||
@ -291,8 +291,8 @@ void WilsonKernels<Impl>::GenericDhopSiteExt(StencilView &st, DoubledGaugeField
|
||||
|
||||
#define DhopDirMacro(Dir,spProj,spRecon) \
|
||||
template <class Impl> accelerator_inline \
|
||||
void WilsonKernels<Impl>::DhopDir##Dir(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, int sF, \
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out, int dir) \
|
||||
void WilsonKernels<Impl>::DhopDir##Dir(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, int sF, \
|
||||
int sU, const FermionFieldView &in, const FermionFieldView &out, int dir) \
|
||||
{ \
|
||||
typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; \
|
||||
typedef decltype(coalescedRead(in[0])) calcSpinor; \
|
||||
@ -302,8 +302,8 @@ void WilsonKernels<Impl>::GenericDhopSiteExt(StencilView &st, DoubledGaugeField
|
||||
StencilEntry *SE; \
|
||||
int ptype; \
|
||||
const int Nsimd = SiteHalfSpinor::Nsimd(); \
|
||||
const int lane=acceleratorSIMTlane(Nsimd); \
|
||||
\
|
||||
const int lane=acceleratorSIMTlane(Nsimd); \
|
||||
\
|
||||
SE = st.GetEntry(ptype, dir, sF); \
|
||||
GENERIC_DHOPDIR_LEG_BODY(Dir,spProj,spRecon); \
|
||||
coalescedWrite(out[sF], result,lane); \
|
||||
@ -319,8 +319,8 @@ DhopDirMacro(Zm,spProjZm,spReconZm);
|
||||
DhopDirMacro(Tm,spProjTm,spReconTm);
|
||||
|
||||
template <class Impl> accelerator_inline
|
||||
void WilsonKernels<Impl>::DhopDirK( StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, FermionFieldView &out, int dir, int gamma)
|
||||
void WilsonKernels<Impl>::DhopDirK(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, int sF,
|
||||
int sU, const FermionFieldView &in, const FermionFieldView &out, int dir, int gamma)
|
||||
{
|
||||
typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;
|
||||
typedef decltype(coalescedRead(in[0])) calcSpinor;
|
||||
@ -345,8 +345,8 @@ void WilsonKernels<Impl>::DhopDirK( StencilView &st, DoubledGaugeFieldView &U,Si
|
||||
}
|
||||
|
||||
template <class Impl>
|
||||
void WilsonKernels<Impl>::DhopDirAll( StencilImpl &st, DoubledGaugeField &U,SiteHalfSpinor *buf, int Ls,
|
||||
int Nsite, const FermionField &in, std::vector<FermionField> &out)
|
||||
void WilsonKernels<Impl>::DhopDirAll(StencilImpl &st, DoubledGaugeField &U,SiteHalfSpinor *buf, int Ls,
|
||||
int Nsite, const FermionField &in, std::vector<FermionField> &out)
|
||||
{
|
||||
autoView(U_v ,U,AcceleratorRead);
|
||||
autoView(in_v ,in,AcceleratorRead);
|
||||
@ -424,6 +424,20 @@ void WilsonKernels<Impl>::DhopDirKernel( StencilImpl &st, DoubledGaugeField &U,S
|
||||
WilsonKernels<Impl>::A(st_v,U_v,buf,sF,sU,in_v,out_v); \
|
||||
});
|
||||
|
||||
#define KERNEL_CALL_TMP(A) \
|
||||
const uint64_t NN = Nsite*Ls; \
|
||||
auto U_p = & U_v[0]; \
|
||||
auto in_p = & in_v[0]; \
|
||||
auto out_p = & out_v[0]; \
|
||||
auto st_p = st_v._entries_p; \
|
||||
auto st_perm = st_v._permute_type; \
|
||||
accelerator_forNB( ss, NN, Simd::Nsimd(), { \
|
||||
int sF = ss; \
|
||||
int sU = ss/Ls; \
|
||||
WilsonKernels<Impl>::A(st_perm,st_p,U_p,buf,sF,sU,in_p,out_p); \
|
||||
}); \
|
||||
accelerator_barrier();
|
||||
|
||||
#define KERNEL_CALL(A) KERNEL_CALLNB(A); accelerator_barrier();
|
||||
|
||||
#define ASM_CALL(A) \
|
||||
@ -446,7 +460,8 @@ void WilsonKernels<Impl>::DhopKernel(int Opt,StencilImpl &st, DoubledGaugeField
|
||||
if( interior && exterior ) {
|
||||
if (Opt == WilsonKernelsStatic::OptGeneric ) { KERNEL_CALL(GenericDhopSite); return;}
|
||||
#ifndef GRID_CUDA
|
||||
if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALL(HandDhopSite); return;}
|
||||
if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALL_TMP(HandDhopSiteSycl); return; }
|
||||
// if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALL(HandDhopSite); return;}
|
||||
if (Opt == WilsonKernelsStatic::OptInlineAsm ) { ASM_CALL(AsmDhopSite); return;}
|
||||
#endif
|
||||
} else if( interior ) {
|
||||
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -32,6 +32,7 @@ directory
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementationSycl.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -1,51 +0,0 @@
|
||||
/*************************************************************************************
|
||||
|
||||
Grid physics library, www.github.com/paboyle/Grid
|
||||
|
||||
Source file: ./lib/qcd/action/fermion/WilsonKernels.cc
|
||||
|
||||
Copyright (C) 2015, 2020
|
||||
|
||||
Author: Peter Boyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Peter Boyle <peterboyle@Peters-MacBook-Pro-2.local>
|
||||
Author: paboyle <paboyle@ph.ed.ac.uk>
|
||||
Author: Nils Meyer <nils.meyer@ur.de> Regensburg University
|
||||
|
||||
This program is free software; you can redistribute it and/or modify
|
||||
it under the terms of the GNU General Public License as published by
|
||||
the Free Software Foundation; either version 2 of the License, or
|
||||
(at your option) any later version.
|
||||
|
||||
This program is distributed in the hope that it will be useful,
|
||||
but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
|
||||
GNU General Public License for more details.
|
||||
|
||||
You should have received a copy of the GNU General Public License along
|
||||
with this program; if not, write to the Free Software Foundation, Inc.,
|
||||
51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
|
||||
|
||||
See the full license in the file "LICENSE" in the top level distribution
|
||||
directory
|
||||
*************************************************************************************/
|
||||
/* END LEGAL */
|
||||
#include <Grid/qcd/action/fermion/FermionCore.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h>
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h>
|
||||
|
||||
#ifndef AVX512
|
||||
#ifndef QPX
|
||||
#ifndef A64FX
|
||||
#ifndef A64FXFIXEDSIZE
|
||||
#include <Grid/qcd/action/fermion/implementation/WilsonKernelsAsmImplementation.h>
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
#endif
|
||||
|
||||
NAMESPACE_BEGIN(Grid);
|
||||
|
||||
#include "impl.h"
|
||||
template class WilsonKernels<IMPLEMENTATION>;
|
||||
|
||||
NAMESPACE_END(Grid);
|
||||
@ -0,0 +1 @@
|
||||
../WilsonKernelsInstantiation.cc.master
|
||||
@ -60,11 +60,25 @@ template<class pair>
|
||||
class GpuComplex {
|
||||
public:
|
||||
pair z;
|
||||
typedef decltype(z.x) real;
|
||||
typedef decltype(z.x) Real;
|
||||
public:
|
||||
accelerator_inline GpuComplex() = default;
|
||||
accelerator_inline GpuComplex(real re,real im) { z.x=re; z.y=im; };
|
||||
accelerator_inline GpuComplex(Real re,Real im) { z.x=re; z.y=im; };
|
||||
accelerator_inline GpuComplex(const GpuComplex &zz) { z = zz.z;};
|
||||
accelerator_inline Real real(void) const { return z.x; };
|
||||
accelerator_inline Real imag(void) const { return z.y; };
|
||||
accelerator_inline GpuComplex &operator*=(const GpuComplex &r) {
|
||||
*this = (*this) * r;
|
||||
return *this;
|
||||
}
|
||||
accelerator_inline GpuComplex &operator+=(const GpuComplex &r) {
|
||||
*this = (*this) + r;
|
||||
return *this;
|
||||
}
|
||||
accelerator_inline GpuComplex &operator-=(const GpuComplex &r) {
|
||||
*this = (*this) - r;
|
||||
return *this;
|
||||
}
|
||||
friend accelerator_inline GpuComplex operator+(const GpuComplex &lhs,const GpuComplex &rhs) {
|
||||
GpuComplex r ;
|
||||
r.z.x = lhs.z.x + rhs.z.x;
|
||||
@ -157,6 +171,11 @@ typedef GpuVector<NSIMD_RealD, double > GpuVectorRD;
|
||||
typedef GpuVector<NSIMD_ComplexD, GpuComplexD > GpuVectorCD;
|
||||
typedef GpuVector<NSIMD_Integer, Integer > GpuVectorI;
|
||||
|
||||
accelerator_inline GpuComplexF timesI(const GpuComplexF &r) { return(GpuComplexF(-r.imag(),r.real()));}
|
||||
accelerator_inline GpuComplexD timesI(const GpuComplexD &r) { return(GpuComplexD(-r.imag(),r.real()));}
|
||||
accelerator_inline GpuComplexF timesMinusI(const GpuComplexF &r){ return(GpuComplexF(r.imag(),-r.real()));}
|
||||
accelerator_inline GpuComplexD timesMinusI(const GpuComplexD &r){ return(GpuComplexD(r.imag(),-r.real()));}
|
||||
|
||||
accelerator_inline float half2float(half h)
|
||||
{
|
||||
float f;
|
||||
|
||||
@ -148,10 +148,14 @@ accelerator_inline void sub (ComplexF * __restrict__ y,const ComplexF * __restri
|
||||
accelerator_inline void add (ComplexF * __restrict__ y,const ComplexF * __restrict__ l,const ComplexF *__restrict__ r){ *y = (*l) + (*r); }
|
||||
|
||||
//conjugate already supported for complex
|
||||
accelerator_inline ComplexF timesI(const ComplexF &r) { return(r*ComplexF(0.0,1.0));}
|
||||
accelerator_inline ComplexD timesI(const ComplexD &r) { return(r*ComplexD(0.0,1.0));}
|
||||
accelerator_inline ComplexF timesMinusI(const ComplexF &r){ return(r*ComplexF(0.0,-1.0));}
|
||||
accelerator_inline ComplexD timesMinusI(const ComplexD &r){ return(r*ComplexD(0.0,-1.0));}
|
||||
accelerator_inline ComplexF timesI(const ComplexF &r) { return(ComplexF(-r.imag(),r.real()));}
|
||||
accelerator_inline ComplexD timesI(const ComplexD &r) { return(ComplexD(-r.imag(),r.real()));}
|
||||
accelerator_inline ComplexF timesMinusI(const ComplexF &r){ return(ComplexF(r.imag(),-r.real()));}
|
||||
accelerator_inline ComplexD timesMinusI(const ComplexD &r){ return(ComplexD(r.imag(),-r.real()));}
|
||||
//accelerator_inline ComplexF timesI(const ComplexF &r) { return(r*ComplexF(0.0,1.0));}
|
||||
//accelerator_inline ComplexD timesI(const ComplexD &r) { return(r*ComplexD(0.0,1.0));}
|
||||
//accelerator_inline ComplexF timesMinusI(const ComplexF &r){ return(r*ComplexF(0.0,-1.0));}
|
||||
//accelerator_inline ComplexD timesMinusI(const ComplexD &r){ return(r*ComplexD(0.0,-1.0));}
|
||||
|
||||
// define projections to real and imaginay parts
|
||||
accelerator_inline ComplexF projReal(const ComplexF &r){return( ComplexF(r.real(), 0.0));}
|
||||
|
||||
@ -7,20 +7,20 @@ template<class vobj>
|
||||
class SimpleCompressor {
|
||||
public:
|
||||
void Point(int) {};
|
||||
accelerator_inline int CommDatumSize(void) { return sizeof(vobj); }
|
||||
accelerator_inline bool DecompressionStep(void) { return false; }
|
||||
template<class cobj> accelerator_inline void Compress(cobj *buf,int o,const cobj &in) { buf[o]=in; }
|
||||
accelerator_inline void Exchange(vobj *mp,vobj *vp0,vobj *vp1,Integer type,Integer o){
|
||||
accelerator_inline int CommDatumSize(void) const { return sizeof(vobj); }
|
||||
accelerator_inline bool DecompressionStep(void) const { return false; }
|
||||
template<class cobj> accelerator_inline void Compress(cobj *buf,int o,const cobj &in) const { buf[o]=in; }
|
||||
accelerator_inline void Exchange(vobj *mp,vobj *vp0,vobj *vp1,Integer type,Integer o) const {
|
||||
exchange(mp[2*o],mp[2*o+1],vp0[o],vp1[o],type);
|
||||
}
|
||||
accelerator_inline void Decompress(vobj *out,vobj *in, int o){ assert(0); }
|
||||
accelerator_inline void Decompress(vobj *out,vobj *in, int o) const { assert(0); }
|
||||
accelerator_inline void CompressExchange(vobj *out0,vobj *out1,const vobj *in,
|
||||
int j,int k, int m,int type){
|
||||
int j,int k, int m,int type) const {
|
||||
exchange(out0[j],out1[j],in[k],in[m],type);
|
||||
}
|
||||
// For cshift. Cshift should drop compressor coupling altogether
|
||||
// because I had to decouple the code from the Stencil anyway
|
||||
accelerator_inline vobj operator() (const vobj &arg) {
|
||||
accelerator_inline vobj operator() (const vobj &arg) const {
|
||||
return arg;
|
||||
}
|
||||
};
|
||||
|
||||
@ -147,16 +147,16 @@ class CartesianStencilAccelerator {
|
||||
cobj* u_recv_buf_p;
|
||||
cobj* u_send_buf_p;
|
||||
|
||||
accelerator_inline cobj *CommBuf(void) { return u_recv_buf_p; }
|
||||
accelerator_inline cobj *CommBuf(void) const { return u_recv_buf_p; }
|
||||
|
||||
accelerator_inline int GetNodeLocal(int osite,int point) {
|
||||
accelerator_inline int GetNodeLocal(int osite,int point) const {
|
||||
return this->_entries_p[point+this->_npoints*osite]._is_local;
|
||||
}
|
||||
accelerator_inline StencilEntry * GetEntry(int &ptype,int point,int osite) {
|
||||
accelerator_inline StencilEntry * GetEntry(int &ptype,int point,int osite) const {
|
||||
ptype = this->_permute_type[point]; return & this->_entries_p[point+this->_npoints*osite];
|
||||
}
|
||||
|
||||
accelerator_inline uint64_t GetInfo(int &ptype,int &local,int &perm,int point,int ent,uint64_t base) {
|
||||
accelerator_inline uint64_t GetInfo(int &ptype,int &local,int &perm,int point,int ent,uint64_t base) const {
|
||||
uint64_t cbase = (uint64_t)&u_recv_buf_p[0];
|
||||
local = this->_entries_p[ent]._is_local;
|
||||
perm = this->_entries_p[ent]._permute;
|
||||
@ -168,14 +168,14 @@ class CartesianStencilAccelerator {
|
||||
}
|
||||
}
|
||||
|
||||
accelerator_inline uint64_t GetPFInfo(int ent,uint64_t base) {
|
||||
accelerator_inline uint64_t GetPFInfo(int ent,uint64_t base) const {
|
||||
uint64_t cbase = (uint64_t)&u_recv_buf_p[0];
|
||||
int local = this->_entries_p[ent]._is_local;
|
||||
if (local) return base + this->_entries_p[ent]._byte_offset;
|
||||
else return cbase + this->_entries_p[ent]._byte_offset;
|
||||
}
|
||||
|
||||
accelerator_inline void iCoorFromIindex(Coordinate &coor,int lane)
|
||||
accelerator_inline void iCoorFromIindex(Coordinate &coor,int lane) const
|
||||
{
|
||||
Lexicographic::CoorFromIndex(coor,lane,this->_simd_layout);
|
||||
}
|
||||
|
||||
@ -64,6 +64,68 @@ void coalescedWriteNonTemporal(vobj & __restrict__ vec,const vobj & __restrict__
|
||||
}
|
||||
#else
|
||||
|
||||
|
||||
#if 0
|
||||
// Use the scalar as our own complex on GPU
|
||||
template<class vsimd,IfSimd<vsimd> = 0> accelerator_inline
|
||||
typename vsimd::scalar_type
|
||||
coalescedRead(const vsimd & __restrict__ vec,int lane=acceleratorSIMTlane(vsimd::Nsimd()))
|
||||
{
|
||||
typedef typename vsimd::scalar_type S;
|
||||
S * __restrict__ p=(S *)&vec;
|
||||
return p[lane];
|
||||
}
|
||||
template<int ptype,class vsimd,IfSimd<vsimd> = 0> accelerator_inline
|
||||
typename vsimd::scalar_type
|
||||
coalescedReadPermute(const vsimd & __restrict__ vec,int doperm,int lane=acceleratorSIMTlane(vsimd::Nsimd()))
|
||||
{
|
||||
typedef typename vsimd::scalar_type S;
|
||||
|
||||
S * __restrict__ p=(S *)&vec;
|
||||
int mask = vsimd::Nsimd() >> (ptype + 1);
|
||||
int plane= doperm ? lane ^ mask : lane;
|
||||
return p[plane];
|
||||
}
|
||||
template<class vsimd,IfSimd<vsimd> = 0> accelerator_inline
|
||||
void coalescedWrite(vsimd & __restrict__ vec,
|
||||
const typename vsimd::scalar_type & __restrict__ extracted,
|
||||
int lane=acceleratorSIMTlane(vsimd::Nsimd()))
|
||||
{
|
||||
typedef typename vsimd::scalar_type S;
|
||||
S * __restrict__ p=(S *)&vec;
|
||||
p[lane]=extracted;
|
||||
}
|
||||
#else
|
||||
template<class vsimd,IfSimd<vsimd> = 0> accelerator_inline
|
||||
typename vsimd::vector_type::datum
|
||||
coalescedRead(const vsimd & __restrict__ vec,int lane=acceleratorSIMTlane(vsimd::Nsimd()))
|
||||
{
|
||||
typedef typename vsimd::vector_type::datum S;
|
||||
S * __restrict__ p=(S *)&vec;
|
||||
return p[lane];
|
||||
}
|
||||
template<int ptype,class vsimd,IfSimd<vsimd> = 0> accelerator_inline
|
||||
typename vsimd::vector_type::datum
|
||||
coalescedReadPermute(const vsimd & __restrict__ vec,int doperm,int lane=acceleratorSIMTlane(vsimd::Nsimd()))
|
||||
{
|
||||
typedef typename vsimd::vector_type::datum S;
|
||||
|
||||
S * __restrict__ p=(S *)&vec;
|
||||
int mask = vsimd::Nsimd() >> (ptype + 1);
|
||||
int plane= doperm ? lane ^ mask : lane;
|
||||
return p[plane];
|
||||
}
|
||||
template<class vsimd,IfSimd<vsimd> = 0> accelerator_inline
|
||||
void coalescedWrite(vsimd & __restrict__ vec,
|
||||
const typename vsimd::vector_type::datum & __restrict__ extracted,
|
||||
int lane=acceleratorSIMTlane(vsimd::Nsimd()))
|
||||
{
|
||||
typedef typename vsimd::vector_type::datum S;
|
||||
S * __restrict__ p=(S *)&vec;
|
||||
p[lane]=extracted;
|
||||
}
|
||||
#endif
|
||||
|
||||
//////////////////////////////////////////
|
||||
// Extract and insert slices on the GPU
|
||||
//////////////////////////////////////////
|
||||
|
||||
@ -221,7 +221,7 @@ accelerator_inline int acceleratorSIMTlane(int Nsimd) {
|
||||
cl::sycl::range<3> global{unum1,unum2,nsimd}; \
|
||||
cgh.parallel_for<class dslash>( \
|
||||
cl::sycl::nd_range<3>(global,local), \
|
||||
[=] (cl::sycl::nd_item<3> item) mutable { \
|
||||
[=] (cl::sycl::nd_item<3> item) /*mutable*/ { \
|
||||
auto iter1 = item.get_global_id(0); \
|
||||
auto iter2 = item.get_global_id(1); \
|
||||
auto lane = item.get_global_id(2); \
|
||||
|
||||
Reference in New Issue
Block a user