CUDA Fortran

NVIDIA推出的运算平台。CUDA是一种由NVIDIA推出的通用并行计算架构，该架构使GPU能够解决复杂的计算问题。2007年以来，以NVIDIA GPU为代表的加速器并行计算开始兴起。目前流行的GPU通用编程语言是CUDA C和OpenCL，它们均是C/C++语言的扩展，因此可以方便地将C/C++代码移植到GPU上。但对于科学与工程计算中的重要编程语言Fortran，无法直接地改写为CUDA C或OpenCL。

为使Fortran应用能够使用GPU加速，The Portland Group设计了CUDA Fortran语言。气象、理论物理等领域的应用代码经过简单的改造，就能够利用GPU的强大计算能力。到目前为止，只有PGI Fortran编译器支持CUDA Fortran架构。

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

program testSaxpy
  use mathOps
  use cudafor
  implicit none
  integer, parameter :: N = 40000
  real :: x(N), y(N), a
  real, device :: x_d(N), y_d(N)
  type(dim3) :: grid, tBlock

  tBlock = dim3(256,1,1)
  grid = dim3(ceiling(real(N)/tBlock%x),1,1)

  x = 1.0; y = 2.0; a = 2.0
  x_d = x
  y_d = y
  call saxpy<<<grid, tBlock>>>(x_d, y_d, a)
  y = y_d
  write(*,*) 'Max error: ', maxval(abs(y-4.0))
end program testSaxpy

如上述代码示例，testSaxpy为主机函数，调用saxpy为GPU侧的函数，saxpy的函数运行将卸载到GPU上进行。real :: x(N), y(N)声明的x和y变量为主机变量，real, device :: x_d(N), y_d(N)声明的x_d和y_d为GPU侧变量。

CUDA Fortran开源通过直接将主机变量和GPU侧变量间进行赋值，进行数据传输。比如：x_d = x就将主机变量x数据赋值给了GPU侧变量x_d。