MigrateOldWindow_DGModel2D Subroutine

  subroutine MigrateOldWindow_DGModel2D(this,winFirst,winLast,wFirst,wLast, &
                                        nBytesRecv,nBytesSent,nElemRemote)
    !! Assemble this rank's window of the pre-regrid old field IN DEVICE MEMORY, replacing the
    !! base implementation's host-memory migration. The rank's own run is copied device-to-device
    !! and the peers' runs are received directly into device memory, so no solution data crosses
    !! the host link and the windowed ApplyTransferPlan that follows runs as a device kernel.
    !!
    !! This hands device pointers to MPI, which requires a GPU-aware MPI. That is not a new
    !! requirement: the per-step aggregated halo exchange has always posted its sends and receives
    !! on device allocations (SideExchangeStart in SELF_MappedScalar_2D), with no host-staged
    !! fallback anywhere, and SELF_REQUIRE_GPU_AWARE_MPI makes the configure-time check fatal by
    !! default. A multi-rank GPU run that lacked it could never have taken a time step.
    !!
    !! Call it BEFORE Regrid, which releases the storage the sends read.
    implicit none
    class(DGModel2D),intent(inout) :: this
    integer,intent(in) :: winFirst(:)
    integer,intent(in) :: winLast(:)
    integer,intent(in) :: wFirst
    integer,intent(in) :: wLast
    integer(int64),intent(inout) :: nBytesRecv
    integer(int64),intent(inout) :: nBytesSent
    integer(int64),intent(inout) :: nElemRemote
    ! Local
    integer :: perElem,nWinElem,nLocalOld,myFirst,a,b,msgCount
    integer(c_size_t) :: nbytes
    integer(int64) :: nWinReals,nLocReals
    integer,allocatable :: requests(:)
    real(prec),pointer,contiguous :: win(:),uloc(:)

    perElem = (this%solution%interp%N+1)*(this%solution%interp%N+1)
    nWinElem = max(wLast-wFirst+1,0)
    nLocalOld = this%solution%nElem
    myFirst = this%mesh%decomp%offsetElem(this%mesh%decomp%rankId+1)+1

    ! Grow-only capacity, with a one-element floor so the descriptor below always has a target
    ! even for an empty window (a rank that owns no new elements still takes part, because its
    ! peers may need old elements it owns).
    nbytes = int(perElem,c_size_t)*int(max(nWinElem,1),c_size_t)*int(this%nvar,c_size_t)*prec
    if(nbytes > this%xferWinBytes) then
      if(c_associated(this%xferWin_gpu)) call gpuCheck(hipFree(this%xferWin_gpu))
      call gpuCheck(hipMalloc(this%xferWin_gpu,nbytes))
      this%xferWinBytes = nbytes
    endif

    ! A flat descriptor cannot span more reals than a default integer can index, and the products
    ! below are what would silently wrap: perElem*nWinElem*nvar is the whole buffer. Fail loudly
    ! instead. The same ceiling already binds the exchange, whose MPI counts are default integers.
    nWinReals = int(perElem,int64)*int(max(nWinElem,1),int64)*int(this%nvar,int64)
    nLocReals = int(perElem,int64)*int(max(nLocalOld,1),int64)*int(this%nvar,int64)
    if(max(nWinReals,nLocReals) > int(huge(1),int64)) then
      print*,__FILE__,':',__LINE__, &
        ' : Error : the migration window exceeds what a default integer can index.'
      stop 1
    endif

    ! Device allocations viewed as flat Fortran arrays, so the migration schedule can index them
    ! and MPI can be handed their addresses: the idiom SideExchangeStart already uses for the
    ! packed halo buffers.
    call c_f_pointer(this%xferWin_gpu,win,[int(nWinReals)])
    if(nLocalOld > 0) then
      call c_f_pointer(this%solution%interior_gpu,uloc,[int(nLocReals)])
    else
      ! This rank owned no old elements, so interior_gpu need not be a valid allocation. It posts
      ! no sends either (its old range intersects nobody's window), so uloc is never dereferenced;
      ! aim it somewhere valid rather than calling c_f_pointer on a possibly null pointer.
      uloc => win
    endif

    ! The run of my window that I already own, device-to-device. This ALSO synchronizes the
    ! device unconditionally - see MigrateWindowLocal_gpu - which is what makes it safe to post
    ! MPI against solution%interior_gpu on the next line. Unlike the host path, which overlaps
    ! its local copy with the messages, the copy happens first: the barrier has to precede the
    ! sends, and once per adapting epoch the lost overlap is not worth a second synchronization.
    call OwnedRun(this%mesh%decomp%offsetElem,this%mesh%decomp%rankId+1,wFirst,wLast,a,b)
    call MigrateWindowLocal_gpu(this%solution%interior_gpu,this%xferWin_gpu, &
                                perElem,this%nvar,max(nLocalOld,1),max(nWinElem,1), &
                                a-wFirst,a-myFirst,max(b-a+1,0))

    allocate(requests(1:2*this%nvar*this%mesh%decomp%nRanks))
    msgCount = 0
    call PostOldWindowExchange(this%mesh%decomp,perElem,this%nvar,nLocalOld,uloc, &
                               winFirst,winLast,wFirst,wLast,win,requests,msgCount, &
                               nBytesRecv,nBytesSent,nElemRemote)
    call FinishOldWindowExchange(requests,msgCount)
    deallocate(requests)

    win => null()
    uloc => null()
    this%xferWinFirst = wFirst
    this%xferNWin = nWinElem

  endsubroutine MigrateOldWindow_DGModel2D