Code conversion method, device, equipment and storage medium

Abstract

The present application discloses a code conversion method, device, equipment and storage medium, the method comprising: transcoding a kernel function of CUDA Fortran language in a code to be converted, obtaining a header file and a source file including a kernel function of CUDA C language, the main code in the code to be converted comprises a first code and a second code, the second code comprises a calling code, a wrapper function is constructed according to the source file, the header file and the kernel function of CUDA C language, the wrapper function is used to call the kernel function of CUDA C language, the second code is modified, a third code compatible with HIP is obtained, the third code comprises a code for calling the wrapper function, and finally, the converted code is obtained according to the modified main code, source file, header file and wrapper function. Thus, code conversion is realized and the efficiency of code conversion is improved.

Description

Technical Field

[0001] The present application relates to the field of computer technology, and in particular to a code conversion method, device, equipment and storage medium. Background Art

[0002] In modern high-performance computers, high-performance computer architectures with heterogeneous accelerators as the main computing units have gradually become the main structure. Heterogeneous accelerators can effectively provide high floating-point computing performance with low power consumption. Graphics Processing Unit (GPU) is a heterogeneous accelerator. In order to use heterogeneous accelerators with structural differences, the industry currently proposes a portable heterogeneous parallel programming model (Heterogeneous-Compute Interface for Portability, HIP) that can run on CUDA and other heterogeneous accelerators. The heterogeneous parallel programming model can explicitly or implicitly control heterogeneous accelerators, such as running computing tasks or copying data with the host. HIP only provides C / C++ syntax and interface support. Among them, CUDA is a heterogeneous parallel programming model that can use NVIDIA GPU for general computing.

[0003] Formula Translation (Fortran) is the main language used in the field of scientific computing, and scientific computing has always been the main field using high-performance computing technology. In some scientific computing fields, such as atmosphere and geology, some of the earliest scientific computing programs were usually developed using Fortran. In order to run these scientific computing programs on NVIDIA GPUs, it is necessary to rewrite the scientific computing programs in Fortran according to the writing requirements of CUDA to obtain heterogeneous parallel programs (referred to as CUDA Fortran codes). Since HIP only provides C / C++ syntax and interface support, in order to run CUDA Fortran codes on other heterogeneous accelerators, CUDA Fortran codes need to be ported to HIP, that is, CUDA Fortran codes are converted into heterogeneous parallel programs written in C according to the writing requirements of HIP (referred to as HIP C codes).

[0004] Currently, the conversion of CUDA Fortran code into HIP C code is done manually, which requires high personnel requirements, takes a long time, and has low code conversion efficiency. Summary of the invention

[0005] The present application provides a code conversion method, apparatus, device and storage medium to solve the problem of low code conversion efficiency when converting CUDA Fortran code to HIP C code.

[0006] In a first aspect, the present application provides a code conversion method, comprising:

[0007] Transcoding a kernel function in the CUDA Fortran language in the code to be converted to obtain a header file and a source file including the kernel function in the CUDA Fortran language, wherein the code to be converted includes a main code and the kernel function in the CUDA Fortran language, wherein the main code includes a first code and a second code, wherein the second code includes a calling code for calling the kernel function in the CUDA Fortran language;

[0008] Constructing a wrapper function according to the source file and the header file, wherein the wrapper function is used to call the kernel function of the CUDA C language;

[0009] Modify the second code to obtain a third code compatible with a heterogeneous parallel programming model (HIP), wherein the third code includes a code for calling the encapsulation function;

[0010] A converted code is obtained according to the modified main code, the source file, the header file and the encapsulation function, wherein the modified main code includes the first code and the third code.

[0011] Optionally, the second code further includes a call code and a reference module of a runtime system interface, and the modifying the second code to obtain a third code compatible with a portable heterogeneous parallel programming model HIP includes:

[0012] Modify the calling code of the runtime system interface to the calling code of the HIP-compatible runtime system interface;

[0013] Modify the reference module to the HIP reference module;

[0014] Modify the calling code into the code for calling the encapsulation function;

[0015] The third code is obtained according to the calling code of the HIP-compatible runtime system interface, the HIP reference module and the code for calling the encapsulation function.

[0016] Another embodiment of the above application has the following advantages or beneficial effects: by modifying the calling code, the reference module and the calling code of the runtime system interface respectively, a third code compatible with HIP can be obtained.

[0017] Optionally, the encapsulation function includes a kernel function encapsulation module in C language and a kernel function calling module in Fortran language, the kernel function encapsulation module in C language is used to perform a calling process of the CUDA C language kernel function and encapsulate the calling process, and the kernel function calling module in Fortran language is used to build a Fortran interface and call the function provided by the kernel function encapsulation module in C language.

[0018] Another embodiment of the above application has the following advantages or beneficial effects: by constructing a C language kernel function encapsulation module and a Fortran language kernel function calling module according to source files and header files, the Fortran language kernel function calling module constructs a Fortran interface and calls the function provided by the C language kernel function encapsulation module, and the C language kernel function encapsulation module calls the CUDA C language kernel function. When running the converted code, the code for calling the encapsulated function can indirectly call the CUDA C language kernel function through the Fortran language kernel function calling module and the C language kernel function encapsulation module, thereby realizing the call of the CUDA C language kernel function, thereby realizing the code conversion of the kernel function calling process.

[0019] Optionally, constructing a package function according to the source file and the header file includes:

[0020] Construct the C language kernel function encapsulation module according to the source file and the header file;

[0021] The Fortran language kernel function calling module is constructed according to the C language kernel function encapsulation module.

[0022] Another embodiment of the above application has the following advantages or beneficial effects: since the kernel function calling module of the Fortran language is used to construct the Fortran interface and call the function provided by the kernel function encapsulation module of the C language, the kernel function encapsulation module of the C language is first constructed according to the source file and the header file, and then the kernel function calling module of the Fortran language is constructed according to the kernel function encapsulation module of the C language, which facilitates the kernel function calling module of the Fortran language to call the function provided by the kernel function encapsulation module of the C language.

[0023] Optionally, constructing the C language kernel function encapsulation module according to the source file and the header file includes:

[0024] Extracting kernel function parameters from the kernel function in the CUDA C language according to the kernel function declaration in the header file;

[0025] According to the structure of the kernel function parameters, construct a parameter declaration of the kernel function encapsulation module in the C language, wherein the parameter declaration of the kernel function encapsulation module in the C language includes the kernel function running parameters and the kernel function parameters;

[0026] According to the parameter declaration of the kernel function encapsulation module in the C language, construct the second calling code, wherein the parameters of the second calling code are the same as the parameters included in the parameter declaration of the kernel function encapsulation module in the C language;

[0027] Copy the code of the source file to the header file to obtain the target header file;

[0028] The C language kernel function encapsulation module is obtained according to the preset header file corresponding to the HIP, the target header file, the parameter declaration of the C language kernel function encapsulation module and the second calling code.

[0029] Another embodiment of the above application has the following advantages or beneficial effects: by first constructing the parameter declaration and the second calling code of the C language kernel function encapsulation module in sequence according to the structure of the kernel function parameters, and then copying the code of the source file into the header file to obtain the target header file, and then obtaining the C language kernel function encapsulation module according to the header file corresponding to the preset HIP, the target header file, the parameter declaration and the second calling code of the C language kernel function encapsulation module, the construction of the C language kernel function encapsulation module is realized, which facilitates the C language kernel function encapsulation module to call the CUDA C language kernel function.

[0030] Optionally, constructing the Fortran language kernel function calling module according to the C language kernel function encapsulation module includes:

[0031] According to the parameter declaration of the kernel function encapsulation module in the C language, construct the first calling code, wherein the parameters of the first calling code are the same as the parameters after removing the matrix dimension description information parameters in the parameter declaration of the kernel function encapsulation module in the C language;

[0032] Constructing a calling function declaration according to the parameters of the first calling code;

[0033] The kernel function calling module of the Fortran language is obtained according to the calling function declaration, the preset reference module and the first calling code.

[0034] Another embodiment of the above application has the following advantages or beneficial effects: by sequentially constructing a first calling code and a calling function declaration according to the parameter declaration of the kernel function encapsulation module of the C language, a kernel function calling module of the Fortran language is obtained according to the calling function declaration, a preset reference module and the first calling code, thereby realizing the construction of the kernel function calling module of the Fortran language, and facilitating the kernel function calling module of the Fortran language to call the function provided by the kernel function encapsulation module of the C language.

[0035] Optionally, the Fortran language kernel function calling module further includes a matrix dimension information description array, and the method further includes:

[0036] When it is determined that the kernel function parameters include matrix dimension description information parameters and the format of the matrix dimension description information parameters is a preset format, constructing a matrix dimension information description array according to the matrix dimension description information parameters;

[0037] The matrix dimension information description array is added to the kernel function calling module of the Fortran language.

[0038] Another embodiment of the above application has the following advantages or beneficial effects: by determining that the kernel function parameters include matrix dimension description information parameters, and the format of the matrix dimension description information parameters is a preset format, a matrix dimension information description array is constructed according to the matrix dimension description information parameters, and the matrix dimension information description array is added to the kernel function calling module of the Fortran language, so that the kernel function calling module of the Fortran language can use the matrix dimension description information array to perform correct address offset calculations and complete value acquisition operations.

[0039] Optionally, constructing a matrix dimension information description array according to the matrix dimension description information parameters includes:

[0040] According to the matrix dimension description information parameter, the matrix dimension description information when the kernel function of the CUDA C language is passed is obtained, wherein the matrix dimension description information includes the matrix dimension and each dimension interval of the matrix;

[0041] A matrix dimension information description array is constructed according to the matrix dimension and each dimension interval of the matrix, wherein the dimensions of the matrix dimension information description array include a first dimension and a second dimension, wherein the first dimension is the same as the matrix dimension, and the second dimension is used to characterize the information of each dimension of the matrix, wherein the information of one dimension includes: a dimension interval starting value, a dimension interval size, a dimension interval ending value, and a total offset value of the dimension relative to the matrix starting address.

[0042] Another embodiment of the above application has the following advantages or beneficial effects: by first obtaining the matrix dimension description information when the kernel function of the CUDA C language passes parameters, a matrix dimension information description array is constructed according to the matrix dimension included in the matrix dimension description information and each dimension interval of the matrix, thereby realizing the construction of the matrix dimension information description array, so that the kernel function calling module of the Fortran language can use the matrix dimension description information array to perform correct address offset calculations and complete value acquisition operations.

[0043] Optionally, the converted code is compiled as follows:

[0044] Use hipcc compiler to compile the encapsulation function, source file and header file to generate a dynamic link library file;

[0045] The Fortran compiler is called to compile the modified main code to generate a compiled file, and the compiled file is linked with the dynamic link library file.

[0046] Another embodiment of the above application has the following advantages or beneficial effects: compiling in this manner facilitates smooth running of the converted code when running on a heterogeneous accelerator.

[0047] Optionally, before modifying the calling code of the runtime system interface to the calling code of the HIP-compatible runtime system interface, the method further includes:

[0048] It is determined that the HIP supports the calling code of the runtime system interface and the function of the kernel function of the CUDA Fortran language.

[0049] Another embodiment of the above application has the following advantages or beneficial effects: by modifying the calling code of the runtime system interface to the calling code of the runtime system interface compatible with the HIP only when it is determined that the HIP supports the calling code of the runtime system interface and the function of the kernel function of the CUDA Fortran language, the efficiency of code conversion can be further improved.

[0050] In a second aspect, the present application provides a code conversion device, comprising:

[0051] a transcoding module, configured to transcode a kernel function in a CUDA Fortran language in a code to be converted, and obtain a header file and a source file including a kernel function in a CUDA C language, wherein the code to be converted includes a main code and the kernel function in the CUDA Fortran language, wherein the main code includes a first code and a second code, and wherein the second code includes a calling code for calling the kernel function in the CUDA Fortran language;

[0052] Processing modules for:

[0053] Constructing a wrapper function according to the source file and the header file, wherein the wrapper function is used to call the kernel function of the CUDA C language;

[0054] Modify the second code to obtain a third code compatible with a heterogeneous parallel programming model (HIP), wherein the third code includes a code for calling the encapsulation function;

[0055] A converted code is obtained according to the modified main code, the source file, the header file and the encapsulation function, wherein the modified main code includes the first code and the third code.

[0056] In a third aspect, the present application provides a code conversion device, including:

[0057] Processor; and

[0058] A memory, configured to store executable instructions of the processor;

[0059] The processor is configured to execute the code conversion method described in any one of the first aspect or possible implementations of the first aspect by executing the executable instructions.

[0060] In a fourth aspect, an embodiment of the present application provides a computer-readable storage medium having a computer program stored thereon, and when the computer program is executed by a processor, the code conversion method described in the first aspect or any possible implementation manner of the first aspect is implemented.

[0061] In a fifth aspect, an embodiment of the present application provides a computer program product, including a computer program, which, when executed by a processor, implements the code conversion method described in the first aspect or any possible implementation manner of the first aspect.

[0062] The code conversion method, device, equipment and storage medium provided by the present application, by transcoding the kernel function of the CUDA Fortran language in the code to be converted, obtain the source file and header file including the kernel function of the CUDA C language, construct the encapsulation function for calling the kernel function of the CUDA C language according to the source file and the header file, modify the second code to the third code compatible with HIP, the third code includes the code for calling the encapsulation function, and obtain the converted code according to the modified main code, source file, header file and encapsulation function. Thus, the conversion of the CUDA Fortran code to the HIP C code is automatically completed, and the efficiency of the code conversion is improved. In the process of converting the programming language of the code to be converted, only the code conversion of the kernel function and the calling process of the kernel function in the code to be converted is performed, and the first code in the code to be converted does not need to be converted in the programming language, which can avoid the overall conversion of the code to be converted, the efficiency of the code conversion is higher, and the dependence on the C language can be reduced. BRIEF DESCRIPTION OF THE DRAWINGS

[0063] Figure 1 A schematic diagram of an application scenario of the code conversion method provided in an embodiment of the present application;

[0064] Figure 2 A flowchart of a code conversion method provided in an embodiment of the present application;

[0065] Figure 3 A flowchart of a code conversion method provided in an embodiment of the present application;

[0066] Figure 4 A flowchart of a code conversion method provided in an embodiment of the present application;

[0067] Figure 5 A flowchart of a code conversion method provided in an embodiment of the present application;

[0068] Figure 6 A flowchart of a code conversion method provided in an embodiment of the present application;

[0069] Figure 7 A flowchart of a code conversion method provided in an embodiment of the present application;

[0070] Figure 8 A schematic diagram of the structure of a code conversion device provided in an embodiment of the present application;

[0071] Fig. 9 A schematic diagram of the structure of a code conversion device provided in an embodiment of the present application. DETAILED DESCRIPTION

[0072] The embodiments of the present application are described in detail below, and examples of the embodiments are shown in the accompanying drawings. The embodiments described below with reference to the accompanying drawings are exemplary and are intended to be used to explain the present application, but cannot be understood as limiting the present application.

[0073] The terms "first" and "second" etc. in the specification, claims and drawings of the embodiments of the present application are used to distinguish similar objects, and are not necessarily used to describe a specific order or sequence. It should be understood that the data used in this way can be interchangeable where appropriate, so that the embodiments of the embodiments of the present application described herein can be implemented in an order other than those illustrated or described herein, for example. In addition, the terms "including" and "having" and any of their variations are intended to cover non-exclusive inclusions, for example, the process, method, system, product or equipment comprising a series of steps or units is not necessarily limited to those steps or units clearly listed, but may include other steps or units that are not clearly listed or inherent to these processes, methods, products or equipment.

[0074] First, some terms in the embodiments of the present application are explained below to facilitate understanding by those skilled in the art.

[0075] 1. Heterogeneous accelerators are GPU-like heterogeneous acceleration hardware, such as graphics processing units (GPUs), integrated many cores (MICs), digital signal processing (DSPs), and field programmable gate arrays (FPGAs).

[0076] 2. HIP is a heterogeneous programming model that can run on CUDA and other heterogeneous accelerators.

[0077] 3. CUDA is a heterogeneous parallel programming model that can use NVIDIA GPUs for general computing.

[0078] 4. CUDA Fortran is a Fortran extension programming model that can run on NVIDIA GPUs.

[0079] In order to solve the problem that when converting CUDA Fortran code to HIP C code, it is done manually in the prior art, which has high requirements for personnel, takes a long time, and has low code conversion efficiency, the embodiment of the present application provides a code conversion method, device, equipment and storage medium. After obtaining the code to be converted (CUDA Fortran code), the code to be converted includes a main code and a kernel function of the CUDA Fortran language. The main code includes a first code and a second code. The second code includes a calling code for calling the kernel function of the CUDA Fortran language. The kernel function of the CUDA Fortran language is first transcoded to obtain a header file and a source file including the kernel function of the CUDA C language, and then the second code is modified to obtain a HIP-compatible third code. The third code includes the modified calling code, and then a wrapper function is constructed according to the source file and the header file. The wrapper function is an intermediate calling layer. Since the calling code is a Fortran language code, the calling code cannot directly call the kernel function of the CUDA C language. Therefore, the calling code is modified to a code that can call the wrapper function. The modified calling code can indirectly call CUDA through the wrapper function. C language kernel function, and finally, according to the modified main code, source file, header file and encapsulation function, the converted code can be obtained. Thus, the conversion of CUDA Fortran code to HIP C code is automatically completed, and the efficiency of code conversion is improved. In the process of converting CUDA Fortran code to HIP C code, only the code conversion of the kernel function and the calling process of the kernel function in the CUDA Fortran code is performed. The first code in the CUDA Fortran code does not need to be converted into a programming language, which can avoid the overall conversion of the CUDA Fortran code. The efficiency of code conversion is higher and the dependence on the C language can be reduced. The technical solution provided by the present application is described in detail below in conjunction with the accompanying drawings.

[0080] Next, the application scenarios involved in the embodiments of the present application are illustrated.

[0081] The code conversion method provided in the embodiment of the present application can be applied at least to the scenario where a scientific computing program developed in the Fortran language is run on a heterogeneous accelerator other than an NVIDIA GPU. Specifically, it can be applied to porting a scientific computing program developed in the Fortran language to a HIP in an efficient manner.

[0082] For example, Figure 1 A schematic diagram of an application scenario of the code conversion method provided in the embodiment of the present application, such as Figure 1As shown, this embodiment involves a CUDA Fortran compiler 10 and a code conversion device 20. During the transplantation process, it is necessary to first rewrite the scientific computing program developed in the Fortran language using the Fortran language according to the writing requirements of CUDA to obtain a heterogeneous parallel program (referred to as CUDA Fortran code). For example, Figure 1 The CUDA Fortran compiler 10 shown in the figure compiles the scientific computing program into CUDA Fortran code, and then inputs the CUDA Fortran code into the code conversion device 20. The code conversion device 20 uses the code conversion method provided in the embodiment of the present application to convert the CUDA Fortran code to obtain the converted code, that is, the heterogeneous parallel program (HIP C code for short) written in C language according to the writing requirements of HIP. Optionally, the code conversion device 20 can output the converted code. Thus, the scientific computing program developed in Fortran language is efficiently transplanted to HIP.

[0083] The technical solution of the present application and how the technical solution of the present application solves the above-mentioned technical problems are described in detail below with specific embodiments. The following specific embodiments can be combined with each other, and the same or similar concepts or processes may not be repeated in some embodiments. The embodiments of the present application will be described below in conjunction with the accompanying drawings.

[0084] Figure 2 The flowchart of a code conversion method provided in an embodiment of the present application, the code conversion method can be performed by a code conversion device, and the code conversion device can be implemented in software and / or hardware. Figure 2 As shown, the method of this embodiment may include:

[0085] S101, transcoding a kernel function in the CUDA Fortran language in the code to be converted to obtain a header file and a source file including the kernel function in the CUDA C language, wherein the code to be converted includes a main code and a kernel function in the CUDA Fortran language, the main code includes a first code and a second code, and the second code includes a calling code for calling the kernel function in the CUDA Fortran language.

[0086] Specifically, the code to be converted is the code input to the code conversion device, and the code to be converted can be the CUDA Fortran code shown above. The code to be converted includes the main code and the kernel function of the CUDA Fortran language. The main code includes a first code and a second code. The second code includes a calling code for calling the kernel function of the CUDA Fortran language, wherein the first code is, for example, a code for managing heterogeneous programs and operating heterogeneous accelerators.

[0087] In an practicable manner, the CUDA Fortran language kernel function in the code to be converted is transcoded by calling a CUDA Fortran compiler (such as pgfortran) to compile the CUDA Fortran language kernel function to generate a source file and a header file, and the file names of the source file and the header file end with ".gpu" and ".h" respectively. Optionally, the CUDA Fortran compiler and the code conversion device of this embodiment can be located on the same device, such as on the same heterogeneous accelerator, so that the code conversion device can call the CUDA Fortran compiler.

[0088] S102: Construct a wrapper function according to the source file and the header file, where the wrapper function is used to call a kernel function of the CUDA C language.

[0089] S103: Modify the second code to obtain a third code compatible with HIP, where the third code includes the modified calling code.

[0090] Optionally, the second code further includes a call code and a reference module of a runtime system interface, and modifying the second code may include:

[0091] S1031. Modify the calling code of the runtime system interface to the calling code of the runtime system interface that is compatible with the HIP.

[0092] Optionally, before modifying the calling code of the runtime system interface to the calling code of the HIP-compatible runtime system interface, the following may also be included:

[0093] Determine whether the HIP supports the calling code of the runtime system interface and the kernel function of the CUDA Fortran language. If it is determined that the HIP does not support the calling code of the runtime system interface and the kernel function of the CUDA Fortran language, it can send an indication of non-support, or it can also send an indication for instructing the user to modify the code to be converted. The efficiency of code conversion can be further improved.

[0094] S1032. Modify the reference module to a HIP reference module.

[0095] S1033. Modify the calling code into a code for calling the encapsulation function.

[0096] S1034: Obtain a third code according to the calling code of the HIP-compatible runtime system interface, the HIP reference module, and the code for calling the encapsulated function.

[0097] S104. Obtain converted code according to the modified main code, source file, header file and encapsulation function.

[0098] Specifically, the modified main code is obtained according to the first code and the third code, and the converted code includes the modified main code, source files, header files and encapsulation functions. Optionally, the modified main code, source files, header files and encapsulation functions can be packaged to obtain the converted code.

[0099] Specifically, since the code to be converted includes the main code and the kernel function of the CUDA Fortran language, the main code (including the first code and the second code, the second code including the calling code of the runtime system interface, the reference module and the calling code for calling the kernel function of the CUDA Fortran language) is a code written in the Fortran language, the kernel function of the CUDA Fortran language in the code to be converted is converted into the kernel function of the CUDA C language after transcoding by S101, and the calling code written in the Fortran language cannot directly call the kernel function of the CUDA C language transcoded by S101. In this embodiment, a wrapper function is first constructed according to the source file and the header file, and the wrapper function is used to call the kernel function of the CUDA C language. The second code is modified to obtain a HIP-compatible third code, and the third code includes a code for calling the wrapper function, wherein the wrapper function (which is an intermediate layer calling process) is constructed according to the kernel function of the CUDAC language, and the calling code is modified to a code for calling the wrapper function. In this way, when the converted code is run, the code for calling the wrapper function can indirectly call the CUDA C language through the wrapper function. The kernel function of the C language, i.e., the code for calling the encapsulation function, calls the encapsulation function, and the encapsulation function calls the kernel function of the CUDAC language, thereby realizing the calling of the kernel function of the CUDA C language.

[0100] The code conversion method provided in this embodiment transcodes the kernel function of the CUDA Fortran language in the code to be converted, obtains the source file and header file including the kernel function of the CUDA C language, constructs the encapsulation function for calling the kernel function of the CUDA C language according to the source file and the header file, modifies the second code into a third code compatible with HIP, the third code includes the code for calling the encapsulation function, and obtains the converted code according to the modified main code, source file, header file and encapsulation function. Thus, the conversion of the CUDA Fortran code to the HIP C code is automatically completed, and the efficiency of the code conversion is improved. In the process of converting the programming language of the code to be converted, only the code conversion of the kernel function and the calling process of the kernel function in the code to be converted is performed, and the first code in the code to be converted does not need to be converted in the programming language, which can avoid the overall conversion of the code to be converted, the efficiency of the code conversion is higher, and the dependence on the C language can be reduced.

[0101] On the basis of the above embodiments, as an implementable method, the encapsulation function includes a kernel function encapsulation module in C language and a kernel function calling module in Fortran language. The kernel function encapsulation module in C language is used to perform the calling process of the kernel function in CUDA C language and encapsulate the calling process. The kernel function calling module in Fortran language is used to build a Fortran interface and call the function provided by the kernel function encapsulation module in C language.

[0102] In this embodiment, a C language kernel function encapsulation module and a Fortran language kernel function calling module are constructed according to source files and header files, the Fortran language kernel function calling module constructs a Fortran interface and calls the function provided by the C language kernel function encapsulation module, and the C language kernel function encapsulation module calls the CUDA C language kernel function. When the converted code is run, the code for calling the encapsulated function can indirectly call the CUDA C language kernel function through the Fortran language kernel function calling module and the C language kernel function encapsulation module, thereby realizing the calling of the CUDA C language kernel function, thereby realizing the code conversion of the kernel function calling process.

[0103] Figure 3 A flowchart of a code conversion method provided in an embodiment of the present application is shown in FIG. Figure 3 As shown, the method of this embodiment is Figure 2 Based on the illustrated embodiment, optionally, S102 may specifically include:

[0104] S1021. Construct a C language kernel function encapsulation module according to the source file and the header file.

[0105] S1022. Construct a Fortran language kernel function calling module according to the C language kernel function encapsulation module.

[0106] In this embodiment, since the Fortran language kernel function calling module is used to construct the Fortran interface and call the function provided by the C language kernel function encapsulation module, the C language kernel function encapsulation module is first constructed according to the source file and the header file, and then the Fortran language kernel function calling module is constructed according to the C language kernel function encapsulation module, so that the Fortran language kernel function calling module can call the function provided by the C language kernel function encapsulation module.

[0107] Figure 4 A flowchart of a code conversion method provided in an embodiment of the present application is shown in FIG. Figure 4 As shown, the method of this embodiment is Figure 3 On the basis of the embodiment shown, further, in S1021, a kernel function encapsulation module in C language is constructed according to the source file and the header file, which can be specifically implemented by the following steps:

[0108] S201. Extract kernel function parameters from a kernel function in CUDA C language according to the kernel function declaration in the header file.

[0109] Specifically, the header file (.h file) has a kernel function declaration (i.e., a C language declaration) of the kernel function. The CUDA C language kernel function includes a header file reference, kernel function parameters, and a kernel function body. Optionally, it may also include a kernel function usage data structure definition. The kernel function parameters can be extracted from the CUDA C language kernel function according to the kernel function declaration in the header file.

[0110] S202 . Construct a parameter declaration of a kernel function encapsulation module in C language according to the structure of the kernel function parameters. The parameter declaration of the kernel function encapsulation module in C language includes kernel function running parameters and kernel function parameters.

[0111] For example, a kernel function parameter is as follows:

[0112] extern"C"__global__void vectoradd_(

[0113] signed char*__restrict___pin,

[0114] signed char*_pout,int n__V_value,

[0115] int n__V_width,

[0116] signed char*__restrict___pin_sd,

[0117] signed char*__restrict___pout_sd)

[0118] The structure of the kernel function parameters includes signed char*__restrict___pin,

[0119] signed char*_pout,int n__V_value,

[0120] int n__V_width,

[0121] signed char*__restrict___pin_sd,

[0122] signed char*__restrict___pout_sd

[0123] According to the structure of kernel function parameters, add kernel function running parameters, which include (four parameters, namely grid dimension (gridSize), block dimension (blockSize), dynamic Shared Memroy size (sharedMemSize) and stream object (stream), and you can get the parameter declaration of the kernel function encapsulation module in C language, for example, as shown below:

[0124] void vectoradd_(

[0125] dim3 gridSize,

[0126] dim3 blockSize,

[0127] size_t sharedMemSize,

[0128] hipStream_t stream,

[0129] signed char*__restrict___pin,

[0130] signed char*_pout,

[0131] int n__V_value,

[0132] int n__V_width,

[0133] signed char*__restrict___pin_sd,

[0134] signed char*__restrict___pout_sd)

[0135] S203. Construct a second calling code according to the parameter declaration of the kernel function encapsulation module in the C language, wherein the parameters of the second calling code are the same as the parameters included in the parameter declaration of the kernel function encapsulation module in the C language.

[0136] Specifically, the parameters of the second calling code are the same as the parameters included in the parameter declaration of the kernel function encapsulation module of the C language. According to the parameter declaration of the kernel function encapsulation module of the C language, the second calling code is constructed according to a preset format. For example, the second calling code constructed according to the parameter declaration of the kernel function encapsulation module of the C language is as follows:

[0137] vectoradd_<<<gridSize,blockSize,sharedMemSize,stream> >>

[0138] (_pin,_pout,n__V_value,n__V_width,_pin_sd,_pout_sd)

[0139] S204, copy the code of the source file into the header file to obtain the target header file.

[0140] For example, copying the source file code into the header file will result in the target header file being "vectorAdd.n001.h".

[0141] S205 , obtaining a C language kernel function encapsulation module according to a header file corresponding to a preset HIP, a target header file, a parameter declaration of the C language kernel function encapsulation module, and a second calling code.

[0142] Specifically, the C language kernel function encapsulation module includes an encapsulation function source code file, which includes a header file corresponding to the HIP, a target header file, a parameter declaration of the C language kernel function encapsulation module, and a second calling code corresponding to the calling process of the C language kernel function encapsulation module calling the CUDA C language kernel function, wherein the target header file is a file obtained by copying the code of the source file into the header file.

[0143] Specifically, according to the preset header file corresponding to the HIP, the target header file, the parameter declaration of the C language kernel function encapsulation module and the second calling code, the codes are combined in a preset order to obtain the C language kernel function encapsulation module. For example, according to the preset header file corresponding to the HIP, the target header file obtained in S202-S204, the parameter declaration of the C language kernel function encapsulation module and the second calling code, the C language kernel function encapsulation module is obtained as follows:

[0144]

[0145] Figure 5 A flowchart of a code conversion method provided in an embodiment of the present application is shown in FIG. Figure 5 As shown, the method of this embodiment is Figure 3 On the basis of the embodiment shown, further, in S1022, a kernel function calling module in Fortran language is constructed according to the kernel function encapsulation module in C language, which can be specifically implemented by the following steps:

[0146] S301. Construct a first calling code according to a parameter declaration of a kernel function encapsulation module in C language, wherein parameters of the first calling code are the same as parameters after removing matrix dimension description information parameters from the parameter declaration of the kernel function encapsulation module in C language.

[0147] Specifically, the parameters of the first calling code are the same as the parameters after removing the matrix dimension description information parameters in the parameter declaration of the kernel function encapsulation module of the C language, Figure 4 Taking the parameter declaration of the kernel function encapsulation module of the C language in the illustrated embodiment as an example, the parameter declaration of the kernel function encapsulation module of the C language includes the parameters gridSize, blockSize, sharedMemSize, stream, _pin, _pout, n__V_value, n__V_width, _pin_sd, _pout_sd, wherein _pin_sd, _pout_sd are matrix dimension description information parameters, _pin_sd, _pout_sd are removed to obtain gridSize, blockSize, sharedMemSize, stream, _pin, _pout, n__V_value, n__V_width, which are the parameters of the first calling code. When constructing the first calling code, it is constructed according to the preset format of the first calling code. For example, the first calling code constructed according to the parameter declaration of the kernel function encapsulation module of the C language is as follows:

[0148] call vectorAdd(gridSize,blockSize,sharedMemSize,streamId,

[0149] in,out,value,width,in_sd,out_sd)

[0150] S302: Construct a calling function declaration according to the parameters of the first calling code.

[0151] Specifically, the parameters of the first calling code are the same as the parameters included in the calling function declaration. For example, the calling function declaration constructed according to the parameter declaration of the first calling code is as follows:

[0152] subroutine vectorAdd(gridSize,blockSize,sharedMemSize,streamId,

[0153] in,out,value,width)

[0154] S303: Obtain a Fortran language kernel function calling module according to the calling function declaration, the preset reference module and the first calling code.

[0155] Specifically, the Fortran language kernel function calling module includes a calling function declaration, a reference module, and a first calling code corresponding to the calling process of the Fortran language kernel function calling module calling the function provided by the C language kernel function encapsulation module. Optionally, it can also include removing the Fortran language definition of the function type corresponding to the function name.

[0156] The preset reference module is, for example, "iso_c_binding". For example, according to the above-mentioned calling function declaration, the preset reference module and the first calling code, they are combined in a preset order to obtain the Fortran language kernel function calling module as follows:

[0157]

[0158] In an embodiment of the present application, when the kernel function parameters include matrix parameters of unknown dimensions, the kernel function calling module of the Fortran language will require the caller to pass in relevant information. An unknown dimension matrix refers to a matrix variable that does not display the declared dimensions in the function area, and displaying the declared dimensions refers to indicating the intervals of each dimension of the matrix using dimension syntax. The kernel function calling module of the Fortran language can only perform correct address offset calculations and complete value-taking operations with the help of a matrix dimension description information array. Therefore, when the kernel function parameters include matrix dimension description information parameters, it is necessary to construct a matrix dimension information description array based on the matrix dimension description information parameters and add it to the kernel function calling module of the Fortran language, so that the kernel function calling module of the Fortran language can perform correct address offset calculations and complete value-taking operations with the help of the matrix dimension description information array.

[0159] Furthermore, in an implementable manner, the kernel function calling module of the Fortran language may also describe the array using matrix dimension information. The method of this embodiment may further include, before S303:

[0160] S304: When it is determined that the kernel function parameters include matrix dimension description information parameters and the format of the matrix dimension description information parameters is a preset format, construct a matrix dimension information description array according to the matrix dimension description information parameters.

[0161] Specifically, the kernel function parameters include array parameters, value parameters and matrix dimension description information parameters. As in the kernel function parameter example above, the matrix dimension description information parameters are located after the array parameters and the value parameters. The matrix dimension description information parameters can be found according to the position settings of the parameters in the kernel function parameters.

[0162] For example, the matrix dimension description information parameter is signed char*__restrict___pin_sd, signedchar*__restrict___pout_sd), and its format is prefixed with signed char*. Then, the matrix dimension information description array is constructed according to the matrix dimension description information parameter, for example, the constructed matrix dimension information description array is generate matrixdimension information in_sd and out_sd.

[0163] S305 , adding the matrix dimension information description array to the kernel function calling module of the Fortran language.

[0164] Specifically, the matrix dimension information description array is added to the kernel function call module of the tran language mentioned above, and the kernel function call module of the Fortran language obtained is, for example:

[0165]

[0166] In this embodiment, by determining that the kernel function parameters include matrix dimension description information parameters and the format of the matrix dimension description information parameters is a preset format, a matrix dimension information description array is constructed according to the matrix dimension description information parameters, and the matrix dimension information description array is added to the kernel function calling module of the Fortran language, so that the kernel function calling module of the Fortran language can use the matrix dimension description information array to perform correct address offset calculation and complete the value acquisition operation.

[0167] Figure 6 A flowchart of a code conversion method provided in an embodiment of the present application is as follows: Figure 6 The method of this embodiment is Figure 5 Based on the exemplary embodiment, further, constructing a matrix dimension information description array according to the matrix dimension description information parameters in S304 may specifically include:

[0168] S3041. Obtain matrix dimension description information when a kernel function of the CUDA C language is passed according to the matrix dimension description information parameter. The matrix dimension description information includes the matrix dimension and each dimension interval of the matrix.

[0169] S3042. Construct a matrix dimension information description array according to the matrix dimension and each dimension interval of the matrix. The dimensions of the matrix dimension information description array include a first dimension and a second dimension. The first dimension is the same as the matrix dimension, and the second dimension is used to characterize the information of each dimension of the matrix. The information of one dimension includes: the starting value of the dimension interval, the size of the dimension interval, the ending value of the dimension interval, and the total offset value of the dimension relative to the starting address of the matrix.

[0170] Specifically, for example, if the matrix is ​​passed in the form of ":", the dimension interval of the matrix is ​​defined as 0. The dimension interval size is calculated according to each dimension interval of the matrix, and a dimension interval size value is defined as the dimension interval end value minus the dimension interval start value plus 1.

[0171] In the above embodiment, the converted code needs to be compiled before running when it is running. As an implementable method, the converted code can be compiled in the following way:

[0172] Use the hipcc compiler to compile the encapsulated functions, source files, and header files to generate dynamic link library files.

[0173] The Fortran compiler is called to compile the modified main code, generate a compiled file, and link the compiled file with the dynamic link library file.

[0174] Compile in the above way to facilitate the smooth running of the converted code.

[0175] The detailed process of the code conversion method provided by the present application is described below in conjunction with a specific embodiment.

[0176] Figure 7 The flowchart of a code conversion method provided in an embodiment of the present application, the code conversion method can be performed by a code conversion device, and the code conversion device can be implemented in software and / or hardware. Figure 7 As shown, the method of this embodiment may include:

[0177] S401: Receive a code to be converted.

[0178] Specifically, the code to be converted is the code input into the code conversion device, and the code to be converted includes the main code and the kernel function of the CUDA Fortran language. The main code includes the first code and the second code. The second code includes the calling code of the runtime system interface, the reference module and the calling code for calling the kernel function of the CUDA Fortran language. The first code is, for example, the code for managing heterogeneous programs and operating heterogeneous accelerators.

[0179] Take the following main code as an example:

[0180] Main code:

[0181]

[0182]

[0183] In the above main code, the calling code of the runtime system interface is not shown. What is shown is the code other than the calling code of the runtime system interface. In this part of the code, the bold black parts "use cudafor" and "callvectorAdd<<<1,n>>>(a_d,b_d,3,n)" are the calling codes for referencing the module and the kernel function for calling the CUDA Fortran language, respectively, and the rest is the first code.

[0184] S402: Determine whether the code to be converted is a code designed and implemented using CUDA Fortran or a code that has been implemented using CUDA Fortran.

[0185] S403 , executing corresponding operations according to whether the HIP supports the calling code of the runtime system interface and the kernel function of the CUDA Fortran language.

[0186] Specifically, the hipifiy-perl tool can be called to scan the code to be converted. The scanning output results include information such as the calling code of the unsupported runtime system interface, the unsupported kernel function function and the kernel function that cannot be identified. Based on the output results, it can be determined whether HIP supports the calling code of the runtime system interface and the kernel function function of the CUDA Fortran language.

[0187] Specifically, if the HIP does not support the call code of the runtime system interface and the kernel function of the CUDA Fortran language, it sends an indication message indicating that it does not support the code, or it can also send an indication message for instructing the user to modify the code to be converted, which can further improve the efficiency of code conversion.

[0188] If the HIP supports the calling code of the runtime system interface and the kernel function of the CUDA Fortran language, the calling code of the runtime system interface is modified to the calling code of the runtime system interface compatible with the HIP, which can be specifically modified by modifying the function name of the calling code of the runtime system interface, or performing format conversion of the data structure of the calling code of the runtime system interface. For example, the calling code of the runtime system interface is the runtime (Application Programming Interface, API) provided by CUDA, and when modifying, the function name of the Runtime API provided by CUDA can be modified, or the format conversion of the data structure of the Runtime API provided by CUDA can be performed.

[0189] S404 , transcode the kernel function of the CUDA Fortran language in the code to be converted to obtain a header file and a source file including the kernel function of the CUDA C language.

[0190] The kernel function of the CUDA C language includes a header file reference, kernel function parameters and a kernel function body, and optionally, may also include a kernel function usage data structure definition.

[0191] Specifically, a CUDA Fortran compiler (such as pgfortran) may be called to compile a kernel function of the CUDA Fortran language to generate a source file and a header file, wherein the file names of the source file and the header file end with ".gpu" and ".h", respectively.

[0192] For example, the following is a transcoding process of a kernel function in the CUDA Fortran language. The following is an example of a code to be converted (named: vectorAdd.n001.cuf):

[0193]

[0194] In the above code to be converted, the CUDA Fortran language kernel function included in the code to be converted is "vectorAdd". When transcoding "vectorAdd", the following instructions are used to generate the transcoded program files vectorAdd.n001.gpu and vectorAdd.n001.h:

[0195] pgfortran-c-Mcuda=keepgpu vectorAdd.cuf

[0196] Among them, vectorAdd.n001.gpu is the source file, vectorAdd.n001.h is the header file, and the source file includes the following contents:

[0197]

[0198]

[0199] Among them, #include "cuda_runtime.h"

[0200] #include "vectorAdd.n001.h" is the header file reference; extern "C" __global__voidvectoradd_(

[0201] signed char*__restrict___pin,

[0202] signed char*_pout,int n__V_value,

[0203] int n__V_width,

[0204] signed char*__restrict___pin_sd,

[0205] signed char*__restrict___pout_sd) is the kernel function parameter;

[0206] The rest is the kernel function body.

[0207] Specifically, the source file vectorAdd.n001.gpu has the C language implementation of the CUDA C language kernel function (vectorAdd C). During the transcoding process, the corresponding C language type identifier will be automatically generated according to the type definition in the CUDA Fortran language kernel function, where the array parameter will be defined in the form of a signed character pointer declaration signed char*, such as the _pin and _pout parameters shown in the above kernel function parameter section; the value parameter will be defined in the form of an immediate value, such as the _value and _width parameters shown in the above kernel function parameter section; the CUDA C language kernel function obtained after transcoding will add a matrix dimension description information parameter to the array parameter, which is the matrix dimension description information when the calling code calls the CUDA Fortran language kernel function, such as the _pin_sd and _pout_sd shown in the above kernel function parameter section, and the format starts with signed char*.

[0208] The header file vectorAdd.n001.h includes the following:

[0209] extern"C"__global__void vectoradd_(

[0210] signed char*__restrict___pin,

[0211] signed char*_pout,

[0212] int n__V_value,

[0213] int n__V_width,

[0214] signed char*__restrict___pin_sd, corresponding to in matrix and out matrix respectively

[0215] signed char*__restrict___pout_sd);

[0216] Among them, vectorAdd.n001.h contains the C language declaration of the kernel function vectorAdd C.

[0217] S405 . Construct a packaging function according to the source file and the header file. The packaging function includes a kernel function packaging module in C language and a kernel function calling module in Fortran language.

[0218] The specific construction process can be found above. Figure 3-Figure 5 As shown, no further description is given here.

[0219] When constructing a kernel function call module in the Fortran language, it is determined that the kernel function parameters include a matrix dimension description information parameter, and when the format of the matrix dimension description information parameter is a preset format, a matrix dimension information description array is constructed according to the matrix dimension description information parameter. For details on how to construct a matrix dimension information description array, see Figure 6 The description of the illustrated embodiment will not be repeated here.

[0220] S406: Modify the reference module to a HIP reference module, and modify the calling code to a code for calling the encapsulated function.

[0221] For example, taking the main code shown above as an example, change the reference module "use cudafor" to "usehipfort", and change the calling code "call vectorAdd<<<1,n>>>(a_d,b_d,3,n)" to "callvectorAdd(1,n,0,0,a_d,b_d,3,n)". Among them, the four parameters "1,n,0,0" are the kernel function running parameters, and the last four parameters "a_d,b_d,3,n" are the kernel function parameters.

[0222] S407: Obtain a third code according to the calling code of the HIP-compatible runtime system interface, the HIP reference module and the code for calling the encapsulated function, and obtain a modified main code according to the first code and the third code.

[0223] For example, taking the main code shown above as an example, except that the calling code of the runtime system interface is not shown, the modified main code is as follows:

[0224]

[0225] S408. Obtain converted code according to the modified main code, source file, header file and encapsulation function.

[0226] In this embodiment, the converted code needs to be compiled as follows:

[0227] First, use the hipcc compiler to compile the encapsulation function, source file, and header file to generate a dynamic link library file. For example, the specific instructions for compiling are as follows, libkernel.so is the dynamic link library file name, and wrapper is the encapsulation function.

[0228] hipcc-fPIC--shared-o libkernel.so wrapper.cpp

[0229] Then, the Fortran compiler is called to compile the first code, generate a compiled file, and link the compiled file with the dynamic link library file.

[0230] In this embodiment, the type of Fortran compiler is not limited. For example, taking the gfortran compiler as an example, the specific instructions used during compilation are as follows:

[0231] gfortran-o test test.F90-lkernel

[0232] The following is an embodiment of the device of the present application, which can be used to execute the above method embodiment of the present application. For details not disclosed in the embodiment of the device of the present application, reference can be made to the above method embodiment of the present application.

[0233] Figure 8 A schematic diagram of the structure of a code conversion device provided in an embodiment of the present application is shown in FIG. Figure 8 As shown, the device of this embodiment may include: a transcoding module 11 and a processing module 12, wherein:

[0234] The transcoding module 11 is used to transcode the kernel function of the CUDA Fortran language in the code to be converted, and obtain a header file and a source file including the kernel function of the CUDA C language, the code to be converted includes a main code and the kernel function of the CUDA Fortran language, the main code includes a first code and a second code, and the second code includes a calling code for calling the kernel function of the CUDA Fortran language;

[0235] The processing module 12 is used for:

[0236] Build a wrapper function based on the source file and header file. The wrapper function is used to call the kernel function of CUDA C language;

[0237] Modify the second code to obtain a third code compatible with a portable heterogeneous parallel programming model HIP, wherein the third code includes a code for calling a package function;

[0238] According to the modified main code, source file, header file and encapsulation function, a converted code is obtained, and the modified main code includes a first code and a third code.

[0239] Optionally, the second code further includes a call code of a runtime system interface and a reference module, and the processing module 12 is used to modify the call code of the runtime system interface to a call code of a HIP-compatible runtime system interface;

[0240] Change the reference module to HIP reference module;

[0241] Modify the calling code to the code used to call the encapsulated function;

[0242] The third code is obtained according to the calling code of the HIP-compatible runtime system interface, the HIP reference module and the code for calling the encapsulated function.

[0243] Optionally, the encapsulation function includes a C language kernel function encapsulation module and a Fortran language kernel function calling module, the C language kernel function encapsulation module is used to perform the calling process of the CUDA C language kernel function and encapsulate the calling process, and the Fortran language kernel function calling module is used to build a Fortran interface and call the function provided by the C language kernel function encapsulation module.

[0244] Optionally, the Fortran language kernel function calling module includes a calling function declaration, a reference module, and a first calling code corresponding to a calling process of the Fortran language kernel function calling module calling a function provided by the C language kernel function encapsulation module;

[0245] The C language kernel function encapsulation module includes an encapsulation function source code file, which includes a header file corresponding to the HIP, a target header file, a parameter declaration of the C language kernel function encapsulation module, and a second calling code corresponding to the calling process of the C language kernel function encapsulation module calling the CUDA C language kernel function. The target header file is a file obtained by copying the code of the source file to the header file.

[0246] Optionally, the processing module 12 is used for:

[0247] Build the C language kernel function encapsulation module according to the source file and header file;

[0248] The kernel function calling module of Fortran language is constructed according to the kernel function encapsulation module of C language.

[0249] Further, the processing module 12 is used to: extract kernel function parameters from the kernel function of CUDA C language according to the kernel function declaration in the header file;

[0250] According to the structure of the kernel function parameters, a parameter declaration of the kernel function encapsulation module in C language is constructed, and the parameter declaration of the kernel function encapsulation module in C language includes the kernel function running parameters and the kernel function parameters;

[0251] According to the parameter declaration of the kernel function encapsulation module in the C language, a second calling code is constructed, wherein the parameters of the second calling code are the same as the parameters included in the parameter declaration of the kernel function encapsulation module in the C language;

[0252] Copy the code of the source file to the header file to get the target header file;

[0253] According to the header file corresponding to the preset HIP, the target header file, the parameter declaration of the C language kernel function encapsulation module and the second calling code, the C language kernel function encapsulation module is obtained.

[0254] Further, the processing module 12 is used to: construct a first calling code according to the parameter declaration of the kernel function encapsulation module of the C language, and the parameters of the first calling code are the same as the parameters after removing the matrix dimension description information parameters in the parameter declaration of the kernel function encapsulation module of the C language;

[0255] Constructing a calling function declaration according to the parameters of the first calling code;

[0256] A kernel function calling module of the Fortran language is obtained according to the calling function declaration, the preset reference module and the first calling code.

[0257] Optionally, the kernel function calling module of the Fortran language further includes a matrix dimension information description array, and the processing module 12 is further used to: when it is determined that the kernel function parameters include a matrix dimension description information parameter and the format of the matrix dimension description information parameter is a preset format, construct a matrix dimension information description array according to the matrix dimension description information parameter;

[0258] Add the matrix dimension information description array to the kernel function call module of the Fortran language.

[0259] Optionally, the processing module 12 is further used to: obtain matrix dimension description information when the kernel function of the CUDA C language passes parameters according to the matrix dimension description information parameter, and the matrix dimension description information includes the matrix dimension and each dimension interval of the matrix;

[0260] A matrix dimension information description array is constructed according to the matrix dimension and each dimension interval of the matrix. The dimensions of the matrix dimension information description array include a first dimension and a second dimension. The first dimension is the same as the matrix dimension, and the second dimension is used to characterize the information of each dimension of the matrix, wherein the information of one dimension includes: a starting value of the dimension interval, a size of the dimension interval, an ending value of the dimension interval, and a total offset value of the dimension relative to the starting address of the matrix.

[0261] Optionally, the converted code is compiled as follows:

[0262] Use hipcc compiler to compile the package functions, source files and header files to generate dynamic link library files;

[0263] The Fortran compiler is called to compile the modified main code, generate a compiled file, and link the compiled file with the dynamic link library file.

[0264] Optionally, the processing module 12 is further configured to: before modifying the runtime system interface calling code to a HIP-compatible runtime system interface calling code, determine whether the HIP supports the runtime system interface calling code and the kernel function of the CUDA Fortran language.

[0265] The device provided in the embodiment of the present application can execute the above method embodiment. Its specific implementation principle and technical effects can be found in the above method embodiment, and this embodiment will not be repeated here.

[0266] It should be noted that it should be understood that the division of the various modules of the above device is only a division of logical functions. In actual implementation, they can be fully or partially integrated into one physical entity, or they can be physically separated. And these modules can all be implemented in the form of software called by processing elements; they can also be all implemented in the form of hardware; some modules can also be implemented in the form of software called by processing elements, and some modules can be implemented in the form of hardware. For example, the processing module can be a separately established processing element, or it can be integrated in a chip of the above device. In addition, it can also be stored in the memory of the above device in the form of program code, and called and executed by a processing element of the above device. The implementation of other modules is similar. In addition, these modules can be fully or partially integrated together, or they can be implemented independently. The processing element here can be an integrated circuit with signal processing capabilities. In the implementation process, each step of the above method or each module above can be completed by an integrated logic circuit of hardware in the processor element or instructions in the form of software.

[0267] For example, the above modules may be one or more integrated circuits configured to implement the above methods, such as one or more application specific integrated circuits (ASIC), or one or more microprocessors (digital signal processors, DSP), or one or more field programmable gate arrays (FPGA), etc. For another example, when a module above is implemented in the form of a processing element scheduling program code, the processing element may be a general-purpose processor, such as a central processing unit (CPU) or other processor that can call program code. For another example, these modules may be integrated together and implemented in the form of a system-on-a-chip (SOC).

[0268] In the above embodiments, it can be implemented in whole or in part by software, hardware, firmware or any combination thereof. When implemented using software, it can be implemented in whole or in part in the form of a computer program product. The computer program product includes one or more computer instructions. When loading and executing computer program instructions on a computer, the process or function according to the embodiment of the present application is generated in whole or in part. The computer can be a general-purpose computer, a special-purpose computer, a computer network, or other programmable devices. The computer instructions can be stored in a computer-readable storage medium, or transmitted from one computer-readable storage medium to another computer-readable storage medium. For example, the computer instructions can be transmitted from a website site, a computer, a server or a data center by wired (e.g., coaxial cable, optical fiber, digital subscriber line (DSL)) or wireless (e.g., infrared, wireless, microwave, etc.) mode to another website site, computer, server or data center. The computer-readable storage medium can be any available medium that a computer can access or a data storage device such as a server or a data center that includes one or more available media integration. The available medium can be a magnetic medium, (e.g., a floppy disk, a hard disk, a tape), an optical medium (e.g., a DVD), or a semiconductor medium (e.g., a solid state disk (SSD)), etc.

[0269] Fig. 9 A schematic diagram of a code conversion device provided in an embodiment of the present application is shown in FIG. Fig. 9 As shown, the code conversion device of this embodiment may include a processor 21 and a memory 22.

[0270] The memory 22 is used to store executable instructions of the processor 21 .

[0271] The processor 21 is configured to perform the code conversion method in the above method embodiment by executing executable instructions.

[0272] Optionally, the memory 22 may be independent or integrated with the processor 21 .

[0273] When the memory 22 is a device independent of the processor 21, the code conversion device of this embodiment may further include:

[0274] The bus 23 is used to connect the memory 22 and the processor 21 .

[0275] Optionally, the code conversion device of this embodiment may further include: a communication interface 24 , and the communication interface 24 may be connected to the processor 21 through a bus 23 .

[0276] The present application also provides a computer-readable storage medium, in which computer-executable instructions are stored. When the computer-readable storage medium is run on a computer, the computer executes the code conversion method as described in the above embodiment.

[0277] An embodiment of the present application further provides a computer program product, including a computer program, which implements the code conversion method in the above embodiment when executed by a processor.

[0278] In the description of this specification, the description with reference to the terms "one embodiment", "some embodiments", "example", "specific example", or "some examples" etc. means that the specific features, structures, materials or characteristics described in conjunction with the embodiment or example are included in at least one embodiment or example of the present application. In this specification, the schematic representations of the above terms do not necessarily refer to the same embodiment or example. Moreover, the specific features, structures, materials or characteristics described may be combined in any one or more embodiments or examples in a suitable manner. In addition, those skilled in the art may combine and combine the different embodiments or examples described in this specification and the features of the different embodiments or examples, without contradiction.

[0279] Although the embodiments of the present application have been shown and described above, it can be understood that the above embodiments are exemplary and cannot be understood as limitations on the present application. Ordinary technicians in this field can change, modify, replace and modify the above embodiments within the scope of the present application.

Claims

1. A code conversion method, characterized in that: include: Transcoding a kernel function in the CUDA Fortran language in the code to be converted to obtain a header file and a source file including the kernel function in the CUDA C language, wherein the code to be converted includes a main code and the kernel function in the CUDA Fortran language, the main code includes a first code and a second code, the second code includes a calling code for calling the kernel function in the CUDA Fortran language, and the first code is all the main codes except the second code; Constructing a wrapper function according to the source file and the header file, wherein the wrapper function is used to call the kernel function of the CUDA C language; Modify the second code to obtain a third code compatible with a heterogeneous parallel programming model (HIP), wherein the third code includes a code for calling the encapsulation function; A converted code is obtained according to the modified main code, the source file, the header file and the encapsulation function, wherein the modified main code includes the first code and the third code.

2. The method according to claim 1, characterized in that The second code also includes a call code and a reference module of a runtime system interface, and the modification of the second code to obtain a third code compatible with a portable heterogeneous parallel programming model HIP includes: Modify the calling code of the runtime system interface to the calling code of the HIP-compatible runtime system interface; Modify the reference module to a HIP reference module; Modify the calling code into the code for calling the encapsulation function; The third code is obtained according to the calling code of the HIP-compatible runtime system interface, the HIP reference module and the code for calling the encapsulation function.

3. The method according to claim 1, characterized in that The encapsulation function includes a kernel function encapsulation module in C language and a kernel function calling module in Fortran language. The kernel function encapsulation module in C language is used to perform a calling process of the kernel function in CUDA C language and encapsulate the calling process, and the kernel function calling module in Fortran language is used to build a Fortran interface and call the function provided by the kernel function encapsulation module in C language.

4. The method according to claim 3, characterized in that The step of constructing a package function according to the source file and the header file includes: Construct the C language kernel function encapsulation module according to the source file and the header file; The Fortran language kernel function calling module is constructed according to the C language kernel function encapsulation module.

5. The method according to claim 4, characterized in that The step of constructing the C language kernel function encapsulation module according to the source file and the header file includes: Extracting kernel function parameters from the kernel function in the CUDA C language according to the kernel function declaration in the header file; According to the structure of the kernel function parameters, construct a parameter declaration of the kernel function encapsulation module in the C language, wherein the parameter declaration of the kernel function encapsulation module in the C language includes the kernel function running parameters and the kernel function parameters; Constructing a second calling code according to the parameter declaration of the kernel function encapsulation module in the C language, wherein the parameters of the second calling code are the same as the parameters included in the parameter declaration of the kernel function encapsulation module in the C language; Copy the code of the source file into the header file to obtain a target header file; The C language kernel function encapsulation module is obtained according to the preset header file corresponding to the HIP, the target header file, the parameter declaration of the C language kernel function encapsulation module and the second calling code.

6. The method according to claim 5, characterized in that The step of constructing the Fortran language kernel function calling module according to the C language kernel function encapsulation module includes: According to the parameter declaration of the kernel function encapsulation module in the C language, a first calling code is constructed, wherein the parameters of the first calling code are the same as the parameters after removing the matrix dimension description information parameters from the parameter declaration of the kernel function encapsulation module in the C language; Constructing a calling function declaration according to the parameters of the first calling code; The kernel function calling module of the Fortran language is obtained according to the calling function declaration, the preset reference module and the first calling code.

7. The method according to claim 6, characterized in that The Fortran language kernel function calling module also includes a matrix dimension information description array, and the method further includes: When it is determined that the kernel function parameters include matrix dimension description information parameters and the format of the matrix dimension description information parameters is a preset format, constructing a matrix dimension information description array according to the matrix dimension description information parameters; The matrix dimension information description array is added to the kernel function calling module of the Fortran language.

8. The method according to claim 7, characterized in that The step of constructing a matrix dimension information description array according to the matrix dimension description information parameters includes: According to the matrix dimension description information parameter, the matrix dimension description information when the kernel function of the CUDA C language is passed is obtained, wherein the matrix dimension description information includes the matrix dimension and each dimension interval of the matrix; A matrix dimension information description array is constructed according to the matrix dimension and each dimension interval of the matrix, wherein the dimensions of the matrix dimension information description array include a first dimension and a second dimension, wherein the first dimension is the same as the matrix dimension, and the second dimension is used to characterize the information of each dimension of the matrix, wherein the information of one dimension includes: a dimension interval starting value, a dimension interval size, a dimension interval ending value, and a total offset value of the dimension relative to the matrix starting address.

9. The method according to any one of claims 1 to 8, characterized in that: The converted code is compiled as follows: Use hipcc compiler to compile the encapsulation function, source file and header file to generate a dynamic link library file; The Fortran compiler is called to compile the modified main code to generate a compiled file, and the compiled file is linked with the dynamic link library file.

10. The method according to claim 2, characterized in that Before modifying the calling code of the runtime system interface to the calling code of the HIP-compatible runtime system interface, the method further includes: It is determined that the HIP supports the calling code of the runtime system interface and the function of the kernel function of the CUDA Fortran language.

11. A code conversion device, characterized in that: include: a transcoding module, configured to transcode a kernel function in a CUDA Fortran language in a code to be converted, and obtain a header file and a source file including a kernel function in a CUDA C language, wherein the code to be converted includes a main code and the kernel function in the CUDA Fortran language, wherein the main code includes a first code and a second code, wherein the second code includes a calling code for calling the kernel function in the CUDA Fortran language, and wherein the first code is all the main codes except the second code; Processing modules for: Constructing a wrapper function according to the source file and the header file, wherein the wrapper function is used to call the kernel function of the CUDA C language; Modify the second code to obtain a third code compatible with a heterogeneous parallel programming model (HIP), wherein the third code includes a code for calling the encapsulation function; A converted code is obtained according to the modified main code, the source file, the header file and the encapsulation function, wherein the modified main code includes the first code and the third code.

12. A code conversion device, characterized in that: include: processor; as well as A memory, configured to store executable instructions of the processor; The processor is configured to perform the code conversion method according to any one of claims 1 to 10 by executing the executable instructions.

13. A computer-readable storage medium having a computer program stored thereon, characterized in that: When the computer program is executed by a processor, the code conversion method according to any one of claims 1 to 10 is implemented.

Images