DownloadOldWindow_DGModel2D Subroutine

  subroutine DownloadOldWindow_DGModel2D(this,wFirst,wLast,uWin)
    !! Copy the device-resident migrated window to the host, for the SELF_AMR_MIGRATE_VERIFY
    !! diagnostic only. This is exactly the host traffic the device path exists to avoid, so it
    !! must never appear on the default path.
    !!
    !! The comparison it feeds stays BITWISE. Everything that touched these values is data
    !! movement - a device-to-device copy, MPI byte transfers, and this download - so no
    !! arithmetic has been applied and exactness is available. The transfer APPLY is a different
    !! matter: it agrees between host and device only to round-off, because the device compiler
    !! contracts its multiply-accumulates into FMAs. The two must not be conflated.
    implicit none
    class(DGModel2D),intent(in) :: this
    integer,intent(in) :: wFirst
    integer,intent(in) :: wLast
    ! target: c_loc below requires it, and contiguous keeps the download a single memcpy.
    real(prec),intent(out),target,contiguous :: uWin(:,:,:,:)
    ! Local
    integer :: perElem,nWinElem
    integer(c_size_t) :: nbytes

    nWinElem = max(wLast-wFirst+1,0)
    if(nWinElem == 0) return

    perElem = (this%solution%interp%N+1)*(this%solution%interp%N+1)
    if(size(uWin,3) /= nWinElem .or. size(uWin,4) /= this%nvar) then
      print*,__FILE__,':',__LINE__, &
        ' : Error : DownloadOldWindow given a buffer that does not match the window.'
      stop 1
    endif
    if(this%xferNWin /= nWinElem .or. this%xferWinFirst /= wFirst) then
      print*,__FILE__,':',__LINE__, &
        ' : Error : DownloadOldWindow called without a matching migrated window.'
      stop 1
    endif

    nbytes = int(perElem,c_size_t)*int(nWinElem,c_size_t)*int(this%nvar,c_size_t)*prec
    call gpuCheck(hipMemcpy(c_loc(uWin),this%xferWin_gpu,nbytes,hipMemcpyDeviceToHost))

  endsubroutine DownloadOldWindow_DGModel2D