Bland SYCL compile

Bugfix/irl convergence

Grid/algorithms/iterative/ImplicitlyRestartedLanczos.h

@@ -419,14 +419,15 @@ until convergence
 	}
       }
 
-      if ( Nconv < Nstop )
+      if ( Nconv < Nstop ) {
 	std::cout << GridLogIRL << "Nconv ("<<Nconv<<") < Nstop ("<<Nstop<<")"<<std::endl;
-
+	std::cout << GridLogIRL << "returning Nstop vectors, the last "<< Nstop-Nconv << "of which might meet convergence criterion only approximately" <<std::endl;
+      }
       eval=eval2;
       
       //Keep only converged
-      eval.resize(Nconv);// Nstop?
-      evec.resize(Nconv,grid);// Nstop?
+      eval.resize(Nstop);// was Nconv
+      evec.resize(Nstop,grid);// was Nconv
       basisSortInPlace(evec,eval,reverse);
       
     }

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

Grid/qcd/action/ActionBase.h

@@ -44,7 +44,7 @@ public:
   ConfigurationBase() {}
   virtual ~ConfigurationBase() {}
   virtual void set_Field(Field& U) =0;
-  virtual void smeared_force(Field&) const = 0;
+  virtual void smeared_force(Field&) = 0;
   virtual Field& get_SmearedU() =0;
   virtual Field &get_U(bool smeared = false) = 0;
 };

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

@@ -423,7 +423,6 @@ void WilsonKernels<Impl>::DhopDirKernel( StencilImpl &st, DoubledGaugeField &U,S
 #define KERNEL_CALL(A) KERNEL_CALLNB(A); accelerator_barrier();
 
 #define KERNEL_CALL_EXT(A)						\
-  const uint64_t    NN = Nsite*Ls;					\
   const uint64_t    sz = st.surface_list.size();			\
   auto ptr = &st.surface_list[0];					\
   accelerator_forNB( ss, sz, Simd::Nsimd(), {				\

Grid/qcd/hmc/HMC.h

@@ -284,11 +284,12 @@ public:
 
       TheIntegrator.print_timer();
       
+      TheIntegrator.Smearer.set_Field(Ucur);
       for (int obs = 0; obs < Observables.size(); obs++) {
       	std::cout << GridLogDebug << "Observables # " << obs << std::endl;
       	std::cout << GridLogDebug << "Observables total " << Observables.size() << std::endl;
       	std::cout << GridLogDebug << "Observables pointer " << Observables[obs] << std::endl;
-        Observables[obs]->TrajectoryComplete(traj + 1, Ucur, sRNG, pRNG);
+        Observables[obs]->TrajectoryComplete(traj + 1, TheIntegrator.Smearer, sRNG, pRNG);
       }
       std::cout << GridLogHMC << ":::::::::::::::::::::::::::::::::::::::::::" << std::endl;
     }

Grid/qcd/hmc/checkpointers/BaseCheckpointer.h

@@ -35,13 +35,16 @@ class CheckpointerParameters : Serializable {
 public:
   GRID_SERIALIZABLE_CLASS_MEMBERS(CheckpointerParameters, 
 				  std::string, config_prefix, 
+				  std::string, smeared_prefix, 
 				  std::string, rng_prefix, 
 				  int, saveInterval, 
+				  bool, saveSmeared, 
 				  std::string, format, );
 
-  CheckpointerParameters(std::string cf = "cfg", std::string rn = "rng",
+  CheckpointerParameters(std::string cf = "cfg", std::string sf="cfg_smr" , std::string rn = "rng",
 			 int savemodulo = 1, const std::string &f = "IEEE64BIG")
     : config_prefix(cf),
+      smeared_prefix(sf),
       rng_prefix(rn),
       saveInterval(savemodulo),
       format(f){};
@@ -61,13 +64,21 @@ template <class Impl>
 class BaseHmcCheckpointer : public HmcObservable<typename Impl::Field> {
 public:
   void build_filenames(int traj, CheckpointerParameters &Params,
-                       std::string &conf_file, std::string &rng_file) {
+                       std::string &conf_file,
+                       std::string &smear_file,
+		       std::string &rng_file) {
     {
       std::ostringstream os;
       os << Params.rng_prefix << "." << traj;
       rng_file = os.str();
     }
 
+    {
+      std::ostringstream os;
+      os << Params.smeared_prefix << "." << traj;
+      smear_file = os.str();
+    }
+
     {
       std::ostringstream os;
       os << Params.config_prefix << "." << traj;
@@ -84,6 +95,11 @@ public:
   }
   virtual void initialize(const CheckpointerParameters &Params) = 0;
 
+  virtual void TrajectoryComplete(int traj,
+                                  typename Impl::Field &U,
+                                  GridSerialRNG &sRNG,
+                                  GridParallelRNG &pRNG) { assert(0); } ; // HMC should pass the smart config with smeared and unsmeared
+  
   virtual void CheckpointRestore(int traj, typename Impl::Field &U,
                                  GridSerialRNG &sRNG,
                                  GridParallelRNG &pRNG) = 0;

Grid/qcd/hmc/checkpointers/BinaryCheckpointer.h

@@ -61,11 +61,14 @@ public:
     fout.close();
   }
 
-  void TrajectoryComplete(int traj, Field &U, GridSerialRNG &sRNG, GridParallelRNG &pRNG) {
+  void TrajectoryComplete(int traj,
+			  ConfigurationBase<Field> &SmartConfig,
+			  GridSerialRNG &sRNG, GridParallelRNG &pRNG)
+  {
 
     if ((traj % Params.saveInterval) == 0) {
-      std::string config, rng;
-      this->build_filenames(traj, Params, config, rng);
+      std::string config, rng, smr;
+      this->build_filenames(traj, Params, config, smr, rng);
 
       uint32_t nersc_csum;
       uint32_t scidac_csuma;
@@ -74,9 +77,15 @@ public:
       BinarySimpleUnmunger<sobj_double, sobj> munge;
       truncate(rng);
       BinaryIO::writeRNG(sRNG, pRNG, rng, 0,nersc_csum,scidac_csuma,scidac_csumb);
-      truncate(config);
+      std::cout << GridLogMessage << "Written Binary RNG " << rng
+                << " checksum " << std::hex 
+		<< nersc_csum   <<"/"
+		<< scidac_csuma   <<"/"
+		<< scidac_csumb 
+		<< std::dec << std::endl;
 
-      BinaryIO::writeLatticeObject<vobj, sobj_double>(U, config, munge, 0, Params.format,
+      truncate(config);
+      BinaryIO::writeLatticeObject<vobj, sobj_double>(SmartConfig.get_U(false), config, munge, 0, Params.format,
 						      nersc_csum,scidac_csuma,scidac_csumb);
 
       std::cout << GridLogMessage << "Written Binary Configuration " << config
@@ -85,6 +94,18 @@ public:
 		<< scidac_csuma   <<"/"
 		<< scidac_csumb 
 		<< std::dec << std::endl;
+
+      if ( Params.saveSmeared ) {
+	truncate(smr);
+	BinaryIO::writeLatticeObject<vobj, sobj_double>(SmartConfig.get_U(true), smr, munge, 0, Params.format,
+							nersc_csum,scidac_csuma,scidac_csumb);
+	std::cout << GridLogMessage << "Written Binary Smeared Configuration " << smr
+                << " checksum " << std::hex 
+		<< nersc_csum   <<"/"
+		<< scidac_csuma   <<"/"
+		<< scidac_csumb 
+		<< std::dec << std::endl;
+      }
     }
 
   };

Grid/qcd/hmc/checkpointers/ILDGCheckpointer.h

@@ -69,17 +69,27 @@ public:
     }
   }
 
-  void TrajectoryComplete(int traj, GaugeField &U, GridSerialRNG &sRNG,
+  void TrajectoryComplete(int traj,
+			  ConfigurationBase<GaugeField> &SmartConfig,
+			  GridSerialRNG &sRNG,
                           GridParallelRNG &pRNG) {
     if ((traj % Params.saveInterval) == 0) {
-      std::string config, rng;
+      std::string config, rng, smr;
       this->build_filenames(traj, Params, config, rng);
-      GridBase *grid = U.Grid();
+      GridBase *grid = SmartConfig.get_U(false).Grid();
       uint32_t nersc_csum,scidac_csuma,scidac_csumb;
       BinaryIO::writeRNG(sRNG, pRNG, rng, 0,nersc_csum,scidac_csuma,scidac_csumb);
+      std::cout << GridLogMessage << "Written BINARY RNG " << rng
+                << " checksum " << std::hex 
+		<< nersc_csum<<"/"
+		<< scidac_csuma<<"/"
+		<< scidac_csumb
+		<< std::dec << std::endl;
+
+      
       IldgWriter _IldgWriter(grid->IsBoss());
       _IldgWriter.open(config);
-      _IldgWriter.writeConfiguration<GaugeStats>(U, traj, config, config);
+      _IldgWriter.writeConfiguration<GaugeStats>(SmartConfig.get_U(false), traj, config, config);
       _IldgWriter.close();
 
       std::cout << GridLogMessage << "Written ILDG Configuration on " << config
@@ -88,6 +98,21 @@ public:
 		<< scidac_csuma<<"/"
 		<< scidac_csumb
 		<< std::dec << std::endl;
+
+      if ( Params.saveSmeared ) { 
+	IldgWriter _IldgWriter(grid->IsBoss());
+	_IldgWriter.open(smr);
+	_IldgWriter.writeConfiguration<GaugeStats>(SmartConfig.get_U(true), traj, config, config);
+	_IldgWriter.close();
+
+	std::cout << GridLogMessage << "Written ILDG Configuration on " << smr
+                << " checksum " << std::hex 
+		<< nersc_csum<<"/"
+		<< scidac_csuma<<"/"
+		<< scidac_csumb
+		<< std::dec << std::endl;
+      }
+
     }
   };
 

Grid/qcd/hmc/checkpointers/NerscCheckpointer.h

@@ -52,23 +52,29 @@ public:
     Params.format = "IEEE64BIG";  // fixed, overwrite any other choice
   }
 
-  void TrajectoryComplete(int traj, GaugeField &U, GridSerialRNG &sRNG,
-                          GridParallelRNG &pRNG) {
+  virtual void TrajectoryComplete(int traj,
+                                  ConfigurationBase<GaugeField> &SmartConfig,
+                                  GridSerialRNG &sRNG,
+                                  GridParallelRNG &pRNG)
+  {
     if ((traj % Params.saveInterval) == 0) {
-      std::string config, rng;
-      this->build_filenames(traj, Params, config, rng);
+      std::string config, rng, smr;
+      this->build_filenames(traj, Params, config, smr, rng);
       
       int precision32 = 1;
       int tworow = 0;
       NerscIO::writeRNGState(sRNG, pRNG, rng);
-      NerscIO::writeConfiguration<GaugeStats>(U, config, tworow, precision32);
+      NerscIO::writeConfiguration<GaugeStats>(SmartConfig.get_U(false), config, tworow, precision32);
+      if ( Params.saveSmeared ) {
+	NerscIO::writeConfiguration<GaugeStats>(SmartConfig.get_U(true), smr, tworow, precision32);
+      }
     }
   };
 
   void CheckpointRestore(int traj, GaugeField &U, GridSerialRNG &sRNG,
                          GridParallelRNG &pRNG) {
-    std::string config, rng;
-    this->build_filenames(traj, Params, config, rng);
+    std::string config, rng, smr;
+    this->build_filenames(traj, Params, config, smr, rng );
     this->check_filename(rng);
     this->check_filename(config);
 

Grid/qcd/hmc/checkpointers/ScidacCheckpointer.h

@@ -70,19 +70,37 @@ class ScidacHmcCheckpointer : public BaseHmcCheckpointer<Implementation> {
     }
   }
 
-  void TrajectoryComplete(int traj, Field &U, GridSerialRNG &sRNG,
+  void TrajectoryComplete(int traj, 
+			  ConfigurationBase<Field> &SmartConfig,
+			  GridSerialRNG &sRNG,
                           GridParallelRNG &pRNG) {
     if ((traj % Params.saveInterval) == 0) {
-      std::string config, rng;
-      this->build_filenames(traj, Params, config, rng);
-      GridBase *grid = U.Grid();
+      std::string config, rng,smr;
+      this->build_filenames(traj, Params, config, smr, rng);
+      GridBase *grid = SmartConfig.get_U(false).Grid();
       uint32_t nersc_csum,scidac_csuma,scidac_csumb;
       BinaryIO::writeRNG(sRNG, pRNG, rng, 0,nersc_csum,scidac_csuma,scidac_csumb);
+      std::cout << GridLogMessage << "Written Binary RNG " << rng
+                << " checksum " << std::hex 
+		<< nersc_csum   <<"/"
+		<< scidac_csuma   <<"/"
+		<< scidac_csumb 
+		<< std::dec << std::endl;
+
+
+      {
 	ScidacWriter _ScidacWriter(grid->IsBoss());
 	_ScidacWriter.open(config);
-      _ScidacWriter.writeScidacFieldRecord(U, MData);
+	_ScidacWriter.writeScidacFieldRecord(SmartConfig.get_U(false), MData);
 	_ScidacWriter.close();
+      }
       
+      if ( Params.saveSmeared ) {
+	ScidacWriter _ScidacWriter(grid->IsBoss());
+	_ScidacWriter.open(smr);
+	_ScidacWriter.writeScidacFieldRecord(SmartConfig.get_U(true), MData);
+	_ScidacWriter.close();
+      }
       std::cout << GridLogMessage << "Written Scidac Configuration on " << config << std::endl;
     }
   };

Grid/qcd/hmc/integrators/Integrator.h

@@ -66,6 +66,7 @@ public:
 template <class FieldImplementation_, class SmearingPolicy, class RepresentationPolicy>
 class Integrator {
 protected:
+public:
   typedef FieldImplementation_ FieldImplementation;
   typedef typename FieldImplementation::Field MomentaField;  //for readability
   typedef typename FieldImplementation::Field Field;
@@ -96,7 +97,6 @@ protected:
   {
     t_P[level] += ep;
     update_P(P, U, level, ep);
-
     std::cout << GridLogIntegrator << "[" << level << "] P " << " dt " << ep << " : t_P " << t_P[level] << std::endl;
   }
 
@@ -130,28 +130,20 @@ protected:
       Field force(U.Grid());
       conformable(U.Grid(), Mom.Grid());
 
-      Field& Us = Smearer.get_U(as[level].actions.at(a)->is_smeared);
       double start_force = usecond();
 
-      std::cout << GridLogMessage << "AuditForce["<<level<<"]["<<a<<"] before"<<std::endl;
-      
       as[level].actions.at(a)->deriv_timer_start();
-      as[level].actions.at(a)->deriv(Us, force);  // deriv should NOT include Ta
+      as[level].actions.at(a)->deriv(Smearer, force);  // deriv should NOT include Ta
       as[level].actions.at(a)->deriv_timer_stop();
 
-      std::cout << GridLogMessage << "AuditForce["<<level<<"]["<<a<<"] after"<<std::endl;
-
-      std::cout << GridLogIntegrator << "Smearing (on/off): " << as[level].actions.at(a)->is_smeared << std::endl;
       auto name = as[level].actions.at(a)->action_name();
-      if (as[level].actions.at(a)->is_smeared) Smearer.smeared_force(force);
 
       force = FieldImplementation::projectForce(force); // Ta for gauge fields
       double end_force = usecond();
       
-      //      DumpSliceNorm("force ",force,Nd-1);
       MomFilter->applyFilter(force);
+
       std::cout << GridLogIntegrator << " update_P : Level [" << level <<"]["<<a <<"] "<<name<<" dt "<<ep<<  std::endl;
-      DumpSliceNorm("force filtered ",force,Nd-1);
       
       Real force_abs   = std::sqrt(norm2(force)/U.Grid()->gSites()); //average per-site norm.  nb. norm2(latt) = \sum_x norm2(latt[x]) 
       Real impulse_abs = force_abs * ep * HMC_MOMENTUM_DENOMINATOR;    
@@ -377,14 +369,9 @@ public:
 	auto name = as[level].actions.at(actionID)->action_name();
         std::cout << GridLogMessage << "refresh [" << level << "][" << actionID << "] "<<name << std::endl;
 
-        Field& Us = Smearer.get_U(as[level].actions.at(actionID)->is_smeared);
-
-	std::cout << GridLogMessage << "AuditRefresh["<<level<<"]["<<actionID<<"] before"<<std::endl;
-
 	as[level].actions.at(actionID)->refresh_timer_start();
-        as[level].actions.at(actionID)->refresh(Us, sRNG, pRNG);
+        as[level].actions.at(actionID)->refresh(Smearer, sRNG, pRNG);
 	as[level].actions.at(actionID)->refresh_timer_stop();
-	std::cout << GridLogMessage << "AuditRefresh["<<level<<"]["<<actionID<<"] after"<<std::endl;
 
       }
 
@@ -425,10 +412,9 @@ 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);
         std::cout << GridLogMessage << "S [" << level << "][" << actionID << "] action eval " << std::endl;
 	        as[level].actions.at(actionID)->S_timer_start();
-        Hterm = as[level].actions.at(actionID)->S(Us);
+        Hterm = as[level].actions.at(actionID)->S(Smearer);
    	        as[level].actions.at(actionID)->S_timer_stop();
         std::cout << GridLogMessage << "S [" << level << "][" << actionID << "] H = " << Hterm << std::endl;
         H += Hterm;
@@ -469,11 +455,10 @@ public:
       for (int actionID = 0; actionID < as[level].actions.size(); ++actionID) {
         // 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);
         std::cout << GridLogMessage << "S [" << level << "][" << actionID << "] action eval " << std::endl;
-	        as[level].actions.at(actionID)->S_timer_start();
 
-        Hterm = as[level].actions.at(actionID)->Sinitial(Us);
+	as[level].actions.at(actionID)->S_timer_start();
+        Hterm = as[level].actions.at(actionID)->S(Smearer);
 	as[level].actions.at(actionID)->S_timer_stop();
 
         std::cout << GridLogMessage << "S [" << level << "][" << actionID << "] H = " << Hterm << std::endl;

Grid/qcd/observables/hmc_observable.h

@@ -34,6 +34,13 @@ NAMESPACE_BEGIN(Grid);
 template <class Field>
 class HmcObservable {
  public:
+  virtual void TrajectoryComplete(int traj,
+                                  ConfigurationBase<Field> &SmartConfig,
+                                  GridSerialRNG &sRNG,
+                                  GridParallelRNG &pRNG)
+  {
+    TrajectoryComplete(traj,SmartConfig.get_U(false),sRNG,pRNG); // Unsmeared observable
+  };
   virtual void TrajectoryComplete(int traj,
                                   Field &U,
                                   GridSerialRNG &sRNG,

Grid/qcd/observables/plaquette.h

@@ -42,6 +42,18 @@ public:
   // necessary for HmcObservable compatibility
   typedef typename Impl::Field Field;
 
+  virtual void TrajectoryComplete(int traj,
+                                  ConfigurationBase<Field> &SmartConfig,
+                                  GridSerialRNG &sRNG,
+                                  GridParallelRNG &pRNG)
+  {
+    std::cout << GridLogMessage << "+++++++++++++++++++"<<std::endl;
+    std::cout << GridLogMessage << "Unsmeared plaquette"<<std::endl;
+    TrajectoryComplete(traj,SmartConfig.get_U(false),sRNG,pRNG); // Unsmeared observable
+    std::cout << GridLogMessage << "Smeared plaquette"<<std::endl;
+    TrajectoryComplete(traj,SmartConfig.get_U(true),sRNG,pRNG); // Unsmeared observable
+    std::cout << GridLogMessage << "+++++++++++++++++++"<<std::endl;
+  };
   void TrajectoryComplete(int traj,
                           Field &U,
                           GridSerialRNG &sRNG,

Grid/qcd/smearing/GaugeConfiguration.h

@@ -19,13 +19,13 @@ public:
 
   NoSmearing(): ThinLinks(NULL) {}
 
-  void set_Field(Field& U) { ThinLinks = &U; }
+  virtual void set_Field(Field& U) { ThinLinks = &U; }
 
-  void smeared_force(Field&) const {}
+  virtual void smeared_force(Field&) {}
 
-  Field& get_SmearedU() { return *ThinLinks; }
+  virtual Field& get_SmearedU() { return *ThinLinks; }
 
-  Field &get_U(bool smeared = false)
+  virtual Field &get_U(bool smeared = false)
   {
     return *ThinLinks;
   }
@@ -235,7 +235,7 @@ public:
     : smearingLevels(0), StoutSmearing(nullptr), SmearedSet(), ThinLinks(NULL) {}
 
   // attach the smeared routines to the thin links U and fill the smeared set
-  void set_Field(GaugeField &U)
+  virtual void set_Field(GaugeField &U)
   {
     double start = usecond();
     fill_smearedSet(U);
@@ -245,7 +245,7 @@ public:
   }
 
   //====================================================================
-  void smeared_force(GaugeField &SigmaTilde) const
+  virtual void smeared_force(GaugeField &SigmaTilde) 
   {
     if (smearingLevels > 0)
     {
@@ -272,14 +272,16 @@ public:
       }
       double end = usecond();
       double time = (end - start)/ 1e3;
-      std::cout << GridLogMessage << "Smearing force in " << time << " ms" << std::endl;  
+      std::cout << GridLogMessage << " GaugeConfiguration: Smeared Force chain rule took " << time << " ms" << std::endl;
     }  // if smearingLevels = 0 do nothing
+    SigmaTilde=Gimpl::projectForce(SigmaTilde); // Ta
+      
   }
   //====================================================================
 
-  GaugeField& get_SmearedU() { return SmearedSet[smearingLevels - 1]; }
+  virtual GaugeField& get_SmearedU() { return SmearedSet[smearingLevels - 1]; }
 
-  GaugeField &get_U(bool smeared = false)
+  virtual GaugeField &get_U(bool smeared = false)
   {
     // get the config, thin links by default
     if (smeared)

Grid/qcd/smearing/GaugeConfigurationMasked.h

@@ -131,6 +131,7 @@ public:
     AdjMatrixField  X(grid);
     Complex ci(0,1);
 
+    RealD t0 = usecond();
     Ident = ComplexD(1.0);
     for(int d=0;d<Nd;d++){
       Umu[d] = peekLorentz(U, d);
@@ -161,6 +162,8 @@ public:
     // Assemble the N matrix
     //////////////////////////////////////////////////////////////////
     // Computes ALL the staples -- could compute one only and do it here
+    RealD time;
+    time=-usecond();
     this->StoutSmearing->BaseSmear(C, U);
     Cmu = peekLorentz(C, mu);
 
@@ -169,7 +172,10 @@ public:
     //////////////////////////////////////////////////////////////////
     // Ta so Z lives in Lie algabra
     Zx  = Ta(Cmu * adj(Umu[mu]));
+    time+=usecond();
+    std::cout << GridLogMessage << "Z took "<<time<< " us"<<std::endl;
 
+    time=-usecond();
     // Move Z to the Adjoint Rep == make_adjoint_representation
     ZxAd = Zero();
     for(int b=0;b<8;b++) {
@@ -180,10 +186,13 @@ public:
       cplx = 2.0*trace(ci*tb*Zx); // my convention 1/2 delta ba
       ZxAd = ZxAd + cplx * TRb; // is this right? YES - Guido used Anti herm Ta's and with bloody wrong sign.
     }
+    time+=usecond();
+    std::cout << GridLogMessage << "ZxAd took "<<time<< " us"<<std::endl;
 
     //////////////////////////////////////
     // J(x) = 1 + Sum_k=1..N (-Zac)^k/(k+1)!
     //////////////////////////////////////
+    time=-usecond();
     X=1.0; 
     JxAd = X;
     mZxAd = (-1.0)*ZxAd; 
@@ -193,10 +202,13 @@ public:
       kpfac = kpfac /(k+1);
       JxAd = JxAd + X * kpfac;
     }
+    time+=usecond();
+    std::cout << GridLogMessage << "Jx took "<<time<< " us"<<std::endl;
 
     //////////////////////////////////////
     // dJ(x)/dxe
     //////////////////////////////////////
+    time=-usecond();
     std::vector<AdjMatrixField>  dJdX;    dJdX.resize(8,grid);
     AdjMatrixField tbXn(grid);
     AdjMatrixField sumXtbX(grid);
@@ -220,12 +232,17 @@ public:
       }
       dJdX[b] = -dt2; 
     }
+    time+=usecond();
+    std::cout << GridLogMessage << "dJx took "<<time<< " us"<<std::endl;
     /////////////////////////////////////////////////////////////////
     // Mask Umu for this link
     /////////////////////////////////////////////////////////////////
+    time=-usecond();
     PlaqL = Ident;
     PlaqR = Utmp*adj(Cmu);
     ComputeNxy(PlaqL,PlaqR,NxxAd);
+    time+=usecond();
+    std::cout << GridLogMessage << "ComputeNxy took "<<time<< " us"<<std::endl;
     
     ////////////////////////////
     // Mab
@@ -236,8 +253,12 @@ public:
     /////////////////////////
     // invert the 8x8
     /////////////////////////
+    time=-usecond();
     MpAdInv = Inverse(MpAd);
+    time+=usecond();
+    std::cout << GridLogMessage << "MpAdInv took "<<time<< " us"<<std::endl;
     
+    RealD t3a = usecond();
     /////////////////////////////////////////////////////////////////
     // Nxx Mp^-1
     /////////////////////////////////////////////////////////////////
@@ -283,6 +304,7 @@ public:
     GaugeField Fdet2(grid);
     GaugeLinkField Fdet_pol(grid); // one polarisation
 
+    RealD t4 = usecond();
     for(int nu=0;nu<Nd;nu++){
 
       if (nu!=mu) {
@@ -291,20 +313,29 @@ public:
 	//    |  |
 	//    x==    // nu polarisation -- clockwise
 
+	time=-usecond();
 	PlaqL=Ident;
 
 	PlaqR=(-rho)*Gimpl::CovShiftForward(Umu[nu], nu,
  	       Gimpl::CovShiftForward(Umu[mu], mu,
 	         Gimpl::CovShiftBackward(Umu[nu], nu,
 		   Gimpl::CovShiftIdentityBackward(Utmp, mu))));
+	time+=usecond();
+	std::cout << GridLogMessage << "PlaqLR took "<<time<< " us"<<std::endl;
 
+	time=-usecond();
 	dJdXe_nMpInv_y =   dJdXe_nMpInv;
 	ComputeNxy(PlaqL,PlaqR,Nxy);
 	Fdet1_nu = transpose(Nxy)*dJdXe_nMpInv_y;
+	time+=usecond();
+	std::cout << GridLogMessage << "ComputeNxy (occurs 6x) took "<<time<< " us"<<std::endl;
 
+	time=-usecond();
 	PlaqR=(-1.0)*PlaqR;
 	Compute_MpInvJx_dNxxdSy(PlaqL,PlaqR,MpInvJx,FdetV);
 	Fdet2_nu = FdetV;
+	time+=usecond();
+	std::cout << GridLogMessage << "Compute_MpInvJx_dNxxSy (occurs 6x) took "<<time<< " us"<<std::endl;
 	
 	//    x==
 	//    |  |
@@ -416,6 +447,7 @@ public:
 	
       }
     }
+    RealD t5 = usecond();
 
     Fdet1_mu = Fdet1_mu + transpose(NxxAd)*dJdXe_nMpInv;
 
@@ -423,6 +455,13 @@ public:
     InsertForce(Fdet2,Fdet2_mu,mu);
 
     force= (-0.5)*( Fdet1 + Fdet2);
+    RealD t1 = usecond();
+    std::cout << GridLogMessage << " logDetJacobianForce level took "<<t1-t0<<" us "<<std::endl;
+    std::cout << GridLogMessage << " logDetJacobianForce t3-t0 "<<t3a-t0<<" us "<<std::endl;
+    std::cout << GridLogMessage << " logDetJacobianForce t4-t3 dJdXe_nMpInv "<<t4-t3a<<" us "<<std::endl;
+    std::cout << GridLogMessage << " logDetJacobianForce t5-t4 mu nu loop "<<t5-t4<<" us "<<std::endl;
+    std::cout << GridLogMessage << " logDetJacobianForce t1-t5 "<<t1-t5<<" us "<<std::endl;
+    std::cout << GridLogMessage << " logDetJacobianForce level took "<<t1-t0<<" us "<<std::endl;
   }
   RealD logDetJacobianLevel(const GaugeField &U,int smr)
   {
@@ -696,10 +735,10 @@ private:
 public:
 
   /* Standard constructor */
-  SmearedConfigurationMasked(GridCartesian* _UGrid, unsigned int Nsmear, Smear_Stout<Gimpl>& Stout,bool domask=false)
+  SmearedConfigurationMasked(GridCartesian* _UGrid, unsigned int Nsmear, Smear_Stout<Gimpl>& Stout)
     : SmearedConfiguration<Gimpl>(_UGrid, Nsmear,Stout)
   {
-    if(domask) assert(Nsmear%(2*Nd)==0); // Or multiply by 8??
+    assert(Nsmear%(2*Nd)==0); // Or multiply by 8??
 
     // was resized in base class
     assert(this->SmearedSet.size()==Nsmear);
@@ -712,7 +751,6 @@ public:
     for (unsigned int i = 0; i < this->smearingLevels; ++i) {
 
       masks.push_back(*(new LatticeLorentzComplex(_UGrid)));
-      if (domask) {
 
       int mu= (i/2) %Nd;
       int cb= (i%2);
@@ -727,11 +765,6 @@ public:
       setCheckerboard(tmp,tmpcb);
       PokeIndex<LorentzIndex>(masks[i],tmp, mu);
 	
-      } else {
-	for(int mu=0;mu<Nd;mu++){
-	  PokeIndex<LorentzIndex>(masks[i],one, mu);
-	}
-      }
     }
     delete UrbGrid;
   }
@@ -764,10 +797,14 @@ public:
         tmp_mu = peekLorentz(*this->ThinLinks, mu) * peekLorentz(force, mu);
         pokeLorentz(SigmaTilde, tmp_mu, mu);
       }
+
+
       double end = usecond();
       double time = (end - start)/ 1e3;
       std::cout << GridLogMessage << " GaugeConfigurationMasked: Smeared Force chain rule took " << time << " ms" << std::endl;
+
     }  // if smearingLevels = 0 do nothing
+    SigmaTilde=Gimpl::projectForce(SigmaTilde); // Ta
   }
 
 };

Grid/qcd/smearing/JacobianAction.h

@@ -2,15 +2,11 @@
 
 Grid physics library, www.github.com/paboyle/Grid
 
-Source file: ./lib/qcd/action/gauge/WilsonGaugeAction.h
+Source file: ./lib/qcd/action/gauge/JacobianAction.h
 
 Copyright (C) 2015
 
-Author: Azusa Yamaguchi <ayamaguc@staffmail.ed.ac.uk>
 Author: Peter Boyle <paboyle@ph.ed.ac.uk>
-Author: neo <cossu@post.kek.jp>
-Author: paboyle <paboyle@ph.ed.ac.uk>
-Author: Guido Cossu <guido.cossu@ed.ac.uk>
 
 This program is free software; you can redistribute it and/or modify
 it under the terms of the GNU General Public License as published by

Grid/qcd/utils/SUn.h

@@ -34,6 +34,61 @@ directory
 
 NAMESPACE_BEGIN(Grid);
 
+template<int N, class Vec>
+Lattice<iScalar<iScalar<iScalar<Vec> > > > Determinant(const Lattice<iScalar<iScalar<iMatrix<Vec, N> > > > &Umu)
+{
+  GridBase *grid=Umu.Grid();
+  auto lvol = grid->lSites();
+  Lattice<iScalar<iScalar<iScalar<Vec> > > > 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<iScalar<iMatrix<scalar, N> > > Us;
+    peekLocalSite(Us, Umu_v, lcoor);
+    for(int i=0;i<N;i++){
+      for(int j=0;j<N;j++){
+	scalar tmp= Us()()(i,j);
+	ComplexD ztmp(real(tmp),imag(tmp));
+	EigenU(i,j)=ztmp;
+      }}
+    ComplexD detD  = EigenU.determinant();
+    typename Vec::scalar_type det(detD.real(),detD.imag());
+    pokeLocalSite(det,ret_v,lcoor);
+  });
+  return ret;
+}
+
+template<int N, class Vec>
+static void ProjectSUn(Lattice<iScalar<iScalar<iMatrix<Vec, N> > > > &Umu)
+{
+  Umu      = ProjectOnGroup(Umu);
+  auto det = Determinant(Umu);
+
+  det = conjugate(det);
+
+  for(int i=0;i<N;i++){
+    auto element = PeekIndex<ColourIndex>(Umu,N-1,i);
+    element = element * det;
+    PokeIndex<ColourIndex>(Umu,element,Nc-1,i);
+  }
+}
+template<int N,class Vec>
+static void ProjectSUn(Lattice<iVector<iScalar<iMatrix<Vec, N> >,Nd> > &U)
+{
+  GridBase *grid=U.Grid();
+  // Reunitarise
+  for(int mu=0;mu<Nd;mu++){
+    auto Umu = PeekIndex<LorentzIndex>(U,mu);
+    Umu      = ProjectOnGroup(Umu);
+    ProjectSUn(Umu);
+    PokeIndex<LorentzIndex>(U,Umu,mu);
+  }
+}
+
 template <int ncolour>
 class SU {
 public:
@@ -741,8 +796,14 @@ public:
     typedef Lattice<vMatrixType> LatticeMatrixType;
 
     LatticeMatrixType Umu(out.Grid());
+    LatticeMatrixType tmp(out.Grid());
     for (int mu = 0; mu < Nd; mu++) {
-      LieRandomize(pRNG, Umu, 1.0);
+      //      LieRandomize(pRNG, Umu, 1.0);
+      //      PokeIndex<LorentzIndex>(out, Umu, mu);
+      gaussian(pRNG,Umu);
+      tmp = Ta(Umu);
+      taExp(tmp,Umu);
+      ProjectSUn(Umu);
       PokeIndex<LorentzIndex>(out, Umu, mu);
     }
   }
@@ -798,30 +859,6 @@ public:
   }
 };
 
-template<int N>
-LatticeComplexD Determinant(const Lattice<iScalar<iScalar<iMatrix<vComplexD, N> > > > &Umu)
-{
-  GridBase *grid=Umu.Grid();
-  auto lvol = grid->lSites();
-  LatticeComplexD 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<iScalar<iMatrix<ComplexD, N> > > Us;
-    peekLocalSite(Us, Umu_v, lcoor);
-    for(int i=0;i<N;i++){
-      for(int j=0;j<N;j++){
-	EigenU(i,j) = Us()()(i,j);
-      }}
-    ComplexD det = EigenU.determinant();
-    pokeLocalSite(det,ret_v,lcoor);
-  });
-  return ret;
-}
 template<int N>
 Lattice<iScalar<iScalar<iMatrix<vComplexD, N> > > > Inverse(const Lattice<iScalar<iScalar<iMatrix<vComplexD, N> > > > &Umu)
 {
@@ -851,32 +888,6 @@ Lattice<iScalar<iScalar<iMatrix<vComplexD, N> > > > Inverse(const Lattice<iScala
   });
   return ret;
 }
-template<int N>
-static void ProjectSUn(Lattice<iScalar<iScalar<iMatrix<vComplexD, N> > > > &Umu)
-{
-  Umu      = ProjectOnGroup(Umu);
-  auto det = Determinant(Umu);
-
-  det = conjugate(det);
-
-  for(int i=0;i<N;i++){
-    auto element = PeekIndex<ColourIndex>(Umu,N-1,i);
-    element = element * det;
-    PokeIndex<ColourIndex>(Umu,element,Nc-1,i);
-  }
-}
-template<int N>
-static void ProjectSUn(Lattice<iVector<iScalar<iMatrix<vComplexD, N> >,Nd> > &U)
-{
-  GridBase *grid=U.Grid();
-  // Reunitarise
-  for(int mu=0;mu<Nd;mu++){
-    auto Umu = PeekIndex<LorentzIndex>(U,mu);
-    Umu      = ProjectOnGroup(Umu);
-    ProjectSUn(Umu);
-    PokeIndex<LorentzIndex>(U,Umu,mu);
-  }
-}
 // Explicit specialisation for SU(3).
 // Explicit specialisation for SU(3).
 static void

Grid/stencil/Stencil.h

@@ -705,7 +705,7 @@ public:
 	}
       }
     }
-    std::cout << "BuildSurfaceList size is "<<surface_list.size()<<std::endl;
+    std::cout << GridLogDebug << "BuildSurfaceList size is "<<surface_list.size()<<std::endl;
   }
   /// Introduce a block structure and switch off comms on boundaries
   void DirichletBlock(const Coordinate &dirichlet_block)

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;
   }

Grid/tensors/Tensor_exp.h

@@ -55,7 +55,7 @@ template<class vtype, int N> accelerator_inline iVector<vtype, N> Exponentiate(c
 
 
 // Specialisation: Cayley-Hamilton exponential for SU(3)
-#ifndef GRID_ACCELERATED
+#if 0
 template<class vtype, typename std::enable_if< GridTypeMapper<vtype>::TensorLevel == 0>::type * =nullptr> 
 accelerator_inline iMatrix<vtype,3> Exponentiate(const iMatrix<vtype,3> &arg, RealD alpha  , Integer Nexp = DEFAULT_MAT_EXP )
 {

HMC/FTHMC2p1f.cc

@@ -0,0 +1,224 @@
+/*************************************************************************************
+
+Grid physics library, www.github.com/paboyle/Grid
+
+Copyright (C) 2023
+
+Author: Peter Boyle <pabobyle@ph.ed.ac.uk>
+
+This program is free software; you can redistribute it and/or modify
+it under the terms of the GNU General Public License as published by
+the Free Software Foundation; either version 2 of the License, or
+(at your option) any later version.
+
+This program is distributed in the hope that it will be useful,
+but WITHOUT ANY WARRANTY; without even the implied warranty of
+MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+GNU General Public License for more details.
+
+You should have received a copy of the GNU General Public License along
+with this program; if not, write to the Free Software Foundation, Inc.,
+51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
+
+See the full license in the file "LICENSE" in the top level distribution
+directory
+*************************************************************************************/
+/*  END LEGAL */
+#include <Grid/Grid.h>
+#include <Grid/qcd/smearing/GaugeConfigurationMasked.h>
+#include <Grid/qcd/smearing/JacobianAction.h>
+
+using namespace Grid;
+
+int main(int argc, char **argv)
+{
+  std::cout << std::setprecision(12);
+  
+  Grid_init(&argc, &argv);
+  int threads = GridThread::GetThreads();
+  // here make a routine to print all the relevant information on the run
+  std::cout << GridLogMessage << "Grid is setup to use " << threads << " threads" << std::endl;
+
+   // Typedefs to simplify notation
+  typedef WilsonImplR FermionImplPolicy;
+  typedef MobiusFermionD FermionAction;
+  typedef typename FermionAction::FermionField FermionField;
+
+  typedef Grid::XmlReader       Serialiser;
+
+  //::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::
+  IntegratorParameters MD;
+  //  typedef GenericHMCRunner<LeapFrog> HMCWrapper;
+  //  MD.name    = std::string("Leap Frog");
+  //  typedef GenericHMCRunner<ForceGradient> HMCWrapper;
+  //  MD.name    = std::string("Force Gradient");
+  typedef GenericHMCRunner<MinimumNorm2> HMCWrapper;
+  MD.name    = std::string("MinimumNorm2");
+  MD.MDsteps = 12;
+  MD.trajL   = 1.0;
+
+  HMCparameters HMCparams;
+  HMCparams.StartTrajectory  = 0;
+  HMCparams.Trajectories     = 200;
+  HMCparams.NoMetropolisUntil=  20;
+  // "[HotStart, ColdStart, TepidStart, CheckpointStart]\n";
+  HMCparams.StartingType     =std::string("HotStart");
+  HMCparams.MD = MD;
+  HMCWrapper TheHMC(HMCparams);
+
+  // Grid from the command line arguments --grid and --mpi
+  TheHMC.Resources.AddFourDimGrid("gauge"); // use default simd lanes decomposition
+
+  CheckpointerParameters CPparams;
+  CPparams.config_prefix = "ckpoint_EODWF_lat";
+  CPparams.smeared_prefix = "ckpoint_EODWF_lat_smr";
+  CPparams.rng_prefix    = "ckpoint_EODWF_rng";
+  CPparams.saveInterval  = 1;
+  CPparams.saveSmeared   = true;
+  CPparams.format        = "IEEE64BIG";
+  TheHMC.Resources.LoadNerscCheckpointer(CPparams);
+
+  RNGModuleParameters RNGpar;
+  RNGpar.serial_seeds = "1 2 3 4 5";
+  RNGpar.parallel_seeds = "6 7 8 9 10";
+  TheHMC.Resources.SetRNGSeeds(RNGpar);
+
+  // Construct observables
+  // here there is too much indirection
+  typedef PlaquetteMod<HMCWrapper::ImplPolicy> PlaqObs;
+  TheHMC.Resources.AddObservable<PlaqObs>();
+  //////////////////////////////////////////////
+
+  const int Ls      = 16;
+  Real beta         = 2.13;
+  Real light_mass   = 0.01;
+  Real strange_mass = 0.04;
+  Real pv_mass      = 1.0;
+  RealD M5  = 1.8;
+  RealD b   = 1.0; // Scale factor two
+  RealD c   = 0.0;
+
+  OneFlavourRationalParams OFRp;
+  OFRp.lo       = 1.0e-2;
+  OFRp.hi       = 64;
+  OFRp.MaxIter  = 10000;
+  OFRp.tolerance= 1.0e-10;
+  OFRp.degree   = 14;
+  OFRp.precision= 40;
+
+  std::vector<Real> hasenbusch({ 0.1 });
+
+  auto GridPtr   = TheHMC.Resources.GetCartesian();
+  auto GridRBPtr = TheHMC.Resources.GetRBCartesian();
+  auto FGrid     = SpaceTimeGrid::makeFiveDimGrid(Ls,GridPtr);
+  auto FrbGrid   = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,GridPtr);
+
+  IwasakiGaugeActionR GaugeAction(beta);
+
+  // temporarily need a gauge field
+  LatticeGaugeField U(GridPtr);
+  LatticeGaugeField Uhot(GridPtr);
+
+  // These lines are unecessary if BC are all periodic
+  std::vector<Complex> boundary = {1,1,1,-1};
+  FermionAction::ImplParams Params(boundary);
+
+  double StoppingCondition = 1e-10;
+  double MaxCGIterations = 30000;
+  ConjugateGradient<FermionField>  CG(StoppingCondition,MaxCGIterations);
+
+  bool ApplySmearing = true;
+  
+  ////////////////////////////////////
+  // Collect actions
+  ////////////////////////////////////
+  ActionLevel<HMCWrapper::Field> Level1(1);
+  ActionLevel<HMCWrapper::Field> Level2(2);
+  ActionLevel<HMCWrapper::Field> Level3(4);
+
+  ////////////////////////////////////
+  // Strange action
+  ////////////////////////////////////
+
+  MobiusEOFAFermionD Strange_Op_L (U , *FGrid , *FrbGrid , *GridPtr , *GridRBPtr , strange_mass, strange_mass, pv_mass, 0.0, -1, M5, b, c);
+  MobiusEOFAFermionD Strange_Op_R (U , *FGrid , *FrbGrid , *GridPtr , *GridRBPtr , pv_mass, strange_mass,      pv_mass, -1.0, 1, M5, b, c);
+  ExactOneFlavourRatioPseudoFermionAction<FermionImplPolicy> 
+    EOFA(Strange_Op_L, Strange_Op_R, 
+	 CG,
+	 CG, CG,
+	 CG, CG, 
+	 OFRp, false);
+
+  EOFA.is_smeared = ApplySmearing;
+  Level1.push_back(&EOFA);
+
+  ////////////////////////////////////
+  // up down action
+  ////////////////////////////////////
+  std::vector<Real> light_den;
+  std::vector<Real> light_num;
+
+  int n_hasenbusch = hasenbusch.size();
+  light_den.push_back(light_mass);
+  for(int h=0;h<n_hasenbusch;h++){
+    light_den.push_back(hasenbusch[h]);
+    light_num.push_back(hasenbusch[h]);
+  }
+  light_num.push_back(pv_mass);
+
+  std::vector<FermionAction *> Numerators;
+  std::vector<FermionAction *> Denominators;
+  std::vector<TwoFlavourEvenOddRatioPseudoFermionAction<FermionImplPolicy> *> Quotients;
+
+  for(int h=0;h<n_hasenbusch+1;h++){
+    std::cout << GridLogMessage << " 2f quotient Action  "<< light_num[h] << " / " << light_den[h]<< std::endl;
+    Numerators.push_back  (new FermionAction(U,*FGrid,*FrbGrid,*GridPtr,*GridRBPtr,light_num[h],M5,b,c, Params));
+    Denominators.push_back(new FermionAction(U,*FGrid,*FrbGrid,*GridPtr,*GridRBPtr,light_den[h],M5,b,c, Params));
+    Quotients.push_back   (new TwoFlavourEvenOddRatioPseudoFermionAction<FermionImplPolicy>(*Numerators[h],*Denominators[h],CG,CG));
+  }
+
+  for(int h=0;h<n_hasenbusch+1;h++){
+    Quotients[h]->is_smeared = ApplySmearing;
+    Level1.push_back(Quotients[h]);
+  }
+
+  /////////////////////////////////////////////////////////////
+  // lnDetJacobianAction
+  /////////////////////////////////////////////////////////////
+  double rho = 0.1;  // smearing parameter
+  int Nsmear = 1;    // number of smearing levels - must be multiple of 2Nd
+  int Nstep  = 8*Nsmear;    // number of smearing levels - must be multiple of 2Nd
+  Smear_Stout<HMCWrapper::ImplPolicy> Stout(rho);
+  SmearedConfigurationMasked<HMCWrapper::ImplPolicy> SmearingPolicy(GridPtr, Nstep, Stout);
+  JacobianAction<HMCWrapper::ImplPolicy> Jacobian(&SmearingPolicy);
+  if( ApplySmearing ) Level2.push_back(&Jacobian);
+  std::cout << GridLogMessage << " Built the Jacobian "<< std::endl;
+
+
+  /////////////////////////////////////////////////////////////
+  // Gauge action
+  /////////////////////////////////////////////////////////////
+  //  GaugeAction.is_smeared = ApplySmearing;
+  GaugeAction.is_smeared = true;
+  Level3.push_back(&GaugeAction);
+
+  std::cout << GridLogMessage << " ************************************************"<< std::endl;
+  std::cout << GridLogMessage << " Action complete -- NO FERMIONS FOR NOW -- FIXME"<< std::endl;
+  std::cout << GridLogMessage << " ************************************************"<< std::endl;
+  std::cout << GridLogMessage <<  std::endl;
+  std::cout << GridLogMessage <<  std::endl;
+
+
+  std::cout << GridLogMessage << " Running the FT HMC "<< std::endl;
+
+  TheHMC.TheAction.push_back(Level1);
+  TheHMC.TheAction.push_back(Level2);
+  TheHMC.TheAction.push_back(Level3);
+
+  TheHMC.Run(SmearingPolicy); // for smearing
+
+  Grid_finalize();
+} // main
+
+
+

HMC/Mobius2p1f_EOFA_96I_hmc.cc

@@ -146,6 +146,8 @@ NAMESPACE_END(Grid);
 int main(int argc, char **argv) {
   using namespace Grid;
 
+  std::cout << " Grid Initialise "<<std::endl;
+  
   Grid_init(&argc, &argv);
 
   CartesianCommunicator::BarrierWorld();
@@ -170,24 +172,24 @@ int main(int argc, char **argv) {
   IntegratorParameters MD;
   //  typedef GenericHMCRunner<LeapFrog> HMCWrapper;
   //  MD.name    = std::string("Leap Frog");
-  typedef GenericHMCRunner<ForceGradient> HMCWrapper;
-  MD.name    = std::string("Force Gradient");
-  //typedef GenericHMCRunner<MinimumNorm2> HMCWrapper;
-  // MD.name    = std::string("MinimumNorm2");
+  //  typedef GenericHMCRunner<ForceGradient> HMCWrapper;
+  //  MD.name    = std::string("Force Gradient");
+  typedef GenericHMCRunner<MinimumNorm2> HMCWrapper;
+  MD.name    = std::string("MinimumNorm2");
   // TrajL = 2
   // 4/2 => 0.6 dH
   // 3/3 => 0.8 dH .. depth 3, slower
   //MD.MDsteps =  4;
-  MD.MDsteps =  12;
+  MD.MDsteps =  14;
   MD.trajL   = 0.5;
 
   HMCparameters HMCparams;
   HMCparams.StartTrajectory  = 1077;
-  HMCparams.Trajectories     = 1;
+  HMCparams.Trajectories     = 20;
   HMCparams.NoMetropolisUntil=  0;
   // "[HotStart, ColdStart, TepidStart, CheckpointStart]\n";
-  //  HMCparams.StartingType     =std::string("ColdStart");
-  HMCparams.StartingType     =std::string("CheckpointStart");
+  HMCparams.StartingType     =std::string("ColdStart");
+  //  HMCparams.StartingType     =std::string("CheckpointStart");
   HMCparams.MD = MD;
   HMCWrapper TheHMC(HMCparams);
 
@@ -223,7 +225,7 @@ int main(int argc, char **argv) {
   Real pv_mass      = 1.0;
   //  std::vector<Real> hasenbusch({ 0.01, 0.045, 0.108, 0.25, 0.51 , pv_mass });
   //  std::vector<Real> hasenbusch({ light_mass, 0.01, 0.045, 0.108, 0.25, 0.51 , pv_mass });
-  std::vector<Real> hasenbusch({ 0.005, 0.0145, 0.045, 0.108, 0.25, 0.51 , pv_mass }); // Updated
+  std::vector<Real> hasenbusch({ 0.005, 0.0145, 0.045, 0.108, 0.25, 0.51 }); // Updated
   //  std::vector<Real> hasenbusch({ light_mass, 0.0145, 0.045, 0.108, 0.25, 0.51 , 0.75 , pv_mass });
 
   auto GridPtr   = TheHMC.Resources.GetCartesian();
@@ -275,10 +277,10 @@ int main(int argc, char **argv) {
 
   //  double StoppingCondition = 1e-14;
   //  double MDStoppingCondition = 1e-9;
-  double StoppingCondition = 1e-8;
-  double MDStoppingCondition = 1e-7;
-  double MDStoppingConditionLoose = 1e-7;
-  double MDStoppingConditionStrange = 1e-7;
+  double StoppingCondition = 1e-9;
+  double MDStoppingCondition = 1e-8;
+  double MDStoppingConditionLoose = 1e-8;
+  double MDStoppingConditionStrange = 1e-8;
   double MaxCGIterations = 300000;
   ConjugateGradient<FermionField>  CG(StoppingCondition,MaxCGIterations);
   ConjugateGradient<FermionField>  MDCG(MDStoppingCondition,MaxCGIterations);

systems/Lumi/benchmarks/bench2.slurm

@@ -0,0 +1,44 @@
+#!/bin/bash -l
+#SBATCH --job-name=bench_lehner
+#SBATCH --partition=small-g
+#SBATCH --nodes=2
+#SBATCH --ntasks-per-node=8
+#SBATCH --cpus-per-task=7
+#SBATCH --gpus-per-node=8
+#SBATCH --time=00:10:00
+#SBATCH --account=project_465000546
+#SBATCH --gpu-bind=none
+#SBATCH --exclusive
+#SBATCH --mem=0
+
+CPU_BIND="map_cpu:48,56,32,40,16,24,1,8"
+echo $CPU_BIND
+
+cat << EOF > select_gpu
+#!/bin/bash
+export GPU_MAP=(0 1 2 3 4 5 6 7)
+export GPU=\${GPU_MAP[\$SLURM_LOCALID]}
+export HIP_VISIBLE_DEVICES=\$GPU
+unset ROCR_VISIBLE_DEVICES
+echo RANK \$SLURM_LOCALID using GPU \$GPU    
+exec \$*
+EOF
+
+chmod +x ./select_gpu
+
+root=/scratch/project_465000546/boylepet/Grid/systems/Lumi
+source ${root}/sourceme.sh
+
+export OMP_NUM_THREADS=7
+export MPICH_GPU_SUPPORT_ENABLED=1
+export MPICH_SMP_SINGLE_COPY_MODE=XPMEM
+
+for vol in 16.16.16.64 32.32.32.64  32.32.32.128
+do
+srun --cpu-bind=${CPU_BIND} ./select_gpu ./Benchmark_dwf_fp32 --mpi 2.2.2.2 --accelerator-threads 8 --comms-overlap --shm 2048 --shm-mpi 0 --grid $vol  > log.shm0.ov.$vol
+#srun --cpu-bind=${CPU_BIND} ./select_gpu ./Benchmark_dwf_fp32 --mpi 2.2.2.2 --accelerator-threads 8 --comms-overlap --shm 2048 --shm-mpi 1 --grid $vol  > log.shm1.ov.$vol
+
+srun --cpu-bind=${CPU_BIND} ./select_gpu ./Benchmark_dwf_fp32 --mpi 2.2.2.2 --accelerator-threads 8 --comms-sequential --shm 2048 --shm-mpi 0 --grid $vol  > log.shm0.seq.$vol
+#srun --cpu-bind=${CPU_BIND} ./select_gpu ./Benchmark_dwf_fp32 --mpi 2.2.2.2 --accelerator-threads 8 --comms-sequential --shm 2048 --shm-mpi 1 --grid $vol > log.shm1.seq.$vol
+done
+

systems/Lumi/config-command

@@ -3,30 +3,28 @@ spack load gmp
 spack load mpfr
 CLIME=`spack find --paths c-lime | grep c-lime| cut -c 15-`
 GMP=`spack find --paths gmp | grep gmp | cut -c 12-`
-MPFR=`spack find --paths mpfr | grep mpfr | cut -c 12-`
-echo clime $CLIME
-echo gmp $GMP
-echo mpfr $MPFR
+MPFR=`spack find --paths mpfr | grep mpfr | cut -c 13-`
+echo clime X$CLIME
+echo gmp X$GMP
+echo mpfr X$MPFR
 
-../../configure --enable-comms=mpi-auto \
+../../configure \
+--enable-comms=mpi-auto \
 --with-lime=$CLIME \
 --enable-unified=no \
 --enable-shm=nvlink \
---enable-tracing=timer \
 --enable-accelerator=hip \
 --enable-gen-simd-width=64 \
 --enable-simd=GPU \
---disable-accelerator-cshift \
---with-gmp=$OLCF_GMP_ROOT \
+--enable-accelerator-cshift \
+--with-gmp=$GMP \
+--with-mpfr=$MPFR \
 --with-fftw=$FFTW_DIR/.. \
---with-mpfr=/opt/cray/pe/gcc/mpfr/3.1.4/ \
 --disable-fermion-reps \
 --disable-gparity \
 CXX=hipcc MPICXX=mpicxx \
-CXXFLAGS="-fPIC -I{$ROCM_PATH}/include/ -std=c++14 -I${MPICH_DIR}/include -L/lib64 --amdgpu-target=gfx90a" \
- LDFLAGS="-L/lib64 -L/opt/rocm-5.2.0/lib/ -L${MPICH_DIR}/lib -lmpi -L${CRAY_MPICH_ROOTDIR}/gtl/lib -lmpi_gtl_hsa -lamdhip64 --amdgpu-target=gfx90a "
+  CXXFLAGS="-fPIC --offload-arch=gfx90a -I/opt/rocm/include/ -std=c++14 -I/opt/cray/pe/mpich/8.1.23/ofi/gnu/9.1/include" \
+  LDFLAGS="-L/opt/cray/pe/mpich/8.1.23/ofi/gnu/9.1/lib -lmpi -L/opt/cray/pe/mpich/8.1.23/gtl/lib -lmpi_gtl_hsa -lamdhip64 -fopenmp" 
 
 
-#--enable-simd=GPU-RRII \
-
 

systems/Lumi/sourceme.sh

@@ -1 +1,5 @@
-module load CrayEnv LUMI/22.12 partition/G  cray-fftw/3.3.10.1
+source ~/spack/share/spack/setup-env.sh
+module load CrayEnv LUMI/22.12 partition/G  cray-fftw/3.3.10.1 rocm
+spack load c-lime
+spack load gmp
+spack load mpfr

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/Sunspot/benchmarks/bench.pbs

@@ -0,0 +1,46 @@
+#!/bin/bash
+
+#PBS -l select=1:system=sunspot,place=scatter
+#PBS -A LatticeQCD_aesp_CNDA
+#PBS -l walltime=01:00:00
+#PBS -N dwf
+#PBS -k doe
+
+HDIR=/home/paboyle/
+module use /soft/testing/modulefiles/
+module load intel-UMD23.05.25593.11/23.05.25593.11
+module load tools/pti-gpu  
+export LD_LIBRARY_PATH=$HDIR/tools/lib64:$LD_LIBRARY_PATH
+export PATH=$HDIR/tools/bin:$PATH
+
+export TZ='/usr/share/zoneinfo/US/Central'
+export OMP_PROC_BIND=spread
+export OMP_NUM_THREADS=3
+unset OMP_PLACES
+
+cd $PBS_O_WORKDIR
+
+qsub jobscript.pbs
+
+echo Jobid: $PBS_JOBID
+echo Running on host `hostname`
+echo Running on nodes `cat $PBS_NODEFILE`
+
+echo NODES
+cat $PBS_NODEFILE
+NNODES=`wc -l < $PBS_NODEFILE`
+NRANKS=12         # Number of MPI ranks per node
+NDEPTH=4          # Number of hardware threads per rank, spacing between MPI ranks on a node
+NTHREADS=$OMP_NUM_THREADS # Number of OMP threads per rank, given to OMP_NUM_THREADS
+
+NTOTRANKS=$(( NNODES * NRANKS ))
+
+echo "NUM_NODES=${NNODES}  TOTAL_RANKS=${NTOTRANKS}  RANKS_PER_NODE=${NRANKS}  THREADS_PER_RANK=${OMP_NUM_THREADS}"
+echo "OMP_PROC_BIND=$OMP_PROC_BIND OMP_PLACES=$OMP_PLACES"
+
+    
+CMD="mpiexec -np ${NTOTRANKS} -ppn ${NRANKS} -d ${NDEPTH} --cpu-bind=depth -envall \
+	     ./gpu_tile_compact.sh \
+	./Benchmark_dwf_fp32 --mpi 1.1.2.6 --grid 16.32.64.192 --comms-overlap \
+	--shm-mpi 0 --shm 2048 --device-mem 32000 --accelerator-threads 32"
+

systems/Sunspot/benchmarks/gpu_tile_compact.sh

@@ -0,0 +1,52 @@
+#!/bin/bash
+
+display_help() {
+  echo " Will map gpu tile to rank in compact and then round-robin fashion"
+  echo " Usage (only work for one node of ATS/PVC):"
+  echo "   mpiexec --np N gpu_tile_compact.sh ./a.out"
+  echo
+  echo " Example 3 GPU of 2 Tiles with 7 Ranks:"
+  echo "   0 Rank 0.0"
+  echo "   1 Rank 0.1"
+  echo "   2 Rank 1.0"
+  echo "   3 Rank 1.1"
+  echo "   4 Rank 2.0"
+  echo "   5 Rank 2.1"
+  echo "   6 Rank 0.0"
+  echo
+  echo " Hacked together by apl@anl.gov, please contact if bug found"
+  exit 1
+}
+
+#This give the exact GPU count i915 knows about and I use udev to only enumerate the devices with physical presence.
+#works? num_gpu=$(/usr/bin/udevadm info /sys/module/i915/drivers/pci\:i915/* |& grep -v Unknown | grep -c "P: /devices")
+num_gpu=6
+num_tile=2
+
+if [ "$#" -eq 0 ] || [ "$1" == "--help" ] || [ "$1" == "-h" ] || [ "$num_gpu" = 0 ]; then
+  display_help
+fi
+
+gpu_id=$(( (PALS_LOCAL_RANKID / num_tile ) % num_gpu ))
+tile_id=$((PALS_LOCAL_RANKID % num_tile))
+
+unset EnableWalkerPartition
+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
+#export SYCL_PI_LEVEL_ZERO_USM_RESIDENT=1
+
+echo "rank $PALS_RANKID ; local rank $PALS_LOCAL_RANKID ; ZE_AFFINITY_MASK=$ZE_AFFINITY_MASK"
+
+if [ $PALS_LOCAL_RANKID = 0 ]
+then
+    onetrace --chrome-device-timeline "$@"
+#    "$@"
+else
+"$@"
+fi

systems/Sunspot/config-command

@@ -0,0 +1,16 @@
+TOOLS=$HOME/tools
+../../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 -lapmidg -L$TOOLS/lib64/" \
+	CXXFLAGS="-fiopenmp -fsycl-unnamed-lambda -fsycl -I$INSTALL/include -Wno-tautological-compare -I$HOME/ -I$TOOLS/include"
+

systems/mac-arm/config-command-mpi

@@ -1,4 +1,4 @@
 BREW=/opt/local/
-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 
+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/core/Test_fft_pf.cc

@@ -0,0 +1,307 @@
+    /*************************************************************************************
+
+    grid` physics library, www.github.com/paboyle/Grid 
+
+    Source file: ./tests/Test_cshift.cc
+
+    Copyright (C) 2015
+
+Author: Azusa Yamaguchi <ayamaguc@staffmail.ed.ac.uk>
+Author: Peter Boyle <paboyle@ph.ed.ac.uk>
+
+    This program is free software; you can redistribute it and/or modify
+    it under the terms of the GNU General Public License as published by
+    the Free Software Foundation; either version 2 of the License, or
+    (at your option) any later version.
+
+    This program is distributed in the hope that it will be useful,
+    but WITHOUT ANY WARRANTY; without even the implied warranty of
+    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+    GNU General Public License for more details.
+
+    You should have received a copy of the GNU General Public License along
+    with this program; if not, write to the Free Software Foundation, Inc.,
+    51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA.
+
+    See the full license in the file "LICENSE" in the top level distribution directory
+    *************************************************************************************/
+    /*  END LEGAL */
+#include <Grid/Grid.h>
+
+using namespace Grid;
+
+int main (int argc, char ** argv)
+{
+  Grid_init(&argc,&argv);
+
+  int threads = GridThread::GetThreads();
+  std::cout<<GridLogMessage << "Grid is setup to use "<<threads<<" threads"<<std::endl;
+
+  Coordinate latt_size   = GridDefaultLatt();
+  Coordinate simd_layout( { vComplexD::Nsimd(),1,1,1});
+  Coordinate mpi_layout  = GridDefaultMpi();
+
+  int vol = 1;
+  for(int d=0;d<latt_size.size();d++){
+    vol = vol * latt_size[d];
+  }
+  GridCartesian         GRID(latt_size,simd_layout,mpi_layout);
+  GridRedBlackCartesian RBGRID(&GRID);
+
+  ComplexD ci(0.0,1.0);
+
+  std::vector<int> seeds({1,2,3,4});
+  GridSerialRNG          sRNG;  sRNG.SeedFixedIntegers(seeds); // naughty seeding
+  GridParallelRNG          pRNG(&GRID);
+  pRNG.SeedFixedIntegers(seeds);
+
+  LatticeGaugeFieldD Umu(&GRID);
+
+  SU<Nc>::ColdConfiguration(pRNG,Umu); // Unit gauge
+
+  ////////////////////////////////////////////////////
+  // PF prop
+  ////////////////////////////////////////////////////
+  LatticeFermionD    src(&GRID);
+
+  gaussian(pRNG,src);
+#if 1
+    Coordinate point(4,0);
+    src=Zero();
+    SpinColourVectorD ferm; gaussian(sRNG,ferm);
+    pokeSite(ferm,src,point);
+#endif
+  
+  {
+    std::cout<<"****************************************"<<std::endl;
+    std::cout << "Testing PartialFraction Hw kernel Mom space 4d propagator \n";
+    std::cout<<"****************************************"<<std::endl;
+
+    //    LatticeFermionD    src(&GRID); gaussian(pRNG,src);
+    LatticeFermionD    tmp(&GRID);
+    LatticeFermionD    ref(&GRID);
+    LatticeFermionD    diff(&GRID);
+
+    const int Ls=48+1;
+    GridCartesian         * FGrid   = SpaceTimeGrid::makeFiveDimGrid(Ls,&GRID);
+    GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,&GRID);
+
+    RealD mass=0.1;
+    RealD M5  =0.8;
+    OverlapWilsonPartialFractionZolotarevFermionD Dov(Umu,*FGrid,*FrbGrid,GRID,RBGRID,mass,M5,0.001,8.0);
+
+    // Momentum space prop
+    std::cout << " Solving by FFT and Feynman rules" <<std::endl;
+    bool fiveD = false; //calculate 4d free propagator
+
+    std::cout << " Free propagator " <<std::endl;
+    Dov.FreePropagator(src,ref,mass) ;
+    std::cout << " Free propagator norm "<< norm2(ref) <<std::endl;
+
+    Gamma G5(Gamma::Algebra::Gamma5);
+
+    LatticeFermionD    src5(FGrid); src5=Zero();
+    LatticeFermionD    tmp5(FGrid); 
+    LatticeFermionD    result5(FGrid); result5=Zero();
+    LatticeFermionD    result4(&GRID); 
+    const int sdir=0;
+
+    ////////////////////////////////////////////////////////////////////////
+    // Import
+    ////////////////////////////////////////////////////////////////////////
+    std::cout << " Free propagator Import "<< norm2(src) <<std::endl;
+    Dov.ImportPhysicalFermionSource  (src,src5);
+    std::cout << " Free propagator Imported "<< norm2(src5) <<std::endl;
+    
+    ////////////////////////////////////////////////////////////////////////
+    // Conjugate gradient on normal equations system
+    ////////////////////////////////////////////////////////////////////////
+    std::cout << " Solving by Conjugate Gradient (CGNE)" <<std::endl;
+    Dov.Mdag(src5,tmp5);
+    src5=tmp5;
+    MdagMLinearOperator<OverlapWilsonPartialFractionZolotarevFermionD,LatticeFermionD> HermOp(Dov);
+    ConjugateGradient<LatticeFermionD> CG(1.0e-8,10000);
+    CG(HermOp,src5,result5);
+    ////////////////////////////////////////////////////////////////////////
+    // Domain wall physical field propagator
+    ////////////////////////////////////////////////////////////////////////
+    Dov.ExportPhysicalFermionSolution(result5,result4);
+
+    // From DWF4d.pdf :
+    //
+    // Dov_pf = 2/(1-m) D_cayley_ovlap  [ Page 43 ]
+    // Dinv_cayley_ovlap = 2/(1-m) Dinv_pf 
+    // Dinv_cayley_surface =1/(1-m) ( Dinv_cayley_ovlap - 1 ) =>  2/(1-m)^2 Dinv_pf - 1/(1-m) * src   [ Eq.2.67 ]
+
+    RealD scale = 2.0/(1.0-mass)/(1.0-mass);
+    result4 = result4 * scale;
+    result4 = result4 - src*(1.0/(1.0-mass)); // Subtract contact term
+    DumpSliceNorm("Src",src);
+    DumpSliceNorm("Grid",result4);
+    DumpSliceNorm("Fourier",ref);
+
+    std::cout << "Dov result4 "<<norm2(result4)<<std::endl;
+    std::cout << "Dov ref     "<<norm2(ref)<<std::endl;
+
+    diff = result4- ref;
+    DumpSliceNorm("diff ",diff);
+    
+  }
+  
+  ////////////////////////////////////////////////////
+  // Dwf prop
+  ////////////////////////////////////////////////////
+  {
+    std::cout<<"****************************************"<<std::endl;
+    std::cout << "Testing Dov(Hw) Mom space 4d propagator \n";
+    std::cout<<"****************************************"<<std::endl;
+
+    LatticeFermionD    tmp(&GRID);
+    LatticeFermionD    ref(&GRID);
+    LatticeFermionD    diff(&GRID);
+
+    const int Ls=48;
+    GridCartesian         * FGrid   = SpaceTimeGrid::makeFiveDimGrid(Ls,&GRID);
+    GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,&GRID);
+
+    RealD mass=0.1;
+    RealD M5  =0.8;
+
+    OverlapWilsonCayleyTanhFermionD Dov(Umu,*FGrid,*FrbGrid,GRID,RBGRID,mass,M5,1.0);
+
+    // Momentum space prop
+    std::cout << " Solving by FFT and Feynman rules" <<std::endl;
+    Dov.FreePropagator(src,ref,mass) ;
+
+    Gamma G5(Gamma::Algebra::Gamma5);
+
+    LatticeFermionD    src5(FGrid); src5=Zero();
+    LatticeFermionD    tmp5(FGrid); 
+    LatticeFermionD    result5(FGrid); result5=Zero();
+    LatticeFermionD    result4(&GRID); 
+    const int sdir=0;
+
+    ////////////////////////////////////////////////////////////////////////
+    // Domain wall physical field source; need D_minus
+    ////////////////////////////////////////////////////////////////////////
+    /*
+	chi_5[0]   = chiralProjectPlus(chi);
+	chi_5[Ls-1]= chiralProjectMinus(chi);
+    */      
+    tmp =   (src + G5*src)*0.5;      InsertSlice(tmp,src5,   0,sdir);
+    tmp =   (src - G5*src)*0.5;      InsertSlice(tmp,src5,Ls-1,sdir);
+    
+    ////////////////////////////////////////////////////////////////////////
+    // Conjugate gradient on normal equations system
+    ////////////////////////////////////////////////////////////////////////
+    std::cout << " Solving by Conjugate Gradient (CGNE)" <<std::endl;
+    Dov.Dminus(src5,tmp5);
+    src5=tmp5;
+    Dov.Mdag(src5,tmp5);
+    src5=tmp5;
+    MdagMLinearOperator<OverlapWilsonCayleyTanhFermionD,LatticeFermionD> HermOp(Dov);
+    ConjugateGradient<LatticeFermionD> CG(1.0e-16,10000);
+    CG(HermOp,src5,result5);
+    
+    ////////////////////////////////////////////////////////////////////////
+    // Domain wall physical field propagator
+    ////////////////////////////////////////////////////////////////////////
+    /*
+      psi  = chiralProjectMinus(psi_5[0]);
+      psi += chiralProjectPlus(psi_5[Ls-1]);
+    */
+    ExtractSlice(tmp,result5,0   ,sdir);   result4 =         (tmp-G5*tmp)*0.5;
+    ExtractSlice(tmp,result5,Ls-1,sdir);   result4 = result4+(tmp+G5*tmp)*0.5;
+    
+    std::cout << " Taking difference" <<std::endl;
+    std::cout << "Dov result4 "<<norm2(result4)<<std::endl;
+    std::cout << "Dov ref     "<<norm2(ref)<<std::endl;
+    DumpSliceNorm("Grid",result4);
+    DumpSliceNorm("Fourier",ref);
+    diff = ref - result4;
+    std::cout << "result - ref     "<<norm2(diff)<<std::endl;
+    
+    DumpSliceNorm("diff",diff);
+
+  }
+
+  
+  {
+    std::cout<<"****************************************"<<std::endl;
+    std::cout << "Testing PartialFraction Hw kernel Mom space 4d propagator with q\n";
+    std::cout<<"****************************************"<<std::endl;
+
+    //    LatticeFermionD    src(&GRID); gaussian(pRNG,src);
+    LatticeFermionD    tmp(&GRID);
+    LatticeFermionD    ref(&GRID);
+    LatticeFermionD    diff(&GRID);
+
+    const int Ls=48+1;
+    GridCartesian         * FGrid   = SpaceTimeGrid::makeFiveDimGrid(Ls,&GRID);
+    GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,&GRID);
+
+    RealD mass=0.1;
+    RealD M5  =0.8;
+    OverlapWilsonPartialFractionZolotarevFermionD Dov(Umu,*FGrid,*FrbGrid,GRID,RBGRID,mass,M5,0.001,8.0);
+
+    // Momentum space prop
+    std::cout << " Solving by FFT and Feynman rules" <<std::endl;
+    bool fiveD = false; //calculate 4d free propagator
+
+    std::cout << " Free propagator " <<std::endl;
+    Dov.FreePropagator(src,ref,mass) ;
+    std::cout << " Free propagator norm "<< norm2(ref) <<std::endl;
+
+    Gamma G5(Gamma::Algebra::Gamma5);
+
+    LatticeFermionD    src5(FGrid); src5=Zero();
+    LatticeFermionD    tmp5(FGrid); 
+    LatticeFermionD    result5(FGrid); result5=Zero();
+    LatticeFermionD    result4(&GRID); 
+    const int sdir=0;
+
+    ////////////////////////////////////////////////////////////////////////
+    // Import
+    ////////////////////////////////////////////////////////////////////////
+    std::cout << " Free propagator Import "<< norm2(src) <<std::endl;
+    Dov.ImportPhysicalFermionSource  (src,src5);
+    std::cout << " Free propagator Imported "<< norm2(src5) <<std::endl;
+    
+    ////////////////////////////////////////////////////////////////////////
+    // Conjugate gradient on normal equations system
+    ////////////////////////////////////////////////////////////////////////
+    std::cout << " Solving by Conjugate Gradient (CGNE)" <<std::endl;
+    Dov.Mdag(src5,tmp5);
+    src5=tmp5;
+    MdagMLinearOperator<OverlapWilsonPartialFractionZolotarevFermionD,LatticeFermionD> HermOp(Dov);
+    ConjugateGradient<LatticeFermionD> CG(1.0e-8,10000);
+    CG(HermOp,src5,result5);
+    ////////////////////////////////////////////////////////////////////////
+    // Domain wall physical field propagator
+    ////////////////////////////////////////////////////////////////////////
+    Dov.ExportPhysicalFermionSolution(result5,result4);
+
+    // From DWF4d.pdf :
+    //
+    // Dov_pf = 2/(1-m) D_cayley_ovlap  [ Page 43 ]
+    // Dinv_cayley_ovlap = 2/(1-m) Dinv_pf 
+    // Dinv_cayley_surface =1/(1-m) ( Dinv_cayley_ovlap - 1 ) =>  2/(1-m)^2 Dinv_pf - 1/(1-m) * src   [ Eq.2.67 ]
+
+    RealD scale = 2.0/(1.0-mass)/(1.0-mass);
+    result4 = result4 * scale;
+    result4 = result4 - src*(1.0/(1.0-mass)); // Subtract contact term
+    DumpSliceNorm("Src",src);
+    DumpSliceNorm("Grid",result4);
+    DumpSliceNorm("Fourier",ref);
+
+    std::cout << "Dov result4 "<<norm2(result4)<<std::endl;
+    std::cout << "Dov ref     "<<norm2(ref)<<std::endl;
+
+    diff = result4- ref;
+    DumpSliceNorm("diff ",diff);
+    
+  }
+
+  
+  Grid_finalize();
+}

tests/forces/Test_fthmc.cc

@@ -32,9 +32,12 @@ Author: Peter Boyle <pboyle@bnl.gov>
 using namespace std;
 using namespace Grid;
 
+typedef MobiusFermionD FermionAction;
+typedef WilsonImplD FimplD;
+typedef WilsonImplD FermionImplPolicy;
 
 template<class Gimpl>
-void ForceTest(Action<LatticeGaugeField> &action,SmearedConfigurationMasked<PeriodicGimplR> & smU,MomentumFilterBase<LatticeGaugeField> &Filter)
+void ForceTest(Action<LatticeGaugeField> &action,ConfigurationBase<LatticeGaugeField> & smU,MomentumFilterBase<LatticeGaugeField> &Filter)
 {
   LatticeGaugeField U = smU.get_U(false); // unsmeared config
   GridBase *UGrid = U.Grid();
@@ -51,14 +54,14 @@ void ForceTest(Action<LatticeGaugeField> &action,SmearedConfigurationMasked<Peri
   std::cout << GridLogMessage << " Force test for "<<action.action_name()<<std::endl;
   std::cout << GridLogMessage << "*********************************************************"<<std::endl;
   
-  RealD eps=0.005;
+  RealD eps=0.01;
 
   std::cout << GridLogMessage << "+++++++++++++++++++++++++++++++++++++++++++++++++++++++++"<<std::endl;
   std::cout << GridLogMessage << " Refresh "<<action.action_name()<<std::endl;
   std::cout << GridLogMessage << "+++++++++++++++++++++++++++++++++++++++++++++++++++++++++"<<std::endl;
   
   Gimpl::generate_momenta(P,sRNG,RNG4);
-  Filter.applyFilter(P);
+  //  Filter.applyFilter(P);
 
   action.refresh(smU,sRNG,RNG4);
 
@@ -76,7 +79,7 @@ void ForceTest(Action<LatticeGaugeField> &action,SmearedConfigurationMasked<Peri
   std::cout << GridLogMessage << "+++++++++++++++++++++++++++++++++++++++++++++++++++++++++"<<std::endl;
   action.deriv(smU,UdSdU);
   UdSdU = Ta(UdSdU);
-  Filter.applyFilter(UdSdU);
+  //  Filter.applyFilter(UdSdU);
 
   DumpSliceNorm("Force",UdSdU,Nd-1);
   
@@ -94,7 +97,7 @@ void ForceTest(Action<LatticeGaugeField> &action,SmearedConfigurationMasked<Peri
   for(int mu=0;mu<Nd;mu++){
     auto UdSdUmu = PeekIndex<LorentzIndex>(UdSdU,mu);
     Pmu= PeekIndex<LorentzIndex>(P,mu);
-    dS = dS - trace(Pmu*UdSdUmu)*eps*2.0*2.0;
+    dS = dS - trace(Pmu*UdSdUmu)*eps*2.0*HMC_MOMENTUM_DENOMINATOR;
   }
   ComplexD dSpred    = sum(dS);
   RealD diff =  S2-S1-dSpred.real();
@@ -125,7 +128,10 @@ int main (int argc, char ** argv)
   const int Ls=12;
   const int Nt = latt_size[3];
   GridCartesian         * UGrid   = SpaceTimeGrid::makeFourDimGrid(GridDefaultLatt(), GridDefaultSimd(Nd,vComplex::Nsimd()),GridDefaultMpi());
+  GridRedBlackCartesian * UrbGrid = SpaceTimeGrid::makeFourDimRedBlackGrid(UGrid);
   GridCartesian         * FGrid   = SpaceTimeGrid::makeFiveDimGrid(Ls,UGrid);
+  GridRedBlackCartesian * FrbGrid = SpaceTimeGrid::makeFiveDimRedBlackGrid(Ls,UGrid);
+
   
   ///////////////////// Gauge Field and Gauge Forces ////////////////////////////
   LatticeGaugeField U(UGrid);
@@ -141,17 +147,40 @@ int main (int argc, char ** argv)
 #endif
 
   
-  RealD beta=6.0;
-  WilsonGaugeActionR  PlaqAction(beta);
-  IwasakiGaugeActionR RectAction(beta);
+  WilsonGaugeActionR  PlaqAction(6.0);
+  IwasakiGaugeActionR RectAction(2.13);
+  PlaqAction.is_smeared = true;  
+  RectAction.is_smeared = true;  
 
+  ////////////////////////////////////
+  // Fermion Action
+  ////////////////////////////////////
+  RealD mass=0.01; 
+  RealD pvmass=1.0; 
+  RealD M5=1.8; 
+  RealD b=1.5;
+  RealD c=0.5;
+  
+  // Double versions
+  std::vector<Complex> boundary = {1,1,1,-1};
+  FermionAction::ImplParams Params(boundary);
+  FermionAction DdwfPeriodic(U,*FGrid,*FrbGrid,*UGrid,*UrbGrid,mass,M5,b,c,Params);
+  FermionAction PVPeriodic  (U,*FGrid,*FrbGrid,*UGrid,*UrbGrid,pvmass,M5,b,c,Params);
+
+  double StoppingCondition = 1.0e-8;
+  double MaxCGIterations = 50000;
+  ConjugateGradient<LatticeFermion>  CG(StoppingCondition,MaxCGIterations);
+
+  TwoFlavourRatioPseudoFermionAction<FimplD> Nf2(PVPeriodic, DdwfPeriodic,CG,CG);
+  Nf2.is_smeared = true;  
   
   ////////////////////////////////////////////////
   // Plaquette only FTHMC smearer
   ////////////////////////////////////////////////
   double rho = 0.1;
   Smear_Stout<PeriodicGimplR> Smearer(rho);
-  SmearedConfigurationMasked<PeriodicGimplR> SmartConfig(UGrid,2*Nd,Smearer,true);
+  SmearedConfigurationMasked<PeriodicGimplR> SmartConfig(UGrid,2*Nd,Smearer);
+  SmearedConfiguration<PeriodicGimplR> StoutConfig(UGrid,1,Smearer);
 
   JacobianAction<PeriodicGimplR> Jacobian(&SmartConfig);
   
@@ -159,12 +188,32 @@ int main (int argc, char ** argv)
   // Run some tests
   ////////////////////////////////////////////////
   MomentumFilterNone<LatticeGaugeField> FilterNone;
+
+  std::cout << " *********  FIELD TRANSFORM SMEARING ***** "<<std::endl;
+
   SmartConfig.set_Field(U);
   ForceTest<GimplTypesR>(PlaqAction,SmartConfig,FilterNone);
+
   SmartConfig.set_Field(U);
   ForceTest<GimplTypesR>(RectAction,SmartConfig,FilterNone);
+
   SmartConfig.set_Field(U);
   ForceTest<GimplTypesR>(Jacobian,SmartConfig,FilterNone);
 
+  SmartConfig.set_Field(U);
+  ForceTest<GimplTypesR>(Nf2,SmartConfig,FilterNone);
+
+  std::cout << " *********    STOUT SMEARING ***** "<<std::endl;
+
+  StoutConfig.set_Field(U);
+  ForceTest<GimplTypesR>(PlaqAction,StoutConfig,FilterNone);
+
+  StoutConfig.set_Field(U);
+  ForceTest<GimplTypesR>(RectAction,StoutConfig,FilterNone);
+  
+  StoutConfig.set_Field(U);
+  ForceTest<GimplTypesR>(Nf2,StoutConfig,FilterNone);
+  
+
   Grid_finalize();
 }