void CommKokkos::reverse_comm_device() { int n; MPI_Request request; AtomVecKokkos *avec = (AtomVecKokkos *) atom->avec; double *buf; // exchange data with another proc // if other proc is self, just copy // if comm_f_only set, exchange or copy directly from f, don't pack k_sendlist.sync<DeviceType>(); atomKK->sync(ExecutionSpaceFromDevice<DeviceType>::space,F_MASK); for (int iswap = nswap-1; iswap >= 0; iswap--) { if (sendproc[iswap] != me) { if (comm_f_only) { if (size_reverse_recv[iswap]) MPI_Irecv(k_buf_recv.view<DeviceType>().data(),size_reverse_recv[iswap],MPI_DOUBLE, sendproc[iswap],0,world,&request); if (size_reverse_send[iswap]) { buf = atomKK->k_f.view<DeviceType>().data() + firstrecv[iswap]*atomKK->k_f.view<DeviceType>().extent(1); MPI_Send(buf,size_reverse_send[iswap],MPI_DOUBLE, recvproc[iswap],0,world); } if (size_reverse_recv[iswap]) { MPI_Wait(&request,MPI_STATUS_IGNORE); atomKK->modified(ExecutionSpaceFromDevice<DeviceType>:: space,F_MASK); } } else { if (size_reverse_recv[iswap]) MPI_Irecv(k_buf_recv.view<DeviceType>().data(), size_reverse_recv[iswap],MPI_DOUBLE, sendproc[iswap],0,world,&request); n = avec->pack_reverse_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_send); DeviceType::fence(); if (n) MPI_Send(k_buf_send.view<DeviceType>().data(),n, MPI_DOUBLE,recvproc[iswap],0,world); if (size_reverse_recv[iswap]) MPI_Wait(&request,MPI_STATUS_IGNORE); } avec->unpack_reverse_kokkos(sendnum[iswap],k_sendlist,iswap, k_buf_recv); DeviceType::fence(); } else { if (sendnum[iswap]) n = avec->unpack_reverse_self(sendnum[iswap],k_sendlist,iswap, firstrecv[iswap]); } } }