From a8c80b8aa2ade26b011a46d070a97d21e64a850e Mon Sep 17 00:00:00 2001 From: George Zagaris Date: Sun, 19 Sep 2021 11:14:45 -0400 Subject: [PATCH 1/5] 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 c4f1646d6261b5ce710c2331a3f329620c3523ec Mon Sep 17 00:00:00 2001 From: George Zagaris Date: Fri, 24 Sep 2021 17:14:36 -0400 Subject: [PATCH 2/5] ENH: add NVTX annotations to UpdateFringes --- src/LocalMesh.C | 30 +++++++++++++++++++++++------- 1 file changed, 23 insertions(+), 7 deletions(-) diff --git a/src/LocalMesh.C b/src/LocalMesh.C index 386312d..4f134ca 100644 --- a/src/LocalMesh.C +++ b/src/LocalMesh.C @@ -337,10 +337,16 @@ void LocalMesh::UpdateFringes(double *qh, double *qd) // doesn't work CUDA-Aware -- debug this void LocalMesh::UpdateFringes(double *qd) { + FVSAND_NVTX_FUNCTION("UpdateFringes"); + nthreads=device2host.size(); if(nthreads == 0) return; - FVSAND_GPU_KERNEL_LAUNCH( updateHost, nthreads, - qbuf_d,qd,device2host_d,nthreads); + + FVSAND_NVTX_SECTION( "pack", + 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 @@ -348,16 +354,26 @@ void LocalMesh::UpdateFringes(double *qd) 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()); + FVSAND_NVTX_SECTION( "copy-to-host", + 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); + // same as above // not doing cuda-aware now - gpu::copy_to_device(qbuf_d2,qbuf2,sizeof(double)*host2device.size()); - + FVSAND_NVTX_SECTION( "copy-to-device", + 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); + + FVSAND_NVTX_SECTION( "unpack", + FVSAND_GPU_KERNEL_LAUNCH( updateDevice, nthreads, + qd,qbuf_d2,host2device_d,nthreads ); + ); + } From ed294302699b54cfe7437f46d0788057e9fe8766 Mon Sep 17 00:00:00 2001 From: George Zagaris Date: Sun, 26 Sep 2021 13:01:30 -0400 Subject: [PATCH 3/5] ENH: comment out Tecplot I/O for mesh partitions --- src/fvsand.C | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/fvsand.C b/src/fvsand.C index e55af82..23a563d 100644 --- a/src/fvsand.C +++ b/src/fvsand.C @@ -143,7 +143,7 @@ int main(int argc, char *argv[]) printf("# ----------------------------------\n"); } - lm->WriteMesh(myid); + // lm->WriteMesh(myid); MPI_Finalize(); From ea3ae3e077a3ff9a788df38bb5332f8a1e91c248 Mon Sep 17 00:00:00 2001 From: George Zagaris Date: Mon, 27 Sep 2021 09:26:34 -0400 Subject: [PATCH 4/5] ENH: sprinkle in some more NVTX annotations --- src/LocalMesh.C | 166 +++++++++++++++++++++++++++--------------------- 1 file changed, 93 insertions(+), 73 deletions(-) diff --git a/src/LocalMesh.C b/src/LocalMesh.C index 4f134ca..89b7d4d 100644 --- a/src/LocalMesh.C +++ b/src/LocalMesh.C @@ -421,88 +421,108 @@ void LocalMesh::Jacobi(double *q, double dt, int nsweep, int istoreJac) exit(0); */ FVSAND_NVTX_FUNCTION("Jacobi"); - nthreads=(ncells+nhalo)*nfields_d; - FVSAND_GPU_KERNEL_LAUNCH(setValues,nthreads, - dq_d, 0.0, nthreads); + + FVSAND_NVTX_SECTION( "setValues", + nthreads=(ncells+nhalo)*nfields_d; + FVSAND_GPU_KERNEL_LAUNCH(setValues,nthreads,dq_d, 0.0, nthreads); + ); + //compute and store Jacobians; if(istoreJac==1){ - nthreads=ncells+nhalo; - FVSAND_GPU_KERNEL_LAUNCH(fillJacobians,nthreads, - q, normals_d, volume_d, - rmatall_d, Dall_d, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); + + FVSAND_NVTX_SECTION( "fillJacobians", + nthreads=ncells+nhalo; + FVSAND_GPU_KERNEL_LAUNCH(fillJacobians,nthreads, + q, normals_d, volume_d, + rmatall_d, Dall_d, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + ); + } if (istoreJac==4) { - nthreads=ncells+nhalo; - FVSAND_GPU_KERNEL_LAUNCH(fillJacobians_diag,nthreads, - q, normals_d, volume_d, - Dall_d, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); + + FVSAND_NVTX_SECTION( "fillJacobians_diag", + nthreads=ncells+nhalo; + FVSAND_GPU_KERNEL_LAUNCH(fillJacobians_diag,nthreads, + q, normals_d, volume_d, + Dall_d, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + ); + } if (istoreJac==5) { - nthreads=ncells+nhalo; - FVSAND_GPU_KERNEL_LAUNCH(fillJacobians_diag_f,nthreads, - q, normals_d, volume_d, - Dall_d_f, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); + + FVSAND_NVTX_SECTION( "fillJacobians_diag_f", + nthreads=ncells+nhalo; + FVSAND_GPU_KERNEL_LAUNCH(fillJacobians_diag_f,nthreads, + q, normals_d, volume_d, + Dall_d_f, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + ); + } // Jacobi Sweeps - for(int m = 0; m < nsweep; m++){ - //printf("Sweep %i\n=================\n",m); - nthreads=ncells+nhalo; - // compute dqtilde for all cells - if(istoreJac==0){ - FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep,nthreads, - q, res_d, dq_d, dqupdate_d, normals_d, volume_d, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); - } - else if(istoreJac==1) { - FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep1,nthreads, - q, res_d, dq_d, dqupdate_d, normals_d, volume_d, - rmatall_d, Dall_d, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); - } - else if(istoreJac==2) { - FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep2,nthreads, - q, res_d, dq_d, dqupdate_d, normals_d, volume_d, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); - } - else if(istoreJac==3) { - FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep3,nthreads, - q, res_d, dq_d, dqupdate_d, normals_d, volume_d, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); - } - else if(istoreJac==4) { - FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep4,nthreads, - q, res_d, dq_d, dqupdate_d, normals_d, volume_d, - Dall_d, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); - } - else if(istoreJac==5) { - FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep5,nthreads, - q, res_d, dq_d, dqupdate_d, normals_d, volume_d, - Dall_d_f, - flovar_d, cell2cell_d, - nccft_d, nfields_d, istor, ncells, facetype_d, dt); - } - // update dq = dqtilde for all cells - - nthreads=(ncells+nhalo)*nfields_d; - FVSAND_GPU_KERNEL_LAUNCH(copyValues,nthreads, - dq_d, dqupdate_d, nthreads); - - // Store final dq in res to be used in update routine - UpdateFringes(dq_d); - } + FVSAND_NVTX_SECTION( "JacobiSweep", + for(int m = 0; m < nsweep; m++){ + //printf("Sweep %i\n=================\n",m); + nthreads=ncells+nhalo; + // compute dqtilde for all cells + if(istoreJac==0){ + FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep,nthreads, + q, res_d, dq_d, dqupdate_d, normals_d, volume_d, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + } + else if(istoreJac==1) { + FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep1,nthreads, + q, res_d, dq_d, dqupdate_d, normals_d, volume_d, + rmatall_d, Dall_d, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + } + else if(istoreJac==2) { + FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep2,nthreads, + q, res_d, dq_d, dqupdate_d, normals_d, volume_d, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + } + else if(istoreJac==3) { + FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep3,nthreads, + q, res_d, dq_d, dqupdate_d, normals_d, volume_d, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + } + else if(istoreJac==4) { + FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep4,nthreads, + q, res_d, dq_d, dqupdate_d, normals_d, volume_d, + Dall_d, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + } + else if(istoreJac==5) { + FVSAND_GPU_KERNEL_LAUNCH(jacobiSweep5,nthreads, + q, res_d, dq_d, dqupdate_d, normals_d, volume_d, + Dall_d_f, + flovar_d, cell2cell_d, + nccft_d, nfields_d, istor, ncells, facetype_d, dt); + } + // update dq = dqtilde for all cells + + FVSAND_NVTX_SECTION( "copyValues", + nthreads=(ncells+nhalo)*nfields_d; + FVSAND_GPU_KERNEL_LAUNCH(copyValues,nthreads, + dq_d, dqupdate_d, nthreads); + ); + + // Store final dq in res to be used in update routine + UpdateFringes(dq_d); + } // END for all jacobi sweeps + + ); } void LocalMesh::Update(double *qdest, double *qsrc, double fscal) From 8661a1e557db72b61da69498383cd44d95b8f71e Mon Sep 17 00:00:00 2001 From: George Zagaris Date: Mon, 27 Sep 2021 11:34:56 -0400 Subject: [PATCH 5/5] ENH: add NVTX section around MPI_Waitall() --- src/LocalMesh.C | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/LocalMesh.C b/src/LocalMesh.C index 89b7d4d..bf95360 100644 --- a/src/LocalMesh.C +++ b/src/LocalMesh.C @@ -359,7 +359,9 @@ void LocalMesh::UpdateFringes(double *qd) ); pc.postSends_direct(qbuf,nfields_d,sndmap,ireq,mycomm,&reqcount); - pc.finish_comm(reqcount,ireq,istatus); + FVSAND_NVTX_SECTION( "waitall", + pc.finish_comm(reqcount,ireq,istatus); + ); // same as above // not doing cuda-aware now