MPIExchangeAsync_MappedVector3D Subroutine

  1. SELF_MappedVector_3D.f90
  2. SELF_MappedVector_3D
  3. MPIExchangeAsync_MappedVector3D

public subroutine MPIExchangeAsync_MappedVector3D(this, mesh)

Post the aggregated halo exchange: one MPI_Irecv/MPI_Isend pair per neighboring rank, carrying every (side,variable,component) boundary trace shared with that rank in a single packed device buffer. The boundary array is laid out with the component index outermost, so the pack/unpack kernels treat the vector as 3*nvar scalar variables. Packed buffers are allocated on first use; the shared side tables are built by SideExchange before this is called.

Arguments

TypeIntentOptionalAttributesName
class(MappedVector3D), intent(inout) :: this
type(Mesh3D), intent(inout) :: mesh

Calls

proc~~mpiexchangeasync_mappedvector3d~~CallsGraph proc~mpiexchangeasync_mappedvector3d MPIExchangeAsync_MappedVector3D interface~hipmalloc hipMalloc proc~mpiexchangeasync_mappedvector3d->interface~hipmalloc mpi_irecv mpi_irecv proc~mpiexchangeasync_mappedvector3d->mpi_irecv proc~gpucheck gpuCheck proc~mpiexchangeasync_mappedvector3d->proc~gpucheck interface~halopack_3d_gpu HaloPack_3D_gpu proc~mpiexchangeasync_mappedvector3d->interface~halopack_3d_gpu mpi_isend mpi_isend proc~mpiexchangeasync_mappedvector3d->mpi_isend

Contents

Source Code

MPIExchangeAsync_MappedVector3D

Source Code

  subroutine MPIExchangeAsync_MappedVector3D(this,mesh)
  !! Post the aggregated halo exchange: one MPI_Irecv/MPI_Isend pair per
  !! neighboring rank, carrying every (side,variable,component) boundary
  !! trace shared with that rank in a single packed device buffer. The
  !! boundary array is laid out with the component index outermost, so the
  !! pack/unpack kernels treat the vector as 3*nvar scalar variables. Packed
  !! buffers are allocated on first use; the shared side tables are built by
  !! SideExchange before this is called.
    implicit none
    class(MappedVector3D),intent(inout) :: this
    type(Mesh3D),intent(inout) :: mesh
    ! Local
    integer :: n,npts,cnt,disp
    integer :: iError
    integer :: msgCount
    integer(c_size_t) :: worksize
    real(prec),pointer :: sendbuf(:)
    real(prec),pointer :: recvbuf(:)

    npts = (this%interp%N+1)*(this%interp%N+1)*3*this%nvar

    if(.not. c_associated(this%halo_sendbuf_gpu)) then
      worksize = int(mesh%decomp%halo_nsides,c_size_t)* &
                 int(npts,c_size_t)*prec
      call gpuCheck(hipMalloc(this%halo_sendbuf_gpu,worksize))
      call gpuCheck(hipMalloc(this%halo_recvbuf_gpu,worksize))
    endif

    call HaloPack_3D_gpu(this%boundary_gpu,this%halo_sendbuf_gpu, &
                         mesh%decomp%halo_sides_gpu,this%interp%N,3*this%nvar, &
                         this%nelem,mesh%decomp%halo_nsides)

    call c_f_pointer(this%halo_sendbuf_gpu,sendbuf,[mesh%decomp%halo_nsides*npts])
    call c_f_pointer(this%halo_recvbuf_gpu,recvbuf,[mesh%decomp%halo_nsides*npts])

    msgCount = 0
    do n = 1,mesh%decomp%halo_nnbr

      cnt = (mesh%decomp%halo_offset(n+1)-mesh%decomp%halo_offset(n))*npts
      disp = mesh%decomp%halo_offset(n)*npts

      msgCount = msgCount+1
      call MPI_IRECV(recvbuf(disp+1),cnt, &
                     mesh%decomp%mpiPrec, &
                     mesh%decomp%halo_rank(n),0, &
                     mesh%decomp%mpiComm, &
                     mesh%decomp%requests(msgCount),iError)

      msgCount = msgCount+1
      call MPI_ISEND(sendbuf(disp+1),cnt, &
                     mesh%decomp%mpiPrec, &
                     mesh%decomp%halo_rank(n),0, &
                     mesh%decomp%mpiComm, &
                     mesh%decomp%requests(msgCount),iError)

    enddo

    mesh%decomp%msgCount = msgCount

  endsubroutine MPIExchangeAsync_MappedVector3D