CUDA (Compute Unified Device Architecture) source program compiling method based on general compiler and related equipment

Abstract

The invention provides a CUDA (Compute Unified Device Architecture) source program compiling method and related equipment based on a universal compiler, and the method comprises the following steps: executing preprocessing operation on source codes of a CUDA source program, and respectively generating a first intermediate file and a second intermediate file; scanning a host side code in the first intermediate file to identify a target function identifier representing equipment side execution semantics, and executing a reconstruction operation on a function containing the target function identifier to generate a proxy function used for triggering equipment side execution; scanning a device side code in the second intermediate file to identify a target function identifier, and performing semantic removal processing on the target function identifier to generate a device side executable function code; generating a host side target file, generating a device side target file, and further generating a device side dynamic link file; and forming a target file. By the adoption of the scheme, the CUDA program can be converted into the target program capable of running on the CPU, and the running environment construction cost is reduced.

Description

Technical Field

[0001] This application relates to the field of computer technology, and more specifically, to a CUDA source program compilation method and related equipment based on a general-purpose compiler. Background Technology

[0002] During algorithm development, it is necessary to build a runtime environment according to development requirements. Specifically, the graphics processing unit (GPU), dedicated high-speed interconnect technology (NVSwitch), and remote direct memory access (RDMA) need to be connected according to different requirements.

[0003] Building an environment using real equipment is time-consuming and costly, while using a central processing unit (CPU) to simulate GPU operation can quickly build an experimental environment.

[0004] However, when using a CPU to emulate a GPU, the CUDA source code needs to be compiled into an object file, but the NVCC compiler (NVIDIA CUDA Compiler) can only compile the CUDA source code into an object file on the GPU side, and cannot execute it on the CPU.

[0005] The information disclosed in the background section is only intended to enhance the understanding of the background of this application, and therefore may contain information that is not part of the prior art known to those skilled in the art. Summary of the Invention

[0006] This application provides a CUDA source program compilation method and related equipment based on a general-purpose compiler, which can convert CUDA programs into target programs that can run on a CPU, thereby reducing the cost of building the runtime environment and shortening the construction cycle.

[0007] The first aspect of this application provides a method for compiling CUDA source programs based on a general-purpose compiler, including: Preprocessing operations are performed on the source code of the CUDA source program to generate a first intermediate file and a second intermediate file, respectively; the preprocessing operations include at least the expansion of macro definition-related instructions; Scan the host-side code in the first intermediate file in the host-side code processing path to identify the target function identifier that represents the device-side execution semantics, and perform a refactoring operation on the function containing the target function identifier to generate a proxy function for triggering the device-side execution; Scan the device-side code in the second intermediate file in the device-side code processing path to identify the target function identifier, and perform semantic removal processing on the target function identifier to generate device-side executable function code; A host-side target file is generated based on the processed first intermediate file, a device-side target file is generated based on the processed second intermediate file, and a device-side dynamic link file is further generated. The device-side dynamic link file and the host-side target file are merged to form a target file.

[0008] In an optional embodiment of this application, the preprocessing operation on the CUDA source code to generate a first intermediate file and a second intermediate file includes: The g++ compiler is used to perform preprocessing operations on source code files with the .cu extension. By calling preprocessing parameters, macro definitions, conditional compilation directives, and header file inclusion directives in the source code file are expanded, and the macro-expanded source code is output to the specified target file as an intermediate file. Specifically, the first intermediate file is generated when the command -E a.cu -o a.cu.E is used, and the second intermediate file is generated when the command -E a.cu -o a.cu.E.dev is used.

[0009] In an optional embodiment of this application, the objective function identifier includes the __global__ keyword; The refactoring of the function containing the target function identifier includes: determining the code between the __global__ keyword and the first left curly brace { as the function header; determining the code between the right curly brace} matching the left curly brace as the function body; and determining the content between the parentheses ( ) in the function header as the parameter list, thereby extracting the function name, parameter list, and function body.

[0010] In an optional embodiment of this application, when generating the proxy function for triggering execution on the device side, a first function, a second function, and a third function are generated at the original location of the function containing the __global__ keyword; The first function is fn_get_mangled, which sends parameters and execution requests to the device side; the second function is fn, which is used as a parameter for calling cudaLaunchKernel; the third function is fn__stub_kernelaunch__, which replaces the kernel boot syntax <<<>>>; and the second function can call the first function, and the third function can call the second function.

[0011] In an optional embodiment of this application, the generation process of the first function includes: using fn_get_mangled as the function name and the parameter list as the formal parameter; using the thread-local variable funcBack to store the file name of the .so file to be executed on the device side, the name of the function to be executed, and the number of parameters; wherein, the thread-local variable funcBack is used when calling cudaLaunchKernel; The second function contains a void Type parameter, and the void The type parameter is assigned to the thread-local variable funcBack; and the parameter list is used as a local variable, and fn_get_mangled is called; The parameter list of the third function includes the parameter list as well as dim3 gridDim, dim3 blockDim, size_t sm_size, and cudaStream_t stream. Then, it constructs an array of parameter pointers and calls cudaLaunchKernel((void )fn, gridDim, blockDim, args, m_size, stream).

[0012] In an optional embodiment of this application, after refactoring the call to the function containing the __global__ keyword, the following is also performed: Replace the __global__ keyword with an empty string; For functions that are only modified by __host__, replace the __host__ modifier with an empty string; For functions that are only modified by __device__, replace the entire function with an empty string; For functions modified by both __host__ and __device__, replace the __host__ and __device__ modifiers with empty strings while preserving the function implementation.

[0013] In an optional embodiment of this application, when generating device-side executable function code, the global object is processed as follows: For global variables modified by __device__, replace __device__ with __attribute__ ((section(.DEVICE))); and for global variables modified by __constant__, replace __constant__ with __attribute__ ((section(.CONST))). Specifically, when determining the type of the object modified by __device__, the presence of a parameter list between __device__ and the first left parenthesis or semicolon is checked to distinguish whether __device__ modifies a function or a variable.

[0014] In an optional embodiment of this application, merging the device-side dynamic link file and the host-side target file to form a target file includes: The binary content of the device-side dynamic link file is sequentially appended to the end of the host-side target file; A fixed-length identifier string is appended to the content of the device-side dynamic link file to identify the existence status of the device-side code; A fixed-length description field describing the byte length of the device-side dynamic link file is appended after the fixed-length identifier string, forming a target file structure that can be parsed and split during runtime.

[0015] A second aspect of this application provides a CUDA source program compilation apparatus based on a general-purpose compiler, comprising: The preprocessing module is configured to perform preprocessing operations on the source code of the CUDA source program to generate a first intermediate file and a second intermediate file, respectively; the preprocessing operations include at least the expansion of macro definition-related instructions; The processing module is configured to scan the host-side code in the first intermediate file in the host-side code processing path to identify the target function identifier representing the device-side execution semantics, and perform a refactoring operation on the function containing the target function identifier to generate a proxy function for triggering the device-side execution; and to scan the device-side code in the second intermediate file in the device-side code processing path to identify the target function identifier, and perform semantic removal processing on the target function identifier to generate device-side executable function code. The merging module is configured to generate a host-side target file based on the processed first intermediate file, generate a device-side target file based on the processed second intermediate file, further generate a device-side dynamic link file, and merge the device-side dynamic link file and the host-side target file to form a target file.

[0016] A third aspect of this application provides a computer device, including: a memory and a processor, wherein the memory stores a computer program, and the processor executes the computer program to implement the steps of any of the methods described in the foregoing embodiments.

[0017] A fourth aspect of this application provides a computer-readable storage medium having a computer program stored thereon, which, when executed by a processor, implements the steps of the method described in any of the foregoing embodiments.

[0018] A fifth aspect of this application provides a computer program product, including a computer program / instructions that, when executed by a processor, implement the steps of the method described in any of the foregoing embodiments.

[0019] In summary, by preprocessing the CUDA source code and expanding macro definitions, clearly structured intermediate files can be generated, making subsequent CUDA semantic recognition more accurate. This allows for targeted refactoring of GPU-related logic. In the host-side processing path, by identifying the target function identifiers of device-side execution semantics, related functions are refactored and proxy functions are generated, allowing the original GPU call flow to be replaced with a CPU-executable call method. In the device-side processing path, by identifying and removing target function identifiers, CPU-executable function code can be generated, enabling computational logic that originally relied on GPU execution to run on the CPU. This allows CUDA programs to be automatically converted into CPU-running target programs without requiring GPU hardware, reducing dependence on specific hardware, improving the program's deployment and running capabilities in more scenarios, reducing runtime environment construction costs, and shortening the development cycle. Attached Figure Description

[0020] The accompanying drawings, which are included to provide a further understanding of this application and form part of this application, illustrate exemplary embodiments and are used to explain this application, but do not constitute an undue limitation of this application. In the drawings: Figure 1 A flowchart illustrating a CUDA source program compilation method based on a general-purpose compiler, provided in this application embodiment; Figure 2 A flowchart illustrating a CUDA source code compilation method in a specific application scenario provided in this application embodiment; Figure 3 A schematic diagram of the structure of a CUDA source program compilation device based on a general-purpose compiler provided in this application embodiment; Figure 4 This is a schematic diagram of the hardware structure of a computer device provided in an embodiment of this application. Detailed Implementation

[0021] In the process of developing this application, the inventors discovered that when using a CPU to simulate a GPU, the CUDA source code needs to be compiled into an object file that can run on the CPU. However, the NVCC compiler can only compile the CUDA source code into an object file for the GPU side and cannot execute it on the CPU. This is because the CUDA source code contains at least two parts: The host-side code, i.e., ordinary C / C++, can be compiled into x86 / ARM and executed on the CPU.

[0022] Device-side code (e.g., __global__ / __device__) is compiled by the NVCC compiler into a GPU instruction set, relying on the GPU's execution model. Therefore, device-side code cannot run on the CPU.

[0023] To address the aforementioned issues, this application provides a CUDA source code compilation method based on a general-purpose compiler. By preprocessing the CUDA source code and expanding macro definitions, clearly structured intermediate files can be generated, making subsequent CUDA semantic recognition more accurate and enabling targeted refactoring of GPU-related logic. In the host-side processing path, by identifying the target function identifiers of device-side execution semantics, related functions are refactored and proxy functions are generated, allowing the original GPU call flow to be replaced with a CPU-executable call method. In the device-side processing path, by identifying and removing the target function identifiers, CPU-executable function code can be generated, enabling computational logic that originally relied on GPU execution to run on the CPU.

[0024] In this way, CUDA programs can be automatically converted into target programs that can run on the CPU without the need for GPU hardware and a CUDA runtime environment. This reduces dependence on specific hardware, improves the ability of programs to be deployed and run in more scenarios, and thus reduces the cost of building the runtime environment and shortens the construction cycle.

[0025] The solutions in this application embodiment can be implemented using various computer languages, such as the object-oriented programming language Java and the interpreted scripting language JavaScript.

[0026] To make the technical solutions and advantages of the embodiments of this application clearer, the exemplary embodiments of this application will be described in further detail below with reference to the accompanying drawings. Obviously, the described embodiments are only a part of the embodiments of this application, and not an exhaustive list of all embodiments. It should be noted that, unless otherwise specified, the embodiments and features in the embodiments of this application can be combined with each other.

[0027] Please see Figure 1The present application provides a flowchart of a CUDA source program compilation method based on a general-purpose compiler, as shown in the embodiments. Figure 1 As shown, steps S101 to S105 can be performed: S101, perform preprocessing operations on the source code of the CUDA source program to generate a first intermediate file and a second intermediate file respectively; the preprocessing operations include at least the expansion of macro definition-related instructions.

[0028] In some embodiments, preprocessing operations may be text-level processing flows performed on CUDA source programs before compilation or analysis, used to unify the syntax elements in the source code that affect subsequent compilation, static analysis, or code transformation, thereby generating intermediate files that can be used in subsequent steps.

[0029] The macro expansion is implemented using the compiler's built-in command parameter -E. By executing macro expansion, only the source code remains after expansion, which is equivalent to a first-level simplification process, thereby generating intermediate files with a more standardized structure and stronger parsability.

[0030] In one embodiment, step S101 may include: The g++ compiler is used to perform preprocessing operations on source code files with the .cu extension. By calling preprocessing parameters, macro definitions, conditional compilation directives, and header file inclusion directives in the source code file are expanded, and the macro-expanded source code is output to the specified object file as an intermediate file.

[0031] As an example, step S101 can be simplified to: macro expansion g++ -E a.cu -o a.cu.E, generating the first intermediate file, and macro expansion g++ -E a.cu -o a.cu.E.dev, generating the second intermediate file.

[0032] Here, g++ is the C++ compiler; -E means to expand macros only, without compiling; a.cu: source code file; oa.cu.E: save the results to the file a.cu.E; o a.cu.E.dev: save the results to the file a.cu.E.dev.

[0033] That is, when the command -E a.cu -o a.cu.E is used, the first intermediate file is generated; when the command -E a.cu -oa.cu.E.dev is used, the second intermediate file is generated. Thus, the result of step 101 is: The preprocessing module processes macro definitions (such as #define), header file inclusions (such as #include), and conditional compilation directives (such as #if and #else) in CUDA source code files to generate code that has expanded macros and includes all header files. The preprocessed content is then saved as a.cu.E and a.cu.E.dev files, thus transforming the source code into a code form that conforms to the general compiler syntax rules during the preprocessing stage.

[0034] S102, scan the host-side code in the first intermediate file in the host-side code processing path to identify the target function identifier that represents the device-side execution semantics, and perform a reconstruction operation on the function containing the target function identifier to generate a proxy function for triggering the device-side execution.

[0035] In some embodiments, by performing macro expansion, a first intermediate file containing source code can be obtained. At this time, a partitioning operation can be performed to generate two logically equivalent intermediate code copies based on the first intermediate file, which are then used as inputs to the host-side code processing path and the device-side code processing path, respectively.

[0036] Among them, logically equivalent intermediate code copies refer to two source code copies that are identical in content in the initial state, and will be processed according to the compilation rules of the host side and the device side respectively.

[0037] Specifically, the intermediate code copy includes: host-side code compiled for the host side (i.e., the first source code copy) and device-side code compiled for the device side (i.e., the second source code copy).

[0038] Thus, since step S101 performs a macro expansion operation, the target function identifier that performs semantics on the device side can be identified when scanning the first intermediate file.

[0039] As mentioned above, GPU-side code cannot be used on the host side. Therefore, a refactoring operation is performed on the function containing the target function identifier so that the host side can call the GPU-side code.

[0040] Meanwhile, the host-side code processing path can also generate a proxy function `post_to_device` to trigger execution on the device side. In practical applications, the `post_to_device` function is a signal sent from the host side to the device side to execute the process on the device side; that is, the device side is the carrier of code execution.

[0041] In one example, the target function identifier includes the keyword __global__. Functions modified by __global__ run on the GPU (i.e., the device side) but can only be called by the CPU (i.e., the host side), and it marks the CUDA Kernel (kernel function).

[0042] It should be noted that, at runtime, although it may appear to be just the string "__global__", in the CUDA compilation context, it actually carries the complete GPU kernel semantics.

[0043] Accordingly, the function containing the target function identifier is refactored, including: The code between the `__global__` keyword and the first left curly brace `{` is defined as the function header; the code between the right curly brace `}` matching the left curly brace is defined as the function body; and the content between parentheses `()` in the function header is defined as the parameter list, thus extracting the function name, parameter list, and function body.

[0044] Specifically, in the first intermediate file, the location of the __global__ keyword is located; starting from the end of the keyword, the scan proceeds until the first appearance of the left curly brace {, and the code text between the two is determined as the function header.

[0045] The function body is defined by matching curly braces starting with the left curly brace {, identifying the matching right curly brace}, and determining the code text between the left curly brace { and the matching right curly brace}.

[0046] Furthermore, locate the pair of parentheses ( ) corresponding to the function declaration in the function header, determine the content between them as the parameter list, and determine the function name based on the identifier to the left of the left parenthesis ( ) of the parameter list, thus completing the extraction of the function name, parameter list, and function body.

[0047] In other words, by parsing functions containing the __global__ keyword, the function header, function body, and parameter list are determined, enabling automatic identification and structured extraction of CUDA Kernel functions. This provides the necessary basic information for converting Kernel functions into CPU-executable functions.

[0048] Furthermore, the bracket matching mechanism can accurately determine the boundaries of the kernel function body, avoiding parsing errors caused by nested structures and improving the accuracy and stability of the CUDA kernel function to CPU execution conversion process.

[0049] Furthermore, based on this, when generating the proxy function to trigger execution on the device side, a first function, a second function, and a third function are generated at the original location of the function containing the __global__ keyword. Among them, the first function is fn_get_mangled, which is used to send parameters and execution requests to the device side; the second function is fn, which is used as a parameter for the cudaLaunchKernel call; and the third function is fn__stub_kernelaunch__, which is used to replace the kernel boot syntax <<<>>>.

[0050] Using the above method, by generating three functions in situ, the conversion from CUDA paradigm <<<>>> to standard C++ paradigm is achieved, as well as the conversion of functions modified by the CUDA keyword __global__ to standard C++ functions. This enables the CUDAKernel function to run in the CPU, that is, to realize the transmission and parsing of device-side function information and the final execution of the CUDAKernel function.

[0051] In some embodiments, the second function can call the first function, and the third function can call the second function. This hierarchical calling mechanism effectively decouples program control logic, task scheduling logic, and computation logic. Thus, when running a CUDA program on the CPU, the second function can uniformly schedule the first function, reducing repetitive parameter configuration and resource initialization operations, lowering function call overhead, and improving the overall execution efficiency of the program.

[0052] Meanwhile, when the execution environment or computational logic of a CUDA program changes, only the first or second function needs to be adjusted, without modifying the control logic of the third function. This improves the compatibility and scalability of CUDA programs on different hardware platforms and enhances the stability of system operation.

[0053] In some embodiments, the generation process of the first function includes: The function is named fn_get_mangled and takes a list of parameters as its parameters. The thread-local variable funcBack stores the filename of the .so file to be executed on the device side, the name of the function to be executed, and the number of parameters. The thread-local variable funcBack is used when calling cudaLaunchKernel.

[0054] Specifically, the generation process of the first function includes: 1) Function definition method: The first function is named fn_get_mangled and uses a parameter list as the formal parameter list, where the parameter list corresponds to the parameter list of the device-side function to be executed.

[0055] 2) Saving the execution context: In the first function, the thread-local variable funcBack is used to save the relevant information of the function to be executed on the device side. The relevant information includes at least: the name of the dynamic link library (.so file) to be executed on the device side, the name of the function to be executed on the device side, and the number of parameters of the function to be executed.

[0056] Among them, the thread-local variable funcBack is a thread-private variable used to ensure that the execution contexts of different threads are isolated from each other in a multi-threaded environment.

[0057] 3) The timing of context usage, that is, the information stored in the thread-local variable funcBack is read and used when the CUDA kernel launch interface cudaLaunchKernel is called to determine the specific device-side function that needs to be launched.

[0058] By centralizing symbol generation and resolution in the first function, cross-platform compatibility can be improved by adapting to symbol rules under different compilers or different compilation options without changing the call entry point and emission process.

[0059] It's important to note that by binding the device-generated .so file to the first function, the semantic characteristics of the device-side function are preserved when the first function executes on the CPU, thus achieving effective simulation. This approach not only improves the compatibility and execution consistency of CUDA programs on non-GPU platforms but also reduces the need to modify existing CUDA code, lowering program porting and maintenance costs. Furthermore, this binding mechanism facilitates independent scheduling and optimization of the device-side function, enhancing system scalability and execution efficiency.

[0060] The second function contains a void Type parameter, and void The type parameter is assigned to the thread-local variable funcBack; and the parameter list is set as a local variable and fn_get_mangled is called.

[0061] Specifically, the second function is used to inject execution context into thread-local variables and trigger the device-side function resolution process, which includes the following: 1) Parameter settings: The second function contains a void... The parameter of type void The type parameter is used to carry the execution context information of the device-side function.

[0062] 2) Context assignment operation: When calling the second function, void... The execution context information pointed to by the type parameter is assigned to the thread-local variable funcBack to complete the initialization of the current thread's execution context.

[0063] 3) Parameter organization and function call: The second function saves the parameter list as a local variable, and after completing the assignment of the thread local variable, it calls the first function fn_get_mangled to obtain the device-side function identifier to be executed based on the execution context information.

[0064] By using void As a parameter form that carries the context, the context structure fields can be expanded (e.g., adding version number, verification information, permission identifier, call source, etc.) while keeping the function signature unchanged, thereby enhancing the evolution capability.

[0065] The parameter list of the third function includes a parameter list as well as dim3 gridDim, dim3 blockDim, size_tsm_size, and cudaStream_t stream. It then constructs an array of parameter pointers and calls cudaLaunchKernel((void )fn, gridDim, blockDim, args, m_size, stream).

[0066] Specifically, the third function is used to complete the parameter assembly of the device-side functions and the execution of the kernel launch, and it includes the following: 1) Function parameter settings: The parameter list of the third function includes: parameter list; grid dimension parameter dim3gridDim; thread block dimension parameter dim3 blockDim; dynamic shared memory size parameter size_t sm_size; CUDA stream parameter cudaStream_t stream.

[0067] 2) Construction of the parameter pointer array: In the third function, the parameter pointer array args is constructed based on the parameter list. Each element in the parameter pointer array is used to point to the address of the corresponding parameter's storage location on the host side.

[0068] 3) Kernel boot operation: The third function calls the CUDA kernel boot interface to start the device-side function to be executed on the GPU, where fn is the device-side function pointer obtained from the first function.

[0069] The above approach utilizes thread-local variables to pass the execution context of device-side functions between different functions, enabling the parsing and invocation of CUDA functions in the device-side dynamic link library. This approach can complete kernel startup without explicitly passing function pointers, exhibiting good thread safety and scalability, and is suitable for dynamic CUDA function invocation scenarios in multi-threaded environments.

[0070] In short, by introducing thread-local variables to store the execution context of device-side functions and combining multi-level function encapsulation, dynamic parsing and unified startup of CUDA functions in device-side dynamic link libraries are achieved.

[0071] In some embodiments, after refactoring the call to the function containing the __global__ keyword, the following are also performed: replacing the __global__ keyword with an empty string; replacing the __host__ modifier with an empty string for functions modified only by __host__; replacing the entire function with an empty string for functions modified only by __device__; and replacing both the __host__ and __device__ modifiers with an empty string while preserving the function implementation for functions modified by both __host__ and __device__.

[0072] By performing the above steps, it is beneficial to convert the host-side code into a host-side object file, and by replacing the original __global__ keyword with an empty string, it is possible to avoid continuing to compile in the CUDA kernel mode on the host side.

[0073] In this embodiment, the device-side .E.dev file is ultimately compiled into a .so file. This .so file is called within the device-side process. The first function returns the name of the .so file, the name of the function, and the number of arguments. The host side needs these three numbers to tell the device-side process which function in which .so file to execute. These three runtime parameters are stored in thread_local funcBack.

[0074] The second function, `fn`, has the same name as the original function. It is the first argument used to call `cudaLaunchKernel`. It is called within `cudaLaunchKernel`, which in turn calls `fn_get_mangled` to obtain the three arguments shown above. `cudaLaunchKernel` then uses these arguments to send them to the device process for execution.

[0075] The third function is used to replace the call to fn in the form <<<>>>, and its content is to call cudaLaunchKernel.

[0076] S103, Scan the device-side code in the second intermediate file in the device-side code processing path to identify the target function identifier, and perform semantic removal processing on the target function identifier to generate device-side executable function code.

[0077] In some embodiments, unlike the host-side processing method, the deletion operation is performed directly in the device-side code processing path.

[0078] Specifically, the reason for this difference is that the device side is the carrier of code execution, that is, the program on the device side is simulated and run on the CPU.

[0079] In the device-side code, the `__global__` keyword is removed while the function body remains unchanged, resulting in a function that can be executed on the device side. In other words, on the device side, the transformed `__global__` function retains its original business logic.

[0080] Specifically, when generating executable function code on the device side, global objects are processed as follows: for global variables modified by __device__, __device__ is replaced with __attribute__ ((section(.DEVICE))), and for global variables modified by __constant__, __constant__ is replaced with __attribute__ ((section(.CONST))). When determining the type of an object modified by __device__, the presence of a parameter list between __device__ and the first left parenthesis or semicolon is checked to distinguish whether __device__ modifies a function or a variable.

[0081] Thus, when processing objects modified by __device__, the system determines whether a parameter list exists between __device__ and the first left parenthesis or semicolon to distinguish whether it modifies a function or a variable, and then adopts different processing methods accordingly.

[0082] S104, generate a host-side target file based on the processed first intermediate file, generate a device-side target file based on the processed second intermediate file, and further generate a device-side dynamic link file.

[0083] In some embodiments, the host-side code is compiled based on the reconstructed host-side code to generate a host-side object file. This host-side object file contains function implementations for encapsulating execution requests and initiating kernel calls.

[0084] The process of generating the target file on the host side includes: After rewriting the __global__ modified functions, replacing the bracketed calls, and replacing or deleting function and variable modifiers, the processed CUDA source code is output as an intermediate source code file .E that has undergone macro expansion and semantic rewriting. The intermediate source code file .E is then compiled using a host-side compiler to generate the corresponding object file .o.

[0085] For example, the command g++ a.cu -o a.out is used to convert the processed first intermediate file into the target file .o.

[0086] Accordingly, after processing, the device-side code is compiled to generate the device-side object file, which specifically includes: The device-side code after removing the __global__ keyword, and the execution of the global object, will generate a device-side intermediate file, which is a .E.dev file or an intermediate representation equivalent to it.

[0087] Then, the .E.dev file is compiled using the device-side compilation tool to generate an object file containing device-side executable code; next, the object file is linked to generate a device-side shared library file .so, which is used to be loaded by the host side or called by the device-side runtime environment at runtime.

[0088] The device-side shared library file's name and function symbol information match the execution request information generated by the host side, enabling the function name and parameter information sent by the host side through the runtime interface to be parsed and executed in the device-side shared library file.

[0089] In short, the rules for generating .E.dev files are relatively simple; just clear __device__ and __global__. This is because: __device__ and __global__ functions run on the GPU, so they should retain their original state; global variables modified by __device__ and __constant__ are handled in the same way as on the host side.

[0090] It should be noted that although the __global__ keyword has been removed, the function body remains unchanged, so the business logic on the device side remains unchanged.

[0091] For example, the command g++ a.cu.E.dev -o a.ou.so is used to convert the processed second intermediate file into the target file .so.

[0092] S105, merge the device-side dynamic link file and the host-side target file to form a target file.

[0093] In some embodiments, the generated host-side object file is linked with the device-side dynamic link file to generate the final object file. Thus, at runtime, the object file can trigger device-side function execution through host-side functions, thereby achieving coordinated operation between the host and device sides.

[0094] In one example, step S105 may include: The binary content of the device-side dynamic link file is sequentially appended to the end of the host-side target file; a fixed-length identifier string is appended after the content of the device-side dynamic link file to identify the existence status of the device-side code; and a fixed-length description field describing the byte length of the device-side dynamic link file is appended after the fixed-length identifier string, forming a target file structure that can be parsed and split during runtime.

[0095] Specifically, the generated .so file is appended directly to the end of the original target file, along with a fixed-length magic string and a fixed-length .so file.

[0096] It's important to note that the reason for performing the merge is that on the host side, the target file generated after compilation with the .E file is the original target file, while the .E.dev file generates a .so file for execution on the device side. This results in an additional .so file compared to the original target file. By performing the merge, the CUDA program can be run.

[0097] Thus, through the compilation operation in the example above, the compiled target file can be run directly, which will call cudart in the simulation environment. This will split the original target file that was merged above and start a new process for the .so part to simulate the execution of the GPU part.

[0098] Specifically, this manifests in: From the generated intermediate file .E, the converted __global__ function can perform parameter packaging and function name extraction operations on the host side; the places where the CUDA paradigm <<<>> calls the __global__ function are correctly converted into the cudaLaunchKernel function; and from the generated .E.dev, the converted __global__ function retains the original business logic unchanged on the device side, so that the generated target program can run normally.

[0099] To better understand and illustrate the CUDA source program compilation method based on a general-purpose compiler in the embodiments of this application, an example is provided.

[0100] See Figure 2 The flowchart shown in this application embodiment illustrates a CUDA source code compilation method in a specific application scenario. Figure 2 As shown, the following can be executed: S201 performs macro expansion on the source code of the CUDA source program.

[0101] The result of macro expansion is that only intermediate files containing the source code are included.

[0102] Specifically, when using the command -E a.cu -o a.cu.E, a first intermediate file is generated on the host side, while when using the command -Ea.cu -o a.cu.E.dev, a second intermediate file is generated on the device side.

[0103] In this embodiment, the execution steps on the host side include: S202, scan the __global__ keyword in the first intermediate file.

[0104] S203 clears the content of the __global__ keyword and generates the first, second, and third functions in place; during the operation, the first, second, and third functions are sent to the device-side process via post_to_device.

[0105] One specific implementation of step S203 is as follows: After scanning __global__, the part between the first "{" after __global__ is the function header, the content between the "}" corresponding to the first "{" is the function body, and the part between "(" and ")" between __global__ and the first "{" is the parameter list. The function name (fn), parameter list (args), and function body (body) can all be extracted according to the above rules.

[0106] Three functions are generated at the original function's location (corresponding to the function whose __global__ keyword was deleted): fn_get_mangled is used to send parameters and execute requests to the GPU.

[0107] Here, fn is used as the fn parameter of cudaLaunchKernel, which is called by cudaLaunchKernel.

[0108] fn__stub_kernelaunch__ is used to replace the function that calls <<<>>>.

[0109] Generation of fn_get_mangled: The function name is fn_get_mangled, and (args) is the parameter list. There is a thread-local variable funcBack to store some variables. The file name of the so file to be executed on the GPU is saved in funcBack[0], the name of the function to be executed is saved in funcBack[1], and the number of parameters is saved in funcBack[2]. funcBack will be used when cudaLaunchKernel is called.

[0110] The generation of fn, where only one parameter is void. The function `back` contains `funcBack = back`, and the parameter list (`args`) is a local variable. Specifically, when calling `fn_get_mangled`, the parameters are obtained from `args`.

[0111] Generation of fn__stub_kernelaunch__: The parameter list is (args, dim3 gridDim, dim blockDim, size_t sm_size, cudaStream_tstream) with void as the empty string. args[number of arguments]; calls cudaLaunchKernel(void )fn, gridDim, blockDim, args, sm_size, stream).

[0112] In addition, __global__ will be replaced with empty; __host__ in functions only modified by __host__ will be directly replaced with empty; __device__ in functions only modified by __device__ will be completely replaced with empty; and __host__ and __device__ in functions modified by both will have both __host__ and __device__ replaced with empty.

[0113] S204 generates a host-side intermediate file (.E) based on the compiled first intermediate file.

[0114] S205, Generate host-side target file.

[0115] Correspondingly, the execution steps on the device side include: S206, Scan the __global__ keyword in the second intermediate file.

[0116] S207, Delete the keyword __global__.

[0117] S208 generates the device-side intermediate file .E.dev based on the compiled second intermediate file.

[0118] S209 generates the device-side target file and further generates a dynamically linked .so file.

[0119] S210 merges the device-side dynamic link file and the host-side target file to form a target file.

[0120] It should be understood that although the steps in the flowchart are shown sequentially according to the arrows, these steps are not necessarily executed in the order indicated by the arrows. Unless explicitly stated herein, there is no strict order constraint on the execution of these steps, and they can be executed in other orders. Moreover, at least some steps in the diagram may include multiple sub-steps or multiple stages. These sub-steps or stages are not necessarily completed at the same time, but can be executed at different times. The execution order of these sub-steps or stages is not necessarily sequential, but can be performed alternately or in turn with other steps or at least some of the sub-steps or stages of other steps.

[0121] Please see Figure 3 This application provides a schematic diagram of the structure of a CUDA source program compilation device based on a general-purpose compiler. The CUDA source program compilation device 300 based on a general-purpose compiler may include: The preprocessing module 310 is configured to perform preprocessing operations on the source code of the CUDA source program to generate a first intermediate file and a second intermediate file, respectively; the preprocessing operations include at least the expansion processing of macro definition-related instructions; The processing module 320 is configured to scan the host-side code in the first intermediate file in the host-side code processing path to identify the target function identifier representing the device-side execution semantics, and perform a refactoring operation on the function containing the target function identifier to generate a proxy function for triggering the device-side execution; and to scan the device-side code in the second intermediate file in the device-side code processing path to identify the target function identifier, and perform semantic removal processing on the target function identifier to generate device-side executable function code. The merging module 330 is configured to generate a host-side target file based on the processed first intermediate file, generate a device-side target file based on the processed second intermediate file, further generate a device-side dynamic link file, and merge the device-side dynamic link file and the host-side target file to form a target file.

[0122] For further details on the preprocessing module 310, the processing module 320, and the compilation module 330, please refer to the aforementioned examples.

[0123] It is understandable that the above division of modules is only a logical functional division. In actual implementation, they can be fully or partially integrated into a single physical entity, or they can be physically separated. Furthermore, the above units can be implemented by the processor calling software.

[0124] This application also provides a computer device, including: a memory and a processor, wherein the memory stores a computer program, and the processor executes the computer program to implement the steps of any of the methods in the foregoing embodiments.

[0125] See Figure 4 , Figure 4 This is a schematic diagram of the hardware structure of a computer device provided in an embodiment of this application.

[0126] Figure 4 The computer device shown includes a memory 41, a processor 42, and a transceiver 43. The processor 42 is coupled to the memory 41 and the transceiver 43. The memory 41 can be located inside or outside the terminal. The memory 41, processor 42, and transceiver 43 can be connected via a communication bus. The transceiver 43 is used to communicate with other devices or communication networks.

[0127] Optionally, the transceiver 43 can be a transmitter. The memory 41 stores a computer program that can run on the processor 42, and when the processor 42 runs the computer program, the transceiver 43 performs the steps of any of the methods in the foregoing embodiments provided in the above embodiments.

[0128] This application also provides a computer-readable storage medium storing a computer program thereon, which, when executed by a processor, implements the steps of any of the methods in the foregoing embodiments.

[0129] Storage media may include read-only memory (ROM), random access memory (RAM), magnetic disks, or optical disks, etc. The storage media may also include non-volatile or non-transitory memory, etc.

[0130] This application also provides a computer program product, including a computer program / instructions that, when executed by a processor, implement the steps of any of the methods in the foregoing embodiments.

[0131] It should be understood that in the embodiments of this application, the processor can be a central processing unit (CPU), or it can be other general-purpose processors, digital signal processors (DSPs), application-specific integrated circuits (ASICs), field-programmable gate arrays (FPGAs), or other programmable logic devices, discrete gate or transistor logic devices, discrete hardware components, etc. The general-purpose processor can be a microprocessor or any conventional processor.

[0132] It should also be understood that the memory in the embodiments of this application can be volatile memory or non-volatile memory, or may include both volatile and non-volatile memory. Non-volatile memory can be ROM, programmable read-only memory (PROM), erasable programmable read-only memory (EPROM), electrically erasable programmable read-only memory (EEPROM), or flash memory. Volatile memory can be random access memory (RAM), which is used as an external cache. By way of example, but not limitation, many forms of RAM are available, such as static random access memory (SRAM), dynamic random access memory (DRAM), synchronous dynamic random access memory (SDRAM), double data rate synchronous dynamic random access memory (DDR SDRAM), enhanced synchronous dynamic random access memory (ESDRAM), synchronous linked dynamic random access memory (SLDRAM), and direct rambus RAM (DRRAM).

[0133] Those skilled in the art will understand that embodiments of this application can be provided as methods, systems, or computer program products. Therefore, this application can take the form of a completely hardware embodiment, a completely software embodiment, or an embodiment combining software and hardware aspects. Furthermore, this application can take the form of a computer program product embodied on one or more computer-usable storage media (including but not limited to disk storage, CD-ROM, optical storage, etc.) containing computer-usable program code.

[0134] This application is described with reference to flowchart illustrations and / or block diagrams of methods, apparatus (systems), and computer program products according to embodiments of this application. It will be understood that each block of the flowchart illustrations and / or block diagrams, and combinations of blocks in the flowchart illustrations and / or block diagrams, can be implemented by computer program instructions. These computer program instructions can be provided to a processor of a general-purpose computer, special-purpose computer, embedded processor, or other programmable data processing apparatus to produce a machine, such that the instructions, which execute via the processor of the computer or other programmable data processing apparatus, generate instructions for implementing the flowchart... Figure 1 One or more processes and / or boxes Figure 1 A device that provides the functions specified in one or more boxes.

[0135] These computer program instructions may also be stored in a computer-readable storage medium that can direct a computer or other programmable data processing device to function in a particular manner, such that the instructions stored in the computer-readable storage medium produce an article of manufacture including instruction means, which are implemented in a process Figure 1 One or more processes and / or boxes Figure 1 The function specified in one or more boxes.

[0136] These computer program instructions may also be loaded onto a computer or other programmable data processing equipment to cause a series of operational steps to be performed on the computer or other programmable equipment to produce a computer-implemented process, thereby providing instructions that execute on the computer or other programmable equipment for implementing the process. Figure 1 One or more processes and / or boxes Figure 1 The steps of the function specified in one or more boxes.

[0137] Although preferred embodiments of this application have been described, those skilled in the art, upon learning the basic inventive concept, can make other changes and modifications to these embodiments. Therefore, the appended claims are intended to be interpreted as including the preferred embodiments as well as all changes and modifications falling within the scope of this application.

[0138] Obviously, those skilled in the art can make various modifications and variations to this application without departing from the spirit and scope of this application. Therefore, if such modifications and variations fall within the scope of the claims of this application and their equivalents, this application also intends to include such modifications and variations.

Claims

1. A method for compiling CUDA source programs based on a general-purpose compiler, characterized in that, include: Preprocessing operations are performed on the source code of the CUDA source program to generate the first intermediate file and the second intermediate file respectively; The preprocessing operation includes at least the expansion of macro definition-related instructions; Scan the host-side code in the first intermediate file in the host-side code processing path to identify the target function identifier that represents the device-side execution semantics, and perform a refactoring operation on the function containing the target function identifier to generate a proxy function for triggering the device-side execution; Scan the device-side code in the second intermediate file in the device-side code processing path to identify the target function identifier, and perform semantic removal processing on the target function identifier to generate device-side executable function code; A host-side target file is generated based on the processed first intermediate file, a device-side target file is generated based on the processed second intermediate file, and a device-side dynamic link file is further generated. The device-side dynamic link file and the host-side target file are merged to form a target file.

2. The compilation method according to claim 1, characterized in that, The preprocessing operation on the CUDA source code to generate a first intermediate file and a second intermediate file includes: The g++ compiler is used to perform preprocessing operations on source code files with the .cu extension. By calling preprocessing parameters, macro definitions, conditional compilation directives, and header file inclusion directives in the source code file are expanded, and the macro-expanded source code is output to the specified target file as an intermediate file. Specifically, the first intermediate file is generated when the command -E a.cu -o a.cu.E is used, and the second intermediate file is generated when the command -E a.cu -oa.cu.E.dev is used.

3. The compilation method according to claim 1, characterized in that, The objective function identifier includes: the __global__ keyword; The refactoring of the function containing the target function identifier includes: determining the code between the __global__ keyword and the first left curly brace { as the function header; determining the code between the right curly brace} matching the left curly brace as the function body; and determining the content between the parentheses ( ) in the function header as the parameter list, thereby extracting the function name, parameter list, and function body.

4. The compilation method according to claim 3, characterized in that, When generating the proxy function to trigger execution on the device side, a first function, a second function, and a third function are generated at the original location of the function containing the __global__ keyword; The first function is fn_get_mangled, which sends parameters and execution requests to the device side; the second function is fn, which is used as a parameter for calling cudaLaunchKernel; the third function is fn__stub_kernelaunch__, which replaces the kernel boot syntax <<<>>>; and the second function can call the first function, and the third function can call the second function.

5. The compilation method according to claim 4, characterized in that, The generation process of the first function includes: using fn_get_mangled as the function name and the parameter list as the formal parameters; using the thread-local variable funcBack to store the file name of the .so file to be executed on the device side, the name of the function to be executed, and the number of parameters; wherein, the thread-local variable funcBack is used when calling cudaLaunchKernel; The second function contains a void Type parameter, and the void The type parameter is assigned to the thread-local variable funcBack; and the parameter list is used as a local variable, and fn_get_mangled is called; The parameter list of the third function includes the parameter list as well as dim3 gridDim, dim3 blockDim, size_t sm_size, and cudaStream_t stream. Then, it constructs an array of parameter pointers and calls cudaLaunchKernel((void )fn, gridDim, blockDim, args, m_size, stream).

6. The compilation method according to any one of claims 3 to 5, characterized in that, After refactoring the call to the function containing the __global__ keyword, the following is also executed: Replace the __global__ keyword with an empty string; For functions that are only modified by __host__, replace the __host__ modifier with an empty string; For functions that are only modified by __device__, replace the entire function with an empty string; For functions modified by both __host__ and __device__, replace the __host__ and __device__ modifiers with empty strings while preserving the function implementation.

7. The compilation method according to claim 1, characterized in that, When generating device-side executable function code, the global object is processed as follows: For global variables modified by __device__, replace __device__ with __attribute__ ((section(.DEVICE))); and for global variables modified by __constant__, replace __constant__ with __attribute__ ((section(.CONST))). Specifically, when determining the type of the object modified by __device__, the presence of a parameter list between __device__ and the first left parenthesis or semicolon is checked to distinguish whether __device__ modifies a function or a variable.

8. The compilation method according to claim 1, characterized in that, The step of merging the device-side dynamic link file and the host-side target file to form a target file includes: The binary content of the device-side dynamic link file is sequentially appended to the end of the host-side target file; A fixed-length identifier string is appended to the content of the device-side dynamic link file to identify the existence status of the device-side code; A fixed-length description field describing the byte length of the device-side dynamic link file is appended after the fixed-length identifier string, forming a target file structure that can be parsed and split during runtime.

9. A CUDA source program compilation device based on a general-purpose compiler, characterized in that, include: The preprocessing module is configured to perform preprocessing operations on the source code of the CUDA source program, generating a first intermediate file and a second intermediate file respectively; The preprocessing operation includes at least the expansion of macro definition-related instructions; The processing module is configured to scan the host-side code in the first intermediate file in the host-side code processing path to identify the target function identifier representing the device-side execution semantics, and perform a refactoring operation on the function containing the target function identifier to generate a proxy function for triggering the device-side execution; and to scan the device-side code in the second intermediate file in the device-side code processing path to identify the target function identifier, and perform semantic removal processing on the target function identifier to generate device-side executable function code. The merging module is configured to generate a host-side target file based on the processed first intermediate file, generate a device-side target file based on the processed second intermediate file, further generate a device-side dynamic link file, and merge the device-side dynamic link file and the host-side target file to form a target file.

10. A computer device, comprising: A memory and a processor, the memory storing a computer program, characterized in that the processor executes the computer program to implement the steps of the method according to any one of claims 1 to 8; And / or, a computer-readable storage medium having a computer program stored thereon, characterized in that, when the computer program is executed by a processor, it implements the steps of the method according to any one of claims 1 to 8; And / or, a computer program product comprising a computer program / instructions, characterized in that, when executed by a processor, the computer program / instructions implement the steps of the method according to any one of claims 1 to 8.

Images