Kernel compiling method, running method, compiling device, programming system and chip

Abstract

Embodiments of the present application provide a kernel compiling method, a running method, a compiling device, a programming system and a chip, and relate to the technical field of artificial intelligence chips. In the compiling stage, kernel code containing a first kernel function and a second kernel function is obtained, two-level processing units are distinguished by using execution space attributes, cross-space calling relationship is automatically recognized according to function calling relationship in the kernel code, cross-space calling processing code based on shared storage is automatically rewritten and generated by a compiler, and parameter calling and execution synchronization across address spaces are automatically completed at a bottom layer, so that a user does not need to manually allocate a shared storage buffer and manage data carrying details, programming complexity and the probability of manual errors are reduced. Meanwhile, target code is generated to adapt to scenarios in which different processing units correspond to different instruction set architectures, so that processing chips with heterogeneous instruction sets can be loaded and cooperatively scheduled at one time during running, and hardware performance of the chips is maximally released.

Description

Technical Field

[0001] This application relates to the field of artificial intelligence chip technology, and in particular to a kernel compilation method, running method, compilation device, programming system and chip. Background Technology

[0002] With the development of heterogeneous computing technology, single-level kernel programming models such as CUDA and OpenCL are widely used in GPU architectures, achieving heterogeneous computing through a unified kernel entry point and hardware automatic caching mechanisms. However, for computing chips with physical topologies featuring separate multi-level processing units, the different processing units often employ different instruction set architectures and reside in independent execution spaces.

[0003] The single-level programming models for GPUs in related technologies cannot be directly adapted to computing chips with multi-level processing unit architectures. Specifically, the orchestration logic and computation logic of single-level kernel models are mixed, making it difficult to perform differential optimizations for the instruction sets corresponding to different processing units. They also lack explicit programming abstractions for shared memory orchestration, requiring users to manually manage data movement and buffers, resulting in high programming complexity. Furthermore, the execution and address spaces of different processing units are independent, lacking automatic cross-level parameter passing mechanisms and failing to support automatic allocation and injection of cross-level barrier synchronization, making it difficult to guarantee the timing correctness of cross-level hardware collaborative execution. Summary of the Invention

[0004] The main objective of this application is to propose a kernel compilation method, a running method, a compilation device, a programming system, and a chip to solve the problems of poor adaptability of programming models, complex development, and unstable cross-level execution in computing chips with multi-level processing units.

[0005] To achieve the above objectives, a first aspect of this application proposes a kernel compilation method, comprising:

[0006] Obtain kernel code, the kernel code including first code of at least one first kernel function and second code of at least one second kernel function, each first kernel function having a first execution space attribute, and each second kernel function having a second execution space attribute different from the first execution space attribute;

[0007] The kernel code is compiled and parsed, and at least one cross-space call relationship is identified based on the function call relationship in the kernel code. Each cross-space call relationship is formed by a first kernel function calling a second kernel function.

[0008] Based on the cross-space call relationships, the first code and the second code corresponding to each cross-space call relationship are compiled and rewritten respectively. A first cross-space call processing code is generated in the corresponding first code, and a second cross-space call processing code is generated in the corresponding second code. A first intermediate representation corresponding to the first kernel function and a second intermediate representation corresponding to the second kernel function are obtained. The first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship are used to pass call parameters and perform execution synchronization when the corresponding first kernel function and second kernel function are executed by the first processing unit and the second processing unit respectively.

[0009] Based on the first execution space attribute and the second execution space attribute, the first intermediate representation and the second intermediate representation are respectively compiled into first target code adapted to the first processing unit and second target code adapted to the second processing unit, and a unified executable file is generated based on the first target code and the second target code.

[0010] In some embodiments, the kernel code includes execution space declaration information for each of the first kernel functions and each of the second kernel functions;

[0011] The first execution space attribute is determined based on the execution space declaration information of each of the first kernel functions, and the second execution space attribute is determined based on the execution space declaration information of each of the second kernel functions;

[0012] The execution space declaration information includes the execution space identifier set in the function declaration of the corresponding kernel function.

[0013] In some embodiments, the step of compiling and rewriting the first code and the second code corresponding to each cross-space call relationship based on each cross-space call relationship includes:

[0014] The parameter layout of the call parameters involved in each of the cross-space call relationships is determined in the corresponding parameter buffer. Each parameter buffer is located in a shared storage area, which can be accessed by the first processing unit and the second processing unit.

[0015] Based on the parameter layout, parameter writing code is generated in the corresponding first code, and parameter retrieval code is generated in the corresponding second code;

[0016] The parameter writing code is used to write the call parameters involved in the cross-space call relationship into the corresponding parameter buffer according to the corresponding parameter layout, and the parameter retrieval code is used to retrieve the call parameters from the corresponding parameter buffer according to the corresponding parameter layout.

[0017] In some embodiments, determining the parameter layout of the call parameters involved in each of the cross-space call relationships in the corresponding parameter buffer includes:

[0018] Identify the dynamic call parameters in the call parameters;

[0019] The parameter byte length of each dynamic call parameter is determined according to the data type of each dynamic call parameter;

[0020] The aligned length of each dynamic calling parameter is determined based on the parameter byte length of each dynamic calling parameter and the preset alignment rules;

[0021] Based on the arrangement order and aligned length of each of the dynamic calling parameters, determine the storage offset of each of the dynamic calling parameters in the parameter buffer;

[0022] The size of the parameter buffer is determined based on the storage offset and aligned length of each of the dynamically invoked parameters.

[0023] In some embodiments, generating parameter writing code in the corresponding first code and parameter retrieval code in the corresponding second code according to each parameter layout includes:

[0024] Based on the size of the parameter buffer, a buffer allocation code is generated in the corresponding first code, and the buffer allocation code is used to allocate the parameter buffer in the shared storage area;

[0025] Based on the storage offset of each of the dynamic call parameters, a parameter storage code is generated in the corresponding first code. The parameter storage code is used to write each of the dynamic call parameters into the corresponding storage location in the parameter buffer.

[0026] In the corresponding second code, a buffer address acquisition code is generated, which is used to obtain the starting address of the parameter buffer in the shared storage area;

[0027] Based on the starting address, the storage offset and data type of each of the dynamic call parameters, parameter reading code is generated in the corresponding second code. The parameter reading code is used to determine the storage address of each of the dynamic call parameters and read the corresponding parameter value from each of the storage addresses.

[0028] The parameter writing code includes the buffer allocation code and the parameter storage code, and the parameter acquisition code includes the buffer address acquisition code and the parameter reading code.

[0029] In some embodiments, the step of compiling and rewriting the first code and the second code corresponding to each cross-space call relationship based on each cross-space call relationship further includes:

[0030] Identify the compile-time constant parameters in the call parameters;

[0031] Based on the parameter values ​​of the compile-time constant parameters, constant assignment code is generated in the corresponding second code.

[0032] The compile-time constant parameters are not written to the parameter buffer.

[0033] In some embodiments, each of the first cross-space call processing codes includes call initiation code and synchronization waiting code;

[0034] Each of the aforementioned call initiation codes is used to cause the first processing unit to initiate a corresponding call request, and each of the aforementioned call requests is used to cause the second processing unit to execute a corresponding second kernel function;

[0035] The synchronous waiting code is used to enable the first processing unit to wait for the corresponding second kernel function to complete execution after initiating the corresponding call request.

[0036] In some embodiments, generating the first cross-space call processing code in the corresponding first code includes:

[0037] Identify at least one entry function that has the first execution space attribute;

[0038] Assign a synchronization object to each of the aforementioned entry functions;

[0039] Based on the cross-space call relationships, call initiation code is generated in the corresponding first code. Each call initiation code is used to generate a call request carrying the identification information of the corresponding synchronization object.

[0040] Determine the calling nodes in the corresponding first code for each of the cross-space calling relationships, and insert synchronization waiting code based on the corresponding synchronization object.

[0041] In some embodiments, multiple cross-space call relationships within the same entry function use the same synchronization object;

[0042] The synchronization wait code is also used to reset the synchronization state of the synchronization object after the corresponding second kernel function is executed, so that the synchronization object can be reused by the next cross-space call relationship in the entry function.

[0043] In some embodiments, second cross-space call processing code is generated in the corresponding second code to obtain a second intermediate representation corresponding to the second kernel function, including:

[0044] Identify multiple second kernel functions that have the second execution space attribute and do not have a mutual calling relationship;

[0045] Determine the function identifiers of the plurality of second kernel functions respectively;

[0046] A unified entry function is generated for the multiple second kernel functions;

[0047] In the unified entry function, identifier acquisition code and call dispatch code are generated. The identifier acquisition code is used to obtain the function identifier carried in the call request, and the call dispatch code is used to dispatch the execution flow to the corresponding second kernel function according to the function identifier.

[0048] Each of the first cross-space call processing codes is also used to generate a call request carrying the function identifier of the second kernel function to be called.

[0049] In some embodiments, compiling and rewriting the first code corresponding to each of the cross-space call relationships includes:

[0050] Identify the data transfer operations in the first code;

[0051] Data transfer code is generated based on the data transfer operation. The data transfer code is used to transfer relevant computational data from the global storage area to the shared storage area.

[0052] Based on the source data layout and target data layout of the computational data, data layout conversion code is generated in the data transport code. The data layout conversion code is used to convert the computational data from the source data layout to the target data layout when transporting the computational data.

[0053] In some embodiments, the source data layout is a row-first layout, and the target data layout is a row-block layout.

[0054] Alternatively, the source data layout is a column-first layout, and the target data layout is a column-block layout;

[0055] The data transport code is also used to transport the computational data commonly used by the multiple second kernel functions to the shared storage area before the execution of the multiple second kernel functions, so that the multiple second kernel functions can reuse it.

[0056] In some embodiments, the first processing unit corresponds to an orchestrator in the computing chip, and the second processing unit corresponds to an execution unit in the computing chip;

[0057] The orchestrator is used to perform shared storage management, data movement, and execution unit task scheduling; the execution unit is used to perform computational processing.

[0058] The first object code is adapted to a 32-bit RISC architecture, and the second object code is adapted to a 64-bit RISC architecture with vector extensions.

[0059] To achieve the above objectives, a second aspect of this application provides a kernel running method, comprising:

[0060] Load a unified executable file, the unified executable file including a first target code adapted to a first processing unit and a second target code adapted to a second processing unit, the first target code including a first cross-space call processing code corresponding to at least one cross-space call relationship, and the second target code including a second cross-space call processing code corresponding to each of the cross-space call relationships;

[0061] The first processing unit runs the first target code. When it reaches the call node corresponding to any cross-space call relationship in the first target code, the corresponding first cross-space call processing code is executed. The call parameters involved in the cross-space call relationship are written into a shared storage area that can be accessed by both the first processing unit and the second processing unit. A call request is initiated and the corresponding second kernel function is waited for to be executed.

[0062] The second processing unit runs the second target code, and in response to the call request, executes the corresponding second cross-space call processing code, obtains the call parameters from the shared storage area, and executes the second kernel function determined by the call request;

[0063] After the second kernel function completes execution, the first processing unit ends its waiting and continues to run the first target code.

[0064] To achieve the above objectives, a third aspect of this application provides a kernel compilation apparatus, comprising:

[0065] The code acquisition module is used to acquire kernel code, the kernel code including first code of at least one first kernel function and second code of at least one second kernel function, each first kernel function having a first execution space attribute, and each second kernel function having a second execution space attribute different from the first execution space attribute;

[0066] The call relationship identification module is used to compile and parse the kernel code, and identify at least one cross-space call relationship based on the function call relationship in the kernel code. Each cross-space call relationship is formed by a first kernel function calling a second kernel function.

[0067] The code rewriting module is used to compile and rewrite the first code and the second code corresponding to each cross-space call relationship based on each cross-space call relationship. It generates a first cross-space call processing code in the corresponding first code and a second cross-space call processing code in the corresponding second code, thereby obtaining a first intermediate representation corresponding to the first kernel function and a second intermediate representation corresponding to the second kernel function. The first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship are used to pass call parameters and perform execution synchronization when the corresponding first kernel function and second kernel function are executed by the first processing unit and the second processing unit, respectively.

[0068] The target code generation module is used to compile the first intermediate representation and the second intermediate representation into first target code adapted to the first processing unit and second target code adapted to the second processing unit, respectively, based on the first execution space attribute and the second execution space attribute, and to generate a unified executable file based on the first target code and the second target code.

[0069] To achieve the above objectives, a fourth aspect of this application proposes a kernel programming system, comprising:

[0070] The programming interface module is used to provide an execution space declaration interface and to set execution space declaration information for kernel functions in the obtained kernel code through the execution space declaration interface;

[0071] The compilation module is used to determine the execution space attributes of each kernel function according to the execution space declaration information, identify at least one cross-space call relationship according to the function call relationship in the kernel code, generate a first cross-space call processing code and a second cross-space call processing code based on each cross-space call relationship, and compile the compiled and rewritten kernel code into target code adapted to the first processing unit and the second processing unit, respectively.

[0072] The runtime module is used to load a unified executable file generated based on each of the target codes, so that the first processing unit and the second processing unit can complete the passing of call parameters and execution synchronization through the first cross-space call processing code and the second cross-space call processing code corresponding to each of the cross-space call relationships.

[0073] To achieve the above objectives, a fifth aspect of the present application provides a chip including a first processing unit, a second processing unit, and a shared memory region;

[0074] The shared storage area can be accessed by the first processing unit and the second processing unit;

[0075] The first processing unit is used to run the first target code, and when it runs to the call node corresponding to any cross-space call relationship, it writes the call parameters involved in the cross-space call relationship into the shared storage area through the corresponding first cross-space call processing code, initiates a call request and waits for the corresponding second kernel function to complete execution;

[0076] The second processing unit is used to run the second target code, and in response to the call request, obtains the call parameters from the shared storage area through the corresponding second cross-space call processing code, and executes the second kernel function determined by the call request;

[0077] The first processing unit is further configured to end the waiting period and continue running the first target code after the second kernel function has been executed;

[0078] The first target code is compiled from at least one first kernel function with a first execution space attribute, and the second target code is compiled from at least one second kernel function with a second execution space attribute, wherein the first execution space attribute is different from the second execution space attribute.

[0079] The kernel compilation method, execution method, compilation device, programming system, and chip proposed in this application include the following: The kernel compilation method includes: obtaining kernel code, which includes first code of at least one first kernel function and second code of at least one second kernel function. Each first kernel function has a first execution space attribute, and each second kernel function has a second execution space attribute different from the first execution space attribute; compiling and parsing the kernel code, identifying at least one cross-space call relationship based on the function call relationships in the kernel code, where each cross-space call relationship is formed by a first kernel function calling a second kernel function; and compiling and rewriting the first code and second code corresponding to each cross-space call relationship based on the cross-space call relationship. The first code generates a first cross-space call processing code, and the corresponding second code generates a second cross-space call processing code, thus obtaining a first intermediate representation corresponding to the first kernel function and a second intermediate representation corresponding to the second kernel function. The first and second cross-space call processing codes corresponding to each cross-space call relationship are used to pass call parameters and synchronize execution when the corresponding first kernel function and second kernel function are executed by the first processing unit and the second processing unit, respectively. According to the first execution space attribute and the second execution space attribute, the first intermediate representation and the second intermediate representation are compiled into a first target code adapted to the first processing unit and a second target code adapted to the second processing unit, respectively, and a unified executable file is generated based on the first target code and the second target code.

[0080] In this embodiment, kernel code containing a first kernel function and a second kernel function is obtained during the compilation stage. The two levels of processing units are distinguished using first and second execution space attributes. Cross-space call relationships are automatically identified based on function call relationships in the kernel code. The compiler automatically rewrites and generates cross-space call processing code based on shared memory, automatically completing parameter calls and execution synchronization across address spaces at the underlying level. This eliminates the need for users to manually allocate shared memory buffers and manage data transfer details, reducing programming complexity and the probability of human error. Simultaneously, based on the first and second execution space attributes, the first and second intermediate representations corresponding to the kernel functions are compiled into first target code adapted to the first processing unit and second target code adapted to the second processing unit, respectively. This adapts to scenarios where different processing units correspond to different instruction set architectures. The compiled first and second target codes are then packaged into a unified executable file, enabling the heterogeneous instruction set processing chip to load and coordinately schedule all at once during runtime, maximizing the release of the chip's hardware performance. Attached Figure Description

[0081] Figure 1 This is a flowchart of the kernel compilation method provided in the embodiments of this application.

[0082] Figure 2 This is a flowchart provided in the embodiments of this application, which describes the compilation and rewriting of the first and second codes to generate cross-level parameter passing logic based on shared storage.

[0083] Figure 3 This is a process provided in the embodiments of this application for determining the parameter layout of the call parameters involved in each cross-space call relationship in the corresponding parameter buffer.

[0084] Figure 4 This is a flowchart provided in this application embodiment, which generates parameter writing code in the corresponding first code and parameter acquisition code in the corresponding second code according to the layout of each parameter.

[0085] Figure 5 This is a flowchart of generating the first cross-space call processing code in the corresponding first code, provided in an embodiment of this application.

[0086] Figure 6 This is a flowchart provided in the embodiments of this application, in which second cross-space call processing code is generated in the corresponding second code to obtain the second intermediate representation corresponding to the second kernel function.

[0087] Figure 7 This is a flowchart of compiling and rewriting the first code corresponding to each cross-space call relationship, provided in an embodiment of this application.

[0088] Figure 8 This is a flowchart provided in an embodiment of the present application, showing how to compile and rewrite the first code and the second code to generate cross-level barrier synchronization logic.

[0089] Figure 9 This is a schematic diagram of the serialization and deserialization process provided in the embodiments of this application.

[0090] Figure 10 This is a flowchart of a kernel running method provided in another embodiment of this application.

[0091] Figure 11 This is a schematic diagram of the architecture of the kernel programming system provided in the embodiments of this application.

[0092] Figure 12 This is a schematic diagram of the kernel programming system provided in the embodiments of this application.

[0093] Figure 13 This is a structural block diagram of a kernel compilation device provided in another embodiment of this application. Detailed Implementation

[0094] To make the objectives, technical solutions, and advantages of this application clearer, the following detailed description is provided in conjunction with the accompanying drawings and embodiments. It should be understood that the specific embodiments described herein are merely illustrative and not intended to limit the scope of this application.

[0095] It should be noted that although functional modules are divided in the device schematic diagram and the logical order is shown in the flowchart, in some cases, the steps shown or described may be performed in a different order than the module division in the device or the order in the flowchart.

[0096] Unless otherwise defined, all technical and scientific terms used herein have the same meaning as commonly understood by one of ordinary skill in the art to which this application belongs. The terminology used herein is for the purpose of describing embodiments of this application only and is not intended to limit this application.

[0097] With the development of heterogeneous computing technology, CUDA and OpenCL are the mainstream heterogeneous computing programming frameworks, widely used in GPU architecture. They both adopt a single-level kernel programming model and achieve heterogeneous computing by means of a unified kernel entry point and hardware automatic caching mechanism.

[0098] In GPU architecture, shared memory management within thread blocks and data movement from global memory to shared memory are handled automatically by the hardware cache controller. Programmers only need to use `__shared__` declarations and explicit synchronization (`__syncthreads()`). GPU hardware provides a unified memory consistency model and automatic caching strategy, requiring no software involvement in data movement orchestration.

[0099] The CUDA programming framework organizes computation at the thread, thread block, and grid levels. Kernel functions are declared with the `__global__` keyword and execute on the device. Shared memory (SM) is declared with `__shared__`, and data movement is automatically handled by hardware, such as the automatic management of L1 / L2 cache and texture cache. Programmers do not need to explicitly orchestrate data movement from DRAM to shared memory. OpenCL uses work-items and work-groups as its levels. Kernel functions execute uniformly, and local memory (LM) is declared with the `__local` qualifier, also relying on hardware caching mechanisms for data movement.

[0100] Specifically, the heterogeneous compilation frameworks for GPUs include Triton and MLIR. Triton is a DSL compiler for GPUs that compiles the Python kernel into PTX. Its programming model remains single-level, with one kernel corresponding to one GPU kernel. Shared memory management and data movement are automatically inserted by the compiler, but there is no spatial separation between the orchestrator and the computation unit. MLIR is a compiler infrastructure that supports multi-layered intermediate representations, such as Affine, Linalg, and Vector Dialects, and can be used to build heterogeneous compilation pipelines, but it does not provide a two-level kernel programming abstraction for CIM hardware.

[0101] For computing chips with a physical topology featuring separate multi-level processing units, these units often employ different instruction set architectures and reside in independent execution spaces. For example, Compute-In-Memory (CIM) chips have a two-level physical topology with separate orchestrators and execution units. Data movement from global memory to shared memory requires explicit software initiation, and the two hardware levels use different instruction set architectures. The orchestrator (or meta-thread) is responsible for data movement, shared memory management, and computation scheduling, and is compiled with a reduced instruction set architecture, such as RISC-V32. The execution unit (or warp-thread) performs vectorized computation and is compiled with an instruction set architecture with vector extensions, such as RISC-V64+RVV. The memory hierarchy includes global memory (DRAM), shared memory (SMEM), scratchpad memory (SPM), and vector local memory (DLM), among others. Data movement from DRAM to SMEM requires explicit software initiation. It is evident that the physical topology of the CIM chip naturally features a two-level separation between the orchestrator and the computing unit, which is fundamentally different from the unified streaming multiprocessor (SM) architecture of the GPU.

[0102] The single-level programming model for GPUs in related technologies cannot be directly adapted to computing chips with multi-level processing unit architectures. This results in several drawbacks.

[0103] Firstly, it cannot map the physical topology of the CIM hardware. The single-level kernel model of CUDA / OpenCL unifies all computations into a single kernel entry point, lacking separation between the orchestrator and the computation unit. This makes it impossible to directly map the two-level physical structure of the CIM chip's orchestrator / execution unit. Forcing the use of a single-level model on the CIM chip leads to a mixture of orchestrator and computation logic, poor code maintainability, and an inability to leverage the unique two-level instruction set differences (RISCv32 vs. RISCv64 + RVV) for targeted optimization.

[0104] Secondly, there is a lack of explicit SMEM orchestration abstraction. On GPUs, data transfer from DRAM to SMEM is automatically handled by hardware cache, while CIM chips require explicit software initiation of the transfer, such as through DMA or a copy engine. The programming models in related technologies do not provide a programming abstraction for "allocating SMEM and transferring data at the orchestrator level and reusing it at the compute unit level," which forces CIM users to manually manage transfer timing, address alignment, and buffer lifetimes, resulting in high programming complexity and a high risk of errors.

[0105] Thirdly, there is a lack of cross-level call and parameter passing mechanisms. In heterogeneous programming frameworks in related technologies, parameter passing between kernels is based on a unified address space or runtime API. In the CIM chip, the orchestrator runs in the RISCv32 address space, while the computing unit runs in the RISCv64 address space. Since they have different address space views, parameter passing requires serialization / deserialization across the SMEM. Frameworks in related technologies do not provide such automatic cross-level parameter passing mechanisms.

[0106] Fourthly, there is a lack of cross-level synchronization automation. On GPUs, __syncthreads() only takes effect within a single core's thread block. The YIS to EU call of the CIM chip requires the boaddr(BarrierObjectAddress) mechanism to make the orchestrator wait for the computation unit to complete, and the programming model in related technologies does not support the automatic allocation and injection of such cross-level barriers.

[0107] Therefore, the single-level programming model for GPUs in related technologies cannot be directly adapted to computing chips with multi-level processing unit architectures. This computing chip may be, for example, an in-memory computing chip, but is not limited to this.

[0108] Based on this, embodiments of this application provide a kernel compilation method, execution method, compilation device, programming system, and chip. During the compilation phase, kernel code containing a first kernel function and a second kernel function is obtained. The two levels of processing units are distinguished using first and second execution space attributes. Cross-space call relationships are automatically identified based on function call relationships in the kernel code. The compiler automatically rewrites and generates cross-space call processing code based on shared memory, automatically completing parameter calls and execution synchronization across address spaces at the underlying level. This eliminates the need for users to manually allocate shared memory buffers and manage data transfer details, reducing programming complexity and the probability of human error. Simultaneously, based on the first and second execution space attributes, the first and second intermediate representations corresponding to the kernel functions are compiled into first target code adapted to the first processing unit and second target code adapted to the second processing unit, respectively. This adapts to scenarios where different processing units correspond to different instruction set architectures. The compiled first and second target codes are then packaged into a unified executable file, enabling this heterogeneous instruction set processing chip to load and coordinate at runtime, maximizing the release of the chip's hardware performance.

[0109] This application provides a kernel compilation method, a running method, a compilation device, a programming system, and a chip, which are specifically described through the following embodiments. First, the chip of this application embodiment is described.

[0110] In one embodiment, the chip includes a first processing unit, a second processing unit, and a shared memory region, which can be accessed by the first processing unit and the second processing unit.

[0111] The first processing unit executes the first target code. Upon reaching a call node corresponding to any cross-space call relationship, it writes the call parameters involved in the cross-space call relationship into the shared memory area using the corresponding first cross-space call processing code, initiates a call request, and waits for the corresponding second kernel function to complete execution. The second processing unit executes the second target code. In response to the call request, it retrieves the call parameters from the shared memory area using the corresponding second cross-space call processing code and executes the second kernel function determined by the call request. The first processing unit also terminates the waiting process and continues executing the first target code after the second kernel function completes execution. The first target code is compiled from at least one first kernel function with a first execution space attribute, and the second target code is compiled from at least one second kernel function with a second execution space attribute, wherein the first execution space attribute and the second execution space attribute are different. The specific execution process will be described in detail below.

[0112] In one embodiment, the first processing unit corresponds to an orchestrator in the computing chip, and the second processing unit corresponds to an execution unit in the computing chip. The orchestrator is used to perform shared memory management, data movement, and execution unit task scheduling, while the execution unit is used to perform computational processing. The first target code is adapted to a 32-bit RISC architecture, and the second target code is adapted to a 64-bit RISC architecture with vector extensions.

[0113] The orchestrator's 32-bit reduced instruction set architecture corresponds to the control hardware hub at the bottom of the computing chip. Since the core responsibilities of this control hardware are only control-related tasks such as memory allocation, data movement, and task scheduling, without requiring complex mathematical calculations, the use of a 32-bit address width and general-purpose register mechanism can meet the control scenario requirements with a smaller chip area and lower power consumption. The execution unit's 64-bit architecture provides more address space to support the memory addressing requirements of computational data. Vector extension technology supports hardware single-instruction multiple-data operations, meaning that dozens or even hundreds of data elements can be processed concurrently with a single underlying instruction, meeting the computational power requirements of the execution unit for executing artificial intelligence operators or large-scale matrix operations. It is important to understand that the instruction set architecture described here is for illustrative purposes only and does not represent a limitation.

[0114] The kernel compilation method in this application embodiment is described below in conjunction with the structure of the aforementioned computing chip.

[0115] Figure 1 This is an optional flowchart of the kernel compilation method provided in the embodiments of this application. Figure 1 The method may include, but is not limited to, steps 110 to 140. It is also understood that this embodiment... Figure 1 The order of steps 110 to 140 is not specifically limited. The order of steps can be adjusted or some steps can be reduced or added according to actual needs.

[0116] Step 110: Obtain the kernel code.

[0117] In one embodiment, the computing chip physically includes an orchestrator responsible for data handling and task scheduling, as well as an execution unit for computing power output. To map this hardware topology, this embodiment first reads two levels of kernel code written by the user. The kernel code includes first code for at least one first kernel function and second code for at least one second kernel function. Each first kernel function has a first execution space attribute, and each second kernel function has a second execution space attribute different from the first execution space attribute. The kernel code includes execution space declaration information for each first and second kernel function. The first execution space attribute is determined based on the execution space declaration information of each first kernel function, and the second execution space attribute is determined based on the execution space declaration information of each second kernel function. The execution space declaration information includes an execution space identifier set in the function declaration of the corresponding kernel function.

[0118] It is understandable that the first kernel function can be the kernel function corresponding to the orchestrator, the first code is the programmer code corresponding to the programmer, and the first execution space attribute indicates that the kernel function corresponds to the programmer space. The second kernel function can be the kernel function corresponding to the execution unit, the second code is the execution unit code corresponding to the execution unit, and the second execution space attribute indicates that the kernel function corresponds to the execution unit space.

[0119] In one embodiment, the first kernel function can be the kernel function corresponding to the orchestrator, which corresponds to the control of the computing chip. Its corresponding first code is the orchestrator code corresponding to the orchestrator. The corresponding code logic mainly includes: partitioning the shared memory buffer through allocation primitives, moving the raw data required for computation from global storage to shared storage through data moving primitives, and initiating scheduling calls to the underlying computing tasks. For example, the first execution space attribute of the first kernel function can be declared through the @ez.jit(space=ez.YIS) decorator, indicating that the first kernel function corresponds to the orchestrator space, and its corresponding code is determined as the first code. Subsequently, it is compiled into a RISCv32 target adapted to the orchestrator.

[0120] The second kernel function can be the kernel function corresponding to the execution unit, which corresponds to the computation layer of the computing chip. Its corresponding second code is the execution unit code, mainly used for the core algorithm implementation. It includes vectorized operation logic such as tensor multiplication and addition, reduction and summation, and activation functions declared using various vector types. For example, the `@ez.jit(space=ez.EU)` decorator can be used to declare the second execution space attribute of the second kernel function, indicating that the second kernel function corresponds to the execution unit space, and its corresponding code is determined as the second code. Subsequently, it is compiled into a RISCv64+RVV target adapted to the execution unit to achieve vectorized computation, which can be converted into 64-bit computation instructions with vector extension characteristics.

[0121] It is understandable that the specific content of the first and second codes can be determined based on the actual application scenario. When writing the corresponding execution logic, users do not need to handle low-level interaction details such as memory allocation and data transfer across levels. This embodiment reduces the programming complexity for users by obtaining the first and second codes and directly mapping the code structure to the two-level physical topology of the computing chip.

[0122] Step 120: Compile and parse the kernel code, and identify at least one cross-space call relationship based on the function call relationships in the kernel code.

[0123] In one embodiment, each cross-space call relationship is formed by a first kernel function calling a second kernel function. The obtained kernel code can be Python code. After obtaining the user-written first and second code, the compiler reads the code, performs preliminary parsing, and obtains an Abstract Syntax Tree (AST). The compiler then converts the AST into an Intermediate Representation (IR), which is determined based on the actual code. During the conversion from source code to the intermediate representation, the compiler performs compilation, parsing, and classification labeling on the kernel code, identifying which parts are the first code corresponding to the first kernel function of the first processing unit and which parts are the second code corresponding to the second kernel function of the second processing unit.

[0124] Because users use decorators as execution space declaration information to indicate code ownership when writing kernel code, the first kernel function (with an execution space identifier corresponding to the first execution space attribute in its function declaration) takes on the role of scheduling and allocation, while the second kernel function (with an execution space identifier corresponding to the second execution space attribute) takes on the role of computing power output. Therefore, the compiler can identify the execution space attributes corresponding to each kernel function based on the execution space declaration information. These execution space attributes include the first execution space attribute and the second execution space attribute. By parsing the corresponding execution space identifier, the compiler tags each generated intermediate representation node with either the first or second execution space attribute.

[0125] After all intermediate representation nodes are assigned execution space attributes, the compiler can construct a complete function call graph and determine cross-space call relationships where a first execution space attribute calls a second execution space attribute. Specifically, the compiler performs a global traversal along the function call graph. When it detects a call from a first kernel function labeled with the first execution space attribute to a second kernel function labeled with the second execution space attribute, the compiler identifies this specific call point based on the function call relationship and marks it as a cross-space call relationship. In this way, it distinguishes between ordinary function calls within the same space and calls across heterogeneous hardware layers.

[0126] Therefore, this embodiment uses abstract syntax tree parsing and spatial labeling mechanisms to convert user-written high-level language code into intermediate representations with clear spatial attributes and well-defined structures, namely, the first intermediate representation and the second intermediate representation. By identifying cross-space call relationships, all task scheduling points across hardware levels are located, providing the compiler with location information for generating the first and second cross-space call handling code for passing call parameters and execution synchronization in the next step. This makes the heterogeneous hardware interaction logic completely transparent to the user.

[0127] In one embodiment, for the intermediate representations obtained above, the compiler performs standardized compilation optimizations and transformations based on execution space attributes to obtain the corresponding first cross-space call handling code and second cross-space call handling code. Compilation optimizations and transformations may include for loop unrolling, data type inference, insertion of YBX termination markers, etc., to ensure that the first and second cross-space call handling codes corresponding to the first and second intermediate representations are semantically correct under their respective target architectures and conform to the corresponding hardware specifications.

[0128] Step 130: Based on each cross-space call relationship, compile and rewrite the first code and the second code corresponding to each cross-space call relationship respectively, generate the first cross-space call processing code in the corresponding first code, and generate the second cross-space call processing code in the corresponding second code, so as to obtain the first intermediate representation corresponding to the first kernel function and the second intermediate representation corresponding to the second kernel function.

[0129] In one embodiment, the first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship are used to pass call parameters and perform execution synchronization when the corresponding first kernel function and second kernel function are executed by the first processing unit and the second processing unit, respectively.

[0130] In one embodiment, reference is made to Figure 2 , Figure 2 This application provides a flowchart of compiling and rewriting the first and second codes to generate cross-level parameter passing logic based on shared memory, specifically including the following steps:

[0131] Step 210: Determine the parameter layout of the call parameters involved in each cross-space call relationship in the corresponding parameter buffer.

[0132] In one embodiment, each parameter buffer is located in a shared memory area, which can be accessed by both the first and second processing units. (See also...) Figure 3 , Figure 3 This is a flowchart illustrating the parameter layout in the corresponding parameter buffer for each cross-space call relationship, provided in an embodiment of this application. The flowchart specifically includes the following steps:

[0133] Step 310: Identify dynamic call parameters in the call parameters.

[0134] Step 320: Determine the parameter byte length of each dynamic call parameter based on the data type of each dynamic call parameter.

[0135] In one embodiment, during the compilation phase, when the compiler locates a cross-hardware-level call instruction from the first processing unit to the second processing unit, it classifies the incoming call parameters, identifies non-constant parameters whose specific values ​​can only be determined at runtime as dynamic call parameters, and treats other constant parameters as compile-time constant parameters.

[0136] Step 330: Determine the aligned length of each dynamic call parameter based on the parameter byte length of each dynamic call parameter and the preset alignment rules.

[0137] In one embodiment, since the underlying hardware corresponding to the first processing unit may run on a 32-bit reduced instruction set, the underlying memory access interface is based on 4-byte units. The compiler iterates through all dynamic call parameters included in this cross-space call interaction, first determining the parameter byte length according to the data type, and then calculating the actual aligned length occupied by each parameter according to a preset alignment rule, such as a 4-byte alignment rule, thereby avoiding hardware errors or performance degradation caused by unaligned memory access.

[0138] Step 340: Determine the storage offset of each dynamic call parameter in the parameter buffer based on the arrangement order and aligned length of each dynamic call parameter.

[0139] Step 350: Determine the size of the parameter buffer based on the storage offset and aligned length of each dynamically invoked parameter.

[0140] In one embodiment, the storage offset of each dynamically invoked parameter within the parameter buffer is calculated based on the original input order and the aligned length of the dynamically invoked parameters, serving as the parameter layout. After calculating the aligned length of all dynamically invoked parameters, the compiler sums them sequentially to determine the size of the parameter buffer that must be allocated for this cross-level pass.

[0141] Step 220: Based on the layout of each parameter, generate parameter writing code in the corresponding first code and parameter retrieval code in the corresponding second code.

[0142] In one embodiment, the parameter writing code is used to write the call parameters involved in the corresponding cross-space call relationship into the corresponding parameter buffer according to the corresponding parameter layout, and the parameter retrieval code is used to retrieve the call parameters from the corresponding parameter buffer according to the corresponding parameter layout. The parameter writing code includes buffer allocation code and parameter storage code, and the parameter retrieval code includes buffer address retrieval code and parameter reading code. (Refer to...) Figure 4 , Figure 4 This application provides a flowchart illustrating how, based on the layout of each parameter, parameter writing code is generated in the corresponding first code, and parameter acquisition code is generated in the corresponding second code. The flowchart specifically includes the following steps:

[0143] Step 410: Generate buffer allocation code in the corresponding first code according to the size of the parameter buffer.

[0144] In one embodiment, the buffer allocation code is used to allocate a parameter buffer in a shared memory area. Specifically, the code can allocate a contiguous parameter buffer in the shared memory area at once, based on the calculated size of the parameter buffer.

[0145] Step 420: Generate parameter storage code in the corresponding first code according to the storage offset of each dynamic call parameter.

[0146] In one embodiment, the parameter storage code is used to write each dynamic call parameter to its corresponding storage location in the parameter buffer. The compiler uses the parameter storage code to generate serialization write instructions, which package the values ​​of all dynamic call parameters one by one according to their calculated storage offsets and write them to the corresponding storage locations in the parameter buffer.

[0147] Step 430: Generate the buffer address retrieval code in the corresponding second code.

[0148] In one embodiment, the buffer address acquisition code in the second code is used to obtain the starting address of the parameter buffer in the shared memory area. The second processing unit uses the buffer address acquisition code to obtain the global variable address of the parameter buffer, which is the starting address, to ensure that the control and calculation processing units can locate the same shared memory area through global symbols.

[0149] Step 440: Generate parameter reading code in the corresponding second code based on the starting address, the storage offset and data type of each dynamic call parameter.

[0150] In one embodiment, parameter reading code is used to determine the storage address of each dynamically invoked parameter and read the corresponding parameter value from each storage address. This parameter reading code uses consistent parameter sorting and alignment rules to locate the underlying storage address of each dynamically invoked parameter by adding the corresponding storage offset to the starting address. Since the hardware address spaces of the first processing unit and the second processing unit are not interconnected, the compiler automatically injects pointer type conversion instructions to forcibly convert the exposed physical storage address into a shared storage pointer with the corresponding data type. Finally, the second processing unit executes data reading instructions based on the converted shared storage pointer, extracts the corresponding parameter value from each storage address in the shared storage area, and restores it to the original variable for direct use by the computational logic.

[0151] In one embodiment, the compilation and rewriting of the first and second code corresponding to each cross-space call relationship further includes: identifying compile-time constant parameters in the call parameters, and generating constant assignment code in the corresponding second code based on the parameter values ​​of the compile-time constant parameters. The compile-time constant parameters are not written to the parameter buffer. In other words, for the identified compile-time constant parameters related to the cross-space call relationship, the compile-time constant parameters are not written to the parameter buffer; instead, the corresponding assignment statements are directly injected into the second code, avoiding participation in the shared memory-based transmission network and saving memory resources and transport overhead.

[0152] In one embodiment, each first cross-space call processing code includes call initiation code and synchronization waiting code. Each call initiation code is used to cause the first processing unit to initiate a corresponding call request. Each call request is used to cause the second processing unit to execute a corresponding second kernel function. Each synchronization waiting code is used to cause the first processing unit to wait for the corresponding second kernel function to complete execution after initiating the corresponding call request.

[0153] In one embodiment, reference is made to Figure 5 , Figure 5 This is a flowchart of generating first cross-space call processing code in the corresponding first code, provided in an embodiment of this application, specifically including the following steps:

[0154] Step 510: Identify at least one entry function that has a first execution space attribute.

[0155] In one embodiment, multiple cross-space call relationships within the same entry function use the same synchronization object. Each synchronization wait code is also used to reset the synchronization state of the synchronization object after the corresponding second kernel function completes execution, so that the synchronization object can be reused by the subsequent cross-space call relationship within the entry function. During the cross-level call rewriting stage, the compiler scans the parsed intermediate representation to identify entry functions with first execution space attributes, such as orchestrator-level entry functions.

[0156] Step 520: Assign synchronization objects to each entry function.

[0157] In one embodiment, after locating the entry point, the compiler injects relevant code statements at the very beginning of the entry function to allocate a globally shared synchronization object, such as a globally shared barrier address. Specifically, the hardware barrier synchronization resource is obtained by injecting a preset low-level allocation interface function, such as the alloc_boaddr function. In this way, all cross-level execution unit calls within the same entry function can reuse this barrier address as the synchronization object, eliminating the need to request and allocate a completely new barrier each time a new call is initiated, thereby saving internal hardware synchronization resources of the chip.

[0158] Step 530: Generate the call initiation code in the corresponding first code according to the cross-space call relationship.

[0159] In one embodiment, each initiating code is used to generate a call request carrying the identification information of the corresponding synchronization object. For dynamic call parameters related to cross-space call relationships, the compiler generates parameter passing logic in the first code. First, the number of bytes of the parameters is calculated. Each dynamic call parameter is traversed, the sizeof function is called, the corresponding dynamic call parameter is passed in, and the size is calculated based on four-byte alignment. The aligned byte sizes of all dynamic call parameters are summed to obtain the total number of bytes required to allocate the parameter buffer. Then, a shared memory buffer is allocated. The alloc_smem function is called, passing in the calculated total number of bytes, and a contiguous parameter buffer is allocated in the shared memory at once. After the parameter buffer is allocated, the first processing unit begins to store each parameter into the shared memory. The store primitive is called, passing in each dynamic call parameter itself and its corresponding offset address in the parameter buffer, and each dynamic call parameter is written into the parameter buffer one by one according to the corresponding offset address to complete the data serialization and storage. Finally, based on the cross-space call relationships, the orchestrator generates call initiation code in the corresponding first code, initiates a call to the execution unit, calls the call_eu_func function and passes in the zero-level parameter 0, the barrier address as the identifier of the synchronization object, and the control information of the specific scenario such as the identifier of the execution unit to be called, and issues the computation task with the call request to the hardware of the execution unit.

[0160] Step 540: Determine the calling nodes of each cross-space call relationship in the corresponding first code, and insert the synchronization wait code based on the corresponding synchronization object respectively.

[0161] In one embodiment, the compiler determines the call node of each cross-space call relationship in the corresponding first code within the code's instruction flow, that is, locates the execution unit call instruction node responsible for issuing tasks to the computing unit. After each such call initiation instruction, a synchronization wait code, i.e., a barrier wait instruction, bound to the inserted and allocated synchronization object (i.e., the shared barrier address) is automatically appended. After the call is initiated, the orchestrator performs the operation of waiting for the execution unit to complete, invoking the `bowait` wait instruction, passing in the workgroup-level parameters and the corresponding barrier address, causing the orchestrator to enter a blocking wait state, ensuring the correct timing of heterogeneous hardware interaction.

[0162] When the orchestrator's underlying hardware executes this barrier wait instruction, it forcibly suspends its subsequent control instruction flow and enters a blocked waiting state. This state continues until the corresponding execution unit completes all assigned computational tasks and actively triggers the barrier to signal completion. Only then does the orchestrator unblock and continue executing subsequent operations. It's understandable that when the compiler automatically injects barrier wait instructions, it flexibly selects the synchronization range based on the parallel requirements of the specific algorithm operators. This synchronization range can be at different granularities, such as workgroup level, cooperative workgroup level, or sub-workgroup level.

[0163] In one embodiment, due to the use of a shared barrier address mechanism, when the orchestrator executes multiple cross-level calls in the order of the code, once the previous barrier wait instruction is completed, the internal hardware state of the barrier as a synchronization state will be automatically reset to zero, so that the next call can continue to reuse the synchronization object for synchronization.

[0164] In one embodiment, reference is made to Figure 6 , Figure 6 This application provides a flowchart of generating second cross-space call processing code in the corresponding second code to obtain a second intermediate representation corresponding to the second kernel function, specifically including the following steps:

[0165] Step 610: Identify multiple second kernel functions that have second execution space attributes and do not have mutual calling relationships.

[0166] In one embodiment, during the compilation and rewriting phase, the compiler comprehensively scans and identifies all second kernel functions defined in the second code that have second execution space attributes and do not have mutual calling relationships. These second kernel functions are typically user-written functional modules responsible for executing specific vectorized multiplication-addition or reduction algorithms. Simultaneously, the compiler also identifies compile-time constant parameters (constexpr parameters) related to cross-space calling relationships. These parameters are not passed through shared memory; the compiler directly injects the corresponding assignment statements into the second code, thereby saving shared memory space and overhead.

[0167] Step 620: Determine the function identifiers of the multiple second kernel functions respectively.

[0168] In one embodiment, after identifying these second kernel functions, the compiler determines the function identifiers of multiple second kernel functions in the underlying compilation process and automatically assigns a corresponding function identifier to each computation function. The function identifier can serve as a digital identity code for each second kernel function, used to point to the computation task when communicating across heterogeneous hardware.

[0169] Step 630: Generate a unified entry function for multiple second kernel functions.

[0170] In one embodiment, the compiler automatically generates a unified entry dispatch function for multiple second kernel functions at the very beginning of the second code. This unified entry function serves as the sole entry point for the execution unit hardware level to receive orchestrator scheduling instructions. When the orchestrator initiates a cross-level call and wakes up the execution unit via hardware instructions, the control flow of the execution unit hardware first enters this unified entry function.

[0171] Step 640: Generate identifier acquisition code and call dispatch code in the unified entry function.

[0172] In one embodiment, the identifier acquisition code is used to acquire the function identifier carried in the call request, the call dispatch code is used to dispatch the execution flow to the corresponding second kernel function according to the function identifier, and each first cross-space call processing code is also used to generate a call request carrying the function identifier of the second kernel function to be called.

[0173] Within the execution unit, the logic triggers the function identifier retrieval operation via identifier acquisition code. For example, it calls the underlying interface function for retrieving the execution unit's function identifier to extract the target function identifier carried in the current orchestrator scheduling request. Then, by calling the dispatch code, based on the received function identifier and combined with the internally generated jump mapping table or branch judgment logic, the execution flow is dispatched to the corresponding second kernel function, i.e., the specific computation function. At this point, the control flow enters the user-written specific computation logic code to perform parameter deserialization and vectorization operations.

[0174] In one embodiment, during parameter deserialization, the execution unit first obtains the address of the global variable in the shared memory by calling the gbuf_access function to obtain the address of the global variable in the parameter buffer allocated in the shared memory. Since different levels of processing units reference the address of the same physical shared memory through global symbols, the execution unit can calculate the internal offset address of each dynamically invoked parameter in the parameter buffer according to the sorting and alignment rules that are completely consistent with those on the orchestrator side.

[0175] After obtaining the internal offset address and combining it with the base address of the global variable to derive the physical address of the corresponding dynamically called parameter, the execution unit performs pointer type conversion. It calls the `cast` instruction to force a type conversion, transforming the raw physical address belonging to the shared memory space into a typed pointer based on the specific data type information. Finally, the execution unit reads the restored parameters, calls the `load` function with the typed pointer obtained above, extracts the underlying data from the shared memory, deserializes it, and restores it to the corresponding restored parameters usable by the computational logic.

[0176] This embodiment reduces the workload of users manually writing complex and error-prone request dispatch and task jump code when programming in heterogeneous systems through this automated design. It allows users to describe business logic as if calling a normal function when using this programming model, thereby reducing the programming threshold and development complexity under the two-level hardware architecture of computing chips.

[0177] In one embodiment, reference is made to Figure 7 , Figure 7 This is a flowchart of the process of compiling and rewriting the first code corresponding to each cross-space call relationship, provided in an embodiment of this application. The process specifically includes the following steps:

[0178] Step 710: Identify the data transfer operation in the first code.

[0179] In one embodiment, when writing control code at the orchestrator level, the user explicitly calls data transfer primitives from global storage to shared storage. During compilation, the compiler recognizes the data transfer operations in the first code. Here, global storage is a large-capacity memory that can store slower raw weights and input / output data, while shared storage is a faster interaction medium with hardware address space numbers for the orchestrator to address.

[0180] Step 720: Generate data transfer code based on the data transfer operation.

[0181] In one embodiment, data movement code is used to move relevant computational data from the global storage area to the shared storage area. The first code includes data movement primitives, such as `alloc_smem`, `copy_g2s`, or `copy_s2g`, which are used to generate data movement code during compilation based on the data movement operations. This code then moves the computational data from the global storage area to the target area in the shared storage area based on these primitives. It is understood that the specific primitive content and nouns can be set according to the actual situation.

[0182] Step 730: Generate data layout transformation code in the data transfer code based on the source data layout and target data layout of the calculated data.

[0183] In one embodiment, the source data layout is a row-major layout, and the target data layout is a row-block layout; or, the source data layout is a column-major layout, and the target data layout is a column-block layout. The data layout conversion code is used to convert the computational data from the source data layout to the target data layout when moving computational data. The data moving code is also used to move computational data shared by multiple second kernel functions to a shared storage area before the execution of multiple second kernel functions, so that the multiple second kernel functions can reuse it.

[0184] Specifically, during the data transfer process, the data layout conversion code transforms the source data layout of the computation data from row-first or column-first to the target data layout, i.e., row-block layout or column-block layout. The entire process requires no additional conversion steps, thus improving data transfer efficiency.

[0185] In one embodiment, since ordinary row-major or column-major storage formats cannot be directly adapted to the efficient vectorized computation of computing chips, this embodiment defines two fixed-size 256-byte block layout formats in the computing chip. The underlying transfer engine will automatically complete the conversion of these layout formats while executing transfer instructions. It is understood that the 256-byte size can be set according to actual needs.

[0186] For example, for row-based block layouts, the block height formula is set to th = 4 rows, and the block width formula is set to tw = 512 divided by nbits columns, where th represents the number of rows in the block, tw represents the number of columns in the block, and nbits represents the bit width precision of a single calculated data element. For example, for 16-bit precision data, the calculated block width is 32 columns, and each row-based block layout is 4 rows and 32 columns, containing 128 data elements and filling a total capacity of 256 bytes. For column-based block layouts, the block width formula is set to tw = 4 columns, and the block height formula is set to th = 512 divided by nbits rows. Similarly, when processing 16-bit precision data, the column-based block layout is automatically divided into 32 rows and 4 columns.

[0187] Once the computational data is preloaded in a chunked layout, multiple execution units can share the computational data within the target area, reducing redundant data transfer overhead. At this point, the orchestrator moves a large amount of computational data, such as weight arrays and input features, from the underlying global storage to high-speed shared storage in the early stages of the computation process. Subsequently, when the orchestrator schedules multiple execution units in parallel for different data slices, these execution units can directly reuse the same computational data from the corresponding target area in the shared storage for reading and processing, without needing to initiate new data retrieval requests from the global storage.

[0188] As can be seen from the above, this embodiment uses the compiler to take over and automatically generate the underlying hardware data interaction logic. Whether it is parameter buffer partitioning and address alignment, serialization writing and deserialization reading across memory, or automatic injection of masked states, these processes do not require manual management by the user. The user only needs to write basic function calls at the high-level language level, thereby reducing the programming threshold and later maintenance cost of computing chips, while avoiding the risk of out-of-bounds crashes caused by manually managing memory offset addresses.

[0189] In one embodiment, reference is made to Figure 8 , Figure 8 This is a flowchart provided in an embodiment of the present application for compiling and rewriting the first code and the second code to generate cross-level barrier synchronization logic, specifically including the following steps:

[0190] Step 810: Identify the orchestrator-level entry functions in the first code and assign barrier addresses to each orchestrator-level entry function.

[0191] In one embodiment, during the cross-level call rewriting phase, the compiler scans the parsed intermediate representation to identify the orchestrator-level entry function, which acts as the control center. Once the entry point is located, the compiler automatically injects relevant code statements at the very beginning of this entry function to allocate a globally shared barrier address, for example, `_shared_eu_boaddr_=ezro.alloc_boaddr()`. By injecting a pre-defined underlying allocation interface function, the compiler obtains the hardware barrier synchronization resource. This allows all cross-level execution unit calls within the same orchestrator-level entry function to reuse this barrier address, eliminating the need to allocate a completely new barrier each time a new call is initiated, thus saving hardware synchronization resources within the underlying computing chip.

[0192] Step 820: Locate the execution unit call instruction node corresponding to each cross-level call, insert the corresponding barrier address barrier waiting instruction after each execution unit call instruction node, and complete the execution synchronization between the orchestrator and the execution unit through the barrier address.

[0193] In one embodiment, the compiler locates the execution unit invocation instruction node responsible for issuing tasks to the computation unit in the code instruction flow. After each such invocation instruction, a barrier wait instruction bound to the allocated shared barrier address is automatically appended, such as bowait(WG,boaddr). When the orchestrator's underlying hardware executes this barrier wait instruction, it forcibly suspends its subsequent control instruction flow and enters a blocked waiting state until the corresponding execution unit completes all allocated computation tasks and actively triggers the barrier to notify completion. Only then does the orchestrator unblock and continue executing subsequent operations.

[0194] In one embodiment, when automatically injecting barrier wait instructions, the compiler can flexibly select the synchronization range according to the parallel requirements of specific algorithm operators. This synchronization range can be at different granularities such as workgroup level, cooperative workgroup level, and sub-workgroup level. Due to the mechanism of using a shared barrier address, when the orchestrator executes multiple cross-level calls sequentially according to the code order, once the previous barrier wait instruction is completed, the internal hardware state of the barrier is automatically reset to zero, allowing subsequent calls to reuse the barrier for synchronization.

[0195] In one embodiment, the passing of call parameters involved in each cross-space call relationship needs to be serialized and deserialized across shared memory regions. (See also...) Figure 9 , Figure 9 This is a schematic diagram of the serialization and deserialization process provided in the embodiments of this application.

[0196] First, during compilation and parsing, the compiler identifies the compile-time constant parameter, namely the `constexpr` parameter, in the call arguments. Since this compile-time constant parameter is not written to the parameter buffer and is not passed through the shared memory area, the compiler directly generates constant assignment code in the corresponding second code based on the parameter value of the compile-time constant parameter, thereby saving shared memory area space and handling overhead.

[0197] Next, for dynamic call parameters related to cross-space call relationships, the compiler generates parameter writing code in the corresponding first code to handle parameter passing. First, based on the data type of each dynamic call parameter, the parameter length in bytes is determined. Specifically, each dynamic call parameter is iterated over, the `sizeof` function is called, the corresponding dynamic call parameter is passed in, and the aligned length of each dynamic call parameter is determined according to a preset alignment rule. The aligned lengths of all dynamic call parameters are accumulated, and the size of the parameter buffer, i.e., the total number of bytes, is determined based on the storage offset and aligned length of each dynamic call parameter. Then, based on the size of the parameter buffer, buffer allocation code is generated in the corresponding first code. For example, the `alloc_smem` function is called, passing in the calculated parameter buffer size, to allocate a contiguous parameter buffer in the shared memory area at once. After the parameter buffer allocation is complete, the first processing unit begins to store the parameters one by one into the shared memory area. Based on the storage offset of each dynamic call parameter, parameter storage code is generated in the corresponding first code. For example, the `store` storage primitive is called, passing in each dynamic call parameter itself and its corresponding storage offset in the parameter buffer, to write each dynamic call parameter to its corresponding storage location in the parameter buffer, completing the data serialization and storage.

[0198] Next, the first cross-space call processing code is generated in the corresponding first code, including call initiation code and synchronization wait code. The first processing unit initiates the corresponding call request through the call initiation code, for example, calling the `call_func` function and passing in the zero-level parameter 0, the barrier address as the synchronization object, and control information specific to the scenario, such as the function identifier of the second kernel function to be called. This call request is used to cause the second processing unit to execute the corresponding second kernel function, for example, to issue a computation task. After the call request is initiated, the call nodes of each cross-space call relationship in the corresponding first code are determined, and synchronization wait code based on the corresponding synchronization object is inserted respectively, for example, calling the `bowait` wait instruction, passing in the workgroup level parameter and the corresponding barrier address, so that the first processing unit enters a blocking wait state after initiating the corresponding call request, waiting for the corresponding second kernel function to complete execution, ensuring the correct timing of heterogeneous hardware interaction.

[0199] Then, a second cross-space call processing code, including parameter retrieval code, is generated in the corresponding second code. The second processing unit first obtains the starting address of the parameter buffer in the shared memory area, i.e., the address of the global variable, by using the buffer address retrieval code generated in the corresponding second code, for example, by calling the gbuf_access function. Since the first and second processing units reference the starting address of the same physical shared memory area through global symbols, the second processing unit can calculate the storage offset of each dynamic call parameter in the parameter buffer, i.e., the internal offset address, according to the arrangement order and alignment rules that are completely consistent with those on the first processing unit side.

[0200] After obtaining the storage offset and combining it with the starting address to derive the storage address of the corresponding dynamically called parameter, the second processing unit performs pointer type conversion, for example, by calling a cast instruction to convert the physical storage address belonging to the shared memory area into a typed pointer of type ptr_t, combined with the data type information of each dynamically called parameter. Finally, based on the starting address, the storage offset and data type of each dynamically called parameter, parameter reading code is generated in the corresponding second code. The second processing unit uses the parameter reading code, for example, by calling the load function and passing in the typed pointer obtained above, to determine the storage address of each dynamically called parameter, read the corresponding parameter value from each storage address, and deserialize it to restore the corresponding parameter value usable by the computational logic.

[0201] In one embodiment, the first code includes data transport primitives, such as alloc_smem(), copy_g2s / copy_s2g, etc. During compilation, computational data in global storage is transported to the target area of ​​shared storage based on these primitives. During data transport, the data format of the computational data is converted from row-major or column-major to a row-block layout ROW<->TROW or a column-block layout COL<->TCOL. This entire process requires no additional conversion steps, improving data transport efficiency. Furthermore, multiple execution units can share computational data within the target area, reducing redundant transport overhead. It is understood that the specific primitive content and terminology can be set according to actual needs.

[0202] When writing control code at the orchestrator level, users explicitly call data transfer primitives from global storage to shared storage. During compilation, the compiler uses these primitives to move computational data from global storage to the target area in shared storage. Global storage here refers to large-capacity memory, suitable for storing slower raw weights and input / output data, while shared storage is a faster interaction medium with hardware address space numbers for precise addressing by the orchestrator.

[0203] During data transfer, the system converts the data format of the computational data from row-major or column-major to either row-block or column-block layout. Since ordinary row-major or column-major storage formats cannot directly adapt to the efficient vectorized computation of the computing chip, this embodiment defines two fixed-size 256-byte block layout formats in the computing chip. The underlying transfer engine automatically completes the conversion of these layout formats while executing transfer instructions. It can be understood that the 245-byte size can be set according to actual needs.

[0204] For example, for row-based block layouts, the block height formula is set to th = 4 rows, and the block width formula is set to tw = 512 / nbits columns, where th represents the number of rows in the block, tw represents the number of columns in the block, and nbits represents the bit width precision of a single calculated data point. For example, for 16-bit precision data, the calculated block width is 32 columns, and each row-based block layout consists of 4 rows. The 32 columns contain exactly 128 data elements, filling a total capacity of 256 bytes. For column-based block layouts, the block width formula is set to tw = 4 columns, and the block height formula is set to th = 512 / nbits rows. When processing 16-bit precision data, the column-based block layout is automatically divided into 32 rows. 4 columns.

[0205] Once the computational data is preloaded in a chunked layout, multiple execution units can share the computational data within the target region. At this point, the orchestrator moves a large amount of computational data, such as weight arrays and input features, from the underlying global storage to high-speed shared storage in the early stages of the computation process. Subsequently, when the orchestrator schedules multiple execution units in parallel for different data slices, these execution units can directly reuse the same computational data from the corresponding target region in the shared storage for reading and processing, without needing to initiate new data retrieval requests from the global storage.

[0206] In one embodiment, compiling and rewriting the first code and the second code further includes: generating a unified entry point distribution function in the second code, identifying the computation functions in the second code, assigning a corresponding function identifier to each computation function, and enabling the unified entry point distribution function to distribute the execution flow to the corresponding computation function according to the received function identifier.

[0207] Specifically, during the compilation and rewriting phase, the compiler comprehensively scans and identifies all computation functions defined in the second-level code. These computation functions are typically user-written functional modules responsible for executing specific vectorized multiplication-addition or reduction algorithms. After identifying these computation functions, the compiler automatically assigns a corresponding function identifier to each computation function in the underlying compilation process. The function identifier is equivalent to an internal digital identity code for each computation function, used to point to the computation task when communicating across heterogeneous hardware.

[0208] The compiler then automatically generates a unified entry point dispatch function at the very beginning of the second code. This unified entry point dispatch function serves as the sole entry point for the execution unit hardware level to receive orchestrator scheduling instructions. When the orchestrator initiates a cross-level call and wakes up the execution unit via hardware instructions, the control flow of the execution unit hardware first enters this unified entry point dispatch function. Subsequently, the logic within the execution unit triggers the operation of obtaining the function identifier, for example, by calling the underlying interface function for obtaining the execution unit function identifier to extract the target function identifier carried in the current orchestrator scheduling request. Then, the unified entry point dispatch function, based on the received function identifier and combined with the internally generated jump mapping table or branch judgment logic, dispatches the execution flow to the corresponding computation function. At this point, the control flow enters the user-written specific computation logic code to perform subsequent parameter deserialization and vectorization operations.

[0209] This embodiment reduces the workload of users manually writing complex and error-prone request dispatch and task jump code when programming in heterogeneous systems through this automated design. It allows users to describe business logic as if calling a normal function when using this programming model, thereby reducing the programming threshold and development complexity under the two-level hardware architecture of computing chips.

[0210] As can be seen from the above, this embodiment uses the compiler to take over and automatically generate the underlying hardware data interaction logic. Whether it is parameter buffer partitioning and address alignment, serialization writing and deserialization reading across memory, or automatic injection of masked states, these processes do not require manual management by the user. The user only needs to write basic function calls at the high-level language level, thereby reducing the programming threshold and later maintenance cost of computing chips, while avoiding the risk of out-of-bounds crashes caused by manually managing memory offset addresses.

[0211] After obtaining the corresponding first cross-space call processing code and second cross-space call processing code through the above process, the code information corresponding to the first intermediate representation of the first kernel function and the second intermediate representation of the second kernel function is updated accordingly, and then the target code is compiled based on the intermediate representation.

[0212] Step 140: Based on the first execution space attribute and the second execution space attribute, compile the first intermediate representation and the second intermediate representation into first target code adapted to the first processing unit and second target code adapted to the second processing unit, respectively, and generate a unified executable file based on the first target code and the second target code.

[0213] In one embodiment, after rewriting and normalizing the underlying cross-level logic of the intermediate representation, the compiler enters the code generation stage. Based on the first and second execution space attributes, it compiles the first and second intermediate representations into first target code adapted to the first processing unit and second target code adapted to the second processing unit, respectively. Specifically, the compiler compiles the first intermediate representation with the first execution space attribute into first target code adapted to a 32-bit RISC architecture, and the second intermediate representation with the second execution space attribute into second target code adapted to a 64-bit RISC architecture with vector extensions. After generating these two completely independent sets of underlying machine code, the compiler's linker component performs memory address relocation and symbol resolution integration on these two types of target code, and generates a unified executable file based on the first and second target codes. This unified executable file internally encapsulates control instruction fragments for the first processing unit to read, and vector calculation instruction fragments for the second processing unit to read.

[0214] In one embodiment, after the executable file is available, the runtime phase begins. (See also...) Figure 10 , Figure 10 This is a flowchart of a kernel running method provided in an embodiment of this application, which specifically includes the following steps:

[0215] Step 1010: Load the unified executable file.

[0216] In one embodiment, the unified executable file includes first target code adapted to a first processing unit and second target code adapted to a second processing unit. The first target code includes first cross-space call processing code corresponding to at least one cross-space call relationship, and the second target code includes second cross-space call processing code corresponding to each cross-space call relationship.

[0217] Step 1020: Execute the first target code using the first processing unit. When the execution reaches the call node corresponding to any cross-space call relationship in the first target code, execute the corresponding first cross-space call processing code, write the call parameters involved in the cross-space call relationship into a shared storage area accessible to both the first and second processing units, initiate a call request, and wait for the corresponding second kernel function to complete execution.

[0218] In one embodiment, the first processing unit runs the first target code, for example, by running the corresponding 32-bit reduced instruction set code using orchestrator hardware, to complete the memory management of shared memory and initiate low-level data transfer operations. This involves moving computational data stored in the slower, large-capacity global storage to the faster shared storage for preloading. After data preparation, when running to any call node corresponding to a cross-space call relationship in the first target code, the first processing unit executes the corresponding first cross-space call processing code. It writes the call parameters involved in the cross-space call relationship into a shared storage area accessible to both the first and second processing units, initiates a call request through corresponding hardware scheduling instructions, and waits for the corresponding second kernel function to complete. For example, it uses a barrier waiting instruction to put itself into a blocked state, thus waiting for the computational task currently assigned to the second processing unit to complete.

[0219] Step 1030: Utilize the second processing unit to run the second target code, respond to the call request, execute the corresponding second cross-space call processing code, obtain the call parameters from the shared memory area, and execute the second kernel function determined by the call request.

[0220] In one embodiment, a second processing unit, i.e., an execution unit hardware, runs second target code, such as corresponding 64-bit execution unit target code. Specifically, in response to a call request, which may be a scheduling instruction issued by the first processing unit, the second processing unit, upon being awakened, executes the corresponding second cross-space call processing code. Through a cross-hardware access mechanism, it retrieves call parameters from the shared memory area. For example, it reads previously stored cross-level dynamic call parameters and computational data, and executes the second kernel function determined by the call request.

[0221] During the execution of the second kernel function, the second processing unit performs corresponding vectorized computation operations and writes the computation results back to the designated cache area of ​​the shared memory region at the end of each computation. After all write operations are completed, the second processing unit triggers a barrier completion signal through the underlying hardware signal path to asynchronously notify the first processing unit, which is in a blocked state, that the task has been completed.

[0222] Step 1040: After the second kernel function is executed, the first processing unit ends its waiting and continues to run the first target code.

[0223] In one embodiment, while continuously monitoring the underlying hardware status, once the first processing unit detects a barrier completion signal, confirming that the issued computation task has been completed (i.e., after the second kernel function has finished executing), the first processing unit ends its waiting and continues running the first target code. At this time, the first processing unit can extract the completed computation result from the shared memory area according to the control logic, and continue to execute subsequent operations of the first target code based on the computation result, such as writing the final result back to the global memory area, or immediately starting a new computation scheduling task for the next data slice.

[0224] In one embodiment, a kernel programming system is provided according to the above compilation process, comprising:

[0225] The programming interface module provides an execution space declaration interface and sets execution space declaration information for kernel functions in the obtained kernel code through the execution space declaration interface.

[0226] The compilation module is used to determine the execution space attributes of each kernel function based on the execution space declaration information, identify at least one cross-space call relationship based on the function call relationship in the kernel code, generate the first cross-space call processing code and the second cross-space call processing code based on each cross-space call relationship, and compile the compiled and rewritten kernel code into target code adapted to the first processing unit and the second processing unit, respectively.

[0227] The runtime module is used to load the unified executable file generated based on each target code, enabling the first processing unit and the second processing unit to complete the passing of call parameters and execution synchronization through the first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship.

[0228] The following describes the specific programming system framework corresponding to the computing chip. In one embodiment, refer to... Figure 11 , Figure 11 This is a schematic diagram of the architecture of the kernel programming system provided in this application embodiment. The entire kernel programming system is divided into three processing layers from top to bottom: programming interface module, compilation module, and runtime module.

[0229] The programming interface module provides an execution space declaration interface, which is used to set execution space declaration information for kernel functions in the obtained kernel code. This module includes an orchestrator-level kernel declaration module and an execution unit-level kernel declaration module.

[0230] The orchestrator-level kernel declaration module primarily provides an execution space declaration interface, which is used to set execution space declaration information for kernel functions in the acquired kernel code. Specifically, this module includes an orchestrator-level kernel declaration module and an execution unit-level kernel declaration module. The orchestrator-level kernel declaration module allows users to declare orchestrator-level kernel functions through the execution space declaration interface, explicitly marking that the current control function will ultimately be compiled into a 32-bit RISC target architecture, undertaking core control responsibilities such as shared memory allocation, global data movement, and execution unit scheduling. Simultaneously, the execution unit-level kernel declaration module allows users to declare execution unit-level kernel functions through the execution space declaration interface, marking that the current computation function will be compiled into a 64-bit RISC target architecture with vector extensions, undertaking complex underlying vectorized computations.

[0231] When the program enters the compilation module, it determines the execution space attributes of each kernel function based on the execution space declaration information, identifies at least one cross-space call relationship based on the function call relationship in the kernel code, generates the first cross-space call processing code and the second cross-space call processing code based on each cross-space call relationship, and compiles the rewritten kernel code into target code adapted to the first processing unit and the second processing unit, respectively.

[0232] The compilation module performs abstract syntax tree (API) parsing and intermediate representation generation. It reads the user's original code written in a high-level language, parses it into an API, and then transforms it into a custom structured intermediate representation. In this conversion and parsing phase, the compilation module determines the execution space attributes of each function based on the execution space declaration information, labels different intermediate representation nodes with orchestrator space or execution unit space, and constructs a complete function call graph. This allows it to distinguish between ordinary function calls within the same space and identify at least one cross-space call relationship spanning different hardware physical spaces.

[0233] After generating the intermediate representation, the compilation module performs two-level intermediate representation rewriting and normalization operations, performing standardized compilation optimization and transformation processes on the intermediate representations of both the orchestrator space and the execution unit space. These compilation optimization and transformation processes include expanding loop structures in the code, inferring and verifying the data types of unknown variables, and inserting specific termination markers at the end of instructions, maintaining the semantic correctness of the code and conforming to the underlying hard specifications of the target chip's instruction set.

[0234] Next, the compilation module performs cross-level call rewriting operations, generating first and second cross-space call handling code based on the various cross-space call relationships. When an orchestrator-level function initiates a cross-level cross-space call relationship to an execution unit-level function, the compilation module automatically completes the hardware interaction conversion logic in the underlying code. When generating the first cross-space call handling code at the first code stage, the compilation module automatically injects instructions to allocate a shared memory buffer at the call point, serializes and packages the dynamic call parameters required for each call, and stores them in the shared memory. Subsequently, it injects the underlying call instruction that initiates the execution unit's computation task, and injects a barrier wait instruction after the call instruction to allow the orchestrator to enter a safe blocking wait state. At the same time, the compilation module also automatically allocates a globally shared barrier address for the call flow specifically for cross-hardware level barrier synchronization. When generating the second cross-space call handling code at the corresponding second code stage, the compilation module automatically inserts parameter deserialization code logic at the beginning of its entry function, and automatically generates a unified entry function thereafter, distributing the execution flow to the corresponding specific computation function according to different underlying execution unit function identifiers. This allows the execution unit's underlying hardware to be woken up, and then first read and extract the corresponding parameters from the shared memory address referenced by the unified reference, restoring them to the original computational variables.

[0235] After the underlying cross-level logic rewriting is completed, the compilation module performs dual-module code generation, linking, and executable file generation operations. Dual-module code generation physically splits the processed intermediate representation, generating two independent underlying compiler intermediate representation modules, which are then compiled into target code adapted to the first and second processing units. Specifically, the intermediate representation belonging to the orchestrator space is directed to generate orchestrator target code adapted to the 32-bit RISC, while the intermediate representation belonging to the execution unit space is generated to generate execution unit target code adapted to the 64-bit RISC with vector extension capabilities. After the two sets of target code are generated, the code linker performs address relocation and symbol integration on these two target codes adapted to different instruction set architectures, merging them into a unified executable file that can be directly burned or loaded.

[0236] Finally, the program enters the runtime module, loads the unified executable file generated based on each target code, and enables the first processing unit and the second processing unit to complete the passing of call parameters and execution synchronization through the first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship.

[0237] Specifically, the unified executable file executes the corresponding runtime interaction logic on the computing chip. The first processing unit, i.e., the orchestrator, starts first and performs memory allocation management of shared memory, transfer of large amounts of data in global storage, and hardware scheduling logic for issuing specific computing tasks to the execution units. The second processing unit, i.e., the execution unit, receives the scheduling instructions sent by the first processing unit, checks the barrier synchronization status, and then performs matrix or vectorized calculations.

[0238] During operation, the first processing unit and the second processing unit use the first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship to realize the passing of call parameters across hardware address spaces, as well as the pre-loading of computing data and multi-core reuse through shared memory, and finally complete the underlying execution synchronization.

[0239] In one embodiment, reference is made to Figure 12 , Figure 12 This is a schematic diagram of the kernel programming system provided in the embodiments of this application.

[0240] First, steps 1 and 2 involve the user writing kernel code containing at least one first kernel function and at least one second kernel function. In the first code, the user can explicitly call shared memory allocation functions, data movement primitives, and initiate cross-space calls to related second kernel functions in the target second code. Simultaneously, the second code internally uses vectorized types with different channel numbers to execute corresponding computational logic, including vector types with 16, 32, 64, 128, 256, and 512 channels. Furthermore, the second processing unit, i.e., the execution unit, is also equipped with a Vector Local Storage (DLM) at its underlying layer, with each computation bundle having a capacity of 192KB and an address space number of 24, for storing high-frequency temporary data.

[0241] Step 3 then involves compiling and parsing the kernel code and generating intermediate representations. The compiler reads the corresponding high-level language code and parses it into an abstract syntax tree, identifying the execution space declaration information bound to each kernel function. It maps the code to low-level intermediate representation nodes marked with either a first execution space attribute (orchestrator space) or a second execution space attribute (execution unit space). When a call is detected from a first kernel function with the first execution space attribute to a second kernel function with the second execution space attribute, the call point spanning heterogeneous hardware layers is identified and marked as the call node corresponding to the cross-space call relationship based on this function call relationship.

[0242] Step 4, based on the cross-space call relationship, rewrites the compilation to generate the first and second cross-space call handling codes, comprising three cooperating sub-steps 4a to 4c. In sub-step 4a, the compiler automatically allocates a globally shared synchronization object (barrier address) for the current entry function with the first execution space attribute, used for execution synchronization of all calls to the second processing unit within that function. In sub-step 4b, the compiler executes the following sequentially on the first code: generating buffer allocation code to allocate a parameter buffer in the shared memory area; calculating the aligned length of each dynamic call parameter according to a preset alignment rule (four-byte alignment rule) to determine the parameter layout; generating parameter storage code to write the parameters one by one to the corresponding storage offset position in the parameter buffer to complete data serialization; generating call initiation code to initiate the call request to the second processing unit; and inserting synchronization waiting code at the call node to put the first processing unit into a blocked waiting state. In sub-step 4c, the compiler generates a unified entry function EU_ENTRY in the second code. The call dispatch code in this function routes and dispatches the execution flow to the corresponding second kernel function (i.e., the specific business calculation function) based on the function identifier number obtained by the identifier acquisition code. Meanwhile, the compiler automatically injects parameter retrieval code at the entry point of each specific second kernel function. It calculates the storage address of the call parameters in the shared memory area, performs type casting operations with pointers, and reads the corresponding parameter values ​​from each storage address through the parameter reading code, and deserializes them to restore the original high-level language variables for use by the computational logic.

[0243] Step 5 involves rewriting and normalizing the intermediate representations obtained after obtaining the first and second intermediate representations. The compiler performs standardized compilation optimizations and logical normalization actions, including type inference and loop unrolling, for each intermediate representation node possessing its own first and second execution space attributes. It is understood that the order of steps 4 and 5 is not fixed.

[0244] Steps 6 and 7 involve compiling the first intermediate representation and the second intermediate representation into object code and generating a unified executable file based on the first and second execution space attributes. The compiler physically separates and isolates the normalized intermediate representations to obtain the first intermediate representation corresponding to the first kernel function and the second intermediate representation corresponding to the second kernel function. Then, the first intermediate representation is directed to generate first object code adapted to the 32-bit RISC architecture of the first processing unit, and the second intermediate representation is generated to generate second object code adapted to the 64-bit RISC architecture with vector extensions of the second processing unit. Finally, these two types of object code undergo memory address relocation and symbol integration packaging to generate a unified executable file, i.e., an ELF file, based on the first and second object code.

[0245] Step 8 involves loading and starting the kernel method. On the actual computing chip physical hardware, the system loads the previously generated unified executable file and starts it. The first processing unit runs the first target code, performing memory allocation management in the shared memory area, moving large amounts of data in the global memory area, and hardware scheduling logic that issues specific computation task call requests to the second processing unit. The second processing unit runs the second target code, responding to the call request, and after verifying the synchronization state maintained based on the synchronization object, executes the second kernel function determined by the call request to complete matrix or vectorized computation.

[0246] The following describes two specific application scenarios of the computing chip.

[0247] In one embodiment, the kernel compilation and execution process of the computing chip can be applied to the grouped RMS normalization and gating activation fusion computation scenario in a large language model. First, the first processing unit, i.e., the orchestrator, performs shared memory area allocation. During the process of moving computational data such as weights and biases from the global storage area to the shared storage area, the data format is synchronously converted from the source data layout to the target data layout, i.e., row-block layout (TROW), through data layout transformation code. Subsequently, the first processing unit initiates a call request to the second processing unit. During this process, the compiler automatically generates first and second cross-space call processing code to complete the passing of call parameters and execution synchronization. Upon responding to the call request, the second processing unit reads the data and executes the second kernel function determined by the call request to complete vectorization computations such as reduction. After completing the gating activation computation, the computation result is written back. This application scenario achieves preloading of computational data in the shared storage area, allowing multiple second kernel functions to share the computational data in this area, effectively reducing the number of accesses to the global storage area. Simultaneously, the row-block layout optimizes the vectorization reading efficiency of the second processing unit.

[0248] In one embodiment, the kernel compilation and execution process of the computing chip can be applied to the Softmax normalization scenario in the attention mechanism. First, the entry function with the first execution space attribute directly calls the corresponding second kernel function. Subsequently, the underlying second processing unit, i.e., the execution unit, uses dual-thread bundles (Warp) to perform vectorized computation in parallel. Each thread bundle is responsible for half of the matrix rows and uses a shared memory area as a physical transfer to complete data exchange such as the maximum value and summation value across thread bundles. Finally, the second processing unit writes the calculation results back to the shared memory area in the target data layout, i.e., a row-block layout. This embodiment improves the computational efficiency of in-row reduction by using dual-thread bundle parallel computation, and at the same time, it uses a shared memory area for data transfer and exchange, saving the overhead of additional calls to underlying synchronization primitives.

[0249] This application also provides a kernel compilation apparatus that can implement the above-described kernel compilation method, see reference. Figure 13 The device includes:

[0250] The code acquisition module 1310 is used to acquire kernel code, which includes first code of at least one first kernel function and second code of at least one second kernel function. Each first kernel function has a first execution space attribute, and each second kernel function has a second execution space attribute that is different from the first execution space attribute.

[0251] The call relationship identification module 1320 is used to compile and parse the kernel code. Based on the function call relationship in the kernel code, it identifies at least one cross-space call relationship. Each cross-space call relationship is formed by a first kernel function calling a second kernel function.

[0252] The code rewriting module 1330 is used to compile and rewrite the first code and the second code corresponding to each cross-space call relationship based on each cross-space call relationship. It generates the first cross-space call processing code in the corresponding first code and the second cross-space call processing code in the corresponding second code, thereby obtaining the first intermediate representation of the first kernel function and the second intermediate representation of the second kernel function. The first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship are used to pass call parameters and synchronize execution when the corresponding first kernel function and the second kernel function are executed by the first processing unit and the second processing unit, respectively.

[0253] The target code generation module 1340 is used to compile the first intermediate representation and the second intermediate representation into first target code adapted to the first processing unit and second target code adapted to the second processing unit, respectively, according to the first execution space attribute and the second execution space attribute, and to generate a unified executable file based on the first target code and the second target code.

[0254] The specific implementation of the kernel compilation device in this embodiment is basically the same as the specific implementation of the kernel compilation method described above, and will not be repeated here.

[0255] In this embodiment, kernel code containing a first kernel function and a second kernel function is obtained during the compilation stage. The two levels of processing units are distinguished using first and second execution space attributes. Cross-space call relationships are automatically identified based on function call relationships in the kernel code. The compiler automatically rewrites and generates cross-space call processing code based on shared memory, automatically completing parameter calls and execution synchronization across address spaces at the underlying level. This eliminates the need for users to manually allocate shared memory buffers and manage data transfer details, reducing programming complexity and the probability of human error. Simultaneously, based on the first and second execution space attributes, the first and second intermediate representations corresponding to the kernel functions are compiled into first target code adapted to the first processing unit and second target code adapted to the second processing unit, respectively. This adapts to scenarios where different processing units correspond to different instruction set architectures. The compiled first and second target codes are then packaged into a unified executable file, enabling the heterogeneous instruction set processing chip to load and coordinately schedule all at once during runtime, maximizing the release of the chip's hardware performance.

[0256] The embodiments described in this application are for the purpose of more clearly illustrating the technical solutions of the embodiments of this application, and do not constitute a limitation on the technical solutions provided by the embodiments of this application. As those skilled in the art will know, with the evolution of technology and the emergence of new application scenarios, the technical solutions provided by the embodiments of this application are also applicable to similar technical problems.

[0257] Those skilled in the art will understand that the technical solutions shown in the figures do not constitute a limitation on the embodiments of this application, and may include more or fewer steps than shown, or combine certain steps, or different steps.

[0258] The device embodiments described above are merely illustrative. The units described as separate components may or may not be physically separate; that is, they may be located in one place or distributed across multiple network units. Some or all of the modules can be selected to achieve the purpose of this embodiment according to actual needs.

[0259] Those skilled in the art will understand that all or some of the steps in the methods disclosed above, as well as the functional modules / units in the systems and devices, can be implemented as software, firmware, hardware, or suitable combinations thereof.

[0260] The terms “first,” “second,” “third,” “fourth,” etc. (if present) in the specification and accompanying drawings of this application are used to distinguish similar objects and are not necessarily used to describe a specific order or sequence. It should be understood that such data can be interchanged where appropriate so that the embodiments of this application described herein can be implemented in orders other than those illustrated or described herein. Furthermore, the terms “comprising” and “having,” and any variations thereof, are intended to cover non-exclusive inclusion; for example, a process, method, system, product, or apparatus that comprises a series of steps or units is not necessarily limited to those steps or units explicitly listed, but may include other steps or units not explicitly listed or inherent to such processes, methods, products, or apparatus.

[0261] The preferred embodiments of the present application have been described above with reference to the accompanying drawings, but this does not limit the scope of the claims of the present application. Any modifications, equivalent substitutions, and improvements made by those skilled in the art without departing from the scope and substance of the embodiments of the present application shall be within the scope of the claims of the present application.

Claims

1. A kernel compilation method, characterized in that, include: Obtain kernel code, the kernel code including first code of at least one first kernel function and second code of at least one second kernel function, each first kernel function having a first execution space attribute, and each second kernel function having a second execution space attribute different from the first execution space attribute; The kernel code is compiled and parsed, and at least one cross-space call relationship is identified based on the function call relationship in the kernel code. Each cross-space call relationship is formed by a first kernel function calling a second kernel function. Based on the cross-space call relationships, the first code and the second code corresponding to each cross-space call relationship are compiled and rewritten respectively. A first cross-space call processing code is generated in the corresponding first code, and a second cross-space call processing code is generated in the corresponding second code. A first intermediate representation corresponding to the first kernel function and a second intermediate representation corresponding to the second kernel function are obtained. The first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship are used to pass call parameters and perform execution synchronization when the corresponding first kernel function and second kernel function are executed by the first processing unit and the second processing unit respectively. Based on the first execution space attribute and the second execution space attribute, the first intermediate representation and the second intermediate representation are respectively compiled into first target code adapted to the first processing unit and second target code adapted to the second processing unit, and a unified executable file is generated based on the first target code and the second target code.

2. The kernel compilation method according to claim 1, characterized in that, The kernel code includes execution space declaration information for each of the first kernel functions and each of the second kernel functions; The first execution space attribute is determined based on the execution space declaration information of each of the first kernel functions, and the second execution space attribute is determined based on the execution space declaration information of each of the second kernel functions; The execution space declaration information includes the execution space identifier set in the function declaration of the corresponding kernel function.

3. The kernel compilation method according to claim 1, characterized in that, The step of compiling and rewriting the first code and the second code corresponding to each of the cross-space call relationships, based on the aforementioned cross-space call relationships, includes: The parameter layout of the call parameters involved in each of the cross-space call relationships is determined in the corresponding parameter buffer. Each parameter buffer is located in a shared storage area, which can be accessed by the first processing unit and the second processing unit. Based on the parameter layout, parameter writing code is generated in the corresponding first code, and parameter retrieval code is generated in the corresponding second code; The parameter writing code is used to write the call parameters involved in the cross-space call relationship into the corresponding parameter buffer according to the corresponding parameter layout, and the parameter retrieval code is used to retrieve the call parameters from the corresponding parameter buffer according to the corresponding parameter layout.

4. The kernel compilation method according to claim 3, characterized in that, The step of determining the parameter layout of the call parameters involved in each of the cross-space call relationships in the corresponding parameter buffer includes: Identify the dynamic call parameters in the call parameters; The parameter byte length of each dynamic call parameter is determined according to the data type of each dynamic call parameter; The aligned length of each dynamic calling parameter is determined based on the parameter byte length of each dynamic calling parameter and the preset alignment rules; Based on the arrangement order and aligned length of each of the dynamic calling parameters, determine the storage offset of each of the dynamic calling parameters in the parameter buffer; The size of the parameter buffer is determined based on the storage offset and aligned length of each of the dynamically invoked parameters.

5. The kernel compilation method according to claim 4, characterized in that, The step of generating parameter writing code in the corresponding first code and parameter retrieval code in the corresponding second code according to the parameter layout includes: Based on the size of the parameter buffer, a buffer allocation code is generated in the corresponding first code, and the buffer allocation code is used to allocate the parameter buffer in the shared storage area; Based on the storage offset of each of the dynamic call parameters, a parameter storage code is generated in the corresponding first code. The parameter storage code is used to write each of the dynamic call parameters into the corresponding storage location in the parameter buffer. In the corresponding second code, a buffer address acquisition code is generated, which is used to obtain the starting address of the parameter buffer in the shared storage area; Based on the starting address, the storage offset and data type of each of the dynamic call parameters, parameter reading code is generated in the corresponding second code. The parameter reading code is used to determine the storage address of each of the dynamic call parameters and read the corresponding parameter value from each of the storage addresses. The parameter writing code includes the buffer allocation code and the parameter storage code, and the parameter acquisition code includes the buffer address acquisition code and the parameter reading code.

6. The kernel compilation method according to claim 3, characterized in that, The step of compiling and rewriting the first code and the second code corresponding to each cross-space call relationship based on each cross-space call relationship also includes: Identify the compile-time constant parameters in the call parameters; Based on the parameter values ​​of the compile-time constant parameters, constant assignment code is generated in the corresponding second code. The compile-time constant parameters are not written to the parameter buffer.

7. The kernel compilation method according to claim 1, characterized in that, Each of the first cross-space call processing codes includes call initiation code and synchronization waiting code; Each of the aforementioned call initiation codes is used to cause the first processing unit to initiate a corresponding call request, and each of the aforementioned call requests is used to cause the second processing unit to execute a corresponding second kernel function; The synchronization waiting code is used to enable the first processing unit to wait for the corresponding second kernel function to complete execution after initiating the corresponding call request.

8. The kernel compilation method according to claim 7, characterized in that, The step of generating the first cross-space call processing code in the corresponding first code includes: Identify at least one entry function that has the first execution space attribute; Assign a synchronization object to each of the aforementioned entry functions; Based on the cross-space call relationships, call initiation code is generated in the corresponding first code. Each call initiation code is used to generate a call request carrying the identification information of the corresponding synchronization object. Determine the calling nodes in the corresponding first code for each of the cross-space calling relationships, and insert synchronization waiting code based on the corresponding synchronization object.

9. The kernel compilation method according to claim 8, characterized in that, Multiple cross-space call relationships within the same entry function use the same synchronization object; The synchronization wait code is also used to reset the synchronization state of the synchronization object after the corresponding second kernel function is executed, so that the synchronization object can be reused by the next cross-space call relationship in the entry function.

10. The kernel compilation method according to claim 1, characterized in that, In the corresponding second code, second cross-space call processing code is generated to obtain the second intermediate representation corresponding to the second kernel function, including: Identify multiple second kernel functions that have the second execution space attribute and do not have a mutual calling relationship; Determine the function identifiers of the plurality of second kernel functions respectively; A unified entry function is generated for the multiple second kernel functions; In the unified entry function, identifier acquisition code and call dispatch code are generated. The identifier acquisition code is used to obtain the function identifier carried in the call request, and the call dispatch code is used to dispatch the execution flow to the corresponding second kernel function according to the function identifier. Each of the first cross-space call processing codes is also used to generate a call request carrying the function identifier of the second kernel function to be called.

11. The kernel compilation method according to claim 1, characterized in that, Compile and rewrite the first code corresponding to each of the aforementioned cross-space call relationships, including: Identify the data transfer operations in the first code; Data transfer code is generated based on the data transfer operation. The data transfer code is used to transfer relevant computational data from the global storage area to the shared storage area. Based on the source data layout and target data layout of the computational data, data layout conversion code is generated in the data transport code. The data layout conversion code is used to convert the computational data from the source data layout to the target data layout when transporting the computational data.

12. The kernel compilation method according to claim 11, characterized in that, The source data is laid out in a row-first layout, and the target data is laid out in a row-block layout. Alternatively, the source data layout is a column-first layout, and the target data layout is a column-block layout; The data transport code is further used to transport the computational data shared by the multiple second kernel functions to the shared storage area before the execution of the multiple second kernel functions, so that the multiple second kernel functions can reuse it.

13. The kernel compilation method according to any one of claims 1 to 12, characterized in that, The first processing unit corresponds to the orchestrator in the computing chip, and the second processing unit corresponds to the execution unit in the computing chip; The orchestrator is used to perform shared storage management, data movement, and execution unit task scheduling; the execution unit is used to perform computational processing. The first object code is adapted to a 32-bit RISC architecture, and the second object code is adapted to a 64-bit RISC architecture with vector extensions.

14. A kernel execution method, characterized in that, include: Load a unified executable file, the unified executable file including a first target code adapted to a first processing unit and a second target code adapted to a second processing unit, the first target code including a first cross-space call processing code corresponding to at least one cross-space call relationship, and the second target code including a second cross-space call processing code corresponding to each of the cross-space call relationships; The first processing unit runs the first target code. When it reaches the call node corresponding to any cross-space call relationship in the first target code, the corresponding first cross-space call processing code is executed. The call parameters involved in the cross-space call relationship are written into a shared storage area that can be accessed by both the first processing unit and the second processing unit. A call request is initiated and the corresponding second kernel function is waited for to be executed. The second processing unit runs the second target code, and in response to the call request, executes the corresponding second cross-space call processing code, obtains the call parameters from the shared storage area, and executes the second kernel function determined by the call request; After the second kernel function completes execution, the first processing unit ends its waiting and continues to run the first target code.

15. A kernel compilation device, characterized in that, include: The code acquisition module is used to acquire kernel code, the kernel code including first code of at least one first kernel function and second code of at least one second kernel function, each first kernel function having a first execution space attribute, and each second kernel function having a second execution space attribute different from the first execution space attribute; The call relationship identification module is used to compile and parse the kernel code, and identify at least one cross-space call relationship based on the function call relationship in the kernel code. Each cross-space call relationship is formed by a first kernel function calling a second kernel function. The code rewriting module is used to compile and rewrite the first code and the second code corresponding to each cross-space call relationship based on each cross-space call relationship. It generates a first cross-space call processing code in the corresponding first code and a second cross-space call processing code in the corresponding second code, thereby obtaining a first intermediate representation corresponding to the first kernel function and a second intermediate representation corresponding to the second kernel function. The first cross-space call processing code and the second cross-space call processing code corresponding to each cross-space call relationship are used to pass call parameters and perform execution synchronization when the corresponding first kernel function and second kernel function are executed by the first processing unit and the second processing unit, respectively. The target code generation module is used to compile the first intermediate representation and the second intermediate representation into first target code adapted to the first processing unit and second target code adapted to the second processing unit, respectively, based on the first execution space attribute and the second execution space attribute, and to generate a unified executable file based on the first target code and the second target code.

16. A kernel programming system, characterized in that, include: The programming interface module is used to provide an execution space declaration interface and to set execution space declaration information for kernel functions in the obtained kernel code through the execution space declaration interface; The compilation module is used to determine the execution space attributes of each kernel function according to the execution space declaration information, identify at least one cross-space call relationship according to the function call relationship in the kernel code, generate a first cross-space call processing code and a second cross-space call processing code based on each cross-space call relationship, and compile the compiled and rewritten kernel code into target code adapted to the first processing unit and the second processing unit, respectively. The runtime module is used to load a unified executable file generated based on each of the target codes, so that the first processing unit and the second processing unit can complete the passing of call parameters and execution synchronization through the first cross-space call processing code and the second cross-space call processing code corresponding to each of the cross-space call relationships.

17. A chip, characterized in that, It includes a first processing unit, a second processing unit, and a shared storage area; The shared storage area can be accessed by the first processing unit and the second processing unit; The first processing unit is used to run the first target code, and when it runs to the call node corresponding to any cross-space call relationship, it writes the call parameters involved in the cross-space call relationship into the shared storage area through the corresponding first cross-space call processing code, initiates a call request and waits for the corresponding second kernel function to complete execution; The second processing unit is used to run the second target code, and in response to the call request, obtains the call parameters from the shared storage area through the corresponding second cross-space call processing code, and executes the second kernel function determined by the call request; The first processing unit is further configured to end the waiting period and continue running the first target code after the second kernel function has been executed; The first target code is compiled from at least one first kernel function with a first execution space attribute, and the second target code is compiled from at least one second kernel function with a second execution space attribute, wherein the first execution space attribute is different from the second execution space attribute.

Images