MortarFluxCollect_MappedVector3D Subroutine

  1. SELF_MappedVector_3D.f90
  2. SELF_MappedVector_3D
  3. MortarFluxCollect_MappedVector3D

public subroutine MortarFluxCollect_MappedVector3D(this, mesh)

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

Arguments

TypeIntentOptionalAttributesName
class(MappedVector3D), intent(inout) :: this
type(Mesh3D), intent(inout) :: mesh

Calls

proc~~mortarfluxcollect_mappedvector3d~~CallsGraph proc~mortarfluxcollect_mappedvector3d MortarFluxCollect_MappedVector3D interface~hipmalloc hipMalloc proc~mortarfluxcollect_mappedvector3d->interface~hipmalloc proc~gpucheck gpuCheck proc~mortarfluxcollect_mappedvector3d->proc~gpucheck interface~mortarflip_3d_gpu MortarFlip_3D_gpu proc~mortarfluxcollect_mappedvector3d->interface~mortarflip_3d_gpu interface~mortargather_3d_gpu MortarGather_3D_gpu proc~mortarfluxcollect_mappedvector3d->interface~mortargather_3d_gpu interface~mortarfluxscatter_3d_gpu MortarFluxScatter_3D_gpu proc~mortarfluxcollect_mappedvector3d->interface~mortarfluxscatter_3d_gpu

Contents

Source Code

MortarFluxCollect_MappedVector3D

Source Code

  subroutine MortarFluxCollect_MappedVector3D(this,mesh)
    !! GPU implementation of MortarFluxCollect (see the base class for the algorithm
    !! and conservation statement). Stages the small faces' boundaryNormal traces in
    !! the mortar buffer, then overwrites the big face's integrand on device.
    implicit none
    class(MappedVector3D),intent(inout) :: this
    type(Mesh3D),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
      ! The MortarFlip_3D kernel stages one face through shared memory
      ! (MORTAR3D_MAXNP in SELF_Mortar.cpp bounds the block size).
      if(this%interp%N+1 > 16) then
        print*,__FILE__,' : Error : 3D mortar kernels support N+1 <= 16.'
        stop 1
      endif
      buffSize = int(this%interp%N+1,c_size_t)*(this%interp%N+1)*8* &
                 mesh%nMortars*this%nvar*3*prec
      call gpuCheck(hipMalloc(this%mortarBuff_gpu,buffSize))
    endif

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

    ! Stage rank-local small-face integrands (the big-face slots are gathered too
    ! but unused by the flux scatter)
    call MortarGather_3D_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_3D_gpu(this%mortarBuff_gpu,mesh%mortarInfo_gpu, &
                             mesh%decomp%elemToRank_gpu,mesh%decomp%rankId, &
                             this%interp%N,this%nvar,mesh%nMortars)
    endif

    call MortarFluxScatter_3D_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_MappedVector3D