CN111966397B

CN111966397B - 一种异构并行程序自动移植和优化方法

Info

Publication number: CN111966397B
Application number: CN202010710022.2A
Authority: CN
Inventors: 张伟哲; 王法瑞; 何慧; 郭浩男; 刘亚维; 张玥
Original assignee: Harbin Institute of Technology Shenzhen
Current assignee: Harbin Institute of Technology Shenzhen
Priority date: 2020-07-22
Filing date: 2020-07-22
Publication date: 2025-01-07
Anticipated expiration: 2040-07-22
Also published as: CN111966397A

Abstract

一种异构并行程序自动移植和优化方法，属于异构并行程序开发技术。本发明是为了实现CPU并行程序自动移植、在减轻开发人员工作负担的同时提高程序性能，从而解决并行指令转换、数据传输管理及优化问题。技术要点：构建异构并行程序自动移植系统的框架，异构并行程序自动移植系统用于将OpenMP CPU并行程序自动翻译为OpenMP Offloading异构并行程序；一致性状态转换形式化，保证数据一致性的前提下，优化传输操作，减少冗余数据传输；运行时库设计，运行时库用于提供自动数据传输管理和优化功能，维护每个变量内存区域一致性状态；源到源翻译器设计，翻译器用于自动转换并行指令及自动插入运行时API。该方法可以自动识别CPU并行指令并转换为加速器并行指令，提高程序性能。

Description

一种异构并行程序自动移植和优化方法

技术领域

本发明涉及一种异构并行程序自动移植和优化方法，属于异构并行程序开发技术。

背景技术

随着人工智能、图像处理、多物理场仿真、量子模拟、气候模拟等不同应用对算力的巨大需求，基于各类加速器的异构平台已经取代CPU(Central Processing Unit，中央处理器)成为算力的主要来源。在高性能计算领域主要使用GPU(Graphics Processing Unit，图形处理器)作为加速器，在移动平台上主要使用GPU、DSP(Digital Signal Processing，数字信号处理器)或FPGA(Field Programmable Gate Array，现场可编程逻辑门阵列)作为加速器。加速器在提供巨大算力的同时，也给应用开发和移植带来了巨大挑战。

CPU并行编程的标准是OpenMP(Open Multi-Processing)模型，而异构并行编程需要使用CUDA(Compute Unified Device Architecture，统一计算架构)、OpenCL(OpenComputing Language，开放计算语言)、OpenACC(OpenAccelerators，开源加速器)、OpenMPOffloading(OpenMP的扩展)等异构并行编程模型。编写高效的异构并行程序往往需要开发人员既熟悉异构平台特性，又掌握并行编程模型；即使开发人员有能力，将CPU并行程序移植到异构平台也是非常耗时且容易出错的工作。因此亟需自动移植工具，自动地将OpenMPCPU并行程序自动移植到异构平台。

文献号为CN104035781A(CN104035781B)的现有技术提供一种快速开发异构并行程序的方法，涉及CPU串行程序的性能分析和异构并行程序的移植：首先对CPU串行程序进行性能和算法分析，定位程序性能瓶颈和其可并行性；然后在原有代码的基础上插入OpenACC预编译指令，得到可在异构并行环境上执行的异构并行代码；将代码根据软硬件平台指定参数编译执行，根据程序运行结果确定是否需要进一步的优化。该现有技术能够高效的对现有程序进行并行化，使程序充分利用异构系统的计算能力，实用性强，易于推广。但是，在移植过程中的两大难题：并行指令转换、数据传输管理及优化仍没有解决。

发明内容

本发明要解决的技术问题为：

本发明的目的是提供一种异构并行程序自动移植和优化方法，以实现CPU并行程序自动移植、在减轻开发人员工作负担的同时提高程序性能，从而解决并行指令转换、数据传输管理及优化问题。

本发明为解决上述技术问题所采用的技术方案为：

一种异构并行程序自动移植和优化方法，所述方法的实现过程为：

步骤1、构建异构并行程序自动移植系统的框架

异构并行程序自动移植系统，简称OAO系统，用于将OpenMP CPU并行程序自动翻译为OpenMP Offloading异构并行程序，并结合运行时系统自动管理并优化CPU和加速器之间的数据传输；OAO系统框架主要包含源到源翻译器、运行时库两部分；

运行时库包含三类API：一致性状态跟踪、数据传输、一致性状态转换；一致性状态跟踪API捕获变量内存区域，并初始化一致性状态；数据传输API根据当前一致性状态和需要满足的一致性约束动态确定传输操作并执行，同时更新一致性状态；状态转换API根据读写操作类型，更新一致性状态；

源到源翻译器将OpenMP CPU并行代码翻译为OpenMP Offloading异构并行代码，并插入合适的运行时API；其由3个模块组成：数据传输API插入模块、状态转换API插入模块、并行指令翻译模块；并行指令翻译模块将OpenMP CPU并行指令翻译为OpenMPOffloading异构并行指令，得到OpenMP Offloading内核；数据传输API插入模块、状态转换API插入模块分别插入两种对应的运行时API；

OAO系统工作过程如下：OpenMP CPU并行代码经过源到源翻译，得到包含运行时API的OpenMP Offloading异构并行代码，OpenMP Offloading代码经过编译后在异构平台上运行，其中OpenMP Offloading内核运行在加速器上，其余程序运行在CPU上；运行时库通过插入的API管理CPU和加速器之间的数据传输，保证数据一致性并动态进行传输优化，减少冗余数据传输；

步骤2、一致性状态转换形式化

对于异构平台，变量在CPU内存和加速器内存中都存在副本，一致性状态用来描述变量的CPU副本和加速器副本的有效性；由当前的一致性状态和需要满足的一致性状态约束推导出最简状态转换函数，该函数通过对应关系确定需要执行的最简传输操作类型，以保证数据一致性的前提下，优化传输操作，减少冗余数据传输；

步骤3、运行时库设计

运行时库用于提供自动数据传输管理和优化功能，维护每个变量内存区域一致性状态，给出运行时库三类API函数及对应的描述；

步骤4源到源翻译器设计

源到源翻译器主要基于Clang/LLVM编译框架的静态分析功能，将OpenMP CPU并行代码翻译为OpenMP Offloading异构并行代码，并以合适的方式插入运行时API，在保证数据一致性的前提下，优化数据传输；

源到源翻译器通过静态分析收集串行域、并行域、变量引用等信息；然后建立程序的串并行控制流图并将变量引用信息与相应的并行域串行域绑定；之后以串行域和并行域为粒度插入数据传输API插入、状态转换API插入；最后将OpenMP CPU并行指令翻译为OpenMP Offloading并行指令。

本发明具有以下有益技术效果：

本发明提出了一种从OpenMP CPU并行程序到OpenMP Offloading异构并行程序的自动移植和优化方法(简称OAO，OpenMPAutomatic Offloading)，基于OAO方法能实现自动移植系统，该系统能在减轻开发人员工作负担的同时提高程序性能。OAO方法能够自动解决移植过程中的两大难题：并行指令转换、数据传输管理及优化。该方法可以自动识别CPU并行指令并转换为加速器并行指令。该方法通过使用运行时系统，能够在程序执行过程中动态优化数据传输，从而在保证数据一致性的同时减少冗余数据传输，提高程序性能。

对于大量现存的OpenMP CPU并行程序，开发人员可以直接使用OAO系统实现自动异构移植，利用异构平台算力获得程序性能提升。对于新应用，开发人员可以继续使用熟悉的OpenMP模型，然后使用OAO系统进行自动异构移植，获得程序性能提升。

附图说明

图1为异构并行程序自动移植系统(OAO系统)框架流程图；

图2为一致性状态转换关系图；

图3为串并行控制流图；

图4为OpenMP CPU并行程序的截图；

图5为人工编写的OpenMP Offloading异构并行程序的截图；

图6为源到源翻译所得OpenMP Offloading异构并行程序的截图；

图7为其他版本相对于OMP版本加速比(K40平台)柱状图；

图8为其他版本相对于OMP版本加速比(2080Ti平台)柱状图；

图9为OAO版本相对于其他版本节省的传输数据量百分比柱状图；

图10为OAO版本相对于其他版本节省的传输时间百分比(2080Ti平台)柱状图。

图7至图10中的英文及缩写为本领域的公知术语。

具体实施方式

下面结合附图1至10对本发明的实现进行如下阐述：

1、异构并行程序自动移植系统框架

异构并行程序自动移植系统(简称OAO系统)，可以将OpenMP CPU并行程序自动翻译为OpenMP Offloading异构并行程序，并结合运行时系统自动管理并优化CPU和加速器之间的数据传输。OAO系统框架如图1所示，主要包含两部分(图中阴影部分)：源到源翻译器，运行时库。

运行时库提供自动数据传输管理和优化功能，维护每个变量内存区域一致性状态，其包含三类API(Application Programming Interface，应用程序接口)：一致性状态跟踪、数据传输、一致性状态转换。一致性状态跟踪API捕获变量内存区域，并初始化一致性状态。数据传输API，根据当前一致性状态和需要满足的一致性约束动态确定传输操作并执行，同时更新一致性状态。状态转换API根据读写操作类型，更新一致性状态。

源到源翻译器将OpenMP CPU并行代码翻译为OpenMP Offloading异构并行代码，并插入合适的运行时API；其由3个模块组成：数据传输API插入模块、状态转换API插入模块、并行指令翻译模块。并行指令翻译模块将OpenMP CPU并行指令翻译为OpenMPOffloading异构并行指令，得到OpenMP Offloading内核。数据传输API插入模块、状态转换API插入模块分别插入两种对应的运行时API。

OAO系统工作过程如下：OpenMP CPU并行代码经过源到源翻译，得到包含运行时API的OpenMP Offloading异构并行代码。OpenMP Offloading代码经过编译后在异构平台上运行，其中OpenMP Offloading内核运行在加速器上，其余程序运行在CPU上。运行时库通过插入的API管理CPU和加速器之间的数据传输，保证数据一致性并进行动态传输优化，减少冗余数据传输。

2、一致性状态转换形式化

数据传输管理及优化是OpenMP CPU程序向异构平台移植过程中的关键，本部分对变量一致性状态、一致性状态转换、一致性状态约束等进行形式化，为之后的运行时库和源到源翻译器设计提供理论支持。

对于异构平台，变量可能在CPU内存和加速器内存中都存在副本，一致性状态用来描述变量的CPU副本和加速器副本的有效性，其定义如下：

定义1.一致性状态State是一个3bit的二进制数。其中Bit0表示变量是(Bit0＝1)/否(Bit0＝0)存在加速器副本；Bit1表示CPU副本是(Bit1＝1)/否(Bit1＝0)有效；Bit2表示加速器副本是(Bit2＝1)/否(Bit2＝0)有效。

表1所有可能一致性状态

根据定义1，给出所有可能的一致性状态如表1所示。CPU和加速器之间的数据传输以及CPU和加速器上的读写操作会改变一致性状态，使用一致性状态转换函数来表示一致性状态的改变，其定义如下：

定义2.一致性状态转换函数TransFunc的形式如式(1)所示，其中InVld和Vld是一对3bit的二进制数，其中运算符是布尔运算符。若将InVld中的某bit置0，则可将inState中对应bit转换为0；若将Vld中的某bit置1则可将inState中对应bit转换为1；因此给二者设置不同的值可以将任意inState转换为任意outState。TransFunc也简记为式(2)形式。

TransFunc＝{InVld,Vld} (2)

根据定义2，给出所有可能的一致性状态转换函数及其对应的数据传输或读写操作如所示。

表2所有可能的一致性状态转换操作

根据定义1和定义2，给出各种一致性状态之间的转换关系图，如图2所示。

为了保证程序正确，CPU和加速器上的读写操作对一致性状态有一定要求，使用一致性状态约束表示这些要求，其定义如下：

定义3.一致性状态约束Constr由一对3bit二进制数组成，其形式如(3)所示。Constr的默认值是{111,000}，表示对一致性状态无要求。可以将ConInVld的某bit置0，或将ConVld的某bit置1表示对一致性状态的不同约束要求，如表3所示。

Constr＝{ConInVld,ConVld} (3)

表3ConVld和ConInVld含义描述

根据定义3和表3，给出CPU和加速器上不同读写操作所需的一致性状态约束，如表4所示。

表4所有可能的一致性约束

由当前的一致性状态State和需要满足的一致性状态约束Constr＝{ConInVld,ConVld}，可以推导出最简状态转换函数如下：

最简状态转换函数的推导公式如式(4)所示。

MinTrFunc(State)＝State·InVld+Vld

其中

最简状态转换函数表示从一致性状态State出发，为了满足一致性状态约束Constr，所必须执行的最简的无冗余的数据传输操作。因此推导出MinTrFunc即可通过表2中的对应关系(前6种)，确定需要执行的最简传输操作类型。从而在保证数据一致性的前提下，优化传输操作，减少冗余数据传输。

3、运行时库设计

运行时库API函数如表5所示，可以分成三类：一致性状态跟踪(前6个)，数据传输(OAODataTrans)，一致性状态转换(OAOStTrans)；下面将分别介绍。表5运行时库API函数

3.1一致性状态跟踪API

运行时库将变量内存区域作为一致性状态跟踪和数据传输的粒度。在C/C++中变量内存区域是一个连续的内存区域，其来源可能是局部变量定义、全局变量定义、malloc操作、new操作等。为了记录并跟踪变量一致性状态，定义变量内存区域和内存环境的形式化表示如下：

定义4.变量内存区域MemBlk是一个四元组，如式(5)所示；其中：Begin表示内存起始地址，Length表示内存区域长度，ElemSize表示元素大小，State表示一致性状态。

MemBlk＝{Begin,Length,ElemSize,State} (5)

定义5.内存环境MemEnv是全部变量内存区域的集合，如式(6)所示。

MemEnv＝{MemBlk₁,…,MemBlk_n} (6)

运行时库将MemEnv定义为全局变量，并在整个OpenMPOffloading异构并行程序执行期间维护MemEnv。当某变量被引用时，可以使用相应的指针ptr在MemEnv中搜索，满足式(7)的即为被引用的MemBlk。

Begin≤ptr≤Begin+Length-1 (7)

一致性状态跟踪API会在源到源翻译过程中被以合适的方式插入源代码。OAOSaveArrayInfo函数被插入到局部变量声明之后或main函数开始处(对于全局变量)，用来收变量信息。OAOMalloc函数替换malloc函数，收集内存分配信息，并进行内存分配。OAONewInfo被插入到new操作之后，收集内存分配信息。以上三个函数利用收集到的信息新建对应MemBlk，并将State初始化为HOST_ONLY。

OAODeleteArrayInfo被插入到变量作用域结束处，main函数结束处(对于全局变量)，删除对应MemBlk。OAOFree替换free函数，释放内存区域同时删除对应MemBlk。OAODeleteInfo被插入到delete操作之后，释放内存区域同时删除对应MemBlk。

此外运行时针对NVIDIA设备进行了专门优化，当需要分配的内存大于128KB时，使用cudaMallocHost()替代malloc()，用来分配内存。

3.2数据传输API

如表4中的分析，变量读写操作需要变量满足某种一致性状态约束，因此在访问某变量之前需要调用数据传输API，即OAODataTrans函数，来执行最简数据传输操作满足对应约束。本部分说明数据传输API原理，其插入方法将在第4部分中说明。OAODataTrans函数使用如下算法1，来确定为了满足一致性状态约束Constr，所需的最简状态转换函数，并执行对应的最简数据传输操作，同时更新一致性状态。算法1的1～2行通过变量指针ptr找到对应的变量内存区域MemBlk；第3行根据需要满足的一致性约束Constr和当前一致性状态State通过公式(4)推导最简状态转换函数MinTrFunc；第4行执行表2中与MinTrFunc对应的数据传输操作；第5行根据公式(1)使用MinTrFunc更新State。

3.3一致性状态转换API

如表2(后三行)中的分析，读写操作可能改变变量一致性状态，因此在某变量被访问之后需要调用一致性状态转换API，即OAOStTrans函数，来更新变量内存区域MemBlk中保存的一致性状态State。本部分说明一致性状态转换API原理，其插入方法将在第4部分中说明。OAOStTrans函数使用如下算法2完成一致性状态转换过程。算法2的1～2行通过变量指针ptr找到对应的变量内存区域MemBlk；第3行根据公式(1)使用状态转换函数StTrans更新MemBlk中包含的一致性状态State。

4源到源翻译器设计

源到源翻译器主要基于Clang/LLVM(一种C/C++语言编译器)编译框架的静态分析功能，将OpenMPCPU并行代码翻译为OpenMPOffloading异构并行代码，并以合适的方式插入运行时API，在保证数据一致性的前提下，优化数据传输。

源到源翻译器通过静态分析收集串行域(定义6)、并行域(定义7)、变量引用等信息；然后建立程序的串并行控制流图(定义8)并将变量引用信息与相应的并行域串行域绑定(定义9-10)；之后以串行域和并行域为粒度插入数据传输API插入、状态转换API插入；最后将OpenMPCPU并行指令翻译为OpenMPOffloading并行指令。

定义6串行域SEQ是一段#pragmaompparallel范围外的、内部无分支的、串行执行的代码片段，SEQ，其在串并行控制流图中也称作串行节点。

定义7并行域OMP是一段#pragmaompparallel范围内的、并行执行的代码片段，其在串并行控制流图中也称作并行节点。

定义8串并行控制流图SPGraph的定义如式(8)所示，是某函数的特殊的控制流图，其中的节点是串行域或并行域；串并行控制流图的示例如图3所示。

SPGraph＝(NodeGrp,CtlEdge)

NodeGrp＝SEQGrp∪OMPGrp

CtlEdge＝{<x,y>|x,y∈NodeGrp} (8)

SEQGrp＝{SEQ₁,…,SEQ_n}

OMPGrp＝{OMP₁,…,OMP_m}

定义9变量引用列表RefList的定义如式(9)所示，是某串行域或并行域中某变量的引用列表。

RefList{Ref₁,...,Ref_p}

定义10变量引用信息表NodeVarRef的定义如式(10)所示，是某串行域或并行域中所有变量信息的集合。每个串行域或并行域都与其对应的NodeVarRef绑定，如所示图3。

NodeVarRef＝{RefList₁,…,RefList_l} (10)

串行域或并行域中的函数调用需要特殊处理。串行域中的函数调用被单独分离成一个独立的串行域。对于传入拷贝的函数实参，其RefList＝{R}。对于传入指针或引用的函数实参，其RefList＝{R}，如果被调函数中没有写操作；或RefList＝{RW}，如果被调函数中有写操作。

对于并行域中的函数调用，其被看成是对函数实参的访问。对于传入拷贝的函数实参，在其RefList中适当位置插入{R}。对于传入指针或引用的函数实参，在其RefList中适当位置插入{R}，如果被调函数中没有写操作；或{RW}，如果被调函数中有写操作。

基于上述程序抽象表示，给出源到源翻译器三个主要功能的设计。

4.1数据传输API插入

如前所述，多数串行域和并行域之前需要插入所需的数据传输API。但是对于在并行域中被调用过的函数(简称OMP调用函数)，其中不能插入运行时API。因为OMP调用函数会在加速器上运行，而运行时API不能在加速器上运行。此类函数的数据一致性，依靠其函数调用前后的运行时API保证。

设计数据传输API插入算法如下。源到源编译器使用如下算法3，处理每个非OMP调用函数，在其中的串行域和并行域之前插入所需的数据传输API。算法3的01～02行是二重循环，将之后的算法步骤施加于控制流图中每个串行域和并行域中的每个被应用变量上；03行使用静态分析获得被引用变量Var的指针ptr。04～15行针对不同情况，分别处理。04行～05行在并行域前插入OAODataTrans(ptr,ConOMPR)语句。06～11行对于从函数调用分离得到的串行域分两种情况处理，如果被调函数是OMP调用函数，则在该串行域前插入OAODataTrans(ptr,ConSEQR)语句(07～08行)；否则对于传入拷贝的函数参数变量，在该串行域前插入OAODataTrans(ptr,ConSEQR)语句。15行对于其他情况，在该串行域前插入OAODataTrans(ptr,ConSEQR)语句。

4.2状态转换API插入

变量经过串行域或并行域访问后，一致性状态可能改变，需要插入所需的状态转换API更新运行时中保存的一致性状态。设计状态转换API插入算法如下。源到源编译器使用如下算法4，处理每个非OMP调用函数，在其中的串行域和并行域之后插入所需的状态转换API。算法4的01～02行是二重循环，将之后的算法步骤施加于控制流图中每个串行域和并行域中的每个被应用变量上；03行使用静态分析获得被引用变量Var的指针ptr。04～16行针对不同情况，分别处理。如果Node是并行域且变量Var对应的RefList中包含W(写操作)，则在该并行域之后插入OAOStTrans(ptr,TrOMPW)语句(04～06行)。如果Node是串行域且从函数调用分离而来，而且被调函数是OMP调用函数且变量Var对应的RefList中包含W(写操作)，则在该串行域之后插入OAOStTrans(ptr,TrSEQW)语句(08～11行)。如果不是以上两种情况，且变量Var对应的RefList中包含W(写操作)，则在该串行域之后插入OAOStTrans(ptr,TrSEQW)语句。

4.3并行指令翻译

本移植框架针对OpenMP工作共享并行模式，该模式在CPU和加速器上的平行指令对应关系如表6所示。源到源编译器使用如下算法5，将OpenMP CPU并行指令翻译成OpenMPOffloading并行指令，从而得到OpenMP Offloading计算内核。算法5的01行是循环，将之后的算法步骤施加于每一个并行域；02行通过静态分析获得OpenMP CPU并行指令；03～04行根据表6，将OpenMP CPU并行指令替换为对应的OpenMP Offloading并行指令。

表6并行指令对应关系

对本发明产生的技术效果进行如下说明：

1、源到源翻译结果示例

使用本专利提出的OAO自动移植系统翻译图4所示OpenMP CPU并行程序，得到如图6所示OpenMP Offloading异构并行程序。基于图4所示OpenMP CPU并行程序，人工编写OpenMP Offloading异构并行程序，如图5所示。观察可知，图5人工编写的程序中以下传输是冗余传输：05行对v1、v2、v3的“from”传输，13行对v3的“to”传输，13行对v4的“from”传输。图5自动翻译的程序使用OAO运行时库避免上述冗余操作。由此示例可知，OAO系统可以成功将OpenMP CPU并行程序翻译为OpenMP Offloading异构并行程序。

2、方法评估

2.1实验方法

Polybench和Rodinia是异构计算中常用的基准测试集。DawnCC是目前生成OpenMPOffloading程序的最先进的源到源翻译器。我们使用Polybench和Rodinia评估OAO系统的性能，并使用DawnCC作为对照。

我们为了对比DawnCC和OAO对过程间数据传输的优化能力，添加了一个测试程序FDTD-2D-FUNC。FDTD-2D-FUNC基于Polybench中的FDTD-2D；FDTD-2D中的每个计算kernel都被封装成一个子函数，kernel所需变量通过函数参数传递；这样就构造出了过程间数据传输。

首先我们生成如下4个版本的程序：

OMP版本：OpenMP CPU并行版本程序；

手动版本：手动翻译OMP版本得到的OpenMP Offloading程序；

DawnCC版本：使用DawnCC翻译OMP版本得到的OpenMP Offloading程序；

OAO版本：使用OAO翻译OMP版本得到的OpenMP Offloading程序。

两个实验软硬件平台如表7实验软硬件平台所示。

表7实验软硬件平台

2.2性能评估

如图7和图8所示是不同版本的OpenMP Offloading程序相对于OMP版本的加速比，图中的“X”表示DawnCC不能正确转换的程序。

可以看出OAO可以处理全部23个测试程序，而DawnCC只能处理15个测试程序。OAO版本在K40平台的9个和2080Ti平台的15个程序上有性能提升，最高加速比32倍；且OAO版本相对于OMP版本和手动版本在所有平台的所有测试程序上都有最好性能表现。

2.3数据传输优化评估

如表8所示是不同版本OpenMP Offloading程序数据传输次数对比，其中“—”表示DawnCC不能正确处理的程序。显然OAO版本在所有测试程序上都有最好的数据传输优化效果，即数据传输次数最少。

观察FDTD-2D和FDTD-2D-FUNC，OAO和DawnCC都能优化FDTD-2D；但是DawnCC不能很好优化FDTD-2D-FUNC；而OAO可以将二者的传输次数都优化为最优的5次。该现象说明OAO能够优化过程间数据传输，而DawnCC不能。

表8数据传输次数对比

如图9所示是OAO版本相对于手动版本和DawnCC版本所节省的传输数据量百分比。OAO版本在所有被测程序中都比手动版本传输了更少的数据，尤其是FDTD-2D-FUNC和FDTD-2D(接近100％)。相对于DawnCC版本，OAO版本在6个测试程序中有显著传输数据量节省，尤其是FDTD-2D-FUNC(接近100％)。

图10展示了OAO版本相对于手动版本和DawnCC版本所节省的传输时间百分比。相对于其他两个版本，OAO版本在所有测试程序上都有显著的传输时间节省。

上述数据和分析证明，OAO系统能够将OpenMP CPU并行程序自动翻译为OpenMPOffloading异构并行程序，同时对数据传输进行优化；与手动版本相比OAO版本取得了显著性能提升。与DawnCC相比，OAO系统能处理更多的输入程序，进行更广泛的数据传输优化，且在所有测试程序上都有更好的性能表现。

Claims

1.一种异构并行程序自动移植和优化方法，其特征在于，所述方法的实现过程为：

步骤1、构建异构并行程序自动移植系统的框架

源到源翻译器将OpenMP CPU并行代码翻译为OpenMP Offloading异构并行代码，并插入运行时API；其由3个模块组成：数据传输API插入模块、状态转换API插入模块、并行指令翻译模块；并行指令翻译模块将OpenMP CPU并行指令翻译为OpenMP Offloading异构并行指令，得到OpenMP Offloading内核；数据传输API插入模块、状态转换API插入模块分别插入两种对应的运行时API；

步骤2、一致性状态转换形式化

步骤3、运行时库设计

步骤4源到源翻译器设计

源到源翻译器主要基于Clang/LLVM编译框架的静态分析功能，将OpenMP CPU并行代码翻译为OpenMP Offloading异构并行代码，并插入运行时API，在保证数据一致性的前提下，优化数据传输；

源到源翻译器通过静态分析收集信息包括：串行域、并行域、变量引用；然后建立程序的串并行控制流图并将变量引用信息与相应的并行域串行域绑定；之后以串行域和并行域为粒度插入数据传输API插入、状态转换API插入；最后将OpenMP CPU并行指令翻译为OpenMP Offloading并行指令，具体过程为：

定义串行域SEQ是一段#pragma omp parallel范围外的、内部无分支的、串行执行的代码片段，SEQ，其在串并行控制流图中也称作串行节点；

定义并行域OMP是一段#pragma omp parallel范围内的、并行执行的代码片段，其在串并行控制流图中也称作并行节点；

定义串并行控制流图SPGraph的定义如式(1)所示，是某函数的特殊的控制流图，其中的节点是串行域或并行域；

SPGraph＝(NodeGrp,CtlEdge)

NodeGrp＝SEQGrp∪OMPGrp

CtlEdge＝{<x,y>|x,y∈NodeGrp}

SEQGrp＝{SEQ₁,…,SEQ_n}

OMPGrp＝{OMP₁,…,OMP_m}

(1)

定义变量引用列表RefList的定义如式(2)所示，是某串行域或并行域中某变量的引用列表，

定义变量引用信息表NodeVarRef的定义如式(3)所示，是某串行域或并行域中所有变量信息的集合，每个串行域或并行域都与其对应的NodeVarRef绑定，

NodeVarRef＝{RefList₁,…,RefList_l} (3)

串行域或并行域中的函数调用需要特殊处理，串行域中的函数调用被单独分离成一个独立的串行域，对于传入拷贝的函数实参，其RefList＝{R}；对于传入指针或引用的函数实参，其RefList＝{R}，当被调函数中没有写操作；或RefList＝{RW}，当被调函数中有写操作；对于并行域中的函数调用，其被看成是对函数实参的访问，对于传入拷贝的函数实参，在其RefList中插入{R}；对于传入指针或引用的函数实参，在其RefList中插入{R}，当被调函数中没有写操作；或{RW}，当被调函数中有写操作。

2.根据权利要求1所述的一种异构并行程序自动移植和优化方法，其特征在于，在步骤2中，一致性状态转换形式化的实现过程为：

给出如下定义：

定义1.一致性状态State是一个3bit的二进制数；其中Bit0＝1表示变量的加速器副本存在，Bit0＝0表示变量的加速器副本不存在；Bit1＝1表示变量的CPU副本有效，Bit1＝0表示变量的CPU副本无效；Bit2＝1表示变量的加速器副本有效，Bit2＝0表示变量的加速器副本无效；根据定义1，得到所有可能的一致性状态为：

HOST_ONLY：Bit2＝0；Bit1＝1；Bit0＝0；

DEVICE_NEW：Bit2＝0；Bit1＝1；Bit0＝1；

DEVICE_NEW：Bit2＝1；Bit1＝0；Bit0＝1；

SYNC：Bit2＝1；Bit1＝1；Bit0＝1；

CPU和加速器之间的数据传输以及CPU和加速器上的读写操作会改变一致性状态，使用一致性状态转换函数来表示一致性状态的改变，其定义如下：

定义2.一致性状态转换函数TransFunc的形式如式(4)所示，其中InVld和Vld是一对3bit的二进制数，其中运算符是布尔运算符；若将InVld中的某bit置0，则将inState中对应bit转换为0；若将Vld中的某bit置1则将inState中对应bit转换为1；因此给二者设置不同的值将任意inState转换为任意outState；TransFunc也简记为式(5)形式；

outState＝TransFunc(inState) (4)

＝inState·InVld+Vld

TransFunc＝{InVld,Vld} (5)

根据定义2，得到所有可能的一致性状态转换函数及其对应的数据传输或读写操作为：

数据传输或读写操作：TransFunc；

无需数据传输：TrNo＝{111,000}；

targetenterdatamap(alloc:…)：TrAlloc＝{111,001}；

targetenterdatamap(to:…)：TrEnTo＝{111,101}；

targetupdateto(…)TrUpTo＝{111,100}；

targetupdatefrom(…)：TrFrom＝{111,010}；

targetexitdatamap(delete:…)：TrDelete＝{010,010}；

CPU写操作：TrSEQW＝{011,010}；

加速器写操作TrOMPW＝{101,100}；

CPU或加速器读操作：TrRead＝{111,000}；

根据定义1和定义2，得到各种一致性状态之间的转换关系；

为保证程序正确，CPU和加速器上的读写操作对一致性状态有一定要求，使用一致性状态约束表示这些要求，其定义如下：

定义3.一致性状态约束Constr由一对3bit二进制数组成，其形式如(6)所示；Constr的默认值是{111,000}，表示对一致性状态无要求；将ConInVld的某bit置0，或将ConVld的某bit置1表示对一致性状态的不同约束要求为：

Constr＝{ConInVld,ConVld} (6)

不同bit及其值的一致性状态约束分别为，ConVldBit0＝1：加速器副本必须存在；ConVldBit1＝1：CPU副本必须有效；ConVldBit2＝1：加速器副本必须有效；ConInVldBit0＝0：加速器副本必须不存在；ConInVldBit＝10：CPU副本必须无效；ConInVldBit2＝0：加速器副本必须无效；

根据定义3和不同bit及其值的一致性状态约束，给出CPU和加速器上不同读写操作所需的一致性状态约束为：

所需的一致性状态约束，读写操作所需的一致性状态约束分别为：

变量没有被使用：ConNo＝{111,000}；

CPU读操作：ConSEQR＝{111,010}；

加速器读操作：ConOMPR＝{111,101}；

CPU写操作：ConSEQW＝{111,000}；

加速器写操作：ConOMPW＝{111,001}；

CPU free操作：ConFREE＝{111,101}；

由当前的一致性状态State和需要满足的一致性状态约束Constr＝{ConInVld,ConVld}，推导出最简状态转换函数如下：

最简状态转换函数的推导公式如式(7)所示：

MinTrFunc(State)＝State·InVld+Vld (7)

其中

MinTrFunc通过所有可能的一致性状态转换操作中的前6种对应关系确定需要执行的最简传输操作类型。

3.根据权利要求2所述的一种异构并行程序自动移植和优化方法，其特征在于，在步骤3中，运行时库设计的具体过程为：

运行时库API函数分成三类：一致性状态跟踪API函数、数据传输API函数OAODataTrans、一致性状态转换API函数OAOStTrans；运行时库API函数中前6个为一致性状态跟踪API函数为：

void OAOSaveArrayInfo(void*ptr,size_t length,size_t ElementSize)：保存数组内存区域信息；

void OAODeleteArrayInfo(void*ptr)：删除数组内存区域信息；

void*OAOMalloc(size_t length)：保存动态内存区域信息；

void OAOFree(void*ptr)：删除动态内存区域信息；

void*OAONewInfo(void*ptr,size_t ElementSize,size_t ElementNum)：保存动态内存区域信息；

void OAODeleteInfo(void*ptr)：删除动态内存区域信息；

void OAODataTrans(void*ptr,STATE_CONSTR Constr)：动态确定最简数据传输操作并执行，同时更新一致性状态；

void OAOStTrans(void*ptr,STATE_CONSTR StTrans)：更新一致性状态；

3.1一致性状态跟踪API

运行时库将变量内存区域作为一致性状态跟踪和数据传输的粒度，在C/C++中变量内存区域是一个连续的内存区域，其来源是局部变量定义、全局变量定义、malloc操作、new操作，为记录并跟踪变量一致性状态，定义变量内存区域和内存环境的形式化表示如下：

定义4.变量内存区域MemBlk是一个四元组，如式(8)所示；其中：Begin表示内存起始地址，Length表示内存区域长度，ElemSize表示元素大小，State表示一致性状态；

MemBlk＝{Begin,Length,ElemSize,State} (8)

定义5.内存环境MemEnv是全部变量内存区域的集合，如式(6)所示，

MemEnv＝{MemBlk₁,…,MemBlk_n} (9)

运行时库将MemEnv定义为全局变量，并在整个OpenMP Offloading异构并行程序执行期间维护MemEnv；当某变量被引用时，使用相应的指针ptr在MemEnv中搜索，满足式(10)的即为被引用的MemBlk：

Begin≤ptr≤Begin+Length-1 (10)

一致性状态跟踪API会在源到源翻译过程中插入源代码，OAOSaveArrayInfo函数被插入到局部变量声明之后或main函数开始处，用来收变量信息；OAOMalloc函数替换malloc函数，收集内存分配信息，并进行内存分配；OAONewInfo被插入到new操作之后，收集内存分配信息，以上三个函数利用收集到的信息新建对应MemBlk，并将State初始化为HOST_ONLY；

OAODeleteArrayInfo被插入到变量作用域结束处，main函数结束处，删除对应MemBlk；OAOFree替换free函数，释放内存区域同时删除对应MemBlk；OAODeleteInfo被插入到delete操作之后，释放内存区域同时删除对应MemBlk；

当需要分配的内存大于128KB时，使用cudaMallocHost()替代malloc()，用来分配内存；

3.2数据传输API

基于变量读写操作需要变量满足某种一致性状态约束，在访问某变量之前需要调用数据传输API，即OAODataTrans函数，来执行最简数据传输操作满足对应约束；OAODataTrans函数使用如下算法1，来确定为了满足一致性状态约束Constr，所需的最简状态转换函数，并执行对应的最简数据传输操作，同时更新一致性状态；

算法1为数据传输算法；输入：变量指针ptr、变量需要满足的一致性状态约约束Constr、变量内存环境集合MemEnv；对于变量内存环境集合MemEnv中的每个变量内存区域MemBlk，执行以下循环：

若ptr满足式(10)则执行：将Constr和MemBlk的State带入式(7)推导MinTrFunc；找到MinTrFunc对应的最简数据传输操作并执行；根据式(4)，使用MinTrFunc更新MemBlk的State；然后跳出此循环；

若ptr不满足式(10)则跳到过当前轮循环，进入下一轮循环；

循环结束；

3.3一致性状态转换API

在某变量被访问之后需要调用一致性状态转换API，即OAOStTrans函数，来更新变量内存区域MemBlk中保存的一致性状态State，OAOStTrans函数使用如下算法2完成一致性状态转换过程；

算法2为一致性状态转换算法，输入：变量指针ptr、状态转换函数StTrans、变量内存区域集合MemEnv；对于变量内存环境集合MemEnv中的每个变量内存区域MemBlk，执行以下循环：

若ptr满足式(10)则执行：根据式(4)，使用StTrans更新MemBlk的State；然后跳出此循环；

若ptr不满足式(10)则跳到过当前轮循环，进入下一轮循环；

循环结束。

4.根据权利要求3所述的一种异构并行程序自动移植和优化方法，其特征在于，基于权利要求3所述程序抽象表示，给出源到源翻译器三个主要功能的设计，具体为：

4.1数据传输API插入

设计数据传输API插入算法如下：源到源编译器使用如下算法3，处理每个非OMP调用函数，在其中的串行域和并行域之前插入所需的数据传输API；

算法3为数据传输API插入算法；输入：非OMP调用函数的串并行控制流图SPGraph；

对于集合NodeGrp中的每个Node，执行以下循环：

对于Node中引用的每个变量Var，执行以下循环：

使用静态分析获得Var的指针ptr；

若Node是并行域，执行：在Node前插入OAODataTrans(ptr,ConOMPR)；

若Node是串行域且Node从函数调用分离而来，执行：若被调函数是OMP调用函数，则在Node前插入OAODataTrans(ptr,ConSEQR)；否则，若传入的是的拷贝则在前插入OAODataTrans(ptr,ConSEQR)；

若Node是串行域且Node不是从函数调用分离而来，执行：在Node前插入OAODataTrans(ptr,ConSEQR)；

结束内层循环；

结束外层循环；

执行结束；

4.2状态转换API插入

设计状态转换API插入算法如下：源到源编译器使用如下算法4，处理每个非OMP调用函数，在其中的串行域和并行域之后插入所需的状态转换API；

算法4为状态转换API插入算法；输入：非OMP调用函数的串并行控制流图SPGraph；

对于集合NodeGrp中的每个Node，执行以下循环：

对于Node中引用的每个变量Var，执行以下循环：

使用静态分析获得Var的指针ptr；

若Node是并行域，则执行：

若Var对应的RefList中包含W，则在Node后插入OAOStTrans(ptr,TrOMPW)；

若Node是串行域且Node从函数调用分离而来，则执行：

若被调函数是OMP调用函数，则执行：若Var对应的RefList中包含W，则在Node后插入OAOStTrans(ptr,TrSEQW)；

若Node是串行域且Node不是从函数调用分离而来，则执行：

若Var对应的RefList中包含W，则在Node后插入OAOStTrans(ptr,TrSEQW)；

结束内层循环；

结束外层循环；

执行结束；

4.3并行指令翻译

本移植框架针对OpenMP工作共享并行模式，该模式在CPU和加速器上的平行指令对应关系根据并行指令对应关系，源到源编译器使用如下算法5，将OpenMP CPU并行指令翻译成OpenMP Offloading并行指令，从而得到OpenMP Offloading计算内核；

将#pragma omp parallel for替换为#pragma omp target teams distributeparallel for；

将#pragma omp parallel loop替换为#pragma omp target teams distributeparallel loop；

将#pragma omp parallel simd替换为#pragma omp target teams distributeparallel simd；

算法5为并行指令翻译算法；输入：串并行控制流图SPGraph；

对于SPGraph的OMPGrp集合中的每个OMP，执行以下循环：使用静态分析获得OpenMPCPU并行指令；根据并行指令对应关系，将OpenMP CPU并行指令替换为对应的OpenMPOffloading并行指令；循环结束。