ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add missing fences in comm_kokkos

Browse Source
This commit is contained in:
Stan Moore
2017-12-15 13:30:17 -07:00
parent 68cf6941e1
commit a7bc3ed391
1 changed files with 2 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
src/KOKKOS/comm_kokkos.cpp
View File

@ -324,6 +324,7 @@ void CommKokkos::reverse_comm_device()
size_reverse_recv[iswap],MPI_DOUBLE, size_reverse_recv[iswap],MPI_DOUBLE,
sendproc[iswap],0,world,&request); sendproc[iswap],0,world,&request);
n = avec->pack_reverse_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send); n = avec->pack_reverse_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send);
DeviceType::fence();
if (n) if (n)
MPI_Send(k_buf_send.view<DeviceType>().ptr_on_device(),n, MPI_Send(k_buf_send.view<DeviceType>().ptr_on_device(),n,
MPI_DOUBLE,recvproc[iswap],0,world); MPI_DOUBLE,recvproc[iswap],0,world);
@ -331,6 +332,7 @@ void CommKokkos::reverse_comm_device()
} }
avec->unpack_reverse_kokkos(sendnum[iswap],k_sendlist,iswap, avec->unpack_reverse_kokkos(sendnum[iswap],k_sendlist,iswap,
k_buf_recv); k_buf_recv);
DeviceType::fence();
} else { } else {
if (sendnum[iswap]) if (sendnum[iswap])
n = avec->unpack_reverse_self(sendnum[iswap],k_sendlist,iswap, n = avec->unpack_reverse_self(sendnum[iswap],k_sendlist,iswap,
@ -933,7 +935,6 @@ void CommKokkos::borders_device() {
"implemented with Kokkos"); "implemented with Kokkos");
n = avec->pack_border_vel(nsend,sendlist[iswap],buf_send, n = avec->pack_border_vel(nsend,sendlist[iswap],buf_send,
pbc_flag[iswap],pbc[iswap]); pbc_flag[iswap],pbc[iswap]);
DeviceType::fence();
} }
else { else {
n = avec-> n = avec->