Bland SYCL compile

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<cl::sycl::backend::level_zero>(theGridAccelerator->get_device());
+    auto zeDevice    = cl::sycl::get_native<cl::sycl::backend::ext_oneapi_level_zero>(theGridAccelerator->get_device());
-    auto zeContext   = cl::sycl::get_native<cl::sycl::backend::level_zero>(theGridAccelerator->get_context());
+    auto zeContext   = cl::sycl::get_native<cl::sycl::backend::ext_oneapi_level_zero>(theGridAccelerator->get_context());
       
     ze_ipc_mem_handle_t ihandle;
     clone_mem_t handle;

Grid/qcd/action/fermion/CayleyFermion5D.h

@@ -124,11 +124,6 @@ public:
   RealD                _b;
   RealD                _c;
 
-  // possible boost
-  std::vector<ComplexD> qmu;
-  void set_qmu(std::vector<ComplexD> _qmu) { qmu=_qmu; assert(qmu.size()==Nd);};
-  void addQmu(const FermionField &in, FermionField &out, int dag);
-  
   // Cayley form Moebius (tanh and zolotarev)
   Vector<Coeff_t> omega;
   Vector<Coeff_t> bs;    // S dependent coeffs

Grid/qcd/action/fermion/ContinuedFractionFermion5D.h

@@ -60,50 +60,6 @@ public:
   //      virtual void   Instantiatable(void)=0;
   virtual void   Instantiatable(void) =0;
 
-  void FreePropagator(const FermionField &in,FermionField &out,RealD mass,std::vector<Complex> boundary, std::vector<double> twist)
-  {
-    std::cout << "Free Propagator for PartialFraction"<<std::endl;
-    FermionField in_k(in.Grid());
-    FermionField prop_k(in.Grid());
-    
-    FFT theFFT((GridCartesian *) in.Grid());
-
-    //phase for boundary condition
-    ComplexField coor(in.Grid());
-    ComplexField ph(in.Grid());  ph = Zero();
-    FermionField in_buf(in.Grid()); in_buf = Zero();
-    typedef typename Simd::scalar_type Scalar;
-    Scalar ci(0.0,1.0);
-    assert(twist.size() == Nd);//check that twist is Nd
-    assert(boundary.size() == Nd);//check that boundary conditions is Nd
-    int shift = 0;
-    for(unsigned int nu = 0; nu < Nd; nu++)
-      {
-	// Shift coordinate lattice index by 1 to account for 5th dimension.
-	LatticeCoordinate(coor, nu + shift);
-	double boundary_phase = ::acos(real(boundary[nu]));
-	ph = ph + boundary_phase*coor*((1./(in.Grid()->_fdimensions[nu+shift])));
-	//momenta for propagator shifted by twist+boundary
-	twist[nu] = twist[nu] + boundary_phase/((2.0*M_PI));
-      }
-    in_buf = exp(ci*ph*(-1.0))*in;
-
-    theFFT.FFT_all_dim(in_k,in,FFT::forward);
-    this->MomentumSpacePropagatorHw(prop_k,in_k,mass,twist);
-    theFFT.FFT_all_dim(out,prop_k,FFT::backward);
-    
-    //phase for boundary condition
-    out = out * exp(ci*ph);
-  };
-
-  virtual void FreePropagator(const FermionField &in,FermionField &out,RealD mass) {
-    std::vector<double> twist(Nd,0.0); //default: periodic boundarys in all directions
-    std::vector<Complex> boundary;
-    for(int i=0;i<Nd;i++) boundary.push_back(1);//default: periodic boundary conditions
-    FreePropagator(in,out,mass,boundary,twist);
-  };
-
-  
   // Efficient support for multigrid coarsening
   virtual void  Mdir (const FermionField &in, FermionField &out,int dir,int disp);
   virtual void  MdirAll(const FermionField &in, std::vector<FermionField> &out);

Grid/qcd/action/fermion/PartialFractionFermion5D.h

@@ -39,7 +39,7 @@ class PartialFractionFermion5D : public WilsonFermion5D<Impl>
 public:
   INHERIT_IMPL_TYPES(Impl);
 
-  const int part_frac_chroma_convention=0;
+  const int part_frac_chroma_convention=1;
 
   void   Meooe_internal(const FermionField &in, FermionField &out,int dag);
   void   Mooee_internal(const FermionField &in, FermionField &out,int dag);
@@ -83,63 +83,12 @@ public:
 			   GridRedBlackCartesian &FourDimRedBlackGrid,
 			   RealD _mass,RealD M5,const ImplParams &p= ImplParams());
 
-  PartialFractionFermion5D(GaugeField &_Umu,
-			   GridCartesian         &FiveDimGrid,
-			   GridRedBlackCartesian &FiveDimRedBlackGrid,
-			   GridCartesian         &FourDimGrid,
-			   GridRedBlackCartesian &FourDimRedBlackGrid,
-			   RealD _mass,RealD M5,std::vector<RealD> &_qmu,const ImplParams &p= ImplParams());
-
-  void FreePropagator(const FermionField &in,FermionField &out,RealD mass,std::vector<Complex> boundary, std::vector<double> twist)
-  {
-    std::cout << "Free Propagator for PartialFraction"<<std::endl;
-    FermionField in_k(in.Grid());
-    FermionField prop_k(in.Grid());
-    
-    FFT theFFT((GridCartesian *) in.Grid());
-
-    //phase for boundary condition
-    ComplexField coor(in.Grid());
-    ComplexField ph(in.Grid());  ph = Zero();
-    FermionField in_buf(in.Grid()); in_buf = Zero();
-    typedef typename Simd::scalar_type Scalar;
-    Scalar ci(0.0,1.0);
-    assert(twist.size() == Nd);//check that twist is Nd
-    assert(boundary.size() == Nd);//check that boundary conditions is Nd
-    int shift = 0;
-    for(unsigned int nu = 0; nu < Nd; nu++)
-      {
-	// Shift coordinate lattice index by 1 to account for 5th dimension.
-	LatticeCoordinate(coor, nu + shift);
-	double boundary_phase = ::acos(real(boundary[nu]));
-	ph = ph + boundary_phase*coor*((1./(in.Grid()->_fdimensions[nu+shift])));
-	//momenta for propagator shifted by twist+boundary
-	twist[nu] = twist[nu] + boundary_phase/((2.0*M_PI));
-      }
-    in_buf = exp(ci*ph*(-1.0))*in;
-
-    theFFT.FFT_all_dim(in_k,in,FFT::forward);
-    this->MomentumSpacePropagatorHw(prop_k,in_k,mass,twist);
-    theFFT.FFT_all_dim(out,prop_k,FFT::backward);
-    
-    //phase for boundary condition
-    out = out * exp(ci*ph);
-  };
-
-  virtual void FreePropagator(const FermionField &in,FermionField &out,RealD mass) {
-    std::vector<double> twist(Nd,0.0); //default: periodic boundarys in all directions
-    std::vector<Complex> boundary;
-    for(int i=0;i<Nd;i++) boundary.push_back(1);//default: periodic boundary conditions
-    FreePropagator(in,out,mass,boundary,twist);
-  };
-  
 protected:
 
   virtual void SetCoefficientsTanh(Approx::zolotarev_data *zdata,RealD scale);
   virtual void SetCoefficientsZolotarev(RealD zolo_hi,Approx::zolotarev_data *zdata);
 
   // Part frac
-  std::vector<RealD> qmu;
   RealD mass;
   RealD dw_diag;
   RealD R;

Grid/qcd/action/fermion/implementation/CayleyFermion5DImplementation.h

@@ -48,8 +48,7 @@ CayleyFermion5D<Impl>::CayleyFermion5D(GaugeField &_Umu,
 			FourDimGrid,
 			FourDimRedBlackGrid,_M5,p),
   mass_plus(_mass), mass_minus(_mass)
-{
+{ 
-  // qmu defaults to zero size;
 }
 
 ///////////////////////////////////////////////////////////////
@@ -271,34 +270,6 @@ void CayleyFermion5D<Impl>::MeooeDag5D    (const FermionField &psi, FermionField
   M5Ddag(psi,psi,Din,lower,diag,upper);
 }
 
-template<class Impl>
-void CayleyFermion5D<Impl>::addQmu(const FermionField &psi,FermionField &chi, int dag)
-{
-  if ( qmu.size() ) {
-
-    Gamma::Algebra Gmu [] = {
-      Gamma::Algebra::GammaX,
-      Gamma::Algebra::GammaY,
-      Gamma::Algebra::GammaZ,
-      Gamma::Algebra::GammaT
-    };
-    std::vector<ComplexD> coeff(Nd);
-    ComplexD ci(0,1);
-
-    assert(qmu.size()==Nd);
-
-    for(int mu=0;mu<Nd;mu++){
-       coeff[mu] = ci*qmu[mu];
-       if ( dag ) coeff[mu] = conjugate(coeff[mu]);
-    }
-
-    chi = chi + Gamma(Gmu[0])*psi*coeff[0];
-    for(int mu=1;mu<Nd;mu++){
-      chi = chi + Gamma(Gmu[mu])*psi*coeff[mu];
-    }
-  }
-}
-
 template<class Impl>
 void CayleyFermion5D<Impl>::M    (const FermionField &psi, FermionField &chi)
 {
@@ -306,12 +277,8 @@ void CayleyFermion5D<Impl>::M    (const FermionField &psi, FermionField &chi)
   
   // Assemble Din
   Meooe5D(psi,Din);
-
-  this->DW(Din,chi,DaggerNo);
-
-  // add i q_mu gamma_mu here
-  addQmu(Din,chi,DaggerNo);
   
+  this->DW(Din,chi,DaggerNo);
   // ((b D_W + D_w hop terms +1) on s-diag
   axpby(chi,1.0,1.0,chi,psi); 
   
@@ -328,9 +295,6 @@ void CayleyFermion5D<Impl>::Mdag (const FermionField &psi, FermionField &chi)
   FermionField Din(psi.Grid());
   // Apply Dw
   this->DW(psi,Din,DaggerYes); 
-
-  // add -i conj(q_mu) gamma_mu here ... if qmu is real, gammm_5 hermitian, otherwise not.
-  addQmu(psi,Din,DaggerYes);
   
   MeooeDag5D(Din,chi);
   

Grid/qcd/action/fermion/implementation/ContinuedFractionFermion5DImplementation.h

@@ -42,13 +42,13 @@ template<class Impl>
 void ContinuedFractionFermion5D<Impl>::SetCoefficientsZolotarev(RealD zolo_hi,Approx::zolotarev_data *zdata)
 {
   // How to check Ls matches??
-  std::cout<<GridLogMessage << zdata->n  << " - n"<<std::endl;
+  //      std::cout<<GridLogMessage << Ls << " Ls"<<std::endl;
-  std::cout<<GridLogMessage << zdata->da << " -da "<<std::endl;
+  //      std::cout<<GridLogMessage << zdata->n  << " - n"<<std::endl;
-  std::cout<<GridLogMessage << zdata->db << " -db"<<std::endl;
+  //      std::cout<<GridLogMessage << zdata->da << " -da "<<std::endl;
-  std::cout<<GridLogMessage << zdata->dn << " -dn"<<std::endl;
+  //      std::cout<<GridLogMessage << zdata->db << " -db"<<std::endl;
-  std::cout<<GridLogMessage << zdata->dd << " -dd"<<std::endl;
+  //      std::cout<<GridLogMessage << zdata->dn << " -dn"<<std::endl;
+  //      std::cout<<GridLogMessage << zdata->dd << " -dd"<<std::endl;
   int Ls = this->Ls;
-  std::cout<<GridLogMessage << Ls << " Ls"<<std::endl;
   assert(zdata->db==Ls);// Beta has Ls coeffs
 
   R=(1+this->mass)/(1-this->mass);
@@ -320,7 +320,7 @@ ContinuedFractionFermion5D<Impl>::ContinuedFractionFermion5D(
       int Ls = this->Ls;
       conformable(solution5d.Grid(),this->FermionGrid());
       conformable(exported4d.Grid(),this->GaugeGrid());
-      ExtractSlice(exported4d, solution5d, Ls-1, 0);
+      ExtractSlice(exported4d, solution5d, Ls-1, Ls-1);
     }
     template<class Impl>
     void ContinuedFractionFermion5D<Impl>::ImportPhysicalFermionSource(const FermionField &input4d,FermionField &imported5d)
@@ -330,7 +330,7 @@ ContinuedFractionFermion5D<Impl>::ContinuedFractionFermion5D(
       conformable(input4d.Grid()   ,this->GaugeGrid());
       FermionField tmp(this->FermionGrid());
       tmp=Zero();
-      InsertSlice(input4d, tmp, Ls-1, 0);
+      InsertSlice(input4d, tmp, Ls-1, Ls-1);
       tmp=Gamma(Gamma::Algebra::Gamma5)*tmp;
       this->Dminus(tmp,imported5d);
     }

Grid/qcd/action/fermion/implementation/PartialFractionFermion5DImplementation.h

@@ -255,76 +255,15 @@ void   PartialFractionFermion5D<Impl>::M_internal(const FermionField &psi, Fermi
   }
 	
   {
-    // The 'conventional' Cayley overlap operator is
-    //
-    // Dov = (1+m)/2 + (1-m)/2 g5 sgn Hw
-    //
-    //
-    // With massless limit 1/2(1+g5 sgnHw)
-    //
-    // Luscher shows quite neatly that 1+g5 sgn Hw has tree level propagator i qslash +O(a^2)
-    //
-    // However, the conventional normalisation has both a leading order factor of 2 in Zq
-    // at tree level AND a mass dependent (1-m) that are convenient to absorb.
-    //
-    // In WilsonFermion5DImplementation.h, the tree level propagator for Hw is
-    //
-    // num = -i sin kmu gmu
-    //
-    // denom ( sqrt(sk^2 + (2shk^2 - 1)^2
-    //    b_k = sk2 - M5;
-    //     
-    //    w_k = sqrt(sk + b_k*b_k);
-    //
-    //    denom= ( w_k + b_k + mass*mass) ;
-    //
-    //    denom= one/denom;
-    //    out = num*denom;
-    //
-    // Chroma, and Grid define partial fraction via 4d operator
-    //
-    //   Dpf = 2/(1-m) x Dov = (1+m)/(1-m) + g5 sgn Hw
-    //
-    // Now since:
-    //
-    //      (1+m)/(1-m) = (1-m)/(1-m) + 2m/(1-m) = 1 + 2m/(1-m)
-    //
-    // This corresponds to a modified mass parameter
-    //
-    // It has an annoying 
-    //
-    // 
     double R=(1+this->mass)/(1-this->mass);
     //R g5 psi[Ls] + p[0] H
     ag5xpbg5y_ssp(chi,R*scale,psi,p[nblock]*scale/amax,D,Ls-1,Ls-1);
-    
+	
     for(int b=0;b<nblock;b++){
       int s = 2*b+1;
       double pp = p[nblock-1-b];
       axpby_ssp(chi,1.0,chi,-sqrt(amax*pp)*scale*sign,psi,Ls-1,s);
     }
-
-    if ( qmu.size() ) {
-
-      FermionField qslash_psi(psi.Grid());
-      
-      Gamma::Algebra Gmu [] = {
-			 Gamma::Algebra::GammaX,
-			 Gamma::Algebra::GammaY,
-			 Gamma::Algebra::GammaZ,
-			 Gamma::Algebra::GammaT
-      };
-      ComplexD ci(0,1);
-      assert(qmu.size()==Nd);
-      qslash_psi = Gamma(Gmu[0])*psi;
-      for(int mu=1;mu<Nd;mu++){
-	qslash_psi = Gamma(Gmu[mu])*psi;
-      }
-      //      RealD coeff = 1.0;
-      qslash_psi = Gamma(Gamma::Algebra::Gamma5)*qslash_psi*ci ; // i g5 qslash -- 1-m factor???
-      axpby_ssp(chi,1.0,chi,1.0, qslash_psi,Ls-1,Ls-1);
-    }
-    
   }
 
 }
@@ -472,7 +411,7 @@ void  PartialFractionFermion5D<Impl>::SetCoefficientsZolotarev(RealD zolo_hi,App
       int Ls = this->Ls;
       conformable(solution5d.Grid(),this->FermionGrid());
       conformable(exported4d.Grid(),this->GaugeGrid());
-      ExtractSlice(exported4d, solution5d, Ls-1, 0);
+      ExtractSlice(exported4d, solution5d, Ls-1, Ls-1);
     }
     template<class Impl>
     void PartialFractionFermion5D<Impl>::ImportPhysicalFermionSource(const FermionField &input4d,FermionField &imported5d)
@@ -482,8 +421,7 @@ void  PartialFractionFermion5D<Impl>::SetCoefficientsZolotarev(RealD zolo_hi,App
       conformable(input4d.Grid()   ,this->GaugeGrid());
       FermionField tmp(this->FermionGrid());
       tmp=Zero();
-      std::cout << " importing to slice " << Ls-1 <<std::endl;
+      InsertSlice(input4d, tmp, Ls-1, Ls-1);
-      InsertSlice(input4d, tmp, Ls-1, 0);
       tmp=Gamma(Gamma::Algebra::Gamma5)*tmp;
       this->Dminus(tmp,imported5d);
     }
@@ -504,7 +442,7 @@ PartialFractionFermion5D<Impl>::PartialFractionFermion5D(GaugeField &_Umu,
 
 {
   int Ls = this->Ls;
-  qmu.resize(0);
+
   assert((Ls&0x1)==1); // Odd Ls required
   int nrational=Ls-1;
 
@@ -522,22 +460,6 @@ PartialFractionFermion5D<Impl>::PartialFractionFermion5D(GaugeField &_Umu,
   Approx::zolotarev_free(zdata);
 
 }
-template<class Impl>
-PartialFractionFermion5D<Impl>::PartialFractionFermion5D(GaugeField &_Umu,
-							 GridCartesian         &FiveDimGrid,
-							 GridRedBlackCartesian &FiveDimRedBlackGrid,
-							 GridCartesian         &FourDimGrid,
-							 GridRedBlackCartesian &FourDimRedBlackGrid,
-							 RealD _mass,RealD M5,
-							 std::vector<RealD> &_qmu,
-							 const ImplParams &p)
-  : PartialFractionFermion5D<Impl>(_Umu,
-			     FiveDimGrid,FiveDimRedBlackGrid,
-			     FourDimGrid,FourDimRedBlackGrid,
-			     _mass,M5,p)
-{
-  qmu=_qmu;
-}
 
 NAMESPACE_END(Grid);
 

Grid/tensors/Tensor_Ta.h

@@ -90,10 +90,12 @@ template<class vtype,int N> accelerator_inline iVector<vtype,N> ProjectOnGroup(c
 template<class vtype,int N, typename std::enable_if< GridTypeMapper<vtype>::TensorLevel == 0 >::type * =nullptr> 
 accelerator_inline iMatrix<vtype,N> ProjectOnGroup(const iMatrix<vtype,N> &arg)
 {
+  typedef typename iMatrix<vtype,N>::scalar_type scalar;
   // need a check for the group type?
   iMatrix<vtype,N> ret(arg);
   vtype nrm;
   vtype inner;
+  scalar one(1.0);
   for(int c1=0;c1<N;c1++){
 
     // Normalises row c1
@@ -102,7 +104,7 @@ accelerator_inline iMatrix<vtype,N> ProjectOnGroup(const iMatrix<vtype,N> &arg)
       inner += innerProduct(ret._internal[c1][c2],ret._internal[c1][c2]);
 
     nrm = sqrt(inner);
-    nrm = 1.0/nrm;
+    nrm = one/nrm;
     for(int c2=0;c2<N;c2++)
       ret._internal[c1][c2]*= nrm;
       
@@ -127,7 +129,7 @@ accelerator_inline iMatrix<vtype,N> ProjectOnGroup(const iMatrix<vtype,N> &arg)
       inner += innerProduct(ret._internal[c1][c2],ret._internal[c1][c2]);
 
     nrm = sqrt(inner);
-    nrm = 1.0/nrm;
+    nrm = one/nrm;
     for(int c2=0;c2<N;c2++)
       ret._internal[c1][c2]*= nrm;
   }

systems/OEM/README

@@ -0,0 +1,53 @@
+1. Prerequisites:
+===================
+Make sure you have the latest Intel ipcx release loaded (via modules or similar)
+Make sure you have SYCL aware MPICH or Intel MPI loaded (assumed as mpicxx)
+
+2. Obtain Grid:
+===================
+
+bash$
+git clone https://github.com/paboyle/Grid
+cd Grid
+./bootstrap.sh
+cd systems/PVC
+
+3. Build Grid:
+===================
+
+Here, configure command is stored in file config-command:
+
+bash$
+../../configure \
+	--enable-simd=GPU \
+	--enable-gen-simd-width=64 \
+	--enable-comms=mpi-auto \
+	--enable-accelerator-cshift \
+	--disable-gparity \
+	--disable-fermion-reps \
+	--enable-shm=nvlink \
+	--enable-accelerator=sycl \
+	--enable-unified=no \
+	MPICXX=mpicxx \
+	CXX=icpx \
+	LDFLAGS="-fiopenmp  -fsycl -fsycl-device-code-split=per_kernel -fsycl-device-lib=all -lze_loader " \
+	CXXFLAGS="-fiopenmp -fsycl-unnamed-lambda -fsycl -Wno-tautological-compare "
+
+make all
+
+4. Run a benchmark:
+===================
+
+*** Assumes interactive access to node. ***
+
+run Benchmark_dwf_fp32 using benchmarks/bench.sh
+
+bash$
+cd benchmarks
+./bench.sh
+
+
+
+
+
+

systems/OEM/benchmarks/bench.sh

@@ -0,0 +1,18 @@
+#!/bin/bash
+
+export EnableImplicitScaling=0
+export ZE_ENABLE_PCI_ID_DEVICE_ORDER=1
+export ZE_AFFINITY_MASK=$gpu_id.$tile_id
+export ONEAPI_DEVICE_FILTER=gpu,level_zero
+export SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS=0
+export SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1
+export SYCL_PI_LEVEL_ZERO_USE_COPY_ENGINE=0:2
+export SYCL_PI_LEVEL_ZERO_USE_COPY_ENGINE_FOR_D2D_COPY=1
+
+mpiexec -launcher ssh -n 1 -host localhost ./select_gpu.sh ./Benchmark_dwf_fp32 --mpi 1.1.1.1 --grid 32.32.32.32 --accelerator-threads 16 --shm-mpi 1 --shm 2048 --device-mem 32768 | tee 1tile.log
+mpiexec -launcher ssh -n 2 -host localhost ./select_gpu.sh ./Benchmark_dwf_fp32 --mpi 1.1.1.2 --grid 32.32.32.64 --accelerator-threads 16 --shm-mpi 1 --shm 2048 --device-mem 32768 | tee 2tile.log
+
+#mpiexec -launcher ssh -n 4 -host localhost ./select_gpu.sh ./Benchmark_dwf_fp32 --mpi 1.1.2.2 --grid 16.16.64.64 --accelerator-threads 16 --shm-mpi 0 --shm 2048 --device-mem 32768 | tee 4tile.log
+#mpiexec -launcher ssh -n 8 -host localhost ./select_gpu.sh ./Benchmark_dwf_fp32 --mpi 1.1.2.4 --grid 16.16.64.128 --accelerator-threads 16 --shm-mpi 0 --shm 2048 --device-mem 32768 | tee 8tile.log
+
+

systems/OEM/benchmarks/select_gpu.sh

@@ -0,0 +1,13 @@
+#!/bin/bash
+
+num_tile=2
+
+gpu_id=$(( (MPI_LOCAL_RANKID % num_tile ) ))
+tile_id=$((MPI_LOCAL_RANKID / num_tile))
+
+export ZE_AFFINITY_MASK=$gpu_id.$tile_id
+
+echo "local rank $MPI_LOCALRANKID ; ZE_AFFINITY_MASK=$ZE_AFFINITY_MASK"
+
+"$@"
+

systems/OEM/config-command

@@ -0,0 +1,15 @@
+../../configure \
+	--enable-simd=GPU \
+	--enable-gen-simd-width=64 \
+	--enable-comms=mpi-auto \
+	--enable-accelerator-cshift \
+	--disable-gparity \
+	--disable-fermion-reps \
+	--enable-shm=nvlink \
+	--enable-accelerator=sycl \
+	--enable-unified=no \
+	MPICXX=mpicxx \
+	CXX=icpx \
+	LDFLAGS="-fiopenmp  -fsycl -fsycl-device-code-split=per_kernel -fsycl-device-lib=all -lze_loader " \
+	CXXFLAGS="-fiopenmp -fsycl-unnamed-lambda -fsycl -Wno-tautological-compare "
+

systems/OEM/setup.sh

@@ -0,0 +1,3 @@
+export https_proxy=http://proxy-chain.intel.com:911
+module load intel-release
+module load intel/mpich

systems/mac-arm/config-command-mpi

@@ -1,4 +1,4 @@
 BREW=/opt/local/
-CXX=mpicxx-openmpi-mp ../../configure --enable-simd=GEN --enable-comms=mpi --enable-unified=yes --prefix $HOME/QCD/GridInstall --with-lime=/Users/peterboyle/QCD/SciDAC/install/ --with-openssl=$BREW --disable-fermion-reps --disable-gparity --disable-debug
+MPICXX=mpicxx CXX=c++-12 ../../configure --enable-simd=GEN --enable-comms=mpi-auto --enable-unified=yes --prefix $HOME/QCD/GridInstall --with-lime=/Users/peterboyle/QCD/SciDAC/install/ --with-openssl=$BREW --disable-fermion-reps --disable-gparity --disable-debug
 
 

tests/qdpxx/Test_qdpxx_munprec.cc

@@ -1,6 +1,7 @@
     /*************************************************************************************
 
     Grid physics library, www.github.com/paboyle/Grid 
+
     Source file: ./tests/qdpxx/Test_qdpxx_munprec.cc
 
     Copyright (C) 2015
@@ -25,17 +26,13 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
     See the full license in the file "LICENSE" in the top level distribution directory
     *************************************************************************************/
     /*  END LEGAL */
-#include <chroma.h>
-#include <actions/ferm/invert/syssolver_linop_cg_array.h>
-#include <actions/ferm/invert/syssolver_linop_aggregate.h>
-
 #include <Grid/Grid.h>
 
 int    Ls=8;
 double M5=1.6;
 double mq=0.01;
-double zolo_lo = 0.01;
+double zolo_lo = 0.1;
-double zolo_hi = 7.0;
+double zolo_hi = 2.0;
 double mobius_scale=2.0;
 
 enum ChromaAction {
@@ -58,6 +55,11 @@ enum ChromaAction {
 void calc_grid      (ChromaAction action,Grid::LatticeGaugeField & lat, Grid::LatticeFermion &src, Grid::LatticeFermion &res,int dag);
 void calc_chroma    (ChromaAction action,Grid::LatticeGaugeField & lat, Grid::LatticeFermion &src, Grid::LatticeFermion &res,int dag);
 
+#include <chroma.h>
+#include <actions/ferm/invert/syssolver_linop_cg_array.h>
+#include <actions/ferm/invert/syssolver_linop_aggregate.h>
+
+
 
 namespace Chroma { 
 
@@ -79,7 +81,7 @@ public:
 
     std::vector<int> x(4);
     QDP::multi1d<int> cx(4);
-    Grid::Coordinate gd = gr.Grid()->GlobalDimensions();
+    std::vector<int> gd= gr.Grid()->GlobalDimensions();
 
     for (x[0]=0;x[0]<gd[0];x[0]++){
     for (x[1]=0;x[1]<gd[1];x[1]++){
@@ -122,7 +124,7 @@ public:
 
     std::vector<int> x(5);
     QDP::multi1d<int> cx(4);
-    Grid::Coordinate gd= gr.Grid()->GlobalDimensions();
+    std::vector<int> gd= gr.Grid()->GlobalDimensions();
 
     for (x[0]=0;x[0]<gd[0];x[0]++){
     for (x[1]=0;x[1]<gd[1];x[1]++){
@@ -164,7 +166,7 @@ public:
 
     std::vector<int> x(5);
     QDP::multi1d<int> cx(4);
-    Grid::Coordinate gd= gr.Grid()->GlobalDimensions();
+    std::vector<int> gd= gr.Grid()->GlobalDimensions();
 
     for (x[0]=0;x[0]<gd[0];x[0]++){
     for (x[1]=0;x[1]<gd[1];x[1]++){
@@ -302,30 +304,7 @@ public:
      //     param.approximation_type=COEFF_TYPE_TANH_UNSCALED;
      //     param.approximation_type=COEFF_TYPE_TANH;
      param.tuning_strategy_xml=
-"<TuningStrategy><Name>OVEXT_CONSTANT_STRATEGY</Name><TuningConstant>1.0</TuningConstant></TuningStrategy>\n";
+"<TuningStrategy><Name>OVEXT_CONSTANT_STRATEGY</Name></TuningStrategy>\n";
-     UnprecOvExtFermActArray S_f(cfs,param);
-     Handle< FermState<T4,U,U> > fs( S_f.createState(u) );
-     Handle< LinearOperatorArray<T4> > M(S_f.linOp(fs));
-     return M;
-   }
-   if ( parms == HwPartFracTanh ) {
-     if ( Ls%2 == 0 ) { 
-       printf("Ls is not odd\n");
-       exit(-1);
-     }
-     UnprecOvExtFermActArrayParams param;
-     param.OverMass=M5; 
-     param.Mass=_mq;
-     param.RatPolyDeg = Ls;
-     param.ApproxMin =eps_lo;
-     param.ApproxMax =eps_hi;
-     param.b5 =1.0;
-     param.c5 =1.0;
-     //     param.approximation_type=COEFF_TYPE_ZOLOTAREV;
-     param.approximation_type=COEFF_TYPE_TANH_UNSCALED;
-     //param.approximation_type=COEFF_TYPE_TANH;
-     param.tuning_strategy_xml=
-       "<TuningStrategy><Name>OVEXT_CONSTANT_STRATEGY</Name><TuningConstant>1.0</TuningConstant></TuningStrategy>\n";
      UnprecOvExtFermActArray S_f(cfs,param);
      Handle< FermState<T4,U,U> > fs( S_f.createState(u) );
      Handle< LinearOperatorArray<T4> > M(S_f.linOp(fs));
@@ -337,35 +316,7 @@ public:
      param.ApproxMin=eps_lo;
      param.ApproxMax=eps_hi;
      param.approximation_type=COEFF_TYPE_ZOLOTAREV;
-     param.RatPolyDeg=Ls-1;
+     param.RatPolyDeg=Ls;
-     // The following is why I think Chroma made some directional errors:
-     param.AuxFermAct= std::string(
-"<AuxFermAct>\n"
-"  <FermAct>UNPRECONDITIONED_WILSON</FermAct>\n"
-"  <Mass>-1.8</Mass>\n"
-"  <b5>1</b5>\n"
-"  <c5>0</c5>\n"
-"  <MaxCG>1000</MaxCG>\n"
-"  <RsdCG>1.0e-9</RsdCG>\n"
-"  <FermionBC>\n"
-"      <FermBC>SIMPLE_FERMBC</FermBC>\n"
-"      <boundary>1 1 1 1</boundary>\n"
-"   </FermionBC> \n"
-"</AuxFermAct>"
-);
-     param.AuxFermActGrp= std::string("");
-     UnprecOvlapContFrac5DFermActArray S_f(fbc,param);
-     Handle< FermState<T4,U,U> > fs( S_f.createState(u) );
-     Handle< LinearOperatorArray<T4> > M(S_f.linOp(fs));
-     return  M;
-   }
-   if ( parms == HwContFracTanh ) {
-     UnprecOvlapContFrac5DFermActParams param;
-     param.Mass=_mq; // How is M5 set? Wilson mass In AuxFermAct
-     param.ApproxMin=eps_lo;
-     param.ApproxMax=eps_hi;
-     param.approximation_type=COEFF_TYPE_TANH_UNSCALED;
-     param.RatPolyDeg=Ls-1;
      // The following is why I think Chroma made some directional errors:
      param.AuxFermAct= std::string(
 "<AuxFermAct>\n"
@@ -427,14 +378,7 @@ int main (int argc,char **argv )
    * Setup QDP
    *********************************************************/
   Chroma::initialize(&argc,&argv);
-  //  Chroma::WilsonTypeFermActs4DEnv::registerAll(); 
+  Chroma::WilsonTypeFermActs4DEnv::registerAll(); 
-  Chroma::WilsonTypeFermActsEnv::registerAll(); 
-  //bool linkageHack(void)
-  //{
-  //  bool foo = true;
-  // Inline Measurements
-  //  InlineAggregateEnv::registerAll();
-  //  GaugeInitEnv::registerAll();
 
   /********************************************************
    * Setup Grid
@@ -444,34 +388,26 @@ int main (int argc,char **argv )
                                                                        Grid::GridDefaultSimd(Grid::Nd,Grid::vComplex::Nsimd()),
                                                                        Grid::GridDefaultMpi());
   
-  Grid::Coordinate gd = UGrid->GlobalDimensions();
+  std::vector<int> gd = UGrid->GlobalDimensions();
   QDP::multi1d<int> nrow(QDP::Nd);
   for(int mu=0;mu<4;mu++) nrow[mu] = gd[mu];
 
   QDP::Layout::setLattSize(nrow);
   QDP::Layout::create();
 
+  Grid::GridCartesian         * FGrid   = Grid::SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
+  Grid::LatticeGaugeField lat(UGrid);
+  Grid::LatticeFermion    src(FGrid);
+  Grid::LatticeFermion    res_chroma(FGrid);
+  Grid::LatticeFermion    res_grid  (FGrid);
+  
   std::vector<ChromaAction> ActionList({
 		 HtCayleyTanh, // Plain old DWF.
 		 HmCayleyTanh,
 		 HwCayleyTanh,
 		 HtCayleyZolo, // Plain old DWF.
 		 HmCayleyZolo,
-		 HwCayleyZolo,
+		 HwCayleyZolo
-		 HwPartFracZolo,
-		 HwContFracZolo,
-		 HwContFracTanh
-  });
-  std::vector<int> LsList({
-      8,//HtCayleyTanh, // Plain old DWF.
-      8,//HmCayleyTanh,
-      8,//HwCayleyTanh,
-      8,//HtCayleyZolo, // Plain old DWF.
-      8,//HmCayleyZolo,
-      8,//HwCayleyZolo,
-      9,//HwPartFracZolo
-      9, //HwContFracZolo
-      9 //HwContFracTanh
   });
   std::vector<std::string> ActionName({
         "HtCayleyTanh",
@@ -479,19 +415,10 @@ int main (int argc,char **argv )
 	"HwCayleyTanh",
 	"HtCayleyZolo",
 	"HmCayleyZolo",
-        "HwCayleyZolo",
+        "HwCayleyZolo"
-	"HwPartFracZolo",
-	"HwContFracZolo",
-	"HwContFracTanh"
   });
 
   for(int i=0;i<ActionList.size();i++) {
-    Ls = LsList[i];
-    Grid::GridCartesian      * FGrid   = Grid::SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
-    Grid::LatticeGaugeField lat(UGrid);
-    Grid::LatticeFermion    src(FGrid);
-    Grid::LatticeFermion    res_chroma(FGrid);
-    Grid::LatticeFermion    res_grid  (FGrid);
     std::cout << "*****************************"<<std::endl;
     std::cout << "Action "<<ActionName[i]<<std::endl;
     std::cout << "*****************************"<<std::endl;
@@ -512,7 +439,6 @@ int main (int argc,char **argv )
       
       std::cout << "Norm of difference "<<Grid::norm2(res_chroma)<<std::endl;
     }
-    delete FGrid;
   }
 
   std::cout << "Finished test "<<std::endl;
@@ -576,7 +502,7 @@ void calc_grid(ChromaAction action,Grid::LatticeGaugeField & Umu, Grid::LatticeF
   Grid::gaussian(RNG5,src);
   Grid::gaussian(RNG5,res);
 
-  Grid::SU<Grid::Nc>::HotConfiguration(RNG4,Umu);
+  Grid::SU<Nc>::HotConfiguration(RNG4,Umu);
 
   /*
   Grid::LatticeColourMatrix U(UGrid);
@@ -593,7 +519,7 @@ void calc_grid(ChromaAction action,Grid::LatticeGaugeField & Umu, Grid::LatticeF
 
   if ( action == HtCayleyTanh ) { 
 
-    Grid::DomainWallFermionD Ddwf(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5);
+    Grid::DomainWallFermionR Ddwf(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5);
 
     std::cout << Grid::GridLogMessage <<" Calling domain wall multiply "<<std::endl;
 
@@ -609,7 +535,7 @@ void calc_grid(ChromaAction action,Grid::LatticeGaugeField & Umu, Grid::LatticeF
 
     Grid::Real _b = 0.5*(mobius_scale +1.0);
     Grid::Real _c = 0.5*(mobius_scale -1.0);
-    Grid::MobiusZolotarevFermionD D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,_b,_c,zolo_lo,zolo_hi);
+    Grid::MobiusZolotarevFermionR D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,_b,_c,zolo_lo,zolo_hi);
 
     std::cout << Grid::GridLogMessage <<" Calling mobius zolo multiply "<<std::endl;
 
@@ -623,7 +549,7 @@ void calc_grid(ChromaAction action,Grid::LatticeGaugeField & Umu, Grid::LatticeF
 
   if ( action == HtCayleyZolo ) {
 
-    Grid::ShamirZolotarevFermionD D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,zolo_lo,zolo_hi);
+    Grid::ShamirZolotarevFermionR D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,zolo_lo,zolo_hi);
 
     std::cout << Grid::GridLogMessage <<" Calling shamir zolo multiply "<<std::endl;
 
@@ -635,60 +561,6 @@ void calc_grid(ChromaAction action,Grid::LatticeGaugeField & Umu, Grid::LatticeF
     return;
   }
 
-  if ( action == HwPartFracTanh ) {
-
-    Grid::OverlapWilsonPartialFractionTanhFermionD Dov(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,1.0);
-
-    std::cout << Grid::GridLogMessage <<" Calling part frac tanh multiply "<<std::endl;
-
-    if ( dag ) 
-      Dov.Mdag(src,res);  
-    else 
-      Dov.M(src,res);  
-
-    return;
-  }
-
-  if ( action == HwContFracTanh ) {
-
-    Grid::OverlapWilsonContFracTanhFermionD Dov(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,1.0);
-
-    std::cout << Grid::GridLogMessage <<" Calling cont frac tanh multiply "<<std::endl;
-
-    if ( dag ) 
-      Dov.Mdag(src,res);  
-    else 
-      Dov.M(src,res);  
-
-    return;
-  }
-  if ( action == HwContFracZolo ) {
-
-    Grid::OverlapWilsonContFracZolotarevFermionD Dov(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,zolo_lo,zolo_hi);
-
-    std::cout << Grid::GridLogMessage <<" Calling cont frac zolo multiply "<<std::endl;
-
-    if ( dag ) 
-      Dov.Mdag(src,res);  
-    else 
-      Dov.M(src,res);  
-
-    return;
-  }
-
-  if ( action == HwPartFracZolo ) {
-
-    Grid::OverlapWilsonPartialFractionZolotarevFermionD Dov(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,zolo_lo,zolo_hi);
-    std::cout << Grid::GridLogMessage <<" Calling part frac zolotarev multiply "<<std::endl;
-
-    if ( dag ) 
-      Dov.Mdag(src,res);  
-    else 
-      Dov.M(src,res);  
-
-    return;
-  }
-  
   /*
   if ( action == HmCayleyTanh ) {
     Grid::Real _b = 0.5*(mobius_scale +1.0);
@@ -709,7 +581,7 @@ void calc_grid(ChromaAction action,Grid::LatticeGaugeField & Umu, Grid::LatticeF
 
   if ( action == HmCayleyTanh ) {
 
-    Grid::ScaledShamirFermionD D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,mobius_scale);
+    Grid::ScaledShamirFermionR D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,mobius_scale);
 
     std::cout << Grid::GridLogMessage <<" Calling scaled shamir multiply "<<std::endl;
 
@@ -723,7 +595,7 @@ void calc_grid(ChromaAction action,Grid::LatticeGaugeField & Umu, Grid::LatticeF
 
   if ( action == HwCayleyTanh ) {
 
-    Grid::OverlapWilsonCayleyTanhFermionD D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,1.0);
+    Grid::OverlapWilsonCayleyTanhFermionR D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,1.0);
 
     if ( dag ) 
       D.Mdag(src,res);  
@@ -735,7 +607,7 @@ void calc_grid(ChromaAction action,Grid::LatticeGaugeField & Umu, Grid::LatticeF
 
   if ( action == HwCayleyZolo ) {
 
-    Grid::OverlapWilsonCayleyZolotarevFermionD D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,zolo_lo,zolo_hi);
+    Grid::OverlapWilsonCayleyZolotarevFermionR D(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,_mass,_M5,zolo_lo,zolo_hi);
 
     if ( dag ) 
       D.Mdag(src,res);  

tests/solver/Test_dwf_cg_prec.cc

@@ -1,4 +1,4 @@
-*************************************************************************************
+/*************************************************************************************
 
 Grid physics library, www.github.com/paboyle/Grid
 
@@ -67,13 +67,7 @@ int main(int argc, char** argv) {
   result = Zero();
   LatticeGaugeField Umu(UGrid);
 
-#if 0
-  FieldMetaData header;
-  std::string file("ckpoint_lat.4000");
-  NerscIO::readConfiguration(Umu,header,file);
-#else  
   SU<Nc>::HotConfiguration(RNG4, Umu);
-#endif
 
   std::cout << GridLogMessage << "Lattice dimensions: " << GridDefaultLatt()
             << "   Ls: " << Ls << std::endl;

tests/solver/Test_dwf_cg_unprec.cc

@@ -54,30 +54,15 @@ int main (int argc, char ** argv)
   GridCartesian         * FGrid   = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
   GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
 
-  std::vector<ComplexD> qmu;
-  qmu.push_back(ComplexD(0.1,0.0));
-  qmu.push_back(ComplexD(0.0,0.0));
-  qmu.push_back(ComplexD(0.0,0.0));
-  qmu.push_back(ComplexD(0.0,0.01));
-  
-
   std::vector<int> seeds4({1,2,3,4});
   std::vector<int> seeds5({5,6,7,8});
   GridParallelRNG          RNG5(FGrid);  RNG5.SeedFixedIntegers(seeds5);
   GridParallelRNG          RNG4(UGrid);  RNG4.SeedFixedIntegers(seeds4);
 
-  LatticeFermion    tmp(FGrid);
   LatticeFermion    src(FGrid); random(RNG5,src);
   LatticeFermion result(FGrid); result=Zero();
-  LatticeGaugeField Umu(UGrid); 
+  LatticeGaugeField Umu(UGrid); SU<Nc>::HotConfiguration(RNG4,Umu);
-#if 0
+
-  FieldMetaData header;
-  std::string file("ckpoint_lat.4000");
-  NerscIO::readConfiguration(Umu,header,file);
-#else  
-  SU<Nc>::HotConfiguration(RNG4,Umu);
-#endif
-  
   std::vector<LatticeColourMatrix> U(4,UGrid);
   for(int mu=0;mu<Nd;mu++){
     U[mu] = PeekIndex<LorentzIndex>(Umu,mu);
@@ -86,15 +71,8 @@ int main (int argc, char ** argv)
   RealD mass=0.1;
   RealD M5=1.8;
   DomainWallFermionD Ddwf(Umu,*FGrid,*FrbGrid,*UGrid,*UrbGrid,mass,M5);
-  Ddwf.qmu = qmu;
 
-  Ddwf.M(src,tmp);
-  std::cout << " |M src|^2 "<<norm2(tmp)<<std::endl;
   MdagMLinearOperator<DomainWallFermionD,LatticeFermion> HermOp(Ddwf);
-  HermOp.HermOp(src,tmp);
-
-  std::cout << " <src|MdagM| src> "<<innerProduct(src,tmp)<<std::endl;
-  
   ConjugateGradient<LatticeFermion> CG(1.0e-6,10000);
   CG(HermOp,src,result);