optimal torus placement for Tofu network (Fujitsu-specific)

Grid/communicator/Communicator_base.cc

@@ -34,6 +34,9 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
 NAMESPACE_BEGIN(Grid);
 
 bool Stencil_force_mpi = true;
+#ifdef GRID_FJMPI
+bool Stencil_persistent = false;
+#endif
 bool Stencil_force_barrier = false;
 
 ///////////////////////////////////////////////////////////////

Grid/communicator/Communicator_base.h

@@ -38,6 +38,10 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
 NAMESPACE_BEGIN(Grid);
 
 extern bool Stencil_force_mpi ;
+#ifdef GRID_FJMPI
+// Fujitsu MPI persistent requests for the halo exchange with other nodes (set by --comms-overlap)
+extern bool Stencil_persistent ;
+#endif
 extern bool Stencil_force_barrier ;
 
 class CartesianCommunicator : public SharedMemory {
@@ -201,6 +205,11 @@ public:
 
   // Could do a PollHtoD and have a CommsMerge dependence
   void StencilSendToRecvFromPollDtoH (std::vector<CommsRequest_t> &list);
+#ifdef A64FX
+  // Make progress on outstanding requests without blocking (MPI only progresses
+  // inside MPI calls on Fugaku); safe to call repeatedly before Complete.
+  void StencilSendToRecvFromProgress(std::vector<CommsRequest_t> &list);
+#endif
   void StencilSendToRecvFromPollIRecv(std::vector<CommsRequest_t> &list);
 
   double StencilSendToRecvFromBegin(std::vector<CommsRequest_t> &list,

Grid/communicator/Communicator_mpi3.cc

@@ -28,6 +28,90 @@ Author: Peter Boyle <paboyle@ph.ed.ac.uk>
 #include <Grid/GridCore.h>
 #include <Grid/communicator/SharedMemory.h>
 
+// Fujitsu MPI extended persistent requests (FJMPI_Prequest_*, Fujitsu MPI User's Guide J2UL-2565 5.3).
+// Fujitsu MPI moves the body of a large non-blocking message only inside MPI calls (6.2), so the halo
+// exchange does not overlap the interior kernel. With these requests the partners exchange buffer
+// information once, at initialisation, and FJMPI_Prequest_start starts the transfer of the message body.
+// GRID_FJMPI is set by configure (--enable-fjmpi) when mpi-ext.h declares the Fujitsu extensions.
+#ifdef GRID_FJMPI
+#include <mpi-ext.h>
+
+NAMESPACE_BEGIN(Grid);
+namespace {
+// Persistent halo requests are used with --comms-overlap (Stencil_persistent); they bring nothing
+// without overlap. Only for peers on other nodes: persistent transfers to peers on the same node delay
+// the transfers to other nodes.
+bool PersistentHaloOn(void) { return Stencil_persistent; }
+bool PersistentHaloPeer(int grank) { return grank==MPI_UNDEFINED; }
+struct PersistentHaloReq {
+  void *buf; int count; MPI_Request req;
+};
+// Requests are keyed by (send, communicator, peer, tag): Fujitsu MPI allows one request per
+// (peer, tag, communicator). Several stencils can use the same communicator (e.g. even and odd
+// checkerboards), so each distinct pair of halo buffers of a direction gets a slot, in order of
+// first use, and the slot enters the tag. All ranks run the same sequence of stencils, so partners
+// agree on the slot.
+struct PersistentHaloCache {
+  std::map<std::pair<MPI_Comm,int>, std::vector<std::pair<void *,void *> > > slots;
+  std::map<std::tuple<int,MPI_Comm,int,int>, PersistentHaloReq> reqs;
+  int tag_ub = -1;
+};
+PersistentHaloCache &PersistentHalo(void) { static PersistentHaloCache c; return c; }
+
+int PersistentHaloSlot(MPI_Comm comm,int dir,void *xmit,void *recv)
+{
+  auto &v = PersistentHalo().slots[std::make_pair(comm,dir)];
+  for(int k=0;k<v.size();k++) if ( v[k].first==xmit && v[k].second==recv ) return k;
+  v.push_back(std::make_pair(xmit,recv));
+  return v.size()-1;
+}
+// Tag of a halo message: the usual tag (dir + 32 * sender rank) shifted by the slot; -1 if it
+// exceeds MPI_TAG_UB (both partners then fall back to MPI_Isend/MPI_Irecv)
+int PersistentHaloTag(int dir,int sender,int nprocs,int slot)
+{
+  auto &c = PersistentHalo();
+  if ( c.tag_ub < 0 ) {
+    int *ub, flag;
+    MPI_Comm_get_attr(MPI_COMM_WORLD,MPI_TAG_UB,&ub,&flag);
+    c.tag_ub = flag ? *ub : 32767;
+  }
+  int64_t tag = dir + 32*((int64_t)sender + (int64_t)nprocs*slot);
+  return (tag <= c.tag_ub) ? (int)tag : -1;
+}
+// The inactive persistent request of one message, created on first use and recreated if the
+// buffer or size changed (a stencil freed and another one allocated at the same address)
+MPI_Request PersistentHaloGet(int send,void *buf,int count,int peer,int tag,MPI_Comm comm)
+{
+  auto key = std::make_tuple(send,comm,peer,tag);
+  auto &reqs = PersistentHalo().reqs;
+  auto it = reqs.find(key);
+  if ( it != reqs.end() ) {
+    if ( it->second.buf==buf && it->second.count==count ) return it->second.req;
+    MPI_Request_free(&it->second.req);
+    reqs.erase(it);
+  }
+  PersistentHaloReq r = {buf,count,MPI_REQUEST_NULL};
+  int ierr = send ? FJMPI_Prequest_send_init(buf,count,MPI_INT32_T,peer,tag,comm,&r.req)
+                  : FJMPI_Prequest_recv_init(buf,count,MPI_INT32_T,peer,tag,comm,&r.req);
+  GRID_ASSERT(ierr==0);
+  reqs[key] = r;
+  return r.req;
+}
+void PersistentHaloFree(MPI_Comm comm)
+{
+  auto &c = PersistentHalo();
+  for(auto it=c.reqs.begin();it!=c.reqs.end();) {
+    if ( std::get<1>(it->first)==comm ) { MPI_Request_free(&it->second.req); it=c.reqs.erase(it); }
+    else ++it;
+  }
+  for(auto it=c.slots.begin();it!=c.slots.end();) {
+    if ( it->first.first==comm ) it=c.slots.erase(it); else ++it;
+  }
+}
+}
+NAMESPACE_END(Grid);
+#endif
+
 void GridAbort(void) { MPI_Abort(MPI_COMM_WORLD,SIGABRT); }
 extern void * Grid_backtrace_buffer[_NBACKTRACE];
 
@@ -262,6 +346,9 @@ CartesianCommunicator::~CartesianCommunicator()
   if (communicator && !MPI_is_finalised) {
     MPI_Comm_free(&communicator);
     for(int i=0;i<communicator_halo.size();i++){
+#ifdef GRID_FJMPI
+      PersistentHaloFree(communicator_halo[i]);
+#endif
       MPI_Comm_free(&communicator_halo[i]);
     }
   }
@@ -421,6 +508,18 @@ int CartesianCommunicator::IsOffNode(int rank)
 #ifdef ACCELERATOR_AWARE_MPI
 void CartesianCommunicator::StencilSendToRecvFromPollIRecv(std::vector<CommsRequest_t> &list) {};
 void CartesianCommunicator::StencilSendToRecvFromPollDtoH(std::vector<CommsRequest_t> &list) {};
+#ifdef A64FX
+void CartesianCommunicator::StencilSendToRecvFromProgress(std::vector<CommsRequest_t> &list)
+{
+  int nreq=list.size();
+  if (nreq==0) return;
+  int flag;
+  // completed requests become MPI_REQUEST_NULL (persistent ones become inactive), which the final
+  // Waitall handles
+  int ierr = MPI_Testall(nreq,&list[0],&flag,MPI_STATUSES_IGNORE);
+  GRID_ASSERT(ierr==0);
+}
+#endif
 double CartesianCommunicator::StencilSendToRecvFromPrepare(std::vector<CommsRequest_t> &list,
 							   void *xmit,
 							   int dest,int dox,
@@ -428,6 +527,34 @@ double CartesianCommunicator::StencilSendToRecvFromPrepare(std::vector<CommsRequ
 							   int from,int dor,
 							   uint64_t xbytes,uint64_t rbytes,int dir)
 {
+#ifdef GRID_FJMPI
+  // Persistent requests. Starting a persistent receive waits until the partner has created the
+  // matching send request, and starting a send waits until the partner has started the receive
+  // (measured; the manual documents the latter). So both requests of this packet are created here
+  // before its receive is started, and all receives are started before any send is started in
+  // StencilSendToRecvFromBegin. Partners process the packets in the same order, so no cycle forms.
+  if ( PersistentHaloOn() ) {
+    int ncomm  =communicator_halo.size();
+    int commdir=dir%ncomm;
+    int gfrom = ShmRanks[from];
+    int gdest = ShmRanks[dest];
+    MPI_Comm comm = communicator_halo[commdir];
+    int slot = PersistentHaloSlot(comm,dir,xmit,recv);
+    int rtag = PersistentHaloTag(dir,from,_Nprocessors,slot);
+    int xtag = PersistentHaloTag(dir,_processor,_Nprocessors,slot);
+    bool mpi_recv = dor && ( (gfrom ==MPI_UNDEFINED) || Stencil_force_mpi ) && rtag >= 0 && PersistentHaloPeer(gfrom);
+    bool mpi_send = dox && ( (gdest ==MPI_UNDEFINED) || Stencil_force_mpi ) && xtag >= 0 && PersistentHaloPeer(gdest);
+    MPI_Request rrq;
+    if ( mpi_recv ) rrq = PersistentHaloGet(0,recv,(int)(rbytes/sizeof(int32_t)),from,rtag,comm);
+    if ( mpi_send ) PersistentHaloGet(1,xmit,(int)(xbytes/sizeof(int32_t)),dest,xtag,comm);
+    if ( mpi_recv ) {
+      int ierr = FJMPI_Prequest_start(&rrq);
+      GRID_ASSERT(ierr==0);
+      list.push_back(rrq);
+      return rbytes;
+    }
+  }
+#endif
   return 0.0; // Do nothing -- no preparation required
 }
 double CartesianCommunicator::StencilSendToRecvFromBegin(std::vector<CommsRequest_t> &list,
@@ -454,7 +581,19 @@ double CartesianCommunicator::StencilSendToRecvFromBegin(std::vector<CommsReques
   double off_node_bytes=0.0;
   int tag;
   
+#ifdef GRID_FJMPI
+  // Persistent requests (see StencilSendToRecvFromPrepare): same slot and tags on both sides
+  int ptag_recv = -1, ptag_send = -1;
+  if ( PersistentHaloOn() ) {
+    MPI_Comm comm = communicator_halo[commdir];
+    int slot = PersistentHaloSlot(comm,dir,xmit_comp,recv_comp);
+    if ( PersistentHaloPeer(gfrom) ) ptag_recv = PersistentHaloTag(dir,from,_Nprocessors,slot);
+    if ( PersistentHaloPeer(gdest) ) ptag_send = PersistentHaloTag(dir,_processor,_Nprocessors,slot);
+  }
+  if ( dor && ptag_recv < 0 ) { // receive not started in StencilSendToRecvFromPrepare
+#else
   if ( dor ) {
+#endif
     if ( (gfrom ==MPI_UNDEFINED) || Stencil_force_mpi ) {
       tag= dir+from*32;
       //      std::cout << " StencilSendToRecvFrom "<<dir<<" MPI_Irecv "<<std::hex<<recv<<std::dec<<std::endl;
@@ -475,8 +614,18 @@ double CartesianCommunicator::StencilSendToRecvFromBegin(std::vector<CommsReques
   // This is a NVLINK PUT  
   if (dox) {
     if ( (gdest == MPI_UNDEFINED) || Stencil_force_mpi ) {
+#ifdef GRID_FJMPI
+      if ( ptag_send >= 0 ) {
+	xrq = PersistentHaloGet(1,xmit_comp,(int)(xbytes/sizeof(int32_t)),dest,ptag_send,communicator_halo[commdir]);
+	ierr = FJMPI_Prequest_start(&xrq);
+      } else {
+	tag= dir+_processor*32;
+	ierr =MPI_Isend(xmit_comp,(int)(xbytes/sizeof(int32_t)), MPI_INT32_T,dest,tag,communicator_halo[commdir],&xrq);
+      }
+#else
       tag= dir+_processor*32;
       ierr =MPI_Isend(xmit_comp,(int)(xbytes/sizeof(int32_t)), MPI_INT32_T,dest,tag,communicator_halo[commdir],&xrq);
+#endif
       GRID_ASSERT(ierr==0);
       list.push_back(xrq);
       off_node_bytes+=xbytes;
@@ -506,6 +655,9 @@ void CartesianCommunicator::StencilSendToRecvFromComplete(std::vector<CommsReque
 }
 
 #else /* NOT     ... ACCELERATOR_AWARE_MPI */
+#ifdef A64FX
+void CartesianCommunicator::StencilSendToRecvFromProgress(std::vector<CommsRequest_t> &list) {};
+#endif
 ///////////////////////////////////////////
 // Pipeline mode through host memory
 ///////////////////////////////////////////

Grid/communicator/Communicator_none.cc

@@ -167,6 +167,9 @@ double CartesianCommunicator::StencilSendToRecvFrom( void *xmit,
 }
 void CartesianCommunicator::StencilSendToRecvFromPollIRecv(std::vector<CommsRequest_t> &list) {};
 void CartesianCommunicator::StencilSendToRecvFromPollDtoH(std::vector<CommsRequest_t> &list) {};
+#ifdef A64FX
+void CartesianCommunicator::StencilSendToRecvFromProgress(std::vector<CommsRequest_t> &list) {};
+#endif
 double CartesianCommunicator::StencilSendToRecvFromPrepare(std::vector<CommsRequest_t> &list,
 							   void *xmit,
 							   int xmit_to_rank,int dox,

Grid/communicator/SharedMemory.cc

@@ -33,6 +33,7 @@ NAMESPACE_BEGIN(Grid);
 // static data
 
 int                 GlobalSharedMemory::HPEhypercube = 1;
+int                 GlobalSharedMemory::TofuPlacement = 0;
 uint64_t            GlobalSharedMemory::MAX_MPI_SHM_BYTES   = 1024LL*1024LL*1024LL; 
 int                 GlobalSharedMemory::Hugepages = 0;
 int                 GlobalSharedMemory::_ShmSetup;

Grid/communicator/SharedMemory.h

@@ -99,6 +99,10 @@ public:
   // HPE 8600 hypercube optimisation
   ///////////////////////////////////////
   static int HPEhypercube;
+  ///////////////////////////////////////
+  // Tofu (Fujitsu MPI) node placement, --tofu
+  ///////////////////////////////////////
+  static int TofuPlacement;
 
   static int      ShmSetup(void)      { return _ShmSetup; }
   static int      ShmAlloc(void)      { return _ShmAlloc; }
@@ -131,6 +135,7 @@ public:
   static void OptimalCommunicator            (const Coordinate &processors,Grid_MPI_Comm & optimal_comm,Coordinate &ShmDims); 
   static void OptimalCommunicatorHypercube   (const Coordinate &processors,Grid_MPI_Comm & optimal_comm,Coordinate &ShmDims); 
   static void OptimalCommunicatorSharedMemory(const Coordinate &processors,Grid_MPI_Comm & optimal_comm,Coordinate &ShmDims); 
+  static void OptimalCommunicatorTofu        (const Coordinate &processors,Grid_MPI_Comm & optimal_comm,Coordinate &ShmDims);
   static void GetShmDims(const Coordinate &WorldDims,Coordinate &ShmDims);
   ///////////////////////////////////////////////////
   // Provide shared memory facilities off comm world

Grid/communicator/SharedMemoryMPI.cc

@@ -32,6 +32,10 @@ Author: Christoph Lehner <christoph@lhnr.de>
 #include <Grid/GridCore.h>
 #include <pwd.h>
 
+#ifdef GRID_FJMPI
+#include <mpi-ext.h>
+#endif
+
 #ifdef GRID_CUDA
 #include <cuda_runtime_api.h>
 #endif
@@ -278,6 +282,18 @@ void GlobalSharedMemory::OptimalCommunicator(const Coordinate &processors,Grid_M
   gethostname(name,namelen);
   int nscan = sscanf(name,"r%di%dn%d",&R,&I,&N) ;
 
+  if ( TofuPlacement ) {
+#ifdef GRID_FJMPI
+    OptimalCommunicatorTofu(processors,optimal_comm,SHM);
+#else
+    if ( WorldRank==0 ) {
+      std::cerr << "Grid --tofu: not built with Fujitsu MPI" << std::endl;
+    }
+    MPI_Barrier(WorldComm);
+    MPI_Abort(WorldComm,EXIT_FAILURE);
+#endif
+    return;
+  }
   if(nscan==3 && HPEhypercube ) OptimalCommunicatorHypercube(processors,optimal_comm,SHM);
   else                          OptimalCommunicatorSharedMemory(processors,optimal_comm,SHM);
 }
@@ -450,6 +466,202 @@ void GlobalSharedMemory::OptimalCommunicatorSharedMemory(const Coordinate &proce
   int ierr= MPI_Comm_split(WorldComm,0,rank,&optimal_comm);
   GRID_ASSERT(ierr==0);
 }
+#ifdef GRID_FJMPI
+////////////////////////////////////////////////////////////////////////////////////////////
+// Tofu placement (--tofu): lay the node grid onto the job's logical torus (pjsub -L node=AxBxC:torus),
+// each direction along a job axis of the same extent. Errors are detected on all ranks alike.
+////////////////////////////////////////////////////////////////////////////////////////////
+static void TofuPlacementError(const std::string &msg)
+{
+  if ( GlobalSharedMemory::WorldRank==0 ) std::cerr << "Grid --tofu: " << msg << std::endl;
+  MPI_Barrier(GlobalSharedMemory::WorldComm);
+  MPI_Abort(GlobalSharedMemory::WorldComm,EXIT_FAILURE);
+}
+static std::string TofuJoin(const std::vector<int> &v,const char *sep)
+{
+  std::ostringstream os;
+  for(int i=0;i<v.size();i++) os << (i ? sep : "") << v[i];
+  return os.str();
+}
+void GlobalSharedMemory::OptimalCommunicatorTofu(const Coordinate &processors,Grid_MPI_Comm & optimal_comm,Coordinate &SHM)
+{
+  int ndimension = processors.size();
+  Coordinate WorldDims = processors; Coordinate ShmDims(ndimension);  Coordinate NodeDims (ndimension);
+  Coordinate ShmCoor(ndimension);    Coordinate NodeCoor(ndimension);   Coordinate WorldCoor(ndimension);
+
+  GetShmDims(WorldDims,ShmDims);
+  SHM=ShmDims;
+  int nodes=1;
+  for(int d=0;d<ndimension;d++){
+    NodeDims[d] = WorldDims[d]/ShmDims[d];
+    nodes *= NodeDims[d];
+  }
+  int Nprocessors=1;
+  for(int i=0;i<ndimension;i++){
+    Nprocessors*=processors[i];
+  }
+  GRID_ASSERT(WorldSize==Nprocessors);
+
+  std::vector<int> pvec(processors.begin(),processors.end());
+  std::vector<int> nvec(NodeDims.begin(),NodeDims.end());
+  std::vector<int> svec(ShmDims.begin(),ShmDims.end());
+  std::string layout = "--mpi " + TofuJoin(pvec,".") + " gives a node grid " + TofuJoin(nvec,".")
+    + " (" + std::to_string(nodes) + " nodes, ranks per node " + TofuJoin(svec,".") + ")";
+
+  ////////////////////////////////////////////////////////////////
+  // Shape of the job's logical node torus and this node's coordinate in it
+  ////////////////////////////////////////////////////////////////
+  int jdim=0;
+  if ( FJMPI_Topology_get_dimension(&jdim)!=FJMPI_SUCCESS || jdim<1 || jdim>3 ) {
+    TofuPlacementError("cannot read the job's node shape (FJMPI_Topology_get_dimension); "
+                       "is this a Fujitsu MPI job on whole nodes?");
+  }
+  std::vector<int> jshape(3,1), jcoor(3,0);
+  if ( FJMPI_Topology_get_shape(&jshape[0],&jshape[1],&jshape[2])!=FJMPI_SUCCESS ) {
+    TofuPlacementError("cannot read the job's node shape (FJMPI_Topology_get_shape)");
+  }
+  for(int a=jdim;a<3;a++) jshape[a]=1;
+  if ( FJMPI_Topology_get_coords(WorldComm,WorldRank,FJMPI_LOGICAL,jdim,&jcoor[0])!=FJMPI_SUCCESS ) {
+    TofuPlacementError("cannot read the node coordinates (FJMPI_Topology_get_coords)");
+  }
+  std::vector<int> jvec(jshape.begin(),jshape.begin()+jdim);
+  int jnodes = jshape[0]*jshape[1]*jshape[2];
+
+  ////////////////////////////////////////////////////////////////
+  // Lay each node grid direction of extent > 1 along an unused job axis of the same extent
+  ////////////////////////////////////////////////////////////////
+  std::vector<int> axis(ndimension,-1);
+  std::vector<int> used(3,0);
+  bool solvable = (nodes==WorldNodes) && (jnodes==WorldNodes);
+  for(int d=0;d<ndimension && solvable;d++){
+    if ( NodeDims[d]==1 ) continue;
+    for(int a=0;a<3;a++){
+      if ( !used[a] && jshape[a]==NodeDims[d] ) { axis[d]=a; used[a]=1; break; }
+    }
+    if ( axis[d]==-1 ) solvable=false;
+  }
+  for(int a=0;a<3;a++) if ( jshape[a]>1 && !used[a] ) solvable=false;
+  if ( !solvable ) {
+    std::vector<int> want;
+    for(int d=0;d<ndimension;d++) if ( NodeDims[d]>1 ) want.push_back(NodeDims[d]);
+    std::ostringstream os;
+    os << layout << ", which cannot be laid onto the job's node shape " << TofuJoin(jvec,"x")
+       << " (" << jnodes << " nodes, " << WorldNodes << " in use): every node grid extent larger than 1"
+       << " must equal a different extent of the job shape. ";
+    if ( want.size()>=1 && want.size()<=3 ) {
+      while ( want.size()<3 ) want.push_back(1);
+      os << "Request a job shape with the extents " << TofuJoin(want,", ") << " in any order"
+         << " (pjsub -L \"node=" << TofuJoin(want,"x") << ":torus\"), or change --mpi.";
+    } else {
+      os << "At most 3 node grid directions may have more than one node; change --mpi.";
+    }
+    TofuPlacementError(os.str());
+  }
+
+  ////////////////////////////////////////////////////////////////
+  // Processor coordinate: node block from the torus, on-node block as in the default layout
+  ////////////////////////////////////////////////////////////////
+  for(int d=0;d<ndimension;d++) NodeCoor[d] = (axis[d]>=0) ? jcoor[axis[d]] : 0;
+  Lexicographic::CoorFromIndexReversed(ShmCoor ,WorldShmRank,ShmDims);
+  for(int d=0;d<ndimension;d++) WorldCoor[d] = NodeCoor[d]*ShmDims[d]+ShmCoor[d];
+  int rank;
+  Lexicographic::IndexFromCoorReversed(WorldCoor,rank,WorldDims);
+
+  ////////////////////////////////////////////////////////////////
+  // Every processor coordinate must be taken exactly once
+  ////////////////////////////////////////////////////////////////
+  std::vector<int> ranks(WorldSize);
+  MPI_Allgather(&rank,1,MPI_INT,&ranks[0],1,MPI_INT,WorldComm);
+  std::vector<int> owner(WorldSize,-1);
+  for(int r=0;r<WorldSize;r++){
+    int k = ranks[r];
+    if ( owner[k]!=-1 ) {
+      std::ostringstream os;
+      os << layout << ": world ranks " << owner[k] << " and " << r << " get the same position; the job's"
+         << " logical node coordinates are not one per node, or the ranks per node differ from " << WorldShmSize;
+      TofuPlacementError(os.str());
+    }
+    owner[k]=r;
+  }
+
+  ////////////////////////////////////////////////////////////////
+  // Log the placement once per processor grid: summary, then (Debug) one line per rank
+  // and the Tofu hop counts between neighbouring ranks on different nodes
+  ////////////////////////////////////////////////////////////////
+  static std::vector<Coordinate> logged;
+  if ( std::find(logged.begin(),logged.end(),processors)==logged.end() ) {
+    logged.push_back(processors);
+    struct TofuRankInfo {
+      int tofu[6], logical[3];
+      char host[32];
+    } me;
+    for(int i=0;i<6;i++) me.tofu[i]=-1;
+    FJMPI_Topology_get_coords(WorldComm,WorldRank,FJMPI_TOFU_SYS,6,me.tofu);
+    for(int i=0;i<3;i++) me.logical[i]=jcoor[i];
+    memset(me.host,0,sizeof(me.host));
+    gethostname(me.host,sizeof(me.host)-1);
+    std::vector<TofuRankInfo> all(WorldSize);
+    MPI_Allgather(&me,sizeof(me),MPI_BYTE,&all[0],sizeof(me),MPI_BYTE,WorldComm);
+
+    std::vector<std::string> along;
+    for(int d=0;d<ndimension;d++){
+      if ( axis[d]>=0 ) along.push_back("dimension " + std::to_string(d) + " along job axis " + std::to_string(axis[d]));
+    }
+    std::ostringstream os;
+    for(int i=0;i<along.size();i++) os << (i ? ", " : "") << along[i];
+    std::cout << GridLogMessage << "Tofu placement: " << layout << ", job shape " << TofuJoin(jvec,"x")
+              << ": " << os.str() << std::endl;
+
+    // Tofu hops: sum over X,Y,Z,A,B,C, B a ring of 3; approximate, X,Y,Z are taken without wrap
+    auto hops = [](const TofuRankInfo &p,const TofuRankInfo &q) {
+      int h=0;
+      for(int i=0;i<6;i++){
+        int dd = std::abs(p.tofu[i]-q.tofu[i]);
+        if ( i==4 ) dd = std::min(dd,3-dd);
+        h += dd;
+      }
+      return h;
+    };
+    for(int k=0;k<WorldSize;k++){
+      const TofuRankInfo &r = all[owner[k]];
+      Coordinate c;
+      Lexicographic::CoorFromIndexReversed(c,k,WorldDims);
+      std::vector<int> cv(c.begin(),c.end());
+      char line[256];
+      snprintf(line,sizeof(line),"rank %5d coor %-14s world %5d host %-16s tofu %2d %2d %2d %d %d %d logical %2d %2d %2d",
+               k,TofuJoin(cv,".").c_str(),owner[k],r.host,r.tofu[0],r.tofu[1],r.tofu[2],r.tofu[3],r.tofu[4],r.tofu[5],
+               r.logical[0],r.logical[1],r.logical[2]);
+      std::cout << GridLogDebug << "Tofu placement: " << line << std::endl;
+    }
+    for(int d=0;d<ndimension;d++){
+      if ( WorldDims[d]==1 ) continue;
+      std::map<int,int> hist;
+      int onnode=0;
+      for(int k=0;k<WorldSize;k++){
+        Coordinate c;
+        Lexicographic::CoorFromIndexReversed(c,k,WorldDims);
+        c[d] = (c[d]+1)%WorldDims[d];
+        int n;
+        Lexicographic::IndexFromCoorReversed(c,n,WorldDims);
+        const TofuRankInfo &p = all[owner[k]], &q = all[owner[n]];
+        if ( strncmp(p.host,q.host,sizeof(p.host))==0 ) onnode++;
+        else hist[hops(p,q)]++;
+      }
+      std::ostringstream hs;
+      for(auto &h : hist) hs << " " << h.second << " at " << h.first << (h.first==1 ? " hop;" : " hops;");
+      std::cout << GridLogDebug << "Tofu placement: dimension " << d << " (" << WorldDims[d] << " ranks): "
+                << onnode << " of " << WorldSize << " +1 neighbours on the same node;" << hs.str()
+                << " (hops approximate, X/Y/Z not wrapped)" << std::endl;
+    }
+  }
+
+  /////////////////////////////////////////////////////////////////
+  // Build the new communicator
+  /////////////////////////////////////////////////////////////////
+  int ierr= MPI_Comm_split(WorldComm,0,rank,&optimal_comm);
+  GRID_ASSERT(ierr==0);
+}
+#endif
 ////////////////////////////////////////////////////////////////////////////////////////////
 // SHMGET
 ////////////////////////////////////////////////////////////////////////////////////////////

Grid/communicator/SharedMemoryNone.cc

@@ -50,6 +50,10 @@ void GlobalSharedMemory::Init(Grid_MPI_Comm comm)
 
 void GlobalSharedMemory::OptimalCommunicator(const Coordinate &processors,Grid_MPI_Comm & optimal_comm,Coordinate &SHM)
 {
+  if ( TofuPlacement ) {
+    std::cerr << "Grid --tofu: using a comms=none build" << std::endl;
+    exit(EXIT_FAILURE);
+  }
   optimal_comm = WorldComm;
   SHM = Coordinate(processors.size(),1);
 }

Grid/qcd/action/fermion/WilsonCompressor.h

@@ -253,6 +253,11 @@ public:
 
     int dag = compress.dag;
     int face_idx=0;
+#ifdef GRID_FUSED_HALO_LOOPS
+    // queue the face gathers and run them together below (not on the first call, which builds the face tables)
+    const bool fused = this->face_table_computed;
+    if ( fused ) FusedHaloLoops::begin();
+#endif
 #define vet_same_node(a,b) \
       { auto tmp = b;  }
     if ( dag ) { 
@@ -274,6 +279,9 @@ public:
       vet_same_node(this->same_node[Zm],this->HaloGatherDir(source,ZpCompress,Zm,face_idx));
       vet_same_node(this->same_node[Tm],this->HaloGatherDir(source,TpCompress,Tm,face_idx));
     }
+#ifdef GRID_FUSED_HALO_LOOPS
+    if ( fused ) FusedHaloLoops::end();
+#endif
     this->face_table_computed=1;
     GRID_ASSERT(this->u_comm_offset==this->_unified_buffer_size);
     accelerator_barrier();

Grid/qcd/action/fermion/WilsonKernels.h

@@ -42,6 +42,12 @@ public:
   enum { CommsAndCompute, CommsThenCompute };
   static int Opt;  
   static int Comms;
+#ifdef A64FX
+  // If set, called by thread 0 every few sites of the interior kernel so that
+  // communications started before it progress during the compute.
+  static void (*CommsProgress)(void *);
+  static void *CommsProgressArg;
+#endif
 };
  
 template<class Impl> class WilsonKernels : public FermionOperator<Impl> , public WilsonKernelsStatic { 

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

@@ -346,6 +346,11 @@ void WilsonFermion5D<Impl>::DhopInternalOverlappedComms(StencilImpl & st,
   // do the compute interior
   /////////////////////////////
   int Opt = WilsonKernelsStatic::Opt; // Why pass this. Kernels should know
+#ifdef A64FX
+  // MPI does not progress outside MPI calls: let the interior kernel poll
+  WilsonKernelsStatic::CommsProgress    = [](void *s) { ((StencilImpl *)s)->CommunicateProgress(); };
+  WilsonKernelsStatic::CommsProgressArg = (void *)&st;
+#endif
   if (dag == DaggerYes) {
     GRID_TRACE("DhopDagInterior");
     Kernels::DhopDagKernel(Opt,st,U,st.CommBuf(),LLs,U.oSites(),in,out,1,0);
@@ -353,6 +358,10 @@ void WilsonFermion5D<Impl>::DhopInternalOverlappedComms(StencilImpl & st,
     GRID_TRACE("DhopInterior");
     Kernels::DhopKernel   (Opt,st,U,st.CommBuf(),LLs,U.oSites(),in,out,1,0);
   }
+#ifdef A64FX
+  WilsonKernelsStatic::CommsProgress    = nullptr;
+  WilsonKernelsStatic::CommsProgressArg = nullptr;
+#endif
   
   //ifdef GRID_ACCELERATED
 #if 0

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

@@ -436,6 +436,11 @@ void WilsonFermion<Impl>::DhopInternalOverlappedComms(StencilImpl &st,
   // do the compute interior
   /////////////////////////////
   int Opt = WilsonKernelsStatic::Opt;
+#ifdef A64FX
+  // MPI does not progress outside MPI calls: let the interior kernel poll
+  WilsonKernelsStatic::CommsProgress    = [](void *s) { ((StencilImpl *)s)->CommunicateProgress(); };
+  WilsonKernelsStatic::CommsProgressArg = (void *)&st;
+#endif
   if (dag == DaggerYes) {
     GRID_TRACE("DhopDagInterior");
     Kernels::DhopDagKernel(Opt,st,U,st.CommBuf(),1,U.oSites(),in,out,1,0);
@@ -443,6 +448,10 @@ void WilsonFermion<Impl>::DhopInternalOverlappedComms(StencilImpl &st,
     GRID_TRACE("DhopInterior");
     Kernels::DhopKernel(Opt,st,U,st.CommBuf(),1,U.oSites(),in,out,1,0);
   }
+#ifdef A64FX
+  WilsonKernelsStatic::CommsProgress    = nullptr;
+  WilsonKernelsStatic::CommsProgressArg = nullptr;
+#endif
 
   /////////////////////////////
   // Complete comms

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

@@ -95,7 +95,6 @@ Author:  Nils Meyer  <nils.meyer@ur.de>  Regensburg University
 #ifdef INTERIOR_AND_EXTERIOR
 
 #define ASM_LEG(Dir,NxtDir,PERMUTE_DIR,PROJ,RECON)			\
-      basep = st.GetPFInfo(nent,plocal); nent++;			\
       if ( local ) {							            \
     LOAD_CHIMU(base);                                       \
     LOAD_TABLE(PERMUTE_DIR);                                \
@@ -106,23 +105,15 @@ Author:  Nils Meyer  <nils.meyer@ur.de>  Regensburg University
       }									                    \
       base = st.GetInfo(ptype,local,perm,NxtDir,ent,plocal); ent++;	\
     MULT_2SPIN_1(Dir);					                    \
-    PREFETCH_CHIMU(base);                                   \
-    PREFETCH_CHIMU_L2(basep);                               \
-    /* PREFETCH_GAUGE_L1(NxtDir); */                        \
     MULT_2SPIN_2;					                        \
-    if (s == 0) {                                           \
-      if ((Dir == 0) || (Dir == 4)) { PREFETCH_GAUGE_L2(Dir); } \
-    }                                                       \
     RECON;								                    \
 
 /*
-NB: picking PREFETCH_GAUGE_L2(Dir+4); here results in performance penalty
-    though I expected that it would improve on performance
+No software prefetch: the hardware prefetcher handles the neighbour and gauge link streams better on A64FX.
 */
 
 #define ASM_LEG_XP(Dir,NxtDir,PERMUTE_DIR,PROJ,RECON)	    \
   base = st.GetInfo(ptype,local,perm,Dir,ent,plocal); ent++; \
-  PREFETCH1_CHIMU(base);						            \
   ASM_LEG(Dir,NxtDir,PERMUTE_DIR,PROJ,RECON)
 
 #define RESULT(base,basep) SAVE_RESULT(base,basep);
@@ -130,12 +121,11 @@ NB: picking PREFETCH_GAUGE_L2(Dir+4); here results in performance penalty
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////
-// Pre comms kernel -- prefetch like normal because it is mostly right
+// Pre comms kernel -- gauge link prefetch into L2 only
 ////////////////////////////////////////////////////////////////////////////////
 #ifdef INTERIOR
 
 #define ASM_LEG(Dir,NxtDir,PERMUTE_DIR,PROJ,RECON)			\
-      basep = st.GetPFInfo(nent,plocal); nent++;			\
       if ( local ) {							\
   LOAD_CHIMU(base);                                       \
   LOAD_TABLE(PERMUTE_DIR);                                \
@@ -148,12 +138,15 @@ NB: picking PREFETCH_GAUGE_L2(Dir+4); here results in performance penalty
   RECON;								\
       }									\
   base = st.GetInfo(ptype,local,perm,NxtDir,ent,plocal); ent++;	\
-  PREFETCH_CHIMU(base);						\
-  PREFETCH_CHIMU_L2(basep);                               \
+  if (s == 0) {                                           \
+    if ((Dir == 0) || (Dir == 4)) { PREFETCH_GAUGE_L2(Dir); } \
+  }                                                       \
+
+/* No software prefetch of the neighbours (the hardware prefetcher handles them better on A64FX);
+   the gauge links of the next site are prefetched into L2. */
 
 #define ASM_LEG_XP(Dir,NxtDir,PERMUTE_DIR,PROJ,RECON)			\
   base = st.GetInfo(ptype,local,perm,Dir,ent,plocal); ent++;		\
-  PREFETCH1_CHIMU(base);						\
   { ZERO_PSI; }								\
   ASM_LEG(Dir,NxtDir,PERMUTE_DIR,PROJ,RECON)
 

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

@@ -479,6 +479,49 @@ extern "C" {
     int sF = ss*Ls;							\
     WilsonKernels<Impl>::A(st_v,U_v,buf,sF,sU,Ls,1,in_v,out_v);		\
   });
+#ifdef A64FX
+// Interior kernel with communication progress: thread 0 calls the progress hook
+// between chunks of sites; dynamic scheduling lets the other threads pick up the
+// sites thread 0 has no time for while it drives the communications. Chunks of
+// 8 4d sites for 5d operators, 32 sites for 4d operators (measured on Fugaku).
+#define ASM_CALL_PROGRESS(A)						\
+  if ( WilsonKernelsStatic::CommsProgress ) {				\
+    auto progress     = WilsonKernelsStatic::CommsProgress;		\
+    auto progress_arg = WilsonKernelsStatic::CommsProgressArg;		\
+    const uint64_t chunk = std::max<uint64_t>(8,32/Ls);			\
+    const uint64_t nchunk = (Nsite + chunk - 1)/chunk;			\
+    thread_region {							\
+      const bool master = (thread_num() == 0);				\
+      DO_PRAGMA(omp for schedule(dynamic,1) nowait)			\
+      for(uint64_t c=0;c<nchunk;c++){					\
+	const uint64_t hi = std::min((c+1)*chunk,(uint64_t)Nsite);	\
+	for(uint64_t sss=c*chunk;sss<hi;sss++){				\
+	  int sU = sss;							\
+	  int sF = sss*Ls;						\
+	  WilsonKernels<Impl>::A(st_v,U_v,buf,sF,sU,Ls,1,in_v,out_v);	\
+	}								\
+	if ( master ) progress(progress_arg);				\
+      }									\
+    }									\
+  } else {								\
+    ASM_CALL(A);							\
+  }
+
+// Exterior kernel on the surface 4d sites only (the surface list holds Ls
+// consecutive 5d entries per surface 4d site)
+#define ASM_CALL_EXT(A)							\
+  const uint64_t nsurf = st.surface_list.size()/Ls;			\
+  auto surf = &st.surface_list[0];					\
+  thread_for( sss, nsurf, {						\
+    int sU = surf[sss*Ls]/Ls;						\
+    int sF = sU*Ls;							\
+    WilsonKernels<Impl>::A(st_v,U_v,buf,sF,sU,Ls,1,in_v,out_v);		\
+  });
+#else
+#define ASM_CALL_PROGRESS(A) ASM_CALL(A)
+#define ASM_CALL_EXT(A)      ASM_CALL(A)
+#endif
+
 #define ASM_CALL_SLICE(A)						\
   auto grid = in.Grid() ;						\
   int nt = grid->LocalDimensions()[4];					\
@@ -514,7 +557,7 @@ void WilsonKernels<Impl>::DhopKernel(int Opt,StencilImpl &st,  DoubledGaugeField
      if (Opt == WilsonKernelsStatic::OptGeneric    ) { KERNEL_CALLNB(GenericDhopSiteInt); return;}
      if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALLNB(HandDhopSiteInt);    return;}
 #ifndef GRID_CUDA
-     if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL(AsmDhopSiteInt);    return;}
+     if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL_PROGRESS(AsmDhopSiteInt);    return;}
 #endif
    } else if( exterior ) {
      //     // dependent on result of merge
@@ -522,7 +565,7 @@ void WilsonKernels<Impl>::DhopKernel(int Opt,StencilImpl &st,  DoubledGaugeField
      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;}
+     if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL_EXT(AsmDhopSiteExt);    return;}
 #endif
    }
    GRID_ASSERT(0 && " Kernel optimisation case not covered ");
@@ -560,7 +603,7 @@ void WilsonKernels<Impl>::DhopKernel(int Opt,StencilImpl &st,  DoubledGaugeField
      if (Opt == WilsonKernelsStatic::OptGeneric    ) { KERNEL_CALLNB(GenericDhopSiteDagInt); return;}
      if (Opt == WilsonKernelsStatic::OptHandUnroll ) { KERNEL_CALLNB(HandDhopSiteDagInt);    return;}
 #ifndef GRID_CUDA
-     if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL(AsmDhopSiteDagInt);     return;}
+     if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL_PROGRESS(AsmDhopSiteDagInt);     return;}
 #endif
    } else if( exterior ) {
      // Dependent on result of merge
@@ -568,7 +611,7 @@ void WilsonKernels<Impl>::DhopKernel(int Opt,StencilImpl &st,  DoubledGaugeField
      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;}
+     if (Opt == WilsonKernelsStatic::OptInlineAsm  ) {  ASM_CALL_EXT(AsmDhopSiteDagExt);     return;}
 #endif
    }
    GRID_ASSERT(0 && " Kernel optimisation case not covered ");

Grid/qcd/action/fermion/instantiation/WilsonKernelsInstantiation.cc

@@ -34,6 +34,10 @@ NAMESPACE_BEGIN(Grid);
 
 // Move these
 int WilsonKernelsStatic::Opt   = WilsonKernelsStatic::OptGeneric;
+#ifdef A64FX
+void (*WilsonKernelsStatic::CommsProgress)(void *) = nullptr;
+void *WilsonKernelsStatic::CommsProgressArg = nullptr;
+#endif
 int WilsonKernelsStatic::Comms = WilsonKernelsStatic::CommsAndCompute;
 
 NAMESPACE_END(Grid);

Grid/simd/Fujitsu_A64FX_intrin_double.h

@@ -39,8 +39,13 @@ Author: Nils Meyer <nils.meyer@ur.de>
 #define UNLOCK_GAUGE(A)  
 #define MASK_REGS                      DECLARATIONS_A64FXd  
 #define SAVE_RESULT(A,B)               RESULT_A64FXd(A);  
-#define MULT_2SPIN_1(Dir)              MULT_2SPIN_1_A64FXd(Dir)  
-#define MULT_2SPIN_2                   MULT_2SPIN_2_A64FXd  
+#ifdef GRID_A64FX_FCMLA
+#define MULT_2SPIN_1(Dir)              MULT_2SPIN_1_FCMLA_A64FXd(Dir)
+#define MULT_2SPIN_2                   MULT_2SPIN_2_FCMLA_A64FXd
+#else
+#define MULT_2SPIN_1(Dir)              MULT_2SPIN_1_FMLA_A64FXd(Dir)
+#define MULT_2SPIN_2                   MULT_2SPIN_2_FMLA_A64FXd
+#endif
 #define LOAD_CHI(base)                 LOAD_CHI_A64FXd(base)  
 #define ZERO_PSI                       ZERO_PSI_A64FXd  
 #define ADD_RESULT(base,basep)         LOAD_CHIMU(base); ADD_RESULT_INTERNAL_A64FXd; RESULT_A64FXd(base)  
@@ -273,8 +278,88 @@ Author: Nils Meyer <nils.meyer@ur.de>
     U_11 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(-2));  \
     U_21 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(1));  \
 }
+// MULT_2SPIN with FMLA instead of FCMLA (3 uops on A64FX):
+//   UChi = sum_b Re(U_ab) Chi_b + Im(U_ab) (i Chi_b)
+// column by column. Re(U)/Im(U) are duplicated over each complex with trn1/trn2,
+// i Chi_b is formed in place with fcadd once Chi_b has been used (Chi is dead after
+// the multiply). U_00, U_10, U_20 hold column b, U_01 is a temporary.
+// Define GRID_A64FX_FCMLA to use the original FCMLA version.
+#define MULT_2SPIN_1_FMLA_A64FXd(A)  \
+{ \
+    const auto & ref(U[sU](A)); baseU = (uint64_t)&ref; \
+    U_00 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(-6));  \
+    U_10 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(-3));  \
+    U_20 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(0));  \
+    U_01 = svtrn1(U_00, U_00); \
+    UChi_00 = svmul_x(pg1, U_01, Chi_00); \
+    UChi_10 = svmul_x(pg1, U_01, Chi_10); \
+    U_01 = svtrn1(U_10, U_10); \
+    UChi_01 = svmul_x(pg1, U_01, Chi_00); \
+    UChi_11 = svmul_x(pg1, U_01, Chi_10); \
+    U_01 = svtrn1(U_20, U_20); \
+    UChi_02 = svmul_x(pg1, U_01, Chi_00); \
+    UChi_12 = svmul_x(pg1, U_01, Chi_10); \
+    Chi_00 = svcadd_x(pg1, zero0, Chi_00, 90); \
+    Chi_10 = svcadd_x(pg1, zero0, Chi_10, 90); \
+    U_01 = svtrn2(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_00); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_10); \
+    U_01 = svtrn2(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_00); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_10); \
+    U_01 = svtrn2(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_00); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_10); \
+    U_00 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(-5));  \
+    U_10 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(-2));  \
+    U_20 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(1));  \
+    U_01 = svtrn1(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_01); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_11); \
+    U_01 = svtrn1(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_01); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_11); \
+    U_01 = svtrn1(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_01); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_11); \
+    Chi_01 = svcadd_x(pg1, zero0, Chi_01, 90); \
+    Chi_11 = svcadd_x(pg1, zero0, Chi_11, 90); \
+    U_01 = svtrn2(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_01); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_11); \
+    U_01 = svtrn2(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_01); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_11); \
+    U_01 = svtrn2(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_01); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_11); \
+    U_00 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(-4));  \
+    U_10 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(-1));  \
+    U_20 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(2));  \
+    U_01 = svtrn1(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_02); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_12); \
+    U_01 = svtrn1(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_02); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_12); \
+    U_01 = svtrn1(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_02); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_12); \
+    Chi_02 = svcadd_x(pg1, zero0, Chi_02, 90); \
+    Chi_12 = svcadd_x(pg1, zero0, Chi_12, 90); \
+    U_01 = svtrn2(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_02); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_12); \
+    U_01 = svtrn2(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_02); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_12); \
+    U_01 = svtrn2(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_02); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_12); \
+}
+#define MULT_2SPIN_2_FMLA_A64FXd
 // MULT_2SPIN
-#define MULT_2SPIN_1_A64FXd(A)  \
+#define MULT_2SPIN_1_FCMLA_A64FXd(A)  \
 { \
     const auto & ref(U[sU](A)); baseU = (uint64_t)&ref; \
     U_00 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(-6));  \
@@ -300,7 +385,7 @@ Author: Nils Meyer <nils.meyer@ur.de>
     U_20 = svld1_vnum(pg1, (float64_t*)(baseU + 2 * 3 * 64), (int64_t)(2));  \
 }
 // MULT_2SPIN_BACKEND
-#define MULT_2SPIN_2_A64FXd  \
+#define MULT_2SPIN_2_FCMLA_A64FXd  \
 { \
     UChi_00 = svcmla_x(pg1, UChi_00, U_01, Chi_01, 0); \
     UChi_10 = svcmla_x(pg1, UChi_10, U_01, Chi_11, 0); \

Grid/simd/Fujitsu_A64FX_intrin_single.h

@@ -39,8 +39,13 @@ Author: Nils Meyer <nils.meyer@ur.de>
 #define UNLOCK_GAUGE(A)  
 #define MASK_REGS                      DECLARATIONS_A64FXf  
 #define SAVE_RESULT(A,B)               RESULT_A64FXf(A);  
-#define MULT_2SPIN_1(Dir)              MULT_2SPIN_1_A64FXf(Dir)  
-#define MULT_2SPIN_2                   MULT_2SPIN_2_A64FXf  
+#ifdef GRID_A64FX_FCMLA
+#define MULT_2SPIN_1(Dir)              MULT_2SPIN_1_FCMLA_A64FXf(Dir)
+#define MULT_2SPIN_2                   MULT_2SPIN_2_FCMLA_A64FXf
+#else
+#define MULT_2SPIN_1(Dir)              MULT_2SPIN_1_FMLA_A64FXf(Dir)
+#define MULT_2SPIN_2                   MULT_2SPIN_2_FMLA_A64FXf
+#endif
 #define LOAD_CHI(base)                 LOAD_CHI_A64FXf(base)  
 #define ZERO_PSI                       ZERO_PSI_A64FXf  
 #define ADD_RESULT(base,basep)         LOAD_CHIMU(base); ADD_RESULT_INTERNAL_A64FXf; RESULT_A64FXf(base)  
@@ -273,8 +278,88 @@ Author: Nils Meyer <nils.meyer@ur.de>
     U_11 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(-2));  \
     U_21 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(1));  \
 }
+// MULT_2SPIN with FMLA instead of FCMLA (3 uops on A64FX):
+//   UChi = sum_b Re(U_ab) Chi_b + Im(U_ab) (i Chi_b)
+// column by column. Re(U)/Im(U) are duplicated over each complex with trn1/trn2,
+// i Chi_b is formed in place with fcadd once Chi_b has been used (Chi is dead after
+// the multiply). U_00, U_10, U_20 hold column b, U_01 is a temporary.
+// Define GRID_A64FX_FCMLA to use the original FCMLA version.
+#define MULT_2SPIN_1_FMLA_A64FXf(A)  \
+{ \
+    const auto & ref(U[sU](A)); baseU = (uint64_t)&ref; \
+    U_00 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(-6));  \
+    U_10 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(-3));  \
+    U_20 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(0));  \
+    U_01 = svtrn1(U_00, U_00); \
+    UChi_00 = svmul_x(pg1, U_01, Chi_00); \
+    UChi_10 = svmul_x(pg1, U_01, Chi_10); \
+    U_01 = svtrn1(U_10, U_10); \
+    UChi_01 = svmul_x(pg1, U_01, Chi_00); \
+    UChi_11 = svmul_x(pg1, U_01, Chi_10); \
+    U_01 = svtrn1(U_20, U_20); \
+    UChi_02 = svmul_x(pg1, U_01, Chi_00); \
+    UChi_12 = svmul_x(pg1, U_01, Chi_10); \
+    Chi_00 = svcadd_x(pg1, zero0, Chi_00, 90); \
+    Chi_10 = svcadd_x(pg1, zero0, Chi_10, 90); \
+    U_01 = svtrn2(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_00); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_10); \
+    U_01 = svtrn2(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_00); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_10); \
+    U_01 = svtrn2(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_00); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_10); \
+    U_00 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(-5));  \
+    U_10 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(-2));  \
+    U_20 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(1));  \
+    U_01 = svtrn1(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_01); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_11); \
+    U_01 = svtrn1(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_01); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_11); \
+    U_01 = svtrn1(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_01); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_11); \
+    Chi_01 = svcadd_x(pg1, zero0, Chi_01, 90); \
+    Chi_11 = svcadd_x(pg1, zero0, Chi_11, 90); \
+    U_01 = svtrn2(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_01); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_11); \
+    U_01 = svtrn2(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_01); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_11); \
+    U_01 = svtrn2(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_01); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_11); \
+    U_00 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(-4));  \
+    U_10 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(-1));  \
+    U_20 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(2));  \
+    U_01 = svtrn1(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_02); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_12); \
+    U_01 = svtrn1(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_02); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_12); \
+    U_01 = svtrn1(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_02); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_12); \
+    Chi_02 = svcadd_x(pg1, zero0, Chi_02, 90); \
+    Chi_12 = svcadd_x(pg1, zero0, Chi_12, 90); \
+    U_01 = svtrn2(U_00, U_00); \
+    UChi_00 = svmla_x(pg1, UChi_00, U_01, Chi_02); \
+    UChi_10 = svmla_x(pg1, UChi_10, U_01, Chi_12); \
+    U_01 = svtrn2(U_10, U_10); \
+    UChi_01 = svmla_x(pg1, UChi_01, U_01, Chi_02); \
+    UChi_11 = svmla_x(pg1, UChi_11, U_01, Chi_12); \
+    U_01 = svtrn2(U_20, U_20); \
+    UChi_02 = svmla_x(pg1, UChi_02, U_01, Chi_02); \
+    UChi_12 = svmla_x(pg1, UChi_12, U_01, Chi_12); \
+}
+#define MULT_2SPIN_2_FMLA_A64FXf
 // MULT_2SPIN
-#define MULT_2SPIN_1_A64FXf(A)  \
+#define MULT_2SPIN_1_FCMLA_A64FXf(A)  \
 { \
     const auto & ref(U[sU](A)); baseU = (uint64_t)&ref; \
     U_00 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(-6));  \
@@ -300,7 +385,7 @@ Author: Nils Meyer <nils.meyer@ur.de>
     U_20 = svld1_vnum(pg1, (float32_t*)(baseU + 2 * 3 * 64), (int64_t)(2));  \
 }
 // MULT_2SPIN_BACKEND
-#define MULT_2SPIN_2_A64FXf  \
+#define MULT_2SPIN_2_FCMLA_A64FXf  \
 { \
     UChi_00 = svcmla_x(pg1, UChi_00, U_01, Chi_01, 0); \
     UChi_10 = svcmla_x(pg1, UChi_10, U_01, Chi_11, 0); \

Grid/stencil/SimpleCompressor.h

@@ -11,6 +11,43 @@ public:
 };
 
 
+// On A64FX (OpenMP CPU builds) the face gathers of one halo exchange, and the merges, run in one parallel
+// region: each loop is queued, then every thread takes its static share of each loop, with one barrier at
+// the end instead of one parallel region per face. The gathers also prefetch the source sites 4 entries
+// ahead. CPU only: the queued loops run after the views of the source field are closed, which is safe
+// only because a CPU view is a plain pointer to host memory.
+#if defined(A64FX) && defined(_OPENMP) && !defined(GRID_SIMT)
+#define GRID_FUSED_HALO_LOOPS
+struct FusedHaloLoops {
+  // a queued loop over sites [0,n): body(lo,hi) runs sites lo..hi-1
+  typedef std::pair<uint64_t, std::function<void(uint64_t,uint64_t)> > Loop;
+  static std::vector<Loop> &queue(void) { static std::vector<Loop> q; return q; }
+  static bool &deferring(void) { static bool d = false; return d; }
+  // between begin() and end(), the face loops are queued (push) instead of run; end() runs them all
+  static void begin(void) { deferring() = true; }
+  static void end(void)   { deferring() = false; run(); }
+  static void push(uint64_t n,std::function<void(uint64_t,uint64_t)> body) { queue().push_back({n,body}); }
+  static const int prefetch_distance = 4;
+  template<class T> static inline void prefetch(const T *p) {
+    for(int o=0;o<(int)sizeof(T);o+=256) __builtin_prefetch((const char *)p+o,0,3); // 256-byte lines
+  }
+  static void run(void) {
+    auto &q = queue();
+    if ( q.empty() ) return;
+    // one chunk per thread, as in the lattice reductions (Lattice_reduction.h)
+    const int nthread = GridThread::GetThreads();
+    thread_for(thr,nthread,{
+      for(auto &loop : q) {
+	int mywork, myoff;
+	GridThread::GetWork((int)loop.first,thr,mywork,myoff);
+	if ( mywork > 0 ) loop.second(myoff,myoff+mywork);
+      }
+    });
+    q.clear();
+  }
+};
+#endif
+
 // Compressors will inherit buffer management policies
 // Standard comms buffer management
 class FaceGatherSimple
@@ -29,6 +66,20 @@ public:
     std::pair<int,int> *table_v = & table[0];
     
     auto rhs_v = rhs.View(AcceleratorRead);
+#ifdef GRID_FUSED_HALO_LOOPS
+    if ( FusedHaloLoops::deferring() ) {
+      const vobj *rp = &rhs_v[0];
+      const int pf = FusedHaloLoops::prefetch_distance;
+      FusedHaloLoops::push((uint64_t)num, [=](uint64_t lo,uint64_t hi) {
+	for(uint64_t i=lo;i<hi;i++){
+	  if ( i+pf<hi ) FusedHaloLoops::prefetch(&rp[so+table_v[i+pf].second]);
+	  compress.Compress(buffer[off+table_v[i].first],rp[so+table_v[i].second]);
+	}
+      });
+      rhs_v.ViewClose();
+      return;
+    }
+#endif
     accelerator_forNB( i,num, vobj::Nsimd(), {
 	compress.Compress(buffer[off+table_v[i].first],rhs_v[so+table_v[i].second]);
     });
@@ -48,6 +99,23 @@ public:
     auto p1=&pointers[1][0];
     auto tp=&table[0];
     auto rhs_p = &rhs_v[0];
+#ifdef GRID_FUSED_HALO_LOOPS
+    if ( FusedHaloLoops::deferring() ) {
+      const vobj *rp = rhs_p;
+      const int pf = FusedHaloLoops::prefetch_distance;
+      FusedHaloLoops::push((uint64_t)num, [=](uint64_t lo,uint64_t hi) {
+	for(uint64_t j=lo;j<hi;j++){
+	  if ( j+pf<hi ) {
+	    FusedHaloLoops::prefetch(&rp[so+tp[2*(j+pf)].second]);
+	    FusedHaloLoops::prefetch(&rp[so+tp[2*(j+pf)+1].second]);
+	  }
+	  compress.CompressExchange(p0[j],p1[j],rp[so+tp[2*j].second],rp[so+tp[2*j+1].second],type);
+	}
+      });
+      rhs_v.ViewClose();
+      return;
+    }
+#endif
     accelerator_forNB(j, num, vobj::Nsimd(), {
 	compress.CompressExchange(p0[j],p1[j],
 				  rhs_p[so+tp[2*j  ].second],
@@ -73,6 +141,14 @@ public:
     auto vp0= &mm.vpointers[0][0];
     auto vp1= &mm.vpointers[1][0];
     auto type= mm.type;
+#ifdef GRID_FUSED_HALO_LOOPS
+    if ( FusedHaloLoops::deferring() ) {
+      FusedHaloLoops::push((uint64_t)(mm.buffer_size/2), [=](uint64_t lo,uint64_t hi) {
+	for(uint64_t o=lo;o<hi;o++) decompress.Exchange(mp[2*o],mp[2*o+1],vp0[o],vp1[o],type);
+      });
+      return;
+    }
+#endif
     accelerator_forNB(o,mm.buffer_size/2,Merger::Nsimd,{
 	decompress.Exchange(mp[2*o],mp[2*o+1],vp0[o],vp1[o],type);
     });

Grid/stencil/Stencil.h

@@ -630,6 +630,13 @@ public:
     }
   }
 
+#ifdef A64FX
+  // Progress outstanding communications while computing (see CommunicateBegin)
+  void CommunicateProgress(void)
+  {
+    _grid->StencilSendToRecvFromProgress(MpiReqs);
+  }
+#endif
   void CommunicateComplete(std::vector<std::vector<CommsRequest_t> > &reqs)
   {
     //    std::cout << "Communicate Complete "<<std::endl;
@@ -879,6 +886,15 @@ public:
   template<class decompressor>
   void CommsMerge(decompressor decompress,std::vector<Merge> &mm,std::vector<Decompress> &dd)
   {
+#ifdef GRID_FUSED_HALO_LOOPS
+    if ( dd.size()==0 ) { // merges only: run them in one parallel region (SimpleCompressor.h)
+      FusedHaloLoops::begin();
+      for(int i=0;i<mm.size();i++) decompressor::MergeFace(decompress,mm[i]);
+      FusedHaloLoops::end();
+      acceleratorFenceComputeStream();
+      return;
+    }
+#endif
     for(int i=0;i<mm.size();i++){
       decompressor::MergeFace(decompress,mm[i]);
     }

Grid/util/Init.cc

@@ -367,17 +367,14 @@ void Grid_init(int *argc,char ***argv)
     GlobalSharedMemory::HPEhypercube = enable;
   }
 
+  if( GridCmdOptionExists(*argv,*argv+*argc,"--tofu") ){
+    GlobalSharedMemory::TofuPlacement = 1;
+  }
+
   if( GridCmdOptionExists(*argv,*argv+*argc,"--shm-hugepages") ){
     GlobalSharedMemory::Hugepages = 1;
   }
 
-#if defined(A64FX)
-  if( GridCmdOptionExists(*argv,*argv+*argc,"--comms-overlap") ){
-    std::cout << "Option --comms-overlap currently not supported on QPACE4. Exiting." << std::endl;
-    exit(EXIT_FAILURE);
-  }
-#endif
-
   //////////////////////////////////////////////////////////
   // Memory manager
   //////////////////////////////////////////////////////////
@@ -541,6 +538,7 @@ void Grid_init(int *argc,char ***argv)
     std::cout<<GridLogMessage<<"Geometry:"<<std::endl;
     std::cout<<GridLogMessage<<std::endl;
     std::cout<<GridLogMessage<<"  --mpi n.n.n.n   : default MPI decomposition"<<std::endl;    
+    std::cout<<GridLogMessage<<"  --tofu          : place ranks on the job's Tofu torus (Fujitsu MPI only)"<<std::endl;
     std::cout<<GridLogMessage<<"  --threads n     : default number of OMP threads"<<std::endl;
     std::cout<<GridLogMessage<<"  --grid n.n.n.n  : default Grid size"<<std::endl;
     std::cout<<GridLogMessage<<"  --shm  M        : allocate M megabytes of shared memory for comms"<<std::endl;
@@ -595,6 +593,9 @@ void Grid_init(int *argc,char ***argv)
   if( GridCmdOptionExists(*argv,*argv+*argc,"--comms-overlap") ){
     WilsonKernelsStatic::Comms = WilsonKernelsStatic::CommsAndCompute;
     StaggeredKernelsStatic::Comms = StaggeredKernelsStatic::CommsAndCompute;
+#ifdef GRID_FJMPI
+    Stencil_persistent = true;
+#endif
   } else {
     WilsonKernelsStatic::Comms = WilsonKernelsStatic::CommsThenCompute;
     StaggeredKernelsStatic::Comms = StaggeredKernelsStatic::CommsThenCompute;

configure.ac

@@ -846,6 +846,33 @@ esac
 AM_CONDITIONAL(BUILD_COMMS_MPI3,  [ test "${comms_type}X" == "mpi3X" ] )
 AM_CONDITIONAL(BUILD_COMMS_NONE,  [ test "${comms_type}X" == "noneX" ])
 
+############### Fujitsu MPI extensions (persistent halo requests, Tofu topology)
+AC_ARG_ENABLE([fjmpi],[AS_HELP_STRING([--enable-fjmpi=auto|yes|no],[Use the Fujitsu MPI extensions (default: auto, if mpi-ext.h declares them)])],[ac_FJMPI=${enable_fjmpi}],[ac_FJMPI=auto])
+fjmpi=no
+case ${ac_FJMPI} in
+     yes|no|auto) ;;
+     *) AC_MSG_ERROR([${ac_FJMPI} unsupported --enable-fjmpi option]);;
+esac
+if test "${comms_type}X" == "mpi3X" && test "${ac_FJMPI}X" != "noX"; then
+  ac_save_CXXFLAGS=$CXXFLAGS
+  CXXFLAGS="$MPI_CXXFLAGS $CXXFLAGS"
+  AC_MSG_CHECKING([for the Fujitsu MPI extensions])
+  AC_COMPILE_IFELSE([AC_LANG_PROGRAM([[#include <mpi.h>
+#include <mpi-ext.h>]],
+    [[MPI_Request r; int c[3];
+      FJMPI_Prequest_start(&r);
+      FJMPI_Topology_get_coords(MPI_COMM_WORLD,0,FJMPI_LOGICAL,3,c);]])],
+    [fjmpi=yes],[fjmpi=no])
+  AC_MSG_RESULT([${fjmpi}])
+  CXXFLAGS=$ac_save_CXXFLAGS
+fi
+if test "${ac_FJMPI}X" == "yesX" && test "${fjmpi}X" != "yesX"; then
+  AC_MSG_ERROR([--enable-fjmpi=yes but the Fujitsu MPI extensions were not found (needs --enable-comms=mpi and Fujitsu MPI)])
+fi
+if test "${fjmpi}X" == "yesX"; then
+  AC_DEFINE([GRID_FJMPI],[1],[Fujitsu MPI extensions])
+fi
+
 ############### RNG selection
 AC_ARG_ENABLE([rng],[AS_HELP_STRING([--enable-rng=ranlux48|mt19937|sitmo],[\
 	            Select Random Number Generator to be used])],\
@@ -955,6 +982,7 @@ Threading                   : ${ac_openmp}
 Acceleration                : ${ac_ACCELERATOR}
 Unified virtual memory      : ${ac_UNIFIED}
 Communications type         : ${comms_type}
+Fujitsu MPI extensions      : ${fjmpi}
 Shared memory allocator     : ${ac_SHM}
 Shared memory mmap path     : ${ac_SHMPATH}
 Default precision           : ${ac_PRECISION}