MortarFluxCollect_MappedVector2D Subroutine

  1. SELF_MappedVector_2D.f90
  2. SELF_MappedVector_2D
  3. MortarFluxCollect_MappedVector2D

public subroutine MortarFluxCollect_MappedVector2D(this, mesh)

GPU implementation of MortarFluxCollect (see the base class for the algorithm and conservation statement). Stages the small sides' boundaryNormal traces in the mortar buffer, then overwrites the big side's integrand on device.

Arguments

TypeIntentOptionalAttributesName
class(MappedVector2D), intent(inout) :: this
type(Mesh2D), intent(inout) :: mesh

Calls

proc~~mortarfluxcollect_mappedvector2d~~CallsGraph proc~mortarfluxcollect_mappedvector2d MortarFluxCollect_MappedVector2D interface~hipmalloc hipMalloc proc~mortarfluxcollect_mappedvector2d->interface~hipmalloc interface~mortarflip_2d_gpu MortarFlip_2D_gpu proc~mortarfluxcollect_mappedvector2d->interface~mortarflip_2d_gpu interface~mortarfluxscatter_2d_gpu MortarFluxScatter_2D_gpu proc~mortarfluxcollect_mappedvector2d->interface~mortarfluxscatter_2d_gpu proc~gpucheck gpuCheck proc~mortarfluxcollect_mappedvector2d->proc~gpucheck interface~mortargather_2d_gpu MortarGather_2D_gpu proc~mortarfluxcollect_mappedvector2d->interface~mortargather_2d_gpu

Contents

Source Code

MortarFluxCollect_MappedVector2D

Source Code

  subroutine MortarFluxCollect_MappedVector2D(this,mesh)
    !! GPU implementation of MortarFluxCollect (see the base class for the algorithm
    !! and conservation statement). Stages the small sides' boundaryNormal traces in
    !! the mortar buffer, then overwrites the big side's integrand on device.
    implicit none
    class(MappedVector2D),intent(inout) :: this
    type(Mesh2D),intent(inout) :: mesh
    ! Local
    integer :: offset
    integer(c_size_t) :: buffSize

    offset = mesh%decomp%offsetElem(mesh%decomp%rankId+1)

    if(.not. c_associated(this%mortarBuff_gpu)) then
      buffSize = int(this%interp%N+1,c_size_t)*4*mesh%nMortars*this%nvar*2*prec
      call gpuCheck(hipMalloc(this%mortarBuff_gpu,buffSize))
    endif

    if(mesh%decomp%mpiEnabled) then
      call this%MPIMortarFluxAsync(mesh)
    endif

    ! Stage rank-local small-side integrands (the big-side slots are gathered too
    ! but unused by the flux scatter)
    call MortarGather_2D_gpu(this%mortarBuff_gpu,this%boundarynormal_gpu, &
                             mesh%mortarInfo_gpu,mesh%decomp%elemToRank_gpu, &
                             mesh%decomp%rankId,offset,this%interp%N,this%nvar, &
                             mesh%nMortars,this%nelem)

    if(mesh%decomp%mpiEnabled) then
      call mesh%decomp%FinalizeMPIExchangeAsync()
      call MortarFlip_2D_gpu(this%mortarBuff_gpu,mesh%mortarInfo_gpu, &
                             mesh%decomp%elemToRank_gpu,mesh%decomp%rankId, &
                             this%interp%N,this%nvar,mesh%nMortars)
    endif

    call MortarFluxScatter_2D_gpu(this%boundarynormal_gpu,this%mortarBuff_gpu, &
                                  this%interp%mortarP_gpu, &
                                  mesh%mortarInfo_gpu,mesh%decomp%elemToRank_gpu, &
                                  mesh%decomp%rankId,offset,this%interp%N,this%nvar, &
                                  mesh%nMortars,this%nelem)

  endsubroutine MortarFluxCollect_MappedVector2D