SELF_SetInterface_3D.f90 Source File

! //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// !
!
! Maintainers : support@fluidnumerics.com
! Official Repository : https://github.com/FluidNumerics/self/
!
! Copyright © 2024 Fluid Numerics LLC
!
! Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:
!
! 1. Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.
!
! 2. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in
!    the documentation and/or other materials provided with the distribution.
!
! 3. Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from
!    this software without specific prior written permission.
!
! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS “AS IS” AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
! LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
! HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
! LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
! THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF
! THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
!
! //////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////////// !

module SELF_SetInterface_3D
!! GPU backend of the element-set interface exchange. When every set with interface
!! faces supplies a device trace (SetTraceRef%trace_gpu), the face traces are gathered on
!! the device and copied to the host in one transfer; the host exchange then runs as on
!! the CPU, and the ghost traces are copied back to the device for
!! FillSetInterfaceBoundary_gpu. Without device traces the host path runs unchanged.

  use SELF_Constants
  use SELF_Mesh_3D
  use SELF_SetInterface_3D_t
  use SELF_GPU
  use iso_c_binding

  implicit none

  type,extends(SetInterface3D_t) :: SetInterface3D
    type(c_ptr) :: faceElem_gpu = c_null_ptr
    type(c_ptr) :: faceSide_gpu = c_null_ptr
    type(c_ptr) :: outbox_gpu = c_null_ptr
    type(c_ptr) :: ghost_gpu = c_null_ptr

  contains

    procedure :: BuildMap => BuildMap_SetInterface3D
    procedure :: Exchange => Exchange_SetInterface3D
    procedure :: Free => Free_SetInterface3D

  endtype SetInterface3D

  interface
    subroutine GatherSetInterfaceOutbox_gpu(outbox,trace,elems,sides,f0,nFaces,nvar,nvarMax,N,nel) &
      bind(c,name="GatherSetInterfaceOutbox_gpu")
      use iso_c_binding
      type(c_ptr),value :: outbox,trace,elems,sides
      integer(c_int),value :: f0,nFaces,nvar,nvarMax,N,nel
    endsubroutine GatherSetInterfaceOutbox_gpu
  endinterface

  interface
    subroutine FillSetInterfaceBoundary_gpu(extBoundary,ghost,elems,sides,f0,nFaces,nvar,nvarMax,N,nel) &
      bind(c,name="FillSetInterfaceBoundary_gpu")
      use iso_c_binding
      type(c_ptr),value :: extBoundary,ghost,elems,sides
      integer(c_int),value :: f0,nFaces,nvar,nvarMax,N,nel
    endsubroutine FillSetInterfaceBoundary_gpu
  endinterface

contains

  subroutine BuildMap_SetInterface3D(this,globalMesh,setOfElem,setGlobalId, &
                                     setOffsetThisRank,nvarOfSet,N)
    !! Build the map on the host (BuildMap_SetInterface3D_t), then copy the face element
    !! and side lists to the device and allocate device outbox and ghost buffers.
    implicit none
    class(SetInterface3D),intent(inout) :: this
    type(Mesh3D),intent(in) :: globalMesh
    integer,intent(in) :: setOfElem(1:)
    integer,intent(in) :: setGlobalId(1:)
    integer,intent(in) :: setOffsetThisRank(1:)
    integer,intent(in) :: nvarOfSet(1:)
    integer,intent(in) :: N
    ! Local
    integer(c_int),allocatable,target :: buf(:)

    call BuildMap_SetInterface3D_t(this,globalMesh,setOfElem,setGlobalId, &
                                   setOffsetThisRank,nvarOfSet,N)
    if(this%nFaces == 0) return

    allocate(buf(1:this%nFaces))
    buf = this%faceElem
    call gpuCheck(hipMalloc(this%faceElem_gpu,sizeof(buf)))
    call gpuCheck(hipMemcpy(this%faceElem_gpu,c_loc(buf),sizeof(buf),hipMemcpyHostToDevice))
    buf = this%faceSide
    call gpuCheck(hipMalloc(this%faceSide_gpu,sizeof(buf)))
    call gpuCheck(hipMemcpy(this%faceSide_gpu,c_loc(buf),sizeof(buf),hipMemcpyHostToDevice))
    deallocate(buf)

    call gpuCheck(hipMalloc(this%outbox_gpu,sizeof(this%outbox)))
    call gpuCheck(hipMalloc(this%ghost_gpu,sizeof(this%ghost)))

  endsubroutine BuildMap_SetInterface3D

  subroutine Exchange_SetInterface3D(this,traces)
    !! As Exchange_SetInterface3D_t. If every set with local interface faces has a device
    !! trace, the outbox is gathered on the device; either way ghost is copied to the
    !! device afterwards, so FillSetInterfaceBoundary_gpu reads the current exchange.
    implicit none
    class(SetInterface3D),intent(inout) :: this
    type(SetTraceRef),intent(in) :: traces(1:)
    ! Local
    integer :: k,nk
    logical :: onDevice

    if(this%nFaces == 0) then
      call this%ExchangeOutbox()
      return
    endif

    onDevice = .true.
    do k = 1,this%nSets
      nk = this%setFaceStart(k+1)-this%setFaceStart(k)
      if(nk > 0 .and. .not. c_associated(traces(k)%trace_gpu)) onDevice = .false.
    enddo

    if(onDevice) then
      do k = 1,this%nSets
        nk = this%setFaceStart(k+1)-this%setFaceStart(k)
        if(nk == 0) cycle
        call GatherSetInterfaceOutbox_gpu(this%outbox_gpu,traces(k)%trace_gpu, &
                                          this%faceElem_gpu,this%faceSide_gpu, &
                                          this%setFaceStart(k)-1,nk,this%nvarOfSet(k), &
                                          this%nvarMax,this%N,size(traces(k)%trace,4))
      enddo
      call gpuCheck(hipMemcpy(c_loc(this%outbox),this%outbox_gpu,sizeof(this%outbox), &
                              hipMemcpyDeviceToHost))
    else
      call this%GatherOutbox(traces)
    endif

    call this%ExchangeOutbox()

    call gpuCheck(hipMemcpy(this%ghost_gpu,c_loc(this%ghost),sizeof(this%ghost), &
                            hipMemcpyHostToDevice))

  endsubroutine Exchange_SetInterface3D

  subroutine Free_SetInterface3D(this)
    !! Free the device buffers, then the host map. Collective.
    implicit none
    class(SetInterface3D),intent(inout) :: this

    if(c_associated(this%faceElem_gpu)) call gpuCheck(hipFree(this%faceElem_gpu))
    if(c_associated(this%faceSide_gpu)) call gpuCheck(hipFree(this%faceSide_gpu))
    if(c_associated(this%outbox_gpu)) call gpuCheck(hipFree(this%outbox_gpu))
    if(c_associated(this%ghost_gpu)) call gpuCheck(hipFree(this%ghost_gpu))
    this%faceElem_gpu = c_null_ptr
    this%faceSide_gpu = c_null_ptr
    this%outbox_gpu = c_null_ptr
    this%ghost_gpu = c_null_ptr

    call Free_SetInterface3D_t(this)

  endsubroutine Free_SetInterface3D

endmodule SELF_SetInterface_3D