[NVIDIA编程教程]OpenACC: Directives for GPUs
NVIDIA已经在过去五年里大力发展CUDA技术，我们估计CUDA开发人员超过15万，很多重要的科学应用正在CUDA的帮助下完成。但是我们仍然有一个很长的路要走，以帮助每个人从GPU计算中享受到好处。有很多开发者没有太多时间来学习和应用的并行编程语言。还有很多科学家和工程师，正在自己的代码上辛勤工作,他们只能对自己的代码做一些改动,以便适应于整个硬件和操作系统的便携式，并且能在多个平台上的计算。
这一类的开发，需要一个更高级别的GPU加速的方法。他们需要的东西就是简单，功能强大，移植方便，并开放。 这就是OpenACC推出的原因，它是一个开放的标准，是一个编译器指令集合，在标准C，C+ +和Fortran语言中指定代码循环和代码区域从主机CPU上卸载到一个加速器上，同时提供跨作业系统、主机CPU和加速器的可移植性。通过使用指令，开发人员可以让相同的代码运行在多核CPU、GPU或任何编译器支持的其他类型的并行硬件上。如果你是一个C或C+ +程序员，你可能熟悉的＃pragma指令。
简单：一个OpenACC例子
这里是一个非常简单的使用OpenACC的例子。此循环执行SAXPY的操作。 SAXPY表示单精度A乘以X加Y。A是一个标值（scalar value ）,而X和Y是向量，所以这是一个vector scale 和增加操作。
下面是一个C编写的简单的SAXPY，并用OpenACC指令并行。
void saxpy_parallel(int n,
float a,
float*x,
float*restrict y)
{
#pragma acc kernels
for(int i =0; i < n;++i)
y[i]= a*x[i]+ y[i];
}
下面是用Fortran写的，
subroutine saxpy(x,y,n,a)
real:: a, x( : ), y( : )
integer:: n, i
!$ acc kernels
do i = 1, n
y(i) = a*x(i)+y(i)
enddo
!$ acc end kernels
end subroutine saxpy
在C语言中的#pragma 行和Fortran语言中的 !$acc 行就是编译器指令:给编译器提示! 在这个程序里,我们只是建议这是个并行循环,编译器会试图生成一个并行内核代码给一个加速器(比如,一个GPU).同时注意我们不需要再对GPU做任何事情.
跟CUDA不同的是:
- 我们不需要在设备上分配或者初试化阵列,
- 我们也不需要将Host(CPU)的数据复制到加速器(GPU)上或者在循环后将加速器的结果返回到Host.
- 我们不需要写一个CUDA内核去执行并行循环主体;
- 我们也不需要明确在GPU上launch内核.
OpenACC编译器都会在幕后做好这些工作!
便捷性
关于编译器指令（也许令人惊讶的）伟大的事情之一是他们可以被忽略不计。这意味着，你可以在不支持OpenACC的平台上编译代码, 而它会工作，就像你从来没有新增指令一样。这也意味着，一开始,亲的的代码就适应于CPU和GPU,不需要特别指定.
当然我也不打算骗你。有时要用OpenACC获得良好的加速，需要改变原代码。例如，一些数据的布局导致访问模式不是并行的。往往明智的做法是将一个结构数组（AOS）表示成阵列结构（SOA）。这是不指定GPU：SOA通常对任何并行处理器，包括CPU都可以用。这种类型的代码变化有助于暴露并行，因此，它往往具有普遍的性能优势。结合这些好处，加上编译器指令的可移植性,你离“性能可移植性”可望而不可及的目标不远了!
开放性：OpenACC的起源
OpenACC波特兰集团（PGI），Cray公司，CAPS和NVIDIA开发。 PGI，Cray，CAPS花了2年多的时间推出商业用的编译器以加速GPU,但他们的编译器之间有一些不同,所以就成立一个组织标准化指令加速方法。
OpenACC规范1.0版即将正式发布, 今天你可以开始使用上面列出的供应商的编译器。在我的下一篇文章中，我将使用的PGI编译器。今天，PGI的编译器实现了OpenACC加速 。如果您有兴趣尝试OpenACC，你可以下载一个免费试用PGI加速编译器尝试一下。
在我的下一篇文章[编程教程]用OpenACC指令将程序速度提升2倍，我将深入一个更有趣的代码示例演示如何使用OpenACC指令,只需几行代码给你一个3 - 4倍的加速!
更多关于OpenACC的资料,可以上 bbs.gpuworld.cn上