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 hipFree proc~ensuredevicebuffer->interface~hipfree interface~hipmalloc hipMalloc proc~ensuredevicebuffer->interface~hipmalloc proc~gpucheck gpuCheck proc~ensuredevicebuffer->proc~gpucheck interface~hipmemset hipMemset proc~ensuredevicebuffer->interface~hipmemset

Called by

proc~~ensuredevicebuffer~~CalledByGraph proc~ensuredevicebuffer EnsureDeviceBuffer proc~ensuredevicebuffers_vector2d EnsureDeviceBuffers_Vector2D proc~ensuredevicebuffers_vector2d->proc~ensuredevicebuffer proc~ensuredevicebuffers_scalar3d EnsureDeviceBuffers_Scalar3D proc~ensuredevicebuffers_scalar3d->proc~ensuredevicebuffer proc~ensuredevicebuffers_vector3d EnsureDeviceBuffers_Vector3D proc~ensuredevicebuffers_vector3d->proc~ensuredevicebuffer proc~ensuredevicebuffers_mappedscalar3d EnsureDeviceBuffers_MappedScalar3D proc~ensuredevicebuffers_mappedscalar3d->proc~ensuredevicebuffer proc~ensuredevicebuffers_mappedscalar3d->proc~ensuredevicebuffers_scalar3d proc~ensuredevicebuffers_tensor2d EnsureDeviceBuffers_Tensor2D proc~ensuredevicebuffers_tensor2d->proc~ensuredevicebuffer proc~ensuredevicebuffers_tensor3d EnsureDeviceBuffers_Tensor3D proc~ensuredevicebuffers_tensor3d->proc~ensuredevicebuffer proc~ensuredevicebuffers_scalar2d EnsureDeviceBuffers_Scalar2D proc~ensuredevicebuffers_scalar2d->proc~ensuredevicebuffer proc~ensuredevicebuffers_mappedscalar2d EnsureDeviceBuffers_MappedScalar2D proc~ensuredevicebuffers_mappedscalar2d->proc~ensuredevicebuffer proc~ensuredevicebuffers_mappedscalar2d->proc~ensuredevicebuffers_scalar2d proc~init_tensor2d Init_Tensor2D proc~init_tensor2d->proc~ensuredevicebuffers_tensor2d proc~init_tensor3d Init_Tensor3D proc~init_tensor3d->proc~ensuredevicebuffers_tensor3d proc~resize_tensor2d Resize_Tensor2D proc~resize_tensor2d->proc~ensuredevicebuffers_tensor2d proc~resize_scalar2d Resize_Scalar2D proc~resize_scalar2d->proc~ensuredevicebuffers_scalar2d proc~resize_mappedscalar2d Resize_MappedScalar2D proc~resize_mappedscalar2d->proc~ensuredevicebuffers_mappedscalar2d proc~init_mappedscalar2d Init_MappedScalar2D proc~init_mappedscalar2d->proc~ensuredevicebuffers_mappedscalar2d proc~resize_tensor3d Resize_Tensor3D proc~resize_tensor3d->proc~ensuredevicebuffers_tensor3d proc~init_vector2d Init_Vector2D proc~init_vector2d->proc~ensuredevicebuffers_vector2d proc~resize_vector2d Resize_Vector2D proc~resize_vector2d->proc~ensuredevicebuffers_vector2d proc~resize_scalar3d Resize_Scalar3D proc~resize_scalar3d->proc~ensuredevicebuffers_scalar3d proc~resize_mappedvector2d Resize_MappedVector2D proc~resize_mappedvector2d->proc~ensuredevicebuffers_vector2d proc~init_scalar3d Init_Scalar3D proc~init_scalar3d->proc~ensuredevicebuffers_scalar3d proc~init_vector3d Init_Vector3D proc~init_vector3d->proc~ensuredevicebuffers_vector3d proc~resize_mappedvector3d Resize_MappedVector3D proc~resize_mappedvector3d->proc~ensuredevicebuffers_vector3d proc~resize_vector3d Resize_Vector3D proc~resize_vector3d->proc~ensuredevicebuffers_vector3d proc~init_mappedscalar3d Init_MappedScalar3D proc~init_mappedscalar3d->proc~ensuredevicebuffers_mappedscalar3d proc~resize_mappedscalar3d Resize_MappedScalar3D proc~resize_mappedscalar3d->proc~ensuredevicebuffers_mappedscalar3d proc~init_scalar2d Init_Scalar2D proc~init_scalar2d->proc~ensuredevicebuffers_scalar2d

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