A method for automatic migration and optimization of heterogeneous parallel programs

Abstract

A method for automatic migration and optimization of heterogeneous parallel programs belongs to heterogeneous parallel program development technology. The invention is to realize automatic migration of CPU parallel programs, improve program performance while reducing the workload of developers, and thus solve the problems of parallel instruction conversion, data transmission management and optimization. Technical points: Construct a framework for an automatic migration system for heterogeneous parallel programs, which is used to automatically translate OpenMP CPU parallel programs into OpenMP Offloading heterogeneous parallel programs; formalize consistency state conversion, optimize transmission operations and reduce redundant data transmission under the premise of ensuring data consistency; design a runtime library, which is used to provide automatic data transmission management and optimization functions and maintain the consistency state of each variable memory area; design a source-to-source translator, which is used to automatically convert parallel instructions and automatically insert runtime APIs. This method can automatically identify CPU parallel instructions and convert them into accelerator parallel instructions, thereby improving program performance.

Description

Technical Field

[0001] The invention relates to a method for automatically transplanting and optimizing a heterogeneous parallel program, belonging to a heterogeneous parallel program development technology. Background Art

[0002] With the huge demand for computing power in different applications such as artificial intelligence, image processing, multi-physics field simulation, quantum simulation, and climate simulation, heterogeneous platforms based on various accelerators have replaced CPU (Central Processing Unit) as the main source of computing power. In the field of high-performance computing, GPU (Graphics Processing Unit) is mainly used as an accelerator, and on mobile platforms, GPU, DSP (Digital Signal Processing) or FPGA (Field Programmable Gate Array) is mainly used as an accelerator. While accelerators provide huge computing power, they also bring huge challenges to application development and transplantation.

[0003] The standard for CPU parallel programming is the OpenMP (Open Multi-Processing) model, while heterogeneous parallel programming requires the use of heterogeneous parallel programming models such as CUDA (Compute Unified Device Architecture), OpenCL (Open Computing Language), OpenACC (Open Accelerators), and OpenMP Offloading (an extension of OpenMP). Writing efficient heterogeneous parallel programs often requires developers to be familiar with the characteristics of heterogeneous platforms and master parallel programming models; even if developers are capable, porting CPU parallel programs to heterogeneous platforms is a very time-consuming and error-prone task. Therefore, there is an urgent need for automatic porting tools to automatically port OpenMP CPU parallel programs to heterogeneous platforms.

[0004] The prior art with document number CN104035781A (CN104035781B) provides a method for rapidly developing heterogeneous parallel programs, which involves the performance analysis of CPU serial programs and the transplantation of heterogeneous parallel programs: first, the performance and algorithm analysis of the CPU serial program is performed to locate the program performance bottleneck and its parallelism; then, OpenACC pre-compilation instructions are inserted on the basis of the original code to obtain heterogeneous parallel code that can be executed in a heterogeneous parallel environment; the code is compiled and executed according to the parameters specified by the software and hardware platform, and whether further optimization is needed is determined based on the program running results. This prior art can efficiently parallelize existing programs, so that the programs can fully utilize the computing power of heterogeneous systems, and it is highly practical and easy to promote. However, the two major problems in the transplantation process: parallel instruction conversion, data transmission management and optimization have not yet been solved. Summary of the invention

[0005] The technical problems to be solved by the present invention are:

[0006] The purpose of the present invention is to provide a method for automatic migration and optimization of heterogeneous parallel programs to achieve automatic migration of CPU parallel programs, improve program performance while reducing the workload of developers, thereby solving the problems of parallel instruction conversion, data transmission management and optimization.

[0007] The technical solution adopted by the present invention to solve the above technical problems is:

[0008] A method for automatic migration and optimization of heterogeneous parallel programs, the implementation process of the method is as follows:

[0009] Step 1: Build a framework for the automatic migration system of heterogeneous parallel programs

[0010] Heterogeneous Parallel Program Automatic Migration System, referred to as OAO system, is used to automatically translate OpenMP CPU parallel programs into OpenMP Offloading heterogeneous parallel programs, and automatically manage and optimize data transmission between CPU and accelerator in combination with the runtime system; the OAO system framework mainly includes two parts: source-to-source translator and runtime library;

[0011] The runtime library contains three types of APIs: consistency state tracking, data transfer, and consistency state conversion. The consistency state tracking API captures the variable memory area and initializes the consistency state. The data transfer API dynamically determines and executes the transfer operation based on the current consistency state and the consistency constraints that need to be met, and updates the consistency state at the same time. The state conversion API updates the consistency state based on the read and write operation types.

[0012] The source-to-source translator translates OpenMP CPU parallel code into OpenMP Offloading heterogeneous parallel code and inserts appropriate runtime APIs. It consists of three modules: data transfer API insertion module, state transition API insertion module, and parallel instruction translation module. The parallel instruction translation module translates OpenMP CPU parallel instructions into OpenMP Offloading heterogeneous parallel instructions to obtain the OpenMP Offloading kernel. The data transfer API insertion module and the state transition API insertion module insert two corresponding runtime APIs respectively.

[0013] The working process of the OAO system is as follows: the OpenMP CPU parallel code is translated from source to source to obtain the OpenMP Offloading heterogeneous parallel code containing the runtime API. The OpenMP Offloading code is compiled and run on the heterogeneous platform, where the OpenMP Offloading kernel runs on the accelerator and the rest of the program runs on the CPU; the runtime library manages the data transmission between the CPU and the accelerator through the inserted API, ensures data consistency and dynamically optimizes the transmission to reduce redundant data transmission;

[0014] Step 2: Formalize the consistency state transition

[0015] For heterogeneous platforms, variables have copies in both CPU memory and accelerator memory. The consistency state is used to describe the validity of the CPU copy and the accelerator copy of the variable. The simplest state transition function is derived from the current consistency state and the consistency state constraints that need to be met. This function determines the simplest transmission operation type that needs to be performed through the corresponding relationship, so as to optimize the transmission operation and reduce redundant data transmission under the premise of ensuring data consistency.

[0016] Step 3: Runtime library design

[0017] The runtime library is used to provide automatic data transfer management and optimization functions, maintain the consistency state of each variable memory area, and provide three types of runtime library API functions and corresponding descriptions;

[0018] Step 4 Source-to-source translator design

[0019] The source-to-source translator is mainly based on the static analysis function of the Clang / LLVM compilation framework, which translates the OpenMP CPU parallel code into the OpenMP Offloading heterogeneous parallel code and inserts it into the runtime API in an appropriate way to optimize data transmission while ensuring data consistency.

[0020] The source-to-source translator collects information such as serial domains, parallel domains, and variable references through static analysis; then it builds the serial-parallel control flow graph of the program and binds the variable reference information to the corresponding parallel domain serial domain; then it inserts data transfer API insertion and state transition API insertion based on the granularity of serial domains and parallel domains; finally, it translates OpenMP CPU parallel instructions into OpenMP Offloading parallel instructions.

[0021] The present invention has the following beneficial technical effects:

[0022] The present invention proposes an automatic migration and optimization method from an OpenMP CPU parallel program to an OpenMP Offloading heterogeneous parallel program (OAO for short). Based on the OAO method, an automatic migration system can be implemented, which can improve program performance while reducing the workload of developers. The OAO method can automatically solve two major problems in the migration process: parallel instruction conversion and data transmission management and optimization. The method can automatically identify CPU parallel instructions and convert them into accelerator parallel instructions. By using a runtime system, the method can dynamically optimize data transmission during program execution, thereby reducing redundant data transmission while ensuring data consistency and improving program performance.

[0023] For a large number of existing OpenMP CPU parallel programs, developers can directly use the OAO system to achieve automatic heterogeneous migration and use the computing power of heterogeneous platforms to improve program performance. For new applications, developers can continue to use the familiar OpenMP model and then use the OAO system for automatic heterogeneous migration to improve program performance. BRIEF DESCRIPTION OF THE DRAWINGS

[0024] Figure 1 The framework flow chart of the heterogeneous parallel program automatic migration system (OAO system);

[0025] Figure 2 It is a consistency state transition relationship diagram;

[0026] Figure 3 It is a serial-parallel control flow graph;

[0027] Figure 4 This is a screenshot of the OpenMP CPU parallel program;

[0028] Figure 5 Screenshot of a manually written OpenMP Offloading heterogeneous parallel program;

[0029] Figure 6 Screenshot of the OpenMP Offloading heterogeneous parallel program obtained by source-to-source translation;

[0030] Figure 7 This is a bar chart of the acceleration ratio of other versions relative to the OMP version (K40 platform);

[0031] Figure 8 This is a bar chart showing the speedup ratio of other versions relative to the OMP version (on a 2080Ti platform);

[0032] Fig. 9 A bar chart showing the percentage of data transmission saved by the OAO version compared to other versions;

[0033] Fig.10 A bar chart showing the percentage of transfer time saved by the OAO version compared to other versions (on the 2080Ti platform).

[0034] Figures 7 to 10 The English terms and abbreviations are well-known terms in the art. DETAILED DESCRIPTION

[0035] The following is combined with Figures 1 to 10 The implementation of the present invention is described as follows:

[0036] 1. Heterogeneous parallel program automatic migration system framework

[0037] The Heterogeneous Parallel Program Automatic Migration System (OAO system for short) can automatically translate OpenMP CPU parallel programs into OpenMP Offloading heterogeneous parallel programs, and automatically manage and optimize the data transmission between the CPU and the accelerator in conjunction with the runtime system. The OAO system framework is as follows Figure 1 As shown, it mainly includes two parts (the shaded part in the figure): source to source translator and runtime library.

[0038] The runtime library provides automatic data transfer management and optimization functions, and maintains the consistency state of each variable memory area. It contains three types of APIs (Application Programming Interface): consistency state tracking, data transfer, and consistency state conversion. The consistency state tracking API captures the variable memory area and initializes the consistency state. The data transfer API dynamically determines and executes the transfer operation based on the current consistency state and the consistency constraints that need to be met, and updates the consistency state. The state conversion API updates the consistency state based on the read and write operation type.

[0039] The source-to-source translator translates OpenMP CPU parallel code into OpenMP Offloading heterogeneous parallel code and inserts the appropriate runtime API; it consists of three modules: data transfer API insertion module, state transition API insertion module, and parallel instruction translation module. The parallel instruction translation module translates OpenMP CPU parallel instructions into OpenMP Offloading heterogeneous parallel instructions to obtain the OpenMP Offloading kernel. The data transfer API insertion module and the state transition API insertion module insert two corresponding runtime APIs respectively.

[0040] The working process of the OAO system is as follows: OpenMP CPU parallel code is translated from source to source to obtain OpenMP Offloading heterogeneous parallel code containing runtime API. After compilation, the OpenMP Offloading code runs on the heterogeneous platform, where the OpenMP Offloading kernel runs on the accelerator and the rest of the program runs on the CPU. The runtime library manages the data transmission between the CPU and the accelerator through the inserted API, ensures data consistency and performs dynamic transmission optimization to reduce redundant data transmission.

[0041] 2. Formalization of Consistency State Transition

[0042] Data transfer management and optimization are key in the process of porting OpenMP CPU programs to heterogeneous platforms. This section formalizes variable consistency states, consistency state transitions, consistency state constraints, etc., to provide theoretical support for the subsequent runtime library and source-to-source translator design.

[0043] For heterogeneous platforms, a variable may have copies in both the CPU memory and the accelerator memory. The consistency state is used to describe the validity of the CPU copy and the accelerator copy of the variable. It is defined as follows:

[0044] Definition 1. The consistency state State is a 3-bit binary number, where Bit0 indicates whether the variable has an accelerator copy (Bit0=1) / no (Bit0=0); Bit1 indicates whether the CPU copy is valid (Bit1=1) / no (Bit1=0); Bit2 indicates whether the accelerator copy is valid (Bit2=1) / no (Bit2=0).

[0045] Table 1 All possible consistency states

[0046]

[0047] According to Definition 1, all possible consistency states are given as shown in Table 1. Data transmission between the CPU and the accelerator and the read and write operations on the CPU and the accelerator will change the consistency state. The consistency state transition function is used to represent the change of the consistency state, which is defined as follows:

[0048] Definition 2. The form of the consistency state transition function TransFunc is shown in formula (1), where InVld and Vld are a pair of 3-bit binary numbers, and the operator is a Boolean operator. If a bit in InVld is set to 0, the corresponding bit in inState can be converted to 0; if a bit in Vld is set to 1, the corresponding bit in inState can be converted to 1; therefore, setting different values ​​for the two can convert any inState to any outState. TransFunc is also abbreviated as formula (2).

[0049]

[0050] TransFunc={InVld,Vld} (2)

[0051] According to Definition 2, all possible consistency state transition functions and their corresponding data transmission or read / write operations are given as shown.

[0052] Table 2 All possible consistency state transition operations

[0053]

[0054]

[0055] According to Definition 1 and Definition 2, the conversion relationship diagram between various consistency states is given, such as Figure 2 shown.

[0056] In order to ensure program correctness, the read and write operations on the CPU and accelerator have certain requirements on the consistency state. These requirements are expressed using consistency state constraints, which are defined as follows:

[0057] Definition 3. The consistency state constraint Constr consists of a pair of 3-bit binary numbers, and its form is shown in (3). The default value of Constr is {111,000}, indicating that there is no requirement for the consistency state. You can set a bit of ConInVld to 0, or set a bit of ConVld to 1 to indicate different constraint requirements for the consistency state, as shown in Table 3.

[0058] Constr={ConInVld,ConVld} (3)

[0059] Table 3 Description of the meaning of ConVld and ConInVld

[0060]

[0061] According to Definition 3 and Table 3, the consistency state constraints required for different read and write operations on the CPU and accelerator are given, as shown in Table 4.

[0062] Table 4 All possible consistency constraints

[0063]

[0064]

[0065] From the current consistency state State and the consistency state constraint Constr = {ConInVld, ConVld} that needs to be satisfied, the simplest state transition function can be derived as follows:

[0066] The derivation formula of the simplest state transition function is shown in formula (4).

[0067] MinTrFunc(State)=State·InVld+Vld

[0068] in

[0069] The simplest state transition function represents the simplest non-redundant data transmission operation that must be performed from the consistency state State to satisfy the consistency state constraint Constr. Therefore, the MinTrFunc can be derived to determine the simplest transmission operation type that needs to be performed through the corresponding relationship in Table 2 (the first 6 types). In this way, the transmission operation is optimized and redundant data transmission is reduced under the premise of ensuring data consistency.

[0070] 3. Runtime library design

[0071] The runtime library API functions are shown in Table 5 and can be divided into three categories: consistency state tracking (the first 6), data transfer (OAODataTrans), and consistency state transition (OAOStTrans); they will be introduced separately below. Table 5 Runtime Library API Functions

[0072]

[0073]

[0074] 3.1 Consistency State Tracking API

[0075] The runtime library uses the variable memory area as the granularity for consistency state tracking and data transfer. In C / C++, the variable memory area is a continuous memory area, which may come from local variable definitions, global variable definitions, malloc operations, new operations, etc. In order to record and track the variable consistency state, the variable memory area and memory environment are formally defined as follows:

[0076] Definition 4. The variable memory area MemBlk is a four-tuple, as shown in formula (5); where Begin represents the memory start address, Length represents the length of the memory area, ElemSize represents the element size, and State represents the consistency state.

[0077] MemBlk={Begin,Length,ElemSize,State} (5)

[0078] Definition 5. The memory environment MemEnv is the set of all variable memory areas, as shown in formula (6).

[0079] MemEnv = {MemBlk 1 ,…,MemBlk n} (6)

[0080] The runtime library defines MemEnv as a global variable and maintains MemEnv during the execution of the entire OpenMP Offloading heterogeneous parallel program. When a variable is referenced, the corresponding pointer ptr can be used to search in MemEnv, and the one that satisfies formula (7) is the referenced MemBlk.

[0081] Begin≤ptr≤Begin+Length-1 (7)

[0082] The consistency state tracking API will be inserted into the source code in an appropriate manner during the source-to-source translation process. The OAOSaveArrayInfo function is inserted after the local variable declaration or at the beginning of the main function (for global variables) to collect variable information. The OAOMalloc function replaces the malloc function, collects memory allocation information, and performs memory allocation. OAONewInfo is inserted after the new operation to collect memory allocation information. The above three functions use the collected information to create the corresponding MemBlk and initialize State to HOST_ONLY.

[0083] OAODeleteArrayInfo is inserted at the end of the variable scope and the end of the main function (for global variables), and the corresponding MemBlk is deleted. OAOFree replaces the free function, releases the memory area and deletes the corresponding MemBlk. OAODeleteInfo is inserted after the delete operation, releases the memory area and deletes the corresponding MemBlk.

[0084] In addition, the runtime is specially optimized for NVIDIA devices. When the memory to be allocated is larger than 128KB, cudaMallocHost() is used instead of malloc() to allocate memory.

[0085] 3.2 Data Transfer API

[0086] As analyzed in Table 4, variable read and write operations require variables to meet certain consistency state constraints. Therefore, before accessing a variable, it is necessary to call the data transfer API, namely the OAODataTrans function, to perform the simplest data transfer operation to meet the corresponding constraints. This section explains the principle of the data transfer API, and its insertion method will be described in Section 4. The OAODataTrans function uses the following Algorithm 1 to determine the simplest state transition function required to meet the consistency state constraint Constr, and executes the corresponding simplest data transfer operation, while updating the consistency state. Lines 1 to 2 of Algorithm 1 find the corresponding variable memory area MemBlk through the variable pointer ptr; Line 3 derives the simplest state transition function MinTrFunc through formula (4) based on the consistency constraint Constr that needs to be met and the current consistency state State; Line 4 executes the data transfer operation corresponding to MinTrFunc in Table 2; Line 5 uses MinTrFunc to update State according to formula (1).

[0087]

[0088] 3.3 Consistency State Transition API

[0089] As analyzed in Table 2 (the last three rows), read and write operations may change the consistency state of the variable. Therefore, after a variable is accessed, the consistency state transition API, namely the OAOStTrans function, needs to be called to update the consistency state State stored in the variable memory area MemBlk. This section explains the principle of the consistency state transition API, and its insertion method will be explained in Section 4. The OAOStTrans function uses the following Algorithm 2 to complete the consistency state transition process. Lines 1 and 2 of Algorithm 2 find the corresponding variable memory area MemBlk through the variable pointer ptr; Line 3 uses the state transition function StTrans according to formula (1) to update the consistency state State contained in MemBlk.

[0090]

[0091] 4 Source-to-Source Translator Design

[0092] The source-to-source translator is mainly based on the static analysis function of the Clang / LLVM (a C / C++ language compiler) compilation framework. It translates OpenMPCPU parallel code into OpenMPOffloading heterogeneous parallel code and inserts it into the runtime API in an appropriate way to optimize data transmission while ensuring data consistency.

[0093] The source-to-source translator collects information such as serial domain (Definition 6), parallel domain (Definition 7), and variable references through static analysis; then it builds the serial-parallel control flow graph of the program (Definition 8) and binds the variable reference information to the corresponding parallel domain serial domain (Definition 9-10); then it inserts data transfer API insertion and state transition API insertion at the granularity of serial domain and parallel domain; finally, it translates OpenMPCPU parallel instructions into OpenMPOffloading parallel instructions.

[0094] Definition 6: The serial domain SEQ is a code fragment outside the scope of #pragmaompparallel, without internal branches, and executed serially, SEQ, which is also called a serial node in the serial-parallel control flow graph.

[0095] Definition 7: A parallel region OMP is a code segment within the scope of #pragmaompparallel that is executed in parallel. It is also called a parallel node in the serial-parallel control flow graph.

[0096] Definition 8 The definition of the serial-parallel control flow graph SPGraph is shown in formula (8), which is a special control flow graph of a function, in which the nodes are serial domains or parallel domains; examples of serial-parallel control flow graphs are as follows Figure 3 shown.

[0097] SPGraph=(NodeGrp,CtlEdge)

[0098] NodeGrp=SEQGrp∪OMPGrp

[0099] CtlEdge={<x,y> |x,y∈NodeGrp} (8)

[0100] SEQGrp = {SEQ 1 ,…,SEQ n}

[0101] OMPGrp = {OMP 1 ,…,OMP m}

[0102] Definition 9: The definition of variable reference list RefList is shown in formula (9), which is a reference list of a variable in a serial domain or parallel domain.

[0103] RefList{Ref 1 ,...,Ref p}

[0104]

[0105] Definition 10 The definition of the variable reference information table NodeVarRef is shown in formula (10), which is the collection of all variable information in a serial domain or parallel domain. Each serial domain or parallel domain is bound to its corresponding NodeVarRef, as shown Figure 3 .

[0106] NodeVarRef = {RefList 1 ,…,RefList l} (10)

[0107] Function calls in serial or parallel domains require special handling. Function calls in serial domains are separated into a separate serial domain. For function arguments passed in as copies, RefList = {R}. For function arguments passed in as pointers or references, RefList = {R} if there is no write operation in the called function, or RefList = {RW} if there is a write operation in the called function.

[0108] For a function call in a parallel domain, it is considered as an access to the function argument. For a function argument passed in as a copy, {R} is inserted at the appropriate position in its RefList. For a function argument passed in as a pointer or reference, {R} is inserted at the appropriate position in its RefList if there is no write operation in the called function; or {RW} if there is a write operation in the called function.

[0109] Based on the above program abstract representation, the design of three main functions of the source-to-source translator is given.

[0110] 4.1 Data Transfer API Insertion

[0111] As mentioned above, the required data transfer APIs need to be inserted between most serial domains and parallel domains. However, for functions that have been called in parallel domains (referred to as OMP call functions), runtime APIs cannot be inserted. This is because OMP call functions will run on accelerators, and runtime APIs cannot run on accelerators. The data consistency of such functions is guaranteed by the runtime APIs before and after the function call.

[0112] The data transfer API insertion algorithm is designed as follows. The source-to-source compiler uses the following Algorithm 3 to process each non-OMP call function and insert the required data transfer API before the serial domain and parallel domain. Lines 01 to 02 of Algorithm 3 are double loops, which apply the subsequent algorithm steps to each applied variable in each serial domain and parallel domain in the control flow graph; Line 03 uses static analysis to obtain the pointer ptr of the referenced variable Var. Lines 04 to 15 are processed separately for different situations. Lines 04 to 05 insert the OAODataTrans(ptr,ConOMPR) statement before the parallel domain. Lines 06 to 11 process the serial domain separated from the function call in two cases. If the called function is an OMP call function, the OAODataTrans(ptr,ConSEQR) statement is inserted before the serial domain (lines 07 to 08); otherwise, for the function parameter variable passed in the copy, the OAODataTrans(ptr,ConSEQR) statement is inserted before the serial domain. Line 15: For other cases, insert the OAODataTrans(ptr,ConSEQR) statement before the serial field.

[0113]

[0114]

[0115] 4.2 State Transition API Insertion

[0116] After a variable is accessed through a serial domain or a parallel domain, the consistency state may change, and the required state transition API needs to be inserted to update the consistency state saved in the runtime. The state transition API insertion algorithm is designed as follows. The source-to-source compiler uses the following Algorithm 4 to process each non-OMP calling function and insert the required state transition API after the serial domain and parallel domain. Lines 01 to 02 of Algorithm 4 are double loops, which apply the subsequent algorithm steps to each applied variable in each serial domain and parallel domain in the control flow graph; Line 03 uses static analysis to obtain the pointer ptr of the referenced variable Var. Lines 04 to 16 are processed separately for different situations. If Node is a parallel domain and the RefList corresponding to the variable Var contains W (write operation), then the OAOStTrans(ptr,TrOMPW) statement is inserted after the parallel domain (lines 04 to 06). If Node is a serial domain and is separated from a function call, and the called function is an OMP calling function and the RefList corresponding to the variable Var contains W (write operation), then insert the OAOStTrans(ptr,TrSEQW) statement after the serial domain (lines 08 to 11). If it is neither of the above two situations, and the RefList corresponding to the variable Var contains W (write operation), then insert the OAOStTrans(ptr,TrSEQW) statement after the serial domain.

[0117]

[0118]

[0119] 4.3 Parallel Instruction Translation

[0120] This porting framework targets the OpenMP work-sharing parallel mode. The correspondence between the parallel instructions of this mode on the CPU and the accelerator is shown in Table 6. The source-to-source compiler uses the following Algorithm 5 to translate the OpenMP CPU parallel instructions into OpenMP Offloading parallel instructions, thereby obtaining the OpenMP Offloading computing kernel. Line 01 of Algorithm 5 is a loop that applies the subsequent algorithm steps to each parallel domain; Line 02 obtains the OpenMP CPU parallel instructions through static analysis; Lines 03-04 replace the OpenMP CPU parallel instructions with the corresponding OpenMP Offloading parallel instructions according to Table 6.

[0121] Table 6 Parallel instruction correspondence

[0122]

[0123] The technical effects produced by the present invention are described as follows:

[0124] 1. Source-to-source translation result example

[0125] Translated using the OAO automatic transplantation system proposed in this patent Figure 4 The OpenMP CPU parallel program shown in Figure 1. Figure 6 As shown in the figure, OpenMP Offloading heterogeneous parallel programs. Figure 4 The OpenMP CPU parallel program shown in the figure, manually write the OpenMP Offloading heterogeneous parallel program, such as Figure 5 As shown. It can be seen from observation that Figure 5 The following transmissions in the manually written program are redundant: line 05 for the "from" transmission of v1, v2, and v3, line 13 for the "to" transmission of v3, and line 13 for the "from" transmission of v4. Figure 5 The automatically translated program uses the OAO runtime library to avoid the above redundant operations. From this example, we can see that the OAO system can successfully translate the OpenMP CPU parallel program into the OpenMP Offloading heterogeneous parallel program.

[0126] 2. Method evaluation

[0127] 2.1 Experimental methods

[0128] Polybench and Rodinia are commonly used benchmarks in heterogeneous computing. DawnCC is currently the most advanced source-to-source translator for generating OpenMP Offloading programs. We use Polybench and Rodinia to evaluate the performance of the OAO system, and use DawnCC as a control.

[0129] In order to compare the optimization capabilities of DawnCC and OAO for inter-procedural data transmission, we added a test program FDTD-2D-FUNC. FDTD-2D-FUNC is based on FDTD-2D in Polybench; each calculation kernel in FDTD-2D is encapsulated into a sub-function, and the variables required by the kernel are passed through function parameters; in this way, inter-procedural data transmission is constructed.

[0130] First we generate the following four versions of the program:

[0131] OMP version: OpenMP CPU parallel version program;

[0132] Manual version: OpenMP Offloading program obtained by manually translating the OMP version;

[0133] DawnCC version: OpenMP Offloading program obtained by translating the OMP version using DawnCC;

[0134] OAO version: OpenMP Offloading program obtained by translating the OMP version using OAO.

[0135] The two experimental software and hardware platforms are shown in Table 7 Experimental software and hardware platforms.

[0136] Table 7 Experimental software and hardware platform

[0137]

[0138] 2.2 Performance Evaluation

[0139] like Figure 7 and Figure 8 The figure shows the speedup of different versions of OpenMP Offloading programs relative to the OMP version. The "X" in the figure indicates a program that DawnCC cannot convert correctly.

[0140] It can be seen that OAO can handle all 23 test programs, while DawnCC can only handle 15 test programs. The OAO version has improved performance on 9 programs on the K40 platform and 15 programs on the 2080Ti platform, with a maximum speedup of 32 times; and the OAO version has the best performance on all test programs on all platforms compared to the OMP version and the manual version.

[0141] 2.3 Data transmission optimization evaluation

[0142] Table 8 shows the comparison of data transfer times of different versions of OpenMP Offloading programs, where “—” indicates programs that DawnCC cannot process correctly. Obviously, the OAO version has the best data transfer optimization effect on all test programs, that is, the number of data transfers is the least.

[0143] Observing FDTD-2D and FDTD-2D-FUNC, both OAO and DawnCC can optimize FDTD-2D, but DawnCC cannot optimize FDTD-2D-FUNC well, while OAO can optimize the number of transmissions of both to the optimal 5. This phenomenon shows that OAO can optimize the data transmission between processes, while DawnCC cannot.

[0144] Table 8 Comparison of data transmission times

[0145]

[0146]

[0147] like Fig. 9The figure shows the percentage of data transmission saved by the OAO version compared to the manual version and the DawnCC version. The OAO version transmits less data than the manual version in all the tested programs, especially FDTD-2D-FUNC and FDTD-2D (close to 100%). Compared to the DawnCC version, the OAO version has significant data transmission savings in 6 test programs, especially FDTD-2D-FUNC (close to 100%).

[0148] Fig.10 The percentage of transfer time saved by the OAO version compared to the manual version and the DawnCC version is shown. The OAO version has significant transfer time savings on all test programs relative to the other two versions.

[0149] The above data and analysis prove that the OAO system can automatically translate OpenMP CPU parallel programs into OpenMP Offloading heterogeneous parallel programs, while optimizing data transfer; the OAO version has achieved significant performance improvements compared to the manual version. Compared with DawnCC, the OAO system can handle more input programs, perform more extensive data transfer optimization, and has better performance on all test programs.

Claims

1. A method for automatic migration and optimization of heterogeneous parallel programs, characterized in that: The implementation process of the method is: Step 1: Build a framework for the automatic migration system of heterogeneous parallel programs Heterogeneous Parallel Program Automatic Migration System, referred to as OAO system, is used to automatically translate OpenMP CPU parallel programs into OpenMP Offloading heterogeneous parallel programs, and automatically manage and optimize data transmission between CPU and accelerator in combination with the runtime system; the OAO system framework mainly includes two parts: source-to-source translator and runtime library; The runtime library contains three types of APIs: consistency state tracking, data transmission, and consistency state conversion; The consistency state tracking API captures the variable memory area and initializes the consistency state; The data transfer API dynamically determines and executes the transfer operation based on the current consistency state and the consistency constraints that need to be met, and updates the consistency state at the same time; The state transition API updates the consistency state based on the read and write operation types; The source-to-source translator translates the OpenMP CPU parallel code into the OpenMP Offloading heterogeneous parallel code and inserts the runtime API; It consists of three modules: data transfer API insertion module, state transition API insertion module, and parallel instruction translation module; the parallel instruction translation module translates OpenMP CPU parallel instructions into OpenMP Offloading heterogeneous parallel instructions to obtain the OpenMP Offloading kernel; the data transfer API insertion module and the state transition API insertion module insert two corresponding runtime APIs respectively; The working process of the OAO system is as follows: the OpenMP CPU parallel code is translated from source to source to obtain the OpenMP Offloading heterogeneous parallel code containing the runtime API. The OpenMP Offloading code is compiled and run on the heterogeneous platform, where the OpenMP Offloading kernel runs on the accelerator and the rest of the program runs on the CPU; the runtime library manages the data transmission between the CPU and the accelerator through the inserted API, ensures data consistency and dynamically optimizes the transmission to reduce redundant data transmission; Step 2: Formalize the consistency state transition For heterogeneous platforms, variables have copies in both CPU memory and accelerator memory. The consistency state is used to describe the validity of the CPU copy and the accelerator copy of the variable. The simplest state transition function is derived from the current consistency state and the consistency state constraints that need to be met. This function determines the simplest transmission operation type that needs to be performed through the corresponding relationship, so as to optimize the transmission operation and reduce redundant data transmission under the premise of ensuring data consistency. Step 3: Runtime library design The runtime library is used to provide automatic data transfer management and optimization functions, maintain the consistency state of each variable memory area, and provide three types of runtime library API functions and corresponding descriptions; Step 4 Source-to-source translator design The source-to-source translator is mainly based on the static analysis function of the Clang / LLVM compilation framework, which translates the OpenMP CPU parallel code into the OpenMP Offloading heterogeneous parallel code and inserts the runtime API to optimize data transmission while ensuring data consistency. The source-to-source translator collects information including serial domain, parallel domain, and variable reference through static analysis; then builds the serial-parallel control flow graph of the program and binds the variable reference information to the corresponding parallel domain and serial domain; then inserts data transfer API insertion and state transition API insertion based on the granularity of serial domain and parallel domain; finally, translates OpenMP CPU parallel instructions into OpenMP Offloading parallel instructions. The specific process is as follows: Define the serial domain SEQ as a code fragment outside the scope of #pragma omp parallel, without internal branches, and executed serially, SEQ, which is also called a serial node in the serial-parallel control flow graph; Define a parallel region OMP as a code fragment that is executed in parallel within the scope of #pragma omp parallel, which is also called a parallel node in the serial-parallel control flow graph; Definition of the serial-parallel control flow graph SPGraph is shown in formula (1), which is a special control flow graph of a function, in which the nodes are serial domains or parallel domains; SPGraph=(NodeGrp,CtlEdge) NodeGrp=SEQGrp∪OMPGrp CtlEdge={<x,y> |x,y∈NodeGrp} SEQGrp={SEQ1,…,SEQ n } OMPGrp={OMP1,…,OMP m } (1) The definition of the variable reference list RefList is shown in formula (2), which is a reference list of a variable in a serial domain or a parallel domain. The definition of the variable reference information table NodeVarRef is shown in formula (3), which is a collection of all variable information in a serial domain or parallel domain. Each serial domain or parallel domain is bound to its corresponding NodeVarRef. NodeVarRef={RefList1,…,RefList l } (3) Function calls in serial domains or parallel domains require special processing. Function calls in serial domains are separated into an independent serial domain. For function arguments passed in as copies, its RefList = {R}; for function arguments passed in as pointers or references, its RefList = {R}, when there is no write operation in the called function; or RefList = {RW}, when there is a write operation in the called function; for function calls in parallel domains, they are regarded as accesses to function arguments. For function arguments passed in as copies, {R} is inserted in its RefList; for function arguments passed in as pointers or references, {R} is inserted in its RefList, when there is no write operation in the called function; or {RW}, when there is a write operation in the called function.

2. The method for automatic migration and optimization of heterogeneous parallel programs according to claim 1, characterized in that: In step 2, the formal implementation process of the consistency state transition is: Given the following definition: Definition 1. The consistency state State is a 3-bit binary number; Where Bit0=1 indicates that the accelerator copy of the variable exists, Bit0=0 indicates that the accelerator copy of the variable does not exist; Bit1=1 indicates that the CPU copy of the variable is valid, Bit1=0 indicates that the CPU copy of the variable is invalid; Bit2=1 indicates that the accelerator copy of the variable is valid, Bit2=0 indicates that the accelerator copy of the variable is invalid; According to Definition 1, all possible consistency states are obtained as follows: 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; Data transmission between the CPU and the accelerator and the read and write operations on the CPU and the accelerator will change the consistency state. The consistency state transition function is used to represent the change of the consistency state, which is defined as follows: Definition 2. The form of the consistency state transition function TransFunc is shown in formula (4), where InVld and Vld are a pair of 3-bit binary numbers, and the operator is a Boolean operator; if a bit in InVld is set to 0, the corresponding bit in inState is converted to 0; if a bit in Vld is set to 1, the corresponding bit in inState is converted to 1; therefore, setting different values ​​for the two will convert any inState to any outState; TransFunc is also abbreviated as formula (5); outState=TransFunc(inState) (4) =inState·InVld+Vld TransFunc={InVld,Vld} (5) According to Definition 2, all possible consistency state transition functions and their corresponding data transmission or read / write operations are obtained as follows: Data transfer or read and write operations: TransFunc; No data transmission required: 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 write operation: TrSEQW = {011,010}; Accelerator write operation TrOMPW = {101, 100}; CPU or accelerator read operation: TrRead = {111,000}; According to Definition 1 and Definition 2, the conversion relationship between various consistency states is obtained; To ensure program correctness, the read and write operations on the CPU and accelerator have certain requirements on the consistency state. These requirements are expressed using consistency state constraints, which are defined as follows: Definition 3. The consistency state constraint Constr consists of a pair of 3-bit binary numbers, and its form is shown in (6). The default value of Constr is {111,000}, indicating that there is no requirement for the consistency state. Setting a bit of ConInVld to 0 or setting a bit of ConVld to 1 indicates different constraint requirements for the consistency state: Constr={ConInVld,ConVld} (6) The consistency state constraints of different bits and their values ​​are: ConVldBit0=1: the accelerator copy must exist; ConVldBit1=1: the CPU copy must be valid; ConVldBit2=1: the accelerator copy must be valid; ConInVldBit0=0: the accelerator copy must not exist; ConInVldBit=10: the CPU copy must be invalid; ConInVldBit2=0: the accelerator copy must be invalid; According to Definition 3 and the consistency state constraints of different bits and their values, the consistency state constraints required for different read and write operations on the CPU and accelerator are given as follows: The required consistency state constraints, the consistency state constraints required for read and write operations are: Variable not used: ConNo = {111,000}; CPU read operation: ConSEQR = {111,010}; Accelerator read operation: ConOMPR = {111, 101}; CPU write operation: ConSEQW = {111,000}; Accelerator write operation: ConOMPW = {111,001}; CPU free operation: ConFREE = {111, 101}; From the current consistency state State and the consistency state constraint Constr = {ConInVld, ConVld} that needs to be satisfied, the simplest state transition function is derived as follows: The derivation formula of the simplest state transition function is shown in formula (7): MinTrFunc(State)=State·InVld+Vld (7) in MinTrFunc determines the simplest type of transfer operation that needs to be performed through the first six corresponding relationships among all possible consistency state transition operations.

3. The method for automatic migration and optimization of heterogeneous parallel programs according to claim 2, characterized in that: In step 3, the specific process of runtime library design is: The runtime library API functions are divided into three categories: consistency state tracking API functions, data transmission API functions OAODataTrans, and consistency state conversion API functions OAOStTrans; the first six of the runtime library API functions are consistency state tracking API functions: void OAOSaveArrayInfo(void*ptr,size_t length,size_t ElementSize): save array memory area information; void OAODeleteArrayInfo(void*ptr): delete array memory area information; void*OAOMalloc(size_t length): save dynamic memory area information; void OAOFree(void*ptr): delete dynamic memory area information; void*OAONewInfo(void*ptr,size_t ElementSize,size_t ElementNum): save dynamic memory area information; void OAODeleteInfo(void*ptr): delete dynamic memory area information; void OAODataTrans(void*ptr,STATE_CONSTR Constr): dynamically determine the simplest data transfer operation and execute it, while updating the consistency state; void OAOStTrans(void*ptr,STATE_CONSTR StTrans): updates the consistency state; 3.1 Consistency State Tracking API The runtime library uses the variable memory area as the granularity for tracking the consistency state and data transmission. In C / C++, the variable memory area is a continuous memory area, which comes from local variable definitions, global variable definitions, malloc operations, and new operations. In order to record and track the consistency state of variables, the formal representation of the variable memory area and memory environment is defined as follows: Definition 4. The variable memory region MemBlk is a four-tuple, as shown in formula (8); where Begin represents the memory start address, Length represents the length of the memory region, ElemSize represents the element size, and State represents the consistency state; MemBlk={Begin,Length,ElemSize,State} (8) Definition 5. The memory environment MemEnv is the set of all variable memory areas, as shown in formula (6). MemEnv={MemBlk1,…,MemBlk n } (9) The runtime library defines MemEnv as a global variable and maintains MemEnv during the execution of the entire OpenMP Offloading heterogeneous parallel program. When a variable is referenced, the corresponding pointer ptr is used to search in MemEnv, and the one that satisfies formula (10) is the referenced MemBlk: Begin≤ptr≤Begin+Length-1 (10) The consistency state tracking API inserts source code during the source-to-source translation process. The OAOSaveArrayInfo function is inserted after the local variable declaration or at the beginning of the main function to collect variable information; the OAOMalloc function replaces the malloc function, collects memory allocation information, and performs memory allocation; OAONewInfo is inserted after the new operation to collect memory allocation information. The above three functions use the collected information to create a corresponding MemBlk and initialize State to HOST_ONLY; OAODeleteArrayInfo is inserted at the end of the variable scope and the end of the main function, and the corresponding MemBlk is deleted; OAOFree replaces the free function, releases the memory area and deletes the corresponding MemBlk; OAODeleteInfo is inserted after the delete operation, releases the memory area and deletes the corresponding MemBlk; When the memory to be allocated is larger than 128KB, use cudaMallocHost() instead of malloc() to allocate memory; 3.2 Data Transfer API Variable read and write operations require variables to meet certain consistency state constraints. Before accessing a variable, the data transfer API, namely the OAODataTrans function, needs to be called to perform the simplest data transfer operation to meet the corresponding constraints. The OAODataTrans function uses the following algorithm 1 to determine the simplest state transition function required to meet the consistency state constraint Constr, and performs the corresponding simplest data transfer operation, while updating the consistency state. Algorithm 1 is a data transmission algorithm; input: variable pointer ptr, consistency state constraint Constr that the variable needs to satisfy, variable memory environment set MemEnv; for each variable memory area MemBlk in the variable memory environment set MemEnv, execute the following loop: If ptr satisfies formula (10), then execute: Substitute Constr and MemBlk's State into formula (7) to derive MinTrFunc; Find the simplest data transfer operation corresponding to MinTrFunc and execute it; According to formula (4), use MinTrFunc to update MemBlk's State; Then jump out of this loop; If ptr does not satisfy formula (10), it will skip the current loop and enter the next loop; The loop ends; 3.3 Consistency State Transition API After a variable is accessed, the consistency state transition API, namely the OAOStTrans function, needs to be called to update the consistency state State stored in the variable memory area MemBlk. The OAOStTrans function uses the following algorithm 2 to complete the consistency state transition process; Algorithm 2 is a consistency state transition algorithm, with inputs: variable pointer ptr, state transition function StTrans, variable memory region set MemEnv; for each variable memory region MemBlk in the variable memory environment set MemEnv, execute the following loop: If ptr satisfies formula (10), then execute: According to formula (4), use StTrans to update the State of MemBlk; then jump out of this loop; If ptr does not satisfy formula (10), it will skip the current loop and enter the next loop; The loop ends.

4. The method for automatic migration and optimization of heterogeneous parallel programs according to claim 3, characterized in that: Based on the program abstract representation described in claim 3, the design of three main functions of the source-to-source translator is given, specifically: 4.1 Data Transfer API Insertion The data transfer API insertion algorithm is designed as follows: The source-to-source compiler uses the following Algorithm 3 to process each non-OMP call function and insert the required data transfer API before the serial domain and parallel domain therein; Algorithm 3 is the data transmission API insertion algorithm; input: the serial-parallel control flow graph SPGraph of the non-OMP calling function; For each Node in the collection NodeGrp, perform the following loop: For each variable Var referenced in Node, the following loop is executed: Use static analysis to obtain the pointer ptr of Var; If Node is a parallel domain, execute: insert OAODataTrans(ptr,ConOMPR) before Node; If Node is a serial domain and Node is separated from a function call, execute: if the called function is an OMP calling function, insert OAODataTrans(ptr,ConSEQR) before Node; otherwise, if the incoming copy is a copy, insert OAODataTrans(ptr,ConSEQR) before it; If Node is a serial domain and Node is not separated from a function call, execute: insert OAODataTrans(ptr,ConSEQR) before Node; End the inner loop; End the outer loop; Execution ends; 4.2 State Transition API Insertion The state transition API insertion algorithm is designed as follows: the source-to-source compiler uses the following Algorithm 4 to process each non-OMP call function and insert the required state transition API after the serial domain and parallel domain therein; Algorithm 4 is the state transition API insertion algorithm; input: the serial-parallel control flow graph SPGraph of the non-OMP calling function; For each Node in the collection NodeGrp, perform the following loop: For each variable Var referenced in Node, the following loop is executed: Use static analysis to obtain the pointer ptr of Var; If Node is a parallel domain, execute: If the RefList corresponding to Var contains W, insert OAOStTrans (ptr, TrOMPW) after Node; If Node is a serial domain and Node is detached from a function call, then execute: If the called function is an OMP calling function, execute: If the RefList corresponding to Var contains W, insert OAOStTrans (ptr, TrSEQW) after Node; If Node is a serial domain and Node is not separated from a function call, then execute: If the RefList corresponding to Var contains W, insert OAOStTrans (ptr, TrSEQW) after Node; End the inner loop; End the outer loop; Execution ends; 4.3 Parallel Instruction Translation This porting framework targets the OpenMP work-sharing parallel mode, in which the parallel instruction correspondence between the CPU and the accelerator is based on the parallel instruction correspondence. The source-to-source compiler uses the following algorithm 5 to translate the OpenMP CPU parallel instructions into the OpenMP Offloading parallel instructions, thereby obtaining the OpenMP Offloading computing kernel; Replace #pragma omp parallel for with #pragma omp target teams distributeparallel for; Replace #pragma omp parallel loop with #pragma omp target teams distributeparallel loop; Replace #pragma omp parallel simd with #pragma omp target teams distributeparallel simd; Algorithm 5 is a parallel instruction translation algorithm; Input: serial-parallel control flow graph SPGraph; For each OMP in the OMPGrp set of SPGraph, execute the following loop: use static analysis to obtain OpenMPCPU parallel instructions; according to the parallel instruction correspondence, replace the OpenMP CPU parallel instructions with the corresponding OpenMPOffloading parallel instructions; the loop ends.

Images