GPU Programming

Directives: parallel, kernels, loop gang vector, data, enter data/exit data, update host/device, routine seq, reduction, collapse, async.

31.1 OpenACC

Directives: parallel, kernels, loop gang vector, data, enter data/exit data, update host/device, routine seq, reduction, collapse, async.

!$acc data copyin(x, a) copy(y)
!$acc parallel loop
do i = 1, n
  y(i) = a * x(i) + y(i)
end do
!$acc end data

nvfortran -acc -Minfo=accel saxpy.f90      # or gfortran -fopenacc (offload-enabled build)

31.2 OpenMP offload

Supported by ifx (Intel GPUs), nvfortran, AMD flang/Cray, LLVM flang, gfortran (offload builds).

!$omp target teams distribute parallel do map(to: x) map(tofrom: y)
do i = 1, n
  y(i) = a * x(i) + y(i)
end do

31.3 Standard parallelism on GPUs

nvfortran -stdpar=gpu offloads do concurrent loops and some array intrinsics (matmul, transpose) automatically, using unified memory.

31.4 CUDA Fortran (NVIDIA)

Libraries: cuBLAS, cuFFT, cuSOLVER via Fortran interfaces; hipfort for AMD GPUs.

module kernels
  use cudafor
contains
  attributes(global) subroutine saxpy(x, y, a, n)
    real :: x(*), y(*)
    real, value :: a
    integer, value :: n
    integer :: i
    i = blockDim%x * (blockIdx%x - 1) + threadIdx%x
    if (i <= n) y(i) = a * x(i) + y(i)
  end subroutine saxpy
end module kernels

program main
  use cudafor; use kernels
  integer, parameter :: n = 1000000
  real, device :: x_d(n), y_d(n)           ! device arrays
  real :: y(n)
  x_d = 1.0; y_d = 2.0
  call saxpy<<<(n + 255) / 256, 256>>>(x_d, y_d, 2.0, n)
  y = y_d                                  ! copy back to host
  print *, maxval(y)
end program main