Exchange_SetInterface3D Subroutine

public subroutine Exchange_SetInterface3D(this, traces)

As Exchange_SetInterface3D_t. If every set with local interface faces has a device trace, the outbox is gathered on the device; either way ghost is copied to the device afterwards, so FillSetInterfaceBoundary_gpu reads the current exchange.

Arguments

TypeIntentOptionalAttributesName
class(SetInterface3D), intent(inout) :: this
type(SetTraceRef), intent(in) :: traces(1:)

Calls

proc~~exchange_setinterface3d~~CallsGraph proc~exchange_setinterface3d Exchange_SetInterface3D interface~gathersetinterfaceoutbox_gpu GatherSetInterfaceOutbox_gpu proc~exchange_setinterface3d->interface~gathersetinterfaceoutbox_gpu proc~gpucheck gpuCheck proc~exchange_setinterface3d->proc~gpucheck interface~hipmemcpy hipMemcpy proc~exchange_setinterface3d->interface~hipmemcpy

Contents


Source Code

  subroutine Exchange_SetInterface3D(this,traces)
    !! As Exchange_SetInterface3D_t. If every set with local interface faces has a device
    !! trace, the outbox is gathered on the device; either way ghost is copied to the
    !! device afterwards, so FillSetInterfaceBoundary_gpu reads the current exchange.
    implicit none
    class(SetInterface3D),intent(inout) :: this
    type(SetTraceRef),intent(in) :: traces(1:)
    ! Local
    integer :: k,nk
    logical :: onDevice

    if(this%nFaces == 0) then
      call this%ExchangeOutbox()
      return
    endif

    onDevice = .true.
    do k = 1,this%nSets
      nk = this%setFaceStart(k+1)-this%setFaceStart(k)
      if(nk > 0 .and. .not. c_associated(traces(k)%trace_gpu)) onDevice = .false.
    enddo

    if(onDevice) then
      do k = 1,this%nSets
        nk = this%setFaceStart(k+1)-this%setFaceStart(k)
        if(nk == 0) cycle
        call GatherSetInterfaceOutbox_gpu(this%outbox_gpu,traces(k)%trace_gpu, &
                                          this%faceElem_gpu,this%faceSide_gpu, &
                                          this%setFaceStart(k)-1,nk,this%nvarOfSet(k), &
                                          this%nvarMax,this%N,size(traces(k)%trace,4))
      enddo
      call gpuCheck(hipMemcpy(c_loc(this%outbox),this%outbox_gpu,sizeof(this%outbox), &
                              hipMemcpyDeviceToHost))
    else
      call this%GatherOutbox(traces)
    endif

    call this%ExchangeOutbox()

    call gpuCheck(hipMemcpy(this%ghost_gpu,c_loc(this%ghost),sizeof(this%ghost), &
                            hipMemcpyHostToDevice))

  endsubroutine Exchange_SetInterface3D