Compilation method and device for converting CUDA language into BANGC language

Abstract

The invention provides a method for compiling a CUDA language into a BANGC language. The method comprises the steps that CUDA original codes are obtained; representing a CUDA kernel of the CUDA original code as a structured loop nested UPA function; performing cyclic fission of inter-thread operation on the UPA function so as to divide the UPA function into one or more UPA units; and performing vectorization optimization on each UPA unit in the UPA function in sequence, and converting the optimized UPA function into a BANGC target code. The invention further provides a device for compiling the CUDA language into the BANGC language. According to the method for compiling the CUDA language to the BANGC language, language conversion from the CUDA language to the BANGC language is achieved through multi-layer intermediate representation.

Description

Technical Field

[0001] This invention relates to the field of high-performance computing technology, specifically to a method and apparatus for compiling CUDA language into BANGC language. Background Technology

[0002] The memory architectures of NVIDIA GPUs and Cambricon MLUs differ significantly, requiring a multi-layered, systematic design for mapping from CUDA to BANGC. Current compilation technologies do not support direct conversion from CUDA to Cambricon BANGC source code via translation or compilation; instead, programmers must write the code manually, making it impossible to reuse existing, rich CUDA programs.

[0003] In the process of converting CUDA to Cambricon's BANGC language, the significant differences in memory architecture between NVIDIA GPUs and Cambricon MLUs (GPUs have registers, shared memory, global memory, etc., while MLUs have registers, GDRAM, LDRAM, NRAM, NRAM, SRAM, etc., indicating different memory hierarchies) necessitate a systematic multi-layered design for the mapping conversion. Furthermore, when mapping CUDA programs to BANGC programs, the differences in parallel computing methods require vectorization optimization. Therefore, a new intermediate representation (IR) needs to be designed to achieve the CUDA to BANGC language conversion. Summary of the Invention

[0004] This invention proposes a method for compiling CUDA language into BANGC language, comprising: retrieving the original CUDA code; representing the CUDA kernel of the original CUDA code as a structured nested UPA function; performing inter-thread operation loop fission on the UPA function to divide it into one or more UPA units; sequentially performing vectorization optimization on each UPA unit in the UPA function; and converting the optimized UPA function into BANGC target code.

[0005] Furthermore, the specific steps for representing the CUDA kernel as a UPA function include: converting the original CUDA code into a UPA function. , ;in, For grid level loops, For block-level loops, For the warp level loop, For thread-level loops, It is the kernel. The connector indicates that the UPA function consists of loops at various levels and their internal loop bodies; the complex arithmetic operations of the CUDA source code's built-in calls and derivation of warp-related indices are converted into direct symbolic references to the parallel loop induction variables in the UPA function.

[0006] Furthermore, the specific steps of dividing the UPA function into UPA units include: constructing and dividing the UPA unit by performing loop fission on inter-thread operations within the UPA function. , , It is the cross-thread block loop level of the UPA function. It is the corresponding part within that thread block. It is the loop body of this UPA unit.

[0007] Preferably, the method further includes: using the CNCC compiler of BANGC to compile the BANGC target code into an executable file; and executing the executable file on the Cambrian MLU to obtain the target result.

[0008] This invention also proposes a compilation device for converting CUDA language to BANGC language, comprising: an initialization module for acquiring CUDA source code; an intermediate representation module for representing the CUDA kernel of the CUDA source code as a structured, nested UPA function; an optimization module for performing inter-thread operation loop fission on the UPA function to divide it into one or more UPA units; performing vectorization optimization on each UPA unit in the UPA function in sequence; and a code conversion module for converting the optimized UPA function into BANGC target code.

[0009] Furthermore, the intermediate representation module includes: converting the CUDA source code into UPA functions. , ;in, For grid level loops, For block-level loops, For the warp level loop, For thread-level loops, It is the kernel. This indicates that the UPA function consists of loops at various levels and their inner loop bodies; the complex arithmetic operations involving built-in calls and warp-related index derivations in the original CUDA code are converted into direct symbolic references to the inductive variables of the parallel loops within the UPA function.

[0010] Furthermore, the optimization module specifically includes: constructing and dividing the UPA unit by performing cyclic fission on inter-thread operations within the UPA function. , , It is the cross-thread block loop level of the UPA function. It is the corresponding part within that thread block. It is the loop body of this UPA unit.

[0011] Preferably, it also includes: a compilation and execution module for using the CNCC compiler of BANGC to compile the BANGC target code into an executable file; and executing the executable file on the Cambricon MLU to obtain the target result.

[0012] The present invention also proposes an electronic device including a CUDA language to BANGC language compilation device as described above.

[0013] The present invention also proposes a computer-readable storage medium storing computer-executable instructions, characterized in that, when the computer-executable instructions are executed, the compilation method from CUDA language to BANGC language as described above is implemented.

[0014] The CUDA language to BANGC language compilation method of the present invention can translate the input CUDA code into BANGC code with the same computational logic. Then, it is compiled by Cambricon CNCC compiler and host program. The resulting executable file can run on Cambricon MLU and obtain the expected results. Attached Figure Description

[0015] Figure 1A This is a diagram of a CUDA code segment.

[0016] Figure 1B This is a schematic diagram of the BANGC code segment.

[0017] Figure 2 This is a flowchart of the compilation method for converting CUDA language to BANGC language according to the present invention.

[0018] Figure 3A This is a schematic diagram of CUDA example code in an embodiment of the present invention.

[0019] Figure 3B This is a schematic diagram of the unified intermediate representation of an embodiment of the present invention.

[0020] Figure 3C This is an optimized intermediate representation diagram of an embodiment of the present invention.

[0021] Figure 3D This is a schematic diagram of BANGC example code in an embodiment of the present invention.

[0022] Figure 4 This is a schematic diagram of the CUDA language to BANGC language compilation device of the present invention.

[0023] Figure 5 This is a schematic diagram of an electronic device according to the present invention.

[0024] Figure 6 This is a schematic diagram of the hardware structure of an electronic device according to the present invention. Detailed Implementation

[0025] To make the objectives, technical solutions, and advantages of this invention clearer, the invention will be further described in detail below with reference to the accompanying drawings. It should be understood that the specific embodiments described herein are merely illustrative of the invention and are not intended to limit the invention.

[0026] It should be noted that, in this application, relational terms such as "first" and "second" are used merely to distinguish one entity or operation from another, and do not necessarily require or imply any such actual relationship or order between these entities or operations. Furthermore, the terms "comprising," "including," or any other variations thereof are intended to cover non-exclusive inclusion, such that a process, method, article, or apparatus that comprises a list of elements includes not only those elements but also other elements not expressly listed, or elements inherent to such a process, method, article, or apparatus.

[0027] In the absence of further restrictions, an element defined by the phrase "comprising a..." does not exclude the presence of other identical elements in the process, method, article, or apparatus that includes said element.

[0028] The core of this invention lies in achieving CUDA ecosystem compatibility from the compilation architecture perspective, providing a way to migrate the CUDA ecosystem to computing power solutions represented by Cambricon, and supporting the efficiency and ecosystem compatibility of MLU intelligent chips. The key to achieving this goal is that the mapping and conversion from CUDA language to BANGC language (cuda2bang) requires a systematic design at multiple levels. Therefore, a new intermediate representation (IR) needs to be designed to support execution on Cambricon MLU.

[0029] Intermediate representations (IRs) of the compiler play an important role in the compilation process. Depending on the purpose of the compiler, IRs can currently be divided into three categories: scalar IRs, vector IRs, and tensor IRs.

[0030] 1) Scalar IR. The most representative scalar IR is LLVM IR. LLVM IR is a low-level intermediate representation designed as a universal intermediate language between the compiler front-end and back-end. Generated by the LLVM compiler, it represents the control flow, data flow, and type information of the source code. LLVM IR features type safety, platform independence, optimization potential, easy extensibility, and high readability.

[0031] 2) Vector IR. MLIR (Multi-Level Intermediate Representation) is a multi-level intermediate representation language used to describe, transform, and optimize programs from various programming languages. Vector Dialect is a branch of MLIR; it is a vector IR primarily used to describe and optimize vector computation programs. Specifically, VectorDialect provides several types of vectors, including fixed-length and variable-length vectors. Fixed-length vectors specify the length of the vector and can perform computations on the vector using built-in and custom operations. Variable-length vectors, on the other hand, allow the length to be dynamically determined at runtime, providing greater flexibility and portability.

[0032] Vector Dialect provides various vector operations, such as vector addition and multiplication. It also offers advanced operations like vector transpose. In addition to built-in operations, Vector Dialect supports custom operations. MLIR's Vector Dialect also provides optimization techniques such as automatic vectorization and loop unrolling. These techniques are automatically applied to vector computation programs during compilation, improving performance and efficiency. Overall, Vector Dialect provides compiler developers with a powerful tool for describing and optimizing vector computation programs.

[0033] 3) Tensor IR. Deep learning frameworks typically use tensors as data structures. Currently, common tensor IRs include Triton IR and TVM's Tensor IR.

[0034] Triton-IR is an intermediate representation language based on LLVM, designed to provide an environment suitable for tile-level program analysis, transformation, and optimization. Triton-IR programs share the same high-level structure as LLVM-IR programs, such as modules, functions, and basic blocks, but Triton-IR also includes extensions necessary for tile-level data flow and control flow analysis. To support tile-level data flow analysis, Triton-IR extends the types and instructions of LLVM-IR.

[0035] Tensor IR is part of TVM. Tensor IR is an intermediate representation language for describing and optimizing tensor computations, providing a general way to represent and optimize computations across different frameworks and hardware platforms. A Tensor IR program mainly consists of three parts: multidimensional buffers, nested loops (which may have thread binding in GPU environments), and blocks. A block is the most important abstraction in Tensor IR; a block can contain one or more nested loops, with sub-blocks or a series of imperative statements corresponding to the computation. This representation allows users to divide computations into corresponding sub-block regions and perform efficient program transformations using dependency information stored in the block signature.

[0036] The core of this invention lies in: the intermediate representation mapping from CUDA language to Cambrian BANGC language, exploring the intermediate representation design method and generation technology of the compiler; and designing and implementing an intermediate representation that supports CUDA-BANGC mapping at the intermediate representation level of the compiler.

[0037] CUDA code snippets (such as...) Figure 1A (as shown) and BANGC code segment (as shown) Figure 1B The kernel code for vector addition (as shown) differs significantly. The inventors have conducted a comparative analysis of the memory management and computational models of CUDA and BANGC programming models as follows:

[0038] The CUDA programming model employs an implicit memory management mechanism, using thread IDs (tid) for automated data access. Specifically, kernel functions can directly use the passed-in global memory pointer, employing the thread index (tid) as the array subscript for element-by-element computation. This design achieves efficient parallel computing through hardware-automated memory access and thread scheduling. In contrast, the BANGC programming model uses an explicit memory management strategy, requiring developers to perform multi-stage memory operations. Its execution flow includes three key steps: first, explicitly allocating temporary storage space in on-chip NRAM; second, actively moving global GDRAM data to NRAM via the __memcpy interface; and finally, calling vector computation interfaces (such as __bang_add) for batch vector operations. After computation, the result data must be explicitly written back from NRAM to global GDRAM. This explicit management approach provides developers with finer-grained control but also increases programming complexity. From the implementation mechanisms, it is clear that the two frameworks differ fundamentally in memory management strategies, data movement methods, and computational granularity: at the programming paradigm level, CUDA adopts a thread-based parallel computing model. Its core features are: achieving computational parallelism through a warp mechanism, with each thread performing scalar operations. This design makes the code structure of CUDA kernel functions highly similar to traditional C language, allowing developers to directly use pointers and arrays for element-wise operations, with thread IDs (tid) acting as implicit indices to automatically manage memory access. The BANGC programming model, on the other hand, is built on a vector computing architecture, its essential characteristics being explicit memory management and batch vector operations.

[0039] In recent years, the inventors have conducted a series of studies on compilation systems and optimization for heterogeneous chips, including smart chips. During this process, the inventors believe that there is a significant gap between the software ecosystem for smart chips, exemplified by Cambricon, and the CUDA ecosystem. This gap is mainly reflected in two aspects:

[0040] (1) The software ecosystem of intelligent chips, represented by Cambricon, is not complete enough.

[0041] Compared to the CUDA ecosystem, the software ecosystem for intelligent chips, exemplified by Cambricon, suffers from a significant deficiency. This deficiency manifests in two aspects. First, the completeness of development tools: CUDA provides a comprehensive set of development tools covering programming languages ​​(CUDA C / C++), high-performance libraries for various domains (cuDNN deep learning acceleration library, cuBLAS linear algebra library, cuFFT Fourier transform library, cuSparse sparse computation library), and code generation tools like CUTLASS. Second, there is a lack of robust support for upper-layer deep learning frameworks. NVIDIA GPUs have become the de facto standard for intelligent chips in the deep learning field, and various open-source deep learning frameworks also provide backends for NVIDIA GPUs, i.e., "application frameworks adapt to chip manufacturers." However, due to a less developed user base and development environment, intelligent chips like those from Cambricon receive less proactive adaptation from deep learning frameworks, requiring "chip manufacturers to adapt application frameworks." This results in the deep learning frameworks supported by intelligent chips like Cambricon lagging behind the pace of framework development.

[0042] (2) The intelligent chip software ecosystem, represented by Cambricon, cannot form a synergy.

[0043] When intelligent chips need to be integrated into future intelligent application software ecosystems, customized development of upper-layer frameworks is usually required. For example, when TensorFlow supports different instruction sets, it needs to provide corresponding operator implementations for different backends. That is, the Cambricon MLU backend and the NVIDIA GPU backend need to provide operator implementations for Cambricon BANGC and CUDA, respectively. However, due to the fragmentation of intelligent chip instruction sets, this customization work is tightly coupled with the instruction set and is difficult to reuse across chips. This means that software ecosystems of different instruction sets are unlikely to form a synergy at the application level.

[0044] The aforementioned gap stems from the "siloed" development strategy of the current chip software ecosystem, exemplified by Cambricon, resulting in a lack of technical means to support the coexistence of the software ecosystem. If a CUDA-compatible compilation system existed that could map the rich intelligent applications within the CUDA ecosystem onto intelligent chips like those from Cambricon, it would have significant practical implications for the coexistence of the intelligent chip ecosystem.

[0045] To address the issue that the significant differences in intelligent chip architectures, exemplified by Cambricon, make it difficult for intelligent applications to fully utilize the parallel computing power of intelligent chips, research is being conducted on key technologies for CUDA-compatible compilation infrastructure for the Cambricon architecture. This research aims to achieve CUDA ecosystem compatibility from the compilation architecture perspective, providing a solution for migrating foreign ecosystems to domestic computing power represented by Cambricon, and supporting the efficiency and ecosystem compatibility of domestically produced intelligent chips.

[0046] The core of this invention, focusing on CUDA-compatible compilation technology, is how the compiler can convert between SIMT and SIMD data parallelism representations, enabling CUDA programs to run not only on SIMT-based GPUs but also efficiently migrate to the Cambricon MLU with a SIMD architecture. To address this, the cuda2bang compiler must extend the expressive capabilities of the intermediate representation: enabling it to express multithreading in SIMT and multichanneling in SIMD, allowing the compiler to understand multithreading semantics and complete the SIMT-SIMD conversion through cross-thread code transformation. Therefore, the following specific technical problems need to be solved:

[0047] 1) Research on intermediate representation mapping from CUDA language to Cambrian BANGC language

[0048] At the intermediate representation level of the compiler, an intermediate representation supporting SIMT-SIMD compilation is designed and implemented. This representation converts thread-level parallelism into loops, enabling the application of classic loop analysis compilation techniques for vectorization, thereby reducing it to vector instructions. Through code transformation, CUDA is automatically mapped to the above intermediate representation, and then mapped from this intermediate representation to BANGC.

[0049] This intermediate representation serves as a bridge for the mapping and conversion from CUDA to the Cambrian BANGC language. During the conversion, the CUDA language must first be converted to this intermediate representation. This intermediate representation is then vectorized and converted into BANGC source code.

[0050] 2) Build a compiler that translates CUDA language to IR, vectorizes IR, and then translates IR to BANGC.

[0051] The Intermediate Representation (IR) of this invention is a multi-level parallel representation, which can be used to more comprehensively characterize the parallel semantics in CUDA programs. Unlike previous representations that primarily focused on thread-level and thread-block-level parallelism, this invention explicitly represents each level of the CUDA hierarchy—block, warp, and thread—with equal importance in the program's Intermediate Representation (IR), thus forming a unified multi-level structure. While existing work like COX introduced the concept of warp-level parallelism, it is only recognized when explicit warp-level functions are used.

[0052] The cuda2bang compiler of this invention is built on MLIR and utilizes MLIR's multidimensional parallel loops (such as affine.parallel) to achieve this multi-level representation. Through program analysis, this invention transforms the originally opaque built-in calls (such as blockIdx and threadIdx) and the complex arithmetic operations (such as division and modulo) used to derive warp-related indices (such as warp_id and lane_id) into direct symbolic references to these parallel loop inductive variables.

[0053] The intermediate representation of this invention is a Unified Parallelism Abstraction (UPA), a unified form for representing CUDA kernels. UPA makes the previously implicit parallel semantics explicit, thus facilitating vectorized execution. This invention builds UPA on nested loops, a practice motivated by the inherent thread-level parallelism of CUDA grids, blocks, and warps.

[0054] First, this invention represents the CUDA kernel (e.g., a __global__ function) as a UPA function.

[0055]

[0056] in Indicates in dimension A cycle, This is the kernel body. This construction transforms parallelism at all levels in CUDA (grids, thread blocks, and warps) into structured nested loops, thereby capturing the global parallel behavior of the entire kernel in a unified representation. It is important to note that this invention explicitly models the warp hierarchy, as warp semantics are crucial for enabling warp-based optimizations.

[0057] Secondly, a UPA function is divided into one or more UPA units. Each UPA unit is defined as the outermost, perfectly nested nest of loops that can be vectorized as a whole. The UPA unit is the basic unit of vectorization: each unit must use one and only one vectorization strategy, while different units can use different strategies. A UPA unit is defined as follows:

[0058]

[0059] in It is an inter-block loop hierarchy inherited from the UPA function. It refers to the portion within the intra-block. It is the loop body of the unit. This invention utilizes the UPA unit. The nested loops within are called its parallel regions (PRs).

[0060] Finally, this invention also proposes a loop fission based on inter-thread operations: the UPA unit is constructed by performing loop fission on inter-thread operations within the UPA function. Inter-thread operations refer to CUDA built-in functions involving communication between threads, including: (i) synchronization operations (e.g., __syncthreads within a thread block, __syncwarp within a warp), or (ii) data exchange (e.g., _shfl_sync). Formally, given a loop body... nested loops ,if Includes inter-thread operations , then Split into And generate two units.

[0061]

[0062] For synchronous operations, the operation will be removed after the cyclic fission. ,make sure exist Previously, it was fully executed, thus maintaining CUDA's "execute-synchronize-execute" semantics. For data exchange operations, It will be retained as and Explicit data transfer between loops. This mechanism ensures that nested loops can faithfully express CUDA's inter-thread semantics in a barrier-free manner and are suitable for independent vectorization.

[0063] Figure 2 This is a flowchart of the compilation method for converting CUDA language to BANGC language according to the present invention. Figure 2 As shown, the following uses the vector addition program vecadd as an example to introduce the specific steps of the CUDA language to BANGC language compilation method of this invention:

[0064] The test environment was a Cambricon MLU 370 environment equipped with an x86 CPU (Intel Xeon Gold 6330 CPU), using the Ubuntu 20.04.6 LTS operating system, and deploying the cuda2bang compiler of this invention.

[0065] Step S1: Prepare the original CUDA-written vector addition program vecadd.cu; for example... Figure 3A As shown.

[0066] Step S2: Open the cuda2bang compiler terminal and use cuda2bang in the command line to translate vecadd.cu into vecadd.mlu. The content of vecadd.mlu is a vector addition operation written in Cambrian BANGC language.

[0067] The specific command is `cuda2bang -i vecadd.cu -o vecadd.mlu --bang-mlu-arch MLU370--parallel-size 256`. During the execution of this command, the CUDA program will first be converted to... Figure 3B The initial IR shown is then vectorized and optimized to obtain the following: Figure 3C The optimized IR is shown, and then the optimized IR is converted to Figure 3D The BANGC source code shown.

[0068] Step S3: Using the host program vecadd.cpp, compile it together with the vecadd.mlu obtained in the previous step using the official BANGC compiler CNCC to create the executable file vecadd for Cambricon MLU. The specific command is: cnccvecadd.cpp vecadd.mlu -o vecadd --bang-mlu-arch=mlu370

[0069] Step S4: Execute the executable file vecadd compiled in step S3 on the Cambrian MLU370 to obtain the expected result.

[0070] It should be noted that, in various embodiments of the present invention, the order of the steps does not imply the order of execution. The execution order of each step should be determined by its function and internal logic, and should not constitute any limitation on the implementation process of the embodiments of the present invention.

[0071] The following are system embodiments corresponding to the above method embodiments. This embodiment can be implemented in conjunction with the above embodiments. The relevant technical details mentioned in the above embodiments are still valid in this embodiment, and will not be repeated here to reduce repetition. Accordingly, the relevant technical details mentioned in this embodiment can also be applied to the above embodiments.

[0072] Figure 4 This is a schematic diagram of the CUDA language to BANGC language compilation device of the present invention. Figure 4As shown, in a second embodiment of the present invention, a compilation apparatus 10 is provided, comprising:

[0073] Initialize module 11 to obtain the CUDA source code;

[0074] Intermediate representation module 12 is used to represent the CUDA kernel of the original CUDA code as a structured loop-nested UPA function;

[0075] Optimization module 13 is used to perform inter-thread operation loop fission on the UPA function to divide it into one or more UPA units; and to perform vectorization optimization on each UPA unit in the UPA function in sequence.

[0076] Code conversion module 14 is used to convert the optimized UPA function into BANGC target code;

[0077] The compilation and execution module 15 is used to compile the BANGC target code into an executable file using the CNCC compiler of BANGC; the executable file is then executed on the Cambricon MLU to obtain the target result.

[0078] In a third embodiment of the present invention, a computer-readable storage medium is proposed. The CUDA-to-BANGC language compilation apparatus of the present invention, if its functionality is implemented as a software functional unit and sold or used as an independent product, can be stored in a computer-readable storage medium. Based on this understanding, the technical solution of the present invention, in essence, or the part that contributes to the prior art, or a part of the technical solution, can be embodied in the form of a software product. This computer software product is stored in a computer-readable storage medium and includes several instructions to cause a computer device (which may be a personal computer, a server, or a network device, etc.) to execute all or part of the steps of the methods described in the various embodiments of the present invention. Therefore, in the third embodiment of the present invention, a computer-readable storage medium is provided for storing a computer program of a CUDA-to-BANGC language compilation method. It should be understood that the computer-readable storage medium in the embodiments of the present invention can be volatile memory and / or non-volatile memory. The non-volatile memory can be read-only memory (ROM), programmable read-only memory (PROM), erasable programmable read-only memory (EPROM), electrically erasable programmable read-only memory (EEPROM), or flash memory. The volatile memory can be random access memory (RAM), which serves as an external cache. By way of example, but not limitation, many forms of random access memory (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).

[0079] Figure 5 This is a schematic diagram of an electronic device according to the present invention. Figure 5As shown, in the fourth embodiment of the present invention, an electronic device 100 is proposed, including a CUDA-to-BANGC language compilation device as described above. Those skilled in the art will understand that all or part of the steps in the above methods can be implemented by a program instructing related hardware (e.g., processor, FPGA, ASIC, etc.). All or part of the steps in the above embodiments can also be implemented using one or more integrated circuits. Accordingly, each module in the above embodiments can be implemented in hardware, such as by an integrated circuit to implement its corresponding function, or it can be implemented as a software functional module, such as by a processor executing a program / instruction stored in memory to implement its corresponding function. The embodiments of the present invention are not limited to any particular combination of hardware and software.

[0080] It should be noted that the structure of the electronic device shown in the accompanying drawings of this invention does not constitute a limitation thereof. The actual knowledge structure recognition device may include more or fewer components than shown, or combine certain components, or have different component arrangements.

[0081] The electronic device of the present invention can be any device with data processing capabilities, such as a computer or other similar device. The device embodiment can be implemented in software, hardware, or a combination of both. Taking software implementation as an example, as a logical device, it is formed by the processor of any data processing device loading the corresponding computer program instructions from non-volatile memory into memory for execution. Figure 6 This is a schematic diagram of the hardware structure of an electronic device according to the present invention. Figure 6 As shown, from a hardware perspective, this is a hardware structure diagram of any data processing-capable device where the CUDA language to BANGC language compilation device of this invention is located, except... Figure 6 In addition to the processor, memory, network interface, and non-volatile memory shown, any data processing device in the embodiment may also include other hardware depending on the actual function of the data processing device, which will not be described in detail here.

[0082] When implemented using software, the above embodiments can be implemented, in whole or in part, as a computer program product. The computer program product includes one or more computer instructions or computer programs. When the computer instructions or computer programs are loaded or executed on a computer, all or part of the processes or functions described in the embodiments of the present invention are generated. The computer can be a general-purpose computer, a special-purpose computer, a computer network, or other programmable device. The computer instructions can be stored in a computer-readable storage medium or transmitted from one computer-readable storage medium to another. For example, the computer instructions can be transmitted from one website, computer, server, or data center to another website, computer, server, or data center via wired (e.g., infrared, wireless, microwave, etc.) means. 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 data center that includes one or more sets of available media. The available medium can be a magnetic medium (e.g., floppy disk, hard disk, magnetic tape), an optical medium (e.g., DVD), or a semiconductor medium. A semiconductor medium can be a solid-state drive.

[0083] The CUDA language to BANGC language compilation method of the present invention, by building a cuda2bang compiler, inputs CUDA code, which can be translated into BANGC code with the same computational logic. Then, the Cambricon CNCC compiler and the host program are used to compile together, and the resulting executable file can run on the Cambricon MLU and obtain the expected results.

[0084] The above embodiments are only used to illustrate the present invention and are not intended to limit the present invention. Those skilled in the art can make various changes and modifications without departing from the spirit and scope of the present invention. Therefore, all equivalent technical solutions also fall within the scope of the present invention, and the patent protection scope of the present invention should be defined by the claims.

Claims

1. A method for compiling CUDA language into BANGC language, characterized in that, include: Obtain the CUDA source code; The CUDA kernel of the original CUDA code is represented as a UPA function with nested structured loops; The UPA function is subjected to inter-thread operation loop fission to divide it into one or more UPA units; each UPA unit in the UPA function is then vectorized and optimized sequentially. The optimized UPA function is converted into BANGC target code.

2. The compilation method as described in claim 1, characterized in that, The specific steps for representing the CUDA kernel as a UPA function include: Convert the CUDA source code into a UPA function , ;in, For grid level loops, For block-level loops, For the warp level loop, For thread-level loops, It is the kernel. The connector indicates that the UPA function consists of loops at various levels and their internal loop bodies; the complex arithmetic operations of the CUDA source code's built-in calls and derivation of warp-related indices are converted into direct symbolic references to the parallel loop induction variables in the UPA function.

3. The compilation method as described in claim 1, characterized in that, The specific steps for dividing the UPA function into UPA units include: The UPA unit is constructed and partitioned by performing cyclic fission on inter-thread operations within the UPA function. , , It is the cross-thread block loop level of the UPA function. It is the corresponding part within that thread block. It is the loop body of this UPA unit.

4. The compilation method as described in claim 1, characterized in that, Also includes: Use the CNCC compiler for BANGC to compile the BANGC object code into an executable file; Execute the executable file on Cambrian MLU to obtain the desired result.

5. A compiler for converting CUDA language to BANGC language, characterized in that, include: An initialization module is used to obtain the CUDA source code; The intermediate representation module is used to represent the CUDA kernel of the original CUDA code as a structured loop-nested UPA function; The optimization module is used to perform inter-thread operation loop fission on the UPA function to divide it into one or more UPA units; and to perform vectorization optimization on each UPA unit in the UPA function in sequence. The code conversion module is used to convert the optimized UPA function into BANGC target code.

6. The compilation apparatus as described in claim 5, characterized in that, The intermediate representation module includes: converting the CUDA source code into UPA functions. , ;in, For grid level loops, For block-level loops, For the warp level loop, For thread-level loops, It is the kernel. The connector indicates that the UPA function consists of loops at various levels and their internal loop bodies; the complex arithmetic operations of the CUDA source code's built-in calls and derivation of warp-related indices are converted into direct symbolic references to the parallel loop induction variables in the UPA function.

7. The compilation apparatus as described in claim 5, characterized in that, The optimization module specifically includes: constructing and dividing the UPA unit by performing cyclic fission on inter-thread operations within the UPA function. , , It is the cross-thread block loop level of the UPA function. It is the corresponding part within that thread block. It is the loop body of this UPA unit.

8. The compilation apparatus as described in claim 5, characterized in that, Also includes: The compiler and run module is used to compile the BANGC object code into an executable file using the CNCC compiler for BANGC. Execute the executable file on Cambrian MLU to obtain the desired result.

9. An electronic device comprising a CUDA language to BANGC language compilation device as described in any one of claims 5 to 8.

10. A computer-readable storage medium storing computer-executable instructions, characterized in that, When the computer-executable instructions are executed, the method for compiling CUDA language into BANGC language as described in any one of claims 1 to 4 is implemented.

Images