ApplyTransferPlan_DGModel3D Subroutine

  subroutine ApplyTransferPlan_DGModel3D(this,plan,interp,eFirst,eLast,uGlobal,oldFirst)
    !! Apply the transfer plan on the device, writing solution%interior_gpu directly and moving
    !! no solution data across the PCIe/xGMI link - on any number of ranks.
    !!
    !! The old field is read from whichever device buffer the pairing routine left it in: the
    !! whole staged local field on a single rank (StageSolutionForTransfer), or this rank's
    !! migrated window on several (MigrateOldWindow), in which case the kernel is given the
    !! window's element stride and the global old index of its first element so it can rebase the
    !! plan's source indices. The two are mutually exclusive and one of them is required.
    !!
    !! uGlobal remains for a caller that supplies host-side old-field data - the Stage-5 v1
    !! allgather migration, and external callers - and takes the portable host path, forwarding
    !! oldFirst so the window is indexed correctly.
    implicit none
    ! target: matches the base declaration, which remaps a pointer onto its window buffer.
    class(DGModel3D),intent(inout),target :: this
    type(TransferPlan3D),intent(in),target :: plan
    type(Lagrange),intent(in) :: interp
    integer,intent(in) :: eFirst
    integer,intent(in) :: eLast
    real(prec),intent(in),optional,contiguous :: uGlobal(:,:,:,:,:)
    integer,intent(in),optional :: oldFirst
    ! Local
    integer :: nLocal,pathStride,li,c,src,o1,o2
    integer :: nOldStride,oldFirst0
    integer(c_size_t) :: nb
    type(c_ptr) :: uOld_gpu

    if(present(uGlobal)) then
      ! Caller-supplied host old field: the Stage-5 v1 allgather migration, and any external
      ! caller. Never combined with a migrated device window - that would be two sources for one
      ! apply, so it is a caller bug rather than something to resolve silently.
      if(this%xferNWin > 0) then
        print*,__FILE__,':',__LINE__, &
          ' : Error : ApplyTransferPlan given both uGlobal and a migrated window.'
        stop 1
      endif
      call ApplyTransferPlan_DGModel3D_t(this,plan,interp,eFirst,eLast,uGlobal,oldFirst)
      return
    endif

    ! Which device buffer holds the old field, and how the kernel indexes it.
    !
    ! The guards are on the element counts, not on c_associated(). Both buffers are persistent
    ! capacity buffers that stay allocated after use, so testing a pointer would only catch a
    ! missing stage before the FIRST epoch; from the second on, an unpaired call would silently
    ! transfer the previous epoch's field. xferNOld and xferNWin are set by
    ! StageSolutionForTransfer and MigrateOldWindow respectively and both cleared below, so they
    ! track the pairing the way the base implementation's allocatable transferStage does.
    if(this%xferNWin > 0 .and. this%xferNOld > 0) then
      print*,__FILE__,':',__LINE__, &
        ' : Error : ApplyTransferPlan has both a staged local field and a migrated window.'
      stop 1
    elseif(this%xferNWin > 0) then
      ! Multi-rank: a migrated window of the old field, first element global old xferWinFirst.
      uOld_gpu = this%xferWin_gpu
      nOldStride = this%xferNWin
      oldFirst0 = this%xferWinFirst-1
    elseif(this%xferNOld > 0) then
      ! Single rank: the whole local old field, so the window base is zero and the kernel's
      ! indexing is exactly what it was before the window form existed.
      uOld_gpu = this%xferOld_gpu
      nOldStride = this%xferNOld
      oldFirst0 = 0
    else
      print*,__FILE__,':',__LINE__, &
        ' : Error : ApplyTransferPlan called without a staged solution or a migrated window.'
      stop 1
    endif

    ! The device kernel holds three (N+1)^3 working buffers in shared memory, sized to a compile
    ! time bound; guard the degree here rather than overrunning them.
    if(interp%N+1 > 12) then
      print*,__FILE__,':',__LINE__, &
        ' : Error : the device solution transfer supports N+1 <= AMR3D_MAXNP (12).'
      stop 1
    endif

    nLocal = eLast-eFirst+1
    pathStride = size(plan%path,1)

    ! Window coverage. On the host path ApplyTransferPlanWindow checks every element's sources
    ! against the window and stops if one falls outside it. The kernel cannot: it has no way to
    ! report, and an out-of-window index there is a silent out-of-bounds device read rather than a
    ! wrong answer you can see. The same conditions are checked here instead - integer comparisons
    ! over this rank's own element range, once per adapting epoch.
    if(this%xferNWin > 0) then
      o1 = this%xferWinFirst
      o2 = o1+this%xferNWin-1
      do li = eFirst,eLast
        if(plan%sourceKind(li) == SELF_TRANSFER_RESTRICT) then
          do c = 1,8
            src = plan%family(c,li)
            if(src < o1 .or. src > o2) then
              print*,__FILE__,':',__LINE__, &
                ' : Error : a restricted child lies outside the migrated window.'
              stop 1
            endif
          enddo
        else
          src = plan%sourceElem(li)
          if(src < o1 .or. src > o2) then
            print*,__FILE__,':',__LINE__, &
              ' : Error : a transfer source lies outside the migrated window.'
            stop 1
          endif
        endif
      enddo
    endif

    ! (Re)size the device copies of the plan's integer arrays, reusing them across epochs.
    if(plan%nNew > this%planAllocElem .or. pathStride > this%planAllocStride) then
      if(c_associated(this%xferKind_gpu)) call gpuCheck(hipFree(this%xferKind_gpu))
      if(c_associated(this%xferElem_gpu)) call gpuCheck(hipFree(this%xferElem_gpu))
      if(c_associated(this%xferFamily_gpu)) call gpuCheck(hipFree(this%xferFamily_gpu))
      if(c_associated(this%xferDepth_gpu)) call gpuCheck(hipFree(this%xferDepth_gpu))
      if(c_associated(this%xferPath_gpu)) call gpuCheck(hipFree(this%xferPath_gpu))

      nb = int(plan%nNew,c_size_t)*c_int
      call gpuCheck(hipMalloc(this%xferKind_gpu,nb))
      call gpuCheck(hipMalloc(this%xferElem_gpu,nb))
      call gpuCheck(hipMalloc(this%xferDepth_gpu,nb))
      call gpuCheck(hipMalloc(this%xferFamily_gpu,8_c_size_t*nb))
      call gpuCheck(hipMalloc(this%xferPath_gpu,int(pathStride,c_size_t)*nb))

      ! All five are freed and re-allocated together, so both capacities describe the same set
      ! of buffers. They track the last plan rather than a high-water mark (as in the 2-D
      ! implementation this mirrors): a later epoch that grows either dimension re-allocates,
      ! and a settled mesh - the case that matters - trips neither condition.
      this%planAllocElem = plan%nNew
      this%planAllocStride = pathStride
    endif

    nb = int(plan%nNew,c_size_t)*c_int
    call gpuCheck(hipMemcpy(this%xferKind_gpu,c_loc(plan%sourceKind), &
                            nb,hipMemcpyHostToDevice))
    call gpuCheck(hipMemcpy(this%xferElem_gpu,c_loc(plan%sourceElem), &
                            nb,hipMemcpyHostToDevice))
    call gpuCheck(hipMemcpy(this%xferDepth_gpu,c_loc(plan%depth), &
                            nb,hipMemcpyHostToDevice))
    call gpuCheck(hipMemcpy(this%xferFamily_gpu,c_loc(plan%family), &
                            8_c_size_t*nb,hipMemcpyHostToDevice))
    call gpuCheck(hipMemcpy(this%xferPath_gpu,c_loc(plan%path), &
                            int(pathStride,c_size_t)*nb,hipMemcpyHostToDevice))

    call TransferSolution_3D_gpu(uOld_gpu,this%solution%interior_gpu, &
                                 this%xferKind_gpu,this%xferElem_gpu,this%xferFamily_gpu, &
                                 this%xferDepth_gpu,this%xferPath_gpu, &
                                 interp%mortarR_gpu,interp%mortarP_gpu, &
                                 pathStride,eFirst-1,oldFirst0,interp%N,this%nvar, &
                                 nOldStride,this%solution%nElem,nLocal)

    ! The staged field has been consumed. The buffer itself is kept (it is the capacity that
    ! makes a settled run allocation-free); only the pairing marker is cleared.
    this%xferNOld = 0
    this%xferNWin = 0

  endsubroutine ApplyTransferPlan_DGModel3D