HIP runtime examples#
The HIP runtime is the
API used to manage devices, memory, streams, events, graphs and kernel
launches. hipFORT exposes it through the hipfort module, with enumerators in
hipfort_enums, derived types in hipfort_types, and the status-checking
helpers in hipfort_check. A few routines live in their own modules:
hipGetDeviceProperties in hipfort_auxiliary, hipMemcpy2DAsync in
hipfort_hipmemcpy and hipHostRegister in hipfort_hiphostregister.
Every program on this page is complete and self-contained, and is built and run
as part of the hipFORT test suite. The sources live in test/f2003/hip, and
the programs that benefit from Fortran array pointers have a Fortran 2008 twin
in test/f2008/hip.
Conventions#
Device pointers. The Fortran 2003 programs hold device memory in a
type(c_ptr)and pass byte counts, as inhipMalloc(dx, nbytes)andhipMemcpy(dx, c_loc(hx(1)), nbytes, hipMemcpyHostToDevice). The Fortran 2008 interfaces instead accept a Fortran array pointer and an element count, as inhipMalloc(dx, n)orhipMalloc(dx, source=hx).Every call returns a status code. The programs wrap calls in
hipCheckfrom thehipfort_checkmodule, which aborts on failure. A call whose non-success return is the thing being tested, such ashipStreamQuery, keeps the status in a variable instead.Enumerators are integers. Declare status variables as
integer(kind(hipSuccess))so they match the kind the interfaces return.Host callbacks and kernel stubs are passed as
c_funlocof a procedure declaredbind(c).
Building and running#
The programs only need the hip hipFORT component:
find_package(hipfort REQUIRED COMPONENTS hip)
add_executable(my_app stream.f03)
target_link_libraries(my_app PRIVATE hipfort::hip)
See Using hipFORT in your application for the full set of build options.
Device management#
The device queries report how many GPUs are visible, select one for the calling thread, and read back its limits and free memory.
!!!!!!!!!!!!!!
! HIP runtime device-management queries
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Exercises hipGetDeviceCount, hipSetDevice/hipGetDevice, hipDeviceGetAttribute,
! hipDeviceGetLimit, hipDeviceTotalMem and hipMemGetInfo, checking basic
! invariants that hold on any working device.
!!!!!!!!!!!!!!
!
program device_management
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
implicit none
integer(c_int) :: ndev, dev, warp, nmp
integer(c_size_t) :: totmem, freemem, total2, stacklimit
write(*,"(a)",advance="no") "-- Running test 'hip device_management' (Fortran 2003 interfaces) - "
call hipCheck(hipGetDeviceCount(ndev))
if (ndev < 1) then
write(*,*) "FAILED! device count = ", ndev
call exit(1)
end if
call hipCheck(hipSetDevice(0))
call hipCheck(hipGetDevice(dev))
if (dev /= 0) then
write(*,*) "FAILED! current device = ", dev, " (expected 0)"
call exit(1)
end if
call hipCheck(hipDeviceGetAttribute(warp, hipDeviceAttributeWarpSize, 0))
if (warp <= 0) then
write(*,*) "FAILED! warp size = ", warp
call exit(1)
end if
call hipCheck(hipDeviceGetAttribute(nmp, hipDeviceAttributeMultiprocessorCount, 0))
if (nmp <= 0) then
write(*,*) "FAILED! multiprocessor count = ", nmp
call exit(1)
end if
call hipCheck(hipDeviceGetLimit(stacklimit, hipLimitStackSize))
if (stacklimit <= 0) then
write(*,*) "FAILED! stack size limit = ", stacklimit
call exit(1)
end if
call hipCheck(hipDeviceTotalMem(totmem, 0))
if (totmem <= 0) then
write(*,*) "FAILED! total device memory = ", totmem
call exit(1)
end if
call hipCheck(hipMemGetInfo(freemem, total2))
if (total2 <= 0 .or. freemem > total2) then
write(*,*) "FAILED! memGetInfo free = ", freemem, " total = ", total2
call exit(1)
end if
write(*,*) "PASSED!"
end program device_management
hipGetDeviceProperties returns the same information in one
hipDeviceProp_t structure; test/f2003/hip/device_properties.f03 reads it
and cross-checks a few fields against hipDeviceGetAttribute.
Memory copies and fills#
Beyond hipMemcpy and hipMemset, the runtime offers pitched two
dimensional operations, typed fills, and asynchronous forms that take a stream.
!!!!!!!!!!!!!!
! HIP runtime 2-D and typed memory operations
! see: https://rocm.docs.amd.com/projects/HIP/en/latest/
!
! Exercises hipMemset2D (pitched 2-D byte fill), hipMemsetD32 (32-bit word fill),
! hipMemcpy2DAsync (async D2D pitched copy on a stream), and hipMemGetInfo.
!!!!!!!!!!!!!!
!
program memory_ops
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
use hipfort_hipmemcpy
implicit none
! matrix: 8 int32 words per row, 4 rows
integer, parameter :: NCOLS = 8, NROWS = 4, COL_BYTES = NCOLS * 4
! DEADBEEF as signed int32 (two's complement: -559038737)
integer(c_int), parameter :: DEADBEEF = int(z'DEADBEEF', c_int)
type(c_ptr) :: dptr_a = c_null_ptr ! pitched src (memset2D + async test)
type(c_ptr) :: dptr_b = c_null_ptr ! flat buffer (memsetD32 test)
type(c_ptr) :: dptr_c = c_null_ptr ! pitched dst (async test)
type(c_ptr) :: stream = c_null_ptr
integer(c_size_t) :: pitch_a, pitch_c
integer(c_size_t) :: free_mem, total_mem
integer(c_int8_t), target :: hbuf8(COL_BYTES * NROWS)
integer(c_int), target :: hbuf_i(NCOLS * NROWS)
integer :: i
write(*,"(a)",advance="no") "-- Running test 'hip memory_ops' (Fortran 2003 interfaces) - "
call hipCheck(hipSetDevice(0))
call hipCheck(hipMemGetInfo(free_mem, total_mem))
if (total_mem == 0_c_size_t) then
write(*,*) "FAILED! hipMemGetInfo total = 0"
call exit(1)
end if
if (free_mem > total_mem) then
write(*,*) "FAILED! hipMemGetInfo free > total:", free_mem, ">", total_mem
call exit(1)
end if
! Pitched buffer COL_BYTES wide by NROWS rows; the driver picks pitch_a.
call hipCheck(hipMallocPitch(dptr_a, pitch_a, int(COL_BYTES, c_size_t), int(NROWS, c_size_t)))
call hipCheck(hipMemset2D(dptr_a, pitch_a, int(z'42', c_int), &
int(COL_BYTES, c_size_t), int(NROWS, c_size_t)))
hbuf8 = 0_c_int8_t
call hipCheck(hipMemcpy2D(c_loc(hbuf8(1)), int(COL_BYTES, c_size_t), &
dptr_a, pitch_a, &
int(COL_BYTES, c_size_t), int(NROWS, c_size_t), &
hipMemcpyDeviceToHost))
do i = 1, COL_BYTES * NROWS
if (hbuf8(i) /= int(z'42', c_int8_t)) then
write(*,"(a,i0,a,i0,a)") "FAILED! hipMemset2D: hbuf8(", i, ") = ", hbuf8(i), &
" (expected 0x42 = 66)"
call exit(1)
end if
end do
call hipCheck(hipMalloc(dptr_b, int(NCOLS * NROWS * 4, c_size_t)))
call hipCheck(hipMemsetD32(dptr_b, DEADBEEF, int(NCOLS * NROWS, c_size_t)))
hbuf_i = 0
call hipCheck(hipMemcpy(c_loc(hbuf_i(1)), dptr_b, &
int(NCOLS * NROWS * 4, c_size_t), hipMemcpyDeviceToHost))
do i = 1, NCOLS * NROWS
if (hbuf_i(i) /= DEADBEEF) then
write(*,"(a,i0,a,z8.8,a,z8.8,a)") "FAILED! hipMemsetD32: hbuf_i(", i, ") = 0x", &
hbuf_i(i), " (expected 0x", DEADBEEF, ")"
call exit(1)
end if
end do
call hipCheck(hipFree(dptr_b))
dptr_b = c_null_ptr
call hipCheck(hipMallocPitch(dptr_c, pitch_c, int(COL_BYTES, c_size_t), int(NROWS, c_size_t)))
call hipCheck(hipStreamCreate(stream))
call hipCheck(hipMemcpy2DAsync(dptr_c, pitch_c, dptr_a, pitch_a, &
int(COL_BYTES, c_size_t), int(NROWS, c_size_t), &
hipMemcpyDeviceToDevice, stream))
call hipCheck(hipStreamSynchronize(stream))
hbuf8 = 0_c_int8_t
call hipCheck(hipMemcpy2D(c_loc(hbuf8(1)), int(COL_BYTES, c_size_t), &
dptr_c, pitch_c, &
int(COL_BYTES, c_size_t), int(NROWS, c_size_t), &
hipMemcpyDeviceToHost))
do i = 1, COL_BYTES * NROWS
if (hbuf8(i) /= int(z'42', c_int8_t)) then
write(*,"(a,i0,a,i0,a)") "FAILED! hipMemcpy2DAsync: hbuf8(", i, ") = ", hbuf8(i), &
" (expected 0x42 = 66)"
call exit(1)
end if
end do
call hipCheck(hipFree(dptr_a))
call hipCheck(hipFree(dptr_c))
call hipCheck(hipStreamDestroy(stream))
write(*,*) "PASSED!"
end program memory_ops
test/f2003/hip/memcpy2d.f03 covers hipMemcpy2D on a column-major matrix,
and test/f2003/hip/memcpy_async.f03 covers hipMemcpyAsync and
hipMemcpyWithStream.
Pinned and managed memory#
hipHostMalloc allocates page-locked host memory, which the GPU can copy
to and from without a staging buffer.
!!!!!!!!!!!!!!
! HIP runtime hipHostMalloc (pinned host memory, Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Allocates pinned host memory, uses it as the source of a host->device copy,
! copies back into a plain host array and verifies the round trip.
!!!!!!!!!!!!!!
!
program host_malloc
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
implicit none
integer(c_int), parameter :: n = 256
type(c_ptr) :: hpinned = c_null_ptr, dptr = c_null_ptr
real(c_double), pointer :: hp(:)
real(c_double), target :: hcheck(n)
integer(c_size_t) :: nbytes
integer :: i
write(*,"(a)",advance="no") "-- Running test 'hip host_malloc' (Fortran 2003 interfaces) - "
nbytes = int(n, c_size_t) * 8
call hipCheck(hipSetDevice(0))
call hipCheck(hipHostMalloc(hpinned, nbytes, hipHostMallocDefault))
call c_f_pointer(hpinned, hp, [n])
do i = 1, n
hp(i) = real(i, c_double)
end do
call hipCheck(hipMalloc(dptr, nbytes))
call hipCheck(hipMemcpy(dptr, hpinned, nbytes, hipMemcpyHostToDevice))
hcheck = 0.0d0
call hipCheck(hipMemcpy(c_loc(hcheck(1)), dptr, nbytes, hipMemcpyDeviceToHost))
do i = 1, n
if (hcheck(i) /= hp(i)) then
write(*,*) "FAILED! hcheck(", i, ") = ", hcheck(i), " expected ", hp(i)
call exit(1)
end if
end do
call hipCheck(hipFree(dptr))
call hipCheck(hipHostFree(hpinned))
write(*,*) "PASSED!"
end program host_malloc
An existing host array can be page-locked in place with hipHostRegister
(test/f2003/hip/host_register.f03). hipMallocManaged
(test/f2003/hip/malloc_managed.f03) allocates memory that both the host and
the device address directly, and test/f2003/hip/mem_advise.f03 adds
migration hints on top of it. test/f2003/hip/pointer_attributes.f03 queries
which of the three kinds a pointer belongs to.
Virtual memory management#
The virtual memory API separates the address range from the physical memory backing it: reserve a range, create a physical allocation, map one onto the other, then grant the device access. This allows an allocation to grow without changing the pointer the application already holds.
!!!!!!!!!!!!!!
! HIP runtime virtual memory management (Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Reserves a virtual address range, backs it with a physical allocation, makes
! it accessible to the device, uses it as ordinary device memory, and unwinds
! the mapping again.
!!!!!!!!!!!!!!
!
program virtual_memory
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
use hipfort_types
implicit none
integer(c_int), parameter :: n = 1024
type(hipMemAllocationProp) :: prop
type(hipMemAccessDesc) :: desc
type(c_ptr) :: vptr = c_null_ptr
type(c_ptr) :: handle = c_null_ptr
integer(c_size_t), target :: granularity
integer(c_size_t) :: nbytes, padded
real(c_double), target :: hx(n), hy(n)
integer :: i
write(*,"(a)",advance="no") "-- Running test 'hip virtual_memory' (Fortran 2003 interfaces) - "
call hipCheck(hipSetDevice(0))
prop%type = hipMemAllocationTypePinned
prop%requestedHandleType = 0
prop%location%type = hipMemLocationTypeDevice
prop%location%id = 0
prop%win32HandleMetaData = c_null_ptr
prop%allocFlags = 0
granularity = 0
call hipCheck(hipMemGetAllocationGranularity(c_loc(granularity), prop, &
hipMemAllocationGranularityMinimum))
if (granularity <= 0) then
write(*,*) "FAILED! allocation granularity =", granularity
call exit(1)
end if
! A physical allocation has to be a whole number of granularity units.
nbytes = int(n, c_size_t) * 8
padded = ((nbytes + granularity - 1) / granularity) * granularity
call hipCheck(hipMemAddressReserve(vptr, padded, 0_c_size_t, c_null_ptr, 0_c_int64_t))
call hipCheck(hipMemCreate(handle, padded, prop, 0_c_int64_t))
call hipCheck(hipMemMap(vptr, padded, 0_c_size_t, handle, 0_c_int64_t))
! Mapped memory starts out inaccessible; the device needs read/write access.
desc%location%type = hipMemLocationTypeDevice
desc%location%id = 0
desc%flags = hipMemAccessFlagsProtReadWrite
call hipCheck(hipMemSetAccess(vptr, padded, desc, 1_c_size_t))
do i = 1, n
hx(i) = real(i, c_double)
hy(i) = 0.0_c_double
end do
call hipCheck(hipMemcpy(vptr, c_loc(hx(1)), nbytes, hipMemcpyHostToDevice))
call hipCheck(hipMemcpy(c_loc(hy(1)), vptr, nbytes, hipMemcpyDeviceToHost))
do i = 1, n
if (hy(i) /= hx(i)) then
write(*,*) "FAILED! hy(", i, ") = ", hy(i), " expected ", hx(i)
call exit(1)
end if
end do
! The device must be able to write the range as well.
call hipCheck(hipMemset(vptr, 0, nbytes))
call hipCheck(hipDeviceSynchronize())
call hipCheck(hipMemcpy(c_loc(hy(1)), vptr, nbytes, hipMemcpyDeviceToHost))
do i = 1, n
if (hy(i) /= 0.0_c_double) then
write(*,*) "FAILED! hy(", i, ") = ", hy(i), " after memset (expected 0)"
call exit(1)
end if
end do
call hipCheck(hipMemUnmap(vptr, padded))
call hipCheck(hipMemRelease(handle))
call hipCheck(hipMemAddressFree(vptr, padded))
write(*,*) "PASSED!"
end program virtual_memory
Streams#
Work queued on the same stream runs in order, and work on different streams may overlap. Streams can be created with flags and with a priority from the range the device reports.
!!!!!!!!!!!!!!
! HIP runtime stream flag and priority API
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Exercises hipDeviceGetStreamPriorityRange, hipStreamCreateWithFlags,
! hipStreamCreateWithPriority, hipStreamGetFlags, hipStreamGetPriority,
! and hipStreamQuery.
!!!!!!!!!!!!!!
!
program stream_flags
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
implicit none
integer(c_int), parameter :: n = 1048576 ! 4 MB of int32
integer(c_int), target :: hsrc(n), hdst(n)
type(c_ptr) :: dptr = c_null_ptr
type(c_ptr) :: strm_nb = c_null_ptr, strm_pri = c_null_ptr
integer(c_int) :: least, greatest, got_flags, got_pri
integer(c_size_t) :: nbytes
integer :: i
integer(kind(hipSuccess)) :: qret
write(*,"(a)",advance="no") &
"-- Running test 'hip stream_flags' (Fortran 2003 interfaces) - "
nbytes = int(n, c_size_t) * 4 ! 4 bytes per int32
! HIP convention: lower number means higher priority, so greatest <= least.
call hipCheck(hipSetDevice(0))
call hipCheck(hipDeviceGetStreamPriorityRange(least, greatest))
write(*,"(a,i0,a,i0,a)",advance="no") &
"[priority range: least=", least, " greatest=", greatest, "] "
if (least < greatest) then
write(*,*) "FAILED! leastPriority (", least, &
") should be >= greatestPriority (", greatest, ")"
call exit(1)
end if
! Non-blocking stream: the flags must read back as they were set.
call hipCheck(hipStreamCreateWithFlags(strm_nb, hipStreamNonBlocking))
if (.not. c_associated(strm_nb)) then
write(*,*) "FAILED! strm_nb is null after hipStreamCreateWithFlags"
call exit(1)
end if
got_flags = -1
call hipCheck(hipStreamGetFlags(strm_nb, got_flags))
if (got_flags /= hipStreamNonBlocking) then
write(*,*) "FAILED! hipStreamGetFlags returned ", got_flags, &
" (expected hipStreamNonBlocking =", hipStreamNonBlocking, ")"
call exit(1)
end if
! Ask for the highest-priority value so the round trip is unambiguous even
! when the range spans a single value.
call hipCheck(hipStreamCreateWithPriority(strm_pri, hipStreamDefault, greatest))
if (.not. c_associated(strm_pri)) then
write(*,*) "FAILED! strm_pri is null after hipStreamCreateWithPriority"
call exit(1)
end if
got_pri = -999
call hipCheck(hipStreamGetPriority(strm_pri, got_pri))
if (got_pri /= greatest) then
write(*,*) "FAILED! hipStreamGetPriority returned ", got_pri, &
" (expected ", greatest, ")"
call exit(1)
end if
! Real work on the priority stream: after sync the stream must report idle
! and the data must be what the copy and memset left behind.
do i = 1, n
hsrc(i) = i
end do
hdst = 0
call hipCheck(hipMalloc(dptr, nbytes))
call hipCheck(hipMemcpyAsync(dptr, c_loc(hsrc(1)), nbytes, hipMemcpyHostToDevice, strm_pri))
! overwrite first half with a known pattern so we can detect stale data
call hipCheck(hipMemsetAsync(dptr, 0, nbytes / 2, strm_pri))
call hipCheck(hipStreamSynchronize(strm_pri))
qret = hipStreamQuery(strm_pri)
if (qret /= hipSuccess) then
write(*,*) "FAILED! hipStreamQuery returned ", qret, " (expected hipSuccess=0)"
call exit(1)
end if
call hipCheck(hipMemcpy(c_loc(hdst(1)), dptr, nbytes, hipMemcpyDeviceToHost))
! first half was memset to 0
do i = 1, n / 2
if (hdst(i) /= 0) then
write(*,*) "FAILED! hdst(", i, ") = ", hdst(i), " (expected 0 from memset)"
call exit(1)
end if
end do
! second half carries the original copy
do i = n / 2 + 1, n
if (hdst(i) /= i) then
write(*,*) "FAILED! hdst(", i, ") = ", hdst(i), " (expected ", i, ")"
call exit(1)
end if
end do
call hipCheck(hipFree(dptr))
call hipCheck(hipStreamDestroy(strm_nb))
call hipCheck(hipStreamDestroy(strm_pri))
write(*,*) "PASSED!"
end program stream_flags
test/f2003/hip/stream.f03 shows the basic create, synchronize and destroy
sequence.
Host functions on a stream#
hipStreamAddCallback and hipLaunchHostFunc run a host procedure once the
work queued before it on the stream has completed.
!!!!!!!!!!!!!!
! HIP runtime host-side stream callbacks
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Exercises hipStreamAddCallback and hipLaunchHostFunc. Queues an async memset
! on a stream, then registers one callback of each kind via c_funloc. After
! synchronize: asserts both callbacks fired (counter == 2) and device data from
! the preceding memset is intact (ordering: callbacks trail enqueued device work).
!!!!!!!!!!!!!!
!
program stream_callback
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
implicit none
! Explicit interfaces for the external bind(c) procedures so c_funloc can see them.
interface
subroutine stream_cb(stream, status, userData) bind(c)
use iso_c_binding
use hipfort_enums
implicit none
type(c_ptr), value :: stream
integer(kind(hipSuccess)), value :: status
type(c_ptr), value :: userData
end subroutine
subroutine host_fn(userData) bind(c)
use iso_c_binding
implicit none
type(c_ptr), value :: userData
end subroutine
end interface
integer(c_int), parameter :: n = 256
integer(c_int8_t), target :: hbuf(n)
integer(c_int), target :: counter = 0 ! incremented by each callback
type(c_ptr) :: stream = c_null_ptr, dptr = c_null_ptr
integer(c_size_t) :: nbytes
write(*,"(a)",advance="no") "-- Running test 'hip stream_callback' (Fortran 2003 interfaces) - "
nbytes = int(n, c_size_t) ! one byte per element
call hipCheck(hipSetDevice(0))
call hipCheck(hipMalloc(dptr, nbytes))
call hipCheck(hipStreamCreate(stream))
! Queue device work; the callbacks below must not fire until this completes.
call hipCheck(hipMemsetAsync(dptr, 7, nbytes, stream))
call hipCheck(hipStreamAddCallback(stream, c_funloc(stream_cb), c_loc(counter), 0_c_int))
call hipCheck(hipLaunchHostFunc(stream, c_funloc(host_fn), c_loc(counter)))
call hipCheck(hipStreamSynchronize(stream))
if (counter /= 2) then
write(*,*) "FAILED! callback counter =", counter, "(expected 2)"
call exit(1)
end if
! Copy device data back; proves memset completed before either callback ran.
hbuf = 0_c_int8_t
call hipCheck(hipMemcpy(c_loc(hbuf(1)), dptr, nbytes, hipMemcpyDeviceToHost))
if (any(hbuf /= 7_c_int8_t)) then
write(*,*) "FAILED! device buffer not filled with 7 after memset"
call exit(1)
end if
call hipCheck(hipFree(dptr))
call hipCheck(hipStreamDestroy(stream))
write(*,*) "PASSED!"
end program stream_callback
! hipStreamCallback_t: void (*)(hipStream_t stream, hipError_t status, void* userData)
subroutine stream_cb(stream, status, userData) bind(c)
use iso_c_binding
use hipfort_enums
implicit none
type(c_ptr), value :: stream
integer(kind(hipSuccess)), value :: status
type(c_ptr), value :: userData
integer(c_int), pointer :: counter
call c_f_pointer(userData, counter)
counter = counter + 1
end subroutine stream_cb
! hipHostFn_t: void (*)(void* userData)
subroutine host_fn(userData) bind(c)
use iso_c_binding
implicit none
type(c_ptr), value :: userData
integer(c_int), pointer :: counter
call c_f_pointer(userData, counter)
counter = counter + 1
end subroutine host_fn
Events#
Events mark a point in a stream. They time device work and make one stream wait for another.
!!!!!!!!!!!!!!
! HIP runtime event timing and cross-stream ordering (Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Exercises hipEventCreateWithFlags, hipEventQuery, hipEventElapsedTime and
! hipStreamWaitEvent. The ordering check poisons a buffer and keeps the
! producing stream busy, so a missing wait reads poison rather than the data.
!!!!!!!!!!!!!!
!
program event_timing
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
implicit none
! dptr: 64 MB for timing + stream_a load dptr2: small buffer for ordering test
integer(c_size_t), parameter :: nbig = 16_c_size_t * 1024 * 1024
integer(c_size_t), parameter :: nsmall = 1024_c_size_t
integer(c_size_t), parameter :: bytes_big = nbig * 4_c_size_t
integer(c_size_t), parameter :: bytes_small = nsmall * 4_c_size_t
type(c_ptr) :: estart = c_null_ptr, estop = c_null_ptr, emark = c_null_ptr
type(c_ptr) :: stream_a = c_null_ptr, stream_b = c_null_ptr
type(c_ptr) :: dptr = c_null_ptr, dptr2 = c_null_ptr
real(c_float), target :: hsrc(nsmall), hdst(nsmall)
real(c_float) :: ms
integer(kind(hipSuccess)) :: istat
integer :: i
write(*,"(a)",advance="no") &
"-- Running test 'hip event_timing' (Fortran 2003 interfaces) - "
call hipCheck(hipSetDevice(0))
call hipCheck(hipStreamCreate(stream_a))
call hipCheck(hipStreamCreate(stream_b))
call hipCheck(hipMalloc(dptr, bytes_big))
call hipCheck(hipMalloc(dptr2, bytes_small))
! estart: default flags; estop: blocking-sync so hipEventSynchronize is efficient
call hipCheck(hipEventCreateWithFlags(estart, hipEventDefault))
call hipCheck(hipEventCreateWithFlags(estop, hipEventBlockingSync))
call hipCheck(hipEventCreateWithFlags(emark, hipEventDefault))
! ---- timing section: bracket four large async memsets on stream_a ----
call hipCheck(hipEventRecord(estart, stream_a))
call hipCheck(hipMemsetAsync(dptr, 0, bytes_big, stream_a))
call hipCheck(hipMemsetAsync(dptr, 1, bytes_big, stream_a))
call hipCheck(hipMemsetAsync(dptr, 2, bytes_big, stream_a))
call hipCheck(hipMemsetAsync(dptr, 3, bytes_big, stream_a))
call hipCheck(hipEventRecord(estop, stream_a))
! before sync: hipErrorNotReady or hipSuccess are both legitimate
istat = hipEventQuery(estop)
if (istat /= hipSuccess .and. istat /= hipErrorNotReady) then
write(*,*) "FAILED! hipEventQuery (pre-sync) unexpected istat = ", istat
call exit(1)
end if
call hipCheck(hipEventSynchronize(estop))
! after sync: must be hipSuccess
istat = hipEventQuery(estop)
if (istat /= hipSuccess) then
write(*,*) "FAILED! hipEventQuery (post-sync) = ", istat, " expected hipSuccess=0"
call exit(1)
end if
call hipCheck(hipEventElapsedTime(ms, estart, estop))
if (ms <= 0.0 .or. ms >= 10000.0) then
write(*,*) "FAILED! elapsed time = ", ms, " ms (expected 0 < ms < 10000)"
call exit(1)
end if
write(*,"(a,f9.4,a)",advance="no") "(elapsed ", ms, " ms) "
! Cross-stream ordering: poison the buffer, then keep stream_a busy so a
! missing wait lets stream_b read poison instead of the pattern.
call hipCheck(hipMemset(dptr2, 255, bytes_small))
call hipCheck(hipDeviceSynchronize())
do i = 1, int(nsmall)
hsrc(i) = real(i, c_float)
end do
hdst = 0.0_c_float
do i = 1, 12
call hipCheck(hipMemsetAsync(dptr, i, bytes_big, stream_a))
end do
call hipCheck(hipMemcpyAsync(dptr2, c_loc(hsrc(1)), bytes_small, &
hipMemcpyHostToDevice, stream_a))
call hipCheck(hipEventRecord(emark, stream_a))
! stream_b must wait for emark before reading dptr2
call hipCheck(hipStreamWaitEvent(stream_b, emark, 0))
call hipCheck(hipMemcpyWithStream(c_loc(hdst(1)), dptr2, bytes_small, &
hipMemcpyDeviceToHost, stream_b))
call hipCheck(hipStreamSynchronize(stream_b))
do i = 1, int(nsmall)
if (hdst(i) /= hsrc(i)) then
write(*,*) "FAILED! cross-stream hdst(", i, ") = ", hdst(i), &
" expected ", hsrc(i)
call exit(1)
end if
end do
call hipCheck(hipEventDestroy(estart))
call hipCheck(hipEventDestroy(estop))
call hipCheck(hipEventDestroy(emark))
call hipCheck(hipStreamDestroy(stream_a))
call hipCheck(hipStreamDestroy(stream_b))
call hipCheck(hipFree(dptr))
call hipCheck(hipFree(dptr2))
write(*,*) "PASSED!"
end program event_timing
test/f2003/hip/event.f03 shows the shorter form: record, synchronize and
read the elapsed time.
Graphs#
A graph records a sequence of operations and their dependencies once, so that repeated executions skip the per-call launch overhead. The simplest way to build one is to capture a stream.
!!!!!!!!!!!!!!
! HIP runtime graphs via stream capture (Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Captures a device memset into a HIP graph, instantiates it, launches the
! executable graph and verifies the buffer was written. Exercises
! hipStreamBeginCapture / hipStreamIsCapturing / hipStreamEndCapture /
! hipGraphInstantiate / hipGraphLaunch / hipGraphExecDestroy / hipGraphDestroy.
!!!!!!!!!!!!!!
!
program test_graph
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
implicit none
integer(c_int), parameter :: n = 256
integer(c_int8_t), target :: hbuf(n)
type(c_ptr) :: stream = c_null_ptr
type(c_ptr) :: graph = c_null_ptr
type(c_ptr) :: gexec = c_null_ptr
type(c_ptr) :: errnode = c_null_ptr
type(c_ptr) :: dptr = c_null_ptr
integer(c_size_t) :: nbytes
integer(kind(hipStreamCaptureStatusNone)), target :: capstat
integer :: i
write(*,"(a)",advance="no") "-- Running test 'hip graph' (Fortran 2003 interfaces) - "
nbytes = int(n, c_size_t) ! one byte per element
call hipCheck(hipSetDevice(0))
call hipCheck(hipStreamCreate(stream))
call hipCheck(hipMalloc(dptr, nbytes))
! Capture a device memset (value 5) into a graph.
call check_capture(hipStreamCaptureStatusNone, "before capture")
call hipCheck(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal))
call check_capture(hipStreamCaptureStatusActive, "during capture")
call hipCheck(hipMemsetAsync(dptr, 5, nbytes, stream))
call hipCheck(hipStreamEndCapture(stream, graph))
call check_capture(hipStreamCaptureStatusNone, "after capture")
if (.not. c_associated(graph)) then
write(*,*) "FAILED! captured graph is null"
call exit(1)
end if
! Instantiate and launch the executable graph.
call hipCheck(hipGraphInstantiate(gexec, graph, errnode, c_null_ptr, 0_c_size_t))
if (.not. c_associated(gexec)) then
write(*,*) "FAILED! instantiated graph is null"
call exit(1)
end if
call hipCheck(hipGraphLaunch(gexec, stream))
call hipCheck(hipStreamSynchronize(stream))
hbuf = 0
call hipCheck(hipMemcpy(c_loc(hbuf(1)), dptr, nbytes, hipMemcpyDeviceToHost))
do i = 1, n
if (hbuf(i) /= 5_c_int8_t) then
write(*,*) "FAILED! hbuf(", i, ") = ", hbuf(i), " (expected 5)"
call exit(1)
end if
end do
call hipCheck(hipGraphExecDestroy(gexec))
call hipCheck(hipGraphDestroy(graph))
call hipCheck(hipFree(dptr))
call hipCheck(hipStreamDestroy(stream))
write(*,*) "PASSED!"
contains
subroutine check_capture(want, what)
integer(kind(hipStreamCaptureStatusNone)), intent(in) :: want
character(len=*), intent(in) :: what
capstat = -1
call hipCheck(hipStreamIsCapturing(stream, c_loc(capstat)))
if (capstat /= want) then
write(*,*) "FAILED! capture status ", what, " = ", capstat, " (expected ", want, ")"
call exit(1)
end if
end subroutine check_capture
end program test_graph
A graph can also be built node by node, with the dependencies stated explicitly.
!!!!!!!!!!!!!!
! HIP runtime graphs built from explicit nodes (Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Builds a graph by hand rather than by stream capture: two 1-D memcpy nodes
! (H2D then D2H) linked with an explicit dependency, then instantiated and
! launched. Exercises hipGraphCreate, hipGraphAddMemcpyNode1D,
! hipGraphAddDependencies, hipGraphGetNodes, hipGraphInstantiate,
! hipGraphLaunch, hipGraphExecDestroy and hipGraphDestroy.
!!!!!!!!!!!!!!
!
program graph_nodes
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
implicit none
integer(c_int), parameter :: n = 256
real(c_double), target :: hsrc(n), hdst(n)
type(c_ptr) :: graph = c_null_ptr, gexec = c_null_ptr, errnode = c_null_ptr
type(c_ptr) :: stream = c_null_ptr, dptr = c_null_ptr
type(c_ptr) :: nodeH2D = c_null_ptr, nodeD2H = c_null_ptr
integer(c_size_t), target :: numnodes
type(c_ptr) :: nodes_out(8) ! capacity buffer for hipGraphGetNodes
integer(c_size_t) :: nbytes
integer :: i
write(*,"(a)",advance="no") "-- Running test 'hip graph_nodes' (Fortran 2003 interfaces) - "
nbytes = int(n, c_size_t) * 8
do i = 1, n
hsrc(i) = real(i, c_double)
end do
hdst = 0.0d0
call hipCheck(hipSetDevice(0))
call hipCheck(hipMalloc(dptr, nbytes))
call hipCheck(hipStreamCreate(stream))
call hipCheck(hipGraphCreate(graph, 0))
! Two 1-D memcpy nodes, initially without dependencies.
call hipCheck(hipGraphAddMemcpyNode1D(nodeH2D, graph, c_null_ptr, 0_c_size_t, &
dptr, c_loc(hsrc(1)), nbytes, hipMemcpyHostToDevice))
call hipCheck(hipGraphAddMemcpyNode1D(nodeD2H, graph, c_null_ptr, 0_c_size_t, &
c_loc(hdst(1)), dptr, nbytes, hipMemcpyDeviceToHost))
! Make the device->host copy depend on the host->device copy.
call hipCheck(hipGraphAddDependencies(graph, nodeH2D, nodeD2H, 1_c_size_t))
! Query the node count: pass an array and its capacity in numnodes; on return
! numnodes holds the actual count. (nodes is a by-reference c_ptr in the
! binding, so a real capacity buffer is used rather than a null query.)
numnodes = size(nodes_out, kind=c_size_t)
call hipCheck(hipGraphGetNodes(graph, nodes_out(1), c_loc(numnodes)))
if (numnodes /= 2) then
write(*,*) "FAILED! graph node count = ", numnodes, " (expected 2)"
call exit(1)
end if
! Instantiate and launch.
call hipCheck(hipGraphInstantiate(gexec, graph, errnode, c_null_ptr, 0_c_size_t))
call hipCheck(hipGraphLaunch(gexec, stream))
call hipCheck(hipStreamSynchronize(stream))
do i = 1, n
if (hdst(i) /= hsrc(i)) then
write(*,*) "FAILED! hdst(", i, ") = ", hdst(i), " expected ", hsrc(i)
call exit(1)
end if
end do
call hipCheck(hipGraphExecDestroy(gexec))
call hipCheck(hipGraphDestroy(graph))
call hipCheck(hipFree(dptr))
call hipCheck(hipStreamDestroy(stream))
write(*,*) "PASSED!"
end program graph_nodes
test/f2003/hip/graph_memset_node.f03 adds a memset node from a
hipMemsetParams structure, and test/f2003/hip/graph_empty_node.f03
builds a diamond shape with an empty node as the join point.
Launching a kernel#
Kernels themselves are written in HIP C++. The Fortran program calls a small
bind(c) launcher that the HIP compiler builds alongside it.
program fortran_hip
use iso_c_binding
use hipfort
use hipfort_check
implicit none
interface
subroutine launch(out,a,b,N) bind(c)
use iso_c_binding
implicit none
type(c_ptr) :: a, b, out
integer, value :: N
end subroutine
end interface
type(c_ptr) :: da = c_null_ptr
type(c_ptr) :: db = c_null_ptr
type(c_ptr) :: dout = c_null_ptr
integer, parameter :: N = 1000000
integer, parameter :: bytes_per_element = 8 !double precision
integer(c_size_t), parameter :: Nbytes = N*bytes_per_element
! Plain real should be equivalent to float
double precision,allocatable,target,dimension(:) :: a, b, out
double precision :: error
double precision, parameter :: error_max = 1.0d-10
integer :: i
type(hipDeviceProp_t),target :: props
!
call hipCheck(hipGetDeviceProperties(props,0))
write(*,"(a)",advance="no") "-- Running test 'vecadd' (Fortran 2003 interfaces)"
write(*,"(a)",advance="no") "- device: "
i=1
do while ( iachar(props%name(i)) .ne. 0 ) ! print till end char
write(*,"(a)",advance="no") props%name(i)
i = i+1
end do
write(*,"(a)",advance="no") " - "
! Allocate host memory
allocate(a(N))
allocate(b(N))
allocate(out(N))
! Initialize host arrays
a(:) = 1.0
b(:) = 2.0
! Allocate array space on the device
call hipCheck(hipMalloc(da,Nbytes))
call hipCheck(hipMalloc(db,Nbytes))
call hipCheck(hipMalloc(dout,Nbytes))
! Transfer data from host to device memory
call hipCheck(hipMemcpy(da, c_loc(a(1)), Nbytes, hipMemcpyHostToDevice))
call hipCheck(hipMemcpy(db, c_loc(b(1)), Nbytes, hipMemcpyHostToDevice))
call launch(dout,da,db,N)
call hipCheck(hipDeviceSynchronize())
! Transfer data back to host memory
call hipCheck(hipMemcpy(c_loc(out(1)), dout, Nbytes, hipMemcpyDeviceToHost))
! Verification
do i = 1,N
error = abs(out(i) - (a(i)+b(i)) )
if( error .gt. error_max ) then
write(*,*) "FAILED! Error bigger than max! Error = ", error, " Out = ", out(i)
call exit
endif
end do
call hipCheck(hipFree(da))
call hipCheck(hipFree(db))
call hipCheck(hipFree(dout))
! Deallocate host memory
deallocate(a)
deallocate(b)
deallocate(out)
write(*,*) "PASSED!"
end program fortran_hip
The kernel and its launcher:
#include <hip/hip_runtime.h>
#include <cstdio>
__global__ void vector_add(double *out, double *a, double *b, int n)
{
size_t index = blockIdx.x * blockDim.x + threadIdx.x;
size_t stride = blockDim.x * gridDim.x;
for (size_t i = index; i < n; i += stride)
out[i] = a[i] + b[i];
}
extern "C"
{
void launch(double **dout, double **da, double **db, int N)
{
//printf("launching kernel\n");
hipLaunchKernelGGL((vector_add), dim3(320), dim3(256), 0, 0, *dout, *da, *db, N);
}
}
Loading a code object#
The module API loads a kernel from a code object at run time, which avoids
linking any HIP C++ into the Fortran program. Build the code object with
hipcc --genco and look the kernel up by its mangled name.
!!!!!!!!!!!!!!
! HIP runtime module / kernel API (Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Loads a separately compiled code object, looks up the kernel it contains,
! queries its attributes and occupancy, then runs it as a graph kernel node,
! as an in-place update of the instantiated graph, and cooperatively.
! CTest passes the code object path in HIPFORT_TEST_CODE_OBJECT.
!!!!!!!!!!!!!!
!
program module_kernel
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
use hipfort_types
implicit none
integer(c_int), parameter :: n = 1024
character(len=22), target :: kname = "_Z10vector_addPdS_S_i"//c_null_char
character(len=4096) :: copath
character(len=4097), target :: cofile
integer :: pathlen, i
real(c_double), target :: ha(n), hb(n), hout(n)
type(c_ptr), target :: da = c_null_ptr, db = c_null_ptr
type(c_ptr), target :: dout = c_null_ptr, dout2 = c_null_ptr
integer(c_int), target :: nn
type(c_ptr), target :: args(4), args2(4)
type(c_ptr) :: hmod = c_null_ptr, kfunc = c_null_ptr
type(c_ptr) :: graph = c_null_ptr, gexec = c_null_ptr, knode = c_null_ptr
type(c_ptr) :: stream = c_null_ptr
type(hipKernelNodeParams) :: kparams
integer(c_size_t) :: nbytes
integer(c_int), target :: maxthreads
integer(c_int) :: gridsize, blocksize, numblocks
write(*,"(a)",advance="no") "-- Running test 'hip module_kernel' (Fortran 2003 interfaces) - "
call get_environment_variable("HIPFORT_TEST_CODE_OBJECT", copath, pathlen)
if (pathlen == 0) then
write(*,*) "FAILED! HIPFORT_TEST_CODE_OBJECT is not set"
call exit(1)
end if
cofile = copath(1:pathlen)//c_null_char
nbytes = int(n, c_size_t) * 8
nn = n
do i = 1, n
ha(i) = real(i, c_double)
hb(i) = real(2*i, c_double)
end do
call hipCheck(hipSetDevice(0))
call hipCheck(hipMalloc(da, nbytes))
call hipCheck(hipMalloc(db, nbytes))
call hipCheck(hipMalloc(dout, nbytes))
call hipCheck(hipMemcpy(da, c_loc(ha(1)), nbytes, hipMemcpyHostToDevice))
call hipCheck(hipMemcpy(db, c_loc(hb(1)), nbytes, hipMemcpyHostToDevice))
call hipCheck(hipModuleLoad(hmod, c_loc(cofile)))
call hipCheck(hipModuleGetFunction(kfunc, hmod, c_loc(kname)))
call hipCheck(hipFuncGetAttribute(c_loc(maxthreads), HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, kfunc))
if (maxthreads <= 0 .or. maxthreads > 1024) then
write(*,*) "FAILED! max threads per block = ", maxthreads
call exit(1)
end if
call hipCheck(hipModuleOccupancyMaxPotentialBlockSize(gridsize, blocksize, kfunc, 0_c_size_t, 0))
if (blocksize <= 0 .or. blocksize > maxthreads .or. gridsize <= 0) then
write(*,*) "FAILED! occupancy block size = ", blocksize, " grid size = ", gridsize
call exit(1)
end if
call hipCheck(hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(numblocks, kfunc, blocksize, 0_c_size_t))
if (numblocks <= 0) then
write(*,*) "FAILED! active blocks per CU = ", numblocks
call exit(1)
end if
! void* args[] = { &dout, &da, &db, &n }
args(1) = c_loc(dout)
args(2) = c_loc(da)
args(3) = c_loc(db)
args(4) = c_loc(nn)
call hipCheck(hipStreamCreate(stream))
call hipCheck(hipGraphCreate(graph, 0))
! Zero the output first so a node that never executes cannot pass.
call hipCheck(hipMemset(dout, 0, nbytes))
kparams%func = kfunc
kparams%gridDim = dim3(4, 1, 1)
kparams%blockDim = dim3(256, 1, 1)
kparams%sharedMemBytes = 0
kparams%kernelParams = c_loc(args(1))
kparams%extra = c_null_ptr
call hipCheck(hipGraphAddKernelNode(knode, graph, c_null_ptr, 0_c_size_t, kparams))
call hipCheck(hipGraphInstantiate(gexec, graph, c_null_ptr, c_null_ptr, 0_c_size_t))
call hipCheck(hipGraphLaunch(gexec, stream))
call hipCheck(hipStreamSynchronize(stream))
call check_result("hipGraphAddKernelNode", dout)
! Retarget the already-instantiated graph at a second output buffer.
call hipCheck(hipMalloc(dout2, nbytes))
call hipCheck(hipMemset(dout2, 0, nbytes))
args2(1) = c_loc(dout2)
args2(2) = c_loc(da)
args2(3) = c_loc(db)
args2(4) = c_loc(nn)
kparams%kernelParams = c_loc(args2(1))
call hipCheck(hipGraphExecKernelNodeSetParams(gexec, knode, kparams))
call hipCheck(hipGraphLaunch(gexec, stream))
call hipCheck(hipStreamSynchronize(stream))
call check_result("hipGraphExecKernelNodeSetParams", dout2)
! Cooperative launch takes no extra pointer, so the argument array goes as-is.
call hipCheck(hipMemset(dout, 0, nbytes))
call hipCheck(hipModuleLaunchCooperativeKernel(kfunc, 4, 1, 1, 256, 1, 1, 0, stream, args(1)))
call hipCheck(hipStreamSynchronize(stream))
call check_result("hipModuleLaunchCooperativeKernel", dout)
call hipCheck(hipGraphExecDestroy(gexec))
call hipCheck(hipGraphDestroy(graph))
call hipCheck(hipStreamDestroy(stream))
call hipCheck(hipModuleUnload(hmod))
call hipCheck(hipFree(da))
call hipCheck(hipFree(db))
call hipCheck(hipFree(dout))
call hipCheck(hipFree(dout2))
write(*,*) "PASSED!"
contains
subroutine check_result(what, buf)
character(len=*), intent(in) :: what
type(c_ptr), intent(in) :: buf
integer :: j
hout = 0.0d0
call hipCheck(hipMemcpy(c_loc(hout(1)), buf, nbytes, hipMemcpyDeviceToHost))
do j = 1, n
if (hout(j) /= ha(j) + hb(j)) then
write(*,*) "FAILED! ", what, " out(", j, ") = ", hout(j), " expected ", ha(j) + hb(j)
call exit(1)
end if
end do
end subroutine check_result
end program module_kernel
Occupancy#
The occupancy calculator reports how many blocks of a given size can be resident on a compute unit, and suggests a block size that maximizes occupancy. Both entry points take the host stub of a kernel.
!!!!!!!!!!!!!!
! HIP runtime occupancy calculator (Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Exercises hipOccupancyMaxPotentialBlockSize and
! hipOccupancyMaxActiveBlocksPerMultiprocessor on the kernel of the existing
! vecadd test, and cross-checks the results against the device properties.
!!!!!!!!!!!!!!
!
program occupancy
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
use hipfort_types
use hipfort_auxiliary
implicit none
! The occupancy entry points take the host stub of a kernel, so the kernel is
! named by the mangled symbol its HIP translation unit exports.
interface
subroutine vector_add(out, a, b, n) bind(c, name="_Z10vector_addPdS_S_i")
use iso_c_binding
implicit none
type(c_ptr), value :: out, a, b
integer(c_int), value :: n
end subroutine vector_add
end interface
type(hipDeviceProp_t) :: prop
integer(c_int), target :: gridsize, blocksize, nblocks, nblocks_small
write(*,"(a)",advance="no") "-- Running test 'hip occupancy' (Fortran 2003 interfaces) - "
call hipCheck(hipSetDevice(0))
call hipCheck(hipGetDeviceProperties(prop, 0))
gridsize = 0
blocksize = 0
call hipCheck(hipOccupancyMaxPotentialBlockSize(c_loc(gridsize), c_loc(blocksize), &
c_funloc(vector_add), 0_c_size_t, 0))
if (blocksize <= 0 .or. blocksize > prop%maxThreadsPerBlock) then
write(*,*) "FAILED! hipOccupancyMaxPotentialBlockSize block size", blocksize, &
" outside 1 ..", prop%maxThreadsPerBlock
call exit(1)
end if
if (gridsize <= 0) then
write(*,*) "FAILED! hipOccupancyMaxPotentialBlockSize grid size", gridsize
call exit(1)
end if
nblocks = 0
call hipCheck(hipOccupancyMaxActiveBlocksPerMultiprocessor(c_loc(nblocks), &
c_funloc(vector_add), &
blocksize, 0_c_size_t))
if (nblocks <= 0) then
write(*,*) "FAILED! hipOccupancyMaxActiveBlocksPerMultiprocessor returned", nblocks
call exit(1)
end if
if (nblocks * blocksize > prop%maxThreadsPerMultiProcessor) then
write(*,*) "FAILED! occupancy", nblocks, "blocks of", blocksize, &
"threads exceeds", prop%maxThreadsPerMultiProcessor
call exit(1)
end if
! Smaller blocks can never fit fewer times on a multiprocessor.
nblocks_small = 0
call hipCheck(hipOccupancyMaxActiveBlocksPerMultiprocessor(c_loc(nblocks_small), &
c_funloc(vector_add), &
blocksize/2, 0_c_size_t))
if (nblocks_small < nblocks) then
write(*,*) "FAILED!", nblocks_small, "blocks of", blocksize/2, "threads but", &
nblocks, "blocks of", blocksize
call exit(1)
end if
! Requesting all of the shared memory leaves room for a single block at most.
nblocks_small = 0
call hipCheck(hipOccupancyMaxActiveBlocksPerMultiprocessor(c_loc(nblocks_small), &
c_funloc(vector_add), &
blocksize, &
prop%sharedMemPerBlock))
if (nblocks_small > 1) then
write(*,*) "FAILED!", nblocks_small, "blocks fit while each claims all shared memory"
call exit(1)
end if
write(*,*) "PASSED!"
end program occupancy
hipModuleOccupancyMaxActiveBlocksPerMultiprocessor answers the same
question for a kernel loaded from a code object.
Cooperative launch#
A cooperative launch guarantees that every block of the grid is resident at the same time, which is what allows a kernel to synchronize across the whole grid. The grid is therefore limited by the occupancy of the kernel times the number of compute units, and a larger grid is rejected.
!!!!!!!!!!!!!!
! HIP runtime cooperative kernel launch (Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Launches the kernel of the existing vecadd test with
! hipLaunchCooperativeKernel on an occupancy-sized grid, checks the result, and
! confirms an oversized grid is rejected.
!!!!!!!!!!!!!!
!
program cooperative_launch
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
use hipfort_types
implicit none
! hipLaunchCooperativeKernel takes the host stub of a kernel, so the kernel is
! named by the mangled symbol its HIP translation unit exports.
interface
subroutine vector_add(out, a, b, n) bind(c, name="_Z10vector_addPdS_S_i")
use iso_c_binding
implicit none
type(c_ptr), value :: out, a, b
integer(c_int), value :: n
end subroutine vector_add
end interface
integer(c_int), parameter :: n = 4096
integer(c_int), parameter :: blocksize = 256
real(c_double), target :: ha(n), hb(n), hout(n)
type(c_ptr), target :: da = c_null_ptr, db = c_null_ptr, dout = c_null_ptr
type(c_ptr), target :: args(4)
integer(c_int), target :: nn, nblocks
integer(c_int) :: coop, ncu, i
integer(kind(hipSuccess)) :: stat
integer(c_size_t) :: nbytes
type(dim3) :: grid, block
write(*,"(a)",advance="no") "-- Running test 'hip cooperative_launch' (Fortran 2003 interfaces) - "
call hipCheck(hipSetDevice(0))
call hipCheck(hipDeviceGetAttribute(coop, hipDeviceAttributeCooperativeLaunch, 0))
if (coop == 0) then
write(*,*) "PASSED! (cooperative launch unsupported on this device)"
stop
end if
nbytes = int(n, c_size_t) * 8
nn = n
do i = 1, n
ha(i) = real(i, c_double)
hb(i) = real(2*i, c_double)
hout(i) = 0.0_c_double
end do
call hipCheck(hipMalloc(da, nbytes))
call hipCheck(hipMalloc(db, nbytes))
call hipCheck(hipMalloc(dout, nbytes))
call hipCheck(hipMemcpy(da, c_loc(ha(1)), nbytes, hipMemcpyHostToDevice))
call hipCheck(hipMemcpy(db, c_loc(hb(1)), nbytes, hipMemcpyHostToDevice))
call hipCheck(hipMemcpy(dout, c_loc(hout(1)), nbytes, hipMemcpyHostToDevice))
args(1) = c_loc(dout)
args(2) = c_loc(da)
args(3) = c_loc(db)
args(4) = c_loc(nn)
! A cooperative grid must be co-resident, so it is capped by the occupancy of
! the kernel times the number of compute units.
call hipCheck(hipOccupancyMaxActiveBlocksPerMultiprocessor(c_loc(nblocks), &
c_funloc(vector_add), &
blocksize, 0_c_size_t))
call hipCheck(hipDeviceGetAttribute(ncu, hipDeviceAttributeMultiprocessorCount, 0))
grid = dim3(nblocks * ncu, 1, 1)
block = dim3(blocksize, 1, 1)
call hipCheck(hipLaunchCooperativeKernel(transfer(c_funloc(vector_add), c_null_ptr), &
grid, block, args(1), 0, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call hipCheck(hipMemcpy(c_loc(hout(1)), dout, nbytes, hipMemcpyDeviceToHost))
do i = 1, n
if (abs(hout(i) - (ha(i) + hb(i))) > 1.0d-12) then
write(*,*) "FAILED! hout(", i, ") = ", hout(i), " expected ", ha(i) + hb(i)
call exit(1)
end if
end do
! More blocks than can be resident at once cannot be launched cooperatively.
grid = dim3(nblocks * ncu * 64, 1, 1)
stat = hipLaunchCooperativeKernel(transfer(c_funloc(vector_add), c_null_ptr), &
grid, block, args(1), 0, c_null_ptr)
if (stat /= hipErrorCooperativeLaunchTooLarge) then
write(*,*) "FAILED! oversized cooperative launch returned", stat, &
" expected", int(hipErrorCooperativeLaunchTooLarge)
call exit(1)
end if
stat = hipGetLastError()
call hipCheck(hipFree(da))
call hipCheck(hipFree(db))
call hipCheck(hipFree(dout))
write(*,*) "PASSED!"
end program cooperative_launch
Error handling and version queries#
HIP records the last error per thread. hipPeekAtLastError reads it and
hipGetLastError reads and clears it, and both a short name and a
description are available for any status code.
!!!!!!!!!!!!!!
! HIP runtime error-handling and version queries
! see: https://rocm.docs.amd.com/projects/HIP/en/latest/
!
! Exercises hipGetLastError, hipPeekAtLastError, hipGetErrorName,
! hipGetErrorString, hipRuntimeGetVersion and hipDriverGetVersion. Deliberately
! provokes hipErrorInvalidDevice via an out-of-range hipSetDevice and validates
! the peek-vs-clear semantics.
!!!!!!!!!!!!!!
!
program error_version
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_enums
implicit none
integer(kind(hipSuccess)) :: stat, stat2
integer(c_int) :: ndev, rver, dver
real(c_float), target :: hval
type(c_ptr) :: dptr = c_null_ptr
integer(c_size_t) :: nbytes
character(len=256) :: ename, estr
write(*,"(a)",advance="no") "-- Running test 'hip error_version' (Fortran 2003 interfaces) - "
call hipCheck(hipSetDevice(0))
! Clear any pre-existing sticky error (ignore its value).
stat = hipGetLastError()
! Provoke a well-defined error: device index far past hipGetDeviceCount.
call hipCheck(hipGetDeviceCount(ndev))
stat = hipSetDevice(ndev + 999)
if (stat /= hipErrorInvalidDevice) then
write(*,*) "FAILED! hipSetDevice(bad) returned", stat, &
" expected", int(hipErrorInvalidDevice)
call exit(1)
end if
! Peek must return the same error without clearing it.
stat = hipPeekAtLastError()
if (stat /= hipErrorInvalidDevice) then
write(*,*) "FAILED! first hipPeekAtLastError returned", stat
call exit(1)
end if
! Second peek: error still set, peek never clears.
stat2 = hipPeekAtLastError()
if (stat2 /= hipErrorInvalidDevice) then
write(*,*) "FAILED! second hipPeekAtLastError returned", stat2, &
"(should still be set)"
call exit(1)
end if
! GetLastError must return the error and then clear it.
stat = hipGetLastError()
if (stat /= hipErrorInvalidDevice) then
write(*,*) "FAILED! hipGetLastError returned", stat, &
" expected", int(hipErrorInvalidDevice)
call exit(1)
end if
! State is now cleared: next GetLastError must return hipSuccess.
stat2 = hipGetLastError()
if (stat2 /= hipSuccess) then
write(*,*) "FAILED! hipGetLastError after clear returned", stat2, &
"(expected hipSuccess)"
call exit(1)
end if
! Version queries: both must return a positive integer.
call hipCheck(hipRuntimeGetVersion(rver))
if (rver <= 0) then
write(*,*) "FAILED! hipRuntimeGetVersion =", rver
call exit(1)
end if
call hipCheck(hipDriverGetVersion(dver))
if (dver <= 0) then
write(*,*) "FAILED! hipDriverGetVersion =", dver
call exit(1)
end if
! Both accessors describe the error the runtime just reported.
ename = c_string(hipGetErrorName(hipErrorInvalidDevice))
if (trim(ename) /= "hipErrorInvalidDevice") then
write(*,*) "FAILED! hipGetErrorName = '", trim(ename), "'"
call exit(1)
end if
estr = c_string(hipGetErrorString(hipErrorInvalidDevice))
if (len_trim(estr) == 0 .or. index(estr, "device") == 0) then
write(*,*) "FAILED! hipGetErrorString = '", trim(estr), "'"
call exit(1)
end if
! Confirm the runtime still accepts work after the error cycle.
nbytes = int(4, c_size_t) ! one real(c_float)
call hipCheck(hipMalloc(dptr, nbytes))
hval = 3.14
call hipCheck(hipMemcpy(dptr, c_loc(hval), nbytes, hipMemcpyHostToDevice))
hval = 0.0
call hipCheck(hipMemcpy(c_loc(hval), dptr, nbytes, hipMemcpyDeviceToHost))
if (abs(hval - 3.14) > 1.0e-6) then
write(*,*) "FAILED! post-error memcpy roundtrip: got", hval
call exit(1)
end if
call hipCheck(hipFree(dptr))
write(*,*) "PASSED!"
contains
! Copy a null-terminated C string returned by reference into a Fortran string.
function c_string(cptr) result(res)
type(c_ptr), intent(in) :: cptr
character(len=256) :: res
character(kind=c_char), pointer :: chars(:)
integer :: i
res = " "
if (.not. c_associated(cptr)) return
call c_f_pointer(cptr, chars, [len(res)])
do i = 1, len(res)
if (chars(i) == c_null_char) exit
res(i:i) = chars(i)
end do
end function c_string
end program error_version
Peer access#
On a multi-GPU host, one device can address another device’s memory once peer access is enabled between them.
!!!!!!!!!!!!!!
! HIP peer access query/enable (Fortran 2003 interfaces)
! see: https:!rocm.docs.amd.com/projects/HIP/en/latest/
!
! Exercises hipDeviceCanAccessPeer (and enable/disable when a second device is
! present). On a single-GPU host a device is not its own peer, so the query must
! return 0 without error.
!!!!!!!!!!!!!!
!
program peer_access
use iso_c_binding
use hipfort
use hipfort_check
implicit none
integer(c_int) :: ndev, canAccess
write(*,"(a)",advance="no") "-- Running test 'hip peer_access' (Fortran 2003 interfaces) - "
call hipCheck(hipGetDeviceCount(ndev))
call hipCheck(hipSetDevice(0))
if (ndev >= 2) then
call hipCheck(hipDeviceCanAccessPeer(canAccess, 0, 1))
if (canAccess == 1) then
! Enable then disable peer access from device 0 to device 1.
call hipCheck(hipDeviceEnablePeerAccess(1, 0))
call hipCheck(hipDeviceDisablePeerAccess(1))
end if
else
! Single device: a device is not its own peer.
call hipCheck(hipDeviceCanAccessPeer(canAccess, 0, 0))
if (canAccess /= 0) then
write(*,*) "FAILED! canAccessPeer(0,0) = ", canAccess, " (expected 0)"
call exit(1)
end if
end if
write(*,*) "PASSED!"
end program peer_access