OpenACC并行编程实战.html.pdf

《OpenACC并行编程实战.html.pdf》由会员分享，可在线阅读，更多相关《OpenACC并行编程实战.html.pdf（87页珍藏版）》请在三一文库上搜索。

1、序 随着计算机科学技术的飞速发展，作为计算机的核心处理器的体系结构也经历了从单核、多核到众核的革命性跨越。如今，“异构+众核”已成为超级计算机主流的体系结构，并将引领未来E级计 算的进一步发展。 然而，这种较为“前卫”的结构设计，也给编程者带来了很大的挑战：他们不得不面对更为复杂的底层结构和更多的存储层次，以往的算法设计和程序代码也不得不随之调整。在这种情况下，诸如 CUDA、OpenCL、OpenACC等新的异构并行编程语言应运而生。其中，OpenACC是一种极具发展前景的编程模型，其具有使用方便、代码改动小、平台适用范围广的特性，将为新时代的编程者们带来极大 的帮助。 OpenACC虽然有

2、着易学易用的特点，但是要想全面掌握其丰富的语法特性和使用技巧，以编写高效的程序，还是需要一本有权威性、实用性的技术书籍来指导。本书作者何沧平博士是华为高级系统设 计工程师，也是OpenACC QQ群的群主，具有多年的程序开发经验，一直致力于OpenACC的发展与推广，在并行计算应用领域有着很深的造诣和独到的见解。 对于想要掌握OpenACC使用技巧和编程精髓的读者来说，本书是一本非常具有参考价值的学习教程。与传统的技术书籍相比，本书更加注重内容的可读性和易用性，逻辑清晰，内容全面准确，且更加 注重编程实践，有大量C/C+/Fortran的完整代码实例，便于读者学习和实践。作为第一本中文的Ope

3、nACC技术书籍，可谓为国内的编程学习者带来了福音。 特别值得一提的是，本书首次引入了OpenACC在“神威太湖之光”超级计算机上应用情况的章节。神威太湖之光是2016年全球TOP500排名第一的超级计算机，配备了完全由国人自主研发的异构 众核处理器。其超强的计算速度成为了高性能计算应用的强力助推。结合应用的特点和处理器独特的结构设计，系统对OpenACC进行了扩展。目前，该系统完成了气候气象、航空航天、船舶工程、药物设 计等十多个领域的大型应用课题，其中三个高性能计算应用入围“戈登贝尔奖”。这是我国近30年来首次入围该奖项。在这些应用的开发过程中，OpenACC起到了关键作用。 漆锋滨 国家

4、并行计算机工程技术研究中心 前言 2010年以来，中国超级计算机建设突飞猛进，欣欣向荣。一个原因是国力强盛，大力投资高新科技；另一个原因是整体科技水平提高，需求旺盛。天气预报、石油物探、工程仿真、基因测序等传统应 用对计算资源的需求持续增长，以深度学习为代表的人工智能大爆发，资金雄厚的互联网公司对计算能力极度渴求。超级计算机的建设、应用主战场正在从教育科研单位转向科技企业。 为什么要写这本书 面对浩如烟海的数据，CPU已经力不从心，因此世界领先的超级计算机都装备大量的加速器或者众核处理器。 目前主流加速器产品是NVIDIA GPU、AMD GPU和Intel至强Phi协处理器。三种加速器使用的

5、编程语言分别为CUDA C/CUDA Fortran、OpenCL和MIC导语。加速器计算有4个困难。 一是CUDA/OpenCL等低级语言编程难度大，且需要深入了解加速器的硬件结构。而大部分的用户不是专业编程人员，学习一门新的编程技术将耗费大量时间。 二是加速器的计算模型与CPU差别很大，移植旧程序需要几乎完全重写。大量的旧程序在性能优化上已经千锤百炼，稳定性上也久经考验，完全重写是不可完成的任务。 三是低级编程语言开发的程序与硬件结构密切相关，硬件升级时必须升级软件，否则将损失性能。而硬件每隔两三年就升级一次，频繁的软件升级将给用户带来巨大负担。 四是投资方向难以选择。三种加速器均有自己独

6、特的编程语言，且互不兼容。用户在投资建设硬件平台、选择软件开发语言时就会陷入困境，不知三种设备中哪个会在竞争中胜出。如果所选加速器将 来落败，将会带来巨大损失；而犹豫不决又将错过技术变革的历史机遇。 OpenACC应运而生，可以克服这4个困难。OpenACC的编程机制是，程序员只在原程序中添加少量编译标识，编译器根据作者的意图自动产生低级语言代码。无须学习新的编程语言和加速器硬件知 识，便能迅速掌握。只添加少量编译标识，不破坏原代码，开发速度快，既可并行执行又可恢复串行执行。在硬件更新时，重新编译一次代码即可，不必手工修改代码。OpenACC标准制定时就考虑了目前 及将来的多种加速器产品，同一

7、份代码可以在多种加速器设备上编译、运行，无成本切换硬件平台。掌握OpenACC后，编写程序省时、省力、省心。 本书特色 笔者学习超算技术时有过苦泪：教材一上来就讲技术细节，只能机械地学习，不清楚这些算法、语法要解决什么问题，花费巨大精力后却发现解决不了自己的难题；新技术的资料往往是英文的，而且 零碎，还不一定准确，收集、学习成本很高，求助无门；科学计算领域的祖师语言Fortran现在有点小众，新技术资料更少。为节省读者时间，本书特别重视易读性、易用性，具体有如下特点。 第一本中文OpenACC技术书籍，方便国内读者系统阅读，快速掌握。 全面、准确：本书覆盖OpenACC的所有特性，重要特性结合

8、例子详细讲解，不常用的特性列出标准定义，方便查找。对特性、功能的描述均来自OpenACC规范，并用实例验证，确保准确。 大量的例子：书中的示例代码超过160个，不是代码片断，而是完整代码，拿来即用。全部由笔者编写、测试，运行无报错。 兼顾C/C+和Fortran：以C/C+为主来讲解，但所有的例子都有Fortran版的代码，Fortran专用特性会详细讲解。物理、气象等领域积累的优秀Fortran代码可以借此移植到新硬件上，重焕生机。 阅读轻松：按照并行程序的编写、调优顺序，先交代各阶段面临的问题，然后自然引入OpenACC提供的解决办法，最后实测验证效果。一切顺理成章，读者不会有云里雾里的感

9、觉。 读者对象 科学家：迅速改造旧程序，快速编写新的原型程序验证算法、理论，将更多的时间投入到高价值的科研创新中去。 企业程序员：一套程序适配多种运行平台，维护简单；性能敏感的部分用CUDA等低层语言编写，性能不敏感部分用OpenACC编写，相互配合，兼得程序性能和开发效率。 本科生、研究生：花最少的时间掌握一门计算工具，省出时间学习更多的知识。 如何阅读本书 第1章介绍超级计算技术的发展趋势和并行编程概况，可以从中了解OpenACC的作用。没有CUDA C基础的读者能够掌握基本概念，便于深入理解OpenACC的并行化技术。第2章介绍OpenACC语言的设 计思路。第34章是本书的核心，将计算

10、部分并行化，并将数据传递时间减到最少。至此，读者已经能够编写性能良好的OpenACC程序。第57章介绍高级并行技术，以进行极致性能优化，以及与CUDA C/CUDA Fortran和各类库的混合编程。第8章给出OpenACC规定的所有运行时例程，不用细读，用到时再参考。第9章指导部署开发环境，以便快速上手。 勘误和支持 本人水平有限，错误与疏漏在所难免，恳请批评指正。联系笔者请发送电子邮件至。更多技术资源请访问。技术交流QQ群284876008（将满）、564520462，欢迎 加入。 致谢 感谢国家并行计算机工程技术研究中心漆锋滨老师撰写第10章，并为本书作序。感谢PGI美女工程师王珍、帅气

11、工程师仰琎歆、Daniel Tian、资深专家Mathew Colgrove的技术支持。感谢机械工业出版社 的策划编辑高婧雅尽心协调与支持。 何沧平 第1章 并行编程概览 对绝大多数人而言，编程语言只是一个工具，讲究简单高效。科学家的主要精力应该用在科研创新活动上，编程工作仅仅是用来验证创新的理论，编程水平再高也不可能获得诺贝尔奖。对学生和企业 程序员而言，技术无穷尽，永远学不完，不用即忘，应该认清技术发展方向，学习有前途的技术，不浪费青春年华。 OpenACC语言专为超级计算机设计，因此读者需要了解超级计算机的技术演进方向，特别是主流加速器的体系架构、编程模型，看清OpenACC的应用场景，

12、有的放矢。普通读者虽然不会用到大型机 群，但小型机群甚至单台服务器、普通显卡的计算模式都是相同的。 最近几年的著名超级计算机（见附录A）均采用加速器作为主要计算部件，可预见未来几年的上层应用仍将围绕加速器展开。 1.1 加速器产品 超级计算机的加速器历史上有多种，本节只介绍当前流行的两种：英伟达GPU和英特尔融核处理器。加速器的物理形态是PCIe板卡，样子大致如图1.1所示，图1.2是拆掉外壳后的样子，正中央的是 GPU芯片，芯片周围的小黑块是显存颗粒，金黄色的边缘处是与PCIe连接的金手指，通过PCIe插槽与CPU相连。图1.3中的机架式服务器左下部装有4块GPU卡，图1.4是服务器的主板俯

13、视图，箭头处就是4 个PCIe插槽。 图1.1 英伟达Tesla GPU1 图1.2 英伟达GPU内部构造 图1.3 GPU服务器2 图1.4 GPU服务器的主板 从逻辑关系上来看，目前市场在售的英特尔服务器CPU中已经集成了内存控制器和PCIe通道。2颗CPU通过1个或2个QPI接口连接，CPU通过内置的内存控制器连接DDR3或DDR4内存条，GPU加速器 通过PCIe总线与CPU相连（图1.5）。2016年6月主流的英特尔至强E5-2600 v3和E5-2600 v4 CPU每颗拥有40个PCIe Lan，而每块GPU的接口需要16个PCIe Lan，因此每颗CPU上最多全速挂载2块 GP

14、U。有些服务器的PCIe16插槽上只安排8，甚至4的信息速率，可以挂载更多的GPU。 图1.5 加速器在服务器上的逻辑位置 英特尔融核处理器在服务器上的位置与GPU一样，不再赘述。 1 图片来自。 2 图片来自。 1.2 并行编程语言 在并行计算发展史上出现过多种并行编程语言，至今仍在使用的只剩几种，它们各有特色。 （1）Pthreads 20世纪70年代，贝尔实验室发明了UNIX，并于20世纪80年代向美国各大高校分发V7版的源码以做研究。加利福尼亚大学伯克利分校在V7的基础上开发了BSD UNIX。后来很多商业厂家意识到UNIX 的价值也纷纷以贝尔实验室的System V或BSD为基础来开

15、发自己的UNIX，较著名的有Sun OS、AIX、VMS。随着操作系统的增多，应用程序的适配性工作越来越繁重。为了提高UNIX环境下应用程序的 可迁移性，电气和电子工程师协会（Institute of Electrical and Electronics Engineers，IEEE）设计了POSIX标准。然而，POSIX并不局限于UNIX，许多其他的操作系统也支持POSIX标准。POSIX.1已经 被国际标准化组织所接受，POSIX已发展成为一个非常庞大的标准族，一直处在发展之中。 POSIX线程（POSIX threads，Pthreads），是线程的POSIX标准。该标准定义了创建和操纵

16、线程的一整套接口。在类UNIX操作系统（UNIX、Linux、Mac OS X等）中，都使用Pthreads作为操作系 统的线程。Pthreads用来开发与操作系统紧密相关的应用程序，管理粒度很细，例如线程的创建与销毁、线程锁、线程属性、线程优先级、线程间通信等琐碎操作均需要程序员安排。对科学与工程类计算 程序来说，程序员的精力应集中在业务模型和代码算法上，不应浪费在底层代码细节上。 虽然Pthreads不适合编写高性能计算程序，但它多线程并发的设计理念启发了其他并行语言。 （2）OpenMP OpenMP是由一些大型IT厂商和一些学术机构组成的非盈利组织，官网是www.openmp.org。

17、永久成员包括AMD、CAPS-Entreprise、Convey Computer、Cray、Fujitsu、HP、IBM、Intel、 NEC、NVIDIA、Oracle Corporation、Red Hat、ST Microelectronics、Texas Instruments；正式成员是对OpenMP标准感兴趣，但不生产销售相关产品的组织，例如ANL、ASC/LLNL、BSC、 cOMPunity、EPCC、LANL、NASA、ORNL、RWTH Aachen University、SNL-Sandia National Lab、Texas Advanced Computing C

18、enter、University of Houston。 计算热点都是在循环上，OpenMP的并行化思路是将循环的迭代步分摊到多个线程上，每个线程只承担一部分计算任务，循环运行的墙上时间（从开始到结束的自然流逝时间）自然也就减少了。分割 方法是在循环上面添加一些预处理标记（图1.18），编译器识别到这些标记以后，将关联的循环翻译成并行代码，然后再与剩余的串行代码合并起来编译、链接成可执行程序。 图1.18 OpenMP的并行化模式 使用OpenMP编译并行程序时，程序员需要先保证串行代码正确，找出热点循环，然后在循环上添加OpenMP预处理标记。打开编译器的openmp选项可以编译成并行版本，

19、否则编译器将忽略预处理 器标记，仍然编译成串行版本。既不破坏原有代码，开发速度又快，省时省力。 （3）CUDA CUDA（Compute Unified Device Architecture）是英伟达公司设计的GPU并行编程语言，一经推出就引发了GPU通用计算研究热潮。CUDA是闭源的，只能运行在英伟达的产品上。CUDA C/C+是对C/C+语言的扩展，添加了一些数据类型、库函数，并定义一种新的函数调用形式。CUDA起初支持C和少量C+特性，后来逐渐提高对C+的支持度。从CUDA 3.0开始与PGI合作支持 Fortran。CUDA语言还可以细分为CUDA C/C+、CUDA Fortran

20、，本书成稿时的CUDA C/C+最新版本是7.5，8.0版本即将发布。CUDA Fortran没有版本号，只是随着PGI编译器的升级而增加新特 性。1.3节会介绍CUDA C/C+的编程模型和一些示例代码，CUDA Fortran的详细情况可以参考官网https:/ Practices for Efficient CUDA Fortran Programming和中文翻译版CUDA-Fortran高效编程实践，此处不展开介绍。 （4）OpenCL OpenCL（Open Computing Language，开放计算语言）是一个面向异构系统并行编程的免费标准，支持多种多样的设备，包括但不限于C

21、PU、GPU、数字信号处理器（DSP）。OpenCL的优势是 一套代码多处运行，只要为新的设备重新编译代码就可以运行，移植方便。 OpenCL由苹果公司首先提出，随后Khronos Group成立相关工作组，以苹果草案为基础，联合业界各大企业共同完成了标准制定工作，工作组的成员来自各行各业，且都是各自领域的领导者，成员 名单请参见官网www.khronos.org/opencl/。 （5）OpenACC 这里不多说，后面会全面、详细讲解。 1.3 CUDA C 本节简要介绍CUDA C编程的相关概念，使读者能够看懂OpenACC编译过程中出现的CUDA内置变量，理解并行线程的组织方式。如果读者

22、已有CUDA编程经验，请跳过。 CPU用得好好的，为什么要费心费力地改写程序去到GPU上运行呢？只有一个理由：跑得更快。小幅的性能提升吸引力不够，必须有大幅提升才值得采购新设备、学习新工具、设计新算法。从图1.19 可以看出，在双精度浮点峰值和内存带宽这两个关键指标上，GPU的性能都达到同时期主力型号CPU的57倍。如果利用得当，可以预期获得57的性能提升。以前只在CPU上运行，计算方法的数学理论 和程序代码实现已经迭代发展多年，花很大力气才能提速10%20%，提速50%已经很厉害了。简单粗暴地更换硬件设备就能立刻提速几倍，全世界的科学家、工程师一拥而上，GPU加速的应用遍地开 花。注意，评价

23、GPU应用性能的时候，至少要和2颗中高端CPU相对，并且两种代码都优化到最好。任何超过硬件潜能的加速结果都是有问题的。 图1.19 同时期主力CPU与GPU的性能对比1 那么问题来了。GPU的芯片面积与CPU差不多，价格也接近，为什么性能这么强悍呢？图1.20是CPU和GPU芯片的组成示意图，左边是一个单核超标量CPU，4个算术逻辑单元（ALU）承担着全部计 算任务，却只占用一小部分芯片面积。“控制”是指分支预测、乱序执行等功能，占用芯片面积大而且很费电。服务器CPU通常有三级缓存，占用的芯片面积最大，有的型号甚至高达70%。ALU、控制、 缓存都在CPU内部，大量内存条插在主板上，与CPU通

24、过排线相连。GPU中绝大部分芯片面积都是计算核心（4行紧挨着的小方块，每行12个），带阴影的水平小块是控制单元，控制单元上面的水平条是 缓存。 图1.20 CPU（左）和GPU（右）的芯片面积占用情况 通用CPU对追踪链表这样拥有复杂逻辑控制的程序运行得很好，但大规模的科学与工程计算程序的流程控制都比较简单，CPU的长处难以施展。为了解释GPU如何获得极高的性能，需要先了解一下 CPU中的控制、缓存、多线程的作用。 ALU承担最终的计算工作，越多越好。“控制”的目标是预取到正确的指令和数据以保证流水线不中断，挖掘指令流里的并行度，让尽量多的部件都在忙碌工作，从而提高性能。缓存的作用是为了填 补

25、CPU频率与内存条频率的差距、减小CPU与内存条之间数据延时。目前中高端CPU的频率在2.03.2GHz，而内存条的频率还处于1600MHz、1866MHz、2133MHz，内存条供应、承接数据的速度赶不 上CPU处理数据的速度。由于ALU到主板上内存条的路径较长，延时高，而如果需要的数据已经在缓存中，那么就能有效降低延时，提高数据处理速度。缓存没有命中怎么办？只能到内存条上取，延时 高。为了进一步降低延时，英特尔CPU有超线程功能，开启后，一个CPU物理核心就变成了两个逻辑核心，两个逻辑核心分时间片轮流占用物理核心资源。当然了，按时间片切换是有代价的：换出时要保 留正在运行的程序的现场，换入

26、时再恢复现场以便接着上次继续运行。在缓存命中率比较低的情况下，超线程功能能够提高性能。 GPU天生是为并行计算设计的：处理图像的大量像素，像素之间相互独立，可以同时计算，而且没有复杂的流程跳转控制。正如图1.19所示，GPU的大部分芯片面积都是计算核心，缓存和控制单元很 小，那么它是怎么解决分支预测、乱序执行、数据供应速度、存取数据延时这些问题的呢？ GPU的设计目标是大批量的简单计算，没有复杂的跳转，因此直接取消分支预测、乱序执行等高级功能。更进一步，多个计算核心（例如32个）共用一个控制单元再次削减控制单元占用的芯片面积。 这样做的效果就是：发射一条指令，例如加法，32个计算核心步调一致地

27、做加法，只是每个计算核心操作不同的数据。如果只让第1个计算核心做加法，那么在第1个计算核心做加法运算的时候，剩余的计 算核心空闲等待。这种情形下资源浪费，性能低下，要尽量避免。让大量计算核心空转的应用程序不适合GPU，用CPU计算性能更好。 计算核心与显存之间的频率差异如何填补？特别简单，降低计算核心的频率。考虑到芯片功耗与频率的平方近似成正比，降低频率不但能解决数据供应速度问题，而且能降低GPU的功耗，一举两得。 从表1.1可以看出GPU产品的频率在562875MHz，远低于主力CPU的2.0GHz3.2GHz。 最重要是延时，GPU的缓存那么小，怎么解决访问显存的巨大延时呢？答案是多线程，

28、每个计算核心分摊10个以上的线程。执行每条指令之前都要从就绪队列中挑选出一组线程，每组线程每次只执行 一条指令，执行完毕立即到后面排队。如果恰巧碰上了延时较多的访存操作，那么该线程进入等待队列，访存操作完成后再转入就绪队列。只要线程足够多，计算核心总是在忙碌，隐藏了访存延时。有人 立刻会问，这么频繁地切换线程、保存现场、恢复现场也需消耗不少时间吧，会不会得不偿失呢？实际上GPU线程切换瞬间完成，这是因为每个线程都有一份独占资源（例如寄存器），不需要保存、恢复 现场，线程切换只是计算核心使用权的转移。 1 图片来自。 第2章 OpenACC概览 2007年出现的CUDA C/C+语言引爆了GPU

29、通用计算热潮，但编程比较麻烦，挖掘硬件性能需要很多高超的优化技巧。为了降低编程门槛，2011年11月，Cray、PGI、CAPS和英伟达4家公司联合推出 OpenACC 1.0编程标准，2012年3月PGI率先推出支持OpenACC的编译器PGI Accelerator with OpenACC。PGI公司创立于1989年，是一家在高性能计算领域很有名望的编译器和工具供应商，属于意法半导体旗 下的全资子公司。2013年6月，OpenACC 2.0标准发布。2013年7月，英伟达收购了PGI公司，但PGI原有品牌和体系得以保留并继续正常运营，OpenACC、CUDA Fortran、CUDA x

30、86、GPGPU等相关技术的 开发工作也将继续。OpenACC 2.0版本功能已经相当完备，直到2015年11月才推出OpenACC 2.5版本。2013年11月，GCC加入OpenACC组织，2016年5月推出的GCC 6.1支持OpenACC 2.0a标准。 OpenACC组织的成员均为知名企业、高校、科研机构，名单见表2.1，OpenACC标准的更新、活动、组织成员请访问其官网http:/www.openacc.org/。 表2.1 OpenACC组织成员名单 虽然GCC也支持OpenACC标准，但对最新标准的支持比较慢，支持的特性也比较少，最大的优势是免费。详情参见其官网https:/

31、gcc.gnu.org/wiki/OpenACC。PGI编译器对OpenACC的支持最快、最完 善，实际上OpenACC标准的制定人员和PGI员工有很大的重叠。该编译器每月更新1次，每个账户每半年有1个月的试用时间，下载地址为http:/ 对企业用户收费较高。目前OpenACC用户中尝鲜者较多，学生用户和科学家购买的动力不足，不利于推广。因此，英伟达于2015年7月发布了全新OpenACC套件，向学术开发者和研究人员免费提供 OpenACC编译器（其实还是PGI编译器包装的），同时向商业用户提供90天免费试用版，详情见官网https:/ 2.1 OpenACC规范的内容 如前所述，OpenAC

32、C并行化的方式不是重写程序，而是在串行C/C+或Fortran代码上添加一些编译标记。支持OpenACC的编译器能够看懂这些标记，并根据标记含义将代码编译成并行程序。对英 伟达GPU来说，编译器将C/C+/Fortran代码翻译成CUDA C/C+代码，然后编译链接成并行程序。对AMD GPU来说，中间代码是OpenCL。在OpenACC语境中，CPU称为主机（host），GPU等加速 器称为设备（device）。这些术语的使用场景没有严格规定，能准确表达含义即可，本书可能会将名词CPU、主机、GPU、加速器、设备混用。 程序并行化主要包含三方面的工作（表2.2）：计算并行化、数据管理、运行时

33、库和环境变量。 并行化的唯一目的是充分利用硬件资源来提高程序运行速度，缩短运行时间。程序中最耗时间的是循环，计算并行化的目标是将循环迭代步分散到多个不同的线程上执行，这些线程运行在多个加速器 核心上，从而将计算任务由CPU转移到加速器上，减轻CPU的负担。循环并行化需要解决的问题有这些：指定将哪个循环并行化，以什么样的方式组织并行线程。OpenACC使用计算构件kernels或parallel 来完成这个工作，看起来是这个样子（此处不必深究语法，后有详述）： #pragma acc kernels for(i=0; i 2 #ifdef _OPENACC 3 #include 4 #endif

34、 5 int main() 6 7 #ifdef _OPENACC 8 printf(“Number of device: %dn“, 9 acc_get_num_devices(acc_device_not_host); 10 #else 11 printf(“OpenACC is not supported.n“); 12 #endif 13 return 0; 14 例3.1第9行中的函数acc_get_num_devices（.）的作用是获取设备的数量，函数功能细节此处不必深究，它的原型包含在头文件openacc.h中，不支持OpenACC的编译器没有该头文件，因此在第2行 使用#if

35、def条件包含这个文件，否则不支持OpenACC的编译器遇到#include语句时将报错。若编译器支持OpenACC则编译第89行，若编译器不支持OpenACC则编译第11行。 使用不支持OpenACC的gcc编译器编译： $ gcc comp1c.c -o comp1c.exe $ ./comp1c.exe OpenACC is not supported. 使用支持OpenACC的PGI编译器编译，并打开支持选项，在拥有一块英伟达显卡的笔记本上运行，正确给出运行环境中的设备数量： $ pgcc -acc comp1c.c -o comp1c.exe $ ./comp1c.exe Numbe

36、r of device: 1 PGI编译器的C语言编译程序是pgcc，选项-acc的作用是打开编译器对OpenACC的支持，从而有了编译器定义的宏_OPENACC，没有选项-acc的话，第3、89行会被忽略。 Fortran代码有个预编译小技巧：Fortran语言本身不支持#ifdef等预处理指令，因此要在正式编译前预处理一下，见例3.2。将Fortran源码文件的扩展名由通常的小写.f90改为大写的.F90，PGI编译器 就会自动添加预处理过程。 【例3.2】comp1f.F90：用宏变量_OPENACC条件编译。 1 program main 2 #ifdef _OPENACC 3 use

37、 openacc 4 #endif 5 implicit none 6 #ifdef _OPENACC 7 print*, “Number of device: “, 9 10 #define Mx 8191 11 #define Ny 1023 12 float uval(float x, float y)return(x*x+y*y); 13 int main() 14 15 float Widthy = 2.0, Heightx = 1.0; 16 float *u0, *u1; 17 float hx, hy, hx2, hy2, fij, c1, c2; 18 float uerr,

38、 errtol, tmp, *tprt; 19 int ix, jy, maxIter, iter; 20 int Nyp1 = Ny+1; 21 timestruct t1, t2; 22 long long telapsed; 23 24 u0 = (float*)malloc(sizeof(float)*(Mx+1)*(Ny+1); 25 u1 = (float*)malloc(sizeof(float)*(Mx+1)*(Ny+1); 26 maxIter = 100; 27 errtol = 0.00f; 28 hx = Heightx/Mx; 29 hy = Widthy/Ny; 3

39、0 31 / Initialize the left/right boundary of u0/u1 32 for(ix = 0; ix uerr ? tmp : uerr; 67 68 printf(“iter = %d uerr = %en“, iter, uerr); 69 if(uerr #define gettime(a) _ftime(a) #define usec(t1,t2) (t2).time-(t1).time)*1000+ (t2).millitm-(t1).millitm)*100) typedef struct _timeb timestruct; #else #in

40、clude #define gettime(a) gettimeofday(a,NULL) #define usec(t1,t2) (t2).tv_sec-(t1).tv_sec)*1000000+ (t2).tv_usec-(t1).tv_usec) typedef struct timeval timestruct; #endif 例3.40第1011行用宏定义网格尺寸Mx和Ny，第16行定义两个一维数组u0和u1，用来保存二维格点上的旧值和新值。为兼顾数组的内存对齐，Mx和Ny均选择较特殊的值8192-1和1024-1。代码中的 变量Widthy和Heightx分别对应（0，W）（0，H

41、）中的W和H。第3242行给数组u0和u1的左右上下四个边界赋值，这些边界值在整个计算过程中保持不变。数组u0的内部元素对应h的内部格 点，第4448行为数组u0的内部元素赋初始值，这些初始值可以随意指定，代码里指定为0。代码中的变量fij、c1、c2、hx2、hy2分别对应公式中的。第5771行是主体迭 代，每一个iter迭代步中，用u0更新u1同时计算两者间的最大误差值uerr，然后将u1中的新值赋给u0，准备进行下一步迭代。 例3.40第26行的maxIter指定最大迭代次数。第27行的errtol指定误差阈值，当两步迭代之间的最大误差小于该阈值时，退出循环。为方便后文比较，将errto

42、l指定为0，从而每个版本的代码均迭代相同 的次数，实际使用的时候可以为errtol指定大于0的阈值。在每次迭代后，第68行输出新值和旧值的最大误差。第74行输出循环主体消耗的时间。 为节约篇幅，代码中略去了错误的检测和处理。 Fortran版代码例3.41与C版代码例3.40基本相同，区别在于：数组u0和u1动态分配；每两步迭代输出1次误差。为保持与C版本一致，数组起始下标指定为0，而不采用默认的1。 编译运行两种语言版本均能得到相同的100步误差值uerr，绘制成对数坐标图。由图3.14可见，误差单调下降，但下降速度逐渐平稳。 图3.14 Jacobi迭代误衰减情况 后续章节将逐个讲解Ope

43、nACC的导语、子语，同时逐步移植优化Jacobi迭代。计算误差的那几行代码的并行化需要较高级的reduction子语，比较麻烦，因此暂时注释掉，稍后会恢复误差代码。 Jacobi迭代的并行化代码有多个版本，每个版本的完整代码都较长，为节省篇幅，只列出有变化的代码，后面的例子都照此办理。例3.42和例3.43是不计算误差的版本，用来与带误差的版本对比，评估误 差消耗的时间。 【例3.42】ja0c2.c：Jacobi串行迭代（不计算误差，在ja0c1.c基础上修改）。 57 for(iter = 1; iter uerr ? tmp : uerr; 67 68 /printf(“iter =

44、%d uerr = %en“, iter, uerr); 69 /if(uerr 中的一个。 binop、binop=、+和-不能是重载运算符。 表达式x binop expr必须在数学上等价于x binop（expr）。如果expr中操作符的优先级高于binop，或者在expr或expr子表达式的两边使用圆括号，就能够满足这个要求。 表达式expr binop x必须在数学上等价于（expr）binop x。如果expr中操作符的优先级等于或高于binop，或者在expr或expr子表达式的两边使用圆括号，就能够满足这个要求。 在那些x多次出现的形式中，x的求值次数没有指定。 Fortran

45、中，atomic构件的语法是： !$acc atomic read capture表达式 !$acc end atomic 或 !$acc atomic write write表达式 !$acc end atomic 或 !$acc atomic update update表达式 !$acc end atomic 或 !$acc atomic capture update表达式 capture表达式 !$acc end atomic 或 !$acc atomic capture capture表达式 update表达式 !$acc end atomic 或 !$acc atomic captu

46、re capture表达式 write表达式 !$acc end atomic 这里的write表达式具有如下形式（如果子语是write或capture）： x = expr 这里的capture表达式具有如下的形式（如果子语是write或capture）： v = x 这里的update表达式具有如下的形式（如果子语是update、capture，或没有子语）： x = x 操作符 expr x = expr 操作符 x x = 内置过程名字(x, 表达式列表) x = 内置过程名字(表达式列表, x) 在前面的表达式中： x和v（如果用到了）都是内置类型的标量变量。 x决不能是动态分配的变

47、量。 在一个原子区域执行期间，在语法上多次出现的x必须指向相同的存储位置。 v和expr（如果用到了）都不能访问x指向的存储位置。 x和expr（如果用到了）都不能访问v指向的存储位置。 expr是一个标量类型的表达式。 表达式列表是一个用逗号分隔的非空标量表达式列表。如果内置过程名字是iand、ior或ieor，那么必须有且只能有一个表达式出现在表达式列表里。 内置过程名字是max、min、iand、ior或ieor中的一个。操作符是+、*、-、/、.and.、.or.、.eqv.或.neqv.中的一个。 表达式x binop expr必须在数学上等价于x binop（expr）。如果exp

48、r中操作符的优先级高于binop，或者在expr或expr子表达式的两边使用圆括号，就能够满足这个要求。 表达式expr binop x必须在数学上等价于（expr）binop x。如果expr中操作符的优先级等于或高于binop，或者在expr或expr子表达式的两边使用圆括号，就能够满足这个要求。 内置过程名字不能是非内置程序。 操作符必须是内存操作符，不能是用户自定义的运算符。所有的赋值必须是内置赋值。 在那些x多次出现的形式中，x的求值次数没有指定。 一个带有read子语的atomic构件对x指向的位置做一次强制原子读。一个带有write子语的atomic构件对x指向的位置做一次强制原子写。 一个带有update子语的原子构件，使用指定的操作符或固有操作符强制原子更新x指向的位置。注意，没有子语出现时，atomic导语等价于atomic update。对x指向的位置，仅有read和write能够相 互原子地实施。对于需要读或写的x指向的位置，表达式或表达式列表不必原子地求值。 一个带有capture子语的原子构件，使用指定的操作符或固有操作符强制原子更新x指向的位置，与atomi