diff --git a/Grid/Makefile.am b/Grid/Makefile.am index 7c3c151b..8472dd71 100644 --- a/Grid/Makefile.am +++ b/Grid/Makefile.am @@ -66,6 +66,10 @@ if BUILD_FERMION_REPS extra_sources+=$(ADJ_FERMION_FILES) extra_sources+=$(TWOIND_FERMION_FILES) endif +if BUILD_SP + extra_sources+=$(SP_FERMION_FILES) + extra_sources+=$(SP_TWOIND_FERMION_FILES) +endif lib_LIBRARIES = libGrid.a diff --git a/Grid/communicator/SharedMemoryMPI.cc b/Grid/communicator/SharedMemoryMPI.cc index 89a9d316..64a86c4b 100644 --- a/Grid/communicator/SharedMemoryMPI.cc +++ b/Grid/communicator/SharedMemoryMPI.cc @@ -604,8 +604,8 @@ void GlobalSharedMemory::SharedMemoryAllocate(uint64_t bytes, int flags) #ifdef GRID_SYCL_LEVEL_ZERO_IPC typedef struct { int fd; pid_t pid ; ze_ipc_mem_handle_t ze; } clone_mem_t; - auto zeDevice = cl::sycl::get_native(theGridAccelerator->get_device()); - auto zeContext = cl::sycl::get_native(theGridAccelerator->get_context()); + auto zeDevice = cl::sycl::get_native(theGridAccelerator->get_device()); + auto zeContext = cl::sycl::get_native(theGridAccelerator->get_context()); ze_ipc_mem_handle_t ihandle; clone_mem_t handle; diff --git a/Grid/lattice/Lattice.h b/Grid/lattice/Lattice.h index c4adf86a..6343db99 100644 --- a/Grid/lattice/Lattice.h +++ b/Grid/lattice/Lattice.h @@ -47,3 +47,4 @@ Author: Peter Boyle #include #include #include +#include diff --git a/Grid/lattice/Lattice_ET.h b/Grid/lattice/Lattice_ET.h index fdd31b28..2ad93d30 100644 --- a/Grid/lattice/Lattice_ET.h +++ b/Grid/lattice/Lattice_ET.h @@ -345,7 +345,9 @@ GridUnopClass(UnaryNot, Not(a)); GridUnopClass(UnaryTrace, trace(a)); GridUnopClass(UnaryTranspose, transpose(a)); GridUnopClass(UnaryTa, Ta(a)); +GridUnopClass(UnarySpTa, SpTa(a)); GridUnopClass(UnaryProjectOnGroup, ProjectOnGroup(a)); +GridUnopClass(UnaryProjectOnSpGroup, ProjectOnSpGroup(a)); GridUnopClass(UnaryTimesI, timesI(a)); GridUnopClass(UnaryTimesMinusI, timesMinusI(a)); GridUnopClass(UnaryAbs, abs(a)); @@ -456,7 +458,9 @@ GRID_DEF_UNOP(operator!, UnaryNot); GRID_DEF_UNOP(trace, UnaryTrace); GRID_DEF_UNOP(transpose, UnaryTranspose); GRID_DEF_UNOP(Ta, UnaryTa); +GRID_DEF_UNOP(SpTa, UnarySpTa); GRID_DEF_UNOP(ProjectOnGroup, UnaryProjectOnGroup); +GRID_DEF_UNOP(ProjectOnSpGroup, UnaryProjectOnSpGroup); GRID_DEF_UNOP(timesI, UnaryTimesI); GRID_DEF_UNOP(timesMinusI, UnaryTimesMinusI); GRID_DEF_UNOP(abs, UnaryAbs); // abs overloaded in cmath C++98; DON'T do the diff --git a/Grid/lattice/Lattice_trace.h b/Grid/lattice/Lattice_trace.h index 5fe5173d..9a0c6f96 100644 --- a/Grid/lattice/Lattice_trace.h +++ b/Grid/lattice/Lattice_trace.h @@ -66,6 +66,65 @@ inline auto TraceIndex(const Lattice &lhs) -> Lattice +Lattice > > > Determinant(const Lattice > > > &Umu) +{ + GridBase *grid=Umu.Grid(); + auto lvol = grid->lSites(); + Lattice > > > ret(grid); + typedef typename Vec::scalar_type scalar; + autoView(Umu_v,Umu,CpuRead); + autoView(ret_v,ret,CpuWrite); + thread_for(site,lvol,{ + Eigen::MatrixXcd EigenU = Eigen::MatrixXcd::Zero(N,N); + Coordinate lcoor; + grid->LocalIndexToLocalCoor(site, lcoor); + iScalar > > Us; + peekLocalSite(Us, Umu_v, lcoor); + for(int i=0;i +Lattice > > > Inverse(const Lattice > > > &Umu) +{ + GridBase *grid=Umu.Grid(); + auto lvol = grid->lSites(); + Lattice > > > ret(grid); + + autoView(Umu_v,Umu,CpuRead); + autoView(ret_v,ret,CpuWrite); + thread_for(site,lvol,{ + Eigen::MatrixXcd EigenU = Eigen::MatrixXcd::Zero(N,N); + Coordinate lcoor; + grid->LocalIndexToLocalCoor(site, lcoor); + iScalar > > Us; + iScalar > > Ui; + peekLocalSite(Us, Umu_v, lcoor); + for(int i=0;i &From,Lattice & To,Coordinate Fro for(int d=0;d_processors[d] == Tg->_processors[d]); } - // the above should guarantee that the operations are local + +#if 1 + + size_t nsite = 1; + for(int i=0;ioIndex(from_coor); + int fiidx = Fg->iIndex(from_coor); + int toidx = Tg->oIndex(to_coor); + int tiidx = Tg->iIndex(to_coor); + int* tt = table + 4*idx; + tt[0] = foidx; + tt[1] = fiidx; + tt[2] = toidx; + tt[3] = tiidx; + }); + + int* table_d = (int*)acceleratorAllocDevice(tbytes); + acceleratorCopyToDevice(table,table_d,tbytes); + + typedef typename vobj::vector_type vector_type; + typedef typename vobj::scalar_type scalar_type; + + autoView(from_v,From,AcceleratorRead); + autoView(to_v,To,AcceleratorWrite); + + accelerator_for(idx,nsite,1,{ + static const int words=sizeof(vobj)/sizeof(vector_type); + int* tt = table_d + 4*idx; + int from_oidx = *tt++; + int from_lane = *tt++; + int to_oidx = *tt++; + int to_lane = *tt; + + const vector_type* from = (const vector_type *)&from_v[from_oidx]; + vector_type* to = (vector_type *)&to_v[to_oidx]; + + scalar_type stmp; + for(int w=0;w_ldimensions; Coordinate rdf = Fg->_rdimensions; Coordinate isf = Fg->_istride; @@ -738,6 +798,8 @@ void localCopyRegion(const Lattice &From,Lattice & To,Coordinate Fro #endif } }); + +#endif } @@ -830,6 +892,8 @@ void ExtractSlice(Lattice &lowDim,const Lattice & higherDim,int slic } +//Insert subvolume orthogonal to direction 'orthog' with slice index 'slice_lo' from 'lowDim' onto slice index 'slice_hi' of higherDim +//The local dimensions of both 'lowDim' and 'higherDim' orthogonal to 'orthog' should be the same template void InsertSliceLocal(const Lattice &lowDim, Lattice & higherDim,int slice_lo,int slice_hi, int orthog) { @@ -851,6 +915,65 @@ void InsertSliceLocal(const Lattice &lowDim, Lattice & higherDim,int } } +#if 1 + size_t nsite = lg->lSites()/lg->LocalDimensions()[orthog]; + size_t tbytes = 4*nsite*sizeof(int); + int *table = (int*)malloc(tbytes); + + thread_for(idx,nsite,{ + Coordinate lcoor(nl); + Coordinate hcoor(nh); + lcoor[orthog] = slice_lo; + hcoor[orthog] = slice_hi; + size_t rem = idx; + for(int mu=0;muLocalDimensions()[mu]; rem /= lg->LocalDimensions()[mu]; + lcoor[mu] = hcoor[mu] = xmu; + } + } + int loidx = lg->oIndex(lcoor); + int liidx = lg->iIndex(lcoor); + int hoidx = hg->oIndex(hcoor); + int hiidx = hg->iIndex(hcoor); + int* tt = table + 4*idx; + tt[0] = loidx; + tt[1] = liidx; + tt[2] = hoidx; + tt[3] = hiidx; + }); + + int* table_d = (int*)acceleratorAllocDevice(tbytes); + acceleratorCopyToDevice(table,table_d,tbytes); + + typedef typename vobj::vector_type vector_type; + typedef typename vobj::scalar_type scalar_type; + + autoView(lowDim_v,lowDim,AcceleratorRead); + autoView(higherDim_v,higherDim,AcceleratorWrite); + + accelerator_for(idx,nsite,1,{ + static const int words=sizeof(vobj)/sizeof(vector_type); + int* tt = table_d + 4*idx; + int from_oidx = *tt++; + int from_lane = *tt++; + int to_oidx = *tt++; + int to_lane = *tt; + + const vector_type* from = (const vector_type *)&lowDim_v[from_oidx]; + vector_type* to = (vector_type *)&higherDim_v[to_oidx]; + + scalar_type stmp; + for(int w=0;w &lowDim, Lattice & higherDim,int pokeLocalSite(s,higherDimv,hcoor); } }); +#endif } diff --git a/Grid/lattice/PaddedCell.h b/Grid/lattice/PaddedCell.h index b9788ad8..500b97de 100644 --- a/Grid/lattice/PaddedCell.h +++ b/Grid/lattice/PaddedCell.h @@ -26,14 +26,32 @@ Author: Peter Boyle pboyle@bnl.gov /* END LEGAL */ #pragma once +#include + NAMESPACE_BEGIN(Grid); +//Allow the user to specify how the C-shift is performed, e.g. to respect the appropriate boundary conditions +template +struct CshiftImplBase{ + virtual Lattice Cshift(const Lattice &in, int dir, int shift) const = 0; + virtual ~CshiftImplBase(){} +}; +template +struct CshiftImplDefault: public CshiftImplBase{ + Lattice Cshift(const Lattice &in, int dir, int shift) const override{ return Grid::Cshift(in,dir,shift); } +}; +template +struct CshiftImplGauge: public CshiftImplBase{ + typename Gimpl::GaugeLinkField Cshift(const typename Gimpl::GaugeLinkField &in, int dir, int shift) const override{ return Gimpl::CshiftLink(in,dir,shift); } +}; + class PaddedCell { public: GridCartesian * unpadded_grid; int dims; int depth; std::vector grids; + ~PaddedCell() { DeleteGrids(); @@ -78,7 +96,7 @@ public: } }; template - inline Lattice Extract(Lattice &in) + inline Lattice Extract(const Lattice &in) const { Lattice out(unpadded_grid); @@ -89,19 +107,19 @@ public: return out; } template - inline Lattice Exchange(Lattice &in) + inline Lattice Exchange(const Lattice &in, const CshiftImplBase &cshift = CshiftImplDefault()) const { GridBase *old_grid = in.Grid(); int dims = old_grid->Nd(); Lattice tmp = in; for(int d=0;d - inline Lattice Expand(int dim,Lattice &in) + inline Lattice Expand(int dim, const Lattice &in, const CshiftImplBase &cshift = CshiftImplDefault()) const { GridBase *old_grid = in.Grid(); GridCartesian *new_grid = grids[dim];//These are new grids @@ -112,21 +130,41 @@ public: if(dim==0) conformable(old_grid,unpadded_grid); else conformable(old_grid,grids[dim-1]); - // std::cout << " dim "< WilsonTwoIndexSymmetricFermi typedef WilsonFermion WilsonTwoIndexAntiSymmetricFermionF; typedef WilsonFermion WilsonTwoIndexAntiSymmetricFermionD; +// Sp(2n) +typedef WilsonFermion SpWilsonFermionF; +typedef WilsonFermion SpWilsonFermionD; + +typedef WilsonFermion SpWilsonTwoIndexAntiSymmetricFermionF; +typedef WilsonFermion SpWilsonTwoIndexAntiSymmetricFermionD; + +typedef WilsonFermion SpWilsonTwoIndexSymmetricFermionF; +typedef WilsonFermion SpWilsonTwoIndexSymmetricFermionD; + // Twisted mass fermion typedef WilsonTMFermion WilsonTMFermionD2; typedef WilsonTMFermion WilsonTMFermionF; diff --git a/Grid/qcd/action/fermion/WilsonImpl.h b/Grid/qcd/action/fermion/WilsonImpl.h index c7180115..07248160 100644 --- a/Grid/qcd/action/fermion/WilsonImpl.h +++ b/Grid/qcd/action/fermion/WilsonImpl.h @@ -261,6 +261,22 @@ typedef WilsonImpl W typedef WilsonImpl WilsonTwoIndexAntiSymmetricImplF; // Float typedef WilsonImpl WilsonTwoIndexAntiSymmetricImplD; // Double +//sp 2n + +typedef WilsonImpl SpWilsonImplR; // Real.. whichever prec +typedef WilsonImpl SpWilsonImplF; // Float +typedef WilsonImpl SpWilsonImplD; // Double + +typedef WilsonImpl SpWilsonTwoIndexAntiSymmetricImplR; // Real.. whichever prec +typedef WilsonImpl SpWilsonTwoIndexAntiSymmetricImplF; // Float +typedef WilsonImpl SpWilsonTwoIndexAntiSymmetricImplD; // Double + +typedef WilsonImpl SpWilsonTwoIndexSymmetricImplR; // Real.. whichever prec +typedef WilsonImpl SpWilsonTwoIndexSymmetricImplF; // Float +typedef WilsonImpl SpWilsonTwoIndexSymmetricImplD; // Double + +typedef WilsonImpl SpWilsonAdjImplR; // Real.. whichever prec // adj = 2indx symmetric for Sp(2N) +typedef WilsonImpl SpWilsonAdjImplF; // Float // adj = 2indx symmetric for Sp(2N) +typedef WilsonImpl SpWilsonAdjImplD; // Double // adj = 2indx symmetric for Sp(2N) NAMESPACE_END(Grid); - diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonCloverFermionInstantiationSpWilsonImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonCloverFermionInstantiationSpWilsonImplD.cc new file mode 120000 index 00000000..9cc05107 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonCloverFermionInstantiationSpWilsonImplD.cc @@ -0,0 +1 @@ +../WilsonCloverFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonFermionInstantiationSpWilsonImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonFermionInstantiationSpWilsonImplD.cc new file mode 120000 index 00000000..5f6ab65e --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonFermionInstantiationSpWilsonImplD.cc @@ -0,0 +1 @@ +../WilsonFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonKernelsInstantiationSpWilsonImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonKernelsInstantiationSpWilsonImplD.cc new file mode 120000 index 00000000..01c35e7b --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonKernelsInstantiationSpWilsonImplD.cc @@ -0,0 +1 @@ +../WilsonKernelsInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonTMFermionInstantiationSpWilsonImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonTMFermionInstantiationSpWilsonImplD.cc new file mode 120000 index 00000000..d5789bcf --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/WilsonTMFermionInstantiationSpWilsonImplD.cc @@ -0,0 +1 @@ +../WilsonTMFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/impl.h b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/impl.h new file mode 100644 index 00000000..5aa843dd --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplD/impl.h @@ -0,0 +1 @@ +#define IMPLEMENTATION SpWilsonImplD diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonCloverFermionInstantiationSpWilsonImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonCloverFermionInstantiationSpWilsonImplF.cc new file mode 120000 index 00000000..9cc05107 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonCloverFermionInstantiationSpWilsonImplF.cc @@ -0,0 +1 @@ +../WilsonCloverFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonFermionInstantiationSpWilsonImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonFermionInstantiationSpWilsonImplF.cc new file mode 120000 index 00000000..5f6ab65e --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonFermionInstantiationSpWilsonImplF.cc @@ -0,0 +1 @@ +../WilsonFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonKernelsInstantiationSpWilsonImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonKernelsInstantiationSpWilsonImplF.cc new file mode 120000 index 00000000..01c35e7b --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonKernelsInstantiationSpWilsonImplF.cc @@ -0,0 +1 @@ +../WilsonKernelsInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonTMFermionInstantiationSpWilsonImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonTMFermionInstantiationSpWilsonImplF.cc new file mode 120000 index 00000000..d5789bcf --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/WilsonTMFermionInstantiationSpWilsonImplF.cc @@ -0,0 +1 @@ +../WilsonTMFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/impl.h b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/impl.h new file mode 100644 index 00000000..761b0c82 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonImplF/impl.h @@ -0,0 +1 @@ +#define IMPLEMENTATION SpWilsonImplF diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonCloverFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonCloverFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc new file mode 120000 index 00000000..9cc05107 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonCloverFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc @@ -0,0 +1 @@ +../WilsonCloverFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc new file mode 120000 index 00000000..5f6ab65e --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc @@ -0,0 +1 @@ +../WilsonFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonKernelsInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonKernelsInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc new file mode 120000 index 00000000..01c35e7b --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonKernelsInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc @@ -0,0 +1 @@ +../WilsonKernelsInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonTMFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonTMFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc new file mode 120000 index 00000000..d5789bcf --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/WilsonTMFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplD.cc @@ -0,0 +1 @@ +../WilsonTMFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/impl.h b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/impl.h new file mode 100644 index 00000000..4c546322 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplD/impl.h @@ -0,0 +1 @@ +#define IMPLEMENTATION SpWilsonTwoIndexAntiSymmetricImplD diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonCloverFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonCloverFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc new file mode 120000 index 00000000..9cc05107 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonCloverFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc @@ -0,0 +1 @@ +../WilsonCloverFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc new file mode 120000 index 00000000..5f6ab65e --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc @@ -0,0 +1 @@ +../WilsonFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonKernelsInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonKernelsInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc new file mode 120000 index 00000000..01c35e7b --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonKernelsInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc @@ -0,0 +1 @@ +../WilsonKernelsInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonTMFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonTMFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc new file mode 120000 index 00000000..d5789bcf --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/WilsonTMFermionInstantiationSpWilsonTwoIndexAntiSymmetricImplF.cc @@ -0,0 +1 @@ +../WilsonTMFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/impl.h b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/impl.h new file mode 100644 index 00000000..8425236e --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexAntiSymmetricImplF/impl.h @@ -0,0 +1 @@ +#define IMPLEMENTATION SpWilsonTwoIndexAntiSymmetricImplF diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonCloverFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonCloverFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc new file mode 120000 index 00000000..9cc05107 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonCloverFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc @@ -0,0 +1 @@ +../WilsonCloverFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc new file mode 120000 index 00000000..5f6ab65e --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc @@ -0,0 +1 @@ +../WilsonFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonKernelsInstantiationSpWilsonTwoIndexSymmetricImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonKernelsInstantiationSpWilsonTwoIndexSymmetricImplD.cc new file mode 120000 index 00000000..01c35e7b --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonKernelsInstantiationSpWilsonTwoIndexSymmetricImplD.cc @@ -0,0 +1 @@ +../WilsonKernelsInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonTMFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonTMFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc new file mode 120000 index 00000000..d5789bcf --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/WilsonTMFermionInstantiationSpWilsonTwoIndexSymmetricImplD.cc @@ -0,0 +1 @@ +../WilsonTMFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/impl.h b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/impl.h new file mode 100644 index 00000000..ac5babca --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplD/impl.h @@ -0,0 +1 @@ +#define IMPLEMENTATION SpWilsonTwoIndexSymmetricImplD diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonCloverFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonCloverFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc new file mode 120000 index 00000000..9cc05107 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonCloverFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc @@ -0,0 +1 @@ +../WilsonCloverFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc new file mode 120000 index 00000000..5f6ab65e --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc @@ -0,0 +1 @@ +../WilsonFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonKernelsInstantiationSpWilsonTwoIndexSymmetricImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonKernelsInstantiationSpWilsonTwoIndexSymmetricImplF.cc new file mode 120000 index 00000000..01c35e7b --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonKernelsInstantiationSpWilsonTwoIndexSymmetricImplF.cc @@ -0,0 +1 @@ +../WilsonKernelsInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonTMFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonTMFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc new file mode 120000 index 00000000..d5789bcf --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/WilsonTMFermionInstantiationSpWilsonTwoIndexSymmetricImplF.cc @@ -0,0 +1 @@ +../WilsonTMFermionInstantiation.cc.master \ No newline at end of file diff --git a/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/impl.h b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/impl.h new file mode 100644 index 00000000..bea7dce7 --- /dev/null +++ b/Grid/qcd/action/fermion/instantiation/SpWilsonTwoIndexSymmetricImplF/impl.h @@ -0,0 +1 @@ +#define IMPLEMENTATION SpWilsonTwoIndexSymmetricImplF diff --git a/Grid/qcd/action/fermion/instantiation/generate_instantiations.sh b/Grid/qcd/action/fermion/instantiation/generate_instantiations.sh index 4ccc01e8..728dc5e7 100755 --- a/Grid/qcd/action/fermion/instantiation/generate_instantiations.sh +++ b/Grid/qcd/action/fermion/instantiation/generate_instantiations.sh @@ -10,12 +10,18 @@ WILSON_IMPL_LIST=" \ WilsonImplF \ WilsonImplD \ WilsonImplD2 \ + SpWilsonImplF \ + SpWilsonImplD \ WilsonAdjImplF \ WilsonAdjImplD \ WilsonTwoIndexSymmetricImplF \ WilsonTwoIndexSymmetricImplD \ WilsonTwoIndexAntiSymmetricImplF \ WilsonTwoIndexAntiSymmetricImplD \ + SpWilsonTwoIndexAntiSymmetricImplF \ + SpWilsonTwoIndexAntiSymmetricImplD \ + SpWilsonTwoIndexSymmetricImplF \ + SpWilsonTwoIndexSymmetricImplD \ GparityWilsonImplF \ GparityWilsonImplD " diff --git a/Grid/qcd/action/gauge/Gauge.h b/Grid/qcd/action/gauge/Gauge.h index b3b34eba..7a72d087 100644 --- a/Grid/qcd/action/gauge/Gauge.h +++ b/Grid/qcd/action/gauge/Gauge.h @@ -39,6 +39,9 @@ NAMESPACE_BEGIN(Grid); typedef WilsonGaugeAction WilsonGaugeActionR; typedef WilsonGaugeAction WilsonGaugeActionF; typedef WilsonGaugeAction WilsonGaugeActionD; +typedef WilsonGaugeAction SpWilsonGaugeActionR; +typedef WilsonGaugeAction SpWilsonGaugeActionF; +typedef WilsonGaugeAction SpWilsonGaugeActionD; typedef PlaqPlusRectangleAction PlaqPlusRectangleActionR; typedef PlaqPlusRectangleAction PlaqPlusRectangleActionF; typedef PlaqPlusRectangleAction PlaqPlusRectangleActionD; diff --git a/Grid/qcd/action/gauge/GaugeImplTypes.h b/Grid/qcd/action/gauge/GaugeImplTypes.h index 5c9405ab..a9af1fae 100644 --- a/Grid/qcd/action/gauge/GaugeImplTypes.h +++ b/Grid/qcd/action/gauge/GaugeImplTypes.h @@ -61,7 +61,7 @@ NAMESPACE_BEGIN(Grid); typedef typename Impl::Field Field; // hardcodes the exponential approximation in the template -template class GaugeImplTypes { +template > class GaugeImplTypes { public: typedef S Simd; typedef typename Simd::scalar_type scalar_type; @@ -78,8 +78,6 @@ public: typedef Lattice LinkField; typedef Lattice Field; - typedef SU Group; - // Guido: we can probably separate the types from the HMC functions // this will create 2 kind of implementations // probably confusing the users @@ -119,6 +117,7 @@ public: // LinkField Pmu(P.Grid()); Pmu = Zero(); + for (int mu = 0; mu < Nd; mu++) { Group::GaussianFundamentalLieAlgebraMatrix(pRNG, Pmu); RealD scale = ::sqrt(HMC_MOMENTUM_DENOMINATOR) ; @@ -126,8 +125,12 @@ public: PokeIndex(P, Pmu, mu); } } - - static inline Field projectForce(Field &P) { return Ta(P); } + + static inline Field projectForce(Field &P) { + Field ret(P.Grid()); + Group::taProj(P, ret); + return ret; + } static inline void update_field(Field& P, Field& U, double ep){ //static std::chrono::duration diff; @@ -137,14 +140,15 @@ public: autoView(P_v,P,AcceleratorRead); accelerator_for(ss, P.Grid()->oSites(),1,{ for (int mu = 0; mu < Nd; mu++) { - U_v[ss](mu) = ProjectOnGroup(Exponentiate(P_v[ss](mu), ep, Nexp) * U_v[ss](mu)); + U_v[ss](mu) = Exponentiate(P_v[ss](mu), ep, Nexp) * U_v[ss](mu); + U_v[ss](mu) = Group::ProjectOnGeneralGroup(U_v[ss](mu)); } }); //auto end = std::chrono::high_resolution_clock::now(); // diff += end - start; // std::cout << "Time to exponentiate matrix " << diff.count() << " s\n"; } - + static inline RealD FieldSquareNorm(Field& U){ LatticeComplex Hloc(U.Grid()); Hloc = Zero(); @@ -157,7 +161,7 @@ public: } static inline void Project(Field &U) { - ProjectSUn(U); + Group::ProjectOnSpecialGroup(U); } static inline void HotConfiguration(GridParallelRNG &pRNG, Field &U) { @@ -171,6 +175,7 @@ public: static inline void ColdConfiguration(GridParallelRNG &pRNG, Field &U) { Group::ColdConfiguration(pRNG, U); } + }; @@ -178,10 +183,17 @@ typedef GaugeImplTypes GimplTypesR; typedef GaugeImplTypes GimplTypesF; typedef GaugeImplTypes GimplTypesD; +typedef GaugeImplTypes > SpGimplTypesR; +typedef GaugeImplTypes > SpGimplTypesF; +typedef GaugeImplTypes > SpGimplTypesD; + typedef GaugeImplTypes::AdjointDimension> GimplAdjointTypesR; typedef GaugeImplTypes::AdjointDimension> GimplAdjointTypesF; typedef GaugeImplTypes::AdjointDimension> GimplAdjointTypesD; + + + NAMESPACE_END(Grid); #endif // GRID_GAUGE_IMPL_TYPES_H diff --git a/Grid/qcd/action/gauge/GaugeImplementations.h b/Grid/qcd/action/gauge/GaugeImplementations.h index f518b236..312e889c 100644 --- a/Grid/qcd/action/gauge/GaugeImplementations.h +++ b/Grid/qcd/action/gauge/GaugeImplementations.h @@ -176,7 +176,7 @@ public: return PeriodicBC::CshiftLink(Link,mu,shift); } - static inline void setDirections(std::vector &conjDirs) { _conjDirs=conjDirs; } + static inline void setDirections(const std::vector &conjDirs) { _conjDirs=conjDirs; } static inline std::vector getDirections(void) { return _conjDirs; } static inline bool isPeriodicGaugeField(void) { return false; } }; @@ -193,6 +193,11 @@ typedef ConjugateGaugeImpl ConjugateGimplR; // Real.. whichever pre typedef ConjugateGaugeImpl ConjugateGimplF; // Float typedef ConjugateGaugeImpl ConjugateGimplD; // Double +typedef PeriodicGaugeImpl SpPeriodicGimplR; // Real.. whichever prec +typedef PeriodicGaugeImpl SpPeriodicGimplF; // Float +typedef PeriodicGaugeImpl SpPeriodicGimplD; // Double + + NAMESPACE_END(Grid); #endif diff --git a/Grid/qcd/action/gauge/PlaqPlusRectangleAction.h b/Grid/qcd/action/gauge/PlaqPlusRectangleAction.h index 7690092d..b9d6ac16 100644 --- a/Grid/qcd/action/gauge/PlaqPlusRectangleAction.h +++ b/Grid/qcd/action/gauge/PlaqPlusRectangleAction.h @@ -43,7 +43,7 @@ public: private: RealD c_plaq; RealD c_rect; - + typename WilsonLoops::StapleAndRectStapleAllWorkspace workspace; public: PlaqPlusRectangleAction(RealD b,RealD c): c_plaq(b),c_rect(c){}; @@ -79,27 +79,18 @@ public: GridBase *grid = Umu.Grid(); std::vector U (Nd,grid); - std::vector U2(Nd,grid); - for(int mu=0;mu(Umu,mu); - WilsonLoops::RectStapleDouble(U2[mu],U[mu],mu); } + std::vector RectStaple(Nd,grid), Staple(Nd,grid); + WilsonLoops::StapleAndRectStapleAll(Staple, RectStaple, U, workspace); GaugeLinkField dSdU_mu(grid); GaugeLinkField staple(grid); for (int mu=0; mu < Nd; mu++){ - - // Staple in direction mu - - WilsonLoops::Staple(staple,Umu,mu); - - dSdU_mu = Ta(U[mu]*staple)*factor_p; - - WilsonLoops::RectStaple(Umu,staple,U2,U,mu); - - dSdU_mu = dSdU_mu + Ta(U[mu]*staple)*factor_r; + dSdU_mu = Ta(U[mu]*Staple[mu])*factor_p; + dSdU_mu = dSdU_mu + Ta(U[mu]*RectStaple[mu])*factor_r; PokeIndex(dSdU, dSdU_mu, mu); } diff --git a/Grid/qcd/hmc/GenericHMCrunner.h b/Grid/qcd/hmc/GenericHMCrunner.h index 727b3e24..1429d848 100644 --- a/Grid/qcd/hmc/GenericHMCrunner.h +++ b/Grid/qcd/hmc/GenericHMCrunner.h @@ -225,6 +225,18 @@ template ; +// sp2n + +template