diff --git a/Grid/perfmon/Tracing.h b/Grid/perfmon/Tracing.h index 10b638dc7..150529037 100644 --- a/Grid/perfmon/Tracing.h +++ b/Grid/perfmon/Tracing.h @@ -25,11 +25,9 @@ class GridTracer { public: GridTracer(const char* name) { roctxRangePushA(name); - std::cout << "roctxRangePush "<StencilBarrier(); #endif } + traceID = traceStart("Stencil::CommunicateBegin"); for(int i=0;iBarrier(); _grid->StencilSendToRecvFromComplete(MpiReqs,0); // MPI is done // if ( this->partialDirichlet ) DslashLogPartial(); + traceStop(traceID); + if ( this->fullDirichlet ) DslashLogDirichlet(); else DslashLogFull(); // acceleratorCopySynchronise();// is in the StencilSendToRecvFromComplete diff --git a/Grid/threads/Accelerator.h b/Grid/threads/Accelerator.h index 15edda1e0..ac8c35735 100644 --- a/Grid/threads/Accelerator.h +++ b/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 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);} +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 acceleratorMemSet(void *base,int value,size_t bytes) { auto discard=hipMemset(base,value,bytes); discard = hipDeviceSynchronize(); } typedef int acceleratorEvent_t; diff --git a/MPI_benchmark/compile-command b/MPI_benchmark/compile-command index efa7ce21c..c256d26e2 100644 --- a/MPI_benchmark/compile-command +++ b/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 diff --git a/MPI_benchmark/halo_mpi.cc b/MPI_benchmark/halo_mpi.cc index b9b08c14a..433e3fcf5 100644 --- a/MPI_benchmark/halo_mpi.cc +++ b/MPI_benchmark/halo_mpi.cc @@ -19,8 +19,8 @@ ************************************************************** */ #undef ACC_CUDA -#undef ACC_HIP -#define ACC_SYCL +#define ACC_HIP +#undef ACC_SYCL #undef ACC_NONE /************************************************************** diff --git a/examples/Example_pvdagm_3level_dense.cc b/examples/Example_pvdagm_3level_dense.cc index 8462d1b71..84db5730f 100644 --- a/examples/Example_pvdagm_3level_dense.cc +++ b/examples/Example_pvdagm_3level_dense.cc @@ -948,7 +948,7 @@ void runMG( std::cout< SmootherGCR(0.01,1,ShiftedPVdagM,simple_fine,16,16); + PrecGeneralisedConjugateResidualNonHermitian 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; - clatt2[1] /= 4; - clatt2[2] /= 3; - clatt2[3] /= 6; + for (int d = 0; d < 4; d++) clatt2[d] /= Block2[d]; 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 TwoLevelMG; PVdagM_t PVdagM(Ddwf,Dpv); - ShiftedPVdagM_t ShiftedPVdagM(0.01,Ddwf,Dpv); + ShiftedPVdagM_t ShiftedPVdagM(0.2,Ddwf,Dpv); NextToNearestStencilGeometry5D geom(Coarse5d); diff --git a/systems/Frontier/config-command b/systems/Frontier/config-command index 0d6db4390..278a64c1a 100644 --- a/systems/Frontier/config-command +++ b/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 \