Better ROCTX tracing

Grid/perfmon/Tracing.h

@@ -25,11 +25,9 @@ class GridTracer {
  public:
   GridTracer(const char* name) {
     roctxRangePushA(name);
-    std::cout << "roctxRangePush "<<name<<std::endl;
   }
   ~GridTracer() {
     roctxRangePop();
-    std::cout << "roctxRangePop "<<std::endl;
   }
 };
 inline void tracePush(const char *name) { roctxRangePushA(name); }

Grid/stencil/Stencil.h

@@ -314,6 +314,7 @@ public:
     return accessor;
   }
 
+  int              traceID;
   int face_table_computed;
   //  int partialDirichlet;
   int fullDirichlet;
@@ -537,6 +538,7 @@ public:
       _grid->StencilBarrier(); 
 #endif
     }
+    traceID = traceStart("Stencil::CommunicateBegin");
     
     for(int i=0;i<Packets.size();i++){
       //      std::cout << "Communicate prepare "<<i<<std::endl;
@@ -588,6 +590,8 @@ public:
     //    _grid->Barrier();
     _grid->StencilSendToRecvFromComplete(MpiReqs,0); // MPI is done
     //    if   ( this->partialDirichlet ) DslashLogPartial();
+    traceStop(traceID);
+    
     if ( this->fullDirichlet ) DslashLogDirichlet();
     else DslashLogFull();
     //    acceleratorCopySynchronise();// is in the StencilSendToRecvFromComplete

Grid/threads/Accelerator.h

@@ -251,7 +251,7 @@ inline void acceleratorFreeDevice(void *ptr){ cudaFree(ptr);};
 inline void acceleratorFreeHost(void *ptr){ cudaFree(ptr);};
 inline void acceleratorCopyToDevice(const void *from,void *to,size_t bytes)  { cudaMemcpy(to,from,bytes, cudaMemcpyHostToDevice);}
 inline void acceleratorCopyFromDevice(const void *from,void *to,size_t bytes){ cudaMemcpy(to,from,bytes, cudaMemcpyDeviceToHost);}
-inline void acceleratorMemSet(void *base,int value,size_t bytes) { cudaMemset(base,value,bytes);}
+inline void acceleratorMemSet(void *base,int value,size_t bytes) { cudaMemset(base,value,bytes); cudaDeviceSynchronize(); }
 inline acceleratorEvent_t acceleratorCopyToDeviceAsynch(void *from, void *to, size_t bytes, cudaStream_t stream = copyStream) {
   acceleratorCopyToDevice(from,to,bytes);
   return 0;
@@ -525,10 +525,9 @@ inline void *acceleratorAllocDevice(size_t bytes)
 inline void acceleratorFreeHost(void *ptr){ auto discard=hipFree(ptr);};
 inline void acceleratorFreeShared(void *ptr){ auto discard=hipFree(ptr);};
 inline void acceleratorFreeDevice(void *ptr){ auto discard=hipFree(ptr);};
-inline void acceleratorCopyToDevice(const void *from,void *to,size_t bytes)  { auto discard=hipMemcpy(to,from,bytes, hipMemcpyHostToDevice);}
+inline void acceleratorCopyToDevice(const void *from,void *to,size_t bytes)  { auto discard=hipMemcpy(to,from,bytes, hipMemcpyHostToDevice); }
-inline void acceleratorCopyFromDevice(const void *from,void *to,size_t bytes){ auto discard=hipMemcpy(to,from,bytes, hipMemcpyDeviceToHost);}
+inline void acceleratorCopyFromDevice(const void *from,void *to,size_t bytes){ auto discard=hipMemcpy(to,from,bytes, hipMemcpyDeviceToHost); }
-
+inline void acceleratorMemSet(void *base,int value,size_t bytes) { auto discard=hipMemset(base,value,bytes); discard = hipDeviceSynchronize(); }
-inline void acceleratorMemSet(void *base,int value,size_t bytes) { auto discard=hipMemset(base,value,bytes);}
 
 typedef int acceleratorEvent_t;
 

MPI_benchmark/compile-command

@@ -1,4 +1,4 @@
-mpicxx  -fsycl halo_mpi.cc -o halo_mpi
+mpicxx -O2 -x hip halo_mpi.cc -o halo_mpi -L${PE_MPICH_GTL_DIR_amd_gfx90a} -lmpi_gtl_hsa
 mpicxx -O2 -std=c++11 io_mpi.cc -o io_mpi
 # Frontier: (hipcc via mpicxx wrapper; ACC_HIP is the default in-file)
-mpicxx -O2 -x hip gather_mpi.cc -o gather_mpi -L${ROCM_PATH}/lib -lamdhip64
+mpicxx -O2 -x hip gather_mpi.cc -o gather_mpi -L${ROCM_PATH}/lib -lamdhip64 -L${PE_MPICH_GTL_DIR_amd_gfx90a} -lmpi_gtl_hsa

MPI_benchmark/halo_mpi.cc

@@ -19,8 +19,8 @@
  **************************************************************
  */
 #undef ACC_CUDA
-#undef  ACC_HIP
+#define  ACC_HIP
-#define  ACC_SYCL
+#undef  ACC_SYCL
 #undef  ACC_NONE
 
 /**************************************************************

examples/Example_pvdagm_3level_dense.cc

@@ -948,7 +948,7 @@ void runMG(
   std::cout<<GridLogMessage<<" Three-level outer solve (dense CC bottom)"<<std::endl;
   std::cout<<GridLogMessage<<"*******************************************"<<std::endl;
 
-  PrecGeneralisedConjugateResidualNonHermitian<LatticeFermionD> SmootherGCR(0.01,1,ShiftedPVdagM,simple_fine,16,16);
+  PrecGeneralisedConjugateResidualNonHermitian<LatticeFermionD> SmootherGCR(0.01,1,ShiftedPVdagM,simple_fine,8,8);
   SmootherGCR.SetZeroGuess(1);        // pre+post smoother slots both zero their guess: saves 2 fine mults/outer
   SmootherGCR.Level(1);
   SmootherGCR.Name("Fsmoother");
@@ -1000,7 +1000,15 @@ int main (int argc, char ** argv)
 
   // Level 1 coarse grid: block 2^4 from fine (48×48×48×96 → 24×24×24×48, Ls=1)
   Coordinate clatt = lat_size;
-  for (int d = 0; d < 4; d++) clatt[d] /= 2;
+  //  Coordinate Block1({2,2,2,2});
+  //  Coordinate Block2({8,4,3,6});
+  //  Coordinate Block1({2,2,3,3});
+  //  Coordinate Block2({8,4,2,4});
+  //  Coordinate Block1({4,2,3,3}); // 144s with Fsmoother 12
+  //  Coordinate Block2({4,4,2,4});
+  Coordinate Block1({4,4,3,3}); 
+  Coordinate Block2({4,2,2,4});
+  for (int d = 0; d < 4; d++) clatt[d] /= Block1[d];
   std::cout << GridLogMessage << "Level 1 coarse lattice: " << clatt << std::endl;
 
   GridCartesian *Coarse4d  = SpaceTimeGrid::makeFourDimGrid(clatt, GridDefaultSimd(Nd,vComplex::Nsimd()),GridDefaultMpi());
@@ -1015,10 +1023,7 @@ int main (int argc, char ** argv)
   //   blocking; the iterative CC solve was the sole remaining cost -- which the
   //   dense inverse removes.
   Coordinate clatt2 = clatt;
-  clatt2[0] /= 8;
+  for (int d = 0; d < 4; d++) clatt2[d] /= Block2[d];
-  clatt2[1] /= 4;
-  clatt2[2] /= 3;
-  clatt2[3] /= 6;
   std::cout << GridLogMessage << "Level 2 supercoarse lattice: " << clatt2 << std::endl;
 
   GridCartesian *CoarseCoarse4d = SpaceTimeGrid::makeFourDimGrid(clatt2, GridDefaultSimd(Nd,vComplex::Nsimd()),GridDefaultMpi());
@@ -1048,7 +1053,7 @@ int main (int argc, char ** argv)
   typedef MGPreconditioner<vSpinColourVector,vTComplex,nbasis>        TwoLevelMG;
 
   PVdagM_t        PVdagM(Ddwf,Dpv);
-  ShiftedPVdagM_t ShiftedPVdagM(0.01,Ddwf,Dpv);
+  ShiftedPVdagM_t ShiftedPVdagM(0.2,Ddwf,Dpv);
 
   NextToNearestStencilGeometry5D geom(Coarse5d);
 

systems/Frontier/config-command

@@ -4,7 +4,7 @@ CLIME=`spack find --paths c-lime@2-3-9 | grep c-lime| cut -c 15-`
 --with-lime=$CLIME \
 --enable-unified=no \
 --enable-shm=nvlink \
---enable-tracing=none \
+--enable-tracing=roctx \
 --enable-accelerator=hip \
 --enable-gen-simd-width=64 \
 --disable-gparity \