AnisoMortarExchange_MappedScalar3D Subroutine

  1. SELF_MappedScalar_3D.f90
  2. SELF_MappedScalar_3D
  3. AnisoMortarExchange_MappedScalar3D

public subroutine AnisoMortarExchange_MappedScalar3D(this, mesh)

GPU implementation of the anisotropic 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(MappedScalar3D), intent(inout) :: this
type(Mesh3D), intent(inout) :: mesh

Calls

proc~~anisomortarexchange_mappedscalar3d~~CallsGraph proc~anisomortarexchange_mappedscalar3d AnisoMortarExchange_MappedScalar3D interface~hipmalloc hipMalloc proc~anisomortarexchange_mappedscalar3d->interface~hipmalloc interface~anisomortargather_3d_gpu AnisoMortarGather_3D_gpu proc~anisomortarexchange_mappedscalar3d->interface~anisomortargather_3d_gpu interface~anisomortarflip_3d_gpu AnisoMortarFlip_3D_gpu proc~anisomortarexchange_mappedscalar3d->interface~anisomortarflip_3d_gpu proc~gpucheck gpuCheck proc~anisomortarexchange_mappedscalar3d->proc~gpucheck interface~anisomortarscatter_3d_gpu AnisoMortarScatter_3D_gpu proc~anisomortarexchange_mappedscalar3d->interface~anisomortarscatter_3d_gpu

Contents

Source Code

AnisoMortarExchange_MappedScalar3D

Source Code

  subroutine AnisoMortarExchange_MappedScalar3D(this,mesh)
    !! GPU implementation of the anisotropic 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(MappedScalar3D),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%anisoMortarBuff_gpu)) then
      ! AnisoMortarFlip_3D stages one face through shared memory (MORTAR3D_MAXNP).
      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)*4* &
                 mesh%nAnisoMortars*this%nvar*prec
      call gpuCheck(hipMalloc(this%anisoMortarBuff_gpu,buffSize))
    endif

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

    call AnisoMortarGather_3D_gpu(this%anisoMortarBuff_gpu,this%boundary_gpu, &
                                  mesh%anisoMortarInfo_gpu,mesh%decomp%elemToRank_gpu, &
                                  mesh%decomp%rankId,offset,this%interp%N,this%nvar, &
                                  mesh%nAnisoMortars,this%nelem)

    if(mesh%decomp%mpiEnabled) then
      call mesh%decomp%FinalizeMPIExchangeAsync()
      call AnisoMortarFlip_3D_gpu(this%anisoMortarBuff_gpu,mesh%anisoMortarInfo_gpu, &
                                  mesh%decomp%elemToRank_gpu,mesh%decomp%rankId, &
                                  this%interp%N,this%nvar,mesh%nAnisoMortars)
    endif

    call AnisoMortarScatter_3D_gpu(this%extBoundary_gpu,this%anisoMortarBuff_gpu, &
                                   this%interp%mortarR_gpu,this%interp%mortarP_gpu, &
                                   mesh%anisoMortarInfo_gpu,mesh%decomp%elemToRank_gpu, &
                                   mesh%decomp%rankId,offset,this%interp%N,this%nvar, &
                                   mesh%nAnisoMortars,this%nelem)

  endsubroutine AnisoMortarExchange_MappedScalar3D