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指令并行。
1 void saxpy_parallel(int n, 2 float a, 3 float*x, 4 float*restrict y) 5 { 6 #pragma acc kernels 7 for(int i =0; i < n;++i) 8 y[i]= a*x[i]+ y[i]; 9 }
下面是用Fortran写的
1 subroutine saxpy(x,y,n,a) 2 real:: a, x( : ), y( : ) 3 integer:: n, i 4 !$ acc kernels 5 do i = 1, n 6 y(i) = a*x(i)+y(i) 7 enddo 8 !$ acc end kernels 9 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倍的加速!