第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, OpenACC标准的更新、活动、组织成员请访问其官网http://www.openacc.org/。
表2.1 OpenACC组织成员名单
虽然GCC也支持OpenACC标准，但对最新标准的支持比较慢，支持的特性也比较少，最大的优势是免费。详情参见其官网https://gcc.gnu.org/wiki/OpenACC。PGI编译器对OpenACC的支持最快、最完善，实际上OpenACC标准的制定人员和PGI员工有很大的重叠。该编译器每月更新1次，每个账户每半年有1个月的试用时间，下载地址为http://www.pgroup.com/resources/accel.htm。对教育用户收费较低，对企业用户收费较高。目前OpenACC用户中尝鲜者较多，学生用户和科学家购买的动力不足，不利于推广。因此，英伟达于2015年7月发布了全新OpenACC套件，向学术开发者和研究人员免费提供OpenACC编译器（其实还是PGI编译器包装的），同时向商业用户提供90天免费试用版，详情见官网https://developer.nvidia.com/openacc-toolkit。
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节。
表2.2 OpenACC规范的主要组成部分
OpenACC将串行程序并行化的手段是添加一些C/C++预处理语句或形式特殊的Fortran注释，预处理语句和特殊注释分为directive和clause两类。
❑ directive表示主要功能，每句有且只能有一个，作用是给编译器一些指导，指出哪些代码需要并行化、需要怎么并行化，编译器根据程序员的指导信息生成最佳的并行代码。
❑ clause表示对directive的修饰，每句可以有零个或多个。
一个directive和若干个clause就构成一个功能模块construct。
为方便阅读、交流，本书将directive翻译为导语，将clause翻译为子语，将construct翻译为构件。网络上有人将directive翻译为编译制导语句、编译指导语句、指令语、指令等，意思都近似，但编译制导语句、编译指导语句太长，使用不便，指令语、指令中的指令一词太普通，易混，且只有强制的含义，没有指导的含义，不太准确。导语一词长度、意思都比较合适。clause是导语的修饰部分，更详细地表明导语的意图。有人将clause翻译为子句，而子句一词含有小句子的意思，实际上OpenACC中的clause都只有一个词，很短，不能称为一个句子。子语一词既说明了它与导语的关系，又有一个相同的语字，读起来顺口。construct翻译为构件，取自建筑名词，实际功能也很相像。
科学和工程计算领域的大量历史遗留程序绝大部分用C/C++和Fortran语言开发，这三种语言也非常适合计算密集型的高性能计算应用，因此OpenACC目前支持C/C++和Fortran。本书中示例代码均给出C和Fortran两种版本，讲述以C版本代码为主，Fortran版只列出代码，特殊情况下才详细讲解。对科学与工程计算而言，大部分的密集计算任务用C语言就可以完成，C++中的复杂类操作只占用少量运行时间，因此只有在绝对必要的时候才用C++。对Fortran而言，Fortran 77固定格式很少用于开发新程序，Fortran 90/95/2003/2008自由格式应用广泛，所以OpenACC对Fortran 90及以后标准支持得更好一些，本书的示例也采用Fortran自由格式。
2.1.1 抽象加速器模型
市面上的加速器产品多种多样，架构设计也有很大差别。为了能兼容尽可能多的加速器，OpenACC定义了一个抽象的加速器模型，以涵盖市场上主流加速器的特点，然后在抽象模型上建立计算执行模型。在抽象模型中（图2.1），主机可以直接访问主机内存，设备可以直接访问设备内存，主机能够分配、释放设备内存，主机能够启动设备上的函数。但是主机不能直接访问设备内存，设备也不能直接访问主机内存。设备内存中的数据需要在设备运算开始之前从主机内存复制到设备内存，设备运算完成后再将结果复制回主机内存。
图2.1 OpenACC加速器模型
可以比照英伟达GPU来理解这个抽象加速器模型，实际上这个抽象模型就是在它的基础上设计的。设备与设备内存通常离得很近（例如在同一块显卡上），带宽最大。主机与主机内存距离也比较近，带宽次大。主机内存与设备内存距离较远，连接带宽最小。英伟达GPU这样的产品称为分离内存设备
AMD APU这样的产品中，主机内存和设备内存共用一块物理空间，它们之间可以共享数据，不需要复制搬迁。这类产品称为共享内存设备。
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区域通常包含一个或多个工作分摊（worksharing）循环，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章会有例子详解。
图2.2 OpenACC的3个并行执行层次
本节接下来的内容可能较难理解，这是因为它们来自OpenACC规范，概括性强，初次阅读时不必深究，可以读完全书以后回过头来慢慢领会。
执行一个计算区域时，设备会启动一个或多个gang，每个gang都包含一个或多个worker，而每个worker可能还有能力执行一个或多个vector通道。开始执行时，多个gang处于gang冗余模式：每个gang中的每一个worker的每一个vector通道都冗余地执行相同的代码。当到达一个标记为gang层次工作分摊的循环或嵌套循环时，程序开始以gang分裂模式执行，这个循环或嵌套循环的所有迭代步都将分裂，然后分配给各个gang，以实现真正的并行执行，但是每个活动的gang内仅有一个worker，并且一个worker内仅有一个vector通道。
当只有一个活动worker时，无论是在gang冗余模式中还是在gang分裂模式中，程序都处于worker单独模式。当只有一个活动vector通道时，程序就处于vector单独模式。当一个gang到达标记为worker层次工作分摊的循环或嵌套循环时，这个gang就转换为worker分裂模式，从而激活这个gang中的所有worker。循环或嵌套循环的所有迭代步分裂分配给这个gang中的各个worker。如果一个循环同时被标记为gang分裂和worker分裂，那么循环里的所有迭代步将分散到所有gang的所有worker上。如果一个worker到达一个标记为vector层次并行的循环或嵌套循环，那么这个worker将转换为vector分裂模式。与worker分裂模式类似，转换为worker分裂模式将激活这个worker中的所有vector通道。使用向量操作或SIMD操作时，这个循环或嵌套循环的所有迭代步将分散给所有vector通道。单个循环可以被标记为gang并行、worker并行、vector并行中的一种、二种或三种，相应地，所有的迭代步会被酌情分散到所有的gang、worker、vector之上。程序员可以手动地指定使用哪些并行层次以及使用多少个gang、worker、vector，但不一定是最优的。程序员不手动指定时，编译器会选择它自认为最优的并行方式。
主机程序以单线程开始执行，这个线程可以用OpenMP编程接口之类的工具衍生出更多线程。在加速器上，单个gang的单个worker的单个vector通道称为一个线程；在设备上执行时，程序会创建一个并行执行上下文，该上下文可能包含很多这样的线程。
程序员不要试图在任何gang之间、worker之间或vector之间使用障碍同步、临界区域或锁。执行模型允许编译器将一些gang执行完后再开始执行其他gang，这意味着在gang之间实施同步操作很可能会失败。特别是，gang之间的障碍操作无法以可移植的方式实现，因为所有的gang可能永远不会在同一时刻处于活动状态。相似地，执行模型允许编译器执行完一个gang中的一些worker或一个worker中的一些vector通道之后，再开始执行其他的worker或vector通道。也允许编译器将一些worker或vector通道挂起，直到其他worker或vector通道执行完毕。这意味着在worker或vector通道之间实施同步操作很可能会失败。如果使用原子操作和一个忙碌-等待循环来在worker或vector通道间实现一个障碍或关键区域可能永远不会成功，这是因为调度器可能将拥有锁的worker或vector通道挂起，导致等待这个锁的worker或vector通道永远无法完成。
在某些设备上，加速器也可以创建和启动并行内核，并允许嵌套并行。这种情况下，OpenACC导语可以被一个主机线程执行，也可以被一个加速器线程执行。OpenACC规范使用术语本地线程和本地内存来表示执行导语的线程和与该线程关联的内存，无论该线程是在主机上还是在加速器上。
相对于主机线程，大多数加速器可以异步操作。对这种设备，加速器有一个或多个活动队列。主机线程将数据搬移、过程执行等操作压入活动队列。压入操作完成后，当设备正在独立、异步地工作时，主机线程继续向后执行。主机线程可以查询活动队列状态并等待某个队列的所有操作完成。某个活动队列上的操作完成后，才会执行同一队列上的下一个操作；不同活动队列上的操作可以同时处于活动状态，并且可以以任意顺序完成。
2.2 OpenACC 2.5规范
本节列出OpenACC的主要构件、导语，读完本书后可以在此处快速查阅语法，不必到正文中寻找零星的介绍。初次阅读请跳过。
1．导语一般格式
C/C++：
#pragma acc 导语名字 [子语列表] 换行
Fortran
!$acc 导语名字 [子语列表]
2. parallel构件
C/C++：
#pragma acc parallel [子语列表] 换行
结构块
Fortran：
!$acc parallel [子语列表]
结构块
!$acc end parallel
parallel构件的子语：
async[( 整数表达式 )]
wait[(整数表达式列表 )]
num_gangs(整数表达式)
num_workers(整数表达式 )
vector_length(整数表达式 )
device_type(设备类型列表)
if(条件)
reduction(操作符：变量列表)
copy(变量列表)
copyin(变量列表)
copyout(变量列表)
create(变量列表)
present(变量列表)
deviceptr(变量列表)
private(变量列表)
firstprivate(变量列表)
default(none｜present)
3. kernels构件
C/C++：
#pragma acc kernels [子语列表] 换行
结构块
Fortran：
!$acc kernels [子语列表]
结构块
!$acc end kernels
kernels构件的子语：
async[(整数表达式)]
wait[(整数表达式列表)]
num_gangs(整数表达式)
num_workers(整数表达式)
vector_length(整数表达式)
device_type(设备类型列表)
if(条件)
copy(变量列表)
copyin(变量列表)
copyout(变量列表)
create(变量列表)
present(变量列表)
deviceptr(变量列表)
default(none｜present)
4. data构件
C/C++：
#pragma acc data [子语列表] 换行
结构块
Fortran：
!$acc data [子语列表]
结构块
!$acc end data
data构件的子语：
if(条件)
copy(变量列表)
copyin(变量列表)
copyout(变量列表)
create(变量列表)
present(变量列表)
deviceptr(变量列表)
5. enter data导语
C/C++：
#pragma acc enter data 子语列表 换行
Fortran：
!$acc enter data 子语列表
enter data的子语：
if(条件)
async[(整数表达式)]
wait[(整数表达式列表)]
copyin(变量列表)
create(变量列表)
6. exit data导语
C/C++：
#pragma acc exit data 子语列表 换行
Fortran：
!$acc exit data 子语列表
exit data的子语：
if(条件)
async[(整数表达式)]
wait[(整数表达式列表)]
copyout(变量列表)
delete(变量列表)
finalize
7. host_data导语
C/C++：
#pragma acc host_data 子语列表 换行
结构化块
Fortran：
!$acc host_data 子语列表
结构化块
!$acc end host_data
host_data的子语：
use_device(变量列表)
8. loop导语
C/C++：
#pragma acc loop [子语表表] 换行
for循环
Fortran：
!$acc loop [子语列表]
do循环
loop的子语：
collapse(n)
gang[(gang参数列表)]
worker[([num:]整数表达式)]
vector[([length:]整数表达式)]
seq
auto
tile(尺寸表达式列表)
device_type(设备类型列表)
independent
private(列表)
reduction(操作符：列表)
9．组合导语
C/ C++：
#pragma acc parallel loop [子语列表] 换行
for 循环
#pragma acc kernels loop [子语列表] 换行
for 循环
Fortran：
!$acc parallel loop [子语列表]
do 循环
[! $acc end parallel loop]
!$acc kernels loop [子语列表]
do 循环
[! $acc end kernels loop]
C/C++：
10. declare导语
#pragma acc declare 子语列表 换行
Fortran：
!$acc declare 子语列表
declare的子语：
copy(变量列表)
copyin(变量列表)
copyout(变量列表)
create(变量列表)
present(变量列表)
deviceptr(变量列表)
device_resident(变量列表)
link(变量列表)
11. init导语
C/C++：
#pragma acc init [子语列表] 换行
Fortran：
!$acc init [子语列表]
init的子语：
device_type( 设备类型列表 )
device_num( 整数表达式 )
12. shutdown导语
C/C++：
#pragma acc shutdown [子语列表] 换行
Fortran：
!$acc shutdown [子语列表]
shutdown的子语：
device_type( 设备类型列表 )
device_num( 整数表达式 )
13. set导语
C/C++：
#pragma acc set [子语列表] 换行
Fortran：
!$acc set [子语列表]
set的子语：
default_async( 整数表达式 )
device_num( 整数表达式 )
device_type( 设备类型列表 )
14. update导语
C/C++：
#pragma acc update 子语列表 换行
Fortran：
!$acc update 子语列表
update的子语：
async[(整数表达式)]
wait[(整数表达式列表)]
device_type(设备类型列表)
if(条件)
if_present
self(变量列表)
host(变量列表)
device(变量列表)
15. routine导语
C/C++：
#pragma acc routine 子语列表 换行
#pragma acc routine(名字) 子语列表 换行
Fortran：
!$acc routine 子语列表
!$acc routine(名字) 子语列表
set的子语：
gang
worker
vector
seq
bind(名字)
bind(字符串)
device_type(设备类型列表)
nohost
16. wait导语
C/C++：
#pragma acc wait [(整数表达式列表)] 子语列表 换行
Fortran：
!$acc wait [(整数表达式列表)] 子语列表
wait的子语：
async [( 整数表达式)]