CN112083956B - Heterogeneous platform-oriented automatic management system for complex pointer data structure - Google Patents

Abstract

An automatic management system for a complex pointer data structure facing a heterogeneous platform relates to the technical field of heterogeneous programming. The invention aims to realize the automatic management of a complex pointer data structure in an OpenMP Offloading program on a heterogeneous computing platform and ensure the data consistency. The invention includes: the information collection module is used for carrying out static analysis on the source program and collecting program information; the automatic conversion module is mainly responsible for modifying the source code at a proper position according to different variable types and inserting a proper runtime API; and the runtime module is mainly responsible for realizing the memory management operation of the C + + standard again by using the cudaMallocManged () and the cudaFree () and providing an interface outwards. The invention can automatically manage the memory allocation, release and data transmission of the complex pointer data structure in the OpenMP offload program between the CPU and the GPU memory, and ensure the data consistency; therefore, convenience is provided for the development of the OpenMP Offloading program.

Description

A complex pointer data structure automatic management system for heterogeneous platforms

technical field

The invention relates to an automatic management system for complex pointer data structures in an OpenMP Offloading program, and relates to the technical field of heterogeneous programming.

Background technique

OpenMP is led by the OpenMP Architecture Review Board and has been widely accepted as a set of guiding compilation processing scheme (CompilerDirective) [1] for multiprocessor programming of shared memory parallel system. The programming languages supported by OpenMP include C, C++, and Fortran; the compilers that support OpenMP include Sun Compiler, GNU Compiler, and Intel Compiler. OpenMP provides a high-level abstract description of parallel algorithms. Programmers specify their intentions by adding special pragma primitives to the source code, so that the compiler can automatically parallelize the program and add synchronization where necessary. Mutex and communication.

In the field of high-performance computing, various accelerators (such as GPU, FPGA, DSP, MLU, etc.) have become an important source of computing power in addition to the CPU. Starting from version 4.0, OpenMP has added the Offloading (unloading) feature. OpenMPOffloading supports the heterogeneous programming model of CPU+accelerator; after the development of versions 4.5 and 5.0, OpenMPOffloading has been gradually improved. OpenMP Offloading makes it possible for OpenMP programs to make full use of the computing power of heterogeneous computing platforms, but it is still a difficult, tedious, and error-prone task to modify existing OpenMP CPU programs to meet the Offloading characteristics, especially when the program contains complex Pointer data structures, such as: class nested pointers, vector containers, multi-level nested pointers, etc.

Although the syntax of OpenMPOffloading is simple, users still need to use relevant pragma primitives to display and manage data transfer operations between the CPU and the accelerator, which causes great inconvenience to developers, especially when encountering complex pointer data structures. For example, the memory allocation of the vector container in C++ is implicit. It is difficult for developers to control memory allocation and data transmission, and it is difficult to use the Offloading feature. For nested pointers or multi-level nested pointers, dealing with the allocation, release, and data transfer of memory pointed to by different levels of pointers in the CPU and accelerator memory space is tedious and error-prone work, which also discourages developers from the Offloading feature. .

The CUDA programming model of the NVIDIA GPU platform supports the Unified Memory (UM) feature from version 6.0; this feature unifies the address spaces of the CPU and GPU and automatically manages data transmission between the CPU and GPU. The unified memory feature provides a possible technical means for the automatic processing of complex pointer data structures in the development of OpenMP Offloading programs.

Contents of the invention

The technical problem to be solved in the present invention is:

The purpose of the present invention is to propose a complex pointer data structure automatic management system in an OpenMP Offloading program oriented to heterogeneous platforms, so as to solve the problem that the prior art cannot automatically modify memory allocation and release statements and automatically insert relevant pragmas on the basis of OpenMP CPU programs. Primitives cannot automatically manage the data transmission between the CPU and the accelerator, and cannot guarantee the data consistency of programs containing complex pointer data structures on heterogeneous computing platforms, thus affecting program performance.

The technical scheme that the present invention adopts for solving the problems of the technologies described above is:

A heterogeneous platform-oriented complex pointer data structure automatic management system, the system is used to realize the automatic management of the complex pointer data structure in the OpenMPOffloading program on the heterogeneous computing platform and ensure data consistency;

The system includes the following three modules:

Information collection module, this module has two functions: 1) Statically analyze the source program to collect program information; 2) Establish an abstract representation of the source program based on the collected information, that is, a serial-parallel control flow graph;

The working process of this module is divided into the following two steps: 1) Through the Clang compiler, generate the corresponding abstract syntax tree (Abstract Syntax Tree, AST) and control flow graph (Control Flow Graph, CFG) from the source code of C or C++, Traverse AST and CFG, distinguish between serial and parallel domains, and obtain program details; 2) Generate serial and parallel control flow graphs based on these information;

The automatic conversion module is mainly responsible for inserting the runtime API into the source code based on the serial-parallel control flow graph to complete the code conversion; firstly, according to the complex pointer variable information saved in the serial-parallel control flow graph, determine the type of the complex pointer data; then according to For different types, insert the appropriate runtime API at the appropriate position in the source code to complete the code conversion; in this way, the memory allocation, release, and data transfer operations involved in complex pointer variables are all taken over by the runtime, so that complex pointer variables can be automatically managed through runtime. Pointer variables to ensure data consistency between CPU and GPU;

The runtime module is mainly responsible for the following operations of complex pointer data types based on unified memory: memory allocation and release operations on the CPU and GPU, and automatic data transfer operations between the CPU and GPU; the runtime module consists of the UMMaper class and the allocator class Composition, the UMMaper class and the allocator class provide an interface for memory allocation and release to the outside in the form of an API interface.

Further, the realization process of the information collection function of the information collection module is:

First, use the Clang compiler to perform lexical analysis and semantic analysis on the target program, and generate an abstract syntax tree (Abstract Syntax Tree, AST) and a control flow graph (Control Flow Graph, CFG);

After that, the following three static analyzes are performed on AST and CFG to obtain function call relationship information, variable related information, serial/parallel domain related information, and provide information support for the subsequent establishment of serial and parallel control flow graphs and code conversion;

Analysis 1, function call relationship analysis: For an AST, perform the following two steps:

1) Recursively traverse each node on the AST;

2) If the current node is a function definition node, then save the function name and the sub-function information of the function call;

Analysis 2, variable information analysis, for an AST, perform the following three steps:

1) Recursively traverse each node on the AST;

2) If the current node involves variable definition or reference, save variable definition information, variable reference information, and variable scope information;

3) If the current node involves memory allocation or release, save memory allocation information and memory release information;

Analysis 3, serial parallel domain analysis: For a CFG, perform the following three steps:

1) Recursively traverse each node on the CFG and save the relationship information between nodes;

2) If the current node is within the scope of the OpenMP#pragma parallel instruction statement, then the current node is marked as a parallel node, and its type information and range information are saved;

3) If the current node is not within the scope of the OpenMP#pragma parallel instruction statement, then mark the current node as a serial node, and save its type information and range information;

Through the above three static analysis, the following information can be obtained: function call relationship information, variable definition information, variable reference information, variable scope information, memory allocation information, memory release information, serial-parallel node type information, serial-parallel node range information, Relationship information between nodes; these information will provide support for the establishment of serial and parallel control flow graphs.

Further, the serial-parallel control flow graph in the information collection module is defined as follows:

Define the serial-parallel control flow graph to consist of serial nodes, parallel nodes, and directed edges between nodes; among them, a serial node represents a piece of OpenMP#pragma parallel instruction statement scope, internal branchless, and serially executed code Fragment; the code fragment corresponding to the serial node is executed on the CPU, and the serial node is also recorded as a SEQ node;

A parallel node represents a code segment that is executed in parallel within the scope of the OpenMP#pragma parallel instruction statement; the code segment corresponding to the parallel node is offloaded to the GPU for execution, and the parallel node is also recorded as an OMP node;

The serial node and the parallel node store function call information and variable related information;

The directed edges between the nodes represent the execution order of the code fragments corresponding to the nodes.

Further, the serial-parallel control flow graph establishment process in the information collection module:

The process of establishing a serial-parallel control flow graph takes functions as the basic processing unit; for the entire source program, if a serial-parallel control flow graph of a function can be established, combined with the collected function call relationship information, the entire source program can be recursively established. Serial-parallel control flow graph;

For a function, based on the collected information, a serial-parallel control flow graph can be built through the following steps:

1) Use the node type information and node range information to create isolated serial nodes and parallel nodes;

2) Using the relationship information between nodes, establish directed edges between nodes, and connect serial nodes and parallel nodes into a graph;

3) Save function call information, variable definition information, variable reference information, variable scope information, memory allocation information, and memory release information to corresponding serial nodes or parallel nodes according to the node range information.

Further, the complex pointer data is divided into three types according to the position of the pointer:

Class nested pointers, vector containers, multi-level nested pointers;

The class nested pointer refers to the pointer contained in the class, and the vector container refers to the vector container provided in the C++ standard library; the multi-level nested pointer refers to the second-level pointer and the pointer above the second level.

Further, the class nested pointer is processed through the following steps:

① Recursively traverse the serial-parallel control flow graph established above, and find out the class nested pointers defined in the serial node and referenced in the parallel node and their The C++ class where it is located;

② For the C++ class found in step ①, modify the type definition of this class in the source code, so that the type of this class inherits the UMMaper base class provided at runtime when it is defined;

③For the class nesting pointer found in step ①, modify the memory allocation and release statement of the pointer in the source code, use cudaMallocManaged() to allocate memory, and use cudaFree() to release memory;

④ For the C++ class instance found in step ①, modify the definition statement of the class instance in the source code, use the overloaded new operator to create a class instance, and pass the memory space address allocated in step ③ to the C++ class instance The corresponding nested pointer.

Further, the vector container is processed through the following steps:

① Recursively traverse the serial-parallel control flow graph established above, and find out the vector container defined in the serial node and referenced in the parallel node according to the variable definition and reference information saved in the serial/parallel node;

② For the vector container instance found in step ①, modify the definition statement of the vector container instance, and insert an explicit call to the custom allocator provided at runtime.

Further, the multi-level nested pointer is processed through the following steps:

① Recursively traverse the serial-parallel control flow graph established above, and find out the multi-level nested pointers defined in the serial node and referenced in the parallel node according to the variable definition and reference information stored in the serial/parallel node, and Its sub-pointers at all levels;

② For all the multi-level nested pointers and sub-pointers at all levels found in step ①, modify the memory allocation and release statements of these pointers in the source code, use cudaMallocManaged() to allocate memory, and use cudaFree() to release memory.

Further, the implementation process of the UMMaper class in the runtime module: for processing class nested pointers, design the UMMaper class to manage the memory allocation and release of the C++ class; use cudaMallocManaged() and cudaFree() in the UMMaper class to C++ default New and delete operators are overloaded; the allocation, release, and data transfer operations of the memory space of the derived classes of the UMMaper class on the CPU and GPU can be automatically managed by the unified memory.

Further, the implementation process of the custom allocator class in the runtime module: in order to process the vector container, design a custom allocator class to manage the memory allocation and release of the vector container; the vector container uses the allocator space configurator of the C++ standard library by default To manage memory allocation and release, so you can implement a custom allocator class based on unified memory, so as to realize the automatic management of vector container memory allocation and release; in the custom allocator class, implement _allocator( in the allocator class based on cudaMallocManaged() ) function, based on cudaFree() to implement the _deallocate() function in the allocator class; when the vector container is declared, it is displayed to call the custom allocator class, and the unified memory can allocate and release the memory space of the vector container on the CPU and GPU , Automatic management of data transfer operations.

The present invention has the following beneficial technical effects:

The system of the present invention can automatically modify the memory allocation and release statements on the basis of the OpenMP CPU program, and automatically insert relevant pragma primitives, and automatically manage the data transmission between the CPU and the accelerator, so as to ensure that the program containing complex pointer data structures Data consistency on heterogeneous computing platforms and improve program performance.

The heterogeneous programming scheme researched by the present invention is mainly aimed at the automatic management of complex pointer data structures in the OpenMP program, specifically refers to automatically modifying memory allocation and release statements, automatically inserting relevant pragma primitives, and automatically managing data transmission between the CPU and the accelerator; thus Ensure data consistency of programs on heterogeneous computing platforms and improve program performance.

Since OpenMPOffloading supports accelerator programming, the OpenMP code on the CPU can be offloaded to the GPU for execution, and objective conditions allow the offloading of the OpenMP code. Under this condition, offloading the OpenMP code to run on the GPU can improve the running efficiency of the program on the one hand, and make full use of the acceleration effect of the GPU on the other hand. However, manual code unloading cannot guarantee the correctness of the conversion program, and at the same time consumes manpower and material resources, and the efficiency is very low. Therefore, the present invention proposes an automatic management scheme for complex pointer data structures to solve the most difficult problems in the unloading process, such as memory allocation and data transmission between the CPU and the accelerator for complex pointer data structures.

Through comparative experiments on general test sets (PolyBench, Rodinia, etc.), it is proved that the present invention can automatically manage complex pointer data structures in OpenMP Offloading programs, ensure program correctness, and improve program performance.

The complex pointer data structure described in the present invention refers to the complex pointer data type.

Description of drawings

Fig. 1 is the overall frame diagram of the system of the present invention; Fig. 2 is the structural block diagram of AST analysis method; Fig. 3 is a function string parallel control flow diagram; Fig. 4 is a schematic diagram of unified memory; Fig. 5 is a processing class nested pointer Schematic diagram of automatic unloading implementation (comparison of programs using unified memory or not); Figure 6 is a schematic diagram of automatic unloading of vector containers (comparison of programs using space configurator or not); Figure 7 is automatic unloading of processing multi-level nested pointers Realize schematic diagram; Fig. 8 is the overall frame diagram of experimental design of the present invention;

Figure 9 is a histogram of running time on the 2080Ti platform under Large data volume. In the figure: Figure 9(a) is a comparison chart of Large-2080Ti running time (data set a), and Figure 9(b) is a comparison chart of Large-2080Ti running time (Data set b), Figure 9(a) and Figure 9(b) are essentially a graph, because there are many data sets, so they are divided into two graphs;

Figure 10 is a comparison chart of the running time of the Rodinia test set on 2080Ti; Figure 11 is a comparison chart of Polybench-OAO different data volume speedup (2080Ti); Figure 12 is a comparison chart of Polybench speedup (2080Ti); Figure 13 is the runtime overhead Figure (2080Ti); Figure 14 is a detailed test chart of complex pointer data structure (K40);

Figure 15 is an example of program conversion with vector; Figure 16 is an example of program conversion with structnest; Figure 17 is an example of program conversion with multilevel. In Fig. 15 to Fig. 17: vector is a vector container in c++; structnest is a class nesting pointer; multilevel is a multilevel nesting pointer.

detailed description

In conjunction with accompanying drawings 1 to 17, a heterogeneous platform-oriented complex pointer data structure automatic management system according to the present invention is described as follows:

The main task of the present invention is to automatically manage complex pointer data structures (class nested pointers, vector containers, multi-level nested pointers, etc.) in the OpenMPOffloading program, that is, to realize automatic modification of allocation and release statements, automatic management of data transmission, and ensure data consistency sex. The present invention mainly comprises following three modules:

Information collection module, this module has two functions: 1) Statically analyze the source program to collect program information; 2) Create an abstract representation of the source program based on the collected information, that is, a serial-parallel control flow graph. The working process of this module is divided into the following two steps: 1) Through the Clang compiler, generate the corresponding abstract syntax tree (Abstract SyntaxTree, AST) and control flow graph (Control Flow Graph, CFG) from the source code of C or C++, and traverse AST and CFG distinguish serial and parallel domains and obtain program details; 2) generate serial and parallel control flow graphs based on these information.

The automatic conversion module is mainly responsible for inserting the runtime API into the source code based on the serial-parallel control flow graph to complete the code conversion. Firstly, according to the complex pointer variable information stored in the serial-parallel control flow graph, the type of the complex pointer variable is determined. Then according to different types, insert the appropriate runtime API at the appropriate position in the source code to complete the code conversion. In this way, the memory allocation, release, and data transfer operations involved in complex pointer variables are all taken over by the runtime, so that complex pointer variables can be automatically managed by runtime to ensure data consistency between the CPU and GPU.

The runtime module is mainly responsible for the following operations of complex pointer data structures based on unified memory: memory allocation and release operations on the CPU and GPU, automatic data transfer operations between the CPU and GPU; among them, cudaMallocManaged() and cudaFree() are used to re- Implement the C++ default interfaces of new, delete, and allocator, and provide external memory allocation and release interfaces in the form of API interfaces.

As a result, the source program is converted into a new program with a runtime API, and the overall framework of the system is shown in Figure 1.

1 Information Collection Module

The information collection module implements two functions: static analysis and collection of program information; establishment of serial and parallel control flow graphs.

1.1 Information collection

First, use the Clang compiler to perform lexical analysis and semantic analysis on the target program, and generate an abstract syntax tree (Abstract Syntax Tree, AST) and a control flow graph (Control Flow Graph, CFG).

Afterwards, the following three static analyzes are performed on AST and CFG to obtain function call relationship information, variable related information, and serial/parallel domain related information, and provide information support for subsequent serial-parallel control flow graph establishment and code conversion.

Analysis 1, function call relationship analysis: For an AST, perform the following two steps:

1) Recursively traverse each node on the AST;

2) If the current node is a function definition node, save the function name and sub-function information called by the function.

Analysis 2, variable information analysis, for an AST, perform the following three steps:

1) Recursively traverse each node on the AST;

2) If the current node involves variable definition or reference, save variable definition information, variable reference information, and variable scope information;

3) If the current node involves memory allocation or release, save memory allocation information and memory release information.

Analysis 3, serial parallel domain analysis: For a CFG, perform the following three steps:

1) Recursively traverse each node on the CFG and save the relationship information between nodes;

2) If the current node is within the scope of the OpenMP#pragma parallel instruction statement, then the current node is marked as a parallel node, and its type information and range information are saved;

3) If the current node is not within the scope of the OpenMP#pragma parallel instruction statement, then mark the current node as a serial node, and save its type information and range information.

Through the above three static analysis, the following information can be obtained: function call relationship information, variable definition information, variable reference information, variable scope information, memory allocation information, memory release information, serial-parallel node type information, serial-parallel node range information, Relationship information between nodes. These information will provide support for the establishment of serial and parallel control flow graphs.

1.2 Serial-parallel control flow graph establishment

Definition Serial-parallel control flow graph consists of serial nodes, parallel nodes and directed edges between nodes. Among them, the serial node represents a code fragment that is outside the scope of the OpenMP#pragma parallel instruction statement, has no internal branches, and is executed serially; the code fragment corresponding to the serial node is executed on the CPU, and the serial node is also recorded as a SEQ node.

A parallel node represents a code fragment that is executed in parallel within the scope of the OpenMP#pragma parallel instruction statement; the code fragment corresponding to the parallel node is offloaded to the GPU for execution, and the parallel node is also recorded as an OMP node. The serial node and the parallel node store function call information and variable related information. The directed edges between the nodes represent the execution order of the code fragments corresponding to the nodes.

The process of establishing a serial-parallel control flow graph takes functions as the basic processing unit; for the entire source program, if a serial-parallel control flow graph of a function can be established, combined with the function call relationship information, the serial-parallel control of the entire source program can be recursively established flow graph.

For a function, based on the collected program information, a serial-parallel control flow graph can be built through the following steps:

1) Use the node type information and node range information to create isolated serial nodes and parallel nodes;

2) Using the relationship information between nodes, establish directed edges between nodes, and connect serial nodes and parallel nodes into a graph;

3) Save function call information, variable definition information, variable reference information, variable scope information, memory allocation information, and memory release information to corresponding serial nodes or parallel nodes according to the node range information.

2 automatic conversion modules

The automatic conversion module divides complex pointer data types into three types: class nested pointer, vector container, and multi-level nested pointer.

The key technology in the processing of complex pointer data structures is Unified Memory (UM), so first introduce the concept and principle of Unified Memory. UM maintains a unified memory pool, shared between CPU and GPU, and only allocates memory once. , this data pointer is available to both the host side (CPU) and the device side (GPU). Using a single pointer for managed memory, data transfer between different devices is done automatically by UM's runtime system and enables the GPU to process data sets that exceed its memory capacity. The principle of unified memory is shown in Figure 4.

Using unified memory technology, different processing methods are designed for the three complex pointer data types.

2.1 Class nested pointers

The class nested pointer refers to the pointer contained in the class, as shown in Figure 5, it is processed through the following steps:

① Traverse the serial-parallel control flow graph recursively, and find out the class nesting pointers defined in the serial node and referenced in the parallel node and the classes they belong to according to the variable definition and reference information saved in the serial/parallel node.

③ For the type of class found in step ①, modify the type definition of this class in the source code, so that the type of this class inherits the UMMaper base class provided at runtime when it is defined (see runtime design);

③For the class nesting pointer found in step ①, modify the memory allocation and release statement of the pointer in the source code, use cudaMallocManaged() to allocate memory, and use cudaFree() to release memory;

④ For the class instance found in step ①, modify the definition statement of the class instance in the source code, use the overloaded new operator to create the class instance, and pass the memory space address allocated in step ③ to the corresponding embedded in the class Set of pointers.

An example of processing class nested pointers is shown in Figure 5.

2.2vector container

The vector container refers to the vector container provided in the C++ standard library. Vector containers need to be processed through the following steps:

① Traverse the serial-parallel control flow graph recursively, and find out the vector container defined in the serial node and referenced in the parallel node according to the variable definition and reference information saved in the serial/parallel node.

② For the vector container instance found in step ①, modify the definition statement of the vector container instance and insert an explicit call to the custom allocator provided at runtime (see runtime design).

An example of vector container processing is shown in Figure 6.

2.3 Multi-level nested pointers

A multi-level nested pointer refers to a second-level pointer and pointers above the second level. Multiple levels of nested pointers need to be handled through the following steps:

① Recursively traverse the serial-parallel control flow graph, and find out the multi-level nested pointers defined in the serial node and referenced in the parallel node, and their levels child pointer;

② For all the multi-level nested pointers and sub-pointers at all levels found in step ①, modify the memory allocation and release statements of these pointers in the source code, use cudaMallocManaged() to allocate memory, and use cudaFree() to release memory.

An example of processing multi-level nested pointers is shown in Figure 7.

3 runtime

In order to make the converted code closer to the format of the source code and minimize the modification of the source code, this paper proposes a solution to insert a small amount of runtime API into the original program, and encapsulate the memory allocation based on unified memory into the runtime.

For classes and their nested pointers, the UMMapper base class responsible for memory allocation and release is implemented at runtime, as shown in Figure 5. Use cudaMallocManaged() and cudaFree() in UMMapper to overload the default new and delete operators of C++. In this way, you only need to derive based on UMMapper when the class is declared, and then use new to create a new class instance to realize automatic management of memory allocation, release, and data transfer operations related to derived classes through unified memory.

The vector container uses the allocator space configurator of the C++ standard library to manage memory allocation and release, so a custom allocator class can be implemented based on unified memory, thereby realizing automatic management of memory allocation and release of the vector container. In the custom allocator class, the _allocator() function in the allocator class is implemented based on cudaMallocManaged(), and the _deallocate() function in the allocator class is implemented based on cudaFree(). When the vector container is declared, it is displayed to call the custom allocator class, so that unified memory can automatically manage the allocation, release, and data transmission operations of the memory space of the vector container on the CPU and GPU, as shown in Figure 6.

The technical effect of the invention is verified as follows:

In order to verify and analyze the effect of the code unloading scheme, we conducted a test on the RTX 2080Ti platform, using PolyBench, Rodinia and other benchmark test data sets, and compared the test results with another well-known source-to-source compiler (DawnCC) and manual conversion ( Manual) and CPU parallel code for comparison. Our solution is named OAO (OpenMP Automated Offloading with complex data structure support).

The test and acquisition of experimental results need to integrate the four stages of runtime, scripting and running, Python extraction and OriginLab analysis. The overall process of the experiment is shown in Figure 8.

1 runtime

The running time of the 2080Ti platform under the Large data volume is shown in Figure 9.

The running time comparison of the Rodinia test set on the 2080Ti is shown in Figure 10.

Through the above collection and comparison of the running time of the dataset, the following conclusions can be drawn:

(1) It can be seen from the running time graph that OAO can handle all 23 test programs, while DawnCC can only handle 15 test programs. OAO has improved performance on 9 programs on the K40 platform and 15 programs on the 2080Ti platform; and compared to manual conversion, OAO has the best performance on all test programs on all platforms.

(2) In Rodinia-2080Ti, the Rodinia data volume is known to be obtained from the pre-experimental test. There are 8 data sets in total, and the OAO running time of 5 data sets is shorter than that of OMP.

(3) In Polybench-Large, the running time of OAO is better than DawnCC, Manual and native; the running time generally shows a trend of OAO<DawnCC<Manual<native. Among them, the running time of native is the longest, because the selected transmission statement has not been optimized in parallel; Manual includes redundant transmission of some data, and the running time is longer than OAO; and OAO is generally better in the processing of data sets. at DawnCC.

2 speedup ratio

It can be seen from Figure 11 that in terms of speedup ratio, among the data greater than 1, the speedup ratio of Large data volume is greater than that of Medium data, which is also in line with the characteristics of parallel acceleration. The larger the amount of data processed, the more obvious the effect of parallel speedup.

In terms of acceleration ratio, the calculation method is the ratio of OMP/Type, Type is other types, and the result is bounded by 1X, the upper part is positive acceleration, and the lower part is negative acceleration. It can be concluded from Figure 12 that regardless of the platform, the results of all test programs show that the speedup ratio of OAO is better than Manual. In the Polybench test program, OAO is slightly better than DawnCC, and because DawnCC cannot handle the Rodinia test program, Therefore, OAO has the best data transmission optimization effect in all test programs, and can handle complex test programs that DawnCC cannot handle, and has a wider scope of application.

3 runtime overhead

In the Polybench test set, the runtime overhead is shown in Figure 13.

Compared with the runtime overhead, the runtime overhead of all test programs is more than 2 orders of magnitude smaller, so the extra overhead of OAO at runtime is very small, almost negligible, so adding runtime has a great impact on the performance of the source program There is no loss.

4 unified memory

Detailed tests of unified memory include vector containers, class nested pointers (structnest) and multilevel nested pointers (multilevel). The running time and data transfer time comparison of the complex pointer data structure is shown in Figure 14.

From the perspective of data transmission time, the time of H2D and D2H is much less than the total running time, and all complex pointer data structures can be processed correctly, so the OAO compiler can correctly support GPU offloading of complex pointer data structures.

Among them, the vector container, the comparison of the simple conversion programs of structnest and multilevel are shown in Figure 15, Figure 16 and Figure 17 respectively.

As shown in Figure 15, the vector variable in the declaration phase is modified to be overloaded by QY::allocator; in Figure 16, the nested pointers x, y, and z inside the class in the declaration phase are overloaded by cudaMallocManaged, and the instantiated class a is defined by The transfer statement indicates that map(tofrom:a[0:1]) is added; in Figure 17, x is the second-level pointer, y is the first-level pointer of x, and x and y are overloaded by cudaMallocManaged in the declaration stage.

Table 1 Unified memory test result table

The unified memory test results are shown in Table 1. In the process of data transmission, a transmission will be given priority before the official transmission. The data size of H2D is 4B, and the data size of D2H is 1B (4+1 mode). For the two-way transmission of vector in the above table, 52 is 3 larger than 49 , The two-way transmission of structnest, 36 is 3 larger than 33, all due to this reason. In addition, the vector is only transmitted twice, namely X and Y, each time the data size is 24B, and the structnest is transmitted once, A is transmitted, and the data size is 32.

In multilevel, z is a one-way transmission, using the data transmission model, x and y are multi-level nested pointers, using unified memory. Except for the 4+1 mode of priority transmission, the reason why there is a difference in the amount of data transmission between H2D and D2H is that z does not return, but x does.

Experimental results prove that OAO can use unified memory to correctly process complex pointer data structures, and realize the support for processing complex pointer data structures.

Claims (10)

1. An automatic management system of a complex pointer data structure facing a heterogeneous platform is characterized in that the system is used for realizing automatic management of the complex pointer data structure in an OpenMPOfflooding program on the heterogeneous computing platform and ensuring data consistency;

the system comprises the following three modules:

the information collection module has two functions: 1) Performing static analysis on the source program to collect program information; 2) Establishing an abstract representation, namely a serial-parallel control flow diagram, for a source program based on the collected information;

the working process of the module comprises the following two steps: 1) Generating a corresponding abstract syntax tree AST and a control flow graph CFG from a C or C + + source code through a Clang compiler, traversing the AST and the CFG, distinguishing a serial-parallel domain, and acquiring detailed program information; 2) Generating a serial-parallel control flow graph according to the information;

the automatic conversion module is mainly responsible for inserting an API (application program interface) in operation into a source code based on a serial-parallel control flow graph so as to complete code conversion; firstly, determining the type of complex pointer data according to complex pointer variable information stored in a serial-parallel control flow diagram; then according to different types, inserting proper runtime API in proper position in the source code to complete code conversion; in this way, the memory allocation, release and data transmission operations related to the complex pointer variable are all taken over by the runtime, so that the complex pointer variable can be automatically managed by the runtime, and the data consistency between the CPU and the GPU is ensured;

the runtime module is mainly responsible for realizing the following operations of the complex pointer data type based on the unified memory: memory allocation and release operations on the CPU and the GPU and automatic data transmission operations between the CPU and the GPU; the runtime module consists of a ummale class and an allocator class, and the ummale class and the allocator class provide interfaces for memory allocation and release to the outside in the form of API interfaces.

2. The system for automatically managing the complex pointer data structure oriented to the heterogeneous platform according to claim 1, wherein the information collection function of the information collection module is implemented by:

firstly, performing lexical analysis and semantic analysis on a target program by using a Clang compiler, and generating an abstract syntax tree AST and a control flow graph CFG;

then, the AST and the CFG are subjected to the following three static analyses to obtain function call relation information, variable related information and serial/parallel domain related information, and information support is provided for the establishment of a serial-parallel control flow graph and code conversion;

analysis one, function call relation analysis: for an AST, the following two steps are performed:

1) Recursively traversing each node on the AST;

2) If the current node is a function definition node, storing the function name and the sub-function information called by the function;

analysis two, variable information analysis, for an AST, the following three steps of work are performed:

1) Recursively traversing each node on the AST;

2) If the current node relates to variable definition or reference, saving variable definition information, variable reference information and variable scope information;

3) If the current node relates to memory allocation or release, storing memory allocation information and memory release information;

analysis three, series-parallel domain analysis: for one CFG, the following three steps are performed:

1) Recursively traversing each node on the CFG, and storing relationship information between the nodes;

2) If the current node is in the action domain of the OpenMP # pragma parallel instruction statement, marking the current node as a parallel node, and storing the type information and the range information of the current node;

3) If the current node is not in the action domain of the OpenMP # pragma parallel instruction statement, marking the current node as a serial node, and storing the type information and the range information of the serial node;

through the three static analyses described above, the following information can be obtained: function call relation information, variable definition information, variable reference information, variable scope information, memory allocation information, memory release information, serial-parallel node type information, serial-parallel node range information and inter-node relation information; this information will provide support for the serial-parallel control flow graph setup.

3. The automatic management system for the complex pointer data structure oriented to the heterogeneous platform according to claim 1 or 2, characterized in that the serial-parallel control flow graph in the information collection module is defined as follows:

defining a serial-parallel control flow graph to be composed of serial nodes, parallel nodes and directed edges among the nodes; the serial node represents a code segment which is outside the scope of action of an OpenMP # pragma parallel instruction statement, has no branch inside and is executed serially; executing the code segments corresponding to the serial nodes on the CPU, wherein the serial nodes are also marked as SEQ nodes;

the parallel node represents a code segment which is in the OpenMP # pragma parallel instruction statement action domain and is executed in parallel; unloading the code segments corresponding to the parallel nodes to a GPU for execution, wherein the parallel nodes are also marked as OMP nodes;

function calling information and variable related information are saved in the serial node and the parallel node;

and the directed edges among the nodes represent the sequential relation of the execution of the code segments corresponding to the nodes.

4. The system according to claim 3, wherein the process of establishing the serial-parallel control flow graph in the information collection module comprises:

the establishing process of the serial-parallel control flow graph takes a function as a basic processing unit; for the whole source program, if a serial-parallel control flow graph of a function can be established, combining the function call relation information collected in claim 2, the serial-parallel control flow graph of the whole source program can be recursively established;

for a function, based on the information gathered in claim 2, a serial-parallel control flow graph can be built by:

1) Establishing a serial node and a parallel node which are isolated one by using the node type information and the node range information;

2) Establishing directed edges among the nodes by using the relationship information among the nodes, and connecting the serial nodes and the parallel nodes into a graph;

3) And storing the function call information, the variable definition information, the variable reference information, the variable scope information, the memory allocation information and the memory release information to the corresponding serial node or parallel node according to the node range information.

5. The automatic management system for the complex pointer data structure oriented to the heterogeneous platform according to claim 4, wherein the complex pointer data is classified into three types according to the position of the pointer:

the device comprises a class nesting pointer, a vector container and a multi-level nesting pointer;

the class nesting pointer refers to a pointer contained in a class, and the vector container refers to a vector container provided in a C + + standard library; the multi-level nested pointers refer to pointers at two levels and pointers above two levels.

6. The heterogeneous platform-oriented automatic management system for the complex pointer data structure of the claim 5 is characterized in that the class nesting pointer is processed by the following steps:

(1) recursively traversing the serial-parallel control flow graph established in the claim 4, and finding out the class nesting pointers defined in the serial nodes and referred in the parallel nodes and the C + + classes thereof according to the variable definitions and reference information stored in the serial/parallel nodes;

(2) modifying the type definition of the type in the source code for the C + + type found in the step (1) so that the type of the type inherits the UMMaper base class provided by the runtime during definition;

(3) modifying the memory allocation and release statement of the pointer in the source code for the class nested pointer found in the step (1), allocating the memory by using cudaMallocManager () and releasing the memory by using cudaFree ();

(4) and (3) modifying the definition statement of the type instance in the source code for the C + + type instance found in the step (1), creating the type instance by using a reloaded new operator, and transmitting the memory space address distributed in the step (3) to a corresponding nested pointer in the C + + type instance.

7. The system according to claim 5, wherein the vector container is processed by the following steps:

(1) recursively traversing the serial-parallel control flow graph established in claim 4, and finding a vector container defined in the serial node and referred to in the parallel node according to the variable definition and reference information stored in the serial/parallel node;

(2) and (3) modifying the definition statement of the vector container instance found in the step (1), and inserting a display call for a custom allocator provided by the runtime.

8. The heterogeneous platform-oriented complex pointer data structure automatic management system according to claim 5, wherein the multi-level nested pointers are processed by the following steps:

(1) recursively traversing the serial-parallel control flow graph established in claim 4, and finding out multi-level nested pointers defined in the serial nodes and referred to in the parallel nodes and sub-pointers of the levels thereof according to variable definitions and reference information stored in the serial/parallel nodes;

(2) and (2) modifying memory allocation and release statements of all the multi-level nested pointers and all the levels of sub pointers found in the step (1), allocating memory by using cudaMallocManager () and releasing memory by using cudaFree ().

9. The system for automatically managing the complex pointer data structure oriented to the heterogeneous platform as claimed in claim 6, wherein the implementation process of the ummale class in the runtime module is as follows:

designing a UMMaper class to manage the memory allocation and release of the C + + class for processing the class nesting pointers; reloading a C + + default new and delete operator by using cudamallmanaged () and cudaFree () in the UMMaper class; the allocation, release and data transmission operations of the derived classes of the UMMaper class on the memory space of the CPU and the GPU can be automatically managed by the unified memory.

10. The system according to claim 7, wherein the implementation process of the custom allocator class in the runtime module is as follows:

designing a self-defined allocator class to manage memory allocation and release of the vector container for processing the vector container; the vector container defaults to use an allocator space configurator of a C + + standard library to manage memory allocation and release, so that a user-defined allocator class can be realized based on a uniform memory, and automatic management of memory allocation and release of the vector container is realized; in the custom allocator class, realizing a function of _ allocator () in the allocator class based on cudaMallocManager (), and realizing a function of _ deallocate () in the allocator class based on cudaFree (); and displaying and calling the custom allocator class when the vector container declares, so that the allocation and release of the memory space of the vector container on the CPU and the GPU and the automatic management of data transmission operation can be realized by the unified memory.

Images