hipRAND examples#
hipRAND is a thin
layer over rocRAND whose API follows cuRAND. hipFORT exposes it through the
hipfort_hiprand module, together with hipfort_hiprand_enums for the
generator enumerations (HIPRAND_RNG_PSEUDO_PHILOX4_32_10 and so on). The
programs below import both explicitly.
Every program on this page is a complete, self-contained example that is built
and run as part of the hipFORT test suite. The Fortran 2008 version of each
program lives in test/f2008/hiprand, and the equivalent Fortran 2003
version, which uses type(c_ptr) device pointers and explicit byte counts
instead of Fortran array pointers, lives in test/f2003/hiprand.
If you want direct access to rocRAND rather than a cuRAND-style interface, see
the rocRAND examples, where the equivalent programs
are written against the hipfort_rocrand module. The two test suites cover
exactly the same generators and distributions, so the pages differ only in the
names of the entry points:
hiprandCreateGeneratorandhiprandDestroyGeneratorreplacerocrand_create_generatorandrocrand_destroy_generator.hiprandSetPseudoRandomGeneratorSeedreplacesrocrand_set_seed, andhiprandSetQuasiRandomGeneratorDimensionsreplacesrocrand_set_quasi_random_generator_dimensions.hiprandGenerateUniform,hiprandGenerateNormaland the rest replace therocrand_generate_*family.The generator enumerators are spelled
HIPRAND_RNG_*rather thanROCRAND_RNG_*, with the same suffixes.
The page is organized by generator, because that is the one choice a program makes up front. Each section shows one worked example and names the sibling programs that pair the same generator with the other distributions.
hipRAND call pattern#
A hipRAND program always follows the same sequence:
Create a generator with
hiprandCreateGenerator, passing the enumerator that selects the algorithm.Configure it: pseudo-random generators take a seed via
hiprandSetPseudoRandomGeneratorSeed, while quasi-random generators take a dimension count viahiprandSetQuasiRandomGeneratorDimensions.Allocate a device buffer to receive the numbers.
Call the
hiprandGenerate*routine for the distribution you want.Call
hipDeviceSynchronizebefore reading the results.Copy the buffer back to the host.
Free the device memory and release the generator with
hiprandDestroyGenerator.
Keep the following conventions in mind:
The programs generate directly into device memory. The C binding takes
type(c_ptr); the programs passc_loc(dx(1))to it explicitly, although an array overload that acceptsdxdirectly also exists. (hipRAND also provides host generators viahiprandCreateGeneratorHost, which are not covered here.)The count argument is the number of elements, typed
integer(c_size_t), not a byte count.A fixed seed makes a pseudo-random sequence reproducible, which is what lets these programs assert on sample statistics.
Every hipRAND call returns a status code. The examples wrap them in
hiprandCheckfrom thehipfort_checkmodule, which aborts on failure.
Building an example#
The examples only need the hiprand and hip hipFORT components:
find_package(hipfort REQUIRED COMPONENTS hip hiprand)
add_executable(my_rand philox_uniform.f08)
target_link_libraries(my_rand PRIVATE hipfort::hiprand hipfort::hip)
See Using hipFORT in your application for the full set of build options.
Output distributions#
Every generator on this page supports the same four combinations, which differ only in the routine called and the type of the output buffer:
hiprandGenerateUniformfills the buffer withreal(c_float)values distributed uniformly over(0, 1], so a correct sample has a mean near 0.5.hiprandGenerateUniformDoubleis thereal(c_double)form.hiprandGenerateNormaltakes a mean and a standard deviation in addition to the buffer and count, andhiprandGenerateNormalDoubleis its double-precision form.
Two further distributions, Poisson and log-normal, are exercised only with the Philox generator and are covered in their own section at the end of the page.
Pseudo-random generators#
The four pseudo-random generators are interchangeable in these programs: each
is created with its own enumerator, seeded with
hiprandSetPseudoRandomGeneratorSeed, and then used identically.
XORWOW#
HIPRAND_RNG_PSEUDO_XORWOW selects the xorshift-based XORWOW generator.
This example generates uniform single-precision values and checks that they
lie in (0, 1] and that the sample mean and standard deviation are close to
0.5 and 1/sqrt(12).
program hiprand_xorwow_uniform_test
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_hiprand
use hipfort_hiprand_enums
implicit none
integer(c_size_t), parameter :: N = 65536
integer(c_int64_t), parameter :: seed = 12345_c_int64_t
real(c_float), parameter :: expected_mean = 0.5, delta = 0.01
real(c_float), parameter :: expected_std = 0.28867513 ! 1/sqrt(12)
type(c_ptr) :: gen = c_null_ptr
real(c_float), allocatable, target, dimension(:) :: hx
real(c_float), pointer, dimension(:) :: dx
real(c_float) :: sample_mean, sample_std
write(*,"(a)",advance="no") "-- Running test 'XORWOW_UNIFORM' (Fortran 2008 interfaces) - "
! Create generator and set a fixed seed for reproducibility
call hiprandCheck(hiprandCreateGenerator(gen, HIPRAND_RNG_PSEUDO_XORWOW))
call hiprandCheck(hiprandSetPseudoRandomGeneratorSeed(gen, seed))
! Allocate host- and device-side memory
allocate(hx(N))
call hipCheck(hipMalloc(dx, mold=hx))
! Generate uniformly distributed floats on the device.
! output_data is type(c_ptr); take the address of the device pointer target.
call hiprandCheck(hiprandGenerateUniform(gen, c_loc(dx(1)), N))
call hipCheck(hipDeviceSynchronize())
! Transfer data back to host memory
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
! Verification: output in (0,1], with mean 1/2 and stddev 1/sqrt(12)
if (any(hx <= 0) .or. any(hx > 1)) then
write(*,*) "FAILED! value outside (0,1]: ", minval(hx), maxval(hx)
call exit(1)
end if
sample_mean = sum(hx) / real(N)
if (abs(sample_mean - expected_mean) > delta) then
write(*,*) "FAILED! mean out of tolerance: ", sample_mean
call exit(1)
end if
sample_std = sqrt(sum((hx - expected_mean)**2) / real(N))
if (abs(sample_std - expected_std) > delta) then
write(*,*) "FAILED! stddev out of tolerance: ", sample_std
call exit(1)
end if
! Cleanup
call hipCheck(hipFree(dx))
deallocate(hx)
call hiprandCheck(hiprandDestroyGenerator(gen))
write(*,*) "PASSED!"
end program hiprand_xorwow_uniform_test
test/f2008/hiprand/xorwow_uniform_double.f08, xorwow_normal.f08 and
xorwow_normal_double.f08 cover the remaining distributions.
Philox#
HIPRAND_RNG_PSEUDO_PHILOX4_32_10 selects the counter-based Philox
generator. This example uses the normal distribution, so it also shows the
extra mean and stddev arguments; it checks both the sample mean and
the sample standard deviation, since a generator producing the right mean with
the wrong spread would otherwise pass.
program hiprand_philox_normal_test
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_hiprand
use hipfort_hiprand_enums
implicit none
integer(c_size_t), parameter :: N = 65536
integer(c_int64_t), parameter :: seed = 12345_c_int64_t
! Not the standard normal, so a binding that drops mean or stddev cannot pass.
real(c_float), parameter :: mean = 2.0, stddev = 0.5, delta = 0.05
type(c_ptr) :: gen = c_null_ptr
real(c_float), allocatable, target, dimension(:) :: hx
real(c_float), pointer, dimension(:) :: dx
real(c_float) :: sample_mean, sample_std
write(*,"(a)",advance="no") "-- Running test 'PHILOX_NORMAL' (Fortran 2008 interfaces) - "
! Create generator and set a fixed seed for reproducibility
call hiprandCheck(hiprandCreateGenerator(gen, HIPRAND_RNG_PSEUDO_PHILOX4_32_10))
call hiprandCheck(hiprandSetPseudoRandomGeneratorSeed(gen, seed))
! Allocate host- and device-side memory
allocate(hx(N))
call hipCheck(hipMalloc(dx, mold=hx))
! Generate normally distributed floats on the device.
! output_data is type(c_ptr); take the address of the device pointer target.
call hiprandCheck(hiprandGenerateNormal(gen, c_loc(dx(1)), N, mean, stddev))
call hipCheck(hipDeviceSynchronize())
! Transfer data back to host memory
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
! Verification: sample mean and stddev should match requested parameters
sample_mean = sum(hx) / real(N)
sample_std = sqrt(sum((hx - mean)**2) / real(N))
if (abs(sample_mean - mean) > delta) then
write(*,*) "FAILED! mean out of tolerance: ", sample_mean
call exit(1)
end if
if (abs(sample_std - stddev) > delta) then
write(*,*) "FAILED! stddev out of tolerance: ", sample_std
call exit(1)
end if
! Cleanup
call hipCheck(hipFree(dx))
deallocate(hx)
call hiprandCheck(hiprandDestroyGenerator(gen))
write(*,*) "PASSED!"
end program hiprand_philox_normal_test
test/f2008/hiprand/philox_uniform.f08, philox_uniform_double.f08 and
philox_normal_double.f08 cover the remaining distributions. Philox is also
the generator used for the Poisson and log-normal programs below.
MRG32K3A#
HIPRAND_RNG_PSEUDO_MRG32K3A selects the combined multiple-recursive
generator.
program hiprand_mrg32k3a_uniform_test
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_hiprand
use hipfort_hiprand_enums
implicit none
integer(c_size_t), parameter :: N = 65536
integer(c_int64_t), parameter :: seed = 12345_c_int64_t
real(c_float), parameter :: expected_mean = 0.5, delta = 0.01
real(c_float), parameter :: expected_std = 0.28867513 ! 1/sqrt(12)
type(c_ptr) :: gen = c_null_ptr
real(c_float), allocatable, target, dimension(:) :: hx
real(c_float), pointer, dimension(:) :: dx
real(c_float) :: sample_mean, sample_std
write(*,"(a)",advance="no") "-- Running test 'MRG32K3A_UNIFORM' (Fortran 2008 interfaces) - "
call hiprandCheck(hiprandCreateGenerator(gen, HIPRAND_RNG_PSEUDO_MRG32K3A))
call hiprandCheck(hiprandSetPseudoRandomGeneratorSeed(gen, seed))
allocate(hx(N))
call hipCheck(hipMalloc(dx, mold=hx))
! Uniform (0,1] samples have mean 1/2 and stddev 1/sqrt(12).
call hiprandCheck(hiprandGenerateUniform(gen, c_loc(dx(1)), N))
call hipCheck(hipDeviceSynchronize())
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
if (any(hx <= 0) .or. any(hx > 1)) then
write(*,*) "FAILED! value outside (0,1]: ", minval(hx), maxval(hx)
call exit(1)
end if
sample_mean = sum(hx) / real(N, kind(sample_mean))
if (abs(sample_mean - expected_mean) > delta) then
write(*,*) "FAILED! mean out of tolerance: ", sample_mean
call exit(1)
end if
sample_std = sqrt(sum((hx - expected_mean)**2) / real(N, kind(sample_std)))
if (abs(sample_std - expected_std) > delta) then
write(*,*) "FAILED! stddev out of tolerance: ", sample_std
call exit(1)
end if
call hipCheck(hipFree(dx))
deallocate(hx)
call hiprandCheck(hiprandDestroyGenerator(gen))
write(*,*) "PASSED!"
end program hiprand_mrg32k3a_uniform_test
test/f2008/hiprand/mrg32k3a_uniform_double.f08, mrg32k3a_normal.f08
and mrg32k3a_normal_double.f08 cover the remaining distributions.
MTGP32#
HIPRAND_RNG_PSEUDO_MTGP32 selects the Mersenne Twister for graphics
processors. This example uses the normal distribution.
program hiprand_mtgp32_normal_test
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_hiprand
use hipfort_hiprand_enums
implicit none
integer(c_size_t), parameter :: N = 65536
integer(c_int64_t), parameter :: seed = 12345_c_int64_t
! Not the standard normal, so a binding that drops mean or stddev cannot pass.
real(c_float), parameter :: mean = 2.0, stddev = 0.5, delta = 0.05
type(c_ptr) :: gen = c_null_ptr
real(c_float), allocatable, target, dimension(:) :: hx
real(c_float), pointer, dimension(:) :: dx
real(c_float) :: sample_mean
real(c_float) :: sample_std
write(*,"(a)",advance="no") "-- Running test 'MTGP32_NORMAL' (Fortran 2008 interfaces) - "
call hiprandCheck(hiprandCreateGenerator(gen, HIPRAND_RNG_PSEUDO_MTGP32))
call hiprandCheck(hiprandSetPseudoRandomGeneratorSeed(gen, seed))
allocate(hx(N))
call hipCheck(hipMalloc(dx, mold=hx))
! Normal(mean, stddev) samples recover both moments.
call hiprandCheck(hiprandGenerateNormal(gen, c_loc(dx(1)), N, mean, stddev))
call hipCheck(hipDeviceSynchronize())
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
sample_mean = sum(hx) / real(N, kind(sample_mean))
sample_std = sqrt(sum((hx - sample_mean)**2) / real(N, kind(sample_mean)))
if (abs(sample_mean - mean) > delta) then
write(*,*) "FAILED! mean out of tolerance: ", sample_mean
call exit(1)
end if
if (abs(sample_std - stddev) > delta) then
write(*,*) "FAILED! stddev out of tolerance: ", sample_std
call exit(1)
end if
call hipCheck(hipFree(dx))
deallocate(hx)
call hiprandCheck(hiprandDestroyGenerator(gen))
write(*,*) "PASSED!"
end program hiprand_mtgp32_normal_test
test/f2008/hiprand/mtgp32_uniform.f08, mtgp32_uniform_double.f08 and
mtgp32_normal_double.f08 cover the remaining distributions.
Quasi-random generators#
Sobol32#
HIPRAND_RNG_QUASI_SOBOL32 produces a low-discrepancy sequence, which fills
the sample space more evenly than a pseudo-random stream. It is configured
with a dimension count instead of a seed, so
hiprandSetQuasiRandomGeneratorDimensions replaces the
hiprandSetPseudoRandomGeneratorSeed call; everything else is unchanged.
program hiprand_sobol32_uniform_test
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_hiprand
use hipfort_hiprand_enums
implicit none
integer(c_size_t), parameter :: N = 65536
integer(c_int), parameter :: dimensions = 1
real(c_float), parameter :: expected_mean = 0.5, delta = 0.01
real(c_float), parameter :: expected_std = 0.28867513 ! 1/sqrt(12)
type(c_ptr) :: gen = c_null_ptr
real(c_float), allocatable, target, dimension(:) :: hx
real(c_float), pointer, dimension(:) :: dx
real(c_float) :: sample_mean, sample_std
write(*,"(a)",advance="no") "-- Running test 'SOBOL32_UNIFORM' (Fortran 2008 interfaces) - "
! Sobol32 is a quasi-random (low-discrepancy) generator: it takes a
! dimension count rather than a seed.
call hiprandCheck(hiprandCreateGenerator(gen, HIPRAND_RNG_QUASI_SOBOL32))
call hiprandCheck(hiprandSetQuasiRandomGeneratorDimensions(gen, dimensions))
allocate(hx(N))
call hipCheck(hipMalloc(dx, mold=hx))
! Uniform (0,1] samples have mean 1/2 and stddev 1/sqrt(12).
call hiprandCheck(hiprandGenerateUniform(gen, c_loc(dx(1)), N))
call hipCheck(hipDeviceSynchronize())
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
if (any(hx <= 0) .or. any(hx > 1)) then
write(*,*) "FAILED! value outside (0,1]: ", minval(hx), maxval(hx)
call exit(1)
end if
sample_mean = sum(hx) / real(N, kind(sample_mean))
if (abs(sample_mean - expected_mean) > delta) then
write(*,*) "FAILED! mean out of tolerance: ", sample_mean
call exit(1)
end if
sample_std = sqrt(sum((hx - expected_mean)**2) / real(N, kind(sample_std)))
if (abs(sample_std - expected_std) > delta) then
write(*,*) "FAILED! stddev out of tolerance: ", sample_std
call exit(1)
end if
call hipCheck(hipFree(dx))
deallocate(hx)
call hiprandCheck(hiprandDestroyGenerator(gen))
write(*,*) "PASSED!"
end program hiprand_sobol32_uniform_test
test/f2008/hiprand/sobol32_uniform_double.f08, sobol32_normal.f08 and
sobol32_normal_double.f08 cover the remaining distributions.
Poisson and log-normal distributions#
These two distributions are exercised with the Philox generator only, so that a failure points at the distribution entry point rather than at the generator.
Poisson#
hiprandGeneratePoisson draws from a Poisson distribution with parameter
lambda. Unlike the other distributions it produces unsigned 32-bit
integers rather than floating-point values, so the host and device buffers
are integer(c_int) while lambda remains real(c_double). A Poisson
distribution has both mean and variance lambda, which is what the program
checks.
program hiprand_philox_poisson_test
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_hiprand
use hipfort_hiprand_enums
implicit none
integer(c_size_t), parameter :: N = 65536
integer(c_int64_t), parameter :: seed = 12345_c_int64_t
real(c_double), parameter :: lambda = 10.0_c_double
real(c_double), parameter :: delta = 0.1_c_double
! Poisson variance is lambda; allow 5% of lambda, but at least 0.5.
real(c_double), parameter :: var_delta = max(0.05_c_double * lambda, 0.5_c_double)
type(c_ptr) :: gen = c_null_ptr
integer(c_int), allocatable, target, dimension(:) :: hx
integer(c_int), pointer, dimension(:) :: dx
real(c_double) :: sample_mean, sample_var
write(*,"(a)",advance="no") "-- Running test 'PHILOX_POISSON' (Fortran 2008 interfaces) - "
call hiprandCheck(hiprandCreateGenerator(gen, HIPRAND_RNG_PSEUDO_PHILOX4_32_10))
call hiprandCheck(hiprandSetPseudoRandomGeneratorSeed(gen, seed))
allocate(hx(N))
call hipCheck(hipMalloc(dx, mold=hx))
! Poisson(lambda) samples are unsigned 32-bit integers with mean and variance lambda.
call hiprandCheck(hiprandGeneratePoisson(gen, c_loc(dx(1)), N, lambda))
call hipCheck(hipDeviceSynchronize())
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
sample_mean = real(sum(hx), c_double) / real(N, c_double)
if (abs(sample_mean - lambda) > delta) then
write(*,*) "FAILED! mean out of tolerance: ", sample_mean
call exit(1)
end if
sample_var = sum((real(hx, c_double) - sample_mean)**2) / real(N, c_double)
if (abs(sample_var - lambda) > var_delta) then
write(*,*) "FAILED! variance out of tolerance: ", sample_var
call exit(1)
end if
call hipCheck(hipFree(dx))
deallocate(hx)
call hiprandCheck(hiprandDestroyGenerator(gen))
write(*,*) "PASSED!"
end program hiprand_philox_poisson_test
This is the only hipRAND Poisson program; its Fortran 2003 version is
test/f2003/hiprand/philox_poisson.f03.
Log-normal#
hiprandGenerateLogNormal produces values whose logarithm is normally
distributed with the given mean and standard deviation. Every sample is
therefore strictly positive, and the program verifies the distribution by
taking logs and checking the resulting mean and standard deviation.
program hiprand_philox_log_normal_test
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_hiprand
use hipfort_hiprand_enums
implicit none
integer(c_size_t), parameter :: N = 65536
integer(c_int64_t), parameter :: seed = 12345_c_int64_t
! Not the standard normal, so a binding that drops mean or stddev cannot pass.
real(c_float), parameter :: mean = 2.0, stddev = 0.5, delta = 0.05
type(c_ptr) :: gen = c_null_ptr
real(c_float), allocatable, target, dimension(:) :: hx
real(c_float), pointer, dimension(:) :: dx
real(c_float) :: sample_mean
real(c_float) :: sample_std
write(*,"(a)",advance="no") "-- Running test 'PHILOX_LOG_NORMAL' (Fortran 2008 interfaces) - "
call hiprandCheck(hiprandCreateGenerator(gen, HIPRAND_RNG_PSEUDO_PHILOX4_32_10))
call hiprandCheck(hiprandSetPseudoRandomGeneratorSeed(gen, seed))
allocate(hx(N))
call hipCheck(hipMalloc(dx, mold=hx))
! If x is log-normal with parameters (mean, stddev) then log(x) is
! normal with those moments. Log-normal samples are strictly
! positive, so the logarithm is always defined.
call hiprandCheck(hiprandGenerateLogNormal(gen, c_loc(dx(1)), N, mean, stddev))
call hipCheck(hipDeviceSynchronize())
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
if (any(hx <= 0.0)) then
write(*,*) "FAILED! log-normal sample was not positive"
call exit(1)
end if
sample_mean = sum(log(hx)) / real(N, kind(sample_mean))
sample_std = sqrt(sum((log(hx) - sample_mean)**2) / real(N, kind(sample_mean)))
if (abs(sample_mean - mean) > delta) then
write(*,*) "FAILED! mean of log(x) out of tolerance: ", sample_mean
call exit(1)
end if
if (abs(sample_std - stddev) > delta) then
write(*,*) "FAILED! stddev of log(x) out of tolerance: ", sample_std
call exit(1)
end if
call hipCheck(hipFree(dx))
deallocate(hx)
call hiprandCheck(hiprandDestroyGenerator(gen))
write(*,*) "PASSED!"
end program hiprand_philox_log_normal_test
hiprandGenerateLogNormalDouble is the double-precision form; see
test/f2008/hiprand/philox_log_normal_double.f08.
Querying the hipRAND version#
hiprandGetVersion writes the library version to a plain integer(c_int)
argument, so the variable is passed directly rather than through c_loc. On
ROCm it reports the version of the underlying rocRAND library, encoded as
major*100000 + minor*100 + patch.
program hiprand_version_test
use iso_c_binding
use hipfort_check
use hipfort_hiprand
use hipfort_hiprand_enums
implicit none
integer(c_int), parameter :: poison = -12345
integer(c_int) :: version, version2
integer(c_int) :: major, minor, patch
write(*,"(a)",advance="no") "-- Running test 'hipRAND version' &
&(Fortran 2008 interfaces) - "
version = poison
version2 = poison
! Output scalar passed directly -- the dummy is 'integer(c_int) :: version'.
call hiprandCheck(hiprandGetVersion(version))
call hiprandCheck(hiprandGetVersion(version2))
if (version == poison) then
write(*,*) "FAILED! hiprandGetVersion did not write the version argument"
STOP 1
end if
if (version <= 0) then
write(*,*) "FAILED! hiprandGetVersion returned a non-positive version code: ", version
STOP 1
end if
if (version2 /= version) then
write(*,*) "FAILED! hiprandGetVersion is not stable: ", version, " then ", version2
STOP 1
end if
! HIPRAND_VERSION / 100000 is major, / 100 % 1000 is minor, % 100 is patch.
! On the CUDA backend this is cuRAND's version, packed major*1000 +
! minor*100 + patch (CURAND_VERSION), so a code below 100000 decodes that way.
if (version >= 100000) then
major = version / 100000
minor = mod(version / 100, 1000)
else
major = version / 1000
minor = mod(version / 100, 10)
end if
patch = mod(version, 100)
if (major < 1 .or. major > 99) then
write(*,*) "FAILED! implausible hipRAND major version: ", major
STOP 1
end if
write(*,"(a,i0,a,i0,a,i0,a,i0,a)") " PASSED! hipRAND version: ", &
major, ".", minor, ".", patch, " (code ", version, ")"
end program hiprand_version_test