From 7085922de06546eefe982b3a737eb95400c5345d Mon Sep 17 00:00:00 2001 From: George Zagaris Date: Sun, 19 Sep 2021 11:14:45 -0400 Subject: [PATCH 1/3] ENH: pass maps by const reference + inline methods Pass maps to the post-send and post receive methods by const reference to avoid deep-copies. Also, mark those methods as "inline", which might help the compiler. --- src/parallelComm.h | 32 +++++++++++++++++--------------- 1 file changed, 17 insertions(+), 15 deletions(-) diff --git a/src/parallelComm.h b/src/parallelComm.h index 3f5caeb..fbed865 100644 --- a/src/parallelComm.h +++ b/src/parallelComm.h @@ -288,6 +288,7 @@ namespace FVSAND { delete [] istatus; } + inline void postRecvs_direct(double *qbuf, int nfields, std::unordered_map > rcvmap, MPI_Request *ireq, @@ -296,17 +297,17 @@ namespace FVSAND { { int offset=0; int rcount=*k; - for(auto r:rcvmap) - { - MPI_Irecv(qbuf+offset, - r.second.size()*nfields, MPI_DOUBLE, - r.first,0,comm,&ireq[rcount++]); - offset+=r.second.size()*nfields; - } + for(const auto& r: rcvmap) + { + MPI_Irecv(qbuf+offset, + r.second.size()*nfields, MPI_DOUBLE, + r.first,0,comm,&ireq[rcount++]); + offset+=r.second.size()*nfields; + } *k=rcount; } - + inline void postSends_direct(double *qbuf, int nfields, std::unordered_map > sndmap, MPI_Request *ireq, @@ -315,16 +316,17 @@ namespace FVSAND { { int offset=0; int scount=*k; - for(auto s:sndmap) - { - MPI_Isend(qbuf+offset, - s.second.size()*nfields, MPI_DOUBLE, - s.first,0,comm,&ireq[scount++]); - offset+=s.second.size()*nfields; - } + for(const auto& s:sndmap) + { + MPI_Isend(qbuf+offset, + s.second.size()*nfields, MPI_DOUBLE, + s.first,0,comm,&ireq[scount++]); + offset+=s.second.size()*nfields; + } *k=scount; } + inline void finish_comm(int nrequests, MPI_Request *ireq, MPI_Status *istatus) { MPI_Waitall(nrequests,ireq,istatus); From 59ae40a1544bc14576879cc33795a30dc3086538 Mon Sep 17 00:00:00 2001 From: George Zagaris Date: Sun, 19 Sep 2021 11:17:13 -0400 Subject: [PATCH 2/3] ENH: use cuda-aware MPI This works on Ascent, but, it requires the `--smpiargs="-gpu"` flag to be passed to `jsrun`. --- src/LocalMesh.C | 22 ++++++++++------------ 1 file changed, 10 insertions(+), 12 deletions(-) diff --git a/src/LocalMesh.C b/src/LocalMesh.C index 386312d..4d3a992 100644 --- a/src/LocalMesh.C +++ b/src/LocalMesh.C @@ -340,24 +340,22 @@ void LocalMesh::UpdateFringes(double *qd) nthreads=device2host.size(); if(nthreads == 0) return; FVSAND_GPU_KERNEL_LAUNCH( updateHost, nthreads, - qbuf_d,qd,device2host_d,nthreads); + qbuf_d, qd,device2host_d,nthreads); // separate sends and receives so that we can overlap comm and calculation // in the residual and iteration loops. // TODO (george) use qbuf2_d and qbuf_d instead of qbuf2 and qbuf for cuda-aware int reqcount=0; - pc.postRecvs_direct(qbuf2,nfields_d,rcvmap,ireq,mycomm,&reqcount); - // TODO (george) with cuda-aware this pull is not required - // but it doesn't work now - gpu::pull_from_device(qbuf,qbuf_d,sizeof(double)*device2host.size()); - pc.postSends_direct(qbuf,nfields_d,sndmap,ireq,mycomm,&reqcount); + pc.postRecvs_direct(qbuf_d2,nfields_d,rcvmap,ireq,mycomm,&reqcount); + + // wait for pack kernel to finish packing buffers? + FVSAND::gpu::synchronize(); + + pc.postSends_direct(qbuf_d,nfields_d,sndmap,ireq,mycomm,&reqcount); pc.finish_comm(reqcount,ireq,istatus); - // same as above - // not doing cuda-aware now - gpu::copy_to_device(qbuf_d2,qbuf2,sizeof(double)*host2device.size()); - nthreads=host2device.size(); - FVSAND_GPU_KERNEL_LAUNCH( updateDevice, nthreads, - qd,qbuf_d2,host2device_d,nthreads); + // unpack + FVSAND_GPU_KERNEL_LAUNCH( updateDevice, nthreads, qd, + qbuf_d2, host2device_d, nthreads ); } From 5d6c4ee6349da15aa081a70e16e24ad514e12daa Mon Sep 17 00:00:00 2001 From: George Zagaris Date: Tue, 28 Sep 2021 17:59:27 -0400 Subject: [PATCH 3/3] ENH: add cmd line option for CUDA-aware MPI Adds cmd line option to enable/disable CUDA-aware MPI. --- src/LocalMesh.C | 61 ++++++++++++++++++++++++++++++++++--------------- src/LocalMesh.h | 5 +++- src/fvsand.C | 23 +++++++++++++++---- 3 files changed, 65 insertions(+), 24 deletions(-) diff --git a/src/LocalMesh.C b/src/LocalMesh.C index 4d3a992..490884d 100644 --- a/src/LocalMesh.C +++ b/src/LocalMesh.C @@ -64,13 +64,14 @@ LocalMesh::~LocalMesh() FVSAND_FREE_DEVICE(dqupdate_d); } -LocalMesh::LocalMesh(GlobalMesh *g, int myid, MPI_Comm comm) +LocalMesh::LocalMesh(GlobalMesh *g, int myid, MPI_Comm comm, bool usecuda) { Timer stopwatch; mycomm=comm; int ierr=MPI_Comm_rank(comm,&myid); ierr=MPI_Comm_size(comm,&ngroup); parallelComm pc; + usecudampi = usecuda; // create communication patterns and ghost cells stopwatch.tick(); @@ -339,23 +340,47 @@ void LocalMesh::UpdateFringes(double *qd) { nthreads=device2host.size(); if(nthreads == 0) return; - FVSAND_GPU_KERNEL_LAUNCH( updateHost, nthreads, - qbuf_d, qd,device2host_d,nthreads); - // separate sends and receives so that we can overlap comm and calculation - // in the residual and iteration loops. - // TODO (george) use qbuf2_d and qbuf_d instead of qbuf2 and qbuf for cuda-aware - int reqcount=0; - pc.postRecvs_direct(qbuf_d2,nfields_d,rcvmap,ireq,mycomm,&reqcount); - - // wait for pack kernel to finish packing buffers? - FVSAND::gpu::synchronize(); - - pc.postSends_direct(qbuf_d,nfields_d,sndmap,ireq,mycomm,&reqcount); - pc.finish_comm(reqcount,ireq,istatus); - - // unpack - FVSAND_GPU_KERNEL_LAUNCH( updateDevice, nthreads, qd, - qbuf_d2, host2device_d, nthreads ); + + if ( usecudampi ) + { + // uses CUDA-aware MPI and passes device pointers to MPI + FVSAND_GPU_KERNEL_LAUNCH( updateHost, nthreads, + qbuf_d, qd,device2host_d,nthreads); + // separate sends and receives so that we can overlap comm and calculation + // in the residual and iteration loops. + int reqcount=0; + pc.postRecvs_direct(qbuf_d2,nfields_d,rcvmap,ireq,mycomm,&reqcount); + + // wait for pack kernel to finish packing buffers? + FVSAND::gpu::synchronize(); + + pc.postSends_direct(qbuf_d,nfields_d,sndmap,ireq,mycomm,&reqcount); + pc.finish_comm(reqcount,ireq,istatus); + + // unpack + FVSAND_GPU_KERNEL_LAUNCH( updateDevice, nthreads, qd, + qbuf_d2, host2device_d, nthreads ); + } + else + { + // does memcpy from/to global memory and only passes host-pointers to MPI + FVSAND_GPU_KERNEL_LAUNCH( updateHost, nthreads, + qbuf_d,qd,device2host_d,nthreads); + + int reqcount=0; + pc.postRecvs_direct(qbuf2,nfields_d,rcvmap,ireq,mycomm,&reqcount); + gpu::pull_from_device(qbuf,qbuf_d,sizeof(double)*device2host.size()); + + pc.postSends_direct(qbuf,nfields_d,sndmap,ireq,mycomm,&reqcount); + pc.finish_comm(reqcount,ireq,istatus); + + gpu::copy_to_device(qbuf_d2,qbuf2,sizeof(double)*host2device.size()); + + nthreads=host2device.size(); + FVSAND_GPU_KERNEL_LAUNCH( updateDevice, nthreads, + qd,qbuf_d2,host2device_d,nthreads); + } + } diff --git a/src/LocalMesh.h b/src/LocalMesh.h index 666d93a..c5fe577 100644 --- a/src/LocalMesh.h +++ b/src/LocalMesh.h @@ -94,6 +94,8 @@ class LocalMesh MPI_Request *ireq{nullptr}; MPI_Status *istatus{nullptr}; + bool usecudampi{false}; // indicates whether cuda-aware MPI is used + public: // solution fields at n+1,n & n-1 @@ -108,7 +110,8 @@ class LocalMesh ~LocalMesh(); LocalMesh(GlobalMesh *g, int myid, - MPI_Comm comm); + MPI_Comm comm, + bool usecudampi=false ); void WriteMesh(int label); void CreateGridMetrics(int); void CreateFaces(); diff --git a/src/fvsand.C b/src/fvsand.C index e55af82..3a87cdf 100644 --- a/src/fvsand.C +++ b/src/fvsand.C @@ -9,6 +9,7 @@ #include #include "NVTXMacros.h" #include // for std::ostringstream +#include // for std::strcmp() #include "inputParser.h" using namespace FVSAND; @@ -76,10 +77,22 @@ int main(int argc, char *argv[]) int nsweep = 2; // Jacobi Sweeps (=0 means explict) int istoreJac =3; // Jacobian storage or not int restype=0; // restype = 0 (cell-based) 1 (face-based) - if (argc > 1) { - parseInputs(argv[1],fname,&dsmin,&stretch,&nlevels, - flovar,&nsteps,&nsave,&dt,reOrderCells,&nsweep, - &istoreJac,&restype); + bool usecudampi = false; // optional argument to fvsand + if (argc > 1) + { + parseInputs(argv[1],fname,&dsmin,&stretch,&nlevels, + flovar,&nsteps,&nsave,&dt,reOrderCells,&nsweep, + &istoreJac,&restype); + + for ( int iarg=2; iarg < argc; ++iarg ) + { + if ( std::strcmp(argv[iarg], "--usecudampi" ) == 0 ) + { + usecudampi = true; + } + + } // end for all command line arguments + } // runge-kutta tableue @@ -95,7 +108,7 @@ int main(int argc, char *argv[]) // create local mesh partitions // and compute grid metrics LocalMesh *lm; - lm= new LocalMesh(sm,myid,MPI_COMM_WORLD); + lm= new LocalMesh(sm,myid,MPI_COMM_WORLD,usecudampi); lm->CreateGridMetrics(istoreJac); // initialize solution