EnsureDeviceBuffer Subroutine

  1. SELF_GPU.f90
  2. SELF_GPU
  3. EnsureDeviceBuffer

public subroutine EnsureDeviceBuffer(ptr, allocBytes, neededBytes)

Uses

  • proc~~ensuredevicebuffer~~UsesGraph proc~ensuredevicebuffer EnsureDeviceBuffer iso_c_binding iso_c_binding proc~ensuredevicebuffer->iso_c_binding

High-water-mark device allocation (AMR Stage 6b). Grows ptr to hold neededBytes, reusing the existing allocation when it already does, and records the capacity in allocBytes. Contents are not preserved across a growth - every caller rewrites the buffer before reading it.

A device pointer carries no shape, so byte capacity is the only thing that has to be tracked; this is what lets the adaptive loop stop calling hipMalloc/hipFree once the element count settles.

Arguments

TypeIntentOptionalAttributesName
type(c_ptr), intent(inout) :: ptr
integer(kind=c_size_t), intent(inout) :: allocBytes
integer(kind=c_size_t), intent(in) :: neededBytes

Calls

proc~~ensuredevicebuffer~~CallsGraph proc~ensuredevicebuffer EnsureDeviceBuffer interface~hipfree~2 hipFree proc~ensuredevicebuffer->interface~hipfree~2 proc~gpucheck~2 gpuCheck proc~ensuredevicebuffer->proc~gpucheck~2 interface~hipmalloc~2 hipMalloc proc~ensuredevicebuffer->interface~hipmalloc~2 interface~hipmemset hipMemset proc~ensuredevicebuffer->interface~hipmemset

Contents

Source Code

EnsureDeviceBuffer

Source Code

                              subroutine EnsureDeviceBuffer(ptr,allocBytes,neededBytes)
                            !! High-water-mark device allocation (AMR Stage 6b). Grows ptr to
                            !! hold neededBytes, reusing the existing allocation when it already
                            !! does, and records the capacity in allocBytes. Contents are not
                            !! preserved across a growth - every caller rewrites the buffer
                            !! before reading it.
                            !!
                            !! A device pointer carries no shape, so byte capacity is the only
                            !! thing that has to be tracked; this is what lets the adaptive loop
                            !! stop calling hipMalloc/hipFree once the element count settles.
                                use iso_c_binding
                                implicit none
                                type(c_ptr),intent(inout) :: ptr
                                integer(c_size_t),intent(inout) :: allocBytes
                                integer(c_size_t),intent(in) :: neededBytes

                                if(.not.(c_associated(ptr) .and. allocBytes >= neededBytes)) then
                                  if(c_associated(ptr)) then
                                    call gpuCheck(hipFree(ptr))
                                    ptr = c_null_ptr
                                  endif
                                  call gpuCheck(hipMalloc(ptr,neededBytes))
                                  allocBytes = neededBytes
                                endif

                                ! Leave the buffer DEFINED (zeroed), which Init used to guarantee by
                                ! uploading its zeroed host arrays. Resize deliberately performs no
                                ! host-to-device copy, so without this the device buffer would keep
                                ! whatever was previously in that memory - and freshly hipMalloc'd
                                ! memory is uninitialized. That matters because the low-storage RK
                                ! update reads its accumulator before writing it
                                ! (UpdateGRK_Model: grk = rk_a*grk + dSdt, with rk_a = 0 on the first
                                ! stage). Multiplying by zero annihilates any finite leftover, but
                                ! 0*NaN and 0*Inf are NaN, so a stale bit pattern that happens to be
                                ! NaN or Inf silently poisons the solution. The observed symptom was
                                ! an intermittent, allocation-history-dependent NaN several adaptation
                                ! epochs into a GPU run, with the CPU build unaffected because there
                                ! the host array IS the storage and Resize zeroes it.
                                !
                                ! A device-side fill costs HBM bandwidth rather than a PCIe transfer,
                                ! so this keeps the point of skipping UpdateDevice.
                                call gpuCheck(hipMemset(ptr,0,neededBytes))

                              endsubroutine EnsureDeviceBuffer