rocFFT examples#
rocFFT is the AMD
implementation of the fast Fourier transform for AMD GPUs. hipFORT exposes it
through the hipfort_rocfft module, which mirrors the rocFFT C API one to
one.
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 sources live in
test/f2008/rocfft and the equivalent Fortran 2003 sources, which use
type(c_ptr) device pointers and explicit byte counts instead of Fortran
array pointers, live in test/f2003/rocfft.
Most of the examples are Fortran counterparts of the C++ samples shipped with
rocFFT in clients/samples/rocfft:
rocFFT sample |
hipFORT example |
|---|---|
|
Complex-to-complex transform, Out-of-place transforms, Managing the work buffer |
|
Real-to-complex and complex-to-real transforms, In-place real transforms |
|
|
|
Not available: load and store callbacks require device functions, which cannot be written in Fortran. |
The remaining examples cover material from the rocFFT how-to guides: Normalizing with a scale factor, Inspecting a plan and Reusing compiled kernels.
Two areas have no Fortran counterpart. Distributed transforms, which the
clients/samples/multi_gpu sample demonstrates, are built on the rocFFT
field and brick API that is still an experimental preview, and MPI transforms
need a rocFFT built with MPI support. The hipFORT interfaces for both
(rocfft_field_create, rocfft_brick_create,
rocfft_plan_description_set_comm and friends) are generated and callable,
but they are not exercised by the test suite.
Transform workflow#
A rocFFT transform always follows the same sequence:
Call
rocfft_setuponce before any other rocFFT call.Optionally create a plan description with
rocfft_plan_description_createto set a data layout, a scale factor, or other advanced properties.Create a plan with
rocfft_plan_create, passing the placement (in-place or not), the transform type, the precision, the rank, the transform lengths, and the batch size.Optionally create an execution info handle with
rocfft_execution_info_createto supply a HIP stream or a work buffer.Run the transform with
rocfft_execute.Release the plan with
rocfft_plan_destroyand callrocfft_cleanupwhen the application is done with rocFFT.
Keep the following conventions in mind:
rocFFT transforms are unnormalized. A forward transform followed by an inverse transform of length
NreturnsNtimes the original data, unless a scale factor is attached to one of the plans.The
lengthsarray passed torocfft_plan_createstarts with the fastest-varying dimension, which matches Fortran’s column-major storage: for anNxbyNytransform,lengths = [Nx, Ny].Real forward transforms produce Hermitian-symmetric output, so only
N/2 + 1complex values are stored. Size the complex buffer accordingly.rocfft_executetakes arrays of buffer pointers. In Fortran you pass atype(c_ptr)expression, such asc_loc(dx), and the compiler passes its address. For in-place transforms the output argument isc_null_ptr.Every rocFFT call returns a status code. The examples wrap them in
rocfftCheckfrom thehipfort_checkmodule, which aborts on failure.
Building an example#
The examples only need the rocfft and hip hipFORT components:
find_package(hipfort REQUIRED COMPONENTS hip rocfft)
add_executable(my_fft rocfft_c2c_1d_z.f08)
target_link_libraries(my_fft PRIVATE hipfort::rocfft hipfort::hip)
See Using hipFORT in your application for the full set of build options.
Complex-to-complex transform#
The simplest case: an in-place, single-batch, one-dimensional complex-to-complex
transform in double precision. The program runs a forward transform followed by
an inverse transform and checks that the result is N times the input, which
demonstrates that rocFFT does not normalize.
program rocfft_c2c_1d_z
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: N = 16
complex(8), allocatable, target, dimension(:) :: hx, hx_input
complex(8), pointer, dimension(:) :: dx => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: one = 1
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
allocate(hx(N))
allocate(hx_input(N))
do i = 1, N
hx(i) = cmplx(dble(i), dble(N - i), kind=8)
end do
hx_input(:) = hx(:)
call hipCheck(hipMalloc(dx, source=hx))
write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D double (z)' (Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Forward transform (in-place).
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
! Inverse transform (in-place). rocFFT is unnormalized, so this yields N*input.
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dx))
! After forward+inverse the data should equal N times the original input.
do i = 1, N
error = abs(hx(i) - N * hx_input(i))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", N * hx_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_c2c_1d_z
Use rocfft_precision_single and complex(4) host data for a single
precision transform, as in test/f2008/rocfft/rocfft_c2c_1d_c.f08.
Real-to-complex and complex-to-real transforms#
Real transforms use rocfft_transform_type_real_forward and
rocfft_transform_type_real_inverse. Because the spectrum of real data is
Hermitian symmetric, the complex buffer holds N/2 + 1 elements.
program rocfft_r2c_c2r_1d_d
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: N = 16
integer(c_size_t), parameter :: Ncomplex = N/2 + 1
real(8), allocatable, target, dimension(:) :: hr, hr_input
complex(8), pointer, dimension(:) :: dc => null()
real(8), pointer, dimension(:) :: dr => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: one = 1
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
allocate(hr(N))
allocate(hr_input(N))
do i = 1, N
hr(i) = dble(i) + dble(mod(i,3)) - dble(mod(i,7))
end do
hr_input(:) = hr(:)
! Device buffers: real input of length N, complex output of length N/2+1.
call hipCheck(hipMalloc(dr, source=hr))
call hipCheck(hipMalloc(dc, Ncomplex))
write(*,"(a)",advance="no") "-- Running test 'rocFFT R2C/C2R 1D double (d)' (Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Forward real-to-complex (out-of-place): dr -> dc.
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_notinplace,&
rocfft_transform_type_real_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dr), c_loc(dc), c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
! Inverse complex-to-real (out-of-place): dc -> dr. rocFFT unnormalized -> N*input.
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_notinplace,&
rocfft_transform_type_real_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dc), c_loc(dr), c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call hipCheck(hipMemcpy(hr, dr, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dr))
call hipCheck(hipFree(dc))
! After forward+inverse the real data should equal N times the original input.
do i = 1, N
error = abs(hr(i) - N * hr_input(i))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hr(i), " expected=", N * hr_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hr)
deallocate(hr_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_r2c_c2r_1d_d
In-place real transforms#
An in-place real transform reads real values and writes N/2 + 1 complex
values into the same allocation, so the real buffer must be padded to
2*(N/2 + 1) reals: two extra reals in the contiguous dimension. The input
and output array types are declared on a plan description with
rocfft_plan_description_set_data_layout.
! Fortran counterpart of the in-place path of the rocFFT sample
! clients/samples/rocfft/rocfft_example_realcomplex.cpp: an in-place real
! transform, where the real buffer is padded so that it can also hold the
! Hermitian-symmetric spectrum.
program rocfft_r2c_c2r_1d_inplace_d
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: N = 16
integer(c_size_t), parameter :: Ncomplex = N/2 + 1
! An in-place real transform needs room for N/2+1 complex values, that is
! two extra reals in the contiguous dimension.
integer(c_size_t), parameter :: Npad = 2 * Ncomplex
real(8), allocatable, target, dimension(:) :: hr, hr_input
real(8), pointer, dimension(:) :: dr => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
type(c_ptr) :: desc_fwd = c_null_ptr
type(c_ptr) :: desc_bwd = c_null_ptr
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), target :: strides(1)
integer(c_size_t), parameter :: one = 1
integer(c_size_t), parameter :: zero = 0
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
strides(1) = 1
allocate(hr(Npad))
allocate(hr_input(Npad))
hr(:) = 0.0d0
do i = 1, N
hr(i) = dble(i) + dble(mod(i,3)) - dble(mod(i,7))
end do
hr_input(:) = hr(:)
call hipCheck(hipMalloc(dr, source=hr))
write(*,"(a)",advance="no") "-- Running test 'rocFFT R2C/C2R 1D in-place double (d)' &
&(Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Forward: real input, Hermitian-interleaved output, both in the same buffer.
call rocfftCheck(rocfft_plan_description_create(desc_fwd))
call rocfftCheck(rocfft_plan_description_set_data_layout(desc_fwd,&
rocfft_array_type_real,&
rocfft_array_type_hermitian_interleaved,&
c_null_ptr, c_null_ptr,&
one, c_loc(strides), zero,&
one, c_loc(strides), zero))
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_real_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
desc_fwd))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dr), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
call rocfftCheck(rocfft_plan_description_destroy(desc_fwd))
! Inverse: the array types are swapped. rocFFT is unnormalized, so the round
! trip yields N*input.
call rocfftCheck(rocfft_plan_description_create(desc_bwd))
call rocfftCheck(rocfft_plan_description_set_data_layout(desc_bwd,&
rocfft_array_type_hermitian_interleaved,&
rocfft_array_type_real,&
c_null_ptr, c_null_ptr,&
one, c_loc(strides), zero,&
one, c_loc(strides), zero))
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_real_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
desc_bwd))
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dr), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call rocfftCheck(rocfft_plan_description_destroy(desc_bwd))
call hipCheck(hipMemcpy(hr, dr, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dr))
! Only the first N reals hold the signal; the padding is scratch space.
do i = 1, N
error = abs(hr(i) - N * hr_input(i))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hr(i), &
" expected=", N * hr_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hr)
deallocate(hr_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_r2c_c2r_1d_inplace_d
Multi-dimensional transforms#
A multi-dimensional transform only differs in the rank argument and the number
of entries in the lengths array. The first entry is the fastest-varying
dimension.
program rocfft_c2c_2d_z
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: Nx = 4, Ny = 8
integer(c_size_t), parameter :: Ntot = Nx * Ny
complex(8), allocatable, target, dimension(:) :: hx, hx_input
complex(8), pointer, dimension(:) :: dx => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: two = 2
integer(c_size_t), parameter :: one = 1
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(2))
lengths(1) = Nx
lengths(2) = Ny
allocate(hx(Ntot))
allocate(hx_input(Ntot))
do i = 1, Ntot
hx(i) = cmplx(dble(i), dble(Ntot - i), kind=8)
end do
hx_input(:) = hx(:)
call hipCheck(hipMalloc(dx, source=hx))
write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 2D double (z)' (Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Forward 2D transform (in-place).
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
two,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
! Inverse 2D transform (in-place). rocFFT unnormalized -> Ntot*input.
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
two,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dx))
! After forward+inverse the data should equal Ntot times the original input.
do i = 1, Ntot
error = abs(hx(i) - Ntot * hx_input(i))
if (error > error_max * Ntot) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", Ntot * hx_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_c2c_2d_z
test/f2008/rocfft/rocfft_c2c_3d_z.f08 extends the same pattern to three
dimensions.
Batched transforms#
To transform many signals with one plan, pass the batch count as the
number_of_transforms argument and describe the memory layout with
rocfft_plan_description_set_data_layout. The distance argument gives the
number of elements between the start of consecutive signals.
program rocfft_c2c_1d_batched_z
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: N = 16
integer(c_size_t), parameter :: Nbatch = 4
integer(c_size_t), parameter :: Ntot = N * Nbatch
complex(8), allocatable, target, dimension(:) :: hx, hx_input
complex(8), pointer, dimension(:) :: dx => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
type(c_ptr) :: desc = c_null_ptr
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), target :: strides(1)
integer(c_size_t), parameter :: one = 1
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
strides(1) = 1
allocate(hx(Ntot))
allocate(hx_input(Ntot))
do i = 1, Ntot
hx(i) = cmplx(dble(i), dble(Ntot - i), kind=8)
end do
hx_input(:) = hx(:)
call hipCheck(hipMalloc(dx, source=hx))
write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D batched double (z)' (Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Describe a batched layout: contiguous, stride 1, distance N between transforms.
call rocfftCheck(rocfft_plan_description_create(desc))
call rocfftCheck(rocfft_plan_description_set_data_layout(desc,&
rocfft_array_type_complex_interleaved,&
rocfft_array_type_complex_interleaved,&
c_null_ptr, c_null_ptr,&
one, c_loc(strides), N,&
one, c_loc(strides), N))
! Forward batched transform (in-place).
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
Nbatch,&
desc))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
! Inverse batched transform (in-place). rocFFT unnormalized -> N*input per transform.
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
Nbatch,&
desc))
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call rocfftCheck(rocfft_plan_description_destroy(desc))
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dx))
! After forward+inverse each batch element should equal N times the original input.
do i = 1, Ntot
error = abs(hx(i) - N * hx_input(i))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), " expected=", N * hx_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_c2c_1d_batched_z
Out-of-place transforms#
With rocfft_placement_notinplace, the result is written to a separate
buffer that is passed as the third argument of rocfft_execute. Note that
rocFFT is allowed to overwrite the input buffer of an out-of-place transform,
so do not rely on its contents afterwards.
! Fortran counterpart of the out-of-place path of the rocFFT sample
! clients/samples/rocfft/rocfft_example_complexcomplex.cpp (--outofplace).
program rocfft_c2c_1d_notinplace_z
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: N = 16
complex(8), allocatable, target, dimension(:) :: hx, hx_input
complex(8), pointer, dimension(:) :: dx => null() ! device input
complex(8), pointer, dimension(:) :: dy => null() ! device output
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: one = 1
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
allocate(hx(N))
allocate(hx_input(N))
do i = 1, N
hx(i) = cmplx(dble(i), dble(N - i), kind=8)
end do
hx_input(:) = hx(:)
! Out-of-place transforms need a separate output buffer of the same size.
call hipCheck(hipMalloc(dx, source=hx))
call hipCheck(hipMalloc(dy, N))
write(*,"(a)",advance="no") "-- Running test 'rocFFT C2C 1D out-of-place double (z)' &
&(Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Forward transform, out-of-place: dx -> dy.
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_notinplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_loc(dy), c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
! Inverse transform, out-of-place: dy -> dx. rocFFT is unnormalized, so the
! round trip yields N*input. Note that rocFFT may overwrite the input buffer
! of an out-of-place transform, so dy must be treated as clobbered afterwards.
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_notinplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dy), c_loc(dx), c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dx))
call hipCheck(hipFree(dy))
do i = 1, N
error = abs(hx(i) - N * hx_input(i))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), &
" expected=", N * hx_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_c2c_1d_notinplace_z
Normalizing with a scale factor#
Instead of scaling the result with a separate kernel, attach a scale factor to
a plan description. rocFFT multiplies every output element by that factor, so a
factor of 1/N on the inverse plan makes the round trip reproduce the input.
! Demonstrates rocfft_plan_description_set_scale_factor, described in the
! rocFFT "Working with rocFFT" guide: fold the 1/N normalization into the
! transform instead of scaling the result with a separate kernel.
program rocfft_scale_factor_z
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: N = 16
complex(8), allocatable, target, dimension(:) :: hx, hx_input
complex(8), pointer, dimension(:) :: dx => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
type(c_ptr) :: desc = c_null_ptr
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: one = 1
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
allocate(hx(N))
allocate(hx_input(N))
do i = 1, N
hx(i) = cmplx(dble(i), dble(N - i), kind=8)
end do
hx_input(:) = hx(:)
call hipCheck(hipMalloc(dx, source=hx))
write(*,"(a)",advance="no") "-- Running test 'rocFFT scale factor double (z)' &
&(Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Forward transform (in-place), unscaled.
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
! rocFFT transforms are unnormalized. Attaching a scale factor of 1/N to the
! inverse plan folds the normalization into the transform itself, so the
! round trip reproduces the input instead of N*input.
call rocfftCheck(rocfft_plan_description_create(desc))
call rocfftCheck(rocfft_plan_description_set_scale_factor(desc, 1.0d0/dble(N)))
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
desc))
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call rocfftCheck(rocfft_plan_description_destroy(desc))
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dx))
do i = 1, N
error = abs(hx(i) - hx_input(i))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), &
" expected=", hx_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_scale_factor_z
Managing the work buffer#
Large transforms need scratch memory. rocFFT allocates and frees it on every
rocfft_execute call unless the application provides a buffer. Query the
requirement with rocfft_plan_get_work_buffer_size and hand a buffer over
with rocfft_execution_info_set_work_buffer to control its lifetime or to
share one allocation between several plans.
! Fortran counterpart of the work-buffer handling in the rocFFT sample
! clients/samples/rocfft/rocfft_example_complexcomplex.cpp: query the scratch
! memory a plan needs and supply it through an execution info handle.
program rocfft_work_buffer_z
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
! Large transforms are split into several kernels and need scratch memory.
integer(c_size_t), parameter :: N = 262144
complex(8), allocatable, target, dimension(:) :: hx, hx_input
complex(8), pointer, dimension(:) :: dx => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
type(c_ptr) :: info = c_null_ptr
type(c_ptr) :: dwork = c_null_ptr
integer(c_size_t) :: work_fwd = 0
integer(c_size_t) :: work_bwd = 0
integer(c_size_t) :: work_bytes = 0
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: one = 1
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
allocate(hx(N))
allocate(hx_input(N))
do i = 1, N
hx(i) = cmplx(dble(mod(i,7)) - 3.0d0, dble(mod(i,5)) - 2.0d0, kind=8)
end do
hx_input(:) = hx(:)
call hipCheck(hipMalloc(dx, source=hx))
write(*,"(a)",advance="no") "-- Running test 'rocFFT work buffer double (z)' &
&(Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
! Ask each plan how much scratch memory it needs. rocFFT allocates the work
! buffer itself when none is supplied; providing one explicitly lets the
! application control its lifetime and share a single buffer between plans.
call rocfftCheck(rocfft_plan_get_work_buffer_size(plan_fwd, work_fwd))
call rocfftCheck(rocfft_plan_get_work_buffer_size(plan_bwd, work_bwd))
work_bytes = max(work_fwd, work_bwd)
call rocfftCheck(rocfft_execution_info_create(info))
if (work_bytes > 0) then
call hipCheck(hipMalloc(dwork, work_bytes))
call rocfftCheck(rocfft_execution_info_set_work_buffer(info, dwork, work_bytes))
end if
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, info))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, info))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_execution_info_destroy(info))
if (work_bytes > 0) then
call hipCheck(hipFree(dwork))
end if
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dx))
! rocFFT is unnormalized, so the round trip yields N*input.
do i = 1, N
error = abs(hx(i) - N * hx_input(i))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), &
" expected=", N * hx_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_work_buffer_z
Running on HIP streams#
By default rocFFT executes on the null stream. Associate an application-owned
stream with an execution info handle to overlap independent transforms. The
handle must be passed to every rocfft_execute call that should use the
stream, and each stream has to be synchronized before its results are read
back. This example runs two independent transforms on two streams.
! Fortran counterpart of the rocFFT sample
! clients/samples/rocfft/rocfft_example_set_stream.cpp: two independent in-place
! transforms, each queued on its own HIP stream.
program rocfft_stream_z
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: N = 16
integer, parameter :: Nfft = 2
! One bundle of handles per transform, mirroring the fft_fixture_t struct
! used by the rocFFT sample.
type fft_t
complex(8), pointer :: buf(:) => null()
type(c_ptr) :: stream = c_null_ptr
type(c_ptr) :: info = c_null_ptr
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
end type fft_t
type(fft_t) :: ffts(Nfft)
complex(8), allocatable, target, dimension(:,:) :: hx, hx_input
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: one = 1
integer(c_size_t) :: work_size
integer :: i, k
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
allocate(hx(N,Nfft))
allocate(hx_input(N,Nfft))
do k = 1, Nfft
do i = 1, N
hx(i,k) = cmplx(dble(i*k), dble(N - i), kind=8)
end do
end do
hx_input(:,:) = hx(:,:)
write(*,"(a)",advance="no") "-- Running test 'rocFFT stream double (z)' &
&(Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Preparation: one buffer, one stream, one execution info and two plans per
! transform. The stream is carried by the execution info handle, which is
! then passed to every rocfft_execute call that should use that stream.
do k = 1, Nfft
call hipCheck(hipMalloc(ffts(k)%buf, N))
call hipCheck(hipMemcpy(ffts(k)%buf, hx(:,k), hipMemcpyHostToDevice))
call hipCheck(hipStreamCreate(ffts(k)%stream))
call rocfftCheck(rocfft_execution_info_create(ffts(k)%info))
! The stream must be a hipStream_t value, not the address of one.
call rocfftCheck(rocfft_execution_info_set_stream(ffts(k)%info, ffts(k)%stream))
call rocfftCheck(rocfft_plan_create(ffts(k)%plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_plan_create(ffts(k)%plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
! A simple 1D in-place transform needs no extra work buffer.
call rocfftCheck(rocfft_plan_get_work_buffer_size(ffts(k)%plan_fwd, work_size))
if (work_size /= 0) then
write(*,*) "FAILED! unexpected work buffer size ", work_size
STOP 1
end if
end do
! Execution: the calls return as soon as the work is queued on the streams,
! so the two transforms can overlap.
do k = 1, Nfft
call rocfftCheck(rocfft_execute(ffts(k)%plan_fwd, c_loc(ffts(k)%buf), &
c_null_ptr, ffts(k)%info))
call rocfftCheck(rocfft_execute(ffts(k)%plan_bwd, c_loc(ffts(k)%buf), &
c_null_ptr, ffts(k)%info))
end do
! Wait for each stream before reading its results back.
do k = 1, Nfft
call hipCheck(hipStreamSynchronize(ffts(k)%stream))
call hipCheck(hipMemcpy(hx(:,k), ffts(k)%buf, hipMemcpyDeviceToHost))
end do
do k = 1, Nfft
call rocfftCheck(rocfft_plan_destroy(ffts(k)%plan_fwd))
call rocfftCheck(rocfft_plan_destroy(ffts(k)%plan_bwd))
call rocfftCheck(rocfft_execution_info_destroy(ffts(k)%info))
call hipCheck(hipStreamDestroy(ffts(k)%stream))
call hipCheck(hipFree(ffts(k)%buf))
end do
! rocFFT is unnormalized, so each round trip yields N*input.
do k = 1, Nfft
do i = 1, N
error = abs(hx(i,k) - N * hx_input(i,k))
if (error > error_max * N) then
write(*,*) "FAILED! k=", k, " i=", i, " error=", error, &
" got=", hx(i,k), " expected=", N * hx_input(i,k)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_stream_z
Inspecting a plan#
rocfft_plan_get_print writes a summary of a plan to stdout: the precision,
the transform type, the placement, the array types, and the strides, offsets
and distances that rocFFT derived from the plan description. It is the quickest
way to confirm that a layout was described as intended. The call writes from C,
so flush the Fortran output unit first to keep both streams in order.
! Demonstrates rocfft_plan_get_print, the plan introspection helper described in
! the rocFFT "Working with rocFFT" guide. It writes a human-readable summary of
! a plan to stdout, which is the quickest way to check how rocFFT interpreted
! the lengths, strides and distances of a plan description.
program rocfft_plan_print_z
use iso_c_binding
use iso_fortran_env, only: output_unit
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
integer(c_size_t), parameter :: Nx = 4
integer(c_size_t), parameter :: Ny = 6
integer(c_size_t), parameter :: N = Nx*Ny
complex(8), allocatable, target, dimension(:,:) :: hx, hx_input
complex(8), pointer, dimension(:,:) :: dx => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: one = 1
integer(c_size_t), parameter :: two = 2
integer :: i, j
double precision :: error
double precision, parameter :: error_max = 1.0d-8
! lengths starts with the fastest-varying dimension, matching Fortran's
! column-major storage.
allocate(lengths(2))
lengths(1) = Nx
lengths(2) = Ny
allocate(hx(Nx,Ny))
allocate(hx_input(Nx,Ny))
do j = 1, Ny
do i = 1, Nx
hx(i,j) = cmplx(dble(i), dble(j), kind=8)
end do
end do
hx_input(:,:) = hx(:,:)
call hipCheck(hipMalloc(dx, source=hx))
call rocfftCheck(rocfft_setup())
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
two,&
c_loc(lengths),&
one,&
c_null_ptr))
! rocfft_plan_get_print writes to stdout from C, so flush the Fortran unit
! first to keep the two output streams in order.
write(*,"(a)") "-- rocfft_plan_get_print output for a 2D complex plan:"
flush(output_unit)
call rocfftCheck(rocfft_plan_get_print(plan_fwd))
write(*,"(a)",advance="no") "-- Running test 'rocFFT plan print double (z)' &
&(Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
two,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dx))
! rocFFT is unnormalized, so the round trip yields Nx*Ny*input.
do j = 1, Ny
do i = 1, Nx
error = abs(hx(i,j) - N * hx_input(i,j))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " j=", j, " error=", error, &
" got=", hx(i,j), " expected=", N * hx_input(i,j)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED!"
end program rocfft_plan_print_z
Reusing compiled kernels#
rocFFT ships kernels for common problems and compiles the rest when a plan is
created. Those runtime-compiled kernels are cached in memory for the lifetime
of the process. rocfft_cache_serialize copies the cache into a buffer that
rocFFT allocates, which the application can store and hand to
rocfft_cache_deserialize in a later process to avoid compiling the same
kernels again. Release the buffer with rocfft_cache_buffer_free.
! Demonstrates the compiled-kernel cache described in the rocFFT
! "Runtime compilation" guide: rocFFT compiles the kernels a plan needs when the
! plan is created, and the resulting cache can be serialized, moved and loaded
! back into another process.
program rocfft_cache_z
use iso_c_binding
use hipfort
use hipfort_check
use hipfort_rocfft
implicit none
! A length that is not covered by the kernels built into rocFFT, so that plan
! creation has to compile kernels at runtime and the cache is not empty.
integer(c_size_t), parameter :: N = 10007
complex(8), allocatable, target, dimension(:) :: hx, hx_input
complex(8), pointer, dimension(:) :: dx => null()
type(c_ptr) :: plan_fwd = c_null_ptr
type(c_ptr) :: plan_bwd = c_null_ptr
type(c_ptr) :: cache = c_null_ptr
integer(c_size_t), target :: cache_bytes = 0
integer(c_size_t), allocatable, target, dimension(:) :: lengths
integer(c_size_t), parameter :: one = 1
integer :: i
double precision :: error
double precision, parameter :: error_max = 1.0d-8
allocate(lengths(1))
lengths(1) = N
allocate(hx(N))
allocate(hx_input(N))
do i = 1, N
hx(i) = cmplx(dble(mod(i,11)) - 5.0d0, dble(mod(i,7)) - 3.0d0, kind=8)
end do
hx_input(:) = hx(:)
call hipCheck(hipMalloc(dx, source=hx))
write(*,"(a)",advance="no") "-- Running test 'rocFFT kernel cache double (z)' &
&(Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! Creating the plan populates the in-memory kernel cache.
call rocfftCheck(rocfft_plan_create(plan_fwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_forward,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
! Copy the cache into a buffer that rocFFT allocates. The buffer address is
! written to the first argument, its size in bytes to the second, which is
! why the length is passed as the address of a size_t variable.
call rocfftCheck(rocfft_cache_serialize(cache, c_loc(cache_bytes)))
if (.not. c_associated(cache) .or. cache_bytes == 0) then
write(*,*) "FAILED! empty kernel cache"
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
! The buffer can be written to a file and loaded in another process. Loading
! it back replaces matching kernels and leaves the rest of the cache alone,
! so plans created afterwards reuse the kernels instead of recompiling them.
call rocfftCheck(rocfft_cache_deserialize(cache, cache_bytes))
! The buffer belongs to the caller once it has been handed out.
call rocfftCheck(rocfft_cache_buffer_free(cache))
cache = c_null_ptr
call rocfftCheck(rocfft_plan_create(plan_bwd,&
rocfft_placement_inplace,&
rocfft_transform_type_complex_inverse,&
rocfft_precision_double,&
one,&
c_loc(lengths),&
one,&
c_null_ptr))
call rocfftCheck(rocfft_execute(plan_fwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_execute(plan_bwd, c_loc(dx), c_null_ptr, c_null_ptr))
call hipCheck(hipDeviceSynchronize())
call rocfftCheck(rocfft_plan_destroy(plan_fwd))
call rocfftCheck(rocfft_plan_destroy(plan_bwd))
call hipCheck(hipMemcpy(hx, dx, hipMemcpyDeviceToHost))
call hipCheck(hipFree(dx))
! rocFFT is unnormalized, so the round trip yields N*input.
do i = 1, N
error = abs(hx(i) - N * hx_input(i))
if (error > error_max * N) then
write(*,*) "FAILED! i=", i, " error=", error, " got=", hx(i), &
" expected=", N * hx_input(i)
call rocfftCheck(rocfft_cleanup())
STOP 1
end if
end do
deallocate(hx)
deallocate(hx_input)
deallocate(lengths)
call rocfftCheck(rocfft_cleanup())
write(*,*) "PASSED! kernel cache size: ", cache_bytes, " bytes"
end program rocfft_cache_z
Setting the ROCFFT_RTC_CACHE_PATH environment variable to a writable file
achieves the same result without any application code: rocFFT then persists
compiled kernels there by itself.
Querying the rocFFT version#
rocfft_get_version_string fills a C string buffer of at least 30
characters. Pass the address of the first element of a character(kind=c_char)
array and the buffer capacity, then copy the result up to the terminating NUL.
! Demonstrates rocfft_get_version_string, the rocFFT library version query.
program rocfft_version
use iso_c_binding
use hipfort_check
use hipfort_rocfft
implicit none
! rocfft_get_version_string requires a buffer of at least 30 characters.
integer(c_size_t), parameter :: buflen = 64
character(kind=c_char), target :: buf(buflen)
character(kind=c_char, len=buflen) :: version
integer :: i
write(*,"(a)",advance="no") "-- Running test 'rocFFT version string' &
&(Fortran 2008 interfaces) - "
call rocfftCheck(rocfft_setup())
! The C API writes a NUL-terminated string into the buffer, so pass the
! address of the first element and the buffer capacity.
buf(:) = c_null_char
call rocfftCheck(rocfft_get_version_string(c_loc(buf(1)), buflen))
call rocfftCheck(rocfft_cleanup())
! Copy the C string into a Fortran character variable, stopping at the NUL.
version = c_char_''
do i = 1, buflen
if (buf(i) == c_null_char) exit
version(i:i) = buf(i)
end do
if (len_trim(version) == 0) then
write(*,*) "FAILED! rocfft_get_version_string returned an empty string"
STOP 1
end if
write(*,*) "PASSED! rocFFT version: ", trim(version)
end program rocfft_version