mirror of
				https://github.com/paboyle/Grid.git
				synced 2025-10-27 10:09:33 +00:00 
			
		
		
		
	Compare commits
	
		
			5 Commits
		
	
	
		
			6d7219b59d
			...
			feature/sy
		
	
	| Author | SHA1 | Date | |
|---|---|---|---|
|  | 229709a980 | ||
|  | 9295eeadfe | ||
|  | 36f471e333 | ||
|  | ca4eadd4ab | ||
|  | d954595922 | 
| @@ -54,7 +54,7 @@ Version.h: version-cache | |||||||
| include Make.inc | include Make.inc | ||||||
| include Eigen.inc | include Eigen.inc | ||||||
|  |  | ||||||
| extra_sources+=$(ZWILS_FERMION_FILES) | #extra_sources+=$(ZWILS_FERMION_FILES) | ||||||
| extra_sources+=$(WILS_FERMION_FILES) | extra_sources+=$(WILS_FERMION_FILES) | ||||||
| extra_sources+=$(STAG_FERMION_FILES) | extra_sources+=$(STAG_FERMION_FILES) | ||||||
| if BUILD_GPARITY | if BUILD_GPARITY | ||||||
|   | |||||||
| @@ -36,7 +36,7 @@ static const int CbBlack=1; | |||||||
| static const int Even   =CbRed; | static const int Even   =CbRed; | ||||||
| static const int Odd    =CbBlack; | 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(); |   int nd=rdim.size(); | ||||||
|   Coordinate coor(nd); |   Coordinate coor(nd); | ||||||
|   | |||||||
| @@ -67,8 +67,13 @@ public: | |||||||
|   accelerator_inline const vobj & operator()(size_t i) const { return this->_odata[i]; } |   accelerator_inline const vobj & operator()(size_t i) const { return this->_odata[i]; } | ||||||
| #endif | #endif | ||||||
|  |  | ||||||
|   accelerator_inline const vobj & operator[](size_t i) const { return this->_odata[i]; }; | #if 1 | ||||||
|   accelerator_inline vobj       & operator[](size_t i)       { return this->_odata[i]; }; |   //  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 begin(void) const { return 0;}; | ||||||
|   accelerator_inline uint64_t end(void)   const { return this->_odata_size; }; |   accelerator_inline uint64_t end(void)   const { return this->_odata_size; }; | ||||||
|   | |||||||
| @@ -88,7 +88,7 @@ public: | |||||||
| 					  const _Spinor &chi,  | 					  const _Spinor &chi,  | ||||||
| 					  int mu,  | 					  int mu,  | ||||||
| 					  StencilEntry *SE, | 					  StencilEntry *SE, | ||||||
| 					  StencilView &St)  | 					  const StencilView &St)  | ||||||
|   { |   { | ||||||
|     int direction = St._directions[mu]; |     int direction = St._directions[mu]; | ||||||
|     int distance  = St._distances[mu]; |     int distance  = St._distances[mu]; | ||||||
|   | |||||||
| @@ -56,8 +56,12 @@ template<class Impl> class StaggeredKernels : public FermionOperator<Impl> , pub | |||||||
| 		 DoubledGaugeField &U, | 		 DoubledGaugeField &U, | ||||||
| 		 const FermionField &in, FermionField &out, int dag, int interior,int exterior); | 		 const FermionField &in, FermionField &out, int dag, int interior,int exterior); | ||||||
|    |    | ||||||
|   void DhopDirKernel(StencilImpl &st, DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, SiteSpinor * buf, |   void DhopDirKernel(StencilImpl &st, | ||||||
| 		     int sF, int sU, const FermionFieldView &in, FermionFieldView &out, int dir,int disp); | 		     const DoubledGaugeFieldView &U, | ||||||
|  | 		     const DoubledGaugeFieldView &UUU, SiteSpinor * buf, | ||||||
|  | 		     int sF, int sU, | ||||||
|  | 		     const FermionFieldView &in, | ||||||
|  | 		     const FermionFieldView &out, int dir,int disp); | ||||||
|  protected:     |  protected:     | ||||||
|  |  | ||||||
|    /////////////////////////////////////////////////////////////////////////////////////// |    /////////////////////////////////////////////////////////////////////////////////////// | ||||||
| @@ -65,53 +69,67 @@ template<class Impl> class StaggeredKernels : public FermionOperator<Impl> , pub | |||||||
|    /////////////////////////////////////////////////////////////////////////////////////// |    /////////////////////////////////////////////////////////////////////////////////////// | ||||||
|    template<int Naik>  |    template<int Naik>  | ||||||
|    static accelerator_inline |    static accelerator_inline | ||||||
|    void DhopSiteGeneric(StencilView &st,  |    void DhopSiteGeneric(const StencilView &st,  | ||||||
| 			DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,  | 			const DoubledGaugeFieldView &U, | ||||||
|  | 			const DoubledGaugeFieldView &UUU,  | ||||||
| 			SiteSpinor * buf, int LLs, int sU,  | 			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 |    template<int Naik> static accelerator_inline | ||||||
|    void DhopSiteGenericInt(StencilView &st,  |    void DhopSiteGenericInt(const StencilView &st,  | ||||||
| 			   DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU,  | 			   const DoubledGaugeFieldView &U, | ||||||
|  | 			   const DoubledGaugeFieldView &UUU,  | ||||||
| 			   SiteSpinor * buf, int LLs, int sU,  | 			   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 |    template<int Naik> static accelerator_inline | ||||||
|    void DhopSiteGenericExt(StencilView &st,  |    void DhopSiteGenericExt(const StencilView &st,  | ||||||
| 			   DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, | 			   const DoubledGaugeFieldView &U, | ||||||
| 			   SiteSpinor * buf, int LLs, int sU,  | 			   const DoubledGaugeFieldView &UUU, | ||||||
| 			   const FermionFieldView &in, FermionFieldView &out,int dag); |  			   SiteSpinor * buf, int LLs, int sU,  | ||||||
|  | 			   const FermionFieldView &in, | ||||||
|  | 			   const FermionFieldView &out,int dag); | ||||||
|  |  | ||||||
|    /////////////////////////////////////////////////////////////////////////////////////// |    /////////////////////////////////////////////////////////////////////////////////////// | ||||||
|    // Nc=3 specific kernels |    // Nc=3 specific kernels | ||||||
|    /////////////////////////////////////////////////////////////////////////////////////// |    /////////////////////////////////////////////////////////////////////////////////////// | ||||||
|     |     | ||||||
|    template<int Naik> static accelerator_inline |    template<int Naik> static accelerator_inline | ||||||
|    void DhopSiteHand(StencilView &st,  |    void DhopSiteHand(const StencilView &st,  | ||||||
| 		     DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,  | 		     const DoubledGaugeFieldView &U, | ||||||
| 		     SiteSpinor * buf, int LLs, int sU,  | 		     const DoubledGaugeFieldView &UUU,  | ||||||
| 		     const FermionFieldView &in, FermionFieldView &out,int dag); |  		     SiteSpinor * buf, int LLs, int sU,  | ||||||
|  | 		     const FermionFieldView &in, | ||||||
|  | 		     const FermionFieldView &out,int dag); | ||||||
|     |     | ||||||
|    template<int Naik> static accelerator_inline |    template<int Naik> static accelerator_inline | ||||||
|    void DhopSiteHandInt(StencilView &st,  |    void DhopSiteHandInt(const StencilView &st,  | ||||||
| 			DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,  | 			const DoubledGaugeFieldView &U, | ||||||
| 			SiteSpinor * buf, int LLs, int sU,  | 			const DoubledGaugeFieldView &UUU,  | ||||||
| 			const FermionFieldView &in, FermionFieldView &out,int dag); |  			SiteSpinor * buf, int LLs, int sU,  | ||||||
|  | 			const FermionFieldView &in, | ||||||
|  | 			const FermionFieldView &out,int dag); | ||||||
|     |     | ||||||
|    template<int Naik> static accelerator_inline |    template<int Naik> static accelerator_inline | ||||||
|    void DhopSiteHandExt(StencilView &st,  |    void DhopSiteHandExt(const StencilView &st,  | ||||||
| 			DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,  | 			const DoubledGaugeFieldView &U, | ||||||
| 			SiteSpinor * buf, int LLs, int sU,  | 			const DoubledGaugeFieldView &UUU,  | ||||||
| 			const FermionFieldView &in, FermionFieldView &out,int dag); |  			SiteSpinor * buf, int LLs, int sU,  | ||||||
|  | 			const FermionFieldView &in, | ||||||
|  | 			const FermionFieldView &out,int dag); | ||||||
|  |  | ||||||
|    /////////////////////////////////////////////////////////////////////////////////////// |    /////////////////////////////////////////////////////////////////////////////////////// | ||||||
|    // Asm Nc=3 specific kernels |    // Asm Nc=3 specific kernels | ||||||
|    /////////////////////////////////////////////////////////////////////////////////////// |    /////////////////////////////////////////////////////////////////////////////////////// | ||||||
|     |     | ||||||
|    void DhopSiteAsm(StencilView &st,  |    void DhopSiteAsm(const StencilView &st,  | ||||||
| 		    DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU,  | 		    const DoubledGaugeFieldView &U, | ||||||
| 		    SiteSpinor * buf, int LLs, int sU,  | 		    const DoubledGaugeFieldView &UUU,  | ||||||
| 		    const FermionFieldView &in, FermionFieldView &out,int dag); |  		    SiteSpinor * buf, int LLs, int sU,  | ||||||
|  | 		    const FermionFieldView &in, | ||||||
|  | 		    const FermionFieldView &out,int dag); | ||||||
|    |    | ||||||
| public: | public: | ||||||
|  |  | ||||||
|   | |||||||
| @@ -61,7 +61,7 @@ public: | |||||||
|   typedef typename SiteHalfSpinor::vector_type     vComplexHigh; |   typedef typename SiteHalfSpinor::vector_type     vComplexHigh; | ||||||
|   constexpr static int Nw=sizeof(SiteHalfSpinor)/sizeof(vComplexHigh); |   constexpr static int Nw=sizeof(SiteHalfSpinor)/sizeof(vComplexHigh); | ||||||
|  |  | ||||||
|   accelerator_inline int CommDatumSize(void) { |   accelerator_inline int CommDatumSize(void) const { | ||||||
|     return sizeof(SiteHalfCommSpinor); |     return sizeof(SiteHalfCommSpinor); | ||||||
|   } |   } | ||||||
|  |  | ||||||
| @@ -69,7 +69,7 @@ public: | |||||||
|   /* Compress includes precision change if mpi data is not same */ |   /* Compress includes precision change if mpi data is not same */ | ||||||
|   /*****************************************************/ |   /*****************************************************/ | ||||||
|   template<class _SiteHalfSpinor, class _SiteSpinor> |   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; |     _SiteHalfSpinor tmp; | ||||||
|     projector::Proj(tmp,in,mu,dag); |     projector::Proj(tmp,in,mu,dag); | ||||||
|     vstream(buf[o],tmp); |     vstream(buf[o],tmp); | ||||||
| @@ -81,7 +81,7 @@ public: | |||||||
|   accelerator_inline void Exchange(SiteHalfSpinor *mp, |   accelerator_inline void Exchange(SiteHalfSpinor *mp, | ||||||
| 				   const SiteHalfSpinor * __restrict__ vp0, | 				   const SiteHalfSpinor * __restrict__ vp0, | ||||||
| 				   const SiteHalfSpinor * __restrict__ vp1, | 				   const SiteHalfSpinor * __restrict__ vp1, | ||||||
| 				   Integer type,Integer o){ | 				   Integer type,Integer o) const { | ||||||
|     SiteHalfSpinor tmp1; |     SiteHalfSpinor tmp1; | ||||||
|     SiteHalfSpinor tmp2; |     SiteHalfSpinor tmp2; | ||||||
|     exchange(tmp1,tmp2,vp0[o],vp1[o],type); |     exchange(tmp1,tmp2,vp0[o],vp1[o],type); | ||||||
| @@ -93,7 +93,7 @@ public: | |||||||
|   /* Have a decompression step if mpi data is not same */ |   /* Have a decompression step if mpi data is not same */ | ||||||
|   /*****************************************************/ |   /*****************************************************/ | ||||||
|   accelerator_inline void Decompress(SiteHalfSpinor * __restrict__ out, |   accelerator_inline void Decompress(SiteHalfSpinor * __restrict__ out, | ||||||
| 				     SiteHalfSpinor * __restrict__ in, Integer o) {     | 				     SiteHalfSpinor * __restrict__ in, Integer o) const {     | ||||||
|     assert(0); |     assert(0); | ||||||
|   } |   } | ||||||
|  |  | ||||||
| @@ -103,7 +103,7 @@ public: | |||||||
|   accelerator_inline void CompressExchange(SiteHalfSpinor * __restrict__ out0, |   accelerator_inline void CompressExchange(SiteHalfSpinor * __restrict__ out0, | ||||||
| 					   SiteHalfSpinor * __restrict__ out1, | 					   SiteHalfSpinor * __restrict__ out1, | ||||||
| 					   const SiteSpinor * __restrict__ in, | 					   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 temp1, temp2; | ||||||
|     SiteHalfSpinor temp3, temp4; |     SiteHalfSpinor temp3, temp4; | ||||||
| @@ -117,7 +117,7 @@ public: | |||||||
|   /*****************************************************/ |   /*****************************************************/ | ||||||
|   /* Pass the info to the stencil */ |   /* 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; |   typedef typename SiteHalfSpinor::vector_type     vComplexHigh; | ||||||
|   constexpr static int Nw=sizeof(SiteHalfSpinor)/sizeof(vComplexHigh); |   constexpr static int Nw=sizeof(SiteHalfSpinor)/sizeof(vComplexHigh); | ||||||
|  |  | ||||||
|   accelerator_inline int CommDatumSize(void) { |   accelerator_inline int CommDatumSize(void) const { | ||||||
|     return sizeof(SiteHalfCommSpinor); |     return sizeof(SiteHalfCommSpinor); | ||||||
|   } |   } | ||||||
|  |  | ||||||
| @@ -150,7 +150,7 @@ public: | |||||||
|   /* Compress includes precision change if mpi data is not same */ |   /* Compress includes precision change if mpi data is not same */ | ||||||
|   /*****************************************************/ |   /*****************************************************/ | ||||||
|   template<class _SiteHalfSpinor, class _SiteSpinor> |   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; |     _SiteHalfSpinor hsp; | ||||||
|     SiteHalfCommSpinor *hbuf = (SiteHalfCommSpinor *)buf; |     SiteHalfCommSpinor *hbuf = (SiteHalfCommSpinor *)buf; | ||||||
|     projector::Proj(hsp,in,mu,dag); |     projector::Proj(hsp,in,mu,dag); | ||||||
| @@ -163,7 +163,7 @@ public: | |||||||
|   accelerator_inline void Exchange(SiteHalfSpinor *mp, |   accelerator_inline void Exchange(SiteHalfSpinor *mp, | ||||||
|                        SiteHalfSpinor *vp0, |                        SiteHalfSpinor *vp0, | ||||||
|                        SiteHalfSpinor *vp1, |                        SiteHalfSpinor *vp1, | ||||||
| 		       Integer type,Integer o){ | 		       Integer type,Integer o) const { | ||||||
|     SiteHalfSpinor vt0,vt1; |     SiteHalfSpinor vt0,vt1; | ||||||
|     SiteHalfCommSpinor *vpp0 = (SiteHalfCommSpinor *)vp0; |     SiteHalfCommSpinor *vpp0 = (SiteHalfCommSpinor *)vp0; | ||||||
|     SiteHalfCommSpinor *vpp1 = (SiteHalfCommSpinor *)vp1; |     SiteHalfCommSpinor *vpp1 = (SiteHalfCommSpinor *)vp1; | ||||||
| @@ -175,7 +175,7 @@ public: | |||||||
|   /*****************************************************/ |   /*****************************************************/ | ||||||
|   /* Have a decompression step if mpi data is not same */ |   /* 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; |     SiteHalfCommSpinor *hin=(SiteHalfCommSpinor *)in; | ||||||
|     precisionChange((vComplexHigh *)&out[o],(vComplexLow *)&hin[o],Nw); |     precisionChange((vComplexHigh *)&out[o],(vComplexLow *)&hin[o],Nw); | ||||||
|   } |   } | ||||||
| @@ -186,7 +186,7 @@ public: | |||||||
|   accelerator_inline void CompressExchange(SiteHalfSpinor *out0, |   accelerator_inline void CompressExchange(SiteHalfSpinor *out0, | ||||||
| 			       SiteHalfSpinor *out1, | 			       SiteHalfSpinor *out1, | ||||||
| 			       const SiteSpinor *in, | 			       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; |     SiteHalfSpinor temp1, temp2,temp3,temp4; | ||||||
|     SiteHalfCommSpinor *hout0 = (SiteHalfCommSpinor *)out0; |     SiteHalfCommSpinor *hout0 = (SiteHalfCommSpinor *)out0; | ||||||
|     SiteHalfCommSpinor *hout1 = (SiteHalfCommSpinor *)out1; |     SiteHalfCommSpinor *hout1 = (SiteHalfCommSpinor *)out1; | ||||||
| @@ -200,7 +200,7 @@ public: | |||||||
|   /*****************************************************/ |   /*****************************************************/ | ||||||
|   /* Pass the info to the stencil */ |   /* 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, | 					  const _Spinor &chi, | ||||||
| 					  int mu, | 					  int mu, | ||||||
| 					  StencilEntry *SE, | 					  StencilEntry *SE, | ||||||
| 					  StencilView &St)  | 					  const StencilView &St)  | ||||||
|   { |   { | ||||||
|     multLink(phi,U,chi,mu); |     multLink(phi,U,chi,mu); | ||||||
|   } |   } | ||||||
|   | |||||||
| @@ -49,6 +49,7 @@ public: | |||||||
|  |  | ||||||
|   INHERIT_IMPL_TYPES(Impl); |   INHERIT_IMPL_TYPES(Impl); | ||||||
|   typedef FermionOperator<Impl> Base; |   typedef FermionOperator<Impl> Base; | ||||||
|  |   typedef AcceleratorVector<int,STENCIL_MAX> StencilVector; | ||||||
|     |     | ||||||
| public: | public: | ||||||
|  |  | ||||||
| @@ -68,73 +69,87 @@ public: | |||||||
|  |  | ||||||
| private: | private: | ||||||
|  |  | ||||||
|   static accelerator_inline void DhopDirK(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor * buf, |   static accelerator_inline void DhopDirK(const StencilView &st, const DoubledGaugeFieldView &U, | ||||||
| 				   int sF, int sU, const FermionFieldView &in, FermionFieldView &out, int dirdisp, int gamma); | 					  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 DhopDirXp(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU, | ||||||
|   static accelerator_inline void DhopDirYp(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp); | 					   const FermionFieldView &in, const 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 DhopDirYp(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU, | ||||||
|   static accelerator_inline void DhopDirTp(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp); | 					   const FermionFieldView &in, const 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 DhopDirZp(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU, | ||||||
|   static accelerator_inline void DhopDirYm(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp); | 					   const FermionFieldView &in, const 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 DhopDirTp(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU, | ||||||
|   static accelerator_inline void DhopDirTm(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf,int sF,int sU,const FermionFieldView &in,FermionFieldView &out,int dirdisp); | 					   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 |   // Specialised variants | ||||||
|   static accelerator void GenericDhopSite(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void GenericDhopSite(const StencilView &st, | ||||||
| 					  int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 					  const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
|  | 					  int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|         |         | ||||||
|   static accelerator void GenericDhopSiteDag(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void GenericDhopSiteDag(const StencilView &st, const  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 						    int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 					     int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|     |     | ||||||
|   static accelerator void GenericDhopSiteInt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void GenericDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 						    int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 					     int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|         |         | ||||||
|   static accelerator void GenericDhopSiteDagInt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void GenericDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 						int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 						int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|     |     | ||||||
|   static accelerator void GenericDhopSiteExt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void GenericDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 					     int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 					     int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|         |         | ||||||
|   static accelerator void GenericDhopSiteDagExt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void GenericDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 						       int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 						int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|  |  | ||||||
|   static void AsmDhopSite(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | // Keep Hand unrolled  | ||||||
| 			  int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,FermionFieldView &out); |   static accelerator void HandDhopSiteSycl(StencilVector st_perm, StencilEntry *st_p,  SiteDoubledGaugeField *U, SiteHalfSpinor * buf, | ||||||
|  | 					   int sF, int sU, const SiteSpinor *in, SiteSpinor *out); | ||||||
|  |  | ||||||
|   static void AsmDhopSiteDag(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void HandDhopSite(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 			     int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, FermionFieldView &out); | 				       int sF, int sU, const FermionFieldView &in,const FermionFieldView &out); | ||||||
|     |     | ||||||
|   static void AsmDhopSiteInt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void HandDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 			     int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,FermionFieldView &out); | 					  int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|     |     | ||||||
|   static void AsmDhopSiteDagInt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void HandDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 				int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, FermionFieldView &out); | 					  int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|    |    | ||||||
|   static void AsmDhopSiteExt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void HandDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 			     int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,FermionFieldView &out); | 					     int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|    |    | ||||||
|   static void AsmDhopSiteDagExt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static accelerator void HandDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 				int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, FermionFieldView &out); | 					  int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|     |     | ||||||
| // Keep Hand unrolled temporarily   |   static accelerator void HandDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
|   static accelerator void HandDhopSite(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | 					     int sF, int sU, const FermionFieldView &in, const FermionFieldView &out); | ||||||
| 				       int sF, int sU, const FermionFieldView &in, 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 HandDhopSiteDag(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static void AsmDhopSiteDag(const StencilView &st,  const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 					  int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 			     int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|    |    | ||||||
|   static accelerator void HandDhopSiteInt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static void AsmDhopSiteInt(const StencilView &st,  const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 					  int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 			     int sF, int sU, int Ls, int Nsite, const FermionFieldView &in,const FermionFieldView &out); | ||||||
|    |    | ||||||
|   static accelerator void HandDhopSiteDagInt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static void AsmDhopSiteDagInt(const StencilView &st,  const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 					     int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 				int sF, int sU, int Ls, int Nsite, const FermionFieldView &in, const FermionFieldView &out); | ||||||
|    |    | ||||||
|   static accelerator void HandDhopSiteExt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |   static void AsmDhopSiteExt(const StencilView &st,  const DoubledGaugeFieldView &U, SiteHalfSpinor * buf, | ||||||
| 					  int sF, int sU, const FermionFieldView &in, FermionFieldView &out); | 			     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); | ||||||
|  |  | ||||||
|   static accelerator void HandDhopSiteDagExt(StencilView &st,  DoubledGaugeFieldView &U, SiteHalfSpinor * buf, |  | ||||||
| 					     int sF, int sU, const FermionFieldView &in, FermionFieldView &out); |  | ||||||
|  public: |  public: | ||||||
|  WilsonKernels(const ImplParams &p = ImplParams()) : Base(p){}; |  WilsonKernels(const ImplParams &p = ImplParams()) : Base(p){}; | ||||||
| }; | }; | ||||||
|   | |||||||
| @@ -618,11 +618,13 @@ Author: paboyle <paboyle@ph.ed.ac.uk> | |||||||
| NAMESPACE_BEGIN(Grid); | NAMESPACE_BEGIN(Grid); | ||||||
|  |  | ||||||
| template <class Impl> | template <class Impl> | ||||||
| void StaggeredKernels<Impl>::DhopSiteAsm(StencilView &st, | void StaggeredKernels<Impl>::DhopSiteAsm(const StencilView &st, | ||||||
| 					 DoubledGaugeFieldView &U, | 					 const DoubledGaugeFieldView &U, | ||||||
| 					 DoubledGaugeFieldView &UUU, | 					 const DoubledGaugeFieldView &UUU, | ||||||
| 					 SiteSpinor *buf, int sF, |  					 SiteSpinor *buf, int sF, | ||||||
| 					 int sU, const FermionFieldView &in, FermionFieldView &out,int dag)  | 					 int sU, | ||||||
|  | 					 const FermionFieldView &in, | ||||||
|  | 					 const FermionFieldView &out,int dag)  | ||||||
| { | { | ||||||
|   assert(0); |   assert(0); | ||||||
| }; | }; | ||||||
| @@ -683,11 +685,13 @@ void StaggeredKernels<Impl>::DhopSiteAsm(StencilView &st, | |||||||
|  |  | ||||||
|   // This is the single precision 5th direction vectorised kernel |   // This is the single precision 5th direction vectorised kernel | ||||||
| #include <Grid/simd/Intel512single.h> | #include <Grid/simd/Intel512single.h> | ||||||
| template <> void StaggeredKernels<StaggeredVec5dImplF>::DhopSiteAsm(StencilView &st, | template <> void StaggeredKernels<StaggeredVec5dImplF>::DhopSiteAsm(const StencilView &st, | ||||||
| 								    DoubledGaugeFieldView &U, | 								    const DoubledGaugeFieldView &U, | ||||||
| 								    DoubledGaugeFieldView &UUU, | 								    const DoubledGaugeFieldView &UUU, | ||||||
| 								    SiteSpinor *buf, int sF, |  								    SiteSpinor *buf, int sF, | ||||||
| 								    int sU, const FermionFieldView &in, FermionFieldView &out,int dag)  | 								    int sU, | ||||||
|  | 								    const FermionFieldView &in, | ||||||
|  | 								    const FermionFieldView &out,int dag)  | ||||||
| { | { | ||||||
| #ifdef AVX512 | #ifdef AVX512 | ||||||
|   uint64_t gauge0,gauge1,gauge2,gauge3; |   uint64_t gauge0,gauge1,gauge2,gauge3; | ||||||
| @@ -738,11 +742,13 @@ template <> void StaggeredKernels<StaggeredVec5dImplF>::DhopSiteAsm(StencilView | |||||||
| } | } | ||||||
|  |  | ||||||
| #include <Grid/simd/Intel512double.h> | #include <Grid/simd/Intel512double.h> | ||||||
| template <> void StaggeredKernels<StaggeredVec5dImplD>::DhopSiteAsm(StencilView &st,  | template <> void StaggeredKernels<StaggeredVec5dImplD>::DhopSiteAsm(const StencilView &st,  | ||||||
| 								    DoubledGaugeFieldView &U, | 								    const DoubledGaugeFieldView &U, | ||||||
| 								    DoubledGaugeFieldView &UUU, | 								    const DoubledGaugeFieldView &UUU, | ||||||
| 								    SiteSpinor *buf, int sF, |  								    SiteSpinor *buf, int sF, | ||||||
| 								    int sU, const FermionFieldView &in, FermionFieldView &out, int dag)  | 								    int sU, | ||||||
|  | 								    const FermionFieldView &in, | ||||||
|  | 								    const FermionFieldView &out, int dag)  | ||||||
| { | { | ||||||
| #ifdef AVX512 | #ifdef AVX512 | ||||||
|   uint64_t gauge0,gauge1,gauge2,gauge3; |   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 |   // This is the single precision 5th direction vectorised kernel | ||||||
|  |  | ||||||
| #include <Grid/simd/Intel512single.h> | #include <Grid/simd/Intel512single.h> | ||||||
| template <> void StaggeredKernels<StaggeredImplF>::DhopSiteAsm(StencilView &st,  | template <> void StaggeredKernels<StaggeredImplF>::DhopSiteAsm(const StencilView &st,  | ||||||
| 							       DoubledGaugeFieldView &U, | 							       const DoubledGaugeFieldView &U, | ||||||
| 							       DoubledGaugeFieldView &UUU, | 							       const DoubledGaugeFieldView &UUU, | ||||||
| 							       SiteSpinor *buf, int sF, |  							       SiteSpinor *buf, int sF, | ||||||
| 							       int sU, const FermionFieldView &in, FermionFieldView &out,int dag)  | 							       int sU, | ||||||
|  | 							       const FermionFieldView &in, | ||||||
|  | 							       const FermionFieldView &out,int dag)  | ||||||
| { | { | ||||||
| #ifdef AVX512 | #ifdef AVX512 | ||||||
|   uint64_t gauge0,gauge1,gauge2,gauge3; |   uint64_t gauge0,gauge1,gauge2,gauge3; | ||||||
| @@ -893,11 +901,13 @@ template <> void StaggeredKernels<StaggeredImplF>::DhopSiteAsm(StencilView &st, | |||||||
| } | } | ||||||
|  |  | ||||||
| #include <Grid/simd/Intel512double.h> | #include <Grid/simd/Intel512double.h> | ||||||
| template <> void StaggeredKernels<StaggeredImplD>::DhopSiteAsm(StencilView &st,  | template <> void StaggeredKernels<StaggeredImplD>::DhopSiteAsm(const StencilView &st,  | ||||||
| 							       DoubledGaugeFieldView &U, | 							       const DoubledGaugeFieldView &U, | ||||||
| 							       DoubledGaugeFieldView &UUU, | 							       const DoubledGaugeFieldView &UUU, | ||||||
| 							       SiteSpinor *buf, int sF, |  							       SiteSpinor *buf, int sF, | ||||||
| 							       int sU, const FermionFieldView &in, FermionFieldView &out,int dag)  | 							       int sU, | ||||||
|  | 							       const FermionFieldView &in, | ||||||
|  | 							       const FermionFieldView &out,int dag)  | ||||||
| { | { | ||||||
| #ifdef AVX512 | #ifdef AVX512 | ||||||
|   uint64_t gauge0,gauge1,gauge2,gauge3; |   uint64_t gauge0,gauge1,gauge2,gauge3; | ||||||
|   | |||||||
| @@ -147,10 +147,12 @@ NAMESPACE_BEGIN(Grid); | |||||||
|  |  | ||||||
| template <class Impl> | template <class Impl> | ||||||
| template <int Naik> accelerator_inline | template <int Naik> accelerator_inline | ||||||
| void StaggeredKernels<Impl>::DhopSiteHand(StencilView &st, | void StaggeredKernels<Impl>::DhopSiteHand(const StencilView &st, | ||||||
| 					  DoubledGaugeFieldView &U,DoubledGaugeFieldView &UUU, | 					  const DoubledGaugeFieldView &U, | ||||||
| 					  SiteSpinor *buf, int sF, int sU,  | 					  const DoubledGaugeFieldView &UUU, | ||||||
| 					  const FermionFieldView &in, FermionFieldView &out,int dag)  |  					  SiteSpinor *buf, int sF, int sU,  | ||||||
|  | 					  const FermionFieldView &in, | ||||||
|  | 					  const FermionFieldView &out,int dag)  | ||||||
| { | { | ||||||
|   typedef typename Simd::scalar_type S; |   typedef typename Simd::scalar_type S; | ||||||
|   typedef typename Simd::vector_type V; |   typedef typename Simd::vector_type V; | ||||||
| @@ -222,10 +224,12 @@ void StaggeredKernels<Impl>::DhopSiteHand(StencilView &st, | |||||||
|  |  | ||||||
| template <class Impl> | template <class Impl> | ||||||
| template <int Naik> accelerator_inline | template <int Naik> accelerator_inline | ||||||
| void StaggeredKernels<Impl>::DhopSiteHandInt(StencilView &st,  | void StaggeredKernels<Impl>::DhopSiteHandInt(const StencilView &st,  | ||||||
| 					     DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, | 					     const DoubledGaugeFieldView &U, | ||||||
| 					     SiteSpinor *buf, int sF, int sU,  | 					     const DoubledGaugeFieldView &UUU, | ||||||
| 					     const FermionFieldView &in, FermionFieldView &out,int dag)  |  					     SiteSpinor *buf, int sF, int sU,  | ||||||
|  | 					     const FermionFieldView &in, | ||||||
|  | 					     const FermionFieldView &out,int dag)  | ||||||
| { | { | ||||||
|   typedef typename Simd::scalar_type S; |   typedef typename Simd::scalar_type S; | ||||||
|   typedef typename Simd::vector_type V; |   typedef typename Simd::vector_type V; | ||||||
| @@ -301,10 +305,12 @@ void StaggeredKernels<Impl>::DhopSiteHandInt(StencilView &st, | |||||||
|  |  | ||||||
| template <class Impl> | template <class Impl> | ||||||
| template <int Naik> accelerator_inline | template <int Naik> accelerator_inline | ||||||
| void StaggeredKernels<Impl>::DhopSiteHandExt(StencilView &st, | void StaggeredKernels<Impl>::DhopSiteHandExt(const StencilView &st, | ||||||
| 					     DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, | 					     const DoubledGaugeFieldView &U, | ||||||
| 					     SiteSpinor *buf, int sF, int sU,  | 					     const DoubledGaugeFieldView &UUU, | ||||||
| 					     const FermionFieldView &in, FermionFieldView &out,int dag)  |  					     SiteSpinor *buf, int sF, int sU,  | ||||||
|  | 					     const FermionFieldView &in, | ||||||
|  | 					     const FermionFieldView &out,int dag)  | ||||||
| { | { | ||||||
|   typedef typename Simd::scalar_type S; |   typedef typename Simd::scalar_type S; | ||||||
|   typedef typename Simd::vector_type V; |   typedef typename Simd::vector_type V; | ||||||
|   | |||||||
| @@ -79,10 +79,10 @@ StaggeredKernels<Impl>::StaggeredKernels(const ImplParams &p) : Base(p){}; | |||||||
| //////////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////////// | ||||||
| template <class Impl> | template <class Impl> | ||||||
| template <int Naik> accelerator_inline | template <int Naik> accelerator_inline | ||||||
| void StaggeredKernels<Impl>::DhopSiteGeneric(StencilView &st,  | void StaggeredKernels<Impl>::DhopSiteGeneric(const StencilView &st,  | ||||||
| 					     DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, | 					     const DoubledGaugeFieldView &U, const DoubledGaugeFieldView &UUU, | ||||||
| 					     SiteSpinor *buf, int sF, int sU,  |  					     SiteSpinor *buf, int sF, int sU,  | ||||||
| 					     const FermionFieldView &in, FermionFieldView &out, int dag)  | 					     const FermionFieldView &in, const FermionFieldView &out, int dag)  | ||||||
| { | { | ||||||
|   const SiteSpinor *chi_p; |   const SiteSpinor *chi_p; | ||||||
|   SiteSpinor chi; |   SiteSpinor chi; | ||||||
| @@ -127,10 +127,11 @@ void StaggeredKernels<Impl>::DhopSiteGeneric(StencilView &st, | |||||||
|   /////////////////////////////////////////////////// |   /////////////////////////////////////////////////// | ||||||
| template <class Impl> | template <class Impl> | ||||||
| template <int Naik> accelerator_inline | template <int Naik> accelerator_inline | ||||||
| void StaggeredKernels<Impl>::DhopSiteGenericInt(StencilView &st,  | void StaggeredKernels<Impl>::DhopSiteGenericInt(const StencilView &st,  | ||||||
| 						DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, | 						const DoubledGaugeFieldView &U, const DoubledGaugeFieldView &UUU, | ||||||
| 						SiteSpinor *buf, int sF, int sU,  |  						SiteSpinor *buf, int sF, int sU,  | ||||||
| 						const FermionFieldView &in, FermionFieldView &out,int dag) { | 						const FermionFieldView &in, const FermionFieldView &out,int dag) | ||||||
|  | { | ||||||
|   const SiteSpinor *chi_p; |   const SiteSpinor *chi_p; | ||||||
|   SiteSpinor chi; |   SiteSpinor chi; | ||||||
|   SiteSpinor Uchi; |   SiteSpinor Uchi; | ||||||
| @@ -175,10 +176,13 @@ void StaggeredKernels<Impl>::DhopSiteGenericInt(StencilView &st, | |||||||
|   /////////////////////////////////////////////////// |   /////////////////////////////////////////////////// | ||||||
| template <class Impl> | template <class Impl> | ||||||
| template <int Naik> accelerator_inline | template <int Naik> accelerator_inline | ||||||
| void StaggeredKernels<Impl>::DhopSiteGenericExt(StencilView &st,  | void StaggeredKernels<Impl>::DhopSiteGenericExt(const StencilView &st,  | ||||||
| 						DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, | 						const DoubledGaugeFieldView &U, | ||||||
| 						SiteSpinor *buf, int sF, int sU, | 						const DoubledGaugeFieldView &UUU, | ||||||
| 						const FermionFieldView &in, FermionFieldView &out,int dag) { |  						SiteSpinor *buf, int sF, int sU, | ||||||
|  | 						const FermionFieldView &in, | ||||||
|  | 						const FermionFieldView &out,int dag) | ||||||
|  | { | ||||||
|   const SiteSpinor *chi_p; |   const SiteSpinor *chi_p; | ||||||
|   //  SiteSpinor chi; |   //  SiteSpinor chi; | ||||||
|   SiteSpinor Uchi; |   SiteSpinor Uchi; | ||||||
| @@ -225,8 +229,13 @@ void StaggeredKernels<Impl>::DhopSiteGenericExt(StencilView &st, | |||||||
| // Driving / wrapping routine to select right kernel | // Driving / wrapping routine to select right kernel | ||||||
| //////////////////////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////////////////////// | ||||||
| template <class Impl>  | template <class Impl>  | ||||||
| void StaggeredKernels<Impl>::DhopDirKernel(StencilImpl &st, DoubledGaugeFieldView &U, DoubledGaugeFieldView &UUU, SiteSpinor * buf, | void StaggeredKernels<Impl>::DhopDirKernel(StencilImpl &st, | ||||||
| 					   int sF, int sU, const FermionFieldView &in, FermionFieldView &out, int dir,int disp) | 					   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 |   // Disp should be either +1,-1,+3,-3 | ||||||
|   // What about "dag" ? |   // What about "dag" ? | ||||||
| @@ -254,7 +263,8 @@ void StaggeredKernels<Impl>::DhopDirKernel(StencilImpl &st, DoubledGaugeFieldVie | |||||||
|   }); |   }); | ||||||
|  |  | ||||||
| template <class Impl>  | template <class Impl>  | ||||||
| void StaggeredKernels<Impl>::DhopImproved(StencilImpl &st, LebesgueOrder &lo,  | void StaggeredKernels<Impl>::DhopImproved(StencilImpl &st, | ||||||
|  | 					  LebesgueOrder &lo,  | ||||||
| 					  DoubledGaugeField &U, DoubledGaugeField &UUU,  | 					  DoubledGaugeField &U, DoubledGaugeField &UUU,  | ||||||
| 					  const FermionField &in, FermionField &out, int dag, int interior,int exterior) | 					  const FermionField &in, FermionField &out, int dag, int interior,int exterior) | ||||||
| { | { | ||||||
|   | |||||||
| @@ -38,46 +38,46 @@ NAMESPACE_BEGIN(Grid); | |||||||
|  |  | ||||||
| /////////////////////////////////////////////////////////// | /////////////////////////////////////////////////////////// | ||||||
| // Default to no assembler implementation | // Default to no assembler implementation | ||||||
| // Will specialise to  | // Will specialise to AVX512 if available | ||||||
| /////////////////////////////////////////////////////////// | /////////////////////////////////////////////////////////// | ||||||
| template<class Impl> void  | template<class Impl> void  | ||||||
| WilsonKernels<Impl >::AsmDhopSite(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | WilsonKernels<Impl >::AsmDhopSite(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 				  int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out) | 				  int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   assert(0); |   assert(0); | ||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl> void  | template<class Impl> void  | ||||||
| WilsonKernels<Impl >::AsmDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | WilsonKernels<Impl >::AsmDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 				     int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out) | 				     int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   assert(0); |   assert(0); | ||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl> void  | template<class Impl> void  | ||||||
| WilsonKernels<Impl >::AsmDhopSiteInt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | WilsonKernels<Impl >::AsmDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 				     int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out) | 				     int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   assert(0); |   assert(0); | ||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl> void  | template<class Impl> void  | ||||||
| WilsonKernels<Impl >::AsmDhopSiteDagInt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | WilsonKernels<Impl >::AsmDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 					int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out) | 					int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   assert(0); |   assert(0); | ||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl> void  | template<class Impl> void  | ||||||
| WilsonKernels<Impl >::AsmDhopSiteExt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | WilsonKernels<Impl >::AsmDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 				     int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out) | 				     int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   assert(0); |   assert(0); | ||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl> void  | template<class Impl> void  | ||||||
| WilsonKernels<Impl >::AsmDhopSiteDagExt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | WilsonKernels<Impl >::AsmDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 					int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, FermionFieldView &out) | 					int ss,int ssU,int Ls,int Ns,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   assert(0); |   assert(0); | ||||||
| } | } | ||||||
|   | |||||||
| @@ -646,9 +646,14 @@ NAMESPACE_BEGIN(Grid); | |||||||
|   HAND_RESULT_EXT(ss,F) |   HAND_RESULT_EXT(ss,F) | ||||||
|  |  | ||||||
| #define HAND_SPECIALISE_GPARITY(IMPL)					\ | #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						\ |   template<> accelerator_inline void						\ | ||||||
|   WilsonKernels<IMPL>::HandDhopSite(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, \ |   WilsonKernels<IMPL>::HandDhopSite(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, \ | ||||||
| 				    int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \ | 				    int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \ | ||||||
|   {									\ |   {									\ | ||||||
|     typedef IMPL Impl;							\ |     typedef IMPL Impl;							\ | ||||||
|     typedef typename Simd::scalar_type S;				\ |     typedef typename Simd::scalar_type S;				\ | ||||||
| @@ -663,8 +668,8 @@ NAMESPACE_BEGIN(Grid); | |||||||
|   }									\ |   }									\ | ||||||
| 									\ | 									\ | ||||||
|   template<> accelerator_inline void						\ |   template<> accelerator_inline void						\ | ||||||
|   WilsonKernels<IMPL>::HandDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \ |   WilsonKernels<IMPL>::HandDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \ | ||||||
| 				       int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \ | 				       int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \ | ||||||
|   {									\ |   {									\ | ||||||
|     typedef IMPL Impl;							\ |     typedef IMPL Impl;							\ | ||||||
|     typedef typename Simd::scalar_type S;				\ |     typedef typename Simd::scalar_type S;				\ | ||||||
| @@ -679,8 +684,8 @@ NAMESPACE_BEGIN(Grid); | |||||||
|   }									\ |   }									\ | ||||||
| 									\ | 									\ | ||||||
|   template<> accelerator_inline void						\ |   template<> accelerator_inline void						\ | ||||||
|   WilsonKernels<IMPL>::HandDhopSiteInt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, \ |   WilsonKernels<IMPL>::HandDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, \ | ||||||
| 				       int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \ | 				       int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \ | ||||||
|   {									\ |   {									\ | ||||||
|     typedef IMPL Impl;							\ |     typedef IMPL Impl;							\ | ||||||
|     typedef typename Simd::scalar_type S;				\ |     typedef typename Simd::scalar_type S;				\ | ||||||
| @@ -695,8 +700,8 @@ NAMESPACE_BEGIN(Grid); | |||||||
|   }									\ |   }									\ | ||||||
| 									\ | 									\ | ||||||
|   template<> accelerator_inline void						\ |   template<> accelerator_inline void						\ | ||||||
|   WilsonKernels<IMPL>::HandDhopSiteDagInt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \ |   WilsonKernels<IMPL>::HandDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \ | ||||||
| 					  int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \ | 					  int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \ | ||||||
|   {									\ |   {									\ | ||||||
|     typedef IMPL Impl;							\ |     typedef IMPL Impl;							\ | ||||||
|     typedef typename Simd::scalar_type S;				\ |     typedef typename Simd::scalar_type S;				\ | ||||||
| @@ -711,8 +716,8 @@ NAMESPACE_BEGIN(Grid); | |||||||
|   }									\ |   }									\ | ||||||
| 									\ | 									\ | ||||||
|   template<> accelerator_inline void							\ |   template<> accelerator_inline void							\ | ||||||
|   WilsonKernels<IMPL>::HandDhopSiteExt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, \ |   WilsonKernels<IMPL>::HandDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, \ | ||||||
| 				       int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \ | 				       int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \ | ||||||
|   {									\ |   {									\ | ||||||
|     typedef IMPL Impl;							\ |     typedef IMPL Impl;							\ | ||||||
|     typedef typename Simd::scalar_type S;				\ |     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); \ |     HAND_DOP_SITE_EXT(1, LOAD_CHI_GPARITY,LOAD_CHIMU_GPARITY,MULT_2SPIN_GPARITY); \ | ||||||
|   }									\ |   }									\ | ||||||
|   template<> accelerator_inline void						\ |   template<> accelerator_inline void						\ | ||||||
|   WilsonKernels<IMPL>::HandDhopSiteDagExt(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \ |   WilsonKernels<IMPL>::HandDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, \ | ||||||
| 					  int ss,int sU,const FermionFieldView &in, FermionFieldView &out) \ | 					  int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) \ | ||||||
|   {									\ |   {									\ | ||||||
|     typedef IMPL Impl;							\ |     typedef IMPL Impl;							\ | ||||||
|     typedef typename Simd::scalar_type S;				\ |     typedef typename Simd::scalar_type S;				\ | ||||||
|   | |||||||
| @@ -496,8 +496,8 @@ Author: paboyle <paboyle@ph.ed.ac.uk> | |||||||
| NAMESPACE_BEGIN(Grid); | NAMESPACE_BEGIN(Grid); | ||||||
|  |  | ||||||
| template<class Impl> accelerator_inline void  | template<class Impl> accelerator_inline void  | ||||||
| WilsonKernels<Impl>::HandDhopSite(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, | WilsonKernels<Impl>::HandDhopSite(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, | ||||||
| 				  int ss,int sU,const FermionFieldView &in, FermionFieldView &out) | 				  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... | // 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::scalar_type S; | ||||||
| @@ -520,8 +520,8 @@ WilsonKernels<Impl>::HandDhopSite(StencilView &st, DoubledGaugeFieldView &U,Site | |||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl>  accelerator_inline | template<class Impl>  accelerator_inline | ||||||
| void WilsonKernels<Impl>::HandDhopSiteDag(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | void WilsonKernels<Impl>::HandDhopSiteDag(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 					  int ss,int sU,const FermionFieldView &in, FermionFieldView &out) | 					  int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef typename Simd::scalar_type S; |   typedef typename Simd::scalar_type S; | ||||||
|   typedef typename Simd::vector_type V; |   typedef typename Simd::vector_type V; | ||||||
| @@ -543,8 +543,8 @@ void WilsonKernels<Impl>::HandDhopSiteDag(StencilView &st,DoubledGaugeFieldView | |||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl>  accelerator_inline void  | template<class Impl>  accelerator_inline void  | ||||||
| WilsonKernels<Impl>::HandDhopSiteInt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, | WilsonKernels<Impl>::HandDhopSiteInt(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, | ||||||
| 					  int ss,int sU,const FermionFieldView &in, FermionFieldView &out) | 					  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... | // 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::scalar_type S; | ||||||
| @@ -567,8 +567,8 @@ WilsonKernels<Impl>::HandDhopSiteInt(StencilView &st,DoubledGaugeFieldView &U,Si | |||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl> accelerator_inline | template<class Impl> accelerator_inline | ||||||
| void WilsonKernels<Impl>::HandDhopSiteDagInt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | void WilsonKernels<Impl>::HandDhopSiteDagInt(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 						  int ss,int sU,const FermionFieldView &in, FermionFieldView &out) | 						  int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef typename Simd::scalar_type S; |   typedef typename Simd::scalar_type S; | ||||||
|   typedef typename Simd::vector_type V; |   typedef typename Simd::vector_type V; | ||||||
| @@ -590,8 +590,8 @@ void WilsonKernels<Impl>::HandDhopSiteDagInt(StencilView &st,DoubledGaugeFieldVi | |||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl>  accelerator_inline void  | template<class Impl>  accelerator_inline void  | ||||||
| WilsonKernels<Impl>::HandDhopSiteExt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, | WilsonKernels<Impl>::HandDhopSiteExt(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor  *buf, | ||||||
| 					  int ss,int sU,const FermionFieldView &in, FermionFieldView &out) | 					  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... | // 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::scalar_type S; | ||||||
| @@ -615,8 +615,8 @@ WilsonKernels<Impl>::HandDhopSiteExt(StencilView &st,DoubledGaugeFieldView &U,Si | |||||||
| } | } | ||||||
|  |  | ||||||
| template<class Impl>  accelerator_inline | template<class Impl>  accelerator_inline | ||||||
| void WilsonKernels<Impl>::HandDhopSiteDagExt(StencilView &st,DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | void WilsonKernels<Impl>::HandDhopSiteDagExt(const StencilView &st,const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, | ||||||
| 						  int ss,int sU,const FermionFieldView &in, FermionFieldView &out) | 						  int ss,int sU,const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef typename Simd::scalar_type S; |   typedef typename Simd::scalar_type S; | ||||||
|   typedef typename Simd::vector_type V; |   typedef typename Simd::vector_type V; | ||||||
| @@ -682,3 +682,4 @@ NAMESPACE_END(Grid); | |||||||
| #undef HAND_RESULT | #undef HAND_RESULT | ||||||
| #undef HAND_RESULT_INT | #undef HAND_RESULT_INT | ||||||
| #undef HAND_RESULT_EXT | #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 |   // All legs kernels ; comms then compute | ||||||
|   //////////////////////////////////////////////////////////////////// |   //////////////////////////////////////////////////////////////////// | ||||||
| template <class Impl> accelerator_inline | template <class Impl> accelerator_inline | ||||||
| void WilsonKernels<Impl>::GenericDhopSiteDag(StencilView &st, DoubledGaugeFieldView &U, | void WilsonKernels<Impl>::GenericDhopSiteDag(const StencilView &st, const DoubledGaugeFieldView &U, | ||||||
| 					     SiteHalfSpinor *buf, int sF, |  					     SiteHalfSpinor *buf, int sF, | ||||||
| 					     int sU, const FermionFieldView &in, FermionFieldView &out) | 					     int sU, const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef decltype(coalescedRead(buf[0]))   calcHalfSpinor; |   typedef decltype(coalescedRead(buf[0]))   calcHalfSpinor; | ||||||
|   typedef decltype(coalescedRead(in[0])) calcSpinor; |   typedef decltype(coalescedRead(in[0])) calcSpinor; | ||||||
| @@ -141,9 +141,9 @@ void WilsonKernels<Impl>::GenericDhopSiteDag(StencilView &st, DoubledGaugeFieldV | |||||||
| }; | }; | ||||||
|  |  | ||||||
| template <class Impl> accelerator_inline | template <class Impl> accelerator_inline | ||||||
| void WilsonKernels<Impl>::GenericDhopSite(StencilView &st, DoubledGaugeFieldView &U, | void WilsonKernels<Impl>::GenericDhopSite(const StencilView &st, const DoubledGaugeFieldView &U, | ||||||
| 					  SiteHalfSpinor *buf, int sF, |  					  SiteHalfSpinor *buf, int sF, | ||||||
| 					  int sU, const FermionFieldView &in, FermionFieldView &out) | 					  int sU, const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; |   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; | ||||||
|   typedef decltype(coalescedRead(in[0]))  calcSpinor; |   typedef decltype(coalescedRead(in[0]))  calcSpinor; | ||||||
| @@ -170,9 +170,9 @@ void WilsonKernels<Impl>::GenericDhopSite(StencilView &st, DoubledGaugeFieldView | |||||||
|   // Interior kernels |   // Interior kernels | ||||||
|   //////////////////////////////////////////////////////////////////// |   //////////////////////////////////////////////////////////////////// | ||||||
| template <class Impl> accelerator_inline | template <class Impl> accelerator_inline | ||||||
| void WilsonKernels<Impl>::GenericDhopSiteDagInt(StencilView &st,  DoubledGaugeFieldView &U, | void WilsonKernels<Impl>::GenericDhopSiteDagInt(const StencilView &st, const DoubledGaugeFieldView &U, | ||||||
| 						SiteHalfSpinor *buf, int sF, |  						SiteHalfSpinor *buf, int sF, | ||||||
| 						int sU, const FermionFieldView &in, FermionFieldView &out) | 						int sU, const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; |   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; | ||||||
|   typedef decltype(coalescedRead(in[0]))  calcSpinor; |   typedef decltype(coalescedRead(in[0]))  calcSpinor; | ||||||
| @@ -198,9 +198,9 @@ void WilsonKernels<Impl>::GenericDhopSiteDagInt(StencilView &st,  DoubledGaugeFi | |||||||
| }; | }; | ||||||
|  |  | ||||||
| template <class Impl> accelerator_inline | template <class Impl> accelerator_inline | ||||||
| void WilsonKernels<Impl>::GenericDhopSiteInt(StencilView &st,  DoubledGaugeFieldView &U, | void WilsonKernels<Impl>::GenericDhopSiteInt(const StencilView &st, const DoubledGaugeFieldView &U, | ||||||
| 							 SiteHalfSpinor *buf, int sF, | 					     SiteHalfSpinor *buf, int sF, | ||||||
| 							 int sU, const FermionFieldView &in, FermionFieldView &out) | 					     int sU, const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; |   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; | ||||||
|   typedef decltype(coalescedRead(in[0]))  calcSpinor; |   typedef decltype(coalescedRead(in[0]))  calcSpinor; | ||||||
| @@ -228,9 +228,9 @@ void WilsonKernels<Impl>::GenericDhopSiteInt(StencilView &st,  DoubledGaugeField | |||||||
| // Exterior kernels | // Exterior kernels | ||||||
| //////////////////////////////////////////////////////////////////// | //////////////////////////////////////////////////////////////////// | ||||||
| template <class Impl> accelerator_inline | template <class Impl> accelerator_inline | ||||||
| void WilsonKernels<Impl>::GenericDhopSiteDagExt(StencilView &st,  DoubledGaugeFieldView &U, | void WilsonKernels<Impl>::GenericDhopSiteDagExt(const StencilView &st, const DoubledGaugeFieldView &U, | ||||||
| 						SiteHalfSpinor *buf, int sF, |  						SiteHalfSpinor *buf, int sF, | ||||||
| 						int sU, const FermionFieldView &in, FermionFieldView &out) | 						int sU, const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; |   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; | ||||||
|   typedef decltype(coalescedRead(in[0]))  calcSpinor; |   typedef decltype(coalescedRead(in[0]))  calcSpinor; | ||||||
| @@ -259,9 +259,9 @@ void WilsonKernels<Impl>::GenericDhopSiteDagExt(StencilView &st,  DoubledGaugeFi | |||||||
| }; | }; | ||||||
|  |  | ||||||
| template <class Impl> accelerator_inline | template <class Impl> accelerator_inline | ||||||
| void WilsonKernels<Impl>::GenericDhopSiteExt(StencilView &st,  DoubledGaugeFieldView &U, | void WilsonKernels<Impl>::GenericDhopSiteExt(const StencilView &st, const DoubledGaugeFieldView &U, | ||||||
| 					     SiteHalfSpinor *buf, int sF, |  					     SiteHalfSpinor *buf, int sF, | ||||||
| 					     int sU, const FermionFieldView &in, FermionFieldView &out) | 					     int sU, const FermionFieldView &in, const FermionFieldView &out) | ||||||
| { | { | ||||||
|   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; |   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; | ||||||
|   typedef decltype(coalescedRead(in[0]))  calcSpinor; |   typedef decltype(coalescedRead(in[0]))  calcSpinor; | ||||||
| @@ -291,8 +291,8 @@ void WilsonKernels<Impl>::GenericDhopSiteExt(StencilView &st,  DoubledGaugeField | |||||||
|  |  | ||||||
| #define DhopDirMacro(Dir,spProj,spRecon)	\ | #define DhopDirMacro(Dir,spProj,spRecon)	\ | ||||||
|   template <class Impl> accelerator_inline				\ |   template <class Impl> accelerator_inline				\ | ||||||
|   void WilsonKernels<Impl>::DhopDir##Dir(StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, int sF, \ |   void WilsonKernels<Impl>::DhopDir##Dir(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, int sF, \ | ||||||
| 					 int sU, const FermionFieldView &in, FermionFieldView &out, int dir) \ | 					 int sU, const FermionFieldView &in, const FermionFieldView &out, int dir) \ | ||||||
|   {									\ |   {									\ | ||||||
|   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;		\ |   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor;		\ | ||||||
|   typedef decltype(coalescedRead(in[0]))  calcSpinor;			\ |   typedef decltype(coalescedRead(in[0]))  calcSpinor;			\ | ||||||
| @@ -302,8 +302,8 @@ void WilsonKernels<Impl>::GenericDhopSiteExt(StencilView &st,  DoubledGaugeField | |||||||
|   StencilEntry *SE;							\ |   StencilEntry *SE;							\ | ||||||
|   int ptype;								\ |   int ptype;								\ | ||||||
|   const int Nsimd = SiteHalfSpinor::Nsimd();				\ |   const int Nsimd = SiteHalfSpinor::Nsimd();				\ | ||||||
|   const int lane=acceleratorSIMTlane(Nsimd);					\ |   const int lane=acceleratorSIMTlane(Nsimd);				\ | ||||||
| 									\ |   									\ | ||||||
|   SE = st.GetEntry(ptype, dir, sF);					\ |   SE = st.GetEntry(ptype, dir, sF);					\ | ||||||
|   GENERIC_DHOPDIR_LEG_BODY(Dir,spProj,spRecon);				\ |   GENERIC_DHOPDIR_LEG_BODY(Dir,spProj,spRecon);				\ | ||||||
|   coalescedWrite(out[sF], result,lane);					\ |   coalescedWrite(out[sF], result,lane);					\ | ||||||
| @@ -319,8 +319,8 @@ DhopDirMacro(Zm,spProjZm,spReconZm); | |||||||
| DhopDirMacro(Tm,spProjTm,spReconTm); | DhopDirMacro(Tm,spProjTm,spReconTm); | ||||||
|  |  | ||||||
| template <class Impl> accelerator_inline | template <class Impl> accelerator_inline | ||||||
| void WilsonKernels<Impl>::DhopDirK( StencilView &st, DoubledGaugeFieldView &U,SiteHalfSpinor *buf, int sF, | void WilsonKernels<Impl>::DhopDirK(const StencilView &st, const DoubledGaugeFieldView &U,SiteHalfSpinor *buf, int sF, | ||||||
| 				    int sU, const FermionFieldView &in, FermionFieldView &out, int dir, int gamma) | 				   int sU, const FermionFieldView &in, const FermionFieldView &out, int dir, int gamma) | ||||||
| { | { | ||||||
|   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; |   typedef decltype(coalescedRead(buf[0])) calcHalfSpinor; | ||||||
|   typedef decltype(coalescedRead(in[0]))  calcSpinor; |   typedef decltype(coalescedRead(in[0]))  calcSpinor; | ||||||
| @@ -345,8 +345,8 @@ void WilsonKernels<Impl>::DhopDirK( StencilView &st, DoubledGaugeFieldView &U,Si | |||||||
| } | } | ||||||
|  |  | ||||||
| template <class Impl> | template <class Impl> | ||||||
| void WilsonKernels<Impl>::DhopDirAll( StencilImpl &st, DoubledGaugeField &U,SiteHalfSpinor *buf, int Ls, | void WilsonKernels<Impl>::DhopDirAll(StencilImpl &st, DoubledGaugeField &U,SiteHalfSpinor *buf, int Ls, | ||||||
| 				      int Nsite, const FermionField &in, std::vector<FermionField> &out) | 				     int Nsite, const FermionField &in, std::vector<FermionField> &out) | ||||||
| { | { | ||||||
|    autoView(U_v  ,U,AcceleratorRead); |    autoView(U_v  ,U,AcceleratorRead); | ||||||
|    autoView(in_v ,in,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);		\ |       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 KERNEL_CALL(A) KERNEL_CALLNB(A); accelerator_barrier(); | ||||||
|  |  | ||||||
| #define ASM_CALL(A)							\ | #define ASM_CALL(A)							\ | ||||||
| @@ -446,7 +460,8 @@ void WilsonKernels<Impl>::DhopKernel(int Opt,StencilImpl &st,  DoubledGaugeField | |||||||
|    if( interior && exterior ) { |    if( interior && exterior ) { | ||||||
|      if (Opt == WilsonKernelsStatic::OptGeneric    ) { KERNEL_CALL(GenericDhopSite); return;} |      if (Opt == WilsonKernelsStatic::OptGeneric    ) { KERNEL_CALL(GenericDhopSite); return;} | ||||||
| #ifndef GRID_CUDA | #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;} |      if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL(AsmDhopSite);    return;} | ||||||
| #endif | #endif | ||||||
|    } else if( interior ) { |    } 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/FermionCore.h> | ||||||
| #include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h> | #include <Grid/qcd/action/fermion/implementation/WilsonKernelsImplementation.h> | ||||||
| #include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h> | #include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h> | ||||||
|  | #include <Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementationSycl.h> | ||||||
|  |  | ||||||
| #ifndef AVX512 | #ifndef AVX512 | ||||||
| #ifndef QPX | #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 { | class GpuComplex { | ||||||
| public: | public: | ||||||
|   pair z; |   pair z; | ||||||
|   typedef decltype(z.x) real; |   typedef decltype(z.x) Real; | ||||||
| public:  | public:  | ||||||
|   accelerator_inline GpuComplex() = default; |   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 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) {  |   friend accelerator_inline  GpuComplex operator+(const GpuComplex &lhs,const GpuComplex &rhs) {  | ||||||
|     GpuComplex r ;  |     GpuComplex r ;  | ||||||
|     r.z.x = lhs.z.x + rhs.z.x;  |     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_ComplexD, GpuComplexD > GpuVectorCD; | ||||||
| typedef GpuVector<NSIMD_Integer,  Integer     > GpuVectorI; | 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) | accelerator_inline float half2float(half h) | ||||||
| { | { | ||||||
|   float f; |   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); } | accelerator_inline void add (ComplexF * __restrict__ y,const ComplexF * __restrict__ l,const ComplexF *__restrict__ r){ *y = (*l) + (*r); } | ||||||
|    |    | ||||||
| //conjugate already supported for complex | //conjugate already supported for complex | ||||||
| accelerator_inline ComplexF timesI(const ComplexF &r)     { return(r*ComplexF(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(r*ComplexD(0.0,1.0));} | accelerator_inline ComplexD timesI(const ComplexD &r)     { return(ComplexD(-r.imag(),r.real()));} | ||||||
| accelerator_inline ComplexF timesMinusI(const ComplexF &r){ return(r*ComplexF(0.0,-1.0));} | accelerator_inline ComplexF timesMinusI(const ComplexF &r){ return(ComplexF(r.imag(),-r.real()));} | ||||||
| accelerator_inline ComplexD timesMinusI(const ComplexD &r){ return(r*ComplexD(0.0,-1.0));} | 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 | // define projections to real and imaginay parts | ||||||
| accelerator_inline ComplexF projReal(const ComplexF &r){return( ComplexF(r.real(), 0.0));} | accelerator_inline ComplexF projReal(const ComplexF &r){return( ComplexF(r.real(), 0.0));} | ||||||
|   | |||||||
| @@ -7,20 +7,20 @@ template<class vobj> | |||||||
| class SimpleCompressor { | class SimpleCompressor { | ||||||
| public: | public: | ||||||
|   void Point(int) {}; |   void Point(int) {}; | ||||||
|   accelerator_inline int  CommDatumSize(void) { return sizeof(vobj); } |   accelerator_inline int  CommDatumSize(void) const { return sizeof(vobj); } | ||||||
|   accelerator_inline bool DecompressionStep(void) { return false; } |   accelerator_inline bool DecompressionStep(void) const { return false; } | ||||||
|   template<class cobj> accelerator_inline void Compress(cobj *buf,int o,const cobj &in) { buf[o]=in; } |   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){ |   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); |     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, |   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); |     exchange(out0[j],out1[j],in[k],in[m],type); | ||||||
|   } |   } | ||||||
|   // For cshift. Cshift should drop compressor coupling altogether  |   // For cshift. Cshift should drop compressor coupling altogether  | ||||||
|   // because I had to decouple the code from the Stencil anyway |   // 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; |     return arg; | ||||||
|   } |   } | ||||||
| }; | }; | ||||||
|   | |||||||
| @@ -147,16 +147,16 @@ class CartesianStencilAccelerator { | |||||||
|   cobj* u_recv_buf_p; |   cobj* u_recv_buf_p; | ||||||
|   cobj* u_send_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; |     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]; |     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]; |     uint64_t cbase = (uint64_t)&u_recv_buf_p[0]; | ||||||
|     local = this->_entries_p[ent]._is_local; |     local = this->_entries_p[ent]._is_local; | ||||||
|     perm  = this->_entries_p[ent]._permute; |     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]; |     uint64_t cbase = (uint64_t)&u_recv_buf_p[0]; | ||||||
|     int local = this->_entries_p[ent]._is_local; |     int local = this->_entries_p[ent]._is_local; | ||||||
|     if (local) return  base + this->_entries_p[ent]._byte_offset; |     if (local) return  base + this->_entries_p[ent]._byte_offset; | ||||||
|     else       return cbase + 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); |     Lexicographic::CoorFromIndex(coor,lane,this->_simd_layout); | ||||||
|   } |   } | ||||||
|   | |||||||
| @@ -64,6 +64,68 @@ void coalescedWriteNonTemporal(vobj & __restrict__ vec,const vobj & __restrict__ | |||||||
| } | } | ||||||
| #else | #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 | // 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};			\ |       cl::sycl::range<3> global{unum1,unum2,nsimd};			\ | ||||||
|       cgh.parallel_for<class dslash>(					\ |       cgh.parallel_for<class dslash>(					\ | ||||||
|       cl::sycl::nd_range<3>(global,local), \ |       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 iter1    = item.get_global_id(0);	      \ | ||||||
|       auto iter2    = item.get_global_id(1);	      \ |       auto iter2    = item.get_global_id(1);	      \ | ||||||
|       auto lane     = item.get_global_id(2);	      \ |       auto lane     = item.get_global_id(2);	      \ | ||||||
|   | |||||||
		Reference in New Issue
	
	Block a user