开发者社区> 华章计算机> 正文

OpenACC并行编程实战》—— 3.4 loop构件

简介: kernels构件让编译器自动分析代码,挖掘代码里的并行性,并实施并行化。但是,编译器毕竟只是个软件,不会知道程序员的真实意图。若想更准确高效地指导编译器的并行化工作,程序员可以使用loop导语。该导语能告诉编译器哪些循环需要并行化,以及用什么方式并行化。
+关注继续查看

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

3.4 loop构件

kernels构件让编译器自动分析代码,挖掘代码里的并行性,并实施并行化。但是,编译器毕竟只是个软件,不会知道程序员的真实意图。若想更准确高效地指导编译器的并行化工作,程序员可以使用loop导语。该导语能告诉编译器哪些循环需要并行化,以及用什么方式并行化。
loop导语可用在kernels构件内,也可以用在parallel构件内。本节会具体讲解loop导语在两种计算构件中的行为,读完3.5节后读者会豁然开朗。
loop导语作用于紧跟该导语的一个循环。loop导语可以描述执行这个循环的并行类型,还可以声明循环的私有变量、数组和归约操作。

C和C++中,loop导语的语法是:

#pragma acc loop [子语列表] 换行
  for循环

Fortran中,loop导语的语法是:

!$acc loop [子语列表] 
  do循环

这里的子语是下列中的一个:

collapse(n)
gang [(gang参数列表)] 
worker[([num:]整数表达式)] 
vector[([length:]整数表达式)] 
seq 
auto 
tile(尺寸表达式列表)
device_type(设备类型列表)
independent
private(列表) 
reduction(操作符:列表)
这里的gang参数是下列中的一个:
[num:] 整数表达式
static: 尺寸表达式

并且gang参数列表至多只能有一个num或static参数,这里的尺寸表达式是下列中的一个:
        *
        整数表达式

一些子语只能用在kernels区域的上下文中,详述见下文。
字面上没有包含在一个parallel构件或kernels构件之内的loop构件称为孤立loop构件。字面上包含一个loop构件的最里层计算构件称为这个loop构件的父层计算构件。
使用限制:

  • 仅有collapse、gang、worker、vector、seq、auto和tile这几个子语可以跟在一个device_type子语后面。
  • worker和vector子语的参数整数表达式必须在kernels区域内保持不变。
  • 不带seq子语的loop构件的关联循环必须写成这样:进入loop构件的时候就能计算出迭代步的数量。

本节把所有子语的含义都罗列出来,但只挑选重要的子语详细讲解。提请注意,子语的含义均从OpenACC规范翻译而来,可能有拗口的地方,需要细细理解。

3.4.1 independent子语

independent子语告诉编译器该循环的迭代步是彼此数据独立的。因此允许生成并行执行这些迭代步的代码,且不需要同步。在一个parallel构件内,所有不带seq子语的loop导语都暗含一个independent子语。
3_4_1_1

数据依赖是指两步操作之间有数据依赖关系,例如要计算得到数组元素a[0]~a[N]的值,计算a[i]的时候需要用到aj的值,那么就称a[i]和a[j]之间有数据依赖。数据独立是两个操作之间没有数据依赖。
kernels构件会参考导语的指示,但并不是完全按照导语进行并行化,它如果发现问题,会拒绝并行化,例3.13和例3.14就是这样的情形。

471ebc8ca8353a66722ea2841e9313b2b88787f76f2996cf1f2eb491c1822efdbc742cec5d590065

例3.13的编译反馈如下:
14, Loop is parallelizable
    Accelerator kernel generated
    Generating Tesla code
    14, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
17, Loop carried dependence of b prevents parallelization
    Loop carried backward dependence of b prevents vectorization
    Accelerator scalar kernel generated

编译器报告,第17行代码中数组b存在后向数据依赖性,不能向量化(并行化)。这个现象证实,导语不是无条件执行的命令,而只是一种指导。串行执行时,数组b的元素是按图3.9这样移动的,从而执行结果就是:
b[2] = 0
符合代码原意。

d18411b450111bf76a644131219bae1d4503901e

但是,编译器毕竟不一定能完全理解程序员的设计思路,可能会误判。此时就可以用independent子语来告诉编译器:不要自行检查数据依赖关系了,程序员保证循环中的各个迭代步都是数据独立的,可以安全地并行计算。拿例3.13试一试效果,将第16行的
16 #pragma acc loop
改写为
16 #pragma acc loop independent
编译反馈变成:
     17, Loop is parallelizable
         Accelerator kernel generated
         Generating Tesla code
         17, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */

编译器虽然有能力检测到第18行的循环中存在数据依赖,但程序员的指示优先级更高,它不再自行检查,直接按照independent子语的要求并行化。计算结果当然是错的:
b[2] = 1
错误形成的机制看图3.10:7次数据移动本来是按照先后顺序一个接一个进行的,并行执行时就变成了同时移动。
编译器相当聪明,能发现常见的数据依赖关系,但它终究没有人聪明,不能完全理解程序员的用意。对例3.15和例3.16,编译器自动检测就判断错了。

14caa5ef7e8a3a6469c62bdc230ed023fc40a58c

例3.15的编译反馈如下:
14, Loop carried dependence of a prevents parallelization
    Loop carried backward dependence of a prevents vectorization
    Loop carried dependence due to exposed use of a[:] prevents parallelization
    Accelerator scalar kernel generated

编译器拿不准第15行的a[a[i]]是否有交叉,谨慎起见,拒绝并行化,生成的内核保持串行执行。实际上,程序员知道a[a[i]]与a[i]等效,没有数据依赖性,但是编译器不知道啊。此时加上independent子语就正确了。
由本节几个例子可知:independent子语优先级高于编译器的自动检测,程序员要自己保证循环里的迭代步之间确实没有依赖性。

3.4.2 reduction归约子语

科学和工程计算中常要计算级数的前n项和,例如_1,在n较大的时候,并行计算能够显著提高速度。并行计算的理论依据是加法的结合律和交换律:

b9cf4dea430239c4ba9c0cbc45aa0b4c44d85fef
这里的S1和S2称为部分和。S1和S2同时计算理论上可以将计算时间减少一半,增加部分和的数量,可以进一步减少计算时间。因为加法满足结合律和交换律,所以S1和S2可以有多种组合方法,例如图3.11和图3.12。
e4e1a0f00d6c3564dda7c2a649f9f89124985c47

串行计算s时,代码行为:
s = 0;
for(int i = 1; i <= n; i++ )
  s = s + a[i];

并行计算s时,代码行为:
3_12_
计算s1和s2的两段代码由两个线程同时运行,都计算结束后,将两个部分和相加,得到最终需要的值s。这里列出的并行代码只是为了说明编译器的并行化原理,真实的实现方法可能有差别。s1和s2的初始值这里为0,这是因为运算符是加法,0初始值不会影响最终结果的正确性。如果是连乘操作,初始值就应为1,其他值将导致错误的结果。
reduction子语含义:reduction子语可以用到parallel构件上。它指定一个归约操作符和一个或多个标量变量。对每个变量,每个并行运行的gang都会创建一个副本,并根据指定的运算符来初始化这个副本。在并行区域结束时,归约运算符使用每一个gang里的值、变量的原始值计算出归约结果,并将该结果存入原始变量。归约结果在并行区域之后可用。以图3.11和图3.12为例,部分和s1和s2由两个gang分别计算。
表3.2列出了可用的操作符及相应的初始值,初始值跟数据类型有关。对max和min归约,初始值分别为变量所属数据类型能表示的最小值、最大值。C与C++(int、float、double、complex)、C++(char、wchar_t、int、float、double)和Fortran(integer、real、double precision、complex)的数值数据类型都可以用在本子语中。

bfa90d4a90195b27c91cdf16f8552b1926b1e1f6

使用限制:
  • 归约变量不能是数组的一个元素。
  • 归约变量不能是一个C结构体成员、C++类或结构体成员,也不能是Fortran派生类型成员。

下面给出几个reduction子语的典型例子(例3.17~例3.29),看看它的具体用法,类似的操作符就不再举例。

d81f7edd9c15f74fb40bc2fa83ceab62b5c3c7e1
在例3.17第9~14行的并行区域内,计算得到所有a[i]的和,在区域出口,与区域外ired的初始值结合起来,存入ired。归约变量在并行区域结束后才可以使用,否则计算结果不正确。
1fef612eacfcc26d37f7d820ed212c680ad9da473c31a3114de5e2d7863d59ad19794fc204df20c0
例3.18第12~13行计算数组a上的加法归约,结果存入ired。在并行区域内的第15行就使用ired的值,计算结果为:
ired = 5050
a[1] = 2

ired的值是正确的,但a[1]的值是错误的,原因就是在区域内就使用了归约结果。在并行区域内部,不同gang中不同线程所读到的ired的值可能是不同的。在并行区域出口处,所有gang中的部分和结合在一起存入主机变量ired之中,然后才能正确使用。

725e66415e47b726bce176f726d564b844b89fd54c6c5b6a0074a63cd6287875d56946bc8642841c

最大值和最小值操作符,赋值语句一定用大于号或者小于号,如例3.23和例3.25的第13行。
2e73fe8914c861ced8fd3a12a9adcc801971b33bbbb64a5ab8a001a70322805b1915ae9c633f0ebc870635a6e2a889c015421b83f80cc98589e5a0a4

C语言中的按位或|、与&&、或||跟按位与&用法类似,Fortran语言中的按位或ior、按位异或ieor跟按位与iand用法类似,Fortran语言的逻辑或.or.、逻辑等价.eqv.、逻辑不等价.neqv.与逻辑与.and.类似,不再一一举例。

3.4.3 不常用的子语

下面的子语常用于深度调优,在本书中应用不是特别多,只列出含义,不再举例。

(1)private子语

loop构件上的private子语指明,要为变量列表中的每一项创建一个副本。如果循环体以向量分裂模式执行,那么与每一个向量通道相关联的每一个线程都会创建一个副本。如果循环体以worker分裂vector单独的模式执行,那么将只创建一个变量副本,并在每一个worker内部的所有向量通道的关联线程间共享。其他情形中,创建一个副本,并在每一个gang中的所有worker中的所有向量通道的关联线程之间共享。详述参见4.2.8节。

(2)collapse子语

collapse子语用来指定与loop构件相关联的紧密嵌套循环有多少层。collapse子语的参数必须是一个常量正整数表达式。如果没有collapse子语,只有紧接着的循环才与loop构件相关联。
如果有一个以上的循环与loop构件关联,那么关联循环中的所有循环体都会按照剩余的子语来调度。与collapse子语相关联的所有循环的迭代步数必须是可计算的,并在所有循环中保持不变。
编译器自行决定是否将导语的gang、worker或vector子语应用到每一个循环、线性化迭代空间。

(3)gang子语

当上一层计算构件是一个kernels构件时,gang子语指明,关联循环的迭代步需要在为内核创建的所有gang上并行执行。如果指定一个不带关键字的参数,或者指定一个放在num关键字后面的参数,那么该参数指明使用多少个gang来执行本循环的迭代步。除非处在一个嵌套的parallel区域内或嵌套的kernels区域内,否则不允许一个带gang子语的loop区域包含另外一个带gang子语的loop区域。
除非出现static参数,否则循环的迭代步对各个gang的调度方式不确定。如果出现一个带星号的static参数,那么编译器将自行决定块的尺寸。循环的所有迭代步以选定的尺寸分割成多块,所有gang从零号块开始轮流获得若干块。同一个parallel区域内,循环步数相同的两个gang循环,如果使用带相同参数的static子语,那么循环步分配到gang上的方式相同。同一个kernels区域内,循环步数相同的两个gang循环,如果使用相同数量的gang,static子语的参数也相同,那么循环步分配到gang上的方式相同。

(4)worker子语

在一个加速器parallel区域内,worker子语指明,关联循环的迭代步将被分配给单个gang的多个worker并行执行。带worker子语的loop构件将一个gang由worker单独模式切换为worker分裂模式。与gang对比,worker子语先激活额外的worker层次并行,然后将循环分摊到这些worker上。不允许使用参数。循环的所有迭代步必须是数据独立的,但reduction子句中指定的变量除外。除非在一个嵌套的计算区域内,否则不允许一个带worker子语的loop区域包含另外一个带gang子语的loop区域。
当上一层计算构件是kernels构件时,worker子语指明,关联循环的迭代步需要在一个gang内的所有worker上并行执行。仅当kernels构件上没有出现num_workers子语时,才允许指定一个参数。这个可选参数用来指定每个gang使用多少个worker来执行这个循环的迭代步。除非位于一个嵌套的计算区域内,否则不允许一个带worker子语的loop区域包含另外一个带gang子语的loop区域。
在任何worker继续处理本循环后面的代码之前,所有worker都已将各自承担的迭代步执行完毕。

(5)vector子语

在一个加速器parallel区域内,vector子语指明,关联循环的迭代步以向量模式或SIMD模式执行。一个带有vector子语的loop构件将一个worker由向量单独模式切换为向量分裂模式。与worker子语相似,vector子语先激活额外的向量层次并行,然后将循环的迭代步分摊到这些向量通道上。向量的长度由本子语指定或由编译器为本parallel区域自动选择。除非在一个嵌套的计算区域内,否则不允许一个带vector子语的loop区域包含另外一个带gang子语、worker子语或vector子语的循环。
在一个加速器kernels区域,vector子语指明,关联循环的迭代步以向量模式或SIMD模式执行。如果指定了一个参数,迭代步将以这个长度的向量执行;如果没有指定参数,编译器将自行选择一个合适的向量长度。除非在一个嵌套的计算区域内,否则不允许一个带vector子语的loop区域包含另外一个带gang子语、worker子语或vector子语的循环。
在任何vector继续处理本循环后面的代码之前,所有vector都已将各自承担的迭代步执行完毕。

(6)seq子语

考虑这样一个场景:几个可并行计算的循环之间夹着一个串行循环,如果拆成两个计算构件,那么将多出几次数据传输,不划算。本书后面章节将详细介绍数据传输,回过头来再看本节的时候就能明白它的意义。
seq子语指明关联循环需要在加速器上串行地执行。本子语将阻止所有的自动并行化、向量化。

(7)auto子语

auto子语指明,编译器必须自动分析循环并查明循环迭代步是否数据独立,如果数据独立,那么选择到底是并行执行这个循环还是串行执行这个循环。受内、外层循环上带有gang、worker或vector子语的loop导语的影响,编译器能够实施的并行类型可能会受到限制。当上一层计算构件是kernels构件的时候,不带independent子语或seq子语的loop构件视为带有auto子语,即auto是缺省子语。

(8)device_type子语

device_type子语的详细描述在4.7节。

(9)tile子语

tile子语指明,编译器应该将嵌套循环里的每一层循环都剖分成两个循环:一个外层瓦片循环,和一个内层元素循环。tile子语的参数是一个瓦片尺寸列表,每一个瓦片尺寸都可以是一个常量正整数表达式或者一个星号。如果列表中有n个瓦片尺寸,那么loop导语后面必须紧跟着n个紧密嵌套的循环。尺寸表达式列表中的第一个参数对应n个关联循环中的最内层循环,最后一个参数对应最外层的关联循环。如果瓦片尺寸被指定为星号,那么编译器将自主选择一个合适的值。嵌套循环中的每一个循环都将被分割为两个循环,一个外层瓦片循环和一个内层元素循环。元素循环的路程步数受限于尺寸表达式列表中的相应瓦片尺寸。编译器会调整瓦片循环和元素循环的嵌套顺序,瓦片循环均处于所有元素循环的外层,元素循环均处于所有瓦片循环的内层。
如果loop导语上带有vector子语,那么这个vector子语会作用在元素循环上。如果loop导语上带有gang子语,那么这个gang子语会作用到瓦片循环上。loop导语上带有worker子语时,如果没有出现vector子语,那么这个worker子语会作用在元素循环上,否则作用在瓦片循环上。

版权声明:本文内容由阿里云实名注册用户自发贡献,版权归原作者所有,阿里云开发者社区不拥有其著作权,亦不承担相应法律责任。具体规则请查看《阿里云开发者社区用户服务协议》和《阿里云开发者社区知识产权保护指引》。如果您发现本社区中有涉嫌抄袭的内容,填写侵权投诉表单进行举报,一经查实,本社区将立刻删除涉嫌侵权内容。

相关文章
SpringBoot ~ AOP切面编程
AOP切面编程 添加pom依赖 <dependency> <groupId>org.springframework.boot</groupId> <artifactId>spring-boot-starter-aop</artifactId> </dependency> 编写切面类 /** * @author wsyjlly * @create 2019.
780 0
《Linux/UNIX OpenLDAP实战指南》——2.7 OpenLDAP用户以及与用户组相关的配置
添加用户和用户组的方式有两种。一种是将系统用户通过migrationtools工具生成LDIF文件并结合ldapadd命令导入OpenLDAP目录树中,生成OpenLDAP用户。另一种通过自定义LDIF文件并通过OpenLDAP命令进行添加或者修改操作。
2746 0
使用OpenApi弹性释放和设置云服务器ECS释放
云服务器ECS的一个重要特性就是按需创建资源。您可以在业务高峰期按需弹性的自定义规则进行资源创建,在完成业务计算的时候释放资源。本篇将提供几个Tips帮助您更加容易和自动化的完成云服务器的释放和弹性设置。
18689 0
10059
文章
0
问答
来源圈子
更多
+ 订阅
文章排行榜
最热
最新
相关电子书
更多
JS零基础入门教程(上册)
立即下载
性能优化方法论
立即下载
手把手学习日志服务SLS,云启实验室实战指南
立即下载