UpdateDevice_Points Subroutine

  1. SELF_Points.f90
  2. SELF_Points
  3. UpdateDevice_Points

public subroutine UpdateDevice_Points(this)

Copy elements, coordinates, and the per-point Lagrange basis cache from host to device. Lazily (re)allocates the cache buffers if their degree changed since the previous call.

Arguments

TypeIntentOptionalAttributesName
class(Points), intent(inout) :: this

Calls

proc~~updatedevice_points~~CallsGraph proc~updatedevice_points UpdateDevice_Points interface~hipmemcpy hipMemcpy proc~updatedevice_points->interface~hipmemcpy interface~hipmalloc hipMalloc proc~updatedevice_points->interface~hipmalloc interface~hipfree hipFree proc~updatedevice_points->interface~hipfree proc~gpucheck gpuCheck proc~updatedevice_points->proc~gpucheck

Contents

Source Code

UpdateDevice_Points

Source Code

  subroutine UpdateDevice_Points(this)
    !! Copy elements, coordinates, and the per-point Lagrange basis cache
    !! from host to device. Lazily (re)allocates the cache buffers if their
    !! degree changed since the previous call.
    implicit none
    class(Points),intent(inout) :: this
    ! Local
    integer(c_size_t) :: nBytes

    if(this%nPoints == 0) return

    nBytes = int(this%nPoints,c_size_t)*c_sizeof(0_c_int)
    call gpuCheck(hipMemcpy(this%elements_gpu,c_loc(this%elements),nBytes, &
                            hipMemcpyHostToDevice))

    nBytes = int(this%nPoints*this%nDim,c_size_t)*int(prec,c_size_t)
    call gpuCheck(hipMemcpy(this%coordinates_gpu,c_loc(this%coordinates),nBytes, &
                            hipMemcpyHostToDevice))

    if(this%nCached <= 0) return
    if(.not. associated(this%lS_cache) .or. .not. associated(this%lT_cache)) return

    ! (Re)allocate per-point basis caches on the device if needed.
    if(this%nCachedAlloc /= this%nCached) then
      if(c_associated(this%lS_cache_gpu)) call gpuCheck(hipFree(this%lS_cache_gpu))
      if(c_associated(this%lT_cache_gpu)) call gpuCheck(hipFree(this%lT_cache_gpu))
      if(c_associated(this%lU_cache_gpu)) call gpuCheck(hipFree(this%lU_cache_gpu))
      this%lS_cache_gpu = c_null_ptr
      this%lT_cache_gpu = c_null_ptr
      this%lU_cache_gpu = c_null_ptr
      nBytes = int((this%nCached+1)*this%nPoints,c_size_t)*int(prec,c_size_t)
      call gpuCheck(hipMalloc(this%lS_cache_gpu,nBytes))
      call gpuCheck(hipMalloc(this%lT_cache_gpu,nBytes))
      if(this%nDim == 3) call gpuCheck(hipMalloc(this%lU_cache_gpu,nBytes))
      this%nCachedAlloc = this%nCached
    endif

    nBytes = int((this%nCached+1)*this%nPoints,c_size_t)*int(prec,c_size_t)
    call gpuCheck(hipMemcpy(this%lS_cache_gpu,c_loc(this%lS_cache),nBytes, &
                            hipMemcpyHostToDevice))
    call gpuCheck(hipMemcpy(this%lT_cache_gpu,c_loc(this%lT_cache),nBytes, &
                            hipMemcpyHostToDevice))
    if(this%nDim == 3 .and. associated(this%lU_cache)) then
      call gpuCheck(hipMemcpy(this%lU_cache_gpu,c_loc(this%lU_cache),nBytes, &
                              hipMemcpyHostToDevice))
    endif

  endsubroutine UpdateDevice_Points