Highly compatible heterogeneous compilation method and heterogeneous compiler

Abstract

The present invention relates to a highly compatible heterogeneous compilation method and a heterogeneous compiler. The method includes compiling and processing host-side functions in the split host-side source code based on the host-side architecture information, so as to generate a corresponding host-side bytecode file based on each host-side source code after the compilation process; in device-side compilation, compiling and processing the device-side functions in the split device-side source code, so as to generate a host-side bytecode file based on all device-side source codes after the compilation process; linking the host-side bytecode file generated in the device-side compilation and all the host-side bytecode files generated in the host-side compilation into a host-side executable file adapted to the host-side architecture information. The present invention supports the use of device-side functions not defined in the source code, can achieve compatibility with general-purpose processors of different architectures, and reduce the programming difficulty and maintenance cost of heterogeneous systems.

Description

Technical Field

[0001] The present invention relates to a heterogeneous compilation method and a heterogeneous compiler, in particular to a highly compatible heterogeneous compilation method and a heterogeneous compiler. Background Art

[0002] With the rapid development of computer technology, the traditional CPU (Central Processing Unit) processor system can no longer meet the demand for computing resources in fields such as weather forecasting, engineering simulation, and artificial intelligence represented by deep learning, thus prompting a revolutionary leap in the computer processor system from single-core to multi-core, and from homogeneous to heterogeneous.

[0003] Multi-core refers to the integration of two or more complete computing engines in one processor, and heterogeneous refers to the combination of two or more different types of microprocessor or microcontroller architectures. Different from traditional homogeneous multi-core processors, heterogeneous multi-core processors are generally composed of general-purpose microprocessors (CPUs) and dedicated acceleration processors (such as GPUs (Graphics Processing Units), MICs (Many Integrated Cores), and DSPs (Digital Signal Processing)), with the general-purpose microprocessors controlling and managing the operation of the dedicated acceleration processors.

[0004] General-purpose microprocessors provide general computing power and control operations, and are suitable for a wide range of applications. Special-purpose acceleration processors are customized for specific types of workloads and provide powerful computing performance. In heterogeneous processors, since different processors often have different instruction set structures, a compilation mechanism that can efficiently and conveniently handle heterogeneous programming is particularly important. It can not only perform different processing for different architectures, but also simplify the compilation process for developers.

[0005] In the early days of heterogeneous programming, programmers were required to write the code for the general-purpose microprocessor (host side) and the code for the dedicated processor (device side) separately, and then manually call the corresponding compiler for compilation. The binary code of the device side also needed to be manually embedded into the binary code generated by the host side. This compilation method can be called separate compilation. The overall compilation process of separate compilation is cumbersome, code writing is difficult, and programming efficiency is low.

[0006] The emergence of heterogeneous compilation has greatly reduced the compilation process of heterogeneous programming and the difficulty of writing heterogeneous programs. Currently, the mainstream heterogeneous programming support mainly includes NVIDIA's CUDA and AMD's HIP, and the corresponding heterogeneous compilers are NVCC and HIPCC. Heterogeneous compilers process the workflow of heterogeneous programming similar to homogeneous compilers, providing great convenience for related developers.

[0007] Heterogeneous compilers have now achieved shielding against the differences caused by processor heterogeneity. They analyze specific keywords in the source code to distinguish between the general-purpose microprocessor (host side) and dedicated acceleration processor (device side) parts of the source code, and automatically call the corresponding compilers to compile them separately, automatically splitting the entire compilation process into host-side compilation and device-side compilation.

[0008] Figure 1 This is a heterogeneous compiler compilation method taking the heterogeneous compiler HIPCC as an example. The compilation process of the heterogeneous compiler NVCC is similar to the compilation method of the heterogeneous compiler HIPCC.

[0009] From the above description, the separate compilation method has cumbersome compilation process steps, low programming efficiency, and is extremely unfriendly to code developers. The current mainstream heterogeneous compilers have effectively reduced the programming difficulty of heterogeneous systems, but their use has great limitations. Specifically, the limitations include:

[0010] 1) Mainstream heterogeneous compilation methods cannot meet the needs of calling device-side functions across source files, which means that the same device-side function needs to be defined multiple times when used in different source files, resulting in a large amount of duplicate code. 2) Mainstream heterogeneous compilers are only applicable to specific GPUs and can only meet the needs of specific architectures. This is because the architectures of different dedicated acceleration processors are different, and different architectures have their own unique instruction sets. Therefore, heterogeneous compilers need to be redesigned and implemented for other dedicated acceleration processors. 3) Some manufacturers' dedicated acceleration processors are paired with general-purpose microprocessors of multiple architectures. Each general-purpose microprocessor requires a separate set of code to be maintained, which is costly in terms of manpower and material resources and difficult to maintain. Summary of the invention

[0011] The purpose of the present invention is to overcome the deficiencies in the prior art and provide a highly compatible heterogeneous compilation method and a heterogeneous compiler, which support the use of device-side functions that are not defined in the source code, can achieve compatibility with general-purpose processors of different architectures, and reduce the programming difficulty and maintenance cost of heterogeneous systems.

[0012] According to the technical solution provided by the present invention, a highly compatible heterogeneous compilation method is provided, the heterogeneous compilation method comprising:

[0013] Providing source code to be compiled, wherein a keyword representing attribute information of the current function is configured for a function in the source code;

[0014] Based on the keywords of each function in the source code, split the source code into a plurality of host-side source codes and a plurality of device-side source codes;

[0015] Based on the architecture information of the host side, in the host side compilation, the host side functions in the split host side source code are compiled, so as to generate a corresponding host side bytecode file based on each host side source code after the compilation process;

[0016] In the device-side compilation, the device-side functions in the split device-side source code are compiled, so as to generate a host-side bytecode file based on all the device-side source codes after the compilation process;

[0017] Link the host-side bytecode files generated in the device-side compilation and all the host-side bytecode files generated in the host-side compilation into a host-side executable file that adapts to the host-side architecture information.

[0018] The keywords configured for the function in the source code include the __host__ keyword, the __global__ keyword, and the __device__ keyword, wherein the function modified by the __global__ keyword is a kernel function and the kernel function contains the call execution entry of the device-side function in the target source code, the function modified by the __host__ keyword is a function running on a general-purpose microprocessor; the function modified by the __device__ keyword is a function that can only be called and run on a dedicated processor;

[0019] When the __host__ keyword and the __device__ keyword are used to modify the same function, the modified function can be called and executed on a general-purpose microprocessor and a dedicated acceleration processor at the same time.

[0020] The host-side architecture information includes architecture name, system model and / or version information.

[0021] In host-side compilation, the process of compiling each host-side source code to generate the corresponding host-side bytecode file includes:

[0022] Preprocessing the host-side function in each host-side source code to generate a host-side preprocessing file corresponding to the host-side source code;

[0023] The host-side preprocessed file is compiled and processed by the front-end to generate a host-side bytecode file after the compilation and front-end processing.

[0024] In device-side compilation, the process of generating host-side bytecode files includes:

[0025] Preprocessing the device-side function in each device-side source code to generate a device-side preprocessing file corresponding to each device-side source code;

[0026] Compile front-end processing for each device-side processing file to generate a device-side bytecode file after the compilation front-end processing;

[0027] Link all device-side bytecode files into a device-side executable file;

[0028] Generate the host-side bytecode file based on the above device-side executable file.

[0029] When generating a host-side bytecode file based on a device-side executable file, the generating method includes:

[0030] Read the contents of the executable file on the device into an array variable and create an intermediate representation file in the format of the human-readable intermediate representation on the host.

[0031] Convert the human-readable intermediate representation file into a bytecode file to generate a host-side bytecode file.

[0032] Preprocessing of the device-side function includes expansion of the header file in the device-side function, replacement of macro definitions, deletion of unused codes and / or deletion of comments.

[0033] A highly compatible heterogeneous compiler, used for executing a heterogeneous compilation method, the heterogeneous compiler comprising:

[0034] The host-side architecture analysis module is used to analyze and obtain the host-side architecture information in heterogeneous processors;

[0035] A source code analysis and processing module, used for splitting the source code into a plurality of host-side source codes and a plurality of device-side source codes based on a keyword of each function in the source code;

[0036] A host-side compilation module compiles the host functions in the split host-side source code based on the host-side architecture information, so as to generate a corresponding host-side bytecode file based on each host-side source code after compilation;

[0037] The device-side compilation module compiles the device functions in the split device-side source code to generate a host-side bytecode file based on all the device-side source code after compilation.

[0038] The host-side compilation module links all host-side bytecode files into host-side executable files that are adapted to the host-side architecture information.

[0039] The advantages of the present invention are as follows: in the host-side compilation, the host-side functions are processed by preprocessing and compilation front-end processing in sequence to generate host bytecode files, and all host bytecode files are linked into host-side executable files, that is, heterogeneous compilation postpones code generation to linking, at which time, a wider range of codes can be seen, more inlining, constant propagation and useless code deletion can be done, and at the same time, the feature of supporting calling device-side functions in other source files can greatly improve the code reuse rate when applied to large engineering projects, which can provide great convenience for project development work and readability.

[0040] When performing host-side compilation, preprocessing and compilation front-end processing, the corresponding processing processes must be adapted to the host-side architecture information obtained above, that is, the selected compilation options must be adapted to the host-side architecture information, thereby achieving compatibility with general-purpose microprocessors of different architectures, which can greatly reduce the maintenance cost of the source code. BRIEF DESCRIPTION OF THE DRAWINGS

[0041] Figure 1 The present invention is a flow chart of an embodiment of existing heterogeneous compilation.

[0042] Figure 2 The present invention is a flowchart of an embodiment of a heterogeneous encoder performing heterogeneous compilation. DETAILED DESCRIPTION

[0043] The present invention will be further described below in conjunction with specific drawings and embodiments.

[0044] In order to achieve compatibility with general-purpose processors of different architectures and reduce the programming difficulty and maintenance cost of heterogeneous systems, a highly compatible heterogeneous compilation method is provided. In one embodiment of the present invention, the heterogeneous compilation method includes:

[0045] Providing source code to be compiled, wherein a keyword representing attribute information of the current function is configured for a function in the source code;

[0046] Based on the keywords of each function in the source code, split the source code into a plurality of host-side source codes and a plurality of device-side source codes;

[0047] Based on the architecture information of the host side, in the host side compilation, the host side functions in the split host side source code are compiled, so as to generate a corresponding host side bytecode file based on each host side source code after the compilation process;

[0048] In the device-side compilation, the device-side functions in the split device-side source code are compiled, so as to generate a host-side bytecode file based on all the device-side source codes after the compilation process;

[0049] Link the host-side bytecode files generated in the device-side compilation and all the host-side bytecode files generated in the host-side compilation into a host-side executable file that adapts to the host-side architecture information.

[0050] As can be seen from the above description, the heterogeneous compilation of the present invention is to compile the source code, and the source code is a heterogeneous programming code for a heterogeneous multi-core processor. The specific form of the source code can be encoded according to the needs of the actual application scenario. The source code generally includes several functions. In one embodiment of the present invention, keywords are configured for the functions in the source code. By configuring corresponding keywords for each function, the attribute information of each function can be characterized. The attribute information at least includes functions belonging to the host side, functions on the device side, or functions that can be called and executed on both the host side and the device side.

[0051] In one embodiment of the present invention, the keywords configured for the function in the source code include a __host__ keyword, a __global__ keyword, and a __device__ keyword, wherein the function modified by the __global__ keyword is a kernel function and the kernel function contains a call execution entry of a device-side function in the target source code, the function modified by the __host__ keyword is a function running on a general-purpose microprocessor; the function modified by the __device__ keyword is a function that can only be called and run on a dedicated processor;

[0052] When the __host__ keyword and the __device__ keyword are used to modify the same function, the modified function can be called and executed on a general-purpose microprocessor and a dedicated acceleration processor at the same time.

[0053] During specific implementation, the __host__ keyword, __global__ keyword, and __device__ keyword commonly used in this technical field are used to configure different functions in the source code. Among them, the method and process of using the __host__ keyword, __global__ keyword, and __device__ keyword to configure the properties of a function in the source code can be consistent with the existing ones, and the specific method shall be based on whether it can meet the compilation and execution of different functions.

[0054] Generally, the function modified with the __global__ keyword is a core function and the core function is the call execution entry of the device-side function in the target source code, wherein the source code to be compiled is the current source code, and the target source code is the current source code or other source code outside the current source code, that is, the function modified with the __global__ keyword can support the device-side function defined in a source code other than the current source code. At this time, the demand for calling the function in the device-side across source codes can be met, and the generation of duplicate code can be reduced. The call execution method and process of the device-side function modified with the __global__ keyword can be consistent with the existing one, and the specific one shall be based on the ability to implement the call execution of the required device-side function.

[0055] The keywords of each function are generally configured during source code programming. Based on the keywords of each function in the source code, the source code is split into host-side functions and device-side functions. For example, functions modified with the __host__ keyword are split into host-side functions, and functions modified with the __global__ keyword and the __device__ keyword are split into device-side functions. The same function can be split into host-side functions and device-side functions at the same time by modifying the __host__ keyword and the __device__ keyword.

[0056] In one embodiment of the present invention, the host-side architecture information includes an architecture name, a system model and / or version information. Specifically, the host-side architecture information can be obtained by using commonly used technical means in the technical field. The main purpose of obtaining the host-side architecture information is to achieve compatibility with general-purpose microprocessors of different architectures. Generally, according to the architecture name in the host-side architecture information, suitable architecture information and corresponding compilation options can be obtained in the compilation stage and subsequent compilation work can be performed.

[0057] In one embodiment of the present invention, in the host-side compilation, the process of compiling each host-side source code to generate a corresponding host-side bytecode file includes:

[0058] Preprocessing the host-side function in each host-side source code to generate a host-side preprocessing file corresponding to the host-side source code;

[0059] The host-side preprocessed file is compiled and processed by the front-end to generate a host-side bytecode file after the compilation and front-end processing.

[0060] The split host source code can be one or more. Figure 2 An embodiment of host-side compilation is shown in FIG. Figure 2 An embodiment of splitting to obtain n host-side source codes is shown in the figure. During preprocessing, the n host-side source codes are all preprocessed. At this time, n host-side preprocessing files are generated respectively. Specifically, the preprocessing function is to expand the header file, replace the macro definition, delete the useless code and / or delete the comment. The method and process of generating the host-side preprocessing file after preprocessing the host function in the host-side source code can be consistent with the existing one. In addition, the number n of split host-side source codes can be determined according to the specific situation of the source code.

[0061] The obtained n host-side preprocessing files are all compiled front-end processed. The function of the compilation front-end processing is to perform syntax analysis and lexical analysis, check whether there are syntax errors in the code in the current host-side preprocessing file, and generate a host-side bytecode file after analysis and inspection. At this time, n corresponding host-side bytecode files can be generated, that is, one host-side bytecode file corresponds to one host-side source code. The specific method and process of implementing the compilation front-end processing and generating the host-side bytecode file can be consistent with the existing ones.

[0062] From the above description, it can be seen that when performing host-side compilation, preprocessing and compilation front-end processing, the corresponding processing processes must be adapted to the host-side architecture information obtained above, that is, the selected compilation options must be adapted to the host-side architecture information, so as to achieve compatibility with general-purpose microprocessors of different architectures.

[0063] In one embodiment of the present invention, in device-side compilation, the process of generating a host-side bytecode file includes:

[0064] Preprocessing the device-side function in each device-side source code to generate a device-side preprocessing file corresponding to each device-side source code;

[0065] Compile front-end processing for each device-side processing file to generate a device-side bytecode file after the compilation front-end processing;

[0066] Link all device-side bytecode files into a device-side executable file;

[0067] Generate the host-side bytecode file based on the above device-side executable file.

[0068] Figure 2 An embodiment of device-side compilation is shown in the figure. In the figure, it is shown that the number of device-side source codes obtained by splitting is also n. Of course, in the specific implementation, the number of device-side source codes obtained by splitting can also be other numbers, which is specifically related to the situation of the source code.

[0069] Specifically, similar to the host-side compilation, each device-side source code is preprocessed, including the expansion of the header file in the device-side function, the replacement of macro definitions, the deletion of useless codes and / or the deletion of comments. Each device-side source code needs to be preprocessed to generate a corresponding device-side preprocessed file after preprocessing, that is, n device-side source codes will generate n corresponding device-side preprocessed files after preprocessing.

[0070] The main function of the compiler front end is to perform syntax analysis and lexical analysis to check whether the device-side code has syntax errors, so as to generate a device-side bytecode file after being processed by the compiler front end. When there are n device-side preprocessing files, n device-side bytecode files will be generated.

[0071] After generating n device-side bytecode files, the n device-side bytecode files are linked into a device-side executable file using a linking method commonly used in the technical field. The specific method and process of linking to form a device-side executable file can be consistent with the existing ones.

[0072] Figure 2 In one embodiment of the present invention, after generating a device-side executable file, the device-side executable file is processed by special processing to generate a host-side bytecode file. When generating a host-side bytecode file based on the device-side executable file, the generation method includes:

[0073] Read the contents of the executable file on the device into an array variable and create an intermediate representation file in the format of the human-readable intermediate representation on the host.

[0074] Convert the human-readable intermediate representation file into a bytecode file to generate a host-side bytecode file.

[0075] Specifically, the special processing is the process of the above-mentioned generation method. The format of the intermediate representation file can be a human-readable intermediate representation file, such as a human-readable format file. After reading the content of the device-side executable file, based on the rules of the human-readable format intermediate representation file, the read content is written into the human-readable format file, and the intermediate representation file is produced. After the intermediate representation file in the human-readable format is produced, a tool is used to convert and generate a machine-recognizable host-side bytecode file. Of course, other methods can also be used to generate a host-side bytecode file based on a device-side executable file.

[0076] Figure 2 The host-side bytecode file generated in the process and all the host-side bytecode files generated in the host-side compilation are linked to generate a host-side executable file adapted to the host-side architecture information through a linking method commonly used in the technical field, wherein the generated host-side executable file is an executable file that can run on the host architecture platform. Figure 2 The Ito mode in is the heterogeneous compilation method mentioned above.

[0077] For the above heterogeneous compilation method, Figure 2 A heterogeneous compiler for executing the heterogeneous compilation method is shown in FIG. 1 . For a highly compatible heterogeneous compiler, the heterogeneous compiler includes:

[0078] The host-side architecture analysis module is used to analyze and obtain the host-side architecture information in heterogeneous processors;

[0079] A source code analysis and processing module, used for splitting the source code into a plurality of host-side source codes and a plurality of device-side source codes based on a keyword of each function in the source code;

[0080] A host-side compilation module compiles the host functions in the split host-side source code based on the host-side architecture information, so as to generate a corresponding host-side bytecode file based on each host-side source code after compilation;

[0081] The device-side compilation module compiles the device functions in the split device-side source code to generate a host-side bytecode file based on all the device-side source code after compilation.

[0082] The host-side compilation module links all host-side bytecode files into host-side executable files that are adapted to the host-side architecture information.

[0083] Specifically, the host-side architecture analysis module and the source code analysis and processing module both implement preparations before compilation. Among them, the host-side architecture analysis module can analyze and obtain the host-side architecture information. The method of obtaining the host-side architecture information and the specific situation of the host-side architecture information can refer to the above description.

[0084] The device-side compilation module is mainly used to compile the device-side functions in the device-side source code to generate the device-side host-side bytecode file. The method and process of the device-side compilation module to generate the device-side host-side bytecode file can be referred to the above description.

[0085] The host-side compilation module compiles the host-side functions in the host-side source code and generates the final executable file that can run on the host architecture platform based on all the host-side bytecode files, that is, generates the host-side executable file. The process of the host-side compilation module compiling the host-side functions in the host-side source code and generating the host-side executable file can refer to the above description.

[0086] In summary, in host-side compilation, the host-side functions in the host-side source code are processed by preprocessing and compilation front-end processing in sequence to generate host bytecode files, and all host bytecode files are linked into host-side executable files. That is, heterogeneous compilation postpones code generation to link time. At this time, a wider range of code can be seen, and more inlining, constant propagation and useless code deletion can be done. At the same time, it supports the feature of calling device-side functions in other source files. When applied to large engineering projects, it can greatly improve the code reuse rate and provide great convenience for project development work and readability.

[0087] When performing host-side compilation, preprocessing and compilation front-end processing, the corresponding processing processes must be adapted to the host-side architecture information obtained above, that is, the selected compilation options must be adapted to the host-side architecture information, thereby achieving compatibility with general-purpose microprocessors of different architectures, which can greatly reduce the maintenance cost of the source code.

Claims

1. A highly compatible heterogeneous compilation method, characterized in that: The heterogeneous compilation method comprises: Providing source code to be compiled, wherein a keyword representing attribute information of the current function is configured for a function in the source code; Based on the keywords of each function in the source code, split the source code into a plurality of host-side source codes and a plurality of device-side source codes; Based on the architecture information of the host side, in the host side compilation, the host side functions in the split host side source code are compiled, so as to generate a corresponding host side bytecode file based on each host side source code after the compilation process; In the device-side compilation, the device-side functions in the split device-side source code are compiled, so as to generate a host-side bytecode file based on all the device-side source codes after the compilation process; Link the host-side bytecode files generated in the device-side compilation and all the host-side bytecode files generated in the host-side compilation into a host-side executable file that is adapted to the host-side architecture information; In host-side compilation, the process of compiling each host-side source code to generate the corresponding host-side bytecode file includes: Preprocessing the host-side function in each host-side source code to generate a host-side preprocessing file corresponding to the host-side source code; Performing compilation front-end processing on the host-side preprocessed file to generate a host-side bytecode file after the compilation front-end processing; In device-side compilation, the process of generating host-side bytecode files includes: Preprocessing the device-side function in each device-side source code to generate a device-side preprocessing file corresponding to each device-side source code; Compile front-end processing for each device-side processing file to generate a device-side bytecode file after the compilation front-end processing; Link all device-side bytecode files into a device-side executable file; Generate host-side bytecode files based on the above device-side executable files; When generating a host-side bytecode file based on a device-side executable file, the generation method includes: Read the contents of the executable file on the device into an array variable and create an intermediate representation file in the format of the human-readable intermediate representation on the host. Convert the human-readable intermediate representation file into a bytecode file to generate a host-side bytecode file.

2. The highly compatible heterogeneous compilation method according to claim 1, characterized in that: The keywords configured for the function in the source code include the __host__ keyword, the __global__ keyword, and the __device__ keyword, wherein the function modified by the __global__ keyword is a kernel function and the kernel function contains the call execution entry of the device-side function in the target source code, the function modified by the __host__ keyword is a function running on a general-purpose microprocessor; the function modified by the __device__ keyword is a function that can only be called and run on a dedicated processor; When the __host__ keyword and the __device__ keyword are used to modify the same function, the modified function can be called and executed on both a general-purpose microprocessor and a dedicated accelerated processor at the same time.

3. The highly compatible heterogeneous compilation method according to claim 1, characterized in that: The host-side architecture information includes architecture name, system model and / or version information.

4. The highly compatible heterogeneous compilation method according to claim 1, characterized in that: Preprocessing of the device-side function includes expansion of the header file in the device-side function, replacement of macro definitions, deletion of unused codes and / or deletion of comments.

5. A highly compatible heterogeneous compiler, characterized in that: For executing the heterogeneous compilation method in claim 1, the heterogeneous compiler comprises: The host-side architecture analysis module is used to analyze and obtain the host-side architecture information in heterogeneous processors; A source code analysis and processing module, used for splitting the source code into a plurality of host-side source codes and a plurality of device-side source codes based on a keyword of each function in the source code; A host-side compilation module compiles the host functions in the split host-side source code based on the host-side architecture information, so as to generate a corresponding host-side bytecode file based on each host-side source code after compilation; The device-side compilation module compiles the device functions in the split device-side source code to generate a host-side bytecode file based on all the device-side source code after compilation. The host-side compilation module links all host-side bytecode files into host-side executable files that are adapted to the host-side architecture information.

Images