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.
| Type | Intent | Optional | Attributes | Name | ||
|---|---|---|---|---|---|---|
| class(SetInterface3D), | intent(inout) | :: | this | |||
| type(SetTraceRef), | intent(in) | :: | traces(1:) |
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