CN112083956A - 一种面向异构平台的复杂指针数据结构自动管理*** - Google Patents

一种面向异构平台的复杂指针数据结构自动管理*** Download PDF

Info

Publication number
CN112083956A
CN112083956A CN202010971038.9A CN202010971038A CN112083956A CN 112083956 A CN112083956 A CN 112083956A CN 202010971038 A CN202010971038 A CN 202010971038A CN 112083956 A CN112083956 A CN 112083956A
Authority
CN
China
Prior art keywords
information
serial
parallel
node
class
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Granted
Application number
CN202010971038.9A
Other languages
English (en)
Other versions
CN112083956B (zh
Inventor
张伟哲
何慧
王法瑞
方滨兴
郝萌
郭浩男
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Harbin Institute of Technology
Original Assignee
Harbin Institute of Technology
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Harbin Institute of Technology filed Critical Harbin Institute of Technology
Priority to CN202010971038.9A priority Critical patent/CN112083956B/zh
Publication of CN112083956A publication Critical patent/CN112083956A/zh
Application granted granted Critical
Publication of CN112083956B publication Critical patent/CN112083956B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30145—Instruction analysis, e.g. decoding, instruction word fields
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G06F8/42—Syntactic analysis
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G06F8/42—Syntactic analysis
    • G06F8/425—Lexical analysis

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • General Engineering & Computer Science (AREA)
  • Software Systems (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Stored Programmes (AREA)
  • Devices For Executing Special Programs (AREA)

Abstract

一种面向异构平台的复杂指针数据结构自动管理***，涉及异构编程技术领域。本发明为了实现OpenMP Offloading程序中复杂指针数据结构在异构计算平台上的自动管理并保证数据一致性。本发明包括：信息收集模块，用于对源程序的静态分析及程序信息收集；自动转换模块，主要负责根据变量类型的不同，在适当位置修改源码并***合适的运行时API；运行时模块，主要负责使用cudaMallocManaged( )和cudaFree( )重新实现C++标准的内存管理操作并向外提供接口。本发明可以自动管理OpenMP Offloading程序中复杂指针数据结构在CPU和GPU内存之间的内存分配、释放以及数据传输，并保证数据一致性；从而为OpenMP Offloading程序开发提供便利。

Description

一种面向异构平台的复杂指针数据结构自动管理***
技术领域
本发明涉及OpenMP Offloading程序中复杂指针数据结构自动管理***，涉及异构编程技术领域。
背景技术
OpenMP是由OpenMP Architecture Review Board牵头提出的，并已被广泛接受，用于共享内存并行***的多处理器程序设计的一套指导性编译处理方案(CompilerDirective)[1]。OpenMP支持的编程语言包括C、C++和Fortran；而支持OpenMP的编译器包括Sun Compiler，GNU Compiler和Intel Compiler等。OpenMP提供了对并行算法的高层的抽象描述，程序员通过在源代码中加入专用的pragma原语来指明自己的意图，由此编译器可以自动将程序进行并行化，并在必要之处加入同步互斥以及通信。
在高性能计算领域，各种加速器(如GPU、FPGA、DSP、MLU等)已经成为除CPU之外的重要算力来源。从4.0版本开始，OpenMP增加了Offloading(卸载)特性，OpenMP Offloading支持CPU+加速器的异构编程模型；经过4.5和5.0版本的发展，OpenMP Offloading逐步完善。OpenMP Offloading为OpenMP程序充分利用异构计算平台算力提供了可能，但将现有的OpenMP CPU程序修改为符合Offloading特性的程序仍然是一件困难、繁琐、易出错的工作，尤其是程序中包含复杂指针数据结构，例如：类嵌套指针、vector容器、多级嵌套指针等。
虽然OpenMP Offloading语法简单，但仍需用户利用相关pragma原语显示管理CPU和加速器之间的数据传输操作，这给开发者造成了极大不便，尤其是遇到复杂指针数据结构时。例如C++中的vector容器，其内存分配是隐式的，开发者不易控制内存分配和数据传输，也就很难使用Offloading特性。对于嵌套指针或多级嵌套指针，处理不同级别指针指向的内存在CPU和加速器内存空间中的分配、释放、数据传输是繁琐且极易出错的工作，这也使开发者对Offloading特性望而却步。
NVIDIA GPU平台的CUDA编程模型从6.0版本开始支持统一内存(Unified Memory,UM)特性；该特性将CPU和GPU的地址空间统一，并自动管理CPU和GPU之间的数据传输。统一内存特性为OpenMP Offloading程序开发中复杂指针数据结构的自动化处理提供了可能的技术手段。
发明内容
本发明要解决的技术问题为：
本发明的目的是提出一种面向异构平台的OpenMP Offloading程序中复杂指针数据结构自动管理***，以解决现有技术在OpenMP CPU程序的基础上无法自动修改内存分配和释放语句并自动***相关pragma原语，不能自动管理CPU和加速器之间的数据传输，无法保证包含复杂指针数据结构的程序在异构计算平台上的数据一致性，从而影响程序性能的问题。
本发明为解决上述技术问题所采用的技术方案为：
一种面向异构平台的复杂指针数据结构自动管理***，所述***用于实现OpenMPOffloading程序中复杂指针数据结构在异构计算平台上的自动管理并保证数据一致性；
该***包括以下三个模块：
信息收集模块，该模块有两项功能：1)对源程序进行静态分析以收集程序信息；2)基于收集到的信息对源程序建立抽象表示即串并行控制流图；
该模块的工作过程分为以下两步：1)通过Clang编译器，将C或者C++的源代码生成对应的抽象语法树(Abstract Syntax Tree,AST)和控制流图(Control Flow Graph,CFG)，遍历AST和CFG，区分串并行域，并获取程序详细信息；2)根据这些信息生成串并行控制流图；
自动转换模块，主要负责基于串并行控制流图，在源码中***运行时API，从而完成代码转换；首先根据串并行控制流图中保存的复杂指针变量信息，确定复杂指针数据的类型；之后根据不同类型，在源码中适当位置，***合适的运行时API，完成代码转换；这样，复杂指针变量所涉及的内存分配、释放以及数据传输操作都由运行时接管，从而可以通过运行时自动管理复杂指针变量，保证CPU和GPU之间的数据一致性；
运行时模块，主要负责基于统一内存实现复杂指针数据类型的以下操作：CPU和GPU上的内存分配和释放操作，以及CPU和GPU之间的自动数据传输操作；运行时模块由UMMaper类和allocator类组成，UMMaper类和allocator类以API接口的形式向外部提供内存分配和释放的接口。
进一步地，所述信息收集模块的信息收集功能的实现过程为：
首先使用Clang编译器对目标程序进行词法分析和语义分析，并生成抽象语法树(Abstract Syntax Tree,AST)和控制流图(Control Flow Graph,CFG)；
之后对AST和CFG进行下述三种静态分析，获得函数调用关系信息、变量相关信息、串/并行域相关信息，为之后的串并行控制流图建立和代码转换提供信息支持；
分析一，函数调用关系分析：对于一个AST，执行以下两步工作：
1)递归遍历AST上的每个节点；
2)如果当前节点是函数定义节点，则保存该函数名和该函数调用的子函数信息；
分析二，变量信息分析，对于一个AST，执行以下三步工作：
1)递归遍历AST上的每个节点；
2)如果当前节点涉及变量定义或引用，则保存变量定义信息、变量引用信息、变量作用域信息；
3)如果当前节点涉及内存分配或释放，则保存内存分配信息、内存释放信息；
分析三，串并行域分析：对于一个CFG，执行以下三步工作：
1)递归遍历CFG上的每个节点，保存节点间关系信息；
2)如果当前节点在OpenMP#pragma并行指导语句的作用域内，则将当前节点标记为并行节点，并保存其类型信息和范围信息；
3)如果当前节点不在OpenMP#pragma并行指导语句的作用域内，则将当前节点标记为串行节点，并保存其类型信息和范围信息；
通过上述三种静态分析，可以获得以下信息：函数调用关系信息、变量定义信息、变量引用信息、变量作用域信息、内存分配信息、内存释放信息、串并行节点类型信息、串并行节点范围信息、节点间关系信息；这些信息将对串并行控制流图建立提供支持。
进一步地，所述信息收集模块中的串并行控制流图定义如下：
定义串并行控制流图由串行节点、并行节点以及节点之间的有向边组成；其中，串行节点表示一段OpenMP#pragma并行指导语句作用域外的、内部无分支的、串行执行的代码片段；串行节点对应的代码片段在CPU上执行，串行节点也记作SEQ节点；
并行节点表示一段在OpenMP#pragma并行指导语句作用域内的，并行执行的代码片段；并行节点对应的代码片段被卸载到GPU上执行，并行节点也记作OMP节点；
串行节点和并行节点中保存有函数调用信息和变量相关信息；
节点之间的有向边表示节点对应的代码片段执行的先后关系。
进一步地，所述信息收集模块中串并行控制流图建立过程：
串并行控制流图建立过程以函数为基本处理单位；对整个源程序而言，如果可以建立一个函数的串并行控制流图，结合所收集的函数调用关系信息，就可以递归建立整个源程序的串并行控制流图；
对于一个函数，基于所收集的信息，可以通过以下步骤建立串并行控制流图：
1)使用节点类型信息、节点范围信息建立一个个孤立的串行节点和并行节点；
2)使用节点间关系信息，在节点间建立有向边，将串行节点和并行节点连成图；
3)根据节点范围信息，将函数调用信息、变量定义信息、变量引用信息、变量作用域信息、内存分配信息、内存释放信息保存到对应的串行节点或并行节点。
进一步地，将所述复杂指针数据按指针所处的位置分为三种类型：
类嵌套指针、vector容器、多级嵌套指针；
类嵌套指针是指类中所包含的指针，vector容器是指C++标准库中提供的vector容器；多级嵌套指针是指二级指针以及二级以上的指针。
进一步地，所述类嵌套指针通过以下步骤进行处理：
①递归遍历上述所建立的串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的类嵌套指针及其所在的C++类；
②对于步骤①中找出的C++类，修改源码中该类的类型定义，使该类的类型在定义时继承运行时提供的UMMaper基类；
③对于步骤①中找出的类嵌套指针，修改源码中该指针的内存分配和释放语句，使用cudaMallocManaged()分配内存，使用cudaFree()释放内存；
④对于步骤①中找出的C++类实例，修改源码中该类实例的定义语句，使用被重载的new运算符创建类实例，并将步骤③中分配的内存空间地址传递给C++类实例中对应的嵌套指针。
进一步地，对所述vector容器通过以下步骤进行处理：
①递归遍历上述所建立的串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的vector容器；
②对于步骤①中找出的vector容器实例，修改该vector容器实例的定义语句，***对运行时提供的自定义allocator的显示调用。
进一步地，对所述多级嵌套指针通过以下步骤进行处理：
①递归遍历上述建立的串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的多级嵌套指针，及其各级子指针；
②对于步骤①中找出的所有多级嵌套指针和各级子指针，修改源码中这些指针的内存分配和释放语句，使用cudaMallocManaged()分配内存，使用cudaFree()释放内存。
进一步地，所述运行时模块中UMMaper类的实现过程：为处理类嵌套指针，设计UMMaper类来管理C++类的内存分配和释放；UMMaper类中使用cudaMallocManaged()和cudaFree()对C++默认的new、delete运算符进行重载；UMMaper类的派生类在CPU和GPU上的内存空间的分配、释放、数据传输操作即可被统一内存自动管理。
进一步地，所述运行时模块中自定义allocator类的实现过程：为处理vector容器，设计自定义的allocator类来管理vector容器的内存分配和释放；vector容器默认使用C++标准库的allocator空间配置器来管理内存分配和释放，因此可以基于统一内存实现自定义的allocator类，从而实现对vector容器内存分配和释放的自动管理；自定义allocator类中，基于cudaMallocManaged()实现allocator类中的_allocator()函数，基于cudaFree()实现allocator类中的deallocate()函数；在vector容器声明时显示调用自定义allocator类，即可实现统一内存对vector容器在CPU和GPU上的内存空间的分配、释放、数据传输操作的自动管理。
本发明具有以下有益技术效果：
本发明所述***可以在OpenMP CPU程序的基础上，自动修改内存分配和释放语句，并自动***相关pragma原语，自动管理CPU和加速器之间的数据传输，以保证包含复杂指针数据结构的程序在异构计算平台上的数据一致性，并提高程序性能。
本发明研究的异构编程方案主要针对OpenMP程序中复杂指针数据结构的自动管理，具体指自动修改内存分配和释放语句，自动***相关pragma原语，自动管理CPU和加速器之间的数据传输；从而保证程序在异构计算平台上的数据一致性，并提高程序性能。
由于OpenMP Offloading支持加速器编程，使得CPU上的OpenMP代码可以卸载到GPU上执行，客观条件允许OpenMP代码的卸载。在此条件下，将OpenMP代码卸载到GPU上运行，一方面可以提升程序的运行效率，另一方面可以充分利用GPU的加速效果。但手动进行代码卸载既无法保证转换程序的正确性，同时又耗费人力物力，效率很低。因此本发明提出复杂指针数据结构自动管理方案，以解决在卸载过程中最困难的复杂指针数据结构在CPU和加速器之间的内存分配和数据传输等问题。
通过在通用测试集(PolyBench，Rodinia等)上的对比实验，证明本发明可以自动管理OpenMP Offloading程序中复杂指针数据结构，保证程序的正确性，并提高程序性能。
本发明所述的复杂指针数据结构是指复杂指针数据类型。
附图说明
图1是本发明所述***的整体框架图；图2是AST分析方法的结构框图；图3为函数串并行控制流图；图4是统一内存原理图；图5为处理类嵌套指针的自动卸载实现原理图(是否使用统一内存的程序对比)；图6为处理vector容器的自动卸载实现原理图(是否使用空间配置器的程序对比)；图7为处理多级嵌套指针的自动卸载实现原理图；图8为本发明实验设计整体框架图；
图9为2080Ti平台Large数据量下的运行时间柱状图，图中：图9(a)为Large-2080Ti运行时间对比图(数据集a)，图9(b)为Large-2080Ti运行时间对比图(数据集b)，图9(a)和图9(b)本质是一个图，因为数据集比较多，所以分成两个图；
图10为Rodinia测试集在2080Ti上的运行时间比较图；图11为Polybench-OAO不同数据量加速比对比图(2080Ti)；图12为Polybench加速比对比图(2080Ti)；图13为运行时开销图(2080Ti)；图14为复杂指针数据结构详细测试图(K40)；
图15为含有vector的程序转换示例；图16为含有structnest的程序程序转换示例；图17为含有multilevel的程序转换示例。图15至图17中：vector为c++中的vector容器；structnest为类嵌套指针；multilevel为多级嵌套指针。
具体实施方式
结合附图1至17，对本发明所述的一种面向异构平台的复杂指针数据结构自动管理***进行如下阐述：
本发明的主要任务是自动管理OpenMP Offloading程序中的复杂指针数据结构(类嵌套指针、vector容器、多级嵌套指针等)，即实现自动修改分配和释放语句，自动管理数据传输，保证数据一致性。本发明主要包括以下三个模块：
信息收集模块，该模块有两项功能：1)对源程序进行静态分析以收集程序信息；2)基于收集到的信息对源程序建立抽象表示即串并行控制流图。该模块的工作过程分为以下两步：1)通过Clang编译器，将C或者C++的源代码生成对应的抽象语法树(Abstract SyntaxTree,AST)和控制流图(Control Flow Graph,CFG)，遍历AST和CFG，区分串并行域，并获取程序详细信息；2)根据这些信息生成串并行控制流图。
自动转换模块，主要负责基于串并行控制流图，在源码中***运行时API，从而完成代码转换。首先根据串并行控制流图中保存的复杂指针变量信息，确定复杂指针变量的类型。之后根据不同类型，在源码中适当位置，***合适的运行时API，完成代码转换。这样，复杂指针变量所涉及的内存分配、释放以及数据传输操作都由运行时接管，从而可以通过运行时自动管理复杂指针变量，保证CPU和GPU之间的数据一致性。
运行时模块，主要负责基于统一内存实现复杂指针数据结构的以下操作：CPU和GPU上的内存分配和释放操作，CPU和GPU之间的自动数据传输操作；其中使用cudaMallocManaged()和cudaFree()重新实现new、delete、allocator这些C++默认接口，并以API接口的形式向外部提供内存分配和释放的接口。
结果会将源程序转换为带有运行时API的新程序，***的整体框架如图1所示。
1信息收集模块
信息收集模块实现两个功能：静态分析并收集程序信息；建立串并行控制流图。
1.1信息收集
首先使用Clang编译器对目标程序进行词法分析和语义分析，并生成抽象语法树(Abstract Syntax Tree,AST)和控制流图(Control Flow Graph,CFG)。
之后对AST和CFG进行下述三种静态分析，获得函数调用关系信息、变量相关信息、串/并行域相关信息，为之后的串并行控制流图建立和代码转换提供信息支持。
分析一，函数调用关系分析：对于一个AST，执行以下两步工作：
1)递归遍历AST上的每个节点；
2)如果当前节点是函数定义节点，则保存该函数名和该函数调用的子函数信息。
分析二，变量信息分析，对于一个AST，执行以下三步工作：
1)递归遍历AST上的每个节点；
2)如果当前节点涉及变量定义或引用，则保存变量定义信息、变量引用信息、变量作用域信息；
3)如果当前节点涉及内存分配或释放，则保存内存分配信息、内存释放信息。
分析三，串并行域分析：对于一个CFG，执行以下三步工作：
1)递归遍历CFG上的每个节点，保存节点间关系信息；
2)如果当前节点在OpenMP#pragma并行指导语句的作用域内，则将当前节点标记为并行节点，并保存其类型信息和范围信息；
3)如果当前节点不在OpenMP#pragma并行指导语句的作用域内，则将当前节点标记为串行节点，并保存其类型信息和范围信息。
通过上述三种静态分析，可以获得以下信息：函数调用关系信息、变量定义信息、变量引用信息、变量作用域信息、内存分配信息、内存释放信息、串并行节点类型信息、串并行节点范围信息、节点间关系信息。这些信息将对串并行控制流图建立提供支持。
1.2串并行控制流图建立
定义串并行控制流图由串行节点、并行节点以及节点之间的有向边组成。其中，串行节点表示一段OpenMP#pragma并行指导语句作用域外的、内部无分支的、串行执行的代码片段；串行节点对应的代码片段在CPU上执行，串行节点也记作SEQ节点。
并行节点表示一段在OpenMP#pragma并行指导语句作用域内的，并行执行的代码片段；并行节点对应的代码片段被卸载到GPU上执行，并行节点也记作OMP节点。串行节点和并行节点中保存有函数调用信息和变量相关信息。节点之间的有向边表示节点对应的代码片段执行的先后关系。
串并行控制流图建立过程以函数为基本处理单位；对整个源程序而言，如果可以建立一个函数的串并行控制流图，结合函数调用关系信息，就可以递归建立整个源程序的串并行控制流图。
对于一个函数，基于收集的程序信息，可以通过以下步骤建立串并行控制流图：
1)使用节点类型信息、节点范围信息建立一个个孤立的串行节点和并行节点；
2)使用节点间关系信息，在节点间建立有向边，将串行节点和并行节点连成图；
3)根据节点范围信息，将函数调用信息、变量定义信息、变量引用信息、变量作用域信息、内存分配信息、内存释放信息保存到对应的串行节点或并行节点。
2自动转换模块
自动转换模块将复杂指针数据类型分为：类嵌套指针、vector容器、多级嵌套指针三种类型。
复杂指针数据结构处理中的关键技术是统一内存(Unified Memory，UM)，故首先介绍统一内存的概念以及原理，UM维护了一个统一的内存池，在CPU与GPU中共享，仅仅只分配一次内存，这个数据指针对于host端(CPU)和device端(GPU)都是可用的。使用了单一指针进行托管内存，这样在不同设备之间进行的数据传输由UM的运行时***来自动完成，并使GPU可以处理超出其内存容量的数据集。统一内存原理如图4所示。
利用统一内存技术，针对三种复杂指针数据类型，分别设计不同处理方法。
2.1类嵌套指针
类嵌套指针是指类中所包含的指针，如图5所示，通过以下步骤进行处理：
①递归遍历串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的类嵌套指针及其所在的类。
③对于步骤①中找出的类的类型，修改源码中该类的类型定义，使该类的类型在定义时继承运行时提供的UMMaper基类(见运行时设计)；
③对于步骤①中找出的类嵌套指针，修改源码中该指针的内存分配和释放语句，使用cudaMallocManaged()分配内存，使用cudaFree()释放内存；
④对于步骤①中找出的类实例，修改源码中该类实例的定义语句，使用被重载的new运算符创建类实例，并将步骤③中分配的内存空间地址传递给类中对应的嵌套指针。
类嵌套指针的处理示例如图5所示。
2.2vector容器
vector容器是指C++标准库中提供的vector容器。vector容器需要通过以下步骤进行处理：
①递归遍历串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的vector容器。
②对于步骤①中找出的vector容器实例，修改该vector容器实例的定义语句，***对运行时提供的自定义allocator(见运行时设计)的显示调用。
vector容器处理示例如图6所示。
2.3多级嵌套指针
多级嵌套指针是指二级指针以及二级以上的指针。多级嵌套指针需要通过以下步骤进行处理：
①递归遍历串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的多级嵌套指针，及其各级子指针；
②对于步骤①中找出的所有多级嵌套指针和各级子指针，修改源码中这些指针的内存分配和释放语句，使用cudaMallocManaged()分配内存，使用cudaFree()释放内存。
多级嵌套指针的处理示例如图7所示。
3运行时
为了使转换后的代码更加接近源代码的格式，尽量减少对源码的修改，本文提出了在原程序中***少量运行时API的方案，将基于统一内存的内存分配都封装到运行时中。
对于类及其嵌套指针，在运行时中实现负责内存分配和释放的UMMapper基类，如图5所示。在UMMapper中使用cudaMallocManaged()和cudaFree()对C++默认的new、delete运算符进行重载。这样只需在类声明时基于UMMapper派生，再使用new来新建类实例就可以实现通过统一内存自动管理派生类相关的内存分配、释放、数据传输操作。
vector容器使用C++标准库的allocator空间配置器来管理内存分配和释放，因此可以基于统一内存实现自定义的allocator类，从而实现对vector容器内存分配和释放的自动管理。自定义allocator类中，基于cudaMallocManaged()实现allocator类中的_allocator()函数，基于cudaFree()实现allocator类中的deallocate()函数。在vector容器声明时显示调用自定义allocator类，即可实现统一内存对vector容器在CPU和GPU上的内存空间的分配、释放、数据传输操作的自动管理，如图6所示。
对发明的技术效果进行如下验证：
为了验证和分析代码卸载方案的效果，我们在RTX 2080Ti平台上进行测验，使用PolyBench，Rodinia等基准测试数据集，将测试结果和另一较为著名的源到源编译器(DawnCC)以及手动转换(Manual)的结果还有CPU并行的代码进行比较。我们的方案命名为OAO(OpenMP Automated Offloading with complex data structure support)。
实验结果的测试和获取需要整合运行时，编写脚本运行，Python提取以及OriginLab分析这四个阶段。实验整体过程如图8所示。
1运行时间
2080Ti平台Large数据量下的运行时间分别如图9所示。
Rodinia测试集在2080Ti上的运行时间比较，如图10所示。
通过以上对数据集运行时间的收集和比对，可以得出以下结论：
(1)通过运行时间图可以看出，OAO可以处理全部23个测试程序，而DawnCC只能处理15个测试程序。OAO在K40平台的9个和2080Ti平台的15个程序上有性能提升；且OAO相对于手动转换在所有平台的所有测试程序上都有最好性能表现。
(2)在Rodinia-2080Ti中，已知Rodinia数据体量由实验前测试获得，共8个数据集中，有5个数据集的OAO运行时间比OMP更短。
(3)在Polybench-Large中，OAO的运行时间优于DawnCC，Manual和native；运行时间总体上呈现OAO<DawnCC<Manual<native的趋势。其中，native运行时间最长，是因为选择的传输语句问题，没有进行并行优化；Manual包含了部分数据的冗余传输，运行时间比OAO要长；而OAO在数据集的处理上总体也要优于DawnCC。
2加速比
由图11可以看出，在加速比方面，在大于1的数据中，Large数据量加速比比Medium更大，这也正符合并行加速的特点，处理的数据量越大，并行加速效果越明显。
在加速比方面，计算方式为OMP/Type的比值，Type为其他类型，结果以1X为界限，上方为正向加速，下方为负向加速。由图12可以得出，不管在哪个平台上，所有测试程序的结果显示，OAO的加速比都优于Manual，在Polybench测试程序中，OAO略优于DawnCC，且由于DawnCC无法处理Rodinia测试程序，所以OAO在所有测试程序上都有最好的数据传输优化效果，且能处理DawnCC处理不了的复杂测试程序，适用范围更广。
3运行时开销
在Polybench测试集中，运行时的开销如图13所示。
所有测试程序的运行时间，与运行时开销相比，运行时开销都要小2个数量级以上，所以OAO在运行时的额外开销很小，几乎可以忽略，所以加入运行时对源程序在性能上并没有损失。
4统一内存
统一内存的详细测试包括vector容器，类嵌套指针(structnest)和多级嵌套指针(multilevel)。复杂指针数据结构的运行时间和数据传输时间对比如图14所示。
从数据传输时间来看，H2D和D2H的时间远小于总运行时间，且全部复杂指针数据结构都可以正确处理，因此OAO编译器可以正确支持复杂指针数据结构的GPU卸载。
其中vector容器，structnest和multilevel的简单转换程序对比分别如图15、图16和图17所示。
其中由图15所示，声明阶段vector变量处修改为QY::allocator方式重载；图16中，声明阶段类内部嵌套指针x，y，z由cudaMallocManaged重载，而实例化的类a由传输语句标明，添加map(tofrom:a[0:1])；图17中，x为二级指针，y为x的一级指针，声明阶段x，y由cudaMallocManaged重载。
表6统一内存测试结果表
Figure BDA0002684028360000111
统一内存测试结果如表6所示。在数据传输过程中，会在正式传输之前，优先进行一次传输，H2D的数据大小为4B，D2H的数据大小为1B(4+1模式)，上表中vector的双向传输，52比49大3，structnest的双向传输，36比33大3，都是由于这个原因。除此之外，vector只传输2次，分别为X和Y，每次的数据量大小为24B，structnest传输1次，传输A，数据量大小为32。
而multilevel中，z是单向传输，使用的数据传输模型，x和y是多级嵌套指针，使用统一内存。除去优先传输的4+1模式，在H2D和D2H之间数据传输量存在差别的原因是，z并不回传，而x会回传。
实验结果证明OAO可以利用统一内存来正确处理复杂指针数据结构，实现对复杂指针数据结构处理的支持。

Claims (10)

1.一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，所述***用于实现OpenMP Offloading程序中复杂指针数据结构在异构计算平台上的自动管理并保证数据一致性；
该***包括以下三个模块：
信息收集模块，该模块有两项功能：1)对源程序进行静态分析以收集程序信息；2)基于收集到的信息对源程序建立抽象表示即串并行控制流图；
该模块的工作过程分为以下两步：1)通过Clang编译器，将C或者C++的源代码生成对应的抽象语法树AST和控制流图CFG，遍历AST和CFG，区分串并行域，并获取程序详细信息；2)根据这些信息生成串并行控制流图；
自动转换模块，主要负责基于串并行控制流图，在源码中***运行时API，从而完成代码转换；首先根据串并行控制流图中保存的复杂指针变量信息，确定复杂指针数据的类型；之后根据不同类型，在源码中适当位置，***合适的运行时API，完成代码转换；这样，复杂指针变量所涉及的内存分配、释放以及数据传输操作都由运行时接管，从而可以通过运行时自动管理复杂指针变量，保证CPU和GPU之间的数据一致性；
运行时模块，主要负责基于统一内存实现复杂指针数据类型的以下操作：CPU和GPU上的内存分配和释放操作，以及CPU和GPU之间的自动数据传输操作；运行时模块由UMMaper类和allocator类组成，UMMaper类和allocator类以API接口的形式向外部提供内存分配和释放的接口。
2.根据权利要求1所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，所述信息收集模块的信息收集功能的实现过程为：
首先使用Clang编译器对目标程序进行词法分析和语义分析，并生成抽象语法树AST和控制流图CFG；
之后对AST和CFG进行下述三种静态分析，获得函数调用关系信息、变量相关信息、串/并行域相关信息，为之后的串并行控制流图建立和代码转换提供信息支持；
分析一，函数调用关系分析：对于一个AST，执行以下两步工作：
1)递归遍历AST上的每个节点；
2)如果当前节点是函数定义节点，则保存该函数名和该函数调用的子函数信息；
分析二，变量信息分析，对于一个AST，执行以下三步工作：
1)递归遍历AST上的每个节点；
2)如果当前节点涉及变量定义或引用，则保存变量定义信息、变量引用信息、变量作用域信息；
3)如果当前节点涉及内存分配或释放，则保存内存分配信息、内存释放信息；
分析三，串并行域分析：对于一个CFG，执行以下三步工作：
1)递归遍历CFG上的每个节点，保存节点间关系信息；
2)如果当前节点在OpenMP#pragma并行指导语句的作用域内，则将当前节点标记为并行节点，并保存其类型信息和范围信息；
3)如果当前节点不在OpenMP#pragma并行指导语句的作用域内，则将当前节点标记为串行节点，并保存其类型信息和范围信息；
通过上述三种静态分析，可以获得以下信息：函数调用关系信息、变量定义信息、变量引用信息、变量作用域信息、内存分配信息、内存释放信息、串并行节点类型信息、串并行节点范围信息、节点间关系信息；这些信息将对串并行控制流图建立提供支持。
3.根据权利要求1或2所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，所述信息收集模块中的串并行控制流图定义如下：
定义串并行控制流图由串行节点、并行节点以及节点之间的有向边组成；其中，串行节点表示一段OpenMP#pragma并行指导语句作用域外的、内部无分支的、串行执行的代码片段；串行节点对应的代码片段在CPU上执行，串行节点也记作SEQ节点；
并行节点表示一段在OpenMP#pragma并行指导语句作用域内的，并行执行的代码片段；并行节点对应的代码片段被卸载到GPU上执行，并行节点也记作OMP节点；
串行节点和并行节点中保存有函数调用信息和变量相关信息；
节点之间的有向边表示节点对应的代码片段执行的先后关系。
4.根据权利要求3所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，所述信息收集模块中串并行控制流图建立过程：
串并行控制流图建立过程以函数为基本处理单位；对整个源程序而言，如果可以建立一个函数的串并行控制流图，结合权利要求2中收集的函数调用关系信息，就可以递归建立整个源程序的串并行控制流图；
对于一个函数，基于权利要求2中收集的信息，可以通过以下步骤建立串并行控制流图：
1)使用节点类型信息、节点范围信息建立一个个孤立的串行节点和并行节点；
2)使用节点间关系信息，在节点间建立有向边，将串行节点和并行节点连成图；
3)根据节点范围信息，将函数调用信息、变量定义信息、变量引用信息、变量作用域信息、内存分配信息、内存释放信息保存到对应的串行节点或并行节点。
5.根据权利要求4所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，将所述复杂指针数据按指针所处的位置分为三种类型：
类嵌套指针、vector容器、多级嵌套指针；
类嵌套指针是指类中所包含的指针，vector容器是指C++标准库中提供的vector容器；多级嵌套指针是指二级指针以及二级以上的指针。
6.根据权利要求5所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，所述类嵌套指针通过以下步骤进行处理：
①递归遍历权利要求4中所建立的串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的类嵌套指针及其所在的C++类；
②对于步骤①中找出的C++类，修改源码中该类的类型定义，使该类的类型在定义时继承运行时提供的UMMaper基类；
③对于步骤①中找出的类嵌套指针，修改源码中该指针的内存分配和释放语句，使用cudaMallocManaged()分配内存，使用cudaFree()释放内存；
④对于步骤①中找出的C++类实例，修改源码中该类实例的定义语句，使用被重载的new运算符创建类实例，并将步骤③中分配的内存空间地址传递给C++类实例中对应的嵌套指针。
7.根据权利要求5所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，对所述vector容器通过以下步骤进行处理：
①递归遍历权利要求4中建立的串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的vector容器；
②对于步骤①中找出的vector容器实例，修改该vector容器实例的定义语句，***对运行时提供的自定义allocator的显示调用。
8.根据权利要求5所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，对所述多级嵌套指针通过以下步骤进行处理：
①递归遍历权利要求4中建立的串并行控制流图，根据串/并行节点中保存的变量定义和引用信息，找出在串行节点中定义、并在并行节点中被引用的多级嵌套指针，及其各级子指针；
②对于步骤①中找出的所有多级嵌套指针和各级子指针，修改源码中这些指针的内存分配和释放语句，使用cudaMallocManaged()分配内存，使用cudaFree()释放内存。
9.根据权利要求6所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，所述运行时模块中UMMaper类的实现过程：
为处理类嵌套指针，设计UMMaper类来管理C++类的内存分配和释放；UMMaper类中使用cudaMallocManaged()和cudaFree()对C++默认的new、delete运算符进行重载；UMMaper类的派生类在CPU和GPU上的内存空间的分配、释放、数据传输操作即可被统一内存自动管理。
10.根据权利要求7所述的一种面向异构平台的复杂指针数据结构自动管理***，其特征在于，所述运行时模块中自定义allocator类的实现过程：
为处理vector容器，设计自定义的allocator类来管理vector容器的内存分配和释放；vector容器默认使用C++标准库的allocator空间配置器来管理内存分配和释放，因此可以基于统一内存实现自定义的allocator类，从而实现对vector容器内存分配和释放的自动管理；自定义allocator类中，基于cudaMallocManaged()实现allocator类中的_allocator()函数，基于cudaFree()实现allocator类中的_deallocate()函数；在vector容器声明时显示调用自定义allocator类，即可实现统一内存对vector容器在CPU和GPU上的内存空间的分配、释放、数据传输操作的自动管理。
CN202010971038.9A 2020-09-15 2020-09-15 一种面向异构平台的复杂指针数据结构自动管理*** Active CN112083956B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202010971038.9A CN112083956B (zh) 2020-09-15 2020-09-15 一种面向异构平台的复杂指针数据结构自动管理***

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202010971038.9A CN112083956B (zh) 2020-09-15 2020-09-15 一种面向异构平台的复杂指针数据结构自动管理***

Publications (2)

Publication Number Publication Date
CN112083956A true CN112083956A (zh) 2020-12-15
CN112083956B CN112083956B (zh) 2022-12-09

Family

ID=73736379

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202010971038.9A Active CN112083956B (zh) 2020-09-15 2020-09-15 一种面向异构平台的复杂指针数据结构自动管理***

Country Status (1)

Country Link
CN (1) CN112083956B (zh)

Cited By (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
WO2023197397A1 (zh) * 2022-04-13 2023-10-19 堡垒科技有限公司 一种去中心化的开源软件可信通证化协议

Citations (10)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US6725448B1 (en) * 1999-11-19 2004-04-20 Fujitsu Limited System to optimally create parallel processes and recording medium
CN101963918A (zh) * 2010-10-26 2011-02-02 上海交通大学 实现cpu/gpu异构平台的虚拟执行环境的方法
CN102707952A (zh) * 2012-05-16 2012-10-03 上海大学 嵌入式异构多核处理器上基于用户描述的程序设计方法
CN104035781A (zh) * 2014-06-27 2014-09-10 北京航空航天大学 一种快速开发异构并行程序的方法
CN106874113A (zh) * 2017-01-19 2017-06-20 国电南瑞科技股份有限公司 一种cpu+多gpu异构模式静态安全分析计算方法
CN106940654A (zh) * 2017-02-15 2017-07-11 南京航空航天大学 源代码中内存错误的自动检测和定位方法
CN108536581A (zh) * 2018-03-08 2018-09-14 华东师范大学 一种针对源代码的运行时形式化验证方法及***
CN109933327A (zh) * 2019-02-02 2019-06-25 中国科学院计算技术研究所 基于代码融合编译框架的OpenCL编译器设计方法和***
CN110383247A (zh) * 2017-04-28 2019-10-25 伊纽迈茨有限公司 由计算机执行的方法、计算机可读介质与异构计算***
CN111611158A (zh) * 2020-05-08 2020-09-01 中国原子能科学研究院 一种应用性能分析***及方法

Patent Citations (10)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US6725448B1 (en) * 1999-11-19 2004-04-20 Fujitsu Limited System to optimally create parallel processes and recording medium
CN101963918A (zh) * 2010-10-26 2011-02-02 上海交通大学 实现cpu/gpu异构平台的虚拟执行环境的方法
CN102707952A (zh) * 2012-05-16 2012-10-03 上海大学 嵌入式异构多核处理器上基于用户描述的程序设计方法
CN104035781A (zh) * 2014-06-27 2014-09-10 北京航空航天大学 一种快速开发异构并行程序的方法
CN106874113A (zh) * 2017-01-19 2017-06-20 国电南瑞科技股份有限公司 一种cpu+多gpu异构模式静态安全分析计算方法
CN106940654A (zh) * 2017-02-15 2017-07-11 南京航空航天大学 源代码中内存错误的自动检测和定位方法
CN110383247A (zh) * 2017-04-28 2019-10-25 伊纽迈茨有限公司 由计算机执行的方法、计算机可读介质与异构计算***
CN108536581A (zh) * 2018-03-08 2018-09-14 华东师范大学 一种针对源代码的运行时形式化验证方法及***
CN109933327A (zh) * 2019-02-02 2019-06-25 中国科学院计算技术研究所 基于代码融合编译框架的OpenCL编译器设计方法和***
CN111611158A (zh) * 2020-05-08 2020-09-01 中国原子能科学研究院 一种应用性能分析***及方法

Non-Patent Citations (7)

* Cited by examiner, † Cited by third party
Title
WANG FARUI等: "AUTOMATIC translation of data parallel programs FOR HETEROGENEOUS parallelism through OPENMP OFFLOADING", 《JOURNAL OF SUPERCOMPUTING》 *
刘晓娴: "面向共享存储结构的并行编译优化技术研究", 《中国博士学位论文全文数据库 信息科技辑》 *
孙大林; 唐好选: "路面点云的并行简化研究", 《智能计算机与应用》 *
孙守航: "异构多核处理器OpenMP编译实现与优化", 《中国优秀硕士学位论文全文数据库 信息科技辑》 *
李雁冰等: "一种面向异构众核处理器的并行编译框架", 《软件学报》 *
江霞等: "OpenACC到MIC平台上并行程序的自动翻译及优化", 《小型微型计算机***》 *
郭浩男: "面向异构平台的OpenMP程序自动卸载及优化", 《中国优秀硕士学位论文全文数据库 信息科技辑》 *

Cited By (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
WO2023197397A1 (zh) * 2022-04-13 2023-10-19 堡垒科技有限公司 一种去中心化的开源软件可信通证化协议

Also Published As

Publication number Publication date
CN112083956B (zh) 2022-12-09

Similar Documents

Publication Publication Date Title
US6505344B1 (en) Object oriented apparatus and method for allocating objects on an invocation stack
US9471291B2 (en) Multi-processor code for modification for storage areas
US7912877B2 (en) Leveraging garbage collection to dynamically infer heap invariants
EP1519273B1 (en) Region-based memory management for object-oriented programs
US8065668B2 (en) Unified data type system and method
US8707278B2 (en) Embedding class hierarchy into object models for multiple class inheritance
US10452409B2 (en) Universal adapter for native calling
JPH06103075A (ja) オブジェクト指向適用業務
US7058943B2 (en) Object oriented apparatus and method for allocating objects on an invocation stack in a partial compilation environment
Vollmer et al. Compiling tree transforms to operate on packed representations
CN112083956B (zh) 一种面向异构平台的复杂指针数据结构自动管理***
US20020062478A1 (en) Compiler for compiling source programs in an object-oriented programming language
Calvert Parallelisation of java for graphics processors
CN111966397A (zh) 一种异构并行程序自动移植和优化方法
JP7140935B1 (ja) リアルタイムアプリケーションのための決定的メモリ割り当て
Chambers et al. Iterative type analysis and extended message splitting: Optimizing dynamically-typed object-oriented programs
Qiu Programming language translation
Wroblewski et al. Accelerating Spark Datasets by inlining deserialization
Skjellum et al. Object‐oriented analysis and design of the Message Passing Interface
Abbas et al. Object oriented parallel programming
Squyres et al. The design and evolution of the MPI-2 C++ interface
Pettersson et al. 9 Simulating Tailcalls in C
CN118132086A (zh) 一种将CUDA代码转为OpenCL代码的方法及***
Pachev GPUMap: A Transparently GPU-Accelerated Map Function
Benoit et al. Runtime support for automatic placement of workloads on heterogeneous processors

Legal Events

Date Code Title Description
PB01 Publication
PB01 Publication
SE01 Entry into force of request for substantive examination
SE01 Entry into force of request for substantive examination
GR01 Patent grant
GR01 Patent grant