Skip to content

Minor Improvements and more NVTX Annotations #10

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
gzagaris wants to merge 5 commits into
base: main
Choose a base branch
from
Open

Minor Improvements and more NVTX Annotations #10

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
a8c80b8
ENH: pass maps by const reference + inline methods
gzagaris Sep 19, 2021
c4f1646
ENH: add NVTX annotations to UpdateFringes
gzagaris Sep 24, 2021
ed29430
ENH: comment out Tecplot I/O for mesh partitions
gzagaris Sep 26, 2021
ea3ae3e
ENH: sprinkle in some more NVTX annotations
gzagaris Sep 27, 2021
8661a1e
ENH: add NVTX section around MPI_Waitall()
gzagaris Sep 27, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
200 changes: 119 additions & 81 deletions src/LocalMesh.C
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -337,27 +337,45 @@ 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
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<double>(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);
FVSAND_NVTX_SECTION( "waitall",
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 );
);

}


Expand Down Expand Up @@ -405,88 +423,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)
Expand Down
2 changes: 1 addition & 1 deletion src/fvsand.C
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -143,7 +143,7 @@ int main(int argc, char *argv[])
printf("# ----------------------------------\n");
}

lm->WriteMesh(myid);
// lm->WriteMesh(myid);

MPI_Finalize();

Expand Down
32 changes: 17 additions & 15 deletions src/parallelComm.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -288,6 +288,7 @@ namespace FVSAND {
delete [] istatus;
}

inline
void postRecvs_direct(double *qbuf, int nfields,
std::unordered_map <int, std::vector<int>> rcvmap,
MPI_Request *ireq,
Expand All @@ -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 <int, std::vector<int>> sndmap,
MPI_Request *ireq,
Expand All @@ -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);
Expand Down