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.
| Type | Intent | Optional | Attributes | Name | ||
|---|---|---|---|---|---|---|
| class(DGModel2D), | intent(inout), | target | :: | this | ||
| type(TransferPlan2D), | intent(in), | target | :: | plan | ||
| type(Lagrange), | intent(in) | :: | interp | |||
| integer, | intent(in) | :: | eFirst | |||
| integer, | intent(in) | :: | eLast | |||
| real(kind=prec), | intent(in), | optional | contiguous | :: | uGlobal(:,:,:,:) | |
| integer, | intent(in), | optional | :: | oldFirst |
subroutine ApplyTransferPlan_DGModel2D(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(DGModel2D),intent(inout),target :: 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,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_DGModel2D_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 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)
! 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,4
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,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(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_DGModel2D