OpenACC编程标准介绍-公开版-小小河-20130109

OpenACC 编程标准介绍 小小河 2013年1月9日 2007 年以来，以 nVidia GPU 为代表的加速器并行计算风起云涌，带有加速器的超级计 算机在 TOP500 中的份额逐年增加，支持加速器的主流应用软件也呈爆炸式增长，研究加速 器计算的技术人员数以百万计，世界范围内的大学、研究机构竞相开设相关课程。 目前有三家厂商提供加速器产品：nVidia GPU、AMD GPU、Intel 至强 Phi 协处理器。 三种加速器使用的编程语言分别为 CUDA C/CUDA Fortran、OpenCL 和 MIC 导语。加速器 计算有四个技术困难：一是 CUDA/OpenCL 等低级语言编程难度大，且需要深入了解加速 器的硬件结构。而大部分的用户不是职业编程人员，学习一门新的编程技术将耗费大量时间、 人力、财力。二是加速器的计算模型与 CPU 差别很大，移植旧程序需要几乎完全重写。大 量的旧程序在性能优化上已经千锤百炼，稳定性上也久经考验，完全重写是不可完成的任务。 三是低级编程语言开发的程序与硬件结构密切相关，硬件升级时必须升级软件，否则将损失 性能。而硬件每隔两三年就升级一次，频繁的软件升级将给用户带来巨大负担。四是投资方 向难以选择。三种加速器均有自己独特的编程语言，且互不兼容。用户在投资建设硬件平台、 选择软件开发语言时就会陷入困境，不知三种设备中哪个会在将来胜出。如果所选加速器将 来落败，将会带来巨大损失；而犹豫不决又将错过技术变革的历史机遇。 为了克服这些困难，PGI、Cray、CAPS 和 nVidia 四家公司联合创立 OpenACC 应用编 程接口标准，并于 2011 年 11 月的 SC11 大会上发布 1.0 版。OpenACC 1.0 标准大部分内容 来自 PGI Accelerator 编程模型，并借用了 CUDA C 的一些术语，标准文档行文上与 OpenMP 相似。 OpenACC 可以克服这四个困难。OpenACC 应用编程接口的机制是，编译器根据作者的 意图自动产生低级语言代码。程序员只需要在原程序中添加少量导语，无须学习新的编程语 言和加速器硬件知识，可以迅速掌握。只在原始程序上添加少量导语，不破坏原代码，开发 速度快，既可并行执行又可恢复串行执行。在硬件更新时，重新编译一次代码即可，不必手 工修改代码。OpenACC 标准制定时就考虑了目前及将来的多种加速器产品，同一份代码可 以在多种加速器设备上编译、运行，零代价切换硬件平台。  不破坏原始代码，保证修改的可逆性；  开发工作量小；  将计算密集的代码段卸载到加速器上并行计算；  在主机与加速器间提供多种形式的数据传递；  支持多种加速器设备； 两种情形下，不能破坏原始代码。移植旧程序时，旧程序的性能和可靠性已经被时间证 明，修改程序原始代码将带来安全隐患。有些程序不但需要在加速器上运行，还需要在没有 加速器的服务器上运行。OpenACC 借鉴 OpenMP 的导语（directive）模式，在原始代码上 OpenACC设计目标 为什么使用 OpenACC 添加一些导语，告诉编译器将哪些代码段卸载到加速器上执行、如何在主机与加速器间移动 数据。如果编译器不支持 OpenACC 标准或支持选项没有打开，编译器将忽略所有的 OpenACC 导语，编译出的程序只在 CPU 上运行。因此，用 OpenACC 移植程序不必对原始 代码作任何修改。 当前的绝大多数算法都是以串行方式描述，程序设计结构也是串行的，使用 CUDA/OpenCL 等低级语言移植旧程序或编写新程序时，需要重新设计算法，将串行算法转 换为并行算法，这就要求程序员的对算法有深入了解，这对职业编程人员是一个挑战，需要 花费大量时间学习本专业以外的算法；使用低级语言需要了解加速硬件结构，程序优化更需 要对线程调度机制有深刻理解，这对非计算机专业的广大用户非常困难，学习成本高昴。而 OpenACC 只要求用户指明哪些代码需要放在加速器上并行计算，生成低级语言代码、根据 设备优化代码等工作完全交由编译器来完成。导语代码非常简洁，工作量小，开发速度快、 周期短。 程序计算密集部分往往是求解线性方程组、求解特征值、最小二乘、矩阵向量乘、矩阵 矩阵乘等数值代数运算，它们通常占用绝大部分的程序运行时间。这些运算的特点是使用的 数据小、运算量大、并行度高。例如矩阵矩阵乘使用的数据量为 2n，计算量为 n2，并行度 为 n，从而计算密度为 n/2。科学和 工程 路基工程安全技术交底工程项目施工成本控制工程量增项单年度零星工程技术标正投影法基本原理 计算中的 n 通常会超过一千万，这就味着只需要将 少量数据复制到加速器上（几秒钟），就能做长时间的运算（几天或几周），保证不频繁进行 低效的数据传递。高并行度保证程序能充分利用加速器动辄上千的计算核心，提高设备利用 率。OpenACC 着力并行化计算密集的循环代码段，一旦将循环并行化，程序性能将有很大 提高。例如一个可并行部分占 95%计算时间的程序，使用 1 万个线程并行（在 GPU 上很常 见）执行，整体计算时间可以降至(1-95% + (1-95%)/10000) = 5.005%，加速 19.98 倍。并行 比重越高，加速效果越好。随着模型规模的扩大，并行比重会越来越高。OpenACC 提供两 个计算构件（parallel,kernel），并能详细指定并行方式(gang, worker, vector 的数值)，通过编 译器将循环转换为高效的低级语言代码（CUDA/OpenCL）。 主机 CPU 与加速器间的数据传输是并行化的瓶颈。加速器通过主板上的 PCIe 2.0 x16 或 PCIe 3.0 x16 插槽与主机相连。与加速器内部超过 300GB/s 的带宽相比，主机与加速器的 带宽仅为 8GB/s 和 16GB/s；主机与加速器间的访问延时也远高于加速器内部的访存延时， 所以，频繁、大量的数据传递将会严重影响并行部分的性能。OpenACC 设计了多种数据子 语，尽量重用数据，让数据尽可能长时间地留在加速器内存中，避免无谓的数据传递开销。 例如，当所需数据已经存在于设备上时，就不再从主机复制过来；当加速器不需要主机上的 数组元素值时，只在设备内存中开辟数据空间，不从主机复制数据；当主机不需要加速器上 数组元素值时，并行区域结束时直接释放加速器数组，跳过通常的复制回主机的步骤。 加速器设备目前有三种， 将来还可能有更多的厂商提供加速器产品，例如 ARM 移动 处理器。OpenACC 在设计时就考虑到了兼容所有的加速器设备，它的两层并行模式照顾到 了各家厂商的产品，术语、语法也全新的，没有采用任何一种产品的现有术语，但它的概念 在目前的 OpenCL/CUDA 语言中都有对应概念。熟悉任何一种加速器并行语言的用户都能 很方便地切换到 OpenACC. OpenACC 的运行时函数库中提供一组例程，用户可以在运行时 选择、切换加速器类型，也可以在多个同类型设备中选择、切换到某个加速器。 目前OpenACC标准文件版本为1.0，2011年11月发布。支持的编程语言为C/C++、Fortran. OpenACC 1.0 标准大部分来自 PGI Accelerator 编程模型，两个标准只在语法上有区别。PGI Accelerator 编译器于 2008 年首次发布。 OpenACC 1.0 标准支持 nVidia Fermi 架构 GPU，不支持 nVidia Kepler 架构 GPU. OpenACC 2.0 草案增加了对 Kepler 架构的支持，但草案还没有成为正式的标准文件。 OpenACC 1.0 标准文件和 OpenACC 2.0 草案可从 OpenACC 组织官网 http://www.openacc-standard.org/ 获取，OpenACC 标准文件的中文 翻译 阿房宫赋翻译下载德汉翻译pdf阿房宫赋翻译下载阿房宫赋翻译下载翻译理论.doc 版可从 http://www.bytes.me/openacc/ 下载。 目前支持 OpenACC 1.0 标准的编译器有三家：PGI，CAPS，Cray. PGI 编译器 12.6 及以 后版本支持OpenACC，操作系统全面支持Linux, Windows和OS x； 硬件支持x86和x64 CPU; 加速器目前只支持 nVidia GPU, 2013 年上半年将支持 2012 年 11 月发布的 Intel Xeon Phi 众 核处理器；编程语言支持 C/C++和 Fortran. CAPS 编译器支持 Linux 和 Windows 操作系统； 硬件支持 x86 和 x64 CPU；加速器支持 nVidia GPU 和 AMD 的 GPU; 编程语言支持 C/C++ 和 Fortran. Cray 编译器据说只能在 Cray 服务器上使用。 支持 OpenACC 的 PGI 编译器可从 http://www.pgroup.com/resources/accel.htm 下载试用 版。CAPS 编译器主推 CAPS 自已的 OpenHMPP 加速器编程标准，试用版只支持 OpenHMPP， 只有正式版支持 OpenACC. CAPS 中文官网为 www.caps-entreprise.com.cn 中国上海超级计算中心、美国橡树岭国家实验室（拥有 2012 年 11 月 TOP500 排名第一 的 Titan 超级计算机）、桑迪亚国家实验室等知名单位已经使用 OpenACC 标准编写程序。 OpenACC 中的关键术语有 directive、clause、construct. 在 C/C++中语法是： #pragma acc directive-name [clause [[,] clause]…] new-line 在 Fortran 语法为： !$acc directive-name [clause [[,] clause]…] 建议将 directive 翻译为导语，将 clause 翻译为子语，将 construct 翻译为构件。directive 的作 用是给编译器一些指导，指出哪些代码需要并行化、需要怎么并行化，编译器根据程序员的 指导信息生成最佳的并行化代码。有人将 directive 翻译为编译制导语句、编译指导语句、指 令语、指令等，意思都对，但编译制导语句、编译指导语句太长，使用不便，指令语、指令 中的指令一词太普通，易混，且没有指导的含义，不太恰当。导语一词长度、意思都比较合 适。clause 是导语的修饰部分，更详细地表明导语的意图。有人将 clause 翻译为子句，而子 句一词含有小句子的意思，实际上 OpenACC 中的 clause 都只有一个词，很短，不能称为一 个句子。子语一词既说明了它与导语的关系，又有一个相同的语字，读起来顺口。construct 翻译为构件，取自建筑名词，实际功能也很相像。 OpenACC 术语 示例程序 OpenACC 发展现状 Fortran 语言版 ! pi_f.f90 1 program pical 2 implicit none 3 integer, parameter :: N = 10000 4 integer i 5 real t, pi 6 pi = 0.0 7 !$acc parallel loop 8 do i= 0, N-1 9 t = (i+0.5)/N 10 pi = pi + 4.0/(1.0 + t*t) 11 end do 12 !$acc end parallel loop 13 print *, 'pi=', pi/N 14 endprogram 15 使用 PGI 编译器 12.10 编译、运行该代码，反馈信息为 PGI$ pgf90 -acc -Minfo pi_f.f90 pical: 8, Accelerator kernel generated 8, CC 1.0 : 14 registers; 32 shared, 80 constant, 0 local memory bytes CC 2.0 : 16 registers; 0 shared, 48 constant, 0 local memory bytes 9, !$acc loop gang, vector(256) ! blockidx%x threadidx%x 11, Sum reduction generated for pi 8, Generating compute capability 1.0 binary Generating compute capability 2.0 binary PGI$ ./pi_f.exe launch kernel file=C:\cygwin\home\he\test\pi_f.f90 function=pical line=8 device=0 grid=40 block=256 queue=0 launch kernel file=C:\cygwin\home\he\test\pi_f.f90 function=pical line=11 device=0 grid=1 block=256 queue=0 pi= 3.141593 Accelerator Kernel Timing data C:\cygwin\home\he\test\pi_f.f90 pical 8: region entered 1 time time(us): total=82,000 init=77,000 region=5,000 kernels=31 w/o init: total=5,000 max=5,000 min=5,000 avg=5,000 8: kernel launched 1 times grid: [40] block: [256] time(us): total=23 max=23 min=23 avg=23 11: kernel launched 1 times grid: [1] block: [256] time(us): total=8 max=8 min=8 avg=8 C 语言版 /* pi_c.c */ 1 #include 2 #define N 10000 3 int main(void) 4 { 5 double pi = 0.0, tmp; 6 int i; 7 #pragma acc parallel loop 8 for (i=0; i