ApplyTransferPlan_DGModel2D Subroutine

  subroutine ApplyTransferPlan_DGModel2D(this,plan,interp,eFirst,eLast,uGlobal)
    !! Apply the transfer plan on the device (Stage 6a), writing solution%interior_gpu directly
    !! and moving no solution data across the PCIe/xGMI link.
    !!
    !! uGlobal is the multi-rank escape hatch: when the caller has assembled a global old field
    !! on the host (the Stage-5 v1 allgather migration), this falls back to the portable host
    !! path. A device transfer only pays off on several ranks together with a point-to-point
    !! (Stage-5 v2) migration; on a single rank it stands alone, which is the case optimized
    !! here.
    implicit none
    class(DGModel2D),intent(inout) :: this
    type(TransferPlan2D),intent(in),target :: plan
    type(Lagrange),intent(in) :: interp
    integer,intent(in) :: eFirst
    integer,intent(in) :: eLast
    real(prec),intent(in),optional :: uGlobal(:,:,:,:)
    ! Local
    integer :: nLocal,pathStride
    integer(c_size_t) :: nb

    if(present(uGlobal)) then
      call ApplyTransferPlan_DGModel2D_t(this,plan,interp,eFirst,eLast,uGlobal)
      return
    endif

    if(.not. c_associated(this%xferOld_gpu)) then
      print*,__FILE__,':',__LINE__, &
        ' : Error : ApplyTransferPlan called without a staged solution.'
      stop 1
    endif

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

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

    ! (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,4_c_size_t*nb))
      call gpuCheck(hipMalloc(this%xferPath_gpu,int(pathStride,c_size_t)*nb))

      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), &
                            4_c_size_t*nb,hipMemcpyHostToDevice))
    call gpuCheck(hipMemcpy(this%xferPath_gpu,c_loc(plan%path), &
                            int(pathStride,c_size_t)*nb,hipMemcpyHostToDevice))

    call TransferSolution_2D_gpu(this%xferOld_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,interp%N,this%nvar, &
                                 this%xferNOld,this%solution%nElem,nLocal)

  endsubroutine ApplyTransferPlan_DGModel2D