module m_omptgt_sendrecv !! Halo exchange of device buffers for the OpenMP target backend, following !! cuda/sendrecv.f90. !! !! With more than one rank the device addresses go straight to MPI, which !! therefore has to be GPU-aware, as it has to be for the CUDA backend. !! UCX tells device from host buffers with a memory type cache fed by !! intercepting CUDA allocations, and it can miss those the OpenMP runtime !! makes: it then reads the buffers as host memory, failing in !! process_vm_readv with "Bad address". Run with UCX_MEMTYPE_CACHE=n so UCX !! queries the driver for every buffer instead. !! The send buffers are filled by a kernel that the runtime may still have !! queued (see device_sync), so the device is synchronised before MPI !! reads them, as the CUDA backend synchronises its stream. use iso_c_binding, only: c_loc use m_mpi, only: MPI_COMM_WORLD, MPI_STATUSES_IGNORE, MPI_Irecv, MPI_Isend, & MPI_Waitall use m_common, only: dp, MPI_X3D2_DP use m_omptgt_common, only: memcpy_d2d, device_sync implicit none contains subroutine sendrecv_fields(f_recv_s, f_recv_e, f_send_s, f_send_e, & n_data, nproc, prev, next) implicit none real(dp), pointer, contiguous, dimension(:, :, :), intent(in) :: & f_recv_s, f_recv_e, f_send_s, f_send_e integer, intent(in) :: n_data, nproc, prev, next integer :: req(4), err(4), ierr, tag = 1234 if (nproc == 1) then call memcpy_d2d(c_loc(f_recv_s), c_loc(f_send_e), n_data) call memcpy_d2d(c_loc(f_recv_e), c_loc(f_send_s), n_data) else call device_sync() #ifdef MPI call MPI_Isend(f_send_s, n_data, MPI_X3D2_DP, & prev, tag, MPI_COMM_WORLD, req(1), err(1)) call MPI_Irecv(f_recv_e, n_data, MPI_X3D2_DP, & next, tag, MPI_COMM_WORLD, req(2), err(2)) call MPI_Isend(f_send_e, n_data, MPI_X3D2_DP, & next, tag, MPI_COMM_WORLD, req(3), err(3)) call MPI_Irecv(f_recv_s, n_data, MPI_X3D2_DP, & prev, tag, MPI_COMM_WORLD, req(4), err(4)) call MPI_Waitall(4, req, MPI_STATUSES_IGNORE, ierr) #endif end if end subroutine sendrecv_fields subroutine sendrecv_3fields( & f1_recv_s, f1_recv_e, f2_recv_s, f2_recv_e, f3_recv_s, f3_recv_e, & f1_send_s, f1_send_e, f2_send_s, f2_send_e, f3_send_s, f3_send_e, & n_data, nproc, prev, next & ) implicit none real(dp), pointer, contiguous, dimension(:, :, :), intent(in) :: & f1_recv_s, f1_recv_e, f2_recv_s, f2_recv_e, f3_recv_s, f3_recv_e, & f1_send_s, f1_send_e, f2_send_s, f2_send_e, f3_send_s, f3_send_e integer, intent(in) :: n_data, nproc, prev, next integer :: req(12), err(12), ierr, tag = 1234 if (nproc == 1) then call memcpy_d2d(c_loc(f1_recv_s), c_loc(f1_send_e), n_data) call memcpy_d2d(c_loc(f1_recv_e), c_loc(f1_send_s), n_data) call memcpy_d2d(c_loc(f2_recv_s), c_loc(f2_send_e), n_data) call memcpy_d2d(c_loc(f2_recv_e), c_loc(f2_send_s), n_data) call memcpy_d2d(c_loc(f3_recv_s), c_loc(f3_send_e), n_data) call memcpy_d2d(c_loc(f3_recv_e), c_loc(f3_send_s), n_data) else call device_sync() #ifdef MPI call MPI_Isend(f1_send_s, n_data, MPI_X3D2_DP, & prev, tag, MPI_COMM_WORLD, req(1), err(1)) call MPI_Irecv(f1_recv_e, n_data, MPI_X3D2_DP, & next, tag, MPI_COMM_WORLD, req(2), err(2)) call MPI_Isend(f1_send_e, n_data, MPI_X3D2_DP, & next, tag, MPI_COMM_WORLD, req(3), err(3)) call MPI_Irecv(f1_recv_s, n_data, MPI_X3D2_DP, & prev, tag, MPI_COMM_WORLD, req(4), err(4)) call MPI_Isend(f2_send_s, n_data, MPI_X3D2_DP, & prev, tag, MPI_COMM_WORLD, req(5), err(5)) call MPI_Irecv(f2_recv_e, n_data, MPI_X3D2_DP, & next, tag, MPI_COMM_WORLD, req(6), err(6)) call MPI_Isend(f2_send_e, n_data, MPI_X3D2_DP, & next, tag, MPI_COMM_WORLD, req(7), err(7)) call MPI_Irecv(f2_recv_s, n_data, MPI_X3D2_DP, & prev, tag, MPI_COMM_WORLD, req(8), err(8)) call MPI_Isend(f3_send_s, n_data, MPI_X3D2_DP, & prev, tag, MPI_COMM_WORLD, req(9), err(9)) call MPI_Irecv(f3_recv_e, n_data, MPI_X3D2_DP, & next, tag, MPI_COMM_WORLD, req(10), err(10)) call MPI_Isend(f3_send_e, n_data, MPI_X3D2_DP, & next, tag, MPI_COMM_WORLD, req(11), err(11)) call MPI_Irecv(f3_recv_s, n_data, MPI_X3D2_DP, & prev, tag, MPI_COMM_WORLD, req(12), err(12)) call MPI_Waitall(12, req, MPI_STATUSES_IGNORE, ierr) #endif end if end subroutine sendrecv_3fields end module m_omptgt_sendrecv