Debug code

Unused statements generating warnings removed

.github/ISSUE_TEMPLATE/bug-report.yml

@@ -0,0 +1,54 @@
+name: Bug report
+description: Report a bug.
+title: "<insert title>"
+labels: [bug]
+
+body:
+  - type: markdown
+    attributes:
+      value: >
+        Thank you for taking the time to file a bug report.
+        Please check that the code is pointing to the HEAD of develop
+        or any commit in master which is tagged with a version number.
+
+  - type: textarea
+    attributes:
+      label: "Describe the issue:"
+      description: >
+        Describe the issue and any previous attempt to solve it.
+    validations:
+      required: true
+
+  - type: textarea
+    attributes:
+      label: "Code example:"
+      description: >
+        If relevant, show how to reproduce the issue using a minimal working
+        example.
+      placeholder: |
+        << your code here >>
+      render: shell
+    validations:
+      required: false
+
+  - type: textarea
+    attributes:
+      label: "Target platform:"
+      description: >
+        Give a description of the target platform (CPU, network, compiler).
+        Please give the full CPU part description, using for example
+        `cat /proc/cpuinfo | grep 'model name' | uniq` (Linux)
+        or `sysctl machdep.cpu.brand_string` (macOS) and the full output
+        the `--version` option of your compiler.
+    validations:
+      required: true
+
+  - type: textarea
+    attributes:
+      label: "Configure options:"
+      description: >
+        Please give the exact configure command used and attach
+        `config.log`, `grid.config.summary` and the output of `make V=1`.
+      render: shell
+    validations:
+      required: true

Grid/algorithms/LinearOperator.h

@@ -542,6 +542,7 @@ public:
       (*this)(in[i], out[i]);
     }
   }
+  virtual ~LinearFunction(){};
 };
 
 template<class Field> class IdentityLinearFunction : public LinearFunction<Field> {

Grid/algorithms/iterative/ConjugateGradientMultiShiftCleanup.h

@@ -166,16 +166,16 @@ public:
       rsqf[s] =rsq[s];
       std::cout<<GridLogMessage<<"ConjugateGradientMultiShiftMixedPrecCleanup: shift "<< s <<" target resid "<<rsq[s]<<std::endl;
       //      ps_d[s] = src_d;
-      precisionChangeFast(ps_f[s],src_d);
+      precisionChange(ps_f[s],src_d);
     }
     // r and p for primary
     p_d = src_d; //primary copy --- make this a reference to ps_d to save axpys
     r_d = p_d;
     
     //MdagM+m[0]
-    precisionChangeFast(p_f,p_d);
+    precisionChange(p_f,p_d);
     Linop_f.HermOpAndNorm(p_f,mmp_f,d,qq); // mmp = MdagM p        d=real(dot(p, mmp)),  qq=norm2(mmp)
-    precisionChangeFast(tmp_d,mmp_f);
+    precisionChange(tmp_d,mmp_f);
     Linop_d.HermOpAndNorm(p_d,mmp_d,d,qq); // mmp = MdagM p        d=real(dot(p, mmp)),  qq=norm2(mmp)
     tmp_d = tmp_d - mmp_d;
     std::cout << " Testing operators match "<<norm2(mmp_d)<<" f "<<norm2(mmp_f)<<" diff "<< norm2(tmp_d)<<std::endl;
@@ -204,7 +204,7 @@ public:
   
     for(int s=0;s<nshift;s++) {
       axpby(psi_d[s],0.,-bs[s]*alpha[s],src_d,src_d);
-      precisionChangeFast(psi_f[s],psi_d[s]);
+      precisionChange(psi_f[s],psi_d[s]);
     }
   
     ///////////////////////////////////////
@@ -225,7 +225,7 @@ public:
       AXPYTimer.Stop();
 
       PrecChangeTimer.Start();
-      precisionChangeFast(r_f, r_d);
+      precisionChange(r_f, r_d);
       PrecChangeTimer.Stop();
 
       AXPYTimer.Start();
@@ -243,13 +243,13 @@ public:
 
       cp=c;
       PrecChangeTimer.Start();
-      precisionChangeFast(p_f, p_d); //get back single prec search direction for linop
+      precisionChange(p_f, p_d); //get back single prec search direction for linop
       PrecChangeTimer.Stop();
       MatrixTimer.Start();  
       Linop_f.HermOp(p_f,mmp_f);
       MatrixTimer.Stop();  
       PrecChangeTimer.Start();
-      precisionChangeFast(mmp_d, mmp_f); // From Float to Double
+      precisionChange(mmp_d, mmp_f); // From Float to Double
       PrecChangeTimer.Stop();
 
       d=real(innerProduct(p_d,mmp_d));    
@@ -311,7 +311,7 @@ public:
 	SolverTimer.Stop();
 
 	for(int s=0;s<nshift;s++){
-	  precisionChangeFast(psi_d[s],psi_f[s]);
+	  precisionChange(psi_d[s],psi_f[s]);
 	}
 
 	

Grid/algorithms/iterative/ConjugateGradientMultiShiftMixedPrec.h

@@ -211,7 +211,7 @@ public:
     Linop_d.HermOpAndNorm(p_d,mmp_d,d,qq); // mmp = MdagM p        d=real(dot(p, mmp)),  qq=norm2(mmp)
     tmp_d = tmp_d - mmp_d;
     std::cout << " Testing operators match "<<norm2(mmp_d)<<" f "<<norm2(mmp_f)<<" diff "<< norm2(tmp_d)<<std::endl;
-    //    assert(norm2(tmp_d)< 1.0e-4);
+    assert(norm2(tmp_d)< 1.0);
 
     axpy(mmp_d,mass[0],p_d,mmp_d);
     RealD rn = norm2(p_d);

Grid/allocator/MemoryManagerCache.cc

@@ -519,7 +519,6 @@ void MemoryManager::Audit(std::string s)
   uint64_t LruBytes1=0;
   uint64_t LruBytes2=0;
   uint64_t LruCnt=0;
-  uint64_t LockedBytes=0;
   
   std::cout << " Memory Manager::Audit() from "<<s<<std::endl;
   for(auto it=LRU.begin();it!=LRU.end();it++){

Grid/communicator/Communicator_none.cc

@@ -128,7 +128,7 @@ double CartesianCommunicator::StencilSendToRecvFromBegin(std::vector<CommsReques
 							 int recv_from_rank,int dor,
 							 int xbytes,int rbytes, int dir)
 {
-  return 2.0*bytes;
+  return xbytes+rbytes;
 }
 void CartesianCommunicator::StencilSendToRecvFromComplete(std::vector<CommsRequest_t> &waitall,int dir)
 {

Grid/communicator/SharedMemory.cc

@@ -91,6 +91,59 @@ void *SharedMemory::ShmBufferSelf(void)
   //std::cerr << "ShmBufferSelf "<<ShmRank<<" "<<std::hex<< ShmCommBufs[ShmRank] <<std::dec<<std::endl;
   return ShmCommBufs[ShmRank];
 }
+static inline int divides(int a,int b)
+{
+  return ( b == ( (b/a)*a ) );
+}
+void GlobalSharedMemory::GetShmDims(const Coordinate &WorldDims,Coordinate &ShmDims)
+{
+  ////////////////////////////////////////////////////////////////
+  // Allow user to configure through environment variable
+  ////////////////////////////////////////////////////////////////
+  char* str = getenv(("GRID_SHM_DIMS_" + std::to_string(ShmDims.size())).c_str());
+  if ( str ) {
+    std::vector<int> IntShmDims;
+    GridCmdOptionIntVector(std::string(str),IntShmDims);
+    assert(IntShmDims.size() == WorldDims.size());
+    long ShmSize = 1;
+    for (int dim=0;dim<WorldDims.size();dim++) {
+      ShmSize *= (ShmDims[dim] = IntShmDims[dim]);
+      assert(divides(ShmDims[dim],WorldDims[dim]));
+    }
+    assert(ShmSize == WorldShmSize);
+    return;
+  }
+  
+  ////////////////////////////////////////////////////////////////
+  // Powers of 2,3,5 only in prime decomposition for now
+  ////////////////////////////////////////////////////////////////
+  int ndimension = WorldDims.size();
+  ShmDims=Coordinate(ndimension,1);
+
+  std::vector<int> primes({2,3,5});
+
+  int dim = 0;
+  int last_dim = ndimension - 1;
+  int AutoShmSize = 1;
+  while(AutoShmSize != WorldShmSize) {
+    int p;
+    for(p=0;p<primes.size();p++) {
+      int prime=primes[p];
+      if ( divides(prime,WorldDims[dim]/ShmDims[dim])
+        && divides(prime,WorldShmSize/AutoShmSize)  ) {
+  AutoShmSize*=prime;
+  ShmDims[dim]*=prime;
+  last_dim = dim;
+  break;
+      }
+    }
+    if (p == primes.size() && last_dim == dim) {
+      std::cerr << "GlobalSharedMemory::GetShmDims failed" << std::endl;
+      exit(EXIT_FAILURE);
+    }
+    dim=(dim+1) %ndimension;
+  }
+}
 
 NAMESPACE_END(Grid); 
 

Grid/communicator/SharedMemoryMPI.cc

@@ -27,9 +27,10 @@ Author: Christoph Lehner <christoph@lhnr.de>
 *************************************************************************************/
 /*  END LEGAL */
 
+#define header "SharedMemoryMpi: "
+
 #include <Grid/GridCore.h>
 #include <pwd.h>
-#include <syscall.h>
 
 #ifdef GRID_CUDA
 #include <cuda_runtime_api.h>
@@ -39,11 +40,118 @@ Author: Christoph Lehner <christoph@lhnr.de>
 #endif
 #ifdef GRID_SYCL
 #define GRID_SYCL_LEVEL_ZERO_IPC
+#include <syscall.h>
+#define SHM_SOCKETS 
+#endif
+
+#include <sys/socket.h>
+#include <sys/un.h>
+
+NAMESPACE_BEGIN(Grid); 
+
+#ifdef SHM_SOCKETS
+
+/*
+ * Barbaric extra intranode communication route in case we need sockets to pass FDs
+ * Forced by level_zero not being nicely designed
+ */
+static int sock;
+static const char *sock_path_fmt = "/tmp/GridUnixSocket.%d";
+static char sock_path[256];
+class UnixSockets {
+public:
+  static void Open(int rank)
+  {
+    int errnum;
+
+    sock = socket(AF_UNIX, SOCK_DGRAM, 0);  assert(sock>0);
+
+    struct sockaddr_un sa_un = { 0 };
+    sa_un.sun_family = AF_UNIX;
+    snprintf(sa_un.sun_path, sizeof(sa_un.sun_path),sock_path_fmt,rank);
+    unlink(sa_un.sun_path);
+    if (bind(sock, (struct sockaddr *)&sa_un, sizeof(sa_un))) {
+      perror("bind failure");
+      exit(EXIT_FAILURE);
+    }
+  }
+
+  static int RecvFileDescriptor(void)
+  {
+    int n;
+    int fd;
+    char buf[1];
+    struct iovec iov;
+    struct msghdr msg;
+    struct cmsghdr *cmsg;
+    char cms[CMSG_SPACE(sizeof(int))];
+
+    iov.iov_base = buf;
+    iov.iov_len = 1;
+
+    memset(&msg, 0, sizeof msg);
+    msg.msg_name = 0;
+    msg.msg_namelen = 0;
+    msg.msg_iov = &iov;
+    msg.msg_iovlen = 1;
+
+    msg.msg_control = (caddr_t)cms;
+    msg.msg_controllen = sizeof cms;
+
+    if((n=recvmsg(sock, &msg, 0)) < 0) {
+      perror("recvmsg failed");
+      return -1;
+    }
+    if(n == 0){
+      perror("recvmsg returned 0");
+      return -1;
+    }
+    cmsg = CMSG_FIRSTHDR(&msg);
+
+    memmove(&fd, CMSG_DATA(cmsg), sizeof(int));
+
+    return fd;
+  }
+
+  static void SendFileDescriptor(int fildes,int xmit_to_rank)
+  {
+    struct msghdr msg;
+    struct iovec iov;
+    struct cmsghdr *cmsg = NULL;
+    char ctrl[CMSG_SPACE(sizeof(int))];
+    char data = ' ';
+
+    memset(&msg, 0, sizeof(struct msghdr));
+    memset(ctrl, 0, CMSG_SPACE(sizeof(int)));
+    iov.iov_base = &data;
+    iov.iov_len = sizeof(data);
+    
+    sprintf(sock_path,sock_path_fmt,xmit_to_rank);
+    
+    struct sockaddr_un sa_un = { 0 };
+    sa_un.sun_family = AF_UNIX;
+    snprintf(sa_un.sun_path, sizeof(sa_un.sun_path),sock_path_fmt,xmit_to_rank);
+
+    msg.msg_name = (void *)&sa_un;
+    msg.msg_namelen = sizeof(sa_un);
+    msg.msg_iov = &iov;
+    msg.msg_iovlen = 1;
+    msg.msg_controllen =  CMSG_SPACE(sizeof(int));
+    msg.msg_control = ctrl;
+
+    cmsg = CMSG_FIRSTHDR(&msg);
+    cmsg->cmsg_level = SOL_SOCKET;
+    cmsg->cmsg_type = SCM_RIGHTS;
+    cmsg->cmsg_len = CMSG_LEN(sizeof(int));
+
+    *((int *) CMSG_DATA(cmsg)) = fildes;
+
+    sendmsg(sock, &msg, 0);
+  };
+};
 #endif
 
 
-NAMESPACE_BEGIN(Grid); 
-#define header "SharedMemoryMpi: "
 /*Construct from an MPI communicator*/
 void GlobalSharedMemory::Init(Grid_MPI_Comm comm)
 {
@@ -170,59 +278,7 @@ void GlobalSharedMemory::OptimalCommunicator(const Coordinate &processors,Grid_M
   if(nscan==3 && HPEhypercube ) OptimalCommunicatorHypercube(processors,optimal_comm,SHM);
   else                          OptimalCommunicatorSharedMemory(processors,optimal_comm,SHM);
 }
-static inline int divides(int a,int b)
-{
-  return ( b == ( (b/a)*a ) );
-}
-void GlobalSharedMemory::GetShmDims(const Coordinate &WorldDims,Coordinate &ShmDims)
-{
-  ////////////////////////////////////////////////////////////////
-  // Allow user to configure through environment variable
-  ////////////////////////////////////////////////////////////////
-  char* str = getenv(("GRID_SHM_DIMS_" + std::to_string(ShmDims.size())).c_str());
-  if ( str ) {
-    std::vector<int> IntShmDims;
-    GridCmdOptionIntVector(std::string(str),IntShmDims);
-    assert(IntShmDims.size() == WorldDims.size());
-    long ShmSize = 1;
-    for (int dim=0;dim<WorldDims.size();dim++) {
-      ShmSize *= (ShmDims[dim] = IntShmDims[dim]);
-      assert(divides(ShmDims[dim],WorldDims[dim]));
-    }
-    assert(ShmSize == WorldShmSize);
-    return;
-  }
-  
-  ////////////////////////////////////////////////////////////////
-  // Powers of 2,3,5 only in prime decomposition for now
-  ////////////////////////////////////////////////////////////////
-  int ndimension = WorldDims.size();
-  ShmDims=Coordinate(ndimension,1);
 
-  std::vector<int> primes({2,3,5});
-
-  int dim = 0;
-  int last_dim = ndimension - 1;
-  int AutoShmSize = 1;
-  while(AutoShmSize != WorldShmSize) {
-    int p;
-    for(p=0;p<primes.size();p++) {
-      int prime=primes[p];
-      if ( divides(prime,WorldDims[dim]/ShmDims[dim])
-        && divides(prime,WorldShmSize/AutoShmSize)  ) {
-	AutoShmSize*=prime;
-	ShmDims[dim]*=prime;
-	last_dim = dim;
-	break;
-      }
-    }
-    if (p == primes.size() && last_dim == dim) {
-      std::cerr << "GlobalSharedMemory::GetShmDims failed" << std::endl;
-      exit(EXIT_FAILURE);
-    }
-    dim=(dim+1) %ndimension;
-  }
-}
 void GlobalSharedMemory::OptimalCommunicatorHypercube(const Coordinate &processors,Grid_MPI_Comm & optimal_comm,Coordinate &SHM)
 {
   ////////////////////////////////////////////////////////////////
@@ -532,8 +588,13 @@ void GlobalSharedMemory::SharedMemoryAllocate(uint64_t bytes, int flags)
   ///////////////////////////////////////////////////////////////////////////////////////////////////////////
   // Loop over ranks/gpu's on our node
   ///////////////////////////////////////////////////////////////////////////////////////////////////////////
+#ifdef SHM_SOCKETS
+  UnixSockets::Open(WorldShmRank);
+#endif
   for(int r=0;r<WorldShmSize;r++){
 
+    MPI_Barrier(WorldShmComm);
+
 #ifndef GRID_MPI3_SHM_NONE
     //////////////////////////////////////////////////
     // If it is me, pass around the IPC access key
@@ -541,24 +602,32 @@ void GlobalSharedMemory::SharedMemoryAllocate(uint64_t bytes, int flags)
     void * thisBuf = ShmCommBuf;
     if(!Stencil_force_mpi) {
 #ifdef GRID_SYCL_LEVEL_ZERO_IPC
-    typedef struct { int fd; pid_t pid ; } clone_mem_t;
+    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());
       
     ze_ipc_mem_handle_t ihandle;
     clone_mem_t handle;
-
+    
     if ( r==WorldShmRank ) { 
       auto err = zeMemGetIpcHandle(zeContext,ShmCommBuf,&ihandle);
       if ( err != ZE_RESULT_SUCCESS ) {
-	std::cout << "SharedMemoryMPI.cc zeMemGetIpcHandle failed for rank "<<r<<" "<<std::hex<<err<<std::dec<<std::endl;
+	std::cerr << "SharedMemoryMPI.cc zeMemGetIpcHandle failed for rank "<<r<<" "<<std::hex<<err<<std::dec<<std::endl;
 	exit(EXIT_FAILURE);
       } else {
 	std::cout << "SharedMemoryMPI.cc zeMemGetIpcHandle succeeded for rank "<<r<<" "<<std::hex<<err<<std::dec<<std::endl;
       }
       memcpy((void *)&handle.fd,(void *)&ihandle,sizeof(int));
       handle.pid = getpid();
+      memcpy((void *)&handle.ze,(void *)&ihandle,sizeof(ihandle));
+#ifdef SHM_SOCKETS
+      for(int rr=0;rr<WorldShmSize;rr++){
+	if(rr!=r){
+	  UnixSockets::SendFileDescriptor(handle.fd,rr);
+	}
+      }
+#endif
     }
 #endif
 #ifdef GRID_CUDA
@@ -586,6 +655,7 @@ void GlobalSharedMemory::SharedMemoryAllocate(uint64_t bytes, int flags)
     // Share this IPC handle across the Shm Comm
     //////////////////////////////////////////////////
     { 
+      MPI_Barrier(WorldShmComm);
       int ierr=MPI_Bcast(&handle,
 			 sizeof(handle),
 			 MPI_BYTE,
@@ -601,6 +671,10 @@ void GlobalSharedMemory::SharedMemoryAllocate(uint64_t bytes, int flags)
 #ifdef GRID_SYCL_LEVEL_ZERO_IPC
     if ( r!=WorldShmRank ) {
       thisBuf = nullptr;
+      int myfd;
+#ifdef SHM_SOCKETS
+      myfd=UnixSockets::RecvFileDescriptor();
+#else
       std::cout<<"mapping seeking remote pid/fd "
 	       <<handle.pid<<"/"
 	       <<handle.fd<<std::endl;
@@ -608,16 +682,22 @@ void GlobalSharedMemory::SharedMemoryAllocate(uint64_t bytes, int flags)
       int pidfd = syscall(SYS_pidfd_open,handle.pid,0);
       std::cout<<"Using IpcHandle pidfd "<<pidfd<<"\n";
       //      int myfd  = syscall(SYS_pidfd_getfd,pidfd,handle.fd,0);
-      int myfd  = syscall(438,pidfd,handle.fd,0);
-
-      std::cout<<"Using IpcHandle myfd "<<myfd<<"\n";
-      
+      myfd  = syscall(438,pidfd,handle.fd,0);
+      int err_t = errno;
+      if (myfd < 0) {
+        fprintf(stderr,"pidfd_getfd returned %d errno was %d\n", myfd,err_t); fflush(stderr);
+	perror("pidfd_getfd failed ");
+	assert(0);
+      }
+#endif
+      std::cout<<"Using IpcHandle mapped remote pid "<<handle.pid <<" FD "<<handle.fd <<" to myfd "<<myfd<<"\n";
+      memcpy((void *)&ihandle,(void *)&handle.ze,sizeof(ihandle));
       memcpy((void *)&ihandle,(void *)&myfd,sizeof(int));
 
       auto err = zeMemOpenIpcHandle(zeContext,zeDevice,ihandle,0,&thisBuf);
       if ( err != ZE_RESULT_SUCCESS ) {
-	std::cout << "SharedMemoryMPI.cc "<<zeContext<<" "<<zeDevice<<std::endl;
-	std::cout << "SharedMemoryMPI.cc zeMemOpenIpcHandle failed for rank "<<r<<" "<<std::hex<<err<<std::dec<<std::endl; 
+	std::cerr << "SharedMemoryMPI.cc "<<zeContext<<" "<<zeDevice<<std::endl;
+	std::cerr << "SharedMemoryMPI.cc zeMemOpenIpcHandle failed for rank "<<r<<" "<<std::hex<<err<<std::dec<<std::endl; 
 	exit(EXIT_FAILURE);
       } else {
 	std::cout << "SharedMemoryMPI.cc zeMemOpenIpcHandle succeeded for rank "<<r<<std::endl;
@@ -652,6 +732,7 @@ void GlobalSharedMemory::SharedMemoryAllocate(uint64_t bytes, int flags)
 #else
     WorldShmCommBufs[r] = ShmCommBuf;
 #endif
+    MPI_Barrier(WorldShmComm);
   }
 
   _ShmAllocBytes=bytes;

Grid/qcd/action/fermion/WilsonCompressor.h

@@ -507,6 +507,7 @@ public:
     }
     this->face_table_computed=1;
     assert(this->u_comm_offset==this->_unified_buffer_size);
+    accelerator_barrier();
   }
 
 };

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

@@ -196,7 +196,6 @@ void WilsonFermion5D<Impl>::DhopDir(const FermionField &in, FermionField &out,in
   
   uint64_t Nsite = Umu.Grid()->oSites();
   Kernels::DhopDirKernel(Stencil,Umu,Stencil.CommBuf(),Ls,Nsite,in,out,dirdisp,gamma);
-
 };
 template<class Impl>
 void WilsonFermion5D<Impl>::DhopDirAll(const FermionField &in, std::vector<FermionField> &out)
@@ -247,10 +246,14 @@ void WilsonFermion5D<Impl>::DerivInternal(StencilImpl & st,
 
     Kernels::DhopDirKernel(st, U, st.CommBuf(), Ls, Usites, B, Btilde, mu,gamma);
 
+    std::cout << " InsertForce Btilde "<< norm2(Btilde)<<std::endl;
+
     ////////////////////////////
     // spin trace outer product
     ////////////////////////////
     Impl::InsertForce5D(mat, Btilde, Atilde, mu);
+
+    std::cout << " InsertForce "<< norm2(mat)<<std::endl;
   }
 }
 
@@ -332,8 +335,7 @@ void WilsonFermion5D<Impl>::DhopInternalOverlappedComms(StencilImpl & st, Lebesg
   /////////////////////////////
   {
     GRID_TRACE("Gather");
-    st.HaloExchangeOptGather(in,compressor);
-    accelerator_barrier();
+    st.HaloExchangeOptGather(in,compressor); // Put the barrier in the routine
   }
   
   std::vector<std::vector<CommsRequest_t> > requests;

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

@@ -428,9 +428,10 @@ void WilsonKernels<Impl>::DhopDirKernel( StencilImpl &st, DoubledGaugeField &U,S
   auto ptr = &st.surface_list[0];					\
   accelerator_forNB( ss, sz, Simd::Nsimd(), {				\
       int sF = ptr[ss];							\
-      int sU = ss/Ls;							\
+      int sU = sF/Ls;							\
       WilsonKernels<Impl>::A(st_v,U_v,buf,sF,sU,in_v,out_v);		\
-    });									
+    });									\
+  accelerator_barrier();
 
 #define ASM_CALL(A)							\
   thread_for( sss, Nsite, {						\
@@ -474,9 +475,10 @@ void WilsonKernels<Impl>::DhopKernel(int Opt,StencilImpl &st,  DoubledGaugeField
      if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL(AsmDhopSiteInt);    return;}
 #endif
    } else if( exterior ) {
+     // dependent on result of merge
      acceleratorFenceComputeStream();
-     if (Opt == WilsonKernelsStatic::OptGeneric    ) { KERNEL_CALL(GenericDhopSiteExt); return;}
-     if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALL(HandDhopSiteExt);    return;}
+     if (Opt == WilsonKernelsStatic::OptGeneric    ) { KERNEL_CALL_EXT(GenericDhopSiteExt); return;}
+     if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALL_EXT(HandDhopSiteExt);    return;}
 #ifndef GRID_CUDA
      if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL(AsmDhopSiteExt);    return;}
 #endif
@@ -506,9 +508,10 @@ void WilsonKernels<Impl>::DhopKernel(int Opt,StencilImpl &st,  DoubledGaugeField
      if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL(AsmDhopSiteDagInt);     return;}
 #endif
    } else if( exterior ) {
+     // Dependent on result of merge
      acceleratorFenceComputeStream();
-     if (Opt == WilsonKernelsStatic::OptGeneric    ) { KERNEL_CALL(GenericDhopSiteDagExt); return;}
-     if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALL(HandDhopSiteDagExt);    return;}
+     if (Opt == WilsonKernelsStatic::OptGeneric    ) { KERNEL_CALL_EXT(GenericDhopSiteDagExt); return;}
+     if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALL_EXT(HandDhopSiteDagExt);    return;}
 #ifndef GRID_CUDA
      if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL(AsmDhopSiteDagExt);     return;}
 #endif

Grid/qcd/action/pseudofermion/EvenOddSchurDifferentiable.h

@@ -119,13 +119,19 @@ public:
     //  X^dag Der_oe MeeInv Meo Y
     // Use Mooee as nontrivial but gauge field indept
     this->_Mat.MeooeDag   (V,tmp1);      // odd->even -- implicit -0.5 factor to be applied
+    std::cout << " tmp 1" << norm2(tmp1)<<std::endl;
     this->_Mat.MooeeInvDag(tmp1,tmp2);   // even->even 
+    std::cout << " tmp 1" << norm2(tmp2)<<std::endl;
     this->_Mat.MoeDeriv(ForceO,U,tmp2,DaggerYes);
+    std::cout << " ForceO " << norm2(ForceO)<<std::endl;
           
     //  Accumulate X^dag M_oe MeeInv Der_eo Y
     this->_Mat.Meooe   (U,tmp1);    // even->odd -- implicit -0.5 factor to be applied
+    std::cout << " tmp 1" << norm2(tmp1)<<std::endl;
     this->_Mat.MooeeInv(tmp1,tmp2); // even->even 
+    std::cout << " tmp 2" << norm2(tmp2)<<std::endl;
     this->_Mat.MeoDeriv(ForceE,tmp2,V,DaggerYes);
+    std::cout << " ForceE " << norm2(ForceE)<<std::endl;
 
     assert(ForceE.Checkerboard()==Even);
     assert(ForceO.Checkerboard()==Odd);

Grid/qcd/action/pseudofermion/GeneralEvenOddRationalRatioMixedPrec.h

@@ -38,91 +38,73 @@ NAMESPACE_BEGIN(Grid);
     // cf. GeneralEvenOddRational.h for details
     /////////////////////////////////////////////////////////////////////////////////////////////////////////////
       
-    template<class ImplD, class ImplF, class ImplD2>
+    template<class ImplD, class ImplF>
     class GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction : public GeneralEvenOddRatioRationalPseudoFermionAction<ImplD> {
     private:
-      typedef typename ImplD2::FermionField FermionFieldD2;
       typedef typename ImplD::FermionField FermionFieldD;
       typedef typename ImplF::FermionField FermionFieldF;
 
       FermionOperator<ImplD> & NumOpD;
       FermionOperator<ImplD> & DenOpD;
 
-      FermionOperator<ImplD2> & NumOpD2;
-      FermionOperator<ImplD2> & DenOpD2;
-     
       FermionOperator<ImplF> & NumOpF;
       FermionOperator<ImplF> & DenOpF;
 
       Integer ReliableUpdateFreq;
     protected:
 
+      //Action evaluation
       //Allow derived classes to override the multishift CG
       virtual void multiShiftInverse(bool numerator, const MultiShiftFunction &approx, const Integer MaxIter, const FermionFieldD &in, FermionFieldD &out){
-#if 0
+#if 1
 	SchurDifferentiableOperator<ImplD> schurOp(numerator ? NumOpD : DenOpD);
 	ConjugateGradientMultiShift<FermionFieldD> msCG(MaxIter, approx);
 	msCG(schurOp,in, out);
 #else
-	SchurDifferentiableOperator<ImplD2> schurOpD2(numerator ? NumOpD2 : DenOpD2);
+	SchurDifferentiableOperator<ImplD> schurOpD(numerator ? NumOpD : DenOpD);
 	SchurDifferentiableOperator<ImplF> schurOpF(numerator ? NumOpF : DenOpF);
-	FermionFieldD2 inD2(NumOpD2.FermionRedBlackGrid());
-	FermionFieldD2 outD2(NumOpD2.FermionRedBlackGrid());
+	FermionFieldD inD(NumOpD.FermionRedBlackGrid());
+	FermionFieldD outD(NumOpD.FermionRedBlackGrid());
 
 	// Action better with higher precision?
-	ConjugateGradientMultiShiftMixedPrec<FermionFieldD2, FermionFieldF> msCG(MaxIter, approx, NumOpF.FermionRedBlackGrid(), schurOpF, ReliableUpdateFreq);
-	precisionChange(inD2,in);
-	std::cout << "msCG single solve "<<norm2(inD2)<<" " <<norm2(in)<<std::endl;
-	msCG(schurOpD2, inD2, outD2);
-	precisionChange(out,outD2);
+	ConjugateGradientMultiShiftMixedPrec<FermionFieldD, FermionFieldF> msCG(MaxIter, approx, NumOpF.FermionRedBlackGrid(), schurOpF, ReliableUpdateFreq);
+	msCG(schurOpD, in, out);
 #endif
       }
+      //Force evaluation
       virtual void multiShiftInverse(bool numerator, const MultiShiftFunction &approx, const Integer MaxIter, const FermionFieldD &in, std::vector<FermionFieldD> &out_elems, FermionFieldD &out){
-	SchurDifferentiableOperator<ImplD2> schurOpD2(numerator ? NumOpD2 : DenOpD2);
-	SchurDifferentiableOperator<ImplF>  schurOpF (numerator ? NumOpF  : DenOpF);
+	SchurDifferentiableOperator<ImplD> schurOpD(numerator ? NumOpD : DenOpD);
+	SchurDifferentiableOperator<ImplF>  schurOpF(numerator ? NumOpF  : DenOpF);
 
-	FermionFieldD2 inD2(NumOpD2.FermionRedBlackGrid());
-	FermionFieldD2 outD2(NumOpD2.FermionRedBlackGrid());
-	std::vector<FermionFieldD2> out_elemsD2(out_elems.size(),NumOpD2.FermionRedBlackGrid());
-	ConjugateGradientMultiShiftMixedPrecCleanup<FermionFieldD2, FermionFieldF> msCG(MaxIter, approx, NumOpF.FermionRedBlackGrid(), schurOpF, ReliableUpdateFreq);
-	precisionChange(inD2,in);
-	std::cout << "msCG in "<<norm2(inD2)<<" " <<norm2(in)<<std::endl;
-	msCG(schurOpD2, inD2, out_elemsD2, outD2);
-	precisionChange(out,outD2);
-	for(int i=0;i<out_elems.size();i++){
-	  precisionChange(out_elems[i],out_elemsD2[i]);
-	}
+	FermionFieldD inD(NumOpD.FermionRedBlackGrid());
+	FermionFieldD outD(NumOpD.FermionRedBlackGrid());
+	std::vector<FermionFieldD> out_elemsD(out_elems.size(),NumOpD.FermionRedBlackGrid());
+	ConjugateGradientMultiShiftMixedPrecCleanup<FermionFieldD, FermionFieldF> msCG(MaxIter, approx, NumOpF.FermionRedBlackGrid(), schurOpF, ReliableUpdateFreq);
+	msCG(schurOpD, in, out_elems, out);
       }
       //Allow derived classes to override the gauge import
       virtual void ImportGauge(const typename ImplD::GaugeField &Ud){
 
 	typename ImplF::GaugeField Uf(NumOpF.GaugeGrid());
-	typename ImplD2::GaugeField Ud2(NumOpD2.GaugeGrid());
 	precisionChange(Uf, Ud);
-	precisionChange(Ud2, Ud);
 
-	std::cout << "Importing "<<norm2(Ud)<<" "<< norm2(Uf)<<" " << norm2(Ud2)<<std::endl;
+	std::cout << "Importing "<<norm2(Ud)<<" "<< norm2(Uf)<<" " <<std::endl;
 	
 	NumOpD.ImportGauge(Ud);
 	DenOpD.ImportGauge(Ud);
 
 	NumOpF.ImportGauge(Uf);
 	DenOpF.ImportGauge(Uf);
-
-	NumOpD2.ImportGauge(Ud2);
-	DenOpD2.ImportGauge(Ud2);
       }
       
     public:
       GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction(FermionOperator<ImplD>  &_NumOpD, FermionOperator<ImplD>  &_DenOpD, 
 							      FermionOperator<ImplF>  &_NumOpF, FermionOperator<ImplF>  &_DenOpF, 
-							      FermionOperator<ImplD2>  &_NumOpD2, FermionOperator<ImplD2>  &_DenOpD2, 
 							      const RationalActionParams & p, Integer _ReliableUpdateFreq
 							      ) : GeneralEvenOddRatioRationalPseudoFermionAction<ImplD>(_NumOpD, _DenOpD, p),
 								  ReliableUpdateFreq(_ReliableUpdateFreq),
 								  NumOpD(_NumOpD), DenOpD(_DenOpD),
-								  NumOpF(_NumOpF), DenOpF(_DenOpF),
-								  NumOpD2(_NumOpD2), DenOpD2(_DenOpD2)
+								  NumOpF(_NumOpF), DenOpF(_DenOpF)
       {}
       
       virtual std::string action_name(){return "GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction";}

Grid/qcd/action/pseudofermion/OneFlavourEvenOddRationalRatio.h

@@ -67,9 +67,9 @@ NAMESPACE_BEGIN(Grid);
       virtual std::string action_name(){return "OneFlavourEvenOddRatioRationalPseudoFermionAction";}      
     };
 
-    template<class Impl,class ImplF,class ImplD2>
+    template<class Impl,class ImplF>
     class OneFlavourEvenOddRatioRationalMixedPrecPseudoFermionAction
-      : public GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<Impl,ImplF,ImplD2> {
+      : public GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<Impl,ImplF> {
     public:
       typedef OneFlavourRationalParams Params;
     private:
@@ -91,11 +91,9 @@ NAMESPACE_BEGIN(Grid);
 								 FermionOperator<Impl>  &_DenOp, 
 								 FermionOperator<ImplF>  &_NumOpF, 
 								 FermionOperator<ImplF>  &_DenOpF, 
-								 FermionOperator<ImplD2>  &_NumOpD2, 
-								 FermionOperator<ImplD2>  &_DenOpD2, 
 								 const Params & p, Integer ReliableUpdateFreq
 							) : 
-	GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<Impl,ImplF,ImplD2>(_NumOp, _DenOp,_NumOpF, _DenOpF,_NumOpD2, _DenOpD2, transcribe(p),ReliableUpdateFreq){}
+	GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<Impl,ImplF>(_NumOp, _DenOp,_NumOpF, _DenOpF, transcribe(p),ReliableUpdateFreq){}
 
       virtual std::string action_name(){return "OneFlavourEvenOddRatioRationalPseudoFermionAction";}      
     };

Grid/qcd/action/pseudofermion/TwoFlavourEvenOddRatio.h

@@ -207,20 +207,27 @@ NAMESPACE_BEGIN(Grid);
         //X = (Mdag M)^-1 V^dag phi
         //Y = (Mdag)^-1 V^dag  phi
         Vpc.MpcDag(PhiOdd,Y);          // Y= Vdag phi
+	std::cout << GridLogMessage <<" Y "<<norm2(Y)<<std::endl;
         X=Zero();
         DerivativeSolver(Mpc,Y,X);     // X= (MdagM)^-1 Vdag phi
+	std::cout << GridLogMessage <<" X "<<norm2(X)<<std::endl;
         Mpc.Mpc(X,Y);                  // Y=  Mdag^-1 Vdag phi
+	std::cout << GridLogMessage <<" Y "<<norm2(Y)<<std::endl;
 
         // phi^dag V (Mdag M)^-1 dV^dag  phi
         Vpc.MpcDagDeriv(force , X, PhiOdd );   dSdU = force;
+	std::cout << GridLogMessage <<" deriv "<<norm2(force)<<std::endl;
   
         // phi^dag dV (Mdag M)^-1 V^dag  phi
         Vpc.MpcDeriv(force , PhiOdd, X );      dSdU = dSdU+force;
+	std::cout << GridLogMessage <<" deriv "<<norm2(force)<<std::endl;
 
         //    -    phi^dag V (Mdag M)^-1 Mdag dM   (Mdag M)^-1 V^dag  phi
         //    -    phi^dag V (Mdag M)^-1 dMdag M   (Mdag M)^-1 V^dag  phi
         Mpc.MpcDeriv(force,Y,X);              dSdU = dSdU-force;
+	std::cout << GridLogMessage <<" deriv "<<norm2(force)<<std::endl;
         Mpc.MpcDagDeriv(force,X,Y);           dSdU = dSdU-force;
+	std::cout << GridLogMessage <<" deriv "<<norm2(force)<<std::endl;
 
         // FIXME No force contribution from EvenEven assumed here
         // Needs a fix for clover.

Grid/simd/Grid_neon.h

@@ -320,7 +320,7 @@ struct Conj{
 
 struct TimesMinusI{
   //Complex single
-  inline float32x4_t operator()(float32x4_t in, float32x4_t ret){
+  inline float32x4_t operator()(float32x4_t in){
     // ar ai br bi -> ai -ar ai -br
     float32x4_t r0, r1;
     r0 = vnegq_f32(in);        // -ar -ai -br -bi
@@ -328,7 +328,7 @@ struct TimesMinusI{
     return vtrn1q_f32(r1, r0); //  ar -ai  br -bi
   }
   //Complex double
-  inline float64x2_t operator()(float64x2_t in, float64x2_t ret){
+  inline float64x2_t operator()(float64x2_t in){
     // a ib -> b -ia
     float64x2_t tmp;
     tmp = vnegq_f64(in);
@@ -338,7 +338,7 @@ struct TimesMinusI{
 
 struct TimesI{
   //Complex single
-  inline float32x4_t operator()(float32x4_t in, float32x4_t ret){
+  inline float32x4_t operator()(float32x4_t in){
     // ar ai br bi -> -ai ar -bi br
     float32x4_t r0, r1;
     r0 = vnegq_f32(in);        // -ar -ai -br -bi
@@ -346,7 +346,7 @@ struct TimesI{
     return vtrn1q_f32(r1, in); // -ai  ar -bi  br
   }
   //Complex double
-  inline float64x2_t operator()(float64x2_t in, float64x2_t ret){
+  inline float64x2_t operator()(float64x2_t in){
     // a ib -> -b ia
     float64x2_t tmp;
     tmp = vnegq_f64(in);

Grid/stencil/Stencil.h

@@ -339,8 +339,8 @@ public:
   // Vectors that live on the symmetric heap in case of SHMEM
   // These are used; either SHM objects or refs to the above symmetric heap vectors
   // depending on comms target
-  Vector<cobj *> u_simd_send_buf;
-  Vector<cobj *> u_simd_recv_buf;
+  std::vector<cobj *> u_simd_send_buf;
+  std::vector<cobj *> u_simd_recv_buf;
 
   int u_comm_offset;
   int _unified_buffer_size;
@@ -348,7 +348,7 @@ public:
   ////////////////////////////////////////
   // Stencil query
   ////////////////////////////////////////
-#ifdef SHM_FAST_PATH
+#if 1
   inline int SameNode(int point) {
 
     int dimension    = this->_directions[point];
@@ -665,11 +665,9 @@ public:
     for(int i=0;i<mm.size();i++){
       decompressor::MergeFace(decompress,mm[i]);
     }
-    if ( mm.size() )    acceleratorFenceComputeStream();
     for(int i=0;i<dd.size();i++){
       decompressor::DecompressFace(decompress,dd[i]);
     }
-    if ( dd.size() )    acceleratorFenceComputeStream();
   }
   ////////////////////////////////////////
   // Set up routines
@@ -707,6 +705,7 @@ public:
 	}
       }
     }
+    std::cout << "BuildSurfaceList size is "<<surface_list.size()<<std::endl;
   }
   /// Introduce a block structure and switch off comms on boundaries
   void DirichletBlock(const Coordinate &dirichlet_block)
@@ -1368,10 +1367,11 @@ public:
 	    int recv_from_rank;
 	    int xmit_to_rank;
 	    int shm_send=0;
-	    int shm_recv=0;
+
 	    _grid->ShiftedRanks(dimension,nbr_proc,xmit_to_rank,recv_from_rank);
 #ifdef SHM_FAST_PATH
   #warning STENCIL SHM FAST PATH SELECTED
+  	  int shm_recv=0;
 	    // shm == receive pointer         if offnode
 	    // shm == Translate[send pointer] if on node -- my view of his send pointer
 	    cobj *shm = (cobj *) _grid->ShmBufferTranslate(recv_from_rank,sp);
@@ -1404,7 +1404,6 @@ public:
 		acceleratorMemSet(rp,0,bytes); // Zero prefill comms buffer to zero
 	      }
 	      int do_send = (comms_send|comms_partial_send) && (!shm_send );
-	      int do_recv = (comms_send|comms_partial_send) && (!shm_recv );
 	      AddPacket((void *)sp,(void *)rp,
 			xmit_to_rank,do_send,
 			recv_from_rank,do_send,

Grid/tensors/Tensor_extract_merge.h

@@ -133,7 +133,6 @@ typename vobj::scalar_object extractLane(int lane, const vobj & __restrict__ vec
   typedef scalar_type * pointer;
 
   constexpr int words=sizeof(vobj)/sizeof(vector_type);
-  constexpr int Nsimd=vector_type::Nsimd();
 
   scalar_object extracted;
   pointer __restrict__  sp = (pointer)&extracted; // Type pun
@@ -153,7 +152,6 @@ void insertLane(int lane, vobj & __restrict__ vec,const typename vobj::scalar_ob
   typedef scalar_type * pointer;
 
   constexpr int words=sizeof(vobj)/sizeof(vector_type);
-  constexpr int Nsimd=vector_type::Nsimd();
 
   pointer __restrict__ sp = (pointer)&extracted;
   vector_type *vp = (vector_type *)&vec;
@@ -178,8 +176,6 @@ void extract(const vobj &vec,const ExtractPointerArray<sobj> &extracted, int off
   const int s = Nsimd/Nextr;
 
   vector_type * vp = (vector_type *)&vec;
-  scalar_type      vtmp;
-  sobj_scalar_type stmp;
   for(int w=0;w<words;w++){
     for(int i=0;i<Nextr;i++){
       sobj_scalar_type * pointer = (sobj_scalar_type *)& extracted[i][offset];
@@ -205,7 +201,6 @@ void merge(vobj &vec,const ExtractPointerArray<sobj> &extracted, int offset)
 
   vector_type * vp = (vector_type *)&vec;
   scalar_type      vtmp;
-  sobj_scalar_type stmp;
   for(int w=0;w<words;w++){
     for(int i=0;i<Nextr;i++){
       sobj_scalar_type * pointer = (sobj_scalar_type *)& extracted[i][offset];
@@ -242,9 +237,6 @@ void copyLane(vobjOut & __restrict__ vecOut, int lane_out, const vobjIn & __rest
   typedef oextract_type * opointer;
   typedef iextract_type * ipointer;
 
-  constexpr int oNsimd=ovector_type::Nsimd();
-  constexpr int iNsimd=ivector_type::Nsimd();
-
   iscalar_type itmp;
   oscalar_type otmp;
 

Grid/threads/Accelerator.h

@@ -458,7 +458,8 @@ inline void acceleratorCopySynchronise(void) { hipStreamSynchronize(copyStream);
 // Common on all GPU targets
 //////////////////////////////////////////////
 #if defined(GRID_SYCL) || defined(GRID_CUDA) || defined(GRID_HIP)
-#define accelerator_forNB( iter1, num1, nsimd, ... ) accelerator_for2dNB( iter1, num1, iter2, 1, nsimd, {__VA_ARGS__} );
+// FIXME -- the non-blocking nature got broken March 30 2023 by PAB
+#define accelerator_forNB( iter1, num1, nsimd, ... ) accelerator_for2dNB( iter1, num1, iter2, 1, nsimd, {__VA_ARGS__} );  
 
 #define accelerator_for( iter, num, nsimd, ... )		\
   accelerator_forNB(iter, num, nsimd, { __VA_ARGS__ } );	\
@@ -525,7 +526,7 @@ inline void acceleratorFreeCpu  (void *ptr){free(ptr);};
 //////////////////////////////////////////////
 
 #ifdef GRID_SYCL
-inline void acceleratorFenceComputeStream(void){ accelerator_barrier();};
+inline void acceleratorFenceComputeStream(void){ theGridAccelerator->ext_oneapi_submit_barrier(); };
 #else
 // Ordering within a stream guaranteed on Nvidia & AMD
 inline void acceleratorFenceComputeStream(void){ };

HMC/Mobius2p1f_DD_EOFA_96I_mshift.cc

@@ -227,7 +227,7 @@ int main(int argc, char **argv) {
   //  std::vector<Real> hasenbusch({ light_mass, 0.005, 0.0145, 0.045, 0.108, 0.25, 0.51 , pv_mass }); // Updated
   //  std::vector<Real> hasenbusch({ light_mass, 0.0145, 0.045, 0.108, 0.25, 0.51 , 0.75 , pv_mass });
 
-  int SP_iters=10000;
+  int SP_iters=9000;
   
   RationalActionParams OFRp; // Up/down
   OFRp.lo       = 6.0e-5;
@@ -362,12 +362,12 @@ int main(int argc, char **argv) {
 
   // Probably dominates the force - back to EOFA.
   OneFlavourRationalParams SFRp;
-  SFRp.lo       = 0.25;
+  SFRp.lo       = 0.1;
   SFRp.hi       = 25.0;
   SFRp.MaxIter  = 10000;
-  SFRp.tolerance= 1.0e-5;
+  SFRp.tolerance= 1.0e-8;
   SFRp.mdtolerance= 2.0e-4;
-  SFRp.degree   = 8;
+  SFRp.degree   = 12;
   SFRp.precision= 50;
   
   MobiusEOFAFermionD Strange_Op_L (U , *FGrid , *FrbGrid , *GridPtr , *GridRBPtr , strange_mass, strange_mass, pv_mass, 0.0, -1, M5, b, c);
@@ -451,7 +451,7 @@ int main(int argc, char **argv) {
   
 #define MIXED_PRECISION
 #ifdef MIXED_PRECISION
-  std::vector<GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<FermionImplPolicy,FermionImplPolicyF,FermionImplPolicy> *> Bdys;
+  std::vector<GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<FermionImplPolicy,FermionImplPolicyF> *> Bdys;
 #else
   std::vector<GeneralEvenOddRatioRationalPseudoFermionAction<FermionImplPolicy> *> Bdys;
 #endif
@@ -526,15 +526,13 @@ int main(int argc, char **argv) {
       Quotients.push_back (new TwoFlavourEvenOddRatioPseudoFermionAction<FermionImplPolicy>(*Numerators[h],*Denominators[h],*MPCG[h],*ActionMPCG[h],CG));
     } else {
 #ifdef MIXED_PRECISION
-      Bdys.push_back( new GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<FermionImplPolicy,FermionImplPolicyF,FermionImplPolicy>(
+      Bdys.push_back( new GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<FermionImplPolicy,FermionImplPolicyF>(
 			   *Numerators[h],*Denominators[h],
 			   *NumeratorsF[h],*DenominatorsF[h],
-			   *Numerators[h],*Denominators[h],
 			   OFRp, SP_iters) );
-      Bdys.push_back( new GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<FermionImplPolicy,FermionImplPolicyF,FermionImplPolicy>(
+      Bdys.push_back( new GeneralEvenOddRatioRationalMixedPrecPseudoFermionAction<FermionImplPolicy,FermionImplPolicyF>(
 			   *Numerators[h],*Denominators[h],
 			   *NumeratorsF[h],*DenominatorsF[h],
-			   *Numerators[h],*Denominators[h],
 			   OFRp, SP_iters) );
 #else
       Bdys.push_back( new GeneralEvenOddRatioRationalPseudoFermionAction<FermionImplPolicy>(*Numerators[h],*Denominators[h],OFRp));

HMC/Mobius2p1f_DD_RHMC_96I.cc

@@ -329,7 +329,6 @@ int main(int argc, char **argv) {
 
     
     auto grid4= GridPtr;
-    auto rbgrid4= GridRBPtr;
     auto rbgrid = StrangeOp.FermionRedBlackGrid();
     auto grid = StrangeOp.FermionGrid();
     if(1){

HMC/Mobius2p1f_EOFA_96I_hmc.cc

@@ -164,11 +164,6 @@ int main(int argc, char **argv) {
   typedef MobiusEOFAFermionF FermionEOFAActionF;
   typedef typename FermionActionF::FermionField FermionFieldF;
 
-  typedef WilsonImplD2 FermionImplPolicyD2;
-  typedef MobiusFermionD2 FermionActionD2;
-  typedef MobiusEOFAFermionD2 FermionEOFAActionD2;
-  typedef typename FermionActionD2::FermionField FermionFieldD2;
-
   typedef Grid::XmlReader       Serialiser;
 
   //::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::
@@ -249,11 +244,6 @@ int main(int argc, char **argv) {
   Coordinate shm;
 
   GlobalSharedMemory::GetShmDims(mpi,shm);
-  
-  Coordinate CommDim(Nd);
-  for(int d=0;d<Nd;d++) CommDim[d]= (mpi[d]/shm[d])>1 ? 1 : 0;
-
-  Coordinate NonDirichlet(Nd+1,0);
 
   //////////////////////////
   // Fermion Grids
@@ -272,7 +262,6 @@ int main(int argc, char **argv) {
   // temporarily need a gauge field
   LatticeGaugeFieldD  U(GridPtr); U=Zero();
   LatticeGaugeFieldF  UF(GridPtrF); UF=Zero();
-  LatticeGaugeFieldD2 UD2(GridPtrF); UD2=Zero();
 
   std::cout << GridLogMessage << " Running the HMC "<< std::endl;
   TheHMC.ReadCommandLine(argc,argv);  // params on CML or from param file
@@ -283,8 +272,6 @@ int main(int argc, char **argv) {
   std::vector<Complex> boundary = {1,1,1,-1};
   FermionAction::ImplParams Params(boundary);
   FermionActionF::ImplParams ParamsF(boundary);
-  Params.dirichlet=NonDirichlet;
-  ParamsF.dirichlet=NonDirichlet;
 
   //  double StoppingCondition = 1e-14;
   //  double MDStoppingCondition = 1e-9;
@@ -311,12 +298,12 @@ int main(int argc, char **argv) {
 
   // Probably dominates the force - back to EOFA.
   OneFlavourRationalParams SFRp;
-  SFRp.lo       = 0.25;
-  SFRp.hi       = 25.0;
+  SFRp.lo       = 0.1;
+  SFRp.hi       = 30.0;
   SFRp.MaxIter  = 10000;
-  SFRp.tolerance= 1.0e-5;
-  SFRp.mdtolerance= 2.0e-4;
-  SFRp.degree   = 8;
+  SFRp.tolerance= 1.0e-8;
+  SFRp.mdtolerance= 2.0e-6;
+  SFRp.degree   = 10;
   SFRp.precision= 50;
   
   MobiusEOFAFermionD Strange_Op_L (U , *FGrid , *FrbGrid , *GridPtr , *GridRBPtr , strange_mass, strange_mass, pv_mass, 0.0, -1, M5, b, c);
@@ -376,33 +363,29 @@ int main(int argc, char **argv) {
   ////////////////////////////////////
   std::vector<Real> light_den;
   std::vector<Real> light_num;
-  std::vector<int> dirichlet_den;
-  std::vector<int> dirichlet_num;
 
   int n_hasenbusch = hasenbusch.size();
-  light_den.push_back(light_mass);  dirichlet_den.push_back(0);
+  light_den.push_back(light_mass); 
   for(int h=0;h<n_hasenbusch;h++){
-    light_den.push_back(hasenbusch[h]); dirichlet_den.push_back(0);
+    light_den.push_back(hasenbusch[h]);
   }
 
   for(int h=0;h<n_hasenbusch;h++){
-    light_num.push_back(hasenbusch[h]); dirichlet_num.push_back(0);
+    light_num.push_back(hasenbusch[h]);
   }
-  light_num.push_back(pv_mass);  dirichlet_num.push_back(0);
+  light_num.push_back(pv_mass);
 
   std::vector<FermionAction *> Numerators;
   std::vector<FermionAction *> Denominators;
   std::vector<FermionActionF *> NumeratorsF;
   std::vector<FermionActionF *> DenominatorsF;
-  std::vector<FermionActionD2 *> NumeratorsD2;
-  std::vector<FermionActionD2 *> DenominatorsD2;
   std::vector<TwoFlavourEvenOddRatioPseudoFermionAction<FermionImplPolicy> *> Quotients;
   std::vector<MxPCG *> ActionMPCG;
   std::vector<MxPCG *> MPCG;
   
 #define MIXED_PRECISION
 #ifdef MIXED_PRECISION
-  std::vector<OneFlavourEvenOddRatioRationalMixedPrecPseudoFermionAction<FermionImplPolicy,FermionImplPolicyF,FermionImplPolicyD2> *> Bdys;
+  std::vector<OneFlavourEvenOddRatioRationalMixedPrecPseudoFermionAction<FermionImplPolicy,FermionImplPolicyF> *> Bdys;
 #else
   std::vector<OneFlavourEvenOddRatioRationalPseudoFermionAction<FermionImplPolicy> *> Bdys;
 #endif
@@ -416,9 +399,7 @@ int main(int argc, char **argv) {
     std::cout << GridLogMessage
 	      << " 2f quotient Action ";
     std::cout << "det D("<<light_den[h]<<")";
-    if ( dirichlet_den[h] ) std::cout << "^dirichlet    ";
     std::cout << "/ det D("<<light_num[h]<<")";
-    if ( dirichlet_num[h] ) std::cout << "^dirichlet    ";
     std::cout << std::endl;
 
     FermionAction::ImplParams ParamsNum(boundary);
@@ -426,21 +407,11 @@ int main(int argc, char **argv) {
     FermionActionF::ImplParams ParamsDenF(boundary);
     FermionActionF::ImplParams ParamsNumF(boundary);
     
-    ParamsNum.dirichlet = NonDirichlet;
-    ParamsDen.dirichlet = NonDirichlet;
-
-    ParamsNum.partialDirichlet = 0;
-    ParamsDen.partialDirichlet = 0;
-    
     Numerators.push_back  (new FermionAction(U,*FGrid,*FrbGrid,*GridPtr,*GridRBPtr,light_num[h],M5,b,c, ParamsNum));
     Denominators.push_back(new FermionAction(U,*FGrid,*FrbGrid,*GridPtr,*GridRBPtr,light_den[h],M5,b,c, ParamsDen));
 
-    ParamsDenF.dirichlet = ParamsDen.dirichlet;
-    ParamsDenF.partialDirichlet = ParamsDen.partialDirichlet;
     DenominatorsF.push_back(new FermionActionF(UF,*FGridF,*FrbGridF,*GridPtrF,*GridRBPtrF,light_den[h],M5,b,c, ParamsDenF));
 
-    ParamsNumF.dirichlet = ParamsNum.dirichlet;
-    ParamsNumF.partialDirichlet = ParamsNum.partialDirichlet;
     NumeratorsF.push_back  (new FermionActionF(UF,*FGridF,*FrbGridF,*GridPtrF,*GridRBPtrF,light_num[h],M5,b,c, ParamsNumF));
 
     LinOpD.push_back(new LinearOperatorD(*Denominators[h]));
@@ -477,7 +448,6 @@ int main(int argc, char **argv) {
   // Gauge action
   /////////////////////////////////////////////////////////////
   Level3.push_back(&GaugeAction);
-  //  TheHMC.TheAction.push_back(Level1);
   TheHMC.TheAction.push_back(Level2);
   TheHMC.TheAction.push_back(Level3);
   std::cout << GridLogMessage << " Action complete "<< std::endl;

README.md

@@ -1,7 +1,8 @@
-# Grid [![Teamcity status](http://ci.cliath.ph.ed.ac.uk/app/rest/builds/aggregated/strob:(buildType:(affectedProject(id:GridBasedSoftware_Grid)),branch:name:develop)/statusIcon.svg)](http://ci.cliath.ph.ed.ac.uk/project.html?projectId=GridBasedSoftware_Grid&tab=projectOverview) 
-
+# Grid 
 **Data parallel C++ mathematical object library.**
 
+[![Teamcity status](https://ci.dev.dirac.ed.ac.uk/guestAuth/app/rest/builds/aggregated/strob:(buildType:(affectedProject(id:GridBasedSoftware_Grid)),branch:default:true)/statusIcon.svg)](https://ci.dev.dirac.ed.ac.uk/project/GridBasedSoftware_Grid?mode=builds) 
+
 License: GPL v2.
 
 Last update June 2017.

benchmarks/Benchmark_dwf_fp32.cc

@@ -425,7 +425,7 @@ void Benchmark(int Ls, Coordinate Dirichlet)
 
   err = r_eo-result;
   n2e= norm2(err);
-  std::cout<<GridLogMessage << "norm diff   "<< n2e<< "  Line "<<__LINE__ <<std::endl;
+  std::cout<<GridLogMessage << "norm diff   "<< n2e<<std::endl;
   assert(n2e<1.0e-4);
 
   pickCheckerboard(Even,src_e,err);

examples/socket_grid.cc

@@ -0,0 +1,133 @@
+#include <sys/socket.h>
+#include <sys/un.h>
+#include <unistd.h>
+#include <stdio.h>
+#include <err.h>
+#include <fcntl.h>
+#include <assert.h>
+#include <string.h>
+#include <stdlib.h>
+
+static int sock;
+static const char *sock_path_fmt = "/tmp/GridUnixSocket.%d";
+static char sock_path[256];
+
+class UnixSockets {
+public:
+  static void Open(int rank)
+  {
+    int errnum;
+
+    sock = socket(AF_UNIX, SOCK_DGRAM, 0);  assert(sock>0);
+    printf("allocated socket %d\n",sock);
+
+    struct sockaddr_un sa_un = { 0 };
+    sa_un.sun_family = AF_UNIX;
+    snprintf(sa_un.sun_path, sizeof(sa_un.sun_path),sock_path_fmt,rank);
+    unlink(sa_un.sun_path);
+    if (bind(sock, (struct sockaddr *)&sa_un, sizeof(sa_un))) {
+      perror("bind failure");
+      exit(EXIT_FAILURE);
+    }
+    printf("bound socket %d to %s\n",sock,sa_un.sun_path);
+  }
+
+  static int RecvFileDescriptor(void)
+  {
+    int n;
+    int fd;
+    char buf[1];
+    struct iovec iov;
+    struct msghdr msg;
+    struct cmsghdr *cmsg;
+    char cms[CMSG_SPACE(sizeof(int))];
+
+    iov.iov_base = buf;
+    iov.iov_len = 1;
+
+    memset(&msg, 0, sizeof msg);
+    msg.msg_name = 0;
+    msg.msg_namelen = 0;
+    msg.msg_iov = &iov;
+    msg.msg_iovlen = 1;
+
+    msg.msg_control = (caddr_t)cms;
+    msg.msg_controllen = sizeof cms;
+
+    if((n=recvmsg(sock, &msg, 0)) < 0) {
+      perror("recvmsg failed");
+      return -1;
+    }
+    if(n == 0){
+      perror("recvmsg returned 0");
+      return -1;
+    }
+    cmsg = CMSG_FIRSTHDR(&msg);
+    memmove(&fd, CMSG_DATA(cmsg), sizeof(int));
+    printf("received fd %d from socket %d\n",fd,sock);
+    return fd;
+  }
+
+  static void SendFileDescriptor(int fildes,int xmit_to_rank)
+  {
+    struct msghdr msg;
+    struct iovec iov;
+    struct cmsghdr *cmsg = NULL;
+    char ctrl[CMSG_SPACE(sizeof(int))];
+    char data = ' ';
+
+    memset(&msg, 0, sizeof(struct msghdr));
+    memset(ctrl, 0, CMSG_SPACE(sizeof(int)));
+    iov.iov_base = &data;
+    iov.iov_len = sizeof(data);
+    
+    sprintf(sock_path,sock_path_fmt,xmit_to_rank);
+    printf("sending FD %d over socket %d to rank %d AF_UNIX path %s\n",fildes,sock,xmit_to_rank,sock_path);fflush(stdout);
+    
+    struct sockaddr_un sa_un = { 0 };
+    sa_un.sun_family = AF_UNIX;
+    snprintf(sa_un.sun_path, sizeof(sa_un.sun_path),sock_path_fmt,xmit_to_rank);
+
+    msg.msg_name = (void *)&sa_un;
+    msg.msg_namelen = sizeof(sa_un);
+    msg.msg_iov = &iov;
+    msg.msg_iovlen = 1;
+    msg.msg_controllen =  CMSG_SPACE(sizeof(int));
+    msg.msg_control = ctrl;
+
+    cmsg = CMSG_FIRSTHDR(&msg);
+    cmsg->cmsg_level = SOL_SOCKET;
+    cmsg->cmsg_type = SCM_RIGHTS;
+    cmsg->cmsg_len = CMSG_LEN(sizeof(int));
+
+    *((int *) CMSG_DATA(cmsg)) = fildes;
+
+    if ( sendmsg(sock, &msg, 0) == -1 ) perror("sendmsg failed");
+  };
+};
+
+int main(int argc, char **argv)
+{
+  int me = fork()?0:1;
+  
+  UnixSockets::Open(me);
+  
+  // need MPI barrier
+  sleep(10);
+  const char * message = "Hello, World\n";
+  if( me ) {
+    int fd = open("foo",O_RDWR|O_CREAT,0666);
+    if ( fd < 0 ) {
+      perror("failed to open file");
+      exit(EXIT_FAILURE);
+    }
+    // rank 1 sends ot rank 0
+    UnixSockets::SendFileDescriptor(fd,0);
+    close(fd);
+  } else {
+    // rank 0 sends receives frmo rank 1
+    int fd = UnixSockets::RecvFileDescriptor();
+    write(fd,(const void *)message,strlen(message));
+    close(fd);
+  }
+}

systems/Crusher/config-command

@@ -1,7 +1,7 @@
 CLIME=`spack find --paths c-lime@2-3-9 | grep c-lime| cut -c 15-`
 ../../configure --enable-comms=mpi-auto \
 --with-lime=$CLIME \
---enable-unified=yes \
+--enable-unified=no \
 --enable-shm=nvlink \
 --enable-tracing=timer \
 --enable-accelerator=hip \

systems/Crusher/sourceme.sh

@@ -5,8 +5,8 @@ module load emacs
 #module load gperftools
 module load PrgEnv-gnu
 module load rocm/5.3.0
-module load cray-mpich/8.1.16
-#module load cray-mpich/8.1.17
+#module load cray-mpich/8.1.16
+module load cray-mpich/8.1.17
 module load gmp
 module load cray-fftw
 module load craype-accel-amd-gfx90a

systems/PVC/setup.sh

@@ -3,8 +3,14 @@ export https_proxy=http://proxy-chain.intel.com:911
 export LD_LIBRARY_PATH=$HOME/prereqs/lib/:$LD_LIBRARY_PATH
 
 module load intel-release
-source /opt/intel/oneapi/PVC_setup.sh
+module load intel-comp-rt/embargo-ci-neo
+
+#source /opt/intel/oneapi/PVC_setup.sh
 #source /opt/intel/oneapi/ATS_setup.sh
+#module load intel-nightly/20230331
+#module load intel-comp-rt/ci-neo-master/026093
+
+#module load intel/mpich
 module load intel/mpich/pvc45.3
 export PATH=~/ATS/pti-gpu/tools/onetrace/:$PATH
 

tests/core/Test_compact_wilson_clover_speedup.cc

@@ -53,7 +53,7 @@ static int readInt(int* argc, char*** argv, std::string&& option, int defaultVal
 
 static float readFloat(int* argc, char*** argv, std::string&& option, float defaultValue) {
   std::string arg;
-  float       ret = defaultValue;
+  double      ret = defaultValue;
   if(checkPresent(argc, argv, option)) {
     arg = getContent(argc, argv, option);
     GridCmdOptionFloat(arg, ret);

tests/core/Test_fft_matt.cc

@@ -1,244 +0,0 @@
-    /*************************************************************************************
-
-Gamma::Algebra Gmu [] = {
-  Gamma::Algebra::GammaX,
-  Gamma::Algebra::GammaY,
-  Gamma::Algebra::GammaZ,
-  Gamma::Algebra::GammaT,
-  Gamma::Algebra::Gamma5
-};
-
-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 = GridDefaultSimd(Nd,vComplexD::Nsimd());
-  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);
-
-  LatticeComplexD    coor(&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
-
-  ////////////////////////////////////////////////////
-  // Wilson test
-  ////////////////////////////////////////////////////
-  {
-    LatticeFermionD    src(&GRID); gaussian(pRNG,src);
-    LatticeFermionD    src_p(&GRID);
-    LatticeFermionD    tmp(&GRID);
-    LatticeFermionD    ref(&GRID);
-    LatticeFermionD    result(&GRID);
-    
-    RealD mass=0.1;
-    WilsonFermionD Dw(Umu,GRID,RBGRID,mass);
-    
-    Dw.M(src,ref);
-    std::cout << "Norm src "<<norm2(src)<<std::endl;
-    std::cout << "Norm Dw x src "<<norm2(ref)<<std::endl;
-    {
-      FFT theFFT(&GRID);
-
-      ////////////////
-      // operator in Fourier space
-      ////////////////
-      tmp =ref;
-      theFFT.FFT_all_dim(result,tmp,FFT::forward);
-      std::cout<<"FFT[ Dw x src ]  "<< norm2(result)<<std::endl;    
-
-      tmp = src;
-      theFFT.FFT_all_dim(src_p,tmp,FFT::forward);
-      std::cout<<"FFT[ src      ]  "<< norm2(src_p)<<std::endl;
-      
-      /////////////////////////////////////////////////////////////////
-      // work out the predicted FT from Fourier
-      /////////////////////////////////////////////////////////////////
-      auto FGrid = &GRID;
-      LatticeFermionD    Kinetic(FGrid); Kinetic = Zero();
-      LatticeComplexD    kmu(FGrid); 
-      LatticeInteger     scoor(FGrid); 
-      LatticeComplexD    sk (FGrid); sk = Zero();
-      LatticeComplexD    sk2(FGrid); sk2= Zero();
-      LatticeComplexD    W(FGrid); W= Zero();
-      LatticeComplexD    one(FGrid); one =ComplexD(1.0,0.0);
-      ComplexD ci(0.0,1.0);
-    
-      for(int mu=0;mu<Nd;mu++) {
-	
-	RealD TwoPiL =  M_PI * 2.0/ latt_size[mu];
-
-	LatticeCoordinate(kmu,mu);
-
-	kmu = TwoPiL * kmu;
-      
-	sk2 = sk2 + 2.0*sin(kmu*0.5)*sin(kmu*0.5);
-	sk  = sk  +     sin(kmu)    *sin(kmu); 
-      
-	// -1/2 Dw ->  1/2 gmu (eip - emip) = i sinp gmu
-	Kinetic = Kinetic + sin(kmu)*ci*(Gamma(Gmu[mu])*src_p);
-	
-      }
-    
-      W = mass + sk2; 
-      Kinetic = Kinetic + W * src_p;
-    
-      std::cout<<"Momentum space src         "<< norm2(src_p)<<std::endl;
-      std::cout<<"Momentum space Dw x src    "<< norm2(Kinetic)<<std::endl;
-      std::cout<<"FT[Coordinate space Dw]    "<< norm2(result)<<std::endl;
-    
-      result = result - Kinetic;
-      std::cout<<"diff "<< norm2(result)<<std::endl;
-      
-    }
-
-    std::cout << " =======================================" <<std::endl;
-    std::cout << " Checking FourierFreePropagator x Dw = 1" <<std::endl;
-    std::cout << " =======================================" <<std::endl;
-    std::cout << "Dw src = " <<norm2(src)<<std::endl;
-    std::cout << "Dw tmp = " <<norm2(tmp)<<std::endl;
-    Dw.M(src,tmp);
-
-    Dw.FreePropagator(tmp,ref,mass);
-
-    std::cout << "Dw ref = " <<norm2(ref)<<std::endl;
-    
-    ref = ref - src;
-    
-    std::cout << "Dw ref-src = " <<norm2(ref)<<std::endl;
-  }
-
-
-  ////////////////////////////////////////////////////
-  // Wilson prop
-  ////////////////////////////////////////////////////
-  {
-    std::cout<<"****************************************"<<std::endl;
-    std::cout << "Wilson Mom space 4d propagator \n";
-    std::cout<<"****************************************"<<std::endl;
-
-    LatticeFermionD    src(&GRID); gaussian(pRNG,src);
-    LatticeFermionD    tmp(&GRID);
-    LatticeFermionD    ref(&GRID);
-    LatticeFermionD    diff(&GRID);
-
-    src=Zero();
-    Coordinate point(4,0); // 0,0,0,0
-    SpinColourVectorD ferm;
-    ferm=Zero();
-    ferm()(0)(0) = ComplexD(1.0);
-    pokeSite(ferm,src,point);
-
-    RealD mass=0.1;
-
-    WilsonFermionD Dw(Umu,GRID,RBGRID,mass);
-
-    // Momentum space prop
-    std::cout << " Solving by FFT and Feynman rules" <<std::endl;
-    Dw.FreePropagator(src,ref,mass) ;
-
-    Gamma G5(Gamma::Algebra::Gamma5);
-
-    LatticeFermionD    result(&GRID); 
-    const int sdir=0;
-    
-    ////////////////////////////////////////////////////////////////////////
-    // Conjugate gradient on normal equations system
-    ////////////////////////////////////////////////////////////////////////
-    std::cout << " Solving by Conjugate Gradient (CGNE)" <<std::endl;
-    Dw.Mdag(src,tmp);
-    src=tmp;
-    MdagMLinearOperator<WilsonFermionD,LatticeFermionD> HermOp(Dw);
-    ConjugateGradient<LatticeFermionD> CG(1.0e-10,10000);
-    CG(HermOp,src,result);
-    
-    ////////////////////////////////////////////////////////////////////////
-    std::cout << " Taking difference" <<std::endl;
-    std::cout << "Dw result "<<norm2(result)<<std::endl;
-    std::cout << "Dw ref     "<<norm2(ref)<<std::endl;
-    
-    diff = ref - result;
-    std::cout << "result - ref     "<<norm2(diff)<<std::endl;
-
-    DumpSliceNorm("Slice Norm Solution ",result,Nd-1);
-  }
-
-  ////////////////////////////////////////////////////
-  //Gauge invariance test
-  ////////////////////////////////////////////////////
-  {
-    std::cout<<"****************************************"<<std::endl;
-    std::cout << "Gauge invariance test \n";
-    std::cout<<"****************************************"<<std::endl;
-    LatticeGaugeField     U_GT(&GRID); // Gauge transformed field
-    LatticeColourMatrix   g(&GRID);    // local Gauge xform matrix
-    U_GT = Umu;
-    // Make a random xform to teh gauge field
-    SU<Nc>::RandomGaugeTransform(pRNG,U_GT,g); // Unit gauge
-
-    LatticeFermionD    src(&GRID);
-    LatticeFermionD    tmp(&GRID);
-    LatticeFermionD    ref(&GRID);
-    LatticeFermionD    diff(&GRID);
-
-    // could loop over colors
-    src=Zero();
-    Coordinate point(4,0); // 0,0,0,0
-    SpinColourVectorD ferm;
-    ferm=Zero();
-    ferm()(0)(0) = ComplexD(1.0);
-    pokeSite(ferm,src,point);
-
-    RealD mass=0.1;
-    WilsonFermionD Dw(U_GT,GRID,RBGRID,mass);
-
-    // Momentum space prop
-    std::cout << " Solving by FFT and Feynman rules" <<std::endl;
-    Dw.FreePropagator(src,ref,mass) ;
-
-    Gamma G5(Gamma::Algebra::Gamma5);
-
-    LatticeFermionD    result(&GRID); 
-    const int sdir=0;
-    
-    ////////////////////////////////////////////////////////////////////////
-    // Conjugate gradient on normal equations system
-    ////////////////////////////////////////////////////////////////////////
-    std::cout << " Solving by Conjugate Gradient (CGNE)" <<std::endl;
-    Dw.Mdag(src,tmp);
-    src=tmp;
-    MdagMLinearOperator<WilsonFermionD,LatticeFermionD> HermOp(Dw);
-    ConjugateGradient<LatticeFermionD> CG(1.0e-10,10000);
-    CG(HermOp,src,result);
-    
-    ////////////////////////////////////////////////////////////////////////
-    std::cout << " Taking difference" <<std::endl;
-    std::cout << "Dw result "<<norm2(result)<<std::endl;
-    std::cout << "Dw ref     "<<norm2(ref)<<std::endl;
-    
-    diff = ref - result;
-    std::cout << "result - ref     "<<norm2(diff)<<std::endl;
-
-    DumpSliceNorm("Slice Norm Solution ",result,Nd-1);
-  }
-  
-  
-  Grid_finalize();
-}

tests/forces/Test_double_ratio.cc

@@ -476,7 +476,9 @@ int main (int argc, char ** argv)
   //  ForceTest<GimplTypesR>(BdyNf2eo,U,DDHMCFilter);
 
   //////////////////// One flavour boundary det  ////////////////////
+  /*
   RationalActionParams OFRp; // Up/down
+  int SP_iters = 3000;
   OFRp.lo       = 6.0e-5;
   OFRp.hi       = 90.0;
   OFRp.inv_pow  = 2;
@@ -489,7 +491,7 @@ int main (int argc, char ** argv)
   //  OFRp.degree   = 16;
   OFRp.precision= 80;
   OFRp.BoundsCheckFreq=0;
-  /*
+  */
   OneFlavourRationalParams OFRp; // Up/down
   OFRp.lo       = 4.0e-5;
   OFRp.hi       = 90.0;
@@ -499,7 +501,6 @@ int main (int argc, char ** argv)
   OFRp.degree   = 18;
   OFRp.precision= 80;
   OFRp.BoundsCheckFreq=0;
-  */
   std::vector<RealD> ActionTolByPole({
       1.0e-7,1.0e-8,1.0e-8,1.0e-8,
       1.0e-8,1.0e-8,1.0e-8,1.0e-8,

tests/hmc/Test_hmc_Mobius2p1f.cc

@@ -85,7 +85,7 @@ int main(int argc, char **argv) {
   TheHMC.Resources.AddObservable<PlaqObs>();
   //////////////////////////////////////////////
 
-  const int Ls      = 4;
+  const int Ls      = 8;
   Real beta         = 2.13;
   Real light_mass   = 0.01;
   Real strange_mass = 0.04;