CUDA Fortran：统一内存

_梦魇花葬 2014-08-20 05:22:19

加精

前面我们讲过：统一内存是一个CUDA6.0的特性,我们说了很多关于并行原则。CUDA 6.0介绍了统一内存，

引用

https://developer.nvidia.com/cuda-toolkit

，这大大简化了GPU编程给程序员一个指针数据从GPU和CPU访问。但这种增强的内存模型不是万能的，它才可以在CUDA C / c++使用。而新的PGI14.7支持统一内存在CUDA Fortran编译器里。
在CUDA Fortran Managed Memory 这篇PGInsider文章中，Brent Leback，PGI应用程序和服务经理布，说过“使用内存管理简化了许多编码任务,使源代码清洁,并支持复杂的数据结构在主机和设备中进行传递，”PGI 14.7添加了关键字“管理”,你可以使用主机代码端使用，类似于设备端的关键字。下面Brent Leback在他的文章中提到的一个例子，显示数组的分配管理。

! matrix data

real, managed, allocatable, dimension(:,:) :: A, B, C

real, allocatable, dimension(:,:) :: gold

完整的实例代码可以查看http://www.pgroup.com/lit/articles/insider/v6n1a2.htm。
Brent Leback解释道：
我们不再需要两个数组的副本,B和C的数据移动仍然会发生,但不是很明确。现在控制统一内存管理的系统在幕后，就像操作系统管理虚拟内存的方法。如果你的CUDA Fortran程序被两个变量的副本，比如likevariable_name和variable_name_d设备，那么管理内存可能会给你很好的答案。
也就是说就像CUDA C++一样，统一内存会简化CPU和GPU之间复杂的数据传输。
CUDA Fortran管理数据的一个亮点在派生类型。在之前，利用设备端数据进行派生是不合适的，因为设备端数据本身是没有属性的。所以一些数据的派生是在主机端进行的，这样在编程和数据管理上变的很繁琐。
现在加入PGI14.7后，可以很好的在管理数据，从而很好的进行派生。也就是说那些类型模块、定义和实例化实体，在主机端和设备端可以直接使用，不需要来回做回参数传递。
附上一些代码
before CUDA managed data：

! matrix data

  real, allocatable, dimension(:,:) :: A, B, C, gold

  real, allocatable, device, dimension(:,:) :: dA, dB, dC

  . . .



  allocate(A(N,N))

  allocate(B(N,N))

  allocate(C(N,N))

  allocate(gold(N,N))



  call random_number(A)

  call random_number(B)



  allocate(dA(N,N))

  allocate(dB(N,N))

  allocate(dC(N,N))



  dA = A

  dB = B

  dC = 0.0



  alpha = 1

  beta = 0

  m = N

  k = N

  blocks = dim3(N/256, N/16, 1)

  threads = dim3(16, 16, 1)



  call sgemm_cpu(A, B, gold, m, N, k, alpha, beta)



  ! timing experiment

  time = 0.0

  istat = cudaEventRecord(start, 0)

  do j = 1, NREPS

    call sgemmNN_16x16<<>>(dA, dB, dC, m, N, k, alpha, beta)

  end do

  istat = cudaEventRecord(stop, 0)

  istat = cudaDeviceSynchronize()

  istat = cudaEventElapsedTime(time, start, stop)

  time = time / (NREPS*1.0e3)



  C = dC

using managed data:

! matrix data

  real, managed, allocatable, dimension(:,:) :: A, B, C

  real, allocatable, dimension(:,:) :: gold

  . . .



  allocate(A(N,N))

  allocate(B(N,N))

  allocate(C(N,N))

  allocate(gold(N,N))



  call random_number(A)

  call random_number(B)

  C = 0.0



  alpha = 1

  beta = 0

  m = N

  k = N

  blocks = dim3(N/256, N/16, 1)

  threads = dim3(16, 16, 1)



  call sgemm_cpu(A, B, gold, m, N, k, alpha, beta)



  ! timing experiment

  time = 0.0

  istat = cudaEventRecord(start, 0)

  do j = 1, NREPS

    call sgemmNN_16x16<<>>(A, B, C, m, N, k, alpha, beta)

  end do

  istat = cudaEventRecord(stop, 0)

  istat = cudaDeviceSynchronize()

  istat = cudaEventElapsedTime(time, start, stop)

  time = time / (NREPS*1.0e3)