SideExchangeStart_MappedScalar2D Subroutine

  subroutine SideExchangeStart_MappedScalar2D(this,mesh,nactive)
  !! Begin the aggregated halo exchange and launch the local (same-rank)
  !! side exchange kernel.
  !!
  !! One persistent MPI_Recv_init/MPI_Send_init pair per neighboring rank is
  !! (re)created the first time this is called (and again if the number of
  !! exchanged variables changes), then re-armed with MPI_Startall on every
  !! subsequent exchange; receives are always started before sends.
  !!
  !! When `nactive` is provided, only the leading `nactive` variables are
  !! exchanged - the variable index is the outermost dimension of the
  !! boundary arrays, so the leading variables form a contiguous prefix.
  !! The first exchange always carries all variables so that the boundary
  !! traces of static (non-stepped) variables are valid; static traces do
  !! not change thereafter, so later exchanges may carry the prefix only.
  !!
  !! The matching SideExchangeFinish must be called before extBoundary is
  !! read on interior faces. Between Start and Finish, callers may perform
  !! any work that does not read interior-face extBoundary entries; the MPI
  !! messages and the local exchange kernel launched here proceed
  !! concurrently with that work.
    implicit none
    class(MappedScalar2D),intent(inout) :: this
    type(Mesh2D),intent(inout) :: mesh
    integer,intent(in),optional :: nactive
    ! Local
    integer :: n,nnbr,npts,cnt,disp,nact
    integer :: iError
    integer :: offset
    integer(c_size_t) :: worksize
    real(prec),pointer :: sendbuf(:)
    real(prec),pointer :: recvbuf(:)

    offset = mesh%decomp%offsetElem(mesh%decomp%rankId+1)

    if(mesh%decomp%mpiEnabled) then

      if(.not. mesh%decomp%halo_built) then
        call mesh%decomp%BuildHaloExchange(mesh%sideInfo,mesh%nElem,4)
      endif

      if(mesh%decomp%halo_nsides > 0) then

        nact = this%nvar
        if(present(nactive)) then
          if(this%halo_static_done .and. nactive >= 1 .and. nactive <= this%nvar) then
            nact = nactive
          endif
        endif

        if(.not. c_associated(this%halo_sendbuf_gpu)) then
          ! Buffers are sized for a full-variable exchange; prefix exchanges
          ! use the leading portion.
          worksize = int(mesh%decomp%halo_nsides,c_size_t)* &
                     int((this%interp%N+1)*this%nvar,c_size_t)*prec
          call gpuCheck(hipMalloc(this%halo_sendbuf_gpu,worksize))
          call gpuCheck(hipMalloc(this%halo_recvbuf_gpu,worksize))
        endif

        ! Pack the send buffer; the pack kernel synchronizes the device
        ! before returning so the buffer is complete before the sends start.
        call HaloPack_2D_gpu(this%boundary_gpu,this%halo_sendbuf_gpu, &
                             mesh%decomp%halo_sides_gpu,this%interp%N,nact, &
                             this%nelem,mesh%decomp%halo_nsides)

        nnbr = mesh%decomp%halo_nnbr

        if(nact /= this%halo_nactive) then
          ! (Re)create the persistent requests for this variable count
          if(allocated(this%halo_reqs)) then
            do n = 1,size(this%halo_reqs)
              call MPI_REQUEST_FREE(this%halo_reqs(n),iError)
            enddo
            deallocate(this%halo_reqs)
          endif
          allocate(this%halo_reqs(1:2*nnbr))
          npts = (this%interp%N+1)*nact
          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])
          do n = 1,nnbr
            cnt = (mesh%decomp%halo_offset(n+1)-mesh%decomp%halo_offset(n))*npts
            disp = mesh%decomp%halo_offset(n)*npts
            call MPI_RECV_INIT(recvbuf(disp+1),cnt,mesh%decomp%mpiPrec, &
                               mesh%decomp%halo_rank(n),0,mesh%decomp%mpiComm, &
                               this%halo_reqs(n),iError)
            call MPI_SEND_INIT(sendbuf(disp+1),cnt,mesh%decomp%mpiPrec, &
                               mesh%decomp%halo_rank(n),0,mesh%decomp%mpiComm, &
                               this%halo_reqs(nnbr+n),iError)
          enddo
          this%halo_nactive = nact
        endif

        ! Arm the receives before the sends
        call MPI_STARTALL(nnbr,this%halo_reqs(1:nnbr),iError)
        call MPI_STARTALL(nnbr,this%halo_reqs(nnbr+1:2*nnbr),iError)

        this%halo_inflight = nact

      endif
    endif

    ! The local (same-rank) side exchange runs on the device while the
    ! aggregated MPI messages are in flight.
    call SideExchange_2D_gpu(this%extboundary_gpu, &
                             this%boundary_gpu,mesh%sideinfo_gpu,mesh%decomp%elemToRank_gpu, &
                             mesh%decomp%rankid,offset,this%interp%N,this%nvar,this%nelem)

  endsubroutine SideExchangeStart_MappedScalar2D