diff --git a/Grid/algorithms/blas/BatchedBlas.h b/Grid/algorithms/blas/BatchedBlas.h index 0dbd17533..c0be5dd91 100644 --- a/Grid/algorithms/blas/BatchedBlas.h +++ b/Grid/algorithms/blas/BatchedBlas.h @@ -148,14 +148,15 @@ public: #endif #ifdef GRID_HIP std::cout << "hipblasCreate"< &vecs, deviceVector &blas) { + GRID_TRACE("ImportFineGridVectors"); int nvec = vecs.size(); typedef typename Field::vector_object vobj; // std::cout << GridLogMessage <<" BlockProjector importing "< &vecs, deviceVector &blas) { + GRID_TRACE("ExportFineGridVectors"); typedef typename Field::vector_object vobj; int nvec = vecs.size(); @@ -622,6 +624,7 @@ public: template void ImportCoarseGridVectors(std::vector > &vecs, deviceVector &blas) { + GRID_TRACE("ImportCoarseGridVectors"); int nvec = vecs.size(); typedef typename vobj::scalar_object coarse_scalar_object; @@ -688,6 +691,7 @@ public: template void ExportCoarseGridVectors(std::vector > &vecs, deviceVector &blas) { + GRID_TRACE("ExportCoarseGridVectors"); int nvec = vecs.size(); typedef typename vobj::scalar_object coarse_scalar_object; // std::cout << GridLogMessage<<" BlockProjector exporting "< Cd(coarse_vol); // std::cout << "BlockProject pointers"< hVd(coarse_vol), hFd(coarse_vol), hCd(coarse_vol); + for(int c=0;c Fd(coarse_vol); deviceVector Cd(coarse_vol); - for(int c=0;c hVd(coarse_vol), hFd(coarse_vol), hCd(coarse_vol); + for(int c=0;c &Fd, deviceVector &Cd) { - for(int c=0;c hVd(coarse_vol), hFd(coarse_vol), hCd(coarse_vol); + for(int c=0;cAllGatherV(send, counts[me], (void *)&dRecvBuf[0], counts, displs, sizeof(ComplexD)); @@ -385,6 +456,7 @@ public: tAllreduce += tar; tARmin = std::min(tARmin, tar); tARmax = std::max(tARmax, tar); + tARall.push_back(tar); bytesAllreduce += panelBytesThis; nAllreduce++; @@ -657,6 +729,7 @@ public: tLeaf = 0.0; tARmin = 1.0e30; tARmax = 0.0; + tARall.clear(); bytesAllreduce = 0; nAllreduce = 0; nGatherGemm = 0; @@ -718,6 +791,29 @@ public: << " effective " << (tAllreduce>0 ? bytesAllreduce/tAllreduce*1.0e6/1.0e9 : 0.0) << " GB/s" << std::endl; + ////////////////////////////////////////////////////////////////////// + // Distribution, not just min/avg/max. A barrier that merely REBOOKS + // skew shifts the median; a barrier that suppresses pathological + // collectives shortens the tail. p50 vs p99 separates the two from a + // single run, which min/avg/max cannot. + ////////////////////////////////////////////////////////////////////// + if ( tARall.size() ) { + std::vector v = tARall; + std::sort(v.begin(),v.end()); + size_t n = v.size(); + auto pc = [&](double f){ size_t i=(size_t)(f*(n-1)); return v[i]/1.0e3; }; + double med = pc(0.50); + // how much time is spent in calls that are gross outliers + double tail=0.0; size_t ntail=0; + for(size_t i=0;i 4.0*med*1.0e3 ) { tail+=v[i]; ntail++; } + std::cout << GridLogMessage << "Schur comms distribution (boss):" + << " p50 " << med + << " p90 " << pc(0.90) + << " p99 " << pc(0.99) + << " ms calls > 4x median: " << ntail << "/" << n + << " carrying " << tail/1.0e6 << " s" + << std::endl; + } } }; diff --git a/Grid/stencil/Stencil.h b/Grid/stencil/Stencil.h index 687b96219..181f10434 100644 --- a/Grid/stencil/Stencil.h +++ b/Grid/stencil/Stencil.h @@ -315,6 +315,27 @@ public: } int traceID; + // Delivered-comms instrumentation, per CommunicateBegin/Complete pair. + // + // OffNodeBytes : bytes handed to MPI, i.e. EXCLUDING intranode traffic -- + // StencilSendToRecvFrom* return off-node bytes only, which is the + // differentiation we want. It is BIDIRECTIONAL (send + receive) under + // ACCELERATOR_AWARE_MPI (Communicator_mpi3.cc:463,481). On the + // host-staged path the send is deferred to PollDtoH and its bytes are + // NOT currently counted, so figures from the two paths are not + // comparable. Prepare() returns 0.0 on the accelerator-aware path. + // + // CommTimer : microseconds between traceStart and traceStop, i.e. exactly + // the "Stencil::CommunicateBegin" roctx range -- transfer only, with the + // StencilBarrier and the compress kernels already excluded. It spans the + // window in which the interior kernel runs, so the derived rate is + // bandwidth delivered CONCURRENT WITH COMPUTE, which is the quantity a + // comms-only benchmark cannot see. + // + // InterNodeBandwidthMBps : the per-call rate. For a reportable figure + // accumulate bytes and time separately across calls and divide once -- + // averaging per-call rates over-weights the fast calls. See + // benchmarks/Benchmark_dwf.cc. double OffNodeBytes; double CommTimer; double InterNodeBandwidthMBps; @@ -955,6 +976,12 @@ public: bool preserve_shm=false) { SloppyComms = 0; + // Never leave the delivered-comms counters uninitialised: they are read + // from outside (benchmarks, drivers) and a stencil that has not yet + // exchanged would otherwise return denormal garbage. + OffNodeBytes = 0; + CommTimer = 0; + InterNodeBandwidthMBps = 0; face_table_computed=0; _grid = grid; this->parameters=p; diff --git a/benchmarks/Benchmark_dwf.cc b/benchmarks/Benchmark_dwf.cc index 2d21976f8..dbda78ed8 100644 --- a/benchmarks/Benchmark_dwf.cc +++ b/benchmarks/Benchmark_dwf.cc @@ -276,10 +276,15 @@ void Benchmark(int Ls, Coordinate Dirichlet,bool sloppy) Dw.Dhop(src,result,0); std::cout<Barrier(); @@ -300,7 +305,34 @@ void Benchmark(int Ls, Coordinate Dirichlet,bool sloppy) std::cout<Barrier(); Dw.DhopEO(src_o,r_e,DaggerNo); double t0=usecond(); + double locBytes = 0, locTime = 0; for(int i=0;iBarrier(); @@ -420,6 +455,13 @@ void Benchmark(int Ls, Coordinate Dirichlet,bool sloppy) std::cout<GlobalSum(sumBytes); + double sumTime = locTime; FGrid->GlobalSum(sumTime); + double avgTime = sumTime/NP; + std::cout< PVdagM_t; typedef ShiftedPVdagMLinearOperator ShiftedPVdagM_t; PVdagM_t PVdagM(Ddwf,Dpv); @@ -815,13 +834,20 @@ int main (int argc, char ** argv) std::cout << GridLogMessage << "V2 3-level solve Nrhs "< counts(P,0), displs(P,0); int64_t total=0; for(int r=0;r P ? P : AGN); std::vector counts(P,0), displs(P,0); int64_t total=0; for(int r=0;r