Multi-device kernel processing method, apparatus, medium, and product

Abstract

This application provides a multi-device kernel processing method, apparatus, medium, and product. The method includes: receiving multiple intermediate expressions generated by a cross-architecture compiler for the same code; converting the multiple intermediate expressions into kernel objects in a preset binary format; encapsulating the multiple kernel objects according to a preset arrangement rule to generate multiple single binary files; and linking the binary files with a host object to generate a native executable file. The solution of this application converts the multiple intermediate expressions into kernel objects in a preset binary format, eliminating the dependency on vendor-specific instruction sets and achieving cross-architecture kernel code universality; encapsulating the kernel objects into single binary files eliminates the need for recompilation during cross-platform distribution, solving the problem of binary non-portability; and linking the binary files with a host object supports loading the optimal kernel object when called by heterogeneous devices.

Description

Technical Field

[0001] This application relates to the field of computer technology, specifically to a multi-device kernel processing method, apparatus, medium, and product. Background Technology

[0002] Heterogeneous computing systems (such as combinations of CPUs and GPUs, or CPUs and FPGAs) have become core architectures for improving computing power and energy efficiency. Currently, mainstream cross-architecture compilers in the industry have partially solved the programming challenges of heterogeneous kernels through virtual instruction set encapsulation and intermediate representation conversion. However, these solutions still have the following key drawbacks:

[0003] (1) The "island effect" caused by hardware vendor binding. The limitations of proprietary virtual instruction sets and intermediate representations: Some compilers embed vendor-specific virtual instruction sets into the generated executable program, which is then compiled into device binary code during runtime. While this design optimizes compatibility with their own hardware, it results in the generated kernel object being unable to run on other vendors' devices, forming a technical barrier of hardware-compiler binding. In addition, the lack of cross-vendor hardware support forces developers to maintain multiple code branches for different hardware, significantly increasing development costs and system complexity. (2) The non-portability of device binary objects. The device binary objects generated by existing compilers only support the target architecture specified in the pre-compilation. If other architecture devices exist in the heterogeneous system, the code needs to be recompiled and the binary object replaced, resulting in a loss of flexibility in switching between heterogeneous devices. The current solution requires static binding of the target device type during compilation, and cannot dynamically select the optimal kernel based on the actual hardware resources of the heterogeneous device. (3) The rigidity of pre-compiled architecture specification and heterogeneous devices. Existing compilers require developers to specify the backend architecture type in advance, and the generated binary code cannot adapt to the dynamic addition or removal of heterogeneous devices. Summary of the Invention

[0004] At least one embodiment of this application provides a multi-device kernel processing method, apparatus, medium, and product to address the problems of hardware vendor lock-in, binary non-portability, and heterogeneous device rigidity in existing technologies.

[0005] To solve the above-mentioned technical problems, this application is implemented as follows:

[0006] In a first aspect, embodiments of this application provide a multi-device kernel processing method, including:

[0007] It can receive multiple intermediate representations generated by cross-architecture compilers for the same code;

[0008] The intermediate expressions are converted into kernel objects in a preset binary format. The kernel objects are then encapsulated according to a preset arrangement rule to generate multiple single binary files.

[0009] The binary file is linked with the host object to generate a native executable file.

[0010] Optionally, the multiple intermediate expressions are converted into kernel objects in a preset binary format, and the multiple kernel objects are encapsulated according to a preset arrangement rule to generate multiple single binary files, including:

[0011] The intermediate expressions generated by multiple vendor toolchains are converted into multiple target intermediate expressions in a preset binary format; the intermediate expressions include device kernel code generated by a first type of vendor toolchain and computer kernel code generated by a second type of vendor toolchain;

[0012] The wrapped kernel object is determined using a kernel wrapper and multiple target intermediate representations;

[0013] According to a preset arrangement rule, multiple kernel objects are grouped and assembled to generate multiple individual binary files, and the multiple binary files are written to the target storage path in order of device architecture priority.

[0014] Optionally, according to a preset arrangement rule, multiple kernel objects are grouped and assembled to generate multiple individual binary files, including:

[0015] Determine the device type identifier, memory alignment rule, and unique symbol identifier for each kernel object;

[0016] The kernel objects are grouped according to a preset arrangement rule, which includes at least one of the following: grouping by device type, assigning multiple kernel objects of the same device type to the same group; grouping by kernel function category, assigning multiple kernel objects of the same device that implement the same computing function to the same group; grouping by optimization level, assigning kernel objects with the same compilation optimization flag to the same group.

[0017] Each group of kernel objects is encapsulated to generate the binary file.

[0018] Optionally, before encapsulating each group of kernel objects and generating the binary file, the method further includes:

[0019] Extract the kernel name and debug symbols of the kernel object to generate a symbol file, and extract the device type, binary format and architecture characteristic metadata to generate an attribute file;

[0020] Based on the multiple attribute files and the symbol file, the header information of the binary file is generated, and the header information is written into the target segment of the binary file.

[0021] Optionally, the binary file includes at least one of the following information:

[0022] The file header information includes a list of dependencies of shared libraries or external symbols that the current binary file depends on, device version information, and a magic number identifier used to identify the format of the binary file.

[0023] Binary data segment;

[0024] The local symbol table is used to record the mapping relationship between symbol names and binary data segment offsets within a group;

[0025] The paths to all binary files and their corresponding device types or function categories;

[0026] The global symbol table records the names of all kernel symbols, their associated binary files, and their offsets within those files.

[0027] The memory alignment rule table stores the memory alignment values ​​for each device type.

[0028] Optionally, the binary file is linked with the host object to generate a native executable file, including:

[0029] A kernel registration function is embedded in the start field of the binary file, and an unregistration function is embedded in the end field of the binary file; the kernel registration function is used to load the kernel object into the target heterogeneous device memory and generate a kernel handle table when called, and the unregistration function is used to release device memory and handle resources;

[0030] The single binary file containing the kernel registration function and the unregistration function is linked with the host object to form a single executable file.

[0031] Optionally, a kernel registration function is embedded in the start field of the binary file, including:

[0032] Parse the kernel object in the binary file;

[0033] Based on the type and instruction set of the target heterogeneous device, the device driver interface is invoked to allocate resources in the device memory and generate a device image.

[0034] Check if the kernel parameter options are compatible with the kernel object;

[0035] If effective compatibility is achieved, parameter bit codes are generated according to the device register rules, the kernel object is loaded into the memory of the target heterogeneous device, and a kernel handle table is generated;

[0036] If ineffective compatibility is found, the kernel name of the kernel object is extracted to generate a unique kernel identifier, which is then added to the global cache list and awaits dynamic adjustment.

[0037] The address of the device image is bound to the parameter bitcode or the unique kernel identifier, so that the host object can call the kernel through the handle table.

[0038] Secondly, embodiments of this application provide a multi-device kernel processing apparatus, including:

[0039] The receiving module is used to receive multiple intermediate representations generated by cross-architecture compilers for the same code;

[0040] The first processing module is used to convert multiple intermediate expressions into kernel objects in a preset binary format, and encapsulate the multiple kernel objects according to a preset arrangement rule to generate multiple single binary files.

[0041] The second processing module is used to link the binary file with the host object to generate a native executable file.

[0042] Thirdly, embodiments of this application provide a computer-readable storage medium storing a program that, when executed by a processor, implements the steps of the method described in the first aspect.

[0043] Fourthly, embodiments of this application provide a computer program product, including computer instructions, which, when executed by a processor, implement the steps of the method described in the first aspect.

[0044] Compared with existing technologies, the multi-device kernel processing method, apparatus, medium, and product provided in this application receive multiple intermediate expressions generated by a cross-architecture compiler for the same code; convert the multiple intermediate expressions into kernel objects in a preset binary format; encapsulate the multiple kernel objects according to a preset arrangement rule to generate multiple single binary files; and link the binary files with a host object to generate a native executable file. The solution of this application converts the multiple intermediate expressions into kernel objects in a preset binary format, eliminating the dependency on vendor-specific instruction sets and achieving cross-architecture kernel code universality; encapsulates the kernel objects into single binary files, eliminating the need for recompilation during cross-platform distribution and solving the problem of binary non-portability; and links the binary files with a host object, supporting the loading of the optimal kernel object when called by heterogeneous devices. Attached Figure Description

[0045] Various other advantages and benefits will become apparent to those skilled in the art upon reading the following detailed description of preferred embodiments. The accompanying drawings are for illustrative purposes only and are not intended to limit the scope of this application. Furthermore, the same reference numerals denote the same parts throughout the drawings. In the drawings:

[0046] Figure 1 A flowchart of a multi-device kernel processing method provided in an embodiment of this application;

[0047] Figure 2 A diagram illustrating the relationship between the solution provided in this application and the cross-architecture compiler, used for embodiments of this application.

[0048] Figure 3 A flowchart illustrating a single kernel packaging process provided in an embodiment of this application;

[0049] Figure 4 A flowchart illustrating kernel file packaging provided in this application embodiment;

[0050] Figure 5 A diagram showing the relationship between the kernel data structure of a binary file provided in this application embodiment;

[0051] Figure 6 A flowchart of the kernel registration function provided in this application embodiment;

[0052] Figure 7 The structure of the multi-device kernel processing method provided in the embodiments of this application is shown. Detailed Implementation

[0053] The terms "first," "second," etc., used in this application are used to distinguish similar objects and not to describe a specific order or sequence. It should be understood that such terms can be used interchangeably where appropriate so that embodiments of this application can be implemented in orders other than those illustrated or described herein, and the objects distinguished by "first" and "second" are generally of the same class, without limiting the number of objects; for example, the first object can be one or more. Furthermore, "or" in this application indicates at least one of the connected objects. For example, "A or B" covers three scenarios: Scenario 1: including A but not B; Scenario 2: including B but not A; Scenario 3: including both A and B. The character " / " generally indicates that the preceding and following objects are in an "or" relationship.

[0054] The term "instruction" in this application can be either a direct instruction (or explicit instruction) or an indirect instruction (or implicit instruction). A direct instruction can be understood as one in which the sender explicitly informs the receiver of specific information, the operation to be performed, or the requested result, etc., in the instruction sent. An indirect instruction can be understood as one in which the receiver determines the corresponding information based on the instruction sent by the sender, or makes a judgment and determines the operation to be performed or the requested result, etc., based on the judgment result.

[0055] Please refer to Figure 1 This application provides a multi-device kernel processing method, including:

[0056] Step 11: Receive multiple intermediate representations generated by the cross-architecture compiler for the same code.

[0057] In this embodiment, a cross-architecture compiler (such as a heterogeneous toolchain based on LLVM) compiles the same source code into a device-specific intermediate representation (IR), ensuring instruction set compatibility across different hardware architectures. The intermediate representation masks hardware differences, allowing the computational logic of devices such as CPUs, GPUs, and NPUs to be parsed within the same framework. The cross-architecture compiler includes, but is not limited to, a heterogeneous toolchain based on LLVM. The LLVM format refers to the format of the intermediate representation it uses. Intermediate representations include, but are not limited to: a hardware-independent virtual instruction set defined by vendor N, which is a virtual intermediate representation (IR); the intermediate representation language used by vendor H's NPU (Neural Network Processing Unit); and the LLVM intermediate representation language based on the RISC-V instruction set, i.e., RISC-VLLVM IR.

[0058] Step 12: Convert the multiple intermediate expressions into kernel objects in a preset binary format, and encapsulate the multiple kernel objects according to a preset arrangement rule to generate multiple single binary files.

[0059] In this embodiment, the compiler frontend generates a device-independent intermediate representation, and the backend optimizes and generates a kernel object in a corresponding preset binary format for the target device. Here, the preset binary format can be SPIR-V. This application further encapsulates multiple kernel objects according to preset arrangement rules. These preset arrangement rules include, but are not limited to: device priority sorting: dynamically adjusting the order of intermediate representations based on computing power density, prioritizing the loading of code segments from high-computing-power devices (GPU > NPU > CPU). Memory layout alignment: dividing binary segments according to device memory access granularity (e.g., 128-byte alignment for GPUs, 64-byte alignment for NPUs) to avoid cross-device access conflicts. Encapsulation mechanism: packaging the converted intermediate representation (SPIR-V) and metadata (device ID, memory mapping table) into a single binary file, supporting the coexistence of multiple device code segments.

[0060] Optionally, this application supports dynamic relocation table embedding to provide address correction basis for subsequent heterogeneous device loading and ensure cross-architecture binary resolvability.

[0061] Step 13: Link the binary file with the host object to generate a native executable file.

[0062] In this embodiment, the host code of the host object is bound to the device code segment in a single binary file through a system call proxy layer, bypassing the overhead of user-mode-kernel-mode switching. Module loading is completed in kernel mode using the target heterogeneous device, avoiding the latency of user-mode file copying in traditional solutions. In the executable file generation process of this application, host logic and device code segments are integrated to generate a cross-architecture native executable file, supporting direct deployment in hybrid device clusters.

[0063] The proposed solution unifies the intermediate representation before encapsulating it into a single binary file, enabling seamless operation of the same binary file on multiple heterogeneous devices and effectively improving deployment efficiency. Dynamically loading the optimal code segment based on device priority reduces latency during task scheduling and effectively improves resource utilization. Direct loading in kernel mode reduces the number of system calls and lowers context switching overhead. The single binary file in this application contains code segments and metadata for multiple devices, avoiding repeated compilation for different hardware and shortening the development cycle.

[0064] Optionally, step 12 above includes:

[0065] The intermediate expressions generated by multiple vendor toolchains are converted into multiple target intermediate expressions in a preset binary format; the intermediate expressions include device kernel code generated by a first type of vendor toolchain and computer kernel code generated by a second type of vendor toolchain;

[0066] The wrapped kernel object is determined using a kernel wrapper and multiple target intermediate representations;

[0067] According to a preset arrangement rule, multiple kernel objects are grouped and assembled to generate multiple individual binary files, and the multiple binary files are written to the target storage path in order of device architecture priority.

[0068] It should be noted that the device kernel code generated by the first type of vendor toolchain (such as vendor N's ​​toolchain) includes, but is not limited to, PTX and HSACO intermediate formats; the computer kernel code generated by the second type of vendor toolchain (such as vendor H's toolchain) includes, but is not limited to, CANN IR and SYCL intermediate formats.

[0069] In this embodiment, for device kernel code: intermediate representations are converted into a preset binary format (such as SPIR-V, device-specific ISA) using a device-specific compiler (such as nvcc, hipcc); for GPU devices, a hybrid binary format containing parallel thread scheduling instructions (such as bar.sync) and memory barrier markers is generated; for NPU devices, tensor instruction set optimization metadata (such as memory bandwidth allocation rules, operator fusion markers) is injected; for computer kernel code: it is compiled into ELF format native machine code using a general-purpose compiler (such as GCC, Clang), and memory access instructions aligned with the device architecture (such as movntdq non-temporary storage instructions) are inserted. This application also generates a globally unique symbolic identifier for each kernel object to avoid symbol conflicts across vendor toolchains.

[0070] For example, device code (GPU / NPU) can be converted to SPIR-V or device-specific binary format (such as .cann→.bin for NPU); computer code (CPU) can be compiled into ELF format native machine code (.o file) using GCC / Clang.

[0071] Optionally, this application utilizes a kernel wrapper and multiple target intermediate representations to encapsulate binary files (such as .spv, .bin, .o) generated by different vendors into a unified kernel object format, including the following fields: recording the magic number identifier, device type, version number, and dependency hash value; storing the device type identifier, memory alignment rules, and debugging information; storing optimized intermediate representation binary data; and integrating the debugging information of the original intermediate representation with device metadata to the target storage address. This application generates a standardized kernel object file, compatible with the symbol resolution requirements of dynamically loaded interfaces.

[0072] Specifically, refer to Figure 2 As shown, this application sets up multiple vendor toolchains (such as vendor toolchains A to D) to package device kernels (such as SYCL kernel A and SYCL kernel B determined by vendor toolchains A and B, respectively) and computer kernels (such as computer kernel A and computer kernel B determined by vendor toolchains C and D, respectively) from different vendors into a single kernel file in a specific order using the SPIR-V format. Simultaneously, it packages the attribute information of the kernel objects to distinguish different kernel objects, and pre-generates integration header files to store the kernel object information, enabling host code to access kernel code through heterogeneous devices. Here, the SYCL kernels A / B determined by vendor toolchains A / B are parallel computing code for GPUs / NPUs; the computer kernels A / B determined by vendor toolchains C / D are general computing logic for CPUs.

[0073] Optionally, the rules for converting the intermediate representation into a preset binary format include, but are not limited to: for the kernel object of the GPU device, converting it into a mixed format of PTX and SPIR-V and inserting device-specific parallel thread scheduling instructions; for the kernel object of the NPU device, converting it into a tensor instruction set (such as CANN IR) and injecting memory bandwidth optimization metadata; for the kernel object of the CPU device, compiling it into native machine code and adapting it by optimizing the instruction set using -march=native.

[0074] Optionally, according to a preset arrangement rule, multiple kernel objects are grouped and assembled to generate multiple individual binary files, including:

[0075] Determine the device type identifier, memory alignment rule, and unique symbol identifier for each kernel object;

[0076] The kernel objects are grouped according to a preset arrangement rule, which includes at least one of the following: grouping by device type, assigning multiple kernel objects of the same device type to the same group; grouping by kernel function category, assigning multiple kernel objects of the same device that implement the same computing function to the same group; grouping by optimization level, assigning kernel objects with the same compilation optimization flag to the same group.

[0077] Each group of kernel objects is encapsulated to generate the binary file.

[0078] In this embodiment, the kernel object attribute extraction process determines the device type identifier, memory alignment rule, and unique symbol identifier for each kernel object. The device type identifier uniquely identifies the target device architecture of the kernel object and is obtained by parsing the metadata segment of the intermediate expression. The memory alignment rule can be preset and is used to record the minimum alignment unit of the device memory for address correction during dynamic loading. The unique symbol identifier is a globally unique symbol name generated for each kernel to avoid cross-vendor symbol conflicts. This application embeds the above attributes into the metadata segment of the kernel object in key-value pair format.

[0079] The grouping and assembly rules of this application implement at least one of the following: grouping by device type: merging kernel objects with the same device type identifier into the same group; grouping by function category: merging across devices based on the function prefix of the symbol name to achieve a unified calling interface for heterogeneous kernels; grouping by optimization level: separating debug and optimized kernels according to compilation optimization flags (such as -O0, -O3) to adapt to different operating scenarios. In this application, if the binary size within a group exceeds a preset threshold, it is split into multiple subgroups according to the symbol hash value; custom grouping labels can be specified via command-line parameters.

[0080] Optionally, before encapsulating each group of kernel objects and generating the binary file, the method further includes:

[0081] Extract the kernel name and debug symbols of the kernel object to generate a symbol file, and extract the device type, binary format and architecture characteristic metadata to generate an attribute file;

[0082] Based on the multiple attribute files and the symbol file, the header information of the binary file is generated, and the header information is written into the target segment of the binary file.

[0083] In this embodiment, the kernel object is used as the input source. The kernel name and debugging symbols are extracted to generate a symbol file, and the device type, binary format and architecture characteristic metadata are extracted to generate an attribute file. The attribute files and symbol files corresponding to multiple kernel files are used to generate header information based on the kernel object metadata and command line parameters. The header information is first written to a specified segment of the target binary file (for example, it can be written to ".kernel_offloading-spirv64"), and then written to the kernel object file in sequence.

[0084] Optionally, the binary file includes at least one of the following information:

[0085] The file header information includes a list of dependencies of shared libraries or external symbols that the current binary file depends on, device version information, and a magic number identifier used to identify the format of the binary file.

[0086] Binary data segment;

[0087] The local symbol table is used to record the mapping relationship between symbol names and binary data segment offsets within a group;

[0088] The paths to all binary files and their corresponding device types or function categories;

[0089] The global symbol table records the names of all kernel symbols, their associated binary files, and their offsets within those files.

[0090] The memory alignment rule table stores the memory alignment values ​​for each device type.

[0091] In this embodiment, the file header information includes a magic number identifier (e.g., 0x7FELF) and a version number (e.g., 1.0.0) for verifying file format compatibility; the device type, number of kernels, and globally unique group identifier (GUID) within the group. Binary data segment: merges the intermediate representation code of all kernel objects within the group, removes duplicate common functions; inserts a relocation table to record the corrected offset from logical address to device physical address. Local symbol table: maps symbol names within the group to data segment offsets (e.g., conv2d_gpu_v1→0x00001000), supporting symbol resolution during dynamic loading.

[0092] The binary file of this application will remove unreferenced functions or data segments through static analysis; it will also extract common code shared by multiple kernels (such as memory copy functions) into independent segments to reduce redundancy.

[0093] The steps described above in this application can achieve the following: reducing the number of binary files through modular assembly, thereby reducing file addressing and loading overhead; reducing storage usage through code merging and dead code elimination; ensuring a unified calling interface for heterogeneous kernels through standardized encapsulation formats and local symbol tables; and supporting user-defined grouping rules and dynamic adjustment strategies to adapt to diverse deployment scenarios.

[0094] In one specific embodiment, the kernel object wrapping and binary file generation implementation logic is as follows:

[0095] First, a unified encapsulation of intermediate representations is performed. Device kernel code generated by multiple vendor toolchains (such as PTX for GPUs and CANN IR for NPUs) and computation graph kernels are uniformly converted to the SPIR-V intermediate format, eliminating the heterogeneity of vendor-specific formats. An example of attribute metadata injection is as follows: kernel attributes (device type, memory alignment rules, version number) are embedded in the SPIR-V file and marked using the OpDecorate instruction.

[0096] Secondly, the kernel wrapper is implemented. Kernel object encapsulation: SPIR-V code and metadata are packaged into a separate kernel file, containing ELF-formatted code segments, metadata segments, and relocation tables. Globally unique symbol identifiers are generated to avoid cross-vendor symbol conflicts. Header file generation: Header files are generated during the pre-compilation stage, declaring kernel symbols, device types, and memory layouts, and mapping symbols to executable file segment addresses via macros.

[0097] Further, grouping, assembly, and storage are performed. Grouping strategy: Grouping is done by device type (GPU / NPU), function category (convolution / matrix operation), or optimization level to reduce loading complexity on heterogeneous devices. Dynamic splitting rule: If the number of files in a group exceeds a threshold, it is split into subgroups based on hash value. Binary file generation: Duplicate common code (such as math libraries) is removed, and objects within the group are merged through a linker script to generate a single binary file. A relocation table is recorded to correct logical addresses to device physical addresses.

[0098] Storage rules include writing specific segments of the executable file according to device architecture priority (GPU→NPU→CPU) and storing them contiguously to optimize I / O efficiency.

[0099] Finally, dynamic loading from heterogeneous devices is utilized. The kernel is queried from the storage segment based on the device type and symbol identifier, loaded into the target device's memory, and address relocation and dependency resolution are performed. The contiguous storage structure reduces addressing overhead, resulting in a significant improvement in loading speed in actual tests.

[0100] The above process in this application achieves unified management and efficient execution of cross-vendor kernel objects through SPIR-V standardization, dynamic metadata encapsulation, and grouped storage strategies.

[0101] Optionally, step 13 above includes:

[0102] A kernel registration function is embedded in the start field of the binary file, and an unregistration function is embedded in the end field of the binary file; the kernel registration function is used to load the kernel object into the target heterogeneous device memory and generate a kernel handle table when called, and the unregistration function is used to release device memory and handle resources;

[0103] The single binary file containing the kernel registration function and the unregistration function is linked with the host object to form a single executable file.

[0104] In this embodiment, a kernel registration / unregistration function embedding mechanism is set up for binary files. This application proposes embedding a kernel registration function at the beginning of the binary file and a kernel unregistration function at the end, forming a self-managed kernel encapsulation structure. Its core purpose is to achieve automated management of the entire lifecycle of kernel objects. The kernel unregistration function in this application is used to: release kernel objects in device memory; destroy the handle table and cache resources to prevent memory leaks. The execution flow of the kernel unregistration function includes: handle table traversal: based on the device address recorded in the handle table, calling the driver interface (such as cuMemFree) to release video memory; cache cleanup: removing the entry corresponding to the unique identifier from the global cache list; resource reclamation: destroying parameter bits, thread pools, and other heterogeneous device temporary resources.

[0105] In this embodiment, the binary file containing the registration / unregistration function is linked with the host object into a single executable file; the segment layout of the binary file is specified through the compilation and linking script. During execution on a heterogeneous device: Startup phase: The registration function at the beginning of the binary file is automatically executed to complete kernel loading and handle table initialization; Exit phase: The unregistration function is automatically executed to ensure resource release.

[0106] Optionally, a kernel registration function is embedded in the start field of the binary file, including:

[0107] Parse the kernel object in the binary file;

[0108] Based on the type and instruction set of the target heterogeneous device, the device driver interface is invoked to allocate resources in the device memory and generate a device image.

[0109] Check if the kernel parameter options are compatible with the kernel object;

[0110] If effective compatibility is achieved, parameter bit codes are generated according to the device register rules, the kernel object is loaded into the memory of the target heterogeneous device, and a kernel handle table is generated;

[0111] If ineffective compatibility is found, the kernel name of the kernel object is extracted to generate a unique kernel identifier, which is then added to the global cache list and awaits dynamic adjustment.

[0112] The address of the device image is bound to the parameter bitcode or the unique kernel identifier, so that the host object can call the kernel through the handle table.

[0113] In this embodiment, the kernel registration function is responsible for loading the kernel object from the binary file into the memory of the target heterogeneous device (such as GPU, NPU); generating a kernel handle table to record the mapping relationship between device memory addresses and kernel symbols for host code to call. It reads the SPIR-V code segment and metadata (such as device type and memory alignment rules) from the binary file; calls the corresponding driver interface (such as CUDADriver API) according to the device type (such as DeviceType=GPU), allocates space in the device's video memory, and generates a device image; verifies whether the kernel parameters (such as input / output tensor dimensions and data types) are compatible with the target device's instruction set and register rules (e.g., the thread block size needs to be checked for GPU's SIMT architecture); if compatible, it generates parameter bitcodes according to the device register rules and loads the kernel object into the device memory; records the mapping relationship between device image address and parameter bitcodes in the handle table; if incompatible, it extracts the kernel name to generate a unique kernel identifier and adds it to the global cache list; and waits for the heterogeneous device to dynamically adjust parameters, downgrading to a compatible version or triggering recompilation. This application binds the physical address of the device image to the parameter bitcode (or unique identifier), and the host code calls the kernel through the handle table index.

[0114] This application's registration / unregistration functions implement a self-closing management system for kernel loading and unloading, avoiding the risk of resource leaks caused by manual API calls. Through a compatibility check and dynamic caching mechanism, it supports lazy loading and adaptive adjustments for incompatible kernels (such as parameter remapping or version rollback). Parameter bitcodes are directly bound to device addresses, reducing data transfer overhead between host and device. A single executable file integrates the kernel and host code, eliminating dependencies on multi-file distribution and improving deployment robustness. This application improves the reliability and execution efficiency of kernel deployment in heterogeneous computing scenarios through embedded self-management functions and a dynamic compatibility mechanism.

[0115] In another specific embodiment, the multi-device kernel packaging process provided in this application includes:

[0116] Step 1: The compiler uses the kernel files generated by the vendor as input, extracting key information such as the kernel name and debug symbols from these files, and generating a detailed symbol file accordingly. Simultaneously, the compiler also organizes and stores kernel-related metadata such as device type, binary format, and architecture characteristics, as well as necessary input parameters from the compilation process, in a property file. Finally, it packages all of this into a single binary file.

[0117] Reference Figure 3As shown, if a kernel file is read, the kernel file is processed to extract the kernel file's own content, the attribute file that records kernel information, and the symbol file that extracts debugging information such as the kernel name. The kernel file's own content, attribute file, and symbol file are then saved as a complete meshing object. Subsequent processing of this kernel object determines the binary file.

[0118] Step 2: The compiler takes the attribute files and symbol files generated by different kernel files, as well as the kernel files themselves, as input and stores them in memory as kernel objects. Then, it creates file header information using the information contained in the kernel objects and the information specified by command-line arguments. Finally, it writes the file header information to the file first, then writes it to the kernel object files one by one, and finally packages them into the ".kernel_offloading-spirv64" segment of a single binary file.

[0119] Reference Figure 4 The single kernel package processing flow is shown as follows: The operation begins by loading the kernel file and saving it as a kernel object. The binary format is determined based on the kernel object format and the kernel file extension. The binary architecture type is determined based on command-line parameters. The kernel object is checked for debug information or whether it has undergone link-time optimization, and these flags are saved as binary flag information. The next kernel file is traversed and processed based on the kernel object's device kernel size and offset information. After traversal, the file header information is written to the file, and the kernel object is written to the file in one go. The operation ends.

[0120] The binary file kernel data structure of this application can be referred to Figure 5 As shown, the header information added to each binary file is as follows. The header information used to identify the binary file ensures compatibility with the LLVM format, allowing the location of all embedded segments even after they have been merged by the linker. The kernel header information, which identifies kernel objects, records important kernel-related information; detailed correspondences can be found in Table 1.

[0121] Table 1. Binary File Header Information Fields:

[0122]

[0123]

[0124] Step 3: The compiler embeds the kernel registration function in the "start" section of the binary file and the unregistered kernel function in the "stop" section. The kernel registration function plays a crucial role when the executable file starts up; it is responsible for loading kernel information, successfully loading all kernels into memory, and registering them with heterogeneous devices.

[0125] Among them, reference Figure 6 The kernel registration function flow shown is as follows: The registration function begins; kernel data is loaded based on the kernel binary information, and a device image object is created; if the obtained kernel parameter option information is valid, a kernel parameter bitcode is created based on the device attribute information and stored in the mapping table corresponding to the information name; otherwise, a unique kernel identifier object is created based on the kernel name and added to the list cache; the registration function ends.

[0126] After these steps, a single binary file is generated. This file is then processed by a code generation tool, transformed into machine code, and linked with a host object generated from the host code to create a native streaming file.

[0127] In summary, this application proposes a multi-device kernel mechanism. By having the compiler wrap device kernel objects in a specific way, it ensures that multi-device kernel objects follow a specific arrangement rule. This allows the host side to locate the required kernel object based on core parameters such as device type and kernel name. Combined with heterogeneous devices, this enables natively streamed executable programs to load kernel information preferentially at startup, ensuring that heterogeneous devices can correctly call these kernels. The kernel object encapsulation process is completely decoupled from the pre- and post-processes. When it is necessary to extend support for new vendor devices, this application only needs to add the corresponding attribute definitions to include them in the support list for heterogeneous devices to use.

[0128] This application uses a compiler to wrap device kernel objects in a specific way, ensuring that multiple device kernel objects can be packaged into a single binary file according to specific arrangement rules, and adds header information to ensure that the header information is compatible with the LLVM format, so that all embedded segments can still be located after they have been merged by the linker. During kernel wrapping, kernel-related attribute information such as device type, binary format, and architecture characteristics, which are necessary for heterogeneous devices to call the kernel, is added, allowing heterogeneous devices to select the appropriate kernel to call based on the backend device.

[0129] This application utilizes a multi-device kernel mechanism, enabling the compiler to uniformly package kernel objects output by all other vendor toolchains and add necessary header information during the packaging process to ensure high compatibility of the generated binary files. It not only provides a seamless autoloading method for the kernel—meaning the kernel can be automatically loaded into the system without explicit invocation or configuration—but also provides necessary attribute information for heterogeneous devices. This information allows heterogeneous devices to select and invoke the appropriate kernel based on the type and characteristics of the backend device. In this way, this application achieves kernel-backend device matching, thereby improving system flexibility and compatibility.

[0130] The various methods of the embodiments of this application have been described above. Apparatus for implementing the above methods will now be provided.

[0131] Please refer to Figure 7 This application also provides a multi-device kernel processing apparatus, including:

[0132] The receiving module 71 is used to receive multiple intermediate representations generated by the cross-architecture compiler for the same code;

[0133] The first processing module 72 is used to convert multiple intermediate expressions into kernel objects in a preset binary format, and encapsulate the multiple kernel objects according to a preset arrangement rule to generate multiple single binary files.

[0134] The second processing module 73 is used to link the binary file with the host object to generate a native executable file.

[0135] Optionally, the first processing module 72 shown includes:

[0136] The first processing unit is configured to convert the intermediate expressions generated by multiple vendor toolchains into multiple target intermediate expressions in a preset binary format; the intermediate expressions include device kernel code generated by a first type of vendor toolchain and computer kernel code generated by a second type of vendor toolchain;

[0137] The first determining unit is used to determine the packaged kernel object using the kernel wrapper and multiple target intermediate representations;

[0138] The second processing unit is used to assemble multiple kernel objects according to a preset arrangement rule, generate multiple single binary files, and write the multiple binary files into the target storage path in order of device architecture priority.

[0139] Optionally, the second processing unit is specifically used for:

[0140] The kernel objects are grouped according to a preset arrangement rule, which includes at least one of the following: grouping by device type, assigning multiple kernel objects of the same device type to the same group; grouping by kernel function category, assigning multiple kernel objects of the same device that implement the same computing function to the same group; grouping by optimization level, assigning kernel objects with the same compilation optimization flag to the same group.

[0141] Each group of kernel objects is encapsulated to generate the binary file.

[0142] Optionally, the multi-device kernel processing apparatus of this application further includes:

[0143] The third processing module is used to extract the kernel name and debugging symbols of the kernel object to generate a symbol file, and to extract the device type, binary format and architecture characteristic metadata to generate an attribute file;

[0144] The fourth processing module is used to generate the header information of the binary file based on the multiple attribute files and the symbol file, and write the header information into the target segment of the binary file.

[0145] It should be noted that the binary file includes at least one of the following pieces of information:

[0146] The file header information includes a list of dependencies of shared libraries or external symbols that the current binary file depends on, device version information, and a magic number identifier used to identify the format of the binary file.

[0147] Binary data segment;

[0148] The local symbol table is used to record the mapping relationship between symbol names and binary data segment offsets within a group;

[0149] The paths to all binary files and their corresponding device types or function categories;

[0150] The global symbol table records the names of all kernel symbols, their associated binary files, and their offsets within those files.

[0151] The memory alignment rule table stores the memory alignment values ​​for each device type.

[0152] Optionally, the second processing module 73 includes:

[0153] The third processing unit is used to embed a kernel registration function in the start field of the binary file and an unregistration function in the end field of the binary file; the kernel registration function is used to load the kernel object into the target heterogeneous device memory and generate a kernel handle table when called, and the unregistration function is used to release the device memory and handle resources;

[0154] The fourth processing unit is used to link a single binary file containing the kernel registration function and the unregistration function with the host object to form a single executable file.

[0155] Optionally, when the kernel registration function is embedded in the start field of the binary file, the third processing unit is specifically used for:

[0156] Parse the kernel object in the binary file;

[0157] Based on the type and instruction set of the target heterogeneous device, the device driver interface is invoked to allocate resources in the device memory and generate a device image.

[0158] Check if the kernel parameter options are compatible with the kernel object;

[0159] If effective compatibility is achieved, parameter bit codes are generated according to the device register rules, the kernel object is loaded into the memory of the target heterogeneous device, and a kernel handle table is generated;

[0160] If ineffective compatibility is found, the kernel name of the kernel object is extracted to generate a unique kernel identifier, which is then added to the global cache list and awaits dynamic adjustment.

[0161] The address of the device image is bound to the parameter bitcode or the unique kernel identifier, so that the host object can call the kernel through the handle table.

[0162] It should be noted that the device in this embodiment corresponds to the method described above, and the implementation methods in each of the above embodiments are applicable to the embodiment of this device, achieving the same technical effect. The device provided in this application embodiment can implement all the method steps implemented in the above method embodiments and achieve the same technical effect. Therefore, the parts and beneficial effects that are the same as those in the method embodiments will not be described in detail here.

[0163] This application also provides a computer-readable storage medium storing a computer program. When executed by a processor, the computer program implements the various processes of the above-described multi-device kernel processing method embodiments and achieves the same technical effects. To avoid repetition, it will not be described again here. The computer-readable storage medium may be a read-only memory (ROM), a random access memory (RAM), a magnetic disk, or an optical disk.

[0164] This application also provides a computer program product, including computer instructions. When the computer instructions are executed by a processor, they implement the various processes of the above-described multi-device kernel processing method embodiments and achieve the same technical effects. To avoid repetition, they will not be described again here.

[0165] It should be noted that, in this document, the terms "comprising," "including," or any other variations thereof are intended to cover non-exclusive inclusion, such that a process, method, article, or apparatus that comprises a list of elements includes not only those elements but also other elements not expressly listed, or elements inherent to such a process, method, article, or apparatus. Unless otherwise specified, an element defined by the phrase "comprising one..." does not exclude the presence of other identical elements in the process, method, article, or apparatus that includes that element.

[0166] Through the above description of the embodiments, those skilled in the art can clearly understand that the methods of the above embodiments can be implemented by means of software plus necessary general-purpose hardware platforms. Of course, they can also be implemented by hardware, but in many cases the former is a better implementation method. Based on this understanding, the technical solution of this application, in essence, or the part that contributes to the prior art, can be embodied in the form of a software product. The computer software product is stored in a storage medium (such as ROM / RAM, magnetic disk, optical disk), and includes several instructions to cause a terminal (which may be a mobile phone, computer, server, air conditioner, or network device, etc.) to execute the methods described in the various embodiments of this application.

[0167] The embodiments of this application have been described above with reference to the accompanying drawings. However, this application is not limited to the specific embodiments described above. The specific embodiments described above are merely illustrative and not restrictive. Those skilled in the art can make many other forms under the guidance of this application without departing from the spirit and scope of the claims, and all of these forms are within the protection scope of this application.

Claims

1. A multi-device kernel processing method, characterized in that, include: It can receive multiple intermediate representations generated by cross-architecture compilers for the same code; The intermediate expressions are converted into kernel objects in a preset binary format. The kernel objects are then encapsulated according to a preset arrangement rule to generate multiple single binary files. The binary file is linked with the host object to generate a native executable file.

2. The method according to claim 1, characterized in that, The multiple intermediate expressions are converted into kernel objects in a preset binary format. These kernel objects are then encapsulated according to a preset arrangement rule to generate multiple single binary files, including: The intermediate expressions generated by multiple vendor toolchains are converted into multiple target intermediate expressions in a preset binary format; the intermediate expressions include device kernel code generated by a first type of vendor toolchain and computer kernel code generated by a second type of vendor toolchain; The wrapped kernel object is determined using a kernel wrapper and multiple target intermediate representations; According to a preset arrangement rule, multiple kernel objects are grouped and assembled to generate multiple individual binary files, and the multiple binary files are written to the target storage path in order of device architecture priority.

3. The method according to claim 2, characterized in that, According to a preset arrangement rule, multiple kernel objects are grouped and assembled to generate multiple individual binary files, including: The kernel objects are grouped according to a preset arrangement rule, which includes at least one of the following: grouping by device type, assigning multiple kernel objects of the same device type to the same group; grouping by kernel function category, assigning multiple kernel objects of the same device that implement the same computing function to the same group; grouping by optimization level, assigning kernel objects with the same compilation optimization flag to the same group. Each group of kernel objects is encapsulated to generate the binary file.

4. The method according to claim 3, characterized in that, Before encapsulating each group of kernel objects to generate the binary file, the method further includes: Extract the kernel name and debug symbols of the kernel object to generate a symbol file, and extract the device type, binary format and architecture characteristic metadata to generate an attribute file; Based on the multiple attribute files and the symbol file, the header information of the binary file is generated, and the header information is written into the target segment of the binary file.

5. The method according to claim 4, characterized in that, The binary file includes at least one of the following information: The file header information includes a list of dependencies of shared libraries or external symbols that the current binary file depends on, device version information, and a magic number identifier used to identify the format of the binary file. Binary data segment; The local symbol table is used to record the mapping relationship between symbol names and binary data segment offsets within a group; The paths to all binary files and their corresponding device types or function categories; The global symbol table records the names of all kernel symbols, their associated binary files, and their offsets within those files. The memory alignment rule table stores the memory alignment values ​​for each device type.

6. The method according to claim 1, characterized in that, Link the binary file with the host object to generate a native executable file, including: A kernel registration function is embedded in the start field of the binary file, and an unregistration function is embedded in the end field of the binary file; the kernel registration function is used to load the kernel object into the target heterogeneous device memory and generate a kernel handle table when called, and the unregistration function is used to release device memory and handle resources; The single binary file containing the kernel registration function and the unregistration function is linked with the host object to form a single executable file.

7. The method according to claim 6, characterized in that, The kernel registration function is embedded in the start field of the binary file, including: Parse the kernel object in the binary file; Based on the type and instruction set of the target heterogeneous device, the device driver interface is invoked to allocate resources in the device memory and generate a device image. Check if the kernel parameter options are compatible with the kernel object; If effective compatibility is achieved, parameter bit codes are generated according to the device register rules, the kernel object is loaded into the memory of the target heterogeneous device, and a kernel handle table is generated; If ineffective compatibility is found, the kernel name of the kernel object is extracted to generate a unique kernel identifier, which is then added to the global cache list and awaits dynamic adjustment. The address of the device image is bound to the parameter bitcode or the unique kernel identifier, so that the host object can call the kernel through the handle table.

8. A multi-device kernel processing apparatus, characterized in that, include: The receiving module is used to receive multiple intermediate representations generated by cross-architecture compilers for the same code; The first processing module is used to convert multiple intermediate expressions into kernel objects in a preset binary format, and encapsulate the multiple kernel objects according to a preset arrangement rule to generate multiple single binary files. The second processing module is used to link the binary file with the host object to generate a native executable file.

9. A computer-readable storage medium, characterized in that, The computer-readable storage medium stores a computer program that, when executed by a processor, implements the steps of the method as described in any one of claims 1 to 7.

10. A computer program product, characterized in that, Includes computer instructions that, when executed by a processor, implement the steps of the method as described in any one of claims 1 to 7.

Images