StageSolutionForTransfer_DGModel2D Subroutine

  1. SELF_DGModel2D.f90
  2. SELF_DGModel2D
  3. StageSolutionForTransfer_DGModel2D

public subroutine StageSolutionForTransfer_DGModel2D(this)

Stage the pre-regrid solution on the DEVICE (Stage 6a), replacing the base implementation's device-to-host copy of the whole field. Regrid subsequently releases solution%interior_gpu, so the field is copied device-to-device into a buffer this model owns; ApplyTransferPlan then reads it from there.

The staging buffer is grown monotonically and reused, so an adapting run allocates here only when the element count exceeds every previous epoch's.

Arguments

TypeIntentOptionalAttributesName
class(DGModel2D), intent(inout) :: this

Calls

proc~~stagesolutionfortransfer_dgmodel2d~~CallsGraph proc~stagesolutionfortransfer_dgmodel2d StageSolutionForTransfer_DGModel2D interface~hipfree hipFree proc~stagesolutionfortransfer_dgmodel2d->interface~hipfree proc~gpucheck gpuCheck proc~stagesolutionfortransfer_dgmodel2d->proc~gpucheck interface~hipmemcpy hipMemcpy proc~stagesolutionfortransfer_dgmodel2d->interface~hipmemcpy interface~hipmalloc hipMalloc proc~stagesolutionfortransfer_dgmodel2d->interface~hipmalloc

Contents

Source Code

StageSolutionForTransfer_DGModel2D

Source Code

  subroutine StageSolutionForTransfer_DGModel2D(this)
    !! Stage the pre-regrid solution on the DEVICE (Stage 6a), replacing the base
    !! implementation's device-to-host copy of the whole field. Regrid subsequently releases
    !! solution%interior_gpu, so the field is copied device-to-device into a buffer this model
    !! owns; ApplyTransferPlan then reads it from there.
    !!
    !! The staging buffer is grown monotonically and reused, so an adapting run allocates here
    !! only when the element count exceeds every previous epoch's.
    implicit none
    class(DGModel2D),intent(inout) :: this
    ! Local
    integer(c_size_t) :: nbytes

    nbytes = int(this%solution%interp%N+1,c_size_t)*(this%solution%interp%N+1)* &
             this%solution%nElem*this%nvar*prec

    if(nbytes > this%xferAllocBytes) then
      if(c_associated(this%xferOld_gpu)) call gpuCheck(hipFree(this%xferOld_gpu))
      call gpuCheck(hipMalloc(this%xferOld_gpu,nbytes))
      this%xferAllocBytes = nbytes
    endif

    call gpuCheck(hipMemcpy(this%xferOld_gpu,this%solution%interior_gpu,nbytes, &
                            hipMemcpyDeviceToDevice))

    ! Remember the staged field's element count: it is the stride the transfer kernel must use
    ! to index the staged buffer, and it is no longer recoverable from the model after Regrid.
    this%xferNOld = this%solution%nElem

  endsubroutine StageSolutionForTransfer_DGModel2D