MortarExchange_MappedScalar2D Subroutine

  1. SELF_MappedScalar_2D.f90
  2. SELF_MappedScalar_2D
  3. MortarExchange_MappedScalar2D

public subroutine MortarExchange_MappedScalar2D(this, mesh)

GPU implementation of the mortar exchange (see the base class for the algorithm) : traces are staged, reoriented, restricted, and projected in device memory with the SELF_Mortar kernels.

Arguments

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

Calls

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

Contents

Source Code

MortarExchange_MappedScalar2D

Source Code

  subroutine MortarExchange_MappedScalar2D(this,mesh)
    !! GPU implementation of the mortar exchange (see the base class for the
    !! algorithm) : traces are staged, reoriented, restricted, and projected in
    !! device memory with the SELF_Mortar kernels.
    implicit none
    class(MappedScalar2D),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*prec
      call gpuCheck(hipMalloc(this%mortarBuff_gpu,buffSize))
    endif

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

    ! Stage rank-local traces in the big side's edge orientation
    call MortarGather_2D_gpu(this%mortarBuff_gpu,this%boundary_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()
      ! Reorient small-side traces received over MPI
      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

    ! Small sides get the restricted big-side trace; the big side gets the L2
    ! projection of the small-side traces
    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,this%nvar, &
                              mesh%nMortars,this%nelem)

  endsubroutine MortarExchange_MappedScalar2D