From 679d1d22f7e920163acf7b88515c30ceb297387b Mon Sep 17 00:00:00 2001 From: u61464 Date: Wed, 3 Mar 2021 11:21:43 -0800 Subject: [PATCH 1/2] Sycl happier --- Grid/Makefile.am | 4 ++- Grid/cartesian/Cartesian_red_black.h | 2 +- Grid/qcd/action/fermion/FermionOperatorImpl.h | 4 +-- Grid/qcd/action/fermion/WilsonCompressor.h | 24 ++++++++-------- Grid/qcd/action/fermion/WilsonImpl.h | 2 +- .../WilsonKernelsHandImplementation.h | 28 +++++++++---------- Grid/simd/Grid_gpu_vec.h | 1 + Grid/stencil/SimpleCompressor.h | 14 +++++----- Grid/stencil/Stencil.h | 14 +++++----- Grid/tensors/Tensor_SIMT.h | 4 ++- Grid/threads/Accelerator.h | 2 +- benchmarks/Benchmark_dwf_fp32.cc | 2 +- configure.ac | 11 ++++++++ 13 files changed, 64 insertions(+), 48 deletions(-) diff --git a/Grid/Makefile.am b/Grid/Makefile.am index ded6d146..7c3c151b 100644 --- a/Grid/Makefile.am +++ b/Grid/Makefile.am @@ -54,9 +54,11 @@ Version.h: version-cache include Make.inc include Eigen.inc -extra_sources+=$(ZWILS_FERMION_FILES) extra_sources+=$(WILS_FERMION_FILES) extra_sources+=$(STAG_FERMION_FILES) +if BUILD_ZMOBIUS + extra_sources+=$(ZWILS_FERMION_FILES) +endif if BUILD_GPARITY extra_sources+=$(GP_FERMION_FILES) endif diff --git a/Grid/cartesian/Cartesian_red_black.h b/Grid/cartesian/Cartesian_red_black.h index b71981f5..092d4910 100644 --- a/Grid/cartesian/Cartesian_red_black.h +++ b/Grid/cartesian/Cartesian_red_black.h @@ -36,7 +36,7 @@ static const int CbBlack=1; static const int Even =CbRed; static const int Odd =CbBlack; -accelerator_inline int RedBlackCheckerBoardFromOindex (int oindex, Coordinate &rdim, Coordinate &chk_dim_msk) +accelerator_inline int RedBlackCheckerBoardFromOindex (int oindex,const Coordinate &rdim,const Coordinate &chk_dim_msk) { int nd=rdim.size(); Coordinate coor(nd); diff --git a/Grid/qcd/action/fermion/FermionOperatorImpl.h b/Grid/qcd/action/fermion/FermionOperatorImpl.h index b444f6dc..9345c0e6 100644 --- a/Grid/qcd/action/fermion/FermionOperatorImpl.h +++ b/Grid/qcd/action/fermion/FermionOperatorImpl.h @@ -153,8 +153,8 @@ public: typedef typename Impl::StencilImpl StencilImpl; \ typedef typename Impl::ImplParams ImplParams; \ typedef typename Impl::StencilImpl::View_type StencilView; \ - typedef typename ViewMap::Type FermionFieldView; \ - typedef typename ViewMap::Type DoubledGaugeFieldView; + typedef const typename ViewMap::Type FermionFieldView; \ + typedef const typename ViewMap::Type DoubledGaugeFieldView; #define INHERIT_IMPL_TYPES(Base) \ INHERIT_GIMPL_TYPES(Base) \ diff --git a/Grid/qcd/action/fermion/WilsonCompressor.h b/Grid/qcd/action/fermion/WilsonCompressor.h index 10e98f33..0760bcba 100644 --- a/Grid/qcd/action/fermion/WilsonCompressor.h +++ b/Grid/qcd/action/fermion/WilsonCompressor.h @@ -61,7 +61,7 @@ public: typedef typename SiteHalfSpinor::vector_type vComplexHigh; constexpr static int Nw=sizeof(SiteHalfSpinor)/sizeof(vComplexHigh); - accelerator_inline int CommDatumSize(void) { + accelerator_inline int CommDatumSize(void) const { return sizeof(SiteHalfCommSpinor); } @@ -69,7 +69,7 @@ public: /* Compress includes precision change if mpi data is not same */ /*****************************************************/ template - accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) { + accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) const { _SiteHalfSpinor tmp; projector::Proj(tmp,in,mu,dag); vstream(buf[o],tmp); @@ -81,7 +81,7 @@ public: accelerator_inline void Exchange(SiteHalfSpinor *mp, const SiteHalfSpinor * __restrict__ vp0, const SiteHalfSpinor * __restrict__ vp1, - Integer type,Integer o){ + Integer type,Integer o) const { SiteHalfSpinor tmp1; SiteHalfSpinor tmp2; exchange(tmp1,tmp2,vp0[o],vp1[o],type); @@ -93,7 +93,7 @@ public: /* Have a decompression step if mpi data is not same */ /*****************************************************/ accelerator_inline void Decompress(SiteHalfSpinor * __restrict__ out, - SiteHalfSpinor * __restrict__ in, Integer o) { + SiteHalfSpinor * __restrict__ in, Integer o) const { assert(0); } @@ -103,7 +103,7 @@ public: accelerator_inline void CompressExchange(SiteHalfSpinor * __restrict__ out0, SiteHalfSpinor * __restrict__ out1, const SiteSpinor * __restrict__ in, - Integer j,Integer k, Integer m,Integer type) + Integer j,Integer k, Integer m,Integer type) const { SiteHalfSpinor temp1, temp2; SiteHalfSpinor temp3, temp4; @@ -117,7 +117,7 @@ public: /*****************************************************/ /* Pass the info to the stencil */ /*****************************************************/ - accelerator_inline bool DecompressionStep(void) { return false; } + accelerator_inline bool DecompressionStep(void) const { return false; } }; @@ -142,7 +142,7 @@ public: typedef typename SiteHalfSpinor::vector_type vComplexHigh; constexpr static int Nw=sizeof(SiteHalfSpinor)/sizeof(vComplexHigh); - accelerator_inline int CommDatumSize(void) { + accelerator_inline int CommDatumSize(void) const { return sizeof(SiteHalfCommSpinor); } @@ -150,7 +150,7 @@ public: /* Compress includes precision change if mpi data is not same */ /*****************************************************/ template - accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) { + accelerator_inline void Compress(_SiteHalfSpinor *buf,Integer o,const _SiteSpinor &in) const { _SiteHalfSpinor hsp; SiteHalfCommSpinor *hbuf = (SiteHalfCommSpinor *)buf; projector::Proj(hsp,in,mu,dag); @@ -163,7 +163,7 @@ public: accelerator_inline void Exchange(SiteHalfSpinor *mp, SiteHalfSpinor *vp0, SiteHalfSpinor *vp1, - Integer type,Integer o){ + Integer type,Integer o) const { SiteHalfSpinor vt0,vt1; SiteHalfCommSpinor *vpp0 = (SiteHalfCommSpinor *)vp0; SiteHalfCommSpinor *vpp1 = (SiteHalfCommSpinor *)vp1; @@ -175,7 +175,7 @@ public: /*****************************************************/ /* Have a decompression step if mpi data is not same */ /*****************************************************/ - accelerator_inline void Decompress(SiteHalfSpinor *out, SiteHalfSpinor *in, Integer o){ + accelerator_inline void Decompress(SiteHalfSpinor *out, SiteHalfSpinor *in, Integer o) const { SiteHalfCommSpinor *hin=(SiteHalfCommSpinor *)in; precisionChange((vComplexHigh *)&out[o],(vComplexLow *)&hin[o],Nw); } @@ -186,7 +186,7 @@ public: accelerator_inline void CompressExchange(SiteHalfSpinor *out0, SiteHalfSpinor *out1, const SiteSpinor *in, - Integer j,Integer k, Integer m,Integer type){ + Integer j,Integer k, Integer m,Integer type) const { SiteHalfSpinor temp1, temp2,temp3,temp4; SiteHalfCommSpinor *hout0 = (SiteHalfCommSpinor *)out0; SiteHalfCommSpinor *hout1 = (SiteHalfCommSpinor *)out1; @@ -200,7 +200,7 @@ public: /*****************************************************/ /* Pass the info to the stencil */ /*****************************************************/ - accelerator_inline bool DecompressionStep(void) { return true; } + accelerator_inline bool DecompressionStep(void) const { return true; } }; diff --git a/Grid/qcd/action/fermion/WilsonImpl.h b/Grid/qcd/action/fermion/WilsonImpl.h index d7941d1f..94676b6b 100644 --- a/Grid/qcd/action/fermion/WilsonImpl.h +++ b/Grid/qcd/action/fermion/WilsonImpl.h @@ -72,7 +72,7 @@ public: typedef WilsonCompressor Compressor; typedef WilsonImplParams ImplParams; typedef WilsonStencil StencilImpl; - typedef typename StencilImpl::View_type StencilView; + typedef const typename StencilImpl::View_type StencilView; ImplParams Params; diff --git a/Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h b/Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h index b867369f..688cb75a 100644 --- a/Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h +++ b/Grid/qcd/action/fermion/implementation/WilsonKernelsHandImplementation.h @@ -93,7 +93,7 @@ Author: paboyle Chimu_32=coalescedReadPermute(ref()(3)(2),perm); } #define PERMUTE_DIR(dir) ; #else -#define LOAD_CHIMU \ +#define LOAD_CHIMU(ptype) \ {const SiteSpinor & ref (in[offset]); \ Chimu_00=ref()(0)(0);\ Chimu_01=ref()(0)(1);\ @@ -482,19 +482,19 @@ Author: paboyle Simd U_11; \ Simd U_21; -#define ZERO_RESULT \ - result_00=S(0.0,0.0); \ - result_01=S(0.0,0.0); \ - result_02=S(0.0,0.0); \ - result_10=S(0.0,0.0); \ - result_11=S(0.0,0.0); \ - result_12=S(0.0,0.0); \ - result_20=S(0.0,0.0); \ - result_21=S(0.0,0.0); \ - result_22=S(0.0,0.0); \ - result_30=S(0.0,0.0); \ - result_31=S(0.0,0.0); \ - result_32=S(0.0,0.0); +#define ZERO_RESULT \ + zeroit(result_00); \ + zeroit(result_01); \ + zeroit(result_02); \ + zeroit(result_10); \ + zeroit(result_11); \ + zeroit(result_12); \ + zeroit(result_20); \ + zeroit(result_21); \ + zeroit(result_22); \ + zeroit(result_30); \ + zeroit(result_31); \ + zeroit(result_32); #define Chimu_00 Chi_00 #define Chimu_01 Chi_01 diff --git a/Grid/simd/Grid_gpu_vec.h b/Grid/simd/Grid_gpu_vec.h index 2c1a38e7..b2c7588f 100644 --- a/Grid/simd/Grid_gpu_vec.h +++ b/Grid/simd/Grid_gpu_vec.h @@ -67,6 +67,7 @@ public: 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 Zero &zz) { z.x = 0; z.y=0; return *this; }; accelerator_inline GpuComplex &operator*=(const GpuComplex &r) { *this = (*this) * r; return *this; diff --git a/Grid/stencil/SimpleCompressor.h b/Grid/stencil/SimpleCompressor.h index be7c89c0..2ce48369 100644 --- a/Grid/stencil/SimpleCompressor.h +++ b/Grid/stencil/SimpleCompressor.h @@ -7,20 +7,20 @@ template class SimpleCompressor { public: void Point(int) {}; - accelerator_inline int CommDatumSize(void) { return sizeof(vobj); } - accelerator_inline bool DecompressionStep(void) { return false; } - template accelerator_inline void Compress(cobj *buf,int o,const cobj &in) { buf[o]=in; } - accelerator_inline void Exchange(vobj *mp,vobj *vp0,vobj *vp1,Integer type,Integer o){ + accelerator_inline int CommDatumSize(void) const { return sizeof(vobj); } + accelerator_inline bool DecompressionStep(void) const { return false; } + template accelerator_inline void Compress(cobj *buf,int o,const cobj &in) const { buf[o]=in; } + accelerator_inline void Exchange(vobj *mp,vobj *vp0,vobj *vp1,Integer type,Integer o) const { exchange(mp[2*o],mp[2*o+1],vp0[o],vp1[o],type); } - accelerator_inline void Decompress(vobj *out,vobj *in, int o){ assert(0); } + accelerator_inline void Decompress(vobj *out,vobj *in, int o) const { assert(0); } accelerator_inline void CompressExchange(vobj *out0,vobj *out1,const vobj *in, - int j,int k, int m,int type){ + int j,int k, int m,int type) const { exchange(out0[j],out1[j],in[k],in[m],type); } // For cshift. Cshift should drop compressor coupling altogether // because I had to decouple the code from the Stencil anyway - accelerator_inline vobj operator() (const vobj &arg) { + accelerator_inline vobj operator() (const vobj &arg) const { return arg; } }; diff --git a/Grid/stencil/Stencil.h b/Grid/stencil/Stencil.h index 23fc8203..58cebed3 100644 --- a/Grid/stencil/Stencil.h +++ b/Grid/stencil/Stencil.h @@ -147,16 +147,16 @@ class CartesianStencilAccelerator { cobj* u_recv_buf_p; cobj* u_send_buf_p; - accelerator_inline cobj *CommBuf(void) { return u_recv_buf_p; } + accelerator_inline cobj *CommBuf(void) const { return u_recv_buf_p; } - accelerator_inline int GetNodeLocal(int osite,int point) { + accelerator_inline int GetNodeLocal(int osite,int point) const { return this->_entries_p[point+this->_npoints*osite]._is_local; } - accelerator_inline StencilEntry * GetEntry(int &ptype,int point,int osite) { + accelerator_inline StencilEntry * GetEntry(int &ptype,int point,int osite) const { ptype = this->_permute_type[point]; return & this->_entries_p[point+this->_npoints*osite]; } - accelerator_inline uint64_t GetInfo(int &ptype,int &local,int &perm,int point,int ent,uint64_t base) { + accelerator_inline uint64_t GetInfo(int &ptype,int &local,int &perm,int point,int ent,uint64_t base) const { uint64_t cbase = (uint64_t)&u_recv_buf_p[0]; local = this->_entries_p[ent]._is_local; perm = this->_entries_p[ent]._permute; @@ -168,14 +168,14 @@ class CartesianStencilAccelerator { } } - accelerator_inline uint64_t GetPFInfo(int ent,uint64_t base) { + accelerator_inline uint64_t GetPFInfo(int ent,uint64_t base) const { uint64_t cbase = (uint64_t)&u_recv_buf_p[0]; int local = this->_entries_p[ent]._is_local; if (local) return base + this->_entries_p[ent]._byte_offset; else return cbase + this->_entries_p[ent]._byte_offset; } - accelerator_inline void iCoorFromIindex(Coordinate &coor,int lane) + accelerator_inline void iCoorFromIindex(Coordinate &coor,int lane) const { Lexicographic::CoorFromIndex(coor,lane,this->_simd_layout); } @@ -221,7 +221,7 @@ public: typedef typename cobj::vector_type vector_type; typedef typename cobj::scalar_type scalar_type; typedef typename cobj::scalar_object scalar_object; - typedef CartesianStencilView View_type; + typedef const CartesianStencilView View_type; typedef typename View_type::StencilVector StencilVector; /////////////////////////////////////////// // Helper structs diff --git a/Grid/tensors/Tensor_SIMT.h b/Grid/tensors/Tensor_SIMT.h index ede24fbe..672f385f 100644 --- a/Grid/tensors/Tensor_SIMT.h +++ b/Grid/tensors/Tensor_SIMT.h @@ -66,7 +66,7 @@ void coalescedWriteNonTemporal(vobj & __restrict__ vec,const vobj & __restrict__ #ifndef GRID_SYCL -// Use the scalar as our own complex on GPU +// Use the scalar as our own complex on GPU ... thrust::complex or std::complex template = 0> accelerator_inline typename vsimd::scalar_type coalescedRead(const vsimd & __restrict__ vec,int lane=acceleratorSIMTlane(vsimd::Nsimd())) @@ -96,6 +96,8 @@ void coalescedWrite(vsimd & __restrict__ vec, p[lane]=extracted; } #else +// For SyCL have option to use GpuComplex from inside the vector type in SIMT loops +// Faster for some reason template = 0> accelerator_inline typename vsimd::vector_type::datum coalescedRead(const vsimd & __restrict__ vec,int lane=acceleratorSIMTlane(vsimd::Nsimd())) diff --git a/Grid/threads/Accelerator.h b/Grid/threads/Accelerator.h index 2b7bf53a..f1a694fb 100644 --- a/Grid/threads/Accelerator.h +++ b/Grid/threads/Accelerator.h @@ -456,7 +456,7 @@ accelerator_inline void acceleratorSynchronise(void) __syncwarp(); #endif #ifdef GRID_SYCL - // No barrier call on SYCL?? // Option get __spir:: stuff to do warp barrier + cl::sycl::detail::workGroupBarrier(); #endif #ifdef GRID_HIP __syncthreads(); diff --git a/benchmarks/Benchmark_dwf_fp32.cc b/benchmarks/Benchmark_dwf_fp32.cc index cb86177e..03f3ee61 100644 --- a/benchmarks/Benchmark_dwf_fp32.cc +++ b/benchmarks/Benchmark_dwf_fp32.cc @@ -53,7 +53,7 @@ int main (int argc, char ** argv) int threads = GridThread::GetThreads(); Coordinate latt4 = GridDefaultLatt(); - int Ls=8; + int Ls=16; for(int i=0;i> Ls; diff --git a/configure.ac b/configure.ac index fb0c78fc..5f165412 100644 --- a/configure.ac +++ b/configure.ac @@ -140,12 +140,23 @@ AC_ARG_ENABLE([gparity], [ac_GPARITY=${enable_gparity}], [ac_GPARITY=yes]) AM_CONDITIONAL(BUILD_GPARITY, [ test "${ac_GPARITY}X" == "yesX" ]) + +AC_ARG_ENABLE([zmobius], + [AC_HELP_STRING([--enable-zmobius=yes|no], [enable Zmobius support])], + [ac_ZMOBIUS=${enable_zmobius}], [ac_ZMOBIUS=yes]) + +AM_CONDITIONAL(BUILD_ZMOBIUS, [ test "${ac_ZMOBIUS}X" == "yesX" ]) + + case ${ac_FERMION_REPS} in yes) AC_DEFINE([ENABLE_FERMION_REPS],[1],[non QCD fermion reps]);; esac case ${ac_GPARITY} in yes) AC_DEFINE([ENABLE_GPARITY],[1],[fermion actions with GPARITY BCs]);; esac +case ${ac_ZMOBIUS} in + yes) AC_DEFINE([ENABLE_ZMOBIUS],[1],[Zmobius fermion actions]);; +esac ############### Nc AC_ARG_ENABLE([Nc], [AC_HELP_STRING([--enable-Nc=2|3|4], [enable number of colours])], From 1eea9d73b984252c69869f65991f2f33d6e93193 Mon Sep 17 00:00:00 2001 From: Peter Boyle Date: Wed, 3 Mar 2021 23:50:01 +0100 Subject: [PATCH 2/2] Pass serial RNG around --- Grid/qcd/action/ActionBase.h | 2 +- Grid/qcd/action/gauge/GaugeImplTypes.h | 2 +- Grid/qcd/action/gauge/PlaqPlusRectangleAction.h | 2 +- Grid/qcd/action/gauge/WilsonGaugeAction.h | 3 +-- .../qcd/action/pseudofermion/ExactOneFlavourRatio.h | 2 +- .../pseudofermion/OneFlavourEvenOddRational.h | 11 ++++++----- .../pseudofermion/OneFlavourEvenOddRationalRatio.h | 7 +++++-- Grid/qcd/action/pseudofermion/OneFlavourRational.h | 7 +++++-- .../action/pseudofermion/OneFlavourRationalRatio.h | 7 +++++-- Grid/qcd/action/pseudofermion/TwoFlavour.h | 2 +- Grid/qcd/action/pseudofermion/TwoFlavourEvenOdd.h | 2 +- .../action/pseudofermion/TwoFlavourEvenOddRatio.h | 2 +- Grid/qcd/action/pseudofermion/TwoFlavourRatio.h | 2 +- Grid/qcd/action/scalar/ScalarAction.h | 2 +- Grid/qcd/action/scalar/ScalarImpl.h | 4 ++-- Grid/qcd/action/scalar/ScalarInteractionAction.h | 2 +- Grid/qcd/hmc/HMC.h | 2 +- Grid/qcd/hmc/integrators/Integrator.h | 13 ++++++------- benchmarks/Benchmark_dwf_fp32.cc | 2 +- 19 files changed, 42 insertions(+), 34 deletions(-) diff --git a/Grid/qcd/action/ActionBase.h b/Grid/qcd/action/ActionBase.h index bff21d1d..17980ee0 100644 --- a/Grid/qcd/action/ActionBase.h +++ b/Grid/qcd/action/ActionBase.h @@ -41,7 +41,7 @@ class Action public: bool is_smeared = false; // Heatbath? - virtual void refresh(const GaugeField& U, GridParallelRNG& pRNG) = 0; // refresh pseudofermions + virtual void refresh(const GaugeField& U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) = 0; // refresh pseudofermions virtual RealD S(const GaugeField& U) = 0; // evaluate the action virtual void deriv(const GaugeField& U, GaugeField& dSdU) = 0; // evaluate the action derivative virtual std::string action_name() = 0; // return the action name diff --git a/Grid/qcd/action/gauge/GaugeImplTypes.h b/Grid/qcd/action/gauge/GaugeImplTypes.h index 55a20eca..2499e0e9 100644 --- a/Grid/qcd/action/gauge/GaugeImplTypes.h +++ b/Grid/qcd/action/gauge/GaugeImplTypes.h @@ -96,7 +96,7 @@ public: /////////////////////////////////////////////////////////// // Move these to another class // HMC auxiliary functions - static inline void generate_momenta(Field &P, GridParallelRNG &pRNG) + static inline void generate_momenta(Field &P, GridSerialRNG & sRNG, GridParallelRNG &pRNG) { // Zbigniew Srocinsky thesis: // diff --git a/Grid/qcd/action/gauge/PlaqPlusRectangleAction.h b/Grid/qcd/action/gauge/PlaqPlusRectangleAction.h index 639aca19..7690092d 100644 --- a/Grid/qcd/action/gauge/PlaqPlusRectangleAction.h +++ b/Grid/qcd/action/gauge/PlaqPlusRectangleAction.h @@ -49,7 +49,7 @@ public: virtual std::string action_name(){return "PlaqPlusRectangleAction";} - virtual void refresh(const GaugeField &U, GridParallelRNG& pRNG) {}; // noop as no pseudoferms + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) {}; // noop as no pseudoferms virtual std::string LogParameters(){ std::stringstream sstream; diff --git a/Grid/qcd/action/gauge/WilsonGaugeAction.h b/Grid/qcd/action/gauge/WilsonGaugeAction.h index 40d600d2..f535b54f 100644 --- a/Grid/qcd/action/gauge/WilsonGaugeAction.h +++ b/Grid/qcd/action/gauge/WilsonGaugeAction.h @@ -54,8 +54,7 @@ public: return sstream.str(); } - virtual void refresh(const GaugeField &U, - GridParallelRNG &pRNG){}; // noop as no pseudoferms + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG &pRNG){}; // noop as no pseudoferms virtual RealD S(const GaugeField &U) { RealD plaq = WilsonLoops::avgPlaquette(U); diff --git a/Grid/qcd/action/pseudofermion/ExactOneFlavourRatio.h b/Grid/qcd/action/pseudofermion/ExactOneFlavourRatio.h index 9fc0a3b0..576a8cf6 100644 --- a/Grid/qcd/action/pseudofermion/ExactOneFlavourRatio.h +++ b/Grid/qcd/action/pseudofermion/ExactOneFlavourRatio.h @@ -124,7 +124,7 @@ NAMESPACE_BEGIN(Grid); // // As a check of rational require \Phi^dag M_{EOFA} \Phi == eta^dag M^-1/2^dag M M^-1/2 eta = eta^dag eta // - virtual void refresh(const GaugeField& U, GridParallelRNG& pRNG) + virtual void refresh(const GaugeField& U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) { Lop.ImportGauge(U); Rop.ImportGauge(U); diff --git a/Grid/qcd/action/pseudofermion/OneFlavourEvenOddRational.h b/Grid/qcd/action/pseudofermion/OneFlavourEvenOddRational.h index 56dff94d..656e9b2f 100644 --- a/Grid/qcd/action/pseudofermion/OneFlavourEvenOddRational.h +++ b/Grid/qcd/action/pseudofermion/OneFlavourEvenOddRational.h @@ -1,4 +1,3 @@ - /************************************************************************************* Grid physics library, www.github.com/paboyle/Grid @@ -43,8 +42,7 @@ NAMESPACE_BEGIN(Grid); // template -class OneFlavourEvenOddRationalPseudoFermionAction - : public Action { +class OneFlavourEvenOddRationalPseudoFermionAction : public Action { public: INHERIT_IMPL_TYPES(Impl); @@ -103,7 +101,7 @@ public: return sstream.str(); } - virtual void refresh(const GaugeField &U, GridParallelRNG &pRNG) { + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG &pRNG) { // P(phi) = e^{- phi^dag (MpcdagMpc)^-1/2 phi} // = e^{- phi^dag (MpcdagMpc)^-1/4 (MpcdagMpc)^-1/4 phi} // Phi = MpcdagMpc^{1/4} eta @@ -156,7 +154,10 @@ public: msCG(Mpc, PhiOdd, Y); - if ( (rand()%param.BoundsCheckFreq)==0 ) { + auto grid = FermOp.FermionGrid(); + auto r=rand(); + grid->Broadcast(0,r); + if ( (r%param.BoundsCheckFreq)==0 ) { FermionField gauss(FermOp.FermionRedBlackGrid()); gauss = PhiOdd; HighBoundCheck(Mpc,gauss,param.hi); diff --git a/Grid/qcd/action/pseudofermion/OneFlavourEvenOddRationalRatio.h b/Grid/qcd/action/pseudofermion/OneFlavourEvenOddRationalRatio.h index e5f0b602..e968b8e4 100644 --- a/Grid/qcd/action/pseudofermion/OneFlavourEvenOddRationalRatio.h +++ b/Grid/qcd/action/pseudofermion/OneFlavourEvenOddRationalRatio.h @@ -101,7 +101,7 @@ NAMESPACE_BEGIN(Grid); } - virtual void refresh(const GaugeField &U, GridParallelRNG& pRNG) { + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) { // S_f = chi^dag* P(V^dag*V)/Q(V^dag*V)* N(M^dag*M)/D(M^dag*M)* P(V^dag*V)/Q(V^dag*V)* chi // @@ -170,7 +170,10 @@ NAMESPACE_BEGIN(Grid); msCG_M(MdagM,X,Y); // Randomly apply rational bounds checks. - if ( (rand()%param.BoundsCheckFreq)==0 ) { + auto grid = NumOp.FermionGrid(); + auto r=rand(); + grid->Broadcast(0,r); + if ( (r%param.BoundsCheckFreq)==0 ) { FermionField gauss(NumOp.FermionRedBlackGrid()); gauss = PhiOdd; HighBoundCheck(MdagM,gauss,param.hi); diff --git a/Grid/qcd/action/pseudofermion/OneFlavourRational.h b/Grid/qcd/action/pseudofermion/OneFlavourRational.h index f6c823c9..aa647445 100644 --- a/Grid/qcd/action/pseudofermion/OneFlavourRational.h +++ b/Grid/qcd/action/pseudofermion/OneFlavourRational.h @@ -98,7 +98,7 @@ NAMESPACE_BEGIN(Grid); - virtual void refresh(const GaugeField &U, GridParallelRNG& pRNG) { + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) { // P(phi) = e^{- phi^dag (MdagM)^-1/2 phi} @@ -142,7 +142,10 @@ NAMESPACE_BEGIN(Grid); msCG(MdagMOp,Phi,Y); - if ( (rand()%param.BoundsCheckFreq)==0 ) { + auto grid = FermOp.FermionGrid(); + auto r=rand(); + grid->Broadcast(0,r); + if ( (r%param.BoundsCheckFreq)==0 ) { FermionField gauss(FermOp.FermionGrid()); gauss = Phi; HighBoundCheck(MdagMOp,gauss,param.hi); diff --git a/Grid/qcd/action/pseudofermion/OneFlavourRationalRatio.h b/Grid/qcd/action/pseudofermion/OneFlavourRationalRatio.h index 5fae2fe9..128c869a 100644 --- a/Grid/qcd/action/pseudofermion/OneFlavourRationalRatio.h +++ b/Grid/qcd/action/pseudofermion/OneFlavourRationalRatio.h @@ -95,7 +95,7 @@ NAMESPACE_BEGIN(Grid); } - virtual void refresh(const GaugeField &U, GridParallelRNG& pRNG) { + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) { // S_f = chi^dag* P(V^dag*V)/Q(V^dag*V)* N(M^dag*M)/D(M^dag*M)* P(V^dag*V)/Q(V^dag*V)* chi // @@ -156,7 +156,10 @@ NAMESPACE_BEGIN(Grid); msCG_M(MdagM,X,Y); // Randomly apply rational bounds checks. - if ( (rand()%param.BoundsCheckFreq)==0 ) { + auto grid = NumOp.FermionGrid(); + auto r=rand(); + grid->Broadcast(0,r); + if ( (r%param.BoundsCheckFreq)==0 ) { FermionField gauss(NumOp.FermionGrid()); gauss = Phi; HighBoundCheck(MdagM,gauss,param.hi); diff --git a/Grid/qcd/action/pseudofermion/TwoFlavour.h b/Grid/qcd/action/pseudofermion/TwoFlavour.h index f905a675..2ac97ddd 100644 --- a/Grid/qcd/action/pseudofermion/TwoFlavour.h +++ b/Grid/qcd/action/pseudofermion/TwoFlavour.h @@ -73,7 +73,7 @@ public: ////////////////////////////////////////////////////////////////////////////////////// // Push the gauge field in to the dops. Assume any BC's and smearing already applied ////////////////////////////////////////////////////////////////////////////////////// - virtual void refresh(const GaugeField &U, GridParallelRNG &pRNG) { + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG &pRNG) { // P(phi) = e^{- phi^dag (MdagM)^-1 phi} // Phi = Mdag eta // P(eta) = e^{- eta^dag eta} diff --git a/Grid/qcd/action/pseudofermion/TwoFlavourEvenOdd.h b/Grid/qcd/action/pseudofermion/TwoFlavourEvenOdd.h index a3cf8f08..2e5208a8 100644 --- a/Grid/qcd/action/pseudofermion/TwoFlavourEvenOdd.h +++ b/Grid/qcd/action/pseudofermion/TwoFlavourEvenOdd.h @@ -77,7 +77,7 @@ public: ////////////////////////////////////////////////////////////////////////////////////// // Push the gauge field in to the dops. Assume any BC's and smearing already applied ////////////////////////////////////////////////////////////////////////////////////// - virtual void refresh(const GaugeField &U, GridParallelRNG& pRNG) { + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) { // P(phi) = e^{- phi^dag (MpcdagMpc)^-1 phi} // Phi = McpDag eta diff --git a/Grid/qcd/action/pseudofermion/TwoFlavourEvenOddRatio.h b/Grid/qcd/action/pseudofermion/TwoFlavourEvenOddRatio.h index d1d6f336..da628c75 100644 --- a/Grid/qcd/action/pseudofermion/TwoFlavourEvenOddRatio.h +++ b/Grid/qcd/action/pseudofermion/TwoFlavourEvenOddRatio.h @@ -84,7 +84,7 @@ NAMESPACE_BEGIN(Grid); } - virtual void refresh(const GaugeField &U, GridParallelRNG& pRNG) { + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) { // P(phi) = e^{- phi^dag Vpc (MpcdagMpc)^-1 Vpcdag phi} // diff --git a/Grid/qcd/action/pseudofermion/TwoFlavourRatio.h b/Grid/qcd/action/pseudofermion/TwoFlavourRatio.h index 4d72faba..f584706d 100644 --- a/Grid/qcd/action/pseudofermion/TwoFlavourRatio.h +++ b/Grid/qcd/action/pseudofermion/TwoFlavourRatio.h @@ -64,7 +64,7 @@ public: return sstream.str(); } - virtual void refresh(const GaugeField &U, GridParallelRNG& pRNG) { + virtual void refresh(const GaugeField &U, GridSerialRNG &sRNG, GridParallelRNG& pRNG) { // P(phi) = e^{- phi^dag V (MdagM)^-1 Vdag phi} // diff --git a/Grid/qcd/action/scalar/ScalarAction.h b/Grid/qcd/action/scalar/ScalarAction.h index 34fc4fac..8b4f4f79 100644 --- a/Grid/qcd/action/scalar/ScalarAction.h +++ b/Grid/qcd/action/scalar/ScalarAction.h @@ -55,7 +55,7 @@ public: } virtual std::string action_name() {return "ScalarAction";} - virtual void refresh(const Field &U, GridParallelRNG &pRNG) {} // noop as no pseudoferms + virtual void refresh(const Field &U, GridSerialRNG &sRNG, GridParallelRNG &pRNG) {} // noop as no pseudoferms virtual RealD S(const Field &p) { return (mass_square * 0.5 + Nd) * ScalarObs::sumphisquared(p) + diff --git a/Grid/qcd/action/scalar/ScalarImpl.h b/Grid/qcd/action/scalar/ScalarImpl.h index 403ea573..13bd6c90 100644 --- a/Grid/qcd/action/scalar/ScalarImpl.h +++ b/Grid/qcd/action/scalar/ScalarImpl.h @@ -27,7 +27,7 @@ public: typedef Field FermionField; typedef Field PropagatorField; - static inline void generate_momenta(Field& P, GridParallelRNG& pRNG){ + static inline void generate_momenta(Field& P, GridSerialRNG &sRNG, GridParallelRNG& pRNG){ RealD scale = ::sqrt(HMC_MOMENTUM_DENOMINATOR); // CPS/UKQCD momentum rescaling gaussian(pRNG, P); P *= scale; @@ -151,7 +151,7 @@ public: out = one / out; } - static inline void generate_momenta(Field &P, GridParallelRNG &pRNG) + static inline void generate_momenta(Field &P, GridSerialRNG & sRNG, GridParallelRNG &pRNG) { RealD scale = ::sqrt(HMC_MOMENTUM_DENOMINATOR); // CPS/UKQCD momentum rescaling #ifndef USE_FFT_ACCELERATION diff --git a/Grid/qcd/action/scalar/ScalarInteractionAction.h b/Grid/qcd/action/scalar/ScalarInteractionAction.h index 5a5f9251..e04dd486 100644 --- a/Grid/qcd/action/scalar/ScalarInteractionAction.h +++ b/Grid/qcd/action/scalar/ScalarInteractionAction.h @@ -77,7 +77,7 @@ public: virtual std::string action_name() { return "ScalarAction"; } - virtual void refresh(const Field &U, GridParallelRNG &pRNG) {} + virtual void refresh(const Field &U, GridSerialRNG & sRNG, GridParallelRNG &pRNG) {} virtual RealD S(const Field &p) { diff --git a/Grid/qcd/hmc/HMC.h b/Grid/qcd/hmc/HMC.h index f168b69a..44674ea5 100644 --- a/Grid/qcd/hmc/HMC.h +++ b/Grid/qcd/hmc/HMC.h @@ -139,7 +139,7 @@ private: // Evolution ///////////////////////////////////////////////////////// RealD evolve_hmc_step(Field &U) { - TheIntegrator.refresh(U, pRNG); // set U and initialize P and phi's + TheIntegrator.refresh(U, sRNG, pRNG); // set U and initialize P and phi's RealD H0 = TheIntegrator.S(U); // initial state action diff --git a/Grid/qcd/hmc/integrators/Integrator.h b/Grid/qcd/hmc/integrators/Integrator.h index 77b7de52..aa28c6c8 100644 --- a/Grid/qcd/hmc/integrators/Integrator.h +++ b/Grid/qcd/hmc/integrators/Integrator.h @@ -236,10 +236,9 @@ public: // over the representations struct _refresh { template - void operator()(std::vector*> repr_set, Repr& Rep, - GridParallelRNG& pRNG) { + void operator()(std::vector*> repr_set, Repr& Rep, GridSerialRNG & sRNG, GridParallelRNG& pRNG) { for (int a = 0; a < repr_set.size(); ++a){ - repr_set.at(a)->refresh(Rep.U, pRNG); + repr_set.at(a)->refresh(Rep.U, sRNG, pRNG); std::cout << GridLogDebug << "Hirep refreshing pseudofermions" << std::endl; } @@ -247,12 +246,12 @@ public: } refresh_hireps{}; // Initialization of momenta and actions - void refresh(Field& U, GridParallelRNG& pRNG) + void refresh(Field& U, GridSerialRNG & sRNG, GridParallelRNG& pRNG) { assert(P.Grid() == U.Grid()); std::cout << GridLogIntegrator << "Integrator refresh\n"; - FieldImplementation::generate_momenta(P, pRNG); + FieldImplementation::generate_momenta(P, sRNG, pRNG); // Update the smeared fields, can be implemented as observer // necessary to keep the fields updated even after a reject @@ -269,11 +268,11 @@ public: // get gauge field from the SmearingPolicy and // based on the boolean is_smeared in actionID Field& Us = Smearer.get_U(as[level].actions.at(actionID)->is_smeared); - as[level].actions.at(actionID)->refresh(Us, pRNG); + as[level].actions.at(actionID)->refresh(Us, sRNG, pRNG); } // Refresh the higher representation actions - as[level].apply(refresh_hireps, Representations, pRNG); + as[level].apply(refresh_hireps, Representations, sRNG, pRNG); } MomFilter->applyFilter(P); diff --git a/benchmarks/Benchmark_dwf_fp32.cc b/benchmarks/Benchmark_dwf_fp32.cc index cb86177e..03f3ee61 100644 --- a/benchmarks/Benchmark_dwf_fp32.cc +++ b/benchmarks/Benchmark_dwf_fp32.cc @@ -53,7 +53,7 @@ int main (int argc, char ** argv) int threads = GridThread::GetThreads(); Coordinate latt4 = GridDefaultLatt(); - int Ls=8; + int Ls=16; for(int i=0;i> Ls;