MortarExchange_MappedVector2D Subroutine

  1. SELF_MappedVector_2D.f90
  2. SELF_MappedVector_2D
  3. MortarExchange_MappedVector2D

public subroutine MortarExchange_MappedVector2D(this, mesh)

GPU implementation of the vector mortar exchange; the kernels treat the (variable, direction) pairs as 2*nvar independent trace lines.

Arguments

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

Calls

proc~~mortarexchange_mappedvector2d~~CallsGraph proc~mortarexchange_mappedvector2d MortarExchange_MappedVector2D interface~hipmalloc hipMalloc proc~mortarexchange_mappedvector2d->interface~hipmalloc interface~mortarflip_2d_gpu MortarFlip_2D_gpu proc~mortarexchange_mappedvector2d->interface~mortarflip_2d_gpu interface~mortarscatter_2d_gpu MortarScatter_2D_gpu proc~mortarexchange_mappedvector2d->interface~mortarscatter_2d_gpu proc~gpucheck gpuCheck proc~mortarexchange_mappedvector2d->proc~gpucheck interface~mortargather_2d_gpu MortarGather_2D_gpu proc~mortarexchange_mappedvector2d->interface~mortargather_2d_gpu

Contents

Source Code

MortarExchange_MappedVector2D

Source Code

  subroutine MortarExchange_MappedVector2D(this,mesh)
    !! GPU implementation of the vector mortar exchange; the kernels treat the
    !! (variable, direction) pairs as 2*nvar independent trace lines.
    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%MPIMortarExchangeAsync(mesh)
    endif

    call MortarGather_2D_gpu(this%mortarBuff_gpu,this%boundary_gpu, &
                             mesh%mortarInfo_gpu,mesh%decomp%elemToRank_gpu, &
                             mesh%decomp%rankId,offset,this%interp%N,2*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,2*this%nvar,mesh%nMortars)
    endif

    call MortarScatter_2D_gpu(this%extBoundary_gpu,this%mortarBuff_gpu, &
                              this%interp%mortarR_gpu,this%interp%mortarP_gpu, &
                              mesh%mortarInfo_gpu,mesh%decomp%elemToRank_gpu, &
                              mesh%decomp%rankId,offset,this%interp%N,2*this%nvar, &
                              mesh%nMortars,this%nelem)

  endsubroutine MortarExchange_MappedVector2D