《OpenACC并行编程实战》—— 第2章 OpenACC概览 2.1 OpenACC规范的内容

本节书摘来自华章出版社《OpenACC并行编程实战》一 书中的第1章，第1.2节，作者何沧平，更多章节内容可以访问云栖社区“华章计算机”公众号查看。

第2章 OpenACC概览

2007年出现的CUDA C/C++语言引爆了GPU通用计算热潮，但编程比较麻烦，挖掘硬件性能需要很多高超的优化技巧。为了降低编程门槛，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 x86、GPGPU等相关技术的开发工作也将继续。OpenACC 2.0版本功能已经相当完备，直到2015年11月才推出OpenACC 2.5版本。2013年11月，GCC加入OpenACC组织，2016年5月推出的GCC 6.1支持OpenACC 2.0a标准。
OpenACC组织的成员均为知名企业、高校、科研机构，名单见表2.1。

<img src="https://yqfile.alicdn.com/9ef8b1591505a96e264d095abbdecd9356c052b6.png" width="" height="">
<img src="https://yqfile.alicdn.com/eb7c9bafbf338d4dbd8f465d19912b7b96087729.png" width="" height="">

2.1　OpenACC规范的内容

如前所述，OpenACC并行化的方式不是重写程序，而是在串行C/C++或Fortran代码上添加一些编译标记。支持OpenACC的编译器能够看懂这些标记，并根据标记含义将代码编译成并行程序。对英伟达GPU来说，编译器将C/C++/Fortran代码翻译成CUDA C/C++代码，然后编译链接成并行程序。对AMD GPU来说，中间代码是OpenCL。在OpenACC语境中，CPU称为主机（host），GPU等加速器称为设备（device）。这些术语的使用场景没有严格规定，能准确表达含义即可，本书可能会将名词CPU、主机、GPU、加速器、设备混用。
程序并行化主要包含三方面的工作（表2.2）：计算并行化、数据管理、运行时库和环境变量。
并行化的唯一目的是充分利用硬件资源来提高程序运行速度，缩短运行时间。程序中最耗时间的是循环，计算并行化的目标是将循环迭代步分散到多个不同的线程上执行，这些线程运行在多个加速器核心上，从而将计算任务由CPU转移到加速器上，减轻CPU的负担。循环并行化需要解决的问题有这些：指定将哪个循环并行化，以什么样的方式组织并行线程。OpenACC使用计算构件kernels或parallel来完成这个工作，看起来是这个样子（此处不必深究语法，后有详述）：

#pragma acc kernels 
for(i=0; i <N; i++)
{
  代码语句
}

数据管理占用OpenACC规范的大量篇幅，语法也多。数据管理解决的问题是：如何在主机内存与设备内存之间传递数据，如何开辟、释放设备内存，如何管理变量的生存期和作用域。
OpenACC运行时库包含几十个函数，这些函数的功能只有在程序运行时才能实现，在编译阶段不能实现。例如，从几个设备中选取当前设备，初始化设备，分配、释放设备内存，在主机内存与设备内存之间复制数据，等待某个操作的完成。OpenACC规范中规定了几个环境变量，用来指定设备类型和设备编号，详见6.6.1节。

<img src="https://yqfile.alicdn.com/237cf9854c9843030378dd03f129fe75b151b447.png" width="" height="">

2.1.1　抽象加速器模型

市面上的加速器产品多种多样，架构设计也有很大差别。为了能兼容尽可能多的加速器，OpenACC定义了一个抽象的加速器模型，以涵盖市场上主流加速器的特点，然后在抽象模型上建立计算执行模型。在抽象模型中（图2.1），主机可以直接访问主机内存，设备可以直接访问设备内存，主机能够分配、释放设备内存，主机能够启动设备上的函数。但是主机不能直接访问设备内存，设备也不能直接访问主机内存。设备内存中的数据需要在设备运算开始之前从主机内存复制到设备内存，设备运算完成后再将结果复制回主机内存。

<img src="https://yqfile.alicdn.com/463d4789aa3a6dae4a7f14a36a418caa3f5512cf.png" width="" height="">

2.1.2　存储模型

本节描述的存储模型非常概括，初次阅读不强求完全理解，等读完第4章后再读本节就会豁然开朗。
一个仅在主机上运行的程序与一个在主机+加速器上运行的程序，它们最大的区别在于加速器上的内存可能与主机内存完全分离。例如目前大多数GPU就是这样。这种情况下，设备内存可能无法被主机线程直接读写，这是因为它没有被映射到主机线程的虚拟存储空间。主机内存与设备内存间的所有数据移动必须由主机线程完成，主机线程通过系统调用在相互分离的内存之间显式地移动数据。数据移动通常采用直接内存访问（Direct Memory Access，DMA）技术。不能假定加速器能读写主机内存，虽然有些加速器设备支持这样的操作，但常常有严重的性能损失。
在CUDA C和OpenCL等低层级加速器编程语言中，主机和加速器存储器分离的概念非常明确，内存间移动数据的语句甚至占据大部分用户代码。在OpenACC模型中，内存间的数据移动可以是隐式的，编译器根据程序员的导语管理这些数据移动。然而程序员必须了解背后这些相互分离的内存，理由包括但不限于以下几方面。
有效加速一个区域的代码需要较高的计算密度，而计算密度的高低取决于主机内存与设备内存的存储带宽；计算密度可以用计算量除以数据量来衡量，这个商值越大，计算密度越高。
与主机内存相比，设备内存空间有限，因此操作大量数据的代码不能卸载到设备上。在高性能集群典型配置下，主机内存为128GB或256GB，而GPU上的设备内存最大24GB，差一个数量级。
主机上的指针里保存的主机地址可能仅在主机上可用；设备上的指针里保存的地址可能仅在设备上可用。建议不要在主机内存与设备内存之间显式地传递指针的值。主机指针在设备上解引用或设备指针在主机上解引用很可能出错。
OpenACC通过设备数据环境来暴露相互分离的内存。设备数据有一个显式生存期，从分配空间直到被删除。如果设备与本地线程共享物理内存或虚拟内存，那么本地线程也能共享设备数据环境。这种情况下，编译器不必为设备创建新的数据副本，也不需要移动数据。如果设备内存与本地线程的内存物理地或虚拟地分离，那么编译器将在设备内存中创建新的数据副本并将数据从本地内存复制到设备环境中。
一些加速器（例如目前的GPU）使用一个较弱的存储模型。这种模型不支持不同线程上操作的内存一致性，甚至，在同一个执行单元上，只有在存储操作语句之间显式地内存栏栅才能保证内存一致性。否则，如果一个线程更新一个内存地址而另一个线程读取同一个地址，或者两个操作向同一个位置存入数据，那么硬件可能不保证每次运行都能得到相同的结果。尽管编译器可以检测到一些这样的潜在错误，但仍有可能编写出一个产生不一致数值结果的加速器parallel区域或kernels区域。
目前，一些加速器有一块软件管理的缓存，一些加速器有多块硬件管理的缓存。大多数加速器具有仅在特定情形下使用的硬件缓存，并且仅限于存放只读数据。在CUDA C和OpenCL等低级语言的编程模型中，这些缓存交由程序员管理。在OpenACC模型中，编译器会根据程序员的指示来管理这些缓存。

2.1.3　计算执行模型

OpenACC编译器的执行模型是主机指导加速器设备（如GPU）的运行。主机执行用户应用的大部分代码，并将计算密集型区域卸载到加速器上执行，这些计算密集区域通常是循环。设备上用计算构件parallel或kernels将这些循环并行化。两个计算构件的行为稍有差别，后文会详述。即使在加速器负责的区域，主机也必须精心安排程序的运行：在加速器设备上分配存储空间、初始化数据传输、将代码发送到加速器上、给计算区域传递参数、为设备端代码排队、等待完成、将结果传回主机、释放存储空间。大多数时候，主机可以将设备上的所有操作排成一队，一个接一个地顺序执行。然后在设备上执行parallel区域和kernels区域，parallel区域通常包含一个或多个工作分摊（work-sharing）循环，kernels区域通常包含一个或多个被作为设备上内核执行的循环。
目前大多数的加速器支持二到三层并行。大部分加速器支持粗粒度并行：在执行单元层次并行执行。加速器可能有限支持粗粒度并行操作之间的同步。许多加速器也支持细粒度并行：在单个执行单元上执行多个线程，这些线程可以快速地切换，从而可以忍受长时间的存储操作延时。大多数加速器也支持每个执行单元内的单指令多数据（Single Instruction Multiple Data, SIMD）操作或向量操作。该执行模型表明设备有多个层次的并行，因此程序员需要理解它们之间的区别。例如，一个完全并行的循环和一个可向量化但要求语句间同步的循环之间的区别。一个完全并行的循环可以用粗粒度并行执行。有依赖关系的循环要么适当分割以允许粗粒度并行，要么在单个执行单元上以细粒度并行、向量并行或串行执行。
与三个并行层次相对应，OpenACC设计了gang、worker和vector，见图2.2。gang并行是粗粒度并行，加速器上将启动许多个gang。每一个gang都将有一个或多个worker，一个worker内的SIMD操作或向量操作是vector并行。gang、worker、vector都是一维的，没有二维或三维形式。
图2.2中，worker里的每一个小方块代表一个vector通道（vector lane），图中的vector长度为4，实际程序中可能会是其他值；若干worker组成一个gang，一个计算构件可以同时启动多个gang。用OpenACC的术语来说，gang对应于英伟达GPU的流式多处理器（SM或SMX），vector通道对应于GPU核心，worker没有明确的对应硬件。用CUDA C/C++术语来说，gang对应block，worker和vector与线程的对应关系不确定，会根据block的一维、二维、三维组织情况而变化，第3章会有例子详
解。

<img src="https://yqfile.alicdn.com/af5c0cfb102b2763d7ede0e77399af43d5b4f80e.png" width="" height="">