Binary code instrumentation method, system and equipment, storage medium and program product

Abstract

The invention provides a binary code instrumentation method, system and equipment, a storage medium and a program product. The binary code instrumentation method can be applied to a combined processing device. The combined processing device comprises a computing device, an interface device, a processing device and a storage device. The computing device is configured to execute a dynamic binary instrumentation task, is mainly implemented as a single-core intelligent processor or a multi-core intelligent processor, is used for executing deep learning or machine learning computation, and can interact with the processing device through the interface device to jointly complete the dynamic binary instrumentation task. The interface device is used for transmitting data and control instructions between the computing device and the processing device. As a general processing device, the processing device performs basic control including but not limited to data handling, starting and / or stopping of the computing device, and the like. The storage device is used for storing data to be processed. According to the scheme, the flexibility of dynamic binary instrumentation is improved, and the application range is expanded.

Description

Technical Field

[0001] This application relates to the field of computer technology, and in particular to a binary code instrumentation method, system, device, storage medium, and program product. Background Technology

[0002] Dynamic Binary Instrumentation (DBI) is a technique used to modify and analyze the binary code of a program at runtime. DBI tools can insert custom code without requiring the source code to monitor, analyze, or modify program behavior. This technique has wide applications in performance analysis, security testing, debugging, and program optimization.

[0003] Dynamic binary instrumentation is typically used in CPU (Central Processing Unit) or GPU (Graphics Processing Unit) programs for performance analysis and debugging. However, current dynamic binary instrumentation methods are not only memory-intensive but also have high hardware and software requirements, resulting in limited applicability. Summary of the Invention

[0004] This application provides a binary code instrumentation method, system, device, storage medium, and program product, which achieves instrumentation by automatically selecting the implementation information used by the inserted function during linking. This not only reduces memory overhead but also makes it applicable to processor software stacks that do not have a dynamic library loading mechanism, reducing the requirements for the software stack and expanding the scope of application.

[0005] In a first aspect, this application provides a binary code instrumentation method, comprising: obtaining target code and a jump function corresponding to the target code; inserting a jump instruction pointing to the jump function into the target code to generate a first executable file; linking the first executable file and a preset dynamic library to generate a second executable file, wherein the preset dynamic library includes implementation information of the inserted function; and completing the jump function in the second executable file according to the implementation information of the inserted function to generate instrumented code.

[0006] Secondly, this application provides a dynamic binary instrumentation device, comprising:

[0007] The acquisition module is used to acquire the target code and the corresponding jump function of the target code;

[0008] The first generation module is used to insert jump instructions pointing to the springboard function into the target code to generate the first executable file;

[0009] The second generation module is used to link the first executable file and the preset dynamic library to generate a second executable file, wherein the preset dynamic library includes the implementation information of the inserted function;

[0010] The third generation module is used to complete the springboard function in the second executable file based on the implementation information of the insertion function, and generate the instrumented code.

[0011] Thirdly, this application provides a binary code instrumentation framework system, including:

[0012] The driver middleware layer is used to obtain the target code;

[0013] A code generator is used to call a pre-encapsulated instrumentation interface to obtain the jump function corresponding to the target code; insert a jump instruction pointing to the jump function into the target code to generate a first executable file; link the first executable file and a preset dynamic library to generate a second executable file, wherein the preset dynamic library includes the implementation information of the inserted function; and complete the jump function in the second executable file according to the implementation information of the inserted function to generate instrumented code.

[0014] Fourthly, embodiments of this application provide an electronic device, including:

[0015] At least one processor; and

[0016] A memory that is communicatively connected to the at least one processor;

[0017] The memory stores instructions that can be executed by the at least one processor, which, when executed by the at least one processor, cause the electronic device to perform the method described in any of the above aspects.

[0018] Fifthly, embodiments of this application provide a computer-readable storage medium storing computer-executable instructions, which, when executed by a processor, implement the method described in any of the above aspects.

[0019] Sixthly, embodiments of this application provide a computer program product, including a computer program that, when executed by a processor, implements the method described in any of the above aspects.

[0020] The binary code instrumentation method, system, device, storage medium, and program product provided in this application obtain the target code to be instrumented and the corresponding jump function. Then, a jump instruction pointing to the jump function is inserted into the target code to call the jump function, resulting in a first executable file. The first executable file is then linked with a preset dynamic library. During the linking process, the implementation information of the inserted function is specifically selected from the preset dynamic library and included in the second executable file. The jump function is further completed in the second executable file based on the implementation information of the inserted function to generate the final instrumented code. In this way, there is no need for a dynamic library loading mechanism to preload all device-side functions. Instead, the implementation information used by the inserted function is automatically selected during linking to complete the instrumentation. This not only reduces memory overhead but also makes it applicable to processor software stacks that do not have a dynamic library loading mechanism, reducing the requirements for the software stack and expanding the scope of application. Attached Figure Description

[0021] The accompanying drawings, which are incorporated in and form a part of this specification, illustrate embodiments consistent with this disclosure and, together with the description, serve to explain the principles of this disclosure.

[0022] Figure 1A A system schematic diagram of an Intel Pin, a CPU-oriented dynamic binary instrumentation framework, provided for embodiments of this application;

[0023] Figure 1B A schematic diagram of a GPU-oriented dynamic binary instrumentation NVBit framework provided in this application embodiment;

[0024] Figure 1C This is a schematic diagram of a code generation scenario after instrumentation, provided in an embodiment of this application.

[0025] Figure 2 This application provides a schematic diagram of the structure of an electronic device.

[0026] Figure 3 A component distribution diagram of a dynamic binary instrumentation frame provided in an embodiment of this application;

[0027] Figure 4 A schematic diagram of a binary code instrumentation framework system provided in this application embodiment;

[0028] Figure 5 A flowchart illustrating a binary code instrumentation method provided in an embodiment of this application;

[0029] Figure 6 This is a schematic diagram illustrating a specific scenario of a binary code instrumentation method according to an embodiment of this application.

[0030] Figure 7 A structural diagram of a circuit board provided in an embodiment of this application;

[0031] Figure 8 A structural diagram of a combined processing apparatus provided in an embodiment of this application;

[0032] Figure 9 This application provides a schematic diagram of the internal structure of a single-core computing device.

[0033] Figure 10 This application provides a schematic diagram of the internal structure of a multi-core computing device.

[0034] Figure 11 This is a schematic diagram of the internal structure of another single-core computing device provided in an embodiment of this application.

[0035] The accompanying drawings have illustrated specific embodiments of this disclosure, which will be described in more detail below. These drawings and descriptions are not intended to limit the scope of the concept in any way, but rather to illustrate the concepts of this disclosure to those skilled in the art through reference to particular embodiments. Detailed Implementation

[0036] The technical solutions in the embodiments of this disclosure will be clearly and completely described below with reference to the accompanying drawings. Obviously, the described embodiments are only a part of the embodiments of this disclosure, not all of them. Based on the embodiments in this disclosure, all other embodiments obtained by those skilled in the art without creative effort are within the scope of protection of this disclosure.

[0037] It should be understood that the terms "first," "second," etc., in the claims, specification, and drawings of this disclosure are used to distinguish different objects, not to describe a specific order. The terms "comprising" and "including" used in the specification and claims of this disclosure indicate the presence of the described features, integrals, steps, operations, elements, and / or components, but do not exclude the presence or addition of one or more other features, integrals, steps, operations, elements, components, and / or sets thereof.

[0038] It should also be understood that the terminology used herein is for the purpose of describing particular embodiments only and is not intended to limit the scope of this disclosure. As used in this disclosure and claims, the singular forms “a,” “an,” and “the” are intended to include the plural forms unless the context clearly indicates otherwise. It should also be understood that the term “and / or” as used in this disclosure and claims refers to any combination and all possible combinations of one or more of the associated listed items, and includes such combinations.

[0039] As used in this specification and claims, the term "if" may be interpreted, depending on the context, as "when," "once," "in response to determination," or "in response to detection." Similarly, the phrase "if determined" or "if [described condition or event] is detected" may be interpreted, depending on the context, as "once determined," "in response to determination," "once [described condition or event] is detected," or "in response to detection of [described condition or event]."

[0040] To clearly describe the technical solutions of the embodiments of this application, the terms involved in this application are first defined as follows:

[0041] CPU: Central Processing Unit.

[0042] NPU: Neural Processing Unit. An NPU is a specialized piece of hardware used to accelerate artificial intelligence and machine learning tasks. An NPU acts as a coprocessor, forming a heterogeneous computer system with the CPU and other host components.

[0043] Intel Pin is a dynamic binary instrumentation tool. It allows developers to insert custom code at runtime for program analysis, performance measurement, debugging, and other purposes.

[0044] Valgrind is an open-source framework for program analysis, primarily used for memory debugging, memory leak detection, and performance analysis. It monitors and analyzes programs through dynamic binary instrumentation, helping developers discover and fix various problems.

[0045] DynamoRIO (Dynamic Instrumentation and Optimization Runtime Infrastructure) is a dynamic binary instrumentation tool primarily used for program analysis, debugging, and optimization. DynamoRIO provides a powerful framework that allows users to insert custom code at runtime to monitor and modify program behavior.

[0046] NVBit: NVIDIA Binary Instrumentation Tool, a binary instrumentation framework for GPUs.

[0047] API: Application Programming Interface.

[0048] Sanitizer API: A set of tools and libraries for detecting and debugging various errors and problems in programs.

[0049] .so libraries: A shared library file, a type of dynamic link library, that allows multiple programs to share the same library code, thus saving memory and disk space. They are typically used in Unix- or Linux-based systems.

[0050] .a files: Static library files.

[0051] CUDA: Compute Unified Device Architecture, a parallel computing platform and programming model.

[0052] nvcc: NVIDIA CUDA Compiler, is a compiler used to compile CUDA programs.

[0053] ABI stands for Application Binary Interface.

[0054] HAL: Hardware Abstraction Layer.

[0055] PC: Program Counter, is a register in a computer system used to store the address of the instruction currently being executed.

[0056] VLIW: Very Long Instruction Word.

[0057] Predicate instructions are instructions whose execution is determined by the value of a predicate register. If the value of the predicate register is true, the instruction will be executed. If it is false, the instruction will be skipped.

[0058] Predicate computation refers to setting the value of the predicate register through comparison or logical operations. For example, comparing the size of two numbers and storing the result in the predicate register.

[0059] ISA: Instruction Set Architecture.

[0060] ID: Identification.

[0061] JIT: Just-In-Time compilation.

[0062] ELF: Executable and Linkable Format, an executable file format.

[0063] CNBit: A binary instrumentation tool for neural network processors.

[0064] Dynamic Binary Instrumentation (DBI) is a technique that analyzes and modifies target binary code during program execution. By inserting specific instructions or code into the target program's binary code, users can monitor, track, and modify program behavior. This technique has wide applications in performance analysis, security testing, debugging, and program optimization. Dynamic binary instrumentation is typically applied to CPU or GPU programs, commonly used for performance analysis and debugging. However, current dynamic binary instrumentation methods not only incur high memory overhead but also have high hardware and software requirements, resulting in limited applicability.

[0065] In traditional technologies, the dynamic binary instrumentation frameworks for CPUs include Intel Pin, Valgrind, and DynamoRIO. The basic principles of the three are quite similar, and Intel Pin will be used as an example for explanation.

[0066] Figure 1AThis is a system overview of Intel Pin, a CPU-oriented dynamic binary instrumentation framework. Pin consists of a Virtual Machine (VM), a Code Cache, and instrumentation APIs called by the Pintool utility. The VM includes a Just-In-Time (JIT) compiler, an Emulation Unit, and a Dispatcher. Once Pin gains control of the user program (Application), the VM coordinates its internal modules to run the user program (here, the instrumented target code). The JIT compiler disassembles, analyzes, and instruments the user program, then hands the instrumented code to the dispatcher for execution. Code generated by the JIT compiler is stored in the Code Cache for future use. Reading code from the Code Cache into the VM and storing code in the Cache and exiting the VM both involve saving and restoring the program's register states. The Emulation Unit simulates instructions that cannot be directly executed by the program, typically used for system calls requiring special handling by the VM. Both the Emulation Unit and the Code Cache can communicate with external hardware through the operating system. Internally, Pintool allows developers to access basic blocks of the target code, even disassembled instruction-level information, and call interfaces to determine the insertion location of the inserted functions. By implementing just-in-time compilation and instrumentation through Pin, and combining it with Pintool to guide the instrumentation location, tool developers can perform dynamic binary instrumentation on user programs to achieve different development objectives. However, this dynamic binary instrumentation framework is not suitable for heterogeneous programming scenarios.

[0067] Figure 1B This is a schematic diagram of an NVBit framework for GPU-based dynamic binary instrumentation. The core of this NVBit framework includes a static library (Libnvbit.a) and a header file (Nvbit.h) that provides API interfaces. Figure 1B The document demonstrates: (a) the compilation process of the NVBit tool and (b) the runtime usage process of the NVBit tool, wherein:

[0068] (a) The compilation process of the NVBit tool includes the following steps: (1) Develop a .cu file using the NVBit API, such as generating the nvbit_tool.cu file. (2) Compile the nvbit_tool.cu file using a compiler, such as compiling the nvbit_tool.cu file using nvcc-shared (this option is used to instruct the compiler to generate a shared library) to generate the corresponding binary code. (3) Link the corresponding binary code with libnvbit.a to generate a .so dynamic library. For example, instruct the compiler to link using nvbit_tool.cu-lnvbit to generate nvbit_tool.so, where "-lnvbit" is a linking option used to instruct the compiler to link the NVBit library. The `-l` option is used to specify the library to be linked, and `nvbit` is the name of the library.

[0069] (b) The runtime usage of the NVBit tool includes the following steps: To use the NVBit tool, developers need to use the LD_PRELOAD mechanism (an environment variable that allows users to preload one or more shared libraries when running a program) to instruct the Linux loader to prioritize loading the corresponding dynamic library nvbit_tool.so. Afterward, the tool uses the CUDA Driver and GPU HW (GPU hardware) by calling CUDA Driver APIs. The CUDA Driver APIs may include functions such as cuMemAlloc and cuLaunchKernel (e.g., cuMemAlloc, cuLaunchKernel, etc.). The cuMemAlloc function is used to allocate memory on the GPU device, and the cuLaunchKernel function is used to start a CUDA kernel function.

[0070] The NVbit framework core can include a code logger, code unloader, utility function loader, code generator, instruction extractor, driver middleware layer, assembler, and disassembler. The utility function loader is responsible for loading all device-side functions from the NVbit tools. After loading the instrumented functions written by the developer and its built-in state saving / restoring utility functions, the utility function loader maintains a mapping table from functions to load addresses for subsequent instruction generation by the code generator. The code generator is responsible for generating the instrumented code.

[0071] like Figure 1C As shown, Figure 1C The top left corner shows the user's original target code that needs instrumentation. Examples may include the following:

[0072] SHL R8,R 0x1

[0073] STS[R15],R8

[0074] LDG[R15+0x8], R12

[0075] Wherein, (STS[R15],R8) is the target instruction that the user wants to instrument.

[0076] Figure 1C The top right corner of the text specifies the insertion function foo (Instrumentation Function), which can include, for example:

[0077] __device__void foo(int arg){...}.

[0078] In this example, suppose the user wants to call `foo` before the instruction STS runs. To instrument the GPU instruction, the NVBit tool performs the following steps:

[0079] (i) Copying the original target code into system memory is called instrumented code.

[0080] (ii) Generate a new code segment in GPU memory, called a trampoline. For example... Figure 1C As shown, the generated springboard function can contain the following instructions:

[0081] 1. Specific utility functions are used to save the current state. Note that NVBit will only save the necessary registers, including general-purpose registers and predicate registers, based on the function's physical resource usage.

[0082] 2. At least one parameter passing instruction is required to pass the parameters needed by the instrumentation function. Figure 1C There is only one MOV32IR4,arg. Parameter passing needs to comply with the ABI constraints of the target device (GPU in this example), which is guaranteed by the Hardware Abstraction Layer (HAL).

[0083] 3. The target address is the jump instruction of the instrumentation function, and the target address is obtained by querying the mapping table recorded in the utility function loader.

[0084] 4. Specific utility functions are used to restore the current state of the scene, corresponding to the state saving mentioned earlier.

[0085] 5. The target instruction to be replaced (STS[R15], R8). Note that if the target instruction is a relative jump, the jump offset needs to be adjusted synchronously.

[0086] 6. A jump instruction with the target address being the instrumentation location is used to return to the next program counter (PC) at the instrumentation location.

[0087] (III) Modify the instrumentation code to replace the target instruction (STS[R15], R8) with a direct jump (JMP L1) to a function whose target address is a jumper. This instrumented code can be integrated into the dynamic library nvbit_tool.so to be loaded onto the device and run on the device.

[0088] While NVBit implements dynamic binary instrumentation for GPUs, NVBit tools rely on a utility function loader to achieve the final binary instrumentation functionality. However, the utility functions included in the utility function loader incur significant memory overhead. Furthermore, the utility function loader loads all device-side functions from the entire utility library, and combined with NVBit's built-in save / restore functions, the memory overhead introduced by this instruction loading is considerable. This is especially problematic when developing system-level tools using the NVBit mechanism. There might be situations where the utility library contains many instrumentation functions, but the actual target code (e.g., kernel functions) only needs a few. For example, the utility library might have 100 instrumentation functions, but the target code only needs 5. In this case, the utility function loader loads all device functions from the entire utility library, resulting in significant redundant memory overhead. If the user program code itself consumes a large amount of device memory, this can have a significant impact or even cause runtime failures. In addition, internal components such as the NVBit tool's utility function loader need to call the Driver API interface to manually load the device utility function library. This requires NVBit framework developers to be very familiar with the functions of the Driver API and how to use it, thus increasing the development difficulty.

[0089] To address at least one of the aforementioned problems, embodiments of this application provide a binary code instrumentation method. The binary code instrumentation framework and method of this application can be applied to heterogeneous programming scenarios for heterogeneous computer systems, thereby enabling the implementation of all or part of the process described in any of the following embodiments. This eliminates the need for a dynamic library loading mechanism to pre-load all device-side functions; instead, instrumentation is automatically performed at link time by selecting the implementation information used by the inserted functions. This not only reduces memory overhead but also makes it applicable to processor software stacks without a dynamic library loading mechanism, lowering the requirements for the software stack and expanding its applicability.

[0090] The heterogeneous computer system can include a general-purpose processor and a neural network processor or graphics processing unit (GPU). The general-purpose processor can act as the host side, and the neural network processor or GPU can act as the device side. The processor on the device side has a different instruction set architecture than the general-purpose processor on the host side. Based on this heterogeneous computer system, users can program for both the host side and the device side separately, and the obtained program code can include both host-side code and device-side code. The binary code instrumentation method provided in this application embodiment can be executed on the host side, the target code to be instrumented can be device-side code, and the code obtained after instrumentation can be executed on the device side.

[0091] like Figure 2 As shown, taking a neural network processor as an example, the electronic device 2 based on a heterogeneous computer system according to an embodiment of this application includes: at least one general-purpose processor 21, at least one neural network processor 22, and a memory 23. Figure 2 The example uses a general-purpose processor 21 and a neural network processor 22. The general-purpose processor 21, the neural network processor 22, and the memory 23 are connected via a bus 20. The general-purpose processor 21 can be a CPU, and the structure of the neural network processor 22 can be found in [reference needed]. Figures 9-11 The neural network processor 22 may include an arithmetic unit (IPU) and a memory unit (MPU). The MPU includes, but is not limited to, global dynamic random access memory (GDRAM), shared memory (SRAM), neuron memory (NRAM), and weight memory (WRAM). The memory 23 stores instructions that can be executed by the general-purpose processor 21 and / or the neural network processor 22.

[0092] like Figure 3As shown in the illustration, this application provides a binary code instrumentation framework system for neural network processors. This system may include: a dynamic binary instrumentation tool, a user-level API for the dynamic binary instrumentation tool, a core for the dynamic binary instrumentation tool, driver APIs for the neural network processor, an NPU driver, and NPU hardware. The user-level API of the dynamic binary instrumentation tool provides an instrumentation API, an inspection API, and a callback API for users to perform custom instrumentation. The core of the dynamic binary instrumentation tool includes a code generator, a driver interposer, and a disassembler; these components can be used to implement binary code instrumentation functionality.

[0093] The driver middleware layer is used to obtain the target code. Here, target code refers to the program code that needs to be instrumented. The driver middleware layer responds to instrumentation tasks, obtaining the target code corresponding to the task. The instrumentation task allows the user to specify the instrumentation location and the insertion function. Unlike NVBit, which requires extensive interaction with the Driver API, this framework's driver middleware layer only intercepts the API that loads device code within the Driver API to instrument the target code that needs instrumentation. This allows the framework to focus more on instrumentation itself and reduces coupling caused by the Driver API. Users do not need a deep understanding of the Driver API to complete the instrumentation of the target program.

[0094] A code generator is used to generate instrumented code. Specifically, the code generator can call a pre-encapsulated instrumentation interface to obtain the jump function corresponding to the target code; insert jump instructions pointing to the jump function into the target code to generate a first executable file; link the first executable file and a preset dynamic library to generate a second executable file, wherein the preset dynamic library includes the implementation information of the inserted function; and complete the jump function in the second executable file according to the implementation information of the inserted function to generate the instrumented code.

[0095] In this embodiment, the binary code instrumentation framework system does not include a utility function loader, and the binary code instrumentation function does not depend on a utility function loader, thereby reducing memory overhead during code loading. Furthermore, the code generator in this application obtains the implementation information of the inserted functions by just-in-time linking the first executable file and the preset dynamic library, significantly reducing memory overhead compared to NVBit's built-in save / restore functions. Simultaneously, compared to NVBit's solution, the framework structure of this application does not require a utility loader, thus reducing dependence on and coupling to the Driver API interface. This not only reduces memory overhead but also makes it applicable to processor software stacks without dynamic library loading mechanisms, lowering the requirements for the software stack and expanding its applicability.

[0096] In this embodiment, the code generator can generate a jump function corresponding to the instrumentation task based on the instrumentation location, the inserted function, and the target code specified in the instrumentation task. Then, a jump instruction pointing to the jump function is inserted at the instrumentation location in the target code to generate a first executable file. The first executable file is linked with a preset dynamic library to generate a second executable file, wherein the preset dynamic library includes the implementation information of the inserted function. In this embodiment, the linking process selectively selects the implementation information of the inserted function to be used from the preset dynamic library, includes the implementation information of the inserted function in the second executable file, and further completes the jump function in the second executable file based on the implementation information of the inserted function to generate the final instrumented code. Thus, there is no need for a dynamic library loading mechanism to preload all device-side functions; instead, the implementation information used by the inserted function is automatically selected during linking to complete the instrumentation.

[0097] In one embodiment, the system further includes a disassembler for calling a pre-packaged disassembly interface to disassemble the target code and obtain the disassembly result of the target code, so that the code generator can generate instrumented code based on the disassembly result.

[0098] like Figure 4 The diagram shown illustrates the principle of a binary code instrumentation framework system provided in this application embodiment. This framework comprises two parts: a dynamic binary instrumentation tool API Core and a dynamic binary instrumentation tool API Tool. The dynamic binary instrumentation tool API Core includes a static library (libcnsanitizer.a) and a header file (cnsanitizer.h) providing the API interface. The basic workflow of the framework is as follows:

[0099] (A) The compilation process of a dynamic binary instrumentation tool is as follows:

[0100] (1) Use the dynamic binary instrumentation tool API to develop device-side code files, such as generating the cnsanitizer_tool.mlu file.

[0101] (2) Compile the cnsanitizer_tool.mlu file using a compiler, for example, by using cncc-shared (this option is used to instruct the compiler to generate a shared library) to compile the cnsanitizer_tool.mlu file and generate the corresponding binary code.

[0102] (3) Link the corresponding binary code with cnsanitizer.a to generate a .so dynamic library. For example, instruct the compiler to link and generate cnsanitizer_tool.so using cnsanitizer_tool.mlu -lcnsanitizer, where "-lcnsanitizer" is a linking option used to instruct the compiler to link the dynamic binary instrumentation tool library. The `-l` option specifies the library to be linked, and `cnsanitizer` is the name of the library. In this embodiment, cnsanitizer_tool.so does not contain a tool function loader, nor does it have built-in functions for saving / restoring the state.

[0103] (B) The runtime usage process of dynamic binary instrumentation tools includes: If dynamic binary instrumentation tools are to be used, developers need to use the LD_PRELOAD mechanism to tell the Linux loader to load the corresponding dynamic library first.

[0104] For example, a pre-compiled application (referring to application code that has been compiled before distribution, specifically pre-compiled target code that needs instrumentation) is used. Then, the `LD_PRELOAD = cnsanitizer_tool.so` instruction tells the Linux loader to prioritize loading the corresponding dynamic library `cnsanitizer_tool.so`. Afterwards, the NPU Driver APIs are used to call the NPU Driver, which then loads the pre-compiled application onto the NPU HW (NPU hardware) for execution.

[0105] In this embodiment, libcnbin.so provides APIs for disassembling, instrumentation, and ELF manipulation to assist in the implementation of framework functions. In the dynamic binary instrumentation tool Core, both the code generator and disassembler call the relevant APIs of libcnbin.so, encapsulating them into more user-friendly methods for exposure to the dynamic binary instrumentation tool framework. By presenting these basic functions as APIs in libcnbin.so, decoupling from the Driver API is maximized. For other system APIs, the interfaces of libcnbin.so can be directly called to complete the corresponding instrumentation, without needing to call the dynamic binary instrumentation tool API and introduce unnecessary coupling.

[0106] In the framework provided in this application embodiment, the disassembler can integrate the functions of a disassembler and an instruction extractor. Its role is to parse the binary data of the instruction segment into an assembly-level instruction abstraction. While this abstraction has hardware specificities, such as differences in instructions across different NPU architectures, it can also provide general interfaces such as obtaining the binary representation of the instruction, obtaining the instruction opcode, and obtaining the instruction operands, to ensure compatibility across hardware generations as much as possible. The framework of this application embodiment provides this layer of abstraction, allowing users to process different instructions as needed.

[0107] The disassembler's capabilities rely on the disassembly API provided in libcnbin.so. To simplify user experience, the disassembler in the dynamic binary instrumentation tool Core encapsulates these APIs. After the user specifies the instrumentation logic (including the location to be instrumented, the instrumentation function to be called, etc.), the code generator is responsible for generating the corresponding instructions to achieve the final functionality. The code generator in the dynamic binary instrumentation tool Core can also utilize the instrumentation API provided by libcnbin.so ( Figure 4 This is accomplished using the CNBit interface in the system.

[0108] The difference between this embodiment and NVBit is that the libcnbin.so library can pre-encapsulate the necessary interfaces, such as the Disasm API (disassembler interface), CNBit API (instrumentation interface), and ELF API. libcnsanitizer.a can provide only basic instrumentation capabilities. The libcnsanitizer.a library in the CNSanitizer API Core will call the separate libcnbin.so library to implement the most basic disassembler and instrumentation APIs for dynamic binary instrumentation. Thus, in cnsanitizer_tool.mlu, tool developers do not need to be aware of the Driver API operations and can directly instrument the target code.

[0109] In one embodiment, such as Figure 3 As shown, this framework provides several user-level interfaces for dynamic binary instrumentation tools, such as the Callback API, Instrumentation API, and Inspection API, among which:

[0110] Callback APIs can be used to allow users to perform custom actions when the application ends, as shown below:

[0111] / / Run a specific action when the application exits

[0112] void cnsanitizer_at_exit();

[0113] The inspect API can be used to retrieve all instructions within a function, as follows:

[0114] / / Retrieve instruction sequence from instruction segment

[0115] const std::list<CNInst*> &cnsanitizer_get_instrs(const ELFText*text)

[0116] CNInst is the instruction abstraction provided by the framework of this embodiment, and may include the following interfaces:

[0117] class CNInst{

[0118] public:

[0119] / / Get disassembled text

[0120] const char* getCNAsm()const;

[0121] / / Get offset

[0122] uint32_t getOffset()const;

[0123] / / Get operands

[0124] CNOperand*getOperand(uint32_t index)const;

[0125] / / Get operation code

[0126] CNOpcode getOpcode()const;

[0127] / / Get source code information

[0128] void getLineInfo(char**file,uint32_t*line);

[0129] }

[0130] CNOperand / CNOpcode are custom data structures. After obtaining CNInst, users can perform special processing on different commands according to their needs.

[0131] The instrumentation API allows users to perform free instrumentation, as shown below:

[0132] typedef enum{IPOINT_BEFORE,IPOINT_AFTER}ipoint_t;

[0133] / / Instrument the target instruction segment

[0134] void cnsanitizer_patch_text(ELFText*text);

[0135] / / Instrument a target instruction

[0136] void cnsanitizer_insert_call(const CNInst*inst,const char*callee,ipoint point);

[0137] / / Add parameters to the function that was inserted previously

[0138] void cnsanitizer_add_call_arg(const CNOperand*arg);

[0139] void cnsanitizer_add_call_arg_const(uint64_t arg);

[0140] Please refer to Figure 5 This is a binary code instrumentation method according to an embodiment of this application. By automatically selecting the implementation information used by the inserted function during linking, instrumentation is completed. This not only reduces memory overhead but also makes it applicable to processor software stacks without dynamic library loading mechanisms, lowering the requirements for the software stack and expanding its applicability. This method can be... Figure 3 The code generator in the process is executed, and the method includes the following steps:

[0141] Step 501: Obtain the target code and the corresponding jumper function.

[0142] In this step, target code refers to the binary program code that needs to be instrumented, such as kernel function code that needs to be debugged based on instrumentation. The code generator can obtain the target code and its corresponding jumper functions.

[0143] In this context, a stepping stone function is an auxiliary function used to instrument the target code. Stepping stone functions can be generated by the code generator based on the instrumentation locations and functions specified by the user in the target code.

[0144] Step 502: Insert a jump instruction pointing to the jump function into the target code to generate the first executable file.

[0145] In this step, the first executable file includes object code and inserted jump instructions pointing to the jump function, so that the first executable file can jump to the jump function and realize the call to the jump function.

[0146] Step 503: Link the first executable file and the preset dynamic library to generate the second executable file, wherein the preset dynamic library includes the implementation information of the inserted function.

[0147] In this step, Just-In-Time (JIT) linking can be used to link the first executable file and the pre-defined dynamic libraries. During the linking process, the implementation information of the inserted functions to be used is specifically selected from the pre-defined dynamic libraries (such as the aforementioned cnsanitizer_tool.so). Thus, the implementation information of the inserted functions can be included in the second executable file through the linking operation. In this way, there is no need for a dynamic library loading mechanism to preload all device-side functions. Instead, the implementation information used by the inserted functions is automatically selected during linking to complete the instrumentation. This not only reduces memory overhead but also makes it applicable to processor software stacks that do not have a dynamic library loading mechanism, reducing the requirements for the software stack and expanding the scope of application.

[0148] Step 504: Based on the implementation information of the insertion function, complete the stepping stone function in the second executable file to generate the instrumented code.

[0149] In this step, after including the implementation information of the inserted function in the second executable file, the stepping stone function is further completed in the second executable file based on the implementation information of the inserted function, generating the final instrumented code. The implementation information of the inserted function here includes, but is not limited to, functions for saving / restoring the state of the inserted function, parameter pushing functions, etc. By completing this implementation information into the stepping stone function, the complete instrumented code is obtained. Thus, during linking, the implementation information used by the inserted function is automatically selected to complete the instrumentation, which not only reduces memory overhead but also makes it applicable to processor software stacks without dynamic library loading mechanisms, reducing the requirements for the software stack and expanding its applicability.

[0150] In the NVBit implementation described above, the utility function loader loads the entire utility function library onto the device side, regardless of the usage of the functions within the library. In some system tools based on dynamic binary instrumentation, to ensure generality, the utility function library needs to include many functions to guarantee complete functionality. However, the target code may only use a few target instructions, leading to redundant memory overhead. For example, if the utility function library contains instrumentation functions for ten types of instructions, but the target code only contains one type of instruction, loading the instrumentation functions for the remaining nine types onto the device side would result in redundant memory overhead.

[0151] In this embodiment, only the instrumentation functions corresponding to the specific type of instructions present in the target code are loaded, saving memory resources. Furthermore, the code generator in this application obtains the implementation information of the inserted functions through just-in-time linking of the first executable file and a preset dynamic library, significantly reducing memory overhead compared to NVBit's built-in save / restore functions. Simultaneously, because a JIT linking process occurs during instrumentation, and the `--gc-sections` option is enabled during linking to perform garbage collection and remove unused code and data segments, useless code segments are deleted, saving memory overhead caused by code segments.

[0152] In one embodiment, step 501 may specifically include: in response to the instrumentation task, hijacking the target code when the target code is loaded into a module and before relocation.

[0153] In this embodiment, the user can specify the instrumentation location and the insertion function to be used to insert the target code in the instrumentation task. In response to this instrumentation task, the system hijacks the corresponding driver interface (which only includes the interface for loading device-side code) during the module loading of the target code and before the target code relocation, thereby enabling the loading and modification of the target code. During module loading, by hijacking the driver interface and modifying the device-side code before relocation, the original kernel call flow can be reused to the greatest extent, reducing the dependency between the driver intermediate layer and the driver interface, thus improving the modularity and flexibility of the system. Here, module loading refers to loading the kernel (kernel function) running on the device side. In heterogeneous programming, the kernel refers to the kernel function running on the device side, such as the entry function. After the module is loaded, the functionality provided by the kernel function module can be used through kernel function calls.

[0154] For example, the driver middleware layer can hijack the corresponding driver interface to modify the target code on the device side during Module Load and before relocation. This eliminates the need for manual dynamic library loading and runtime address maintenance, and maximizes the reuse of the kernel invoke process without instrumentation, such as reusing runtime relocation and reducing coupling with the Driver API.

[0155] In one embodiment, the method further includes: calling a pre-encapsulated disassembly interface to disassemble the target code and obtain the disassembly result of the target code.

[0156] In this embodiment, the target code can be disassembled before generating the jump function. This disassembly interface can be pre-encapsulated; when disassembly of the target code is required, the interface can be directly called. For example, the method disclosed herein can automatically call the disassembly API provided in libcnbin.so to disassemble the target code, parsing the relevant instruction segment binary data into assembly-level instruction abstractions (the level of which is higher than the level of binary instructions), without requiring developers to understand the DriverAPI. This disclosure encapsulates some or all of the key basic functions using the independent dynamic library libcnbin.so, enabling components tightly coupled with the Driver API to also use it, thus offering high flexibility and expanding its applicability.

[0157] In one embodiment, the above method further includes the process of obtaining the jump function corresponding to the target code, which may specifically include: determining the target instruction corresponding to the instrumentation position specified by the instrumentation task based on the disassembly result; determining the identification information of the jump function based on the target instruction, the program counter of the instrumentation position and the insertion function specified by the instrumentation task, wherein the jump function includes at least a call to the insertion function.

[0158] In this embodiment, the target instruction refers to the instruction at the instrumentation location in the target code. The identification information of the jump function can be the jump function's ID (Identification), used to uniquely identify the jump function. In this embodiment, the identification information of the jump function can be jointly determined by the target instruction corresponding to the insertion location, the PC of the insertion location, and the function to be inserted, thereby improving the accuracy of the jump function ID. Here, the jump function can only generate calls to the inserted function, enough to allow the linker to function normally. The jump function ID can identify the uniqueness of the jump function and can be used for subsequent information tracing and identification of the jump function. For example, a mapping relationship between the jump function and its corresponding ID can be generated. When using it, the jump function ID can be used as an index to accurately find the corresponding jump function according to the mapping relationship, improving the information security of the jump function during use, storage, and transmission.

[0159] This disclosure determines the target instruction corresponding to the instrumentation location specified by the instrumentation task based on the disassembly results, enabling precise location of the instrumentation point and ensuring the accuracy and effectiveness of the instrumentation operation. Furthermore, by determining the identifier information of the jumper function based on the target instruction, the program counter at the instrumentation location, and the insertion function specified by the instrumentation task, effective invocation of the insertion function can be achieved. The jumper function includes at least a call to the insertion function, allowing for the flexible insertion of additional functions or monitoring code without affecting the original program logic. This enhances the program's functionality and debuggability, improves the accuracy and flexibility of the instrumentation operation, and ensures the reliability of the instrumentation process and the correct execution of the insertion function.

[0160] In real-world scenarios, the NVBit solution imposes strong ISA constraints on the target heterogeneous hardware. For example, for hardware types such as VLIW architecture NPUs and NPUs with delay slots in jump instructions, the NVBit code generation scheme may encounter problems. An example is provided below:

[0161] For VLIW architecture NPUs, to achieve instruction parallelism, multiple instructions are packaged into a single very long instruction slot (VLIW) to maximize the compiler's optimization capabilities. Common VLIW instructions are shown below:

[0162] |load|compute|store|

[0163] A very long instruction slot can hold three instructions: load, compute, and store. Typically, these three instructions are independent of each other. After the compiler generates the instructions, the hardware automatically sends them to different processing units for parallel processing.

[0164] For binary instruction instrumentation, whether to jump depends on the predicate of the original instruction. For example, in NVBit's example, STS[R15], R8. If the instruction itself carries a predicate, such as @P1 STS[R15], R8, then the jump after instrumentation should ideally also include the corresponding predicate (@P1 JMP L1). However, for VLIW instructions, if different instruction slots support predicates, such as:

[0165] |@P1 load|@P2 compute|@P3 store|

[0166] For VLIW instructions of this type, whether to perform a jump requires predicate operations, which alters the original instruction segment size and cannot guarantee that the original instruction segment code size required by NVBit remains unchanged. Changes in code size often lead to code execution errors. In this case, a second-best approach is to omit the predicate in JMP L1 itself and then process the predicate in the jump function. This also maintains the original instruction segment code size. However, this approach introduces jump instructions even when the predicate is not executed, increasing instrumentation overhead.

[0167] Generally speaking, VLIW architecture does not allow instructions to have dependencies in its very long instructions, but this is not absolute.

[0168] Some VLIW architectures support instruction slots containing dependent instructions, such as:

[0169] |load|compute0|compute1|store|

[0170] Some architectures allow dependencies between compute0 and compute1. Hardware execution follows the order of compute0 and compute1, ensuring the dependency is maintained if they exist; otherwise, they execute sequentially. However, for such NPUs, when compute0 and compute1 are dependent and the user wants to perform instrumentation after compute0 has finished executing, this involves splitting VLIW instructions, making the NVBit approach unsuitable.

[0171] For heterogeneous NPUs with deferred slots for jump instructions, when a processor encounters a branch instruction, it typically needs time to determine whether the branch will be taken and what the target address is. During this time, the pipeline may stall while waiting for the branch decision, creating a cavitation effect. To address this cavitation, CPUs commonly introduce branch prediction to improve execution efficiency. However, for some NPUs, the hardware overhead of branch prediction can impact maximum computing power, so deferred slots are used to reduce jump cavitation. Deferred slots insert an instruction after the branch instruction; this instruction is always executed, regardless of whether the branch is taken. For NPUs with deferred slots, NVBit's approach is also unsuitable because jump instructions introduce deferred slots and change the size of the original code segment, potentially leading to errors in code execution after instrumentation.

[0172] To address the aforementioned deficiencies, one embodiment of this application further includes: determining the target instruction corresponding to the instrumentation location specified by the instrumentation task based on the disassembly results of the target code; and determining whether the target instruction is a preset type instruction. If the target instruction is not a preset type instruction, the target instruction is replaced with a call instruction to a jump function in the target code to generate a first executable file, wherein the jump function includes at least instructions related to executing the target instruction and a call to the inserted function. If the target instruction is a preset type instruction, a jump instruction pointing to the jump function is inserted into the basic block to which the target instruction belongs, based on the disassembly results of the target code, to generate a first executable file.

[0173] In this embodiment, the preset type instruction can refer to an instruction that, after instrumentation, will cause a change in the size of the target code. For example, the preset type instruction can be a VLIW instruction, a jump instruction with a delay slot on the device side, or a jump instruction that is not supported by the device side (such as hardware-level instructions like CALL / RETURN (function call and return mechanism)). The preset type instruction can be set according to actual needs.

[0174] In this embodiment, for target instructions that do not introduce code segment size expansion, the target instructions can be directly replaced with instructions calling the jump function to generate the first executable file. For preset type instructions that will introduce code segment size expansion, the instructions calling the jump function need to be modified. Instead of directly modifying the instruction list (inst_list) in the target code, the instruction list (bb_inst_list) of the basic block where the target instruction is located is modified. This allows for targeted selection of instrumentation methods based on the target instruction type, improving flexibility and adaptability.

[0175] In real-world scenarios, if instrumenting the target code causes changes in its size, such as code size expansion, it will prevent the target code from executing correctly. For example, if the device-side NPU uses a VLIW architecture and the default instruction type is VLIW, the two situations mentioned in the aforementioned defects may occur:

[0176] Case 1: Only instructions without dependencies are supported between instruction slots.

[0177] When each instruction slot in VLIW only supports dependency-free instructions, different predicate registers may be used for each instruction slot. A typical example is shown below:

[0178] |@P1 load|@P2 compute|@P3 store|

[0179] If it's necessary to determine whether the entire VLIW instruction has been executed, then if it has (indicated by a predicate register value of 1), the counter is incremented; otherwise, the counter remains unchanged. In this case, the AND operation of P1 / P2 / P3 needs to be calculated.

[0180] In this case, assuming the insertion function is the foo function, this application provides two solutions:

[0181] The first approach is to unconditionally replace the corresponding VLIW instruction with a jump to a straddle function at the target address. After the predicate is evaluated in the straddle function, the result is used to determine whether to jump to the instrumented function `foo`. The advantage of this approach is that it doesn't change the size of the original target code segment. The disadvantage is that even if the VLIW instruction isn't executed, a jump to the straddle function will still be introduced, inevitably incurring the overhead of a jump.

[0182] The second approach involves performing predicate computation at the VLIW instruction instead of within the jump function. After the predicate computation, the result is used to determine whether to jump to the jump function. The advantage of this approach is that no additional jump overhead is introduced when the VLIW instruction is not executed. The disadvantage is that the target code size increases after instrumentation. To address this issue, this embodiment of the disclosure can insert jump instructions pointing to the jump function into the basic block containing the target instruction based on the disassembly results of the target code, generating a first executable file. This approach allows for targeted selection of instrumentation methods based on the target instruction type, and modification of the instrumentation process for preset instruction types, improving flexibility and adaptability.

[0183] In one embodiment, if VLIW supports dependencies between instruction slots, and when the device-side architecture supports VLIW and dependencies between instruction slots within a single very long instruction are allowed, the following applies:

[0184] |load|compute0|compute1|store|

[0185] The device-side architecture may allow dependencies between load, compute0, compute1, and store. When dependencies exist, the target hardware will execute them in the order of dependency. A typical example is shown below:

[0186] |ld.gpr R1,[addr]|add.s32 R1,R1,R1|add.s32 R1,R1,4|st.gpr[addr],R1|

[0187] At this point, the hardware will automatically ensure the dependency order between instruction slots. For the VLIW instructions mentioned above, the user might need to use the value of R1 as input to the instrumentation function for custom operations after executing add.s32 R1,R1,R1. This requires inserting a jump after add.s32, which involves splitting the VLIW instructions. Assuming that CALL in the target architecture must occupy the entire VLIW instruction slot, then the following instruction sequence needs to be generated:

[0188] |ld.gpr R1,[addr]|add.s32 R1,R1,R1|nop|nop|

[0189] call trampoline_tid

[0190] |nop|nop|add.s32 R1,R1,4|st.gpr[addr],R1|

[0191] This will inevitably lead to an increase in the size of the target code after instrumentation. To solve the problem of increased target code size after instrumentation, the embodiments of this disclosure can insert jump instructions pointing to the jump function in the basic block to which the target instruction belongs, based on the disassembly results of the target code, to generate a first executable file, thereby realizing the targeted selection of instrumentation methods according to the type of target instruction, improving flexibility and adaptability.

[0192] In one embodiment, if the device-side jump has a delay slot, that is, the preset type instruction is the corresponding jump instruction with a delay slot. In actual scenarios, some devices only support jumps with delay slots, meaning that after the jump, there must be a delay slot of several bytes (bits). Regardless of whether the jump reaches a specified target address, the instruction in the delay slot will be executed. An example of the target code is as follows:

[0193] memcpy

[0194] sadd.s32

[0195] memcpy

[0196] `memcpy` is a function provided by the device-side code, used to copy a specified number of bytes from a source address to a specified destination address. It is typically used for copying memory blocks. The definition of `memcpy` is in the header file. If the target instruction `sadd.s32` is replaced with a function call with a delay slot (`call` can be considered a special type of jump), the generated instruction sequence is as follows:

[0197] memcpy

[0198] call trampoline_tid

[0199] nop

[0200] nop

[0201] memcpy

[0202] The `nop` (No Operation) instruction is used to fill the delay slots of function calls. When the `nop` instruction is executed, the processor does not perform any operation; it simply consumes one or more clock cycles and then continues to execute the next instruction. It is evident that the existence of delay slots inevitably leads to changes in the size of the target code after instrumentation. To ensure accurate execution of the instrumented code, jump instructions pointing to the jumper function can be inserted into the basic block to which the target instruction belongs, based on the disassembly results of the target code, generating the first executable file. This embodiment allows for targeted selection of instrumentation methods according to the type of target instruction, with minimal constraints on the device side. For example, it is applicable to architectures that support hardware features such as VLIW and delay slots, improving flexibility and adaptability.

[0203] In one embodiment, the device side may not support hardware instruction-level call / return, but instead implements the corresponding functionality through special jumps. For example, the jump instruction contains a specific field to indicate whether to update linkreg (LinkRegister stores the return address of a subroutine call). When the flag for updating linkreg is enabled, if a jump actually occurs, the PC of the next instruction after the jump instruction will be updated in linkreg. In this case, at the beginning of the insertion function foo, linkreg is first read and stored on the stack. After the insertion function finishes execution, the saved address is read from the stack and used as the jump destination address to achieve the return function. In this situation, a specific general-purpose register is usually needed to assist in implementing the function call functionality. The instruction generation for the function call to the insertion function foo is shown below:

[0204] read.linkreg R1,LINKREG

[0205] st.gpr R1[addr]

[0206] / / Function body implementation ...

[0208] ld.gpr R1[addr]

[0209] jmp[R1] (This jumps without updating the linkreg)

[0210] In the example above, R1 is used as a special register to assist in implementing function calls. At this time, the target instruction that needs instrumentation may also use R1. For example, if a function needs to be inserted before `add.s32 R1,R1,R1` (target instruction), the target instruction cannot be moved to the jump function for execution; a call to the jump function must be inserted before the target instruction in the target code. This will increase the size of the target code after instrumentation. To ensure accurate execution of the instrumented code, a jump instruction pointing to the jump function can be inserted into the basic block to which the target instruction belongs, based on the disassembly results of the target code, generating the first executable file. This embodiment of the disclosure implements targeted selection of instrumentation methods according to the type of target instruction, improving flexibility and adaptability.

[0211] In one embodiment, the method described above, which inserts jump instructions pointing to a jump function into the basic block to which the target instructions belong based on the disassembly results of the target code, and generates a first executable file, may include the following steps:

[0212] 1. Based on the disassembly results of the target code, determine the mapping relationship between instructions and basic blocks in the target code;

[0213] 2. Determine the target basic block to which the target instruction belongs based on the mapping relationship;

[0214] 3. Determine the target instruction list corresponding to the target basic block based on the mapping relationship;

[0215] 4. Insert jump instructions pointing to the jump function into the target instruction list to obtain the modified instruction list of the target basic block;

[0216] 5. Traverse the list of basic blocks in the target code, update the program counter corresponding to each basic block according to the modified instruction list corresponding to each basic block, and generate the first executable file.

[0217] In one embodiment, after traversing the list of basic blocks in the target code and updating the program counter corresponding to each basic block based on the modified instruction list corresponding to each basic block, the method further includes: traversing the list of basic blocks in the target code, and for basic blocks ending with jump instructions based on instruction offsets, recalculating the corresponding instruction offsets based on the updated program counters. In this embodiment, offset correction can be further performed on basic blocks ending with jump instructions based on instruction offsets to ensure the accuracy of instruction offsets. Here, a basic block refers to a sequentially executed sequence of statements, with only one entry point and one exit point. The entry point is the first statement, and the exit point is the last statement. For a basic block, execution only enters from its entry point and exits from its exit point.

[0218] As previously mentioned, there may be unavoidable issues with the size of the target code after instrumentation on the device side. In this case, it is necessary to modify the jump instructions in the user code segment to ensure that the instrumented code can be executed accurately.

[0219] First, the disassembler interface can be called to parse the instruction segments of the target code and obtain the disassembled result. The disassembled result contains, but is not limited to, the following information:

[0220] cfg: Control Flow Graph, is a directed graph that includes multiple nodes and edges. Each node represents a basic block, and edges represent the control flow between basic blocks. cfg is used to represent the execution path of a program, aiding the compiler in optimization and analysis, and recording the interconnection information between all basic blocks.

[0221] bb_list: Basic Block List, a list of basic blocks obtained after disassembling the instruction segment of the object code. It can be sorted from low to high by the PC of the instruction at the beginning of the basic block.

[0222] inst_list: Instruction List, a list of instructions obtained after disassembling the instruction segment of the target code, sorted from low to high according to the PC of the instruction.

[0223] inst2bb: A mapping table of instructions to basic blocks in the object code obtained after disassembly, used to represent the mapping relationship between instructions and basic blocks in the object code.

[0224] Each basic block also maintains its own `bb_inst_list` to record the instruction list within that block. When instrumenting after determining the required instrumentation location, the `inst_list` is not directly modified; instead, the following steps are performed:

[0225] (1) Determine the target instruction inst(list) corresponding to the stub position.

[0226] (2) Determine the target basic block where each target instruction is located by inst2bb, and then obtain the target instruction list bb_inst_list corresponding to the target basic block.

[0227] (3) Modify the target instruction list bb_inst_list. Add jump instructions pointing to the jump function in the target instruction list bb_inst_list according to user needs, or adjust the order of instructions.

[0228] (4) Traverse the list of basic blocks bb_list in the target code, update the PC corresponding to each basic block according to the modified bb_inst_list in each basic block, and generate the first executable file.

[0229] (5) It can iterate through bb_list. For jump instructions at the end of bb_inst_list that jump based on instruction offsets, the updated PC information of the target basic block of the jump instruction can be obtained through the control flow graph cfg of the target code, and the corresponding instruction offset can be recalculated. In addition, for jumps corresponding to function calls, since there are relocation entries, it is not necessary to update the offset.

[0230] The embodiments disclosed herein can correct the changes in code volume after instrumentation, ensuring that the instrumented code can be executed accurately, thereby improving flexibility and adaptability.

[0231] This embodiment of the present disclosure thus completes the correction of offset-based jump instructions within the instruction segment caused by the expansion of the code segment size. It is worth noting that, since the target code is modified in this embodiment before runtime relocation during Module Load, it is only necessary to ensure that the target code is generated correctly, and then proceed with the normal runtime relocation and execution process.

[0232] In one embodiment, if the target instruction is a preset type instruction, after inserting a jump instruction pointing to a jump function into the basic block to which the target instruction belongs based on the disassembly result of the target code to generate the first executable file, the method further includes: if the device side has already allocated memory for the target code, releasing the allocated memory and reallocating memory according to the instrumented code.

[0233] In this embodiment, the target code may have already been allocated memory on the device side. However, due to the expansion of the code size after instrumentation to generate the first executable file, the originally allocated memory may not be sufficient to meet the demand. Therefore, the original memory can be released first, and then the modified binary can be relocated at runtime, with memory reallocated for the target code. This ensures that modifications to the target binary code are supported at any time. For example, device-side malloc / free ("malloc" and "free" are the NPU device-side memory malloc / free functions, used for dynamic memory allocation and release, respectively) can be used to update the memory.

[0234] The state saving / restoration functions mentioned in the aforementioned prior art are pre-written device-side functions based on different resource usage levels and compiled into a static library. This is indeed a good approach when there are few types of CPU resources. However, when the target NPU has a wide variety of physical resources, combinatorial explosion may occur, resulting in an excessive number of functions that need to be generated.

[0235] In this embodiment, a suitable code generation method can be selected for different scenarios. For example, if the target NPU has few resource types, the above-mentioned solution of pre-writing different device-side functions based on resource usage and compiling them into a static library can be used. However, if the NPU has too many physical register resource types, the JIT generation method of this application embodiment can be used to generate the instrumented code. This achieves flexible support for different hardware architectures.

[0236] In one embodiment, step 504 may specifically include: determining the set of physical registers used by the insertion function based on the implementation information of the insertion function; generating save code and restore code based on the set of physical registers used by the insertion function; and completing the jumper function in the second executable file based on the save code and restore code to generate instrumented code.

[0237] In this embodiment, assuming the inserted function is the foo function, if a call to the function foo needs to be inserted before a certain target instruction, first determine the set of physical registers R1 / R2 / R3 that will be used in foo, then generate the save code and restore code of physical resources according to the set of physical registers R1 / R2 / R3 used, complete the jump board function based on the save and restore code, generate the instrumented code, and optimize the instrumentation resource overhead.

[0238] In an optional embodiment, if a call to function foo needs to be inserted before a certain instruction, and foo uses the physical register set R1 / R2 / R3, R1 / R2 / R3 can be saved and restored.

[0239] In the aforementioned embodiments, if R1 / R2 / R3 are all saved and restored, but the instrumentation point in the user code only uses R1, then saving and restoring R2 / R3 is redundant and will still lead to redundancy overhead. To address this issue, the embodiments of this application make further improvements by performing an activity analysis on the physical register set; based on the activity analysis results, the target physical register used by the instrumentation position specified by the instrumentation task is determined from the physical register set, thereby generating save and restore code for the target physical register based on the implementation information of the insertion function.

[0240] In this embodiment, the initial instrumentation time overhead optimization has already been implemented in the previous embodiments. This is reflected in the fact that during state saving / restoration, only the necessary physical register resources are saved based on the physical resource usage of the target instrumentation function, thereby reducing instrumentation overhead. In this embodiment, further overhead optimization can be performed based on this optimization.

[0241] Suppose a call to function `foo` needs to be inserted before a certain instruction, and `foo` uses physical register sets R1 / R2 / R3. Before generating code to save and restore physical resources, an activity analysis of the physical register resources can be performed. Then, the resource utilization of the target instrumented function can be considered to save and restore the physical resources. If the activity analysis finds that instrumentation points R2 / R3 are not active, meaning that the instrumentation point in the user code only uses R1, then only R1 will be saved and restored, reducing runtime overhead.

[0242] Please refer to Figure 6 This is a schematic diagram illustrating a specific scenario of a binary code instrumentation method according to an embodiment of this application. This embodiment aims to provide a solution for heterogeneous computing where the device-side software stack is incomplete (e.g., lacking a function library loading mechanism). The dynamic binary instrumentation tool Core may not include a tool function loader. After the user specifies the instrumentation logic in the instrumentation task, in response to the instrumentation task, the target code is obtained. Specifically, a source file (origin.CNBin) capable of independently running target code can be generated during the target code generation stage. origin.CNBin may include a kernel function file (.text.kernel), for example, it may include the following target code:

[0243] memcpy

[0244] sadd.s32

[0245] memcpy

[0246] The origin.CNBin file is in binary format. The code generator first modifies the original origin.CNBin. Assume the target user wants to insert a call to the insertion function foo before the target instruction sadd.s32. The entire process can be divided into the following three stages:

[0247] Phase 1: Constructing the minimum CNBin capable of correct linking. The first executable file, Phase1.cnbin, is generated based on origin.CNBin. Phase1.cnbin may include the kernel function file (.text.kernel), .text.trampoline_tid (trampoline function identifier file), and other related files. In this first phase, a trampoline function (corresponding to trampoline_tid) is generated, and the original target instruction sadd.s32 is replaced with the call instruction calltrampoline_tid. Here, tid represents the ID of the corresponding trampoline function, determined by the target code at the insertion point, the PC at the insertion point, and the target function foo to be called. In this phase, the trampoline function can only insert instructions related to function calls, execution of target instructions, and return operations. At this time, the trampoline function may include the following:

[0248] / / Call the insert function

[0249] Call foo

[0250] / / Execute instruction + return

[0251] sadd.s32

[0252] return

[0253] It should be noted that since the operations of saving the context, pushing parameters onto the stack, and restoring the context depend on the implementation information of the target function foo, but at this time foo exists in phase1.cnbin as an extern function (a function that accesses files across files), its function implementation is unknown. Therefore, it is not necessary to insert instructions for saving the context, pushing parameters onto the stack, and restoring the context.

[0254] On the other hand, it is not appropriate to parse cnsanitizer_tool.so at this time, because the tool library may contain multiple implementations of the foo function with the weak attribute (a weak attribute used to indicate a weak reference or weak binding), and which one to choose depends entirely on the implementation of the JIT linker. Therefore, in this stepping stone function, it is only necessary to generate function calls that allow the linker to work properly.

[0255] The linker is an important tool responsible for linking compiled object files and library files together to generate the final executable file or library file. During this process, the linker resolves symbols, such as function names and variable names, and translates their references into actual memory addresses.

[0256] The second stage is the Just-In-Time (JIT) linking stage. The implementation of the function `foo` is inserted into the pre-defined dynamic library `cnsanitizer_tool.so`. In this second stage, the JIT linker is invoked to link the first executable file `phase1.cnbin` and the pre-defined dynamic library `cnsanitizer_tool.so`. After linking, a second executable file `phase2.cnbin` is generated, which includes the implementation information of `foo` in `.text.foo` (the `foo` function file).

[0257] Phase 3: Completing the jump function. Based on the implementation information of the linked `foo` function and the resource utilization information (such as register resource utilization) recorded in its `config` (configuration) information, the jump function generates save code for saving the current state, parameter push code for pushing parameters onto the stack, and restore code for restoring the current state. This generates the final instrumented code `phase3.cnbin`. The jump function in `phase3.cnbin` can contain the following:

[0258] / / Save the scene

[0259] ...

[0260] / / Parameters are pushed onto the stack

[0261] ...

[0262] / / Call the insert function

[0263] Call foo

[0264] / / Restore the scene

[0265] ...

[0266] / / Execute instruction + return

[0267] sadd.s32

[0268] return

[0269] The above embodiment is just a typical code generation example. In actual scenarios, code generation can be adjusted according to different target NPU architectures.

[0270] In one embodiment, for an instrumentation task scenario that requires obtaining the specific number of instructions executed, the solution provided by the embodiments of this application can be implemented through the following code example:

[0271] __managed__long counter=0;

[0272] extern "C"__device____noinline__void incr_counter(){

[0273] atomicAdd(&counter,1);

[0274] }CNSANITIZER_EXPORT_DEV_FUNC(incr_counter);

[0275] void cnsanitizer_patch_text(ELFText*text){

[0276] for(auto&i:cnsanitizer_get_instrs(text)){

[0277] cnsanitizer_insert_call(i,"incr_counter",IPOINT_BEFORE);

[0278] }

[0279] }

[0280] void cnsanitizer_at_exit(){

[0281] cout<<“Total instructions“< <counter<<“\n”;

[0282] }

[0283] Here, incr_counter() is the specified insertion function, and the meaning of the above code is as follows:

[0284] Declare a managed variable `counter` that can be shared between the host and the device, and initialize it to 0.

[0285] Declare a device function `incr_counter`, specify C language linking using `extern "C"`, and prevent the compiler from inlining this function using `__noinline__`.

[0286] The function body uses `atomicAdd(&counter,1)` to perform an atomic addition operation, incrementing `counter` by 1.

[0287] Export the device function `incr_counter` so that it can be used by other files.

[0288] Declare a function to insert a call to `incr_counter` before each instruction.

[0289] Returns a list of instructions, which can be iterated over.

[0290] Insert a call to `incr_counter` before each instruction `i`.

[0291] Declare a function that is called when the program exits.

[0292] Use `cout` to output the total number of instructions `counter`.

[0293] As can be seen from this example, compared to the instrumentation task example based on NVBit implemented in the previous embodiment to obtain the specific number of instructions to be executed, this embodiment shields the user from the details related to the Driver API, allowing the user to focus more on the instrumentation logic itself, without needing to have too much understanding of the abstraction of the Driver API layer.

[0294] For details of each step of the above method, please refer to the relevant descriptions in the above embodiments, which will not be repeated here.

[0295] It should be noted that the above embodiments can provide an approach to support dynamic binary instrumentation for NPUs and corresponding software stacks that do not meet the prerequisites for instrumentation schemes based on NVBit. This does not mean that the functions supported by NVBit are not applicable in this embodiment. If the target hardware supports instrumentation without modifying the target code instruction segment, or the target NPU's software stack supports a dynamic library loading mechanism, or the tool library user wants to be aware of the details of the Driver API and wants to perform corresponding callback operations in different Driver APIs, the solutions in this application can all provide corresponding support.

[0296] In one possible implementation, a chip is also disclosed that includes the dynamic binary instrumentation device provided in any of the above embodiments.

[0297] In one possible implementation, a board is also disclosed. Figure 7 This is a structural diagram of a board according to an embodiment of this application, such as... Figure 7As shown, board 500 includes chip 501, which is a system-on-chip (SoC) for artificial intelligence. It integrates one or more combined processing units, which are artificial intelligence computing units used to support various deep learning and machine learning algorithms, meeting the intelligent processing needs of complex scenarios in fields such as computer vision, speech, natural language processing, and data mining. In particular, deep learning technology is widely used in cloud intelligence. A significant characteristic of cloud intelligence applications is the large volume of input data, placing high demands on the platform's storage and computing capabilities. Board 500 in this embodiment is suitable for cloud intelligence applications, possessing massive off-chip storage, on-chip storage, and substantial computing power.

[0298] Chip 501 is connected to external device 503 via external interface device 502. External device 503 may be, for example, a server, computer, camera, monitor, mouse, keyboard, network card, or Wi-Fi interface. Optionally, in this embodiment, it may be various vision sensors. Data to be processed can be transmitted from external device 503 to chip 501 via external interface device 502. The calculation results of chip 501 can be transmitted back to external device 503 via external interface device 502. Depending on the application scenario, external interface device 502 may have different interface forms, such as a PCIe interface.

[0299] The board 500 also includes a storage device 504 for storing data, which includes one or more memory cells 505. The storage device 504 is connected to and transmits data with the controller 506 and the artificial intelligence chip 501 via a bus. The controller 506 in the board 500 is configured to regulate the state of the artificial intelligence chip 501. Therefore, in one application scenario, the controller 506 may include a microcontroller (MCU).

[0300] In one possible implementation, a combined processing device is also provided. Figure 8 A structural diagram of the combined processing device in chip 501 of this embodiment is shown. Figure 8 As shown, the combined processing device 60 includes a computing device 601, an interface device 602, a processing device 603, and a storage device 604.

[0301] The computing device 601 is configured to perform user-specified operations. It is mainly implemented as a single-core intelligent processor or a multi-core intelligent processor to perform deep learning or machine learning calculations. It can interact with the processing device 603 through the interface device 602 to jointly complete the user-specified operations.

[0302] Interface device 602 is used to transmit data and control commands between computing device 601 and processing device 603. For example, computing device 601 can obtain input data from processing device 603 via interface device 602 and write it to on-chip storage device of computing device 601. Further, computing device 601 can obtain control commands from processing device 603 via interface device 602 and write them to on-chip control cache of computing device 601. Alternatively or optionally, interface device 602 can also read data from storage device of computing device 601 and transmit it to processing device 603.

[0303] The processing device 603, as a general-purpose processing device, performs basic controls including but not limited to data transfer and starting / stopping the computing device 601. Depending on the implementation, the processing device 603 may be one or more types of processors, such as a central processing unit (CPU), a graphics processing unit (GPU), or other general-purpose and / or special-purpose processors. These processors include, but are not limited to, digital signal processors (DSPs), application-specific integrated circuits (ASICs), field-programmable gate arrays (FPGAs), or other programmable logic devices, discrete gate or transistor logic devices, discrete hardware components, etc., and their number can be determined according to actual needs. As mentioned above, the computing device 601 of this application can be considered as having a single-core structure or a homogeneous multi-core structure. However, when the computing device 601 and the processing device 603 are considered together, they are considered to form a heterogeneous multi-core structure.

[0304] Storage device 604 is used to store data to be processed. It may be DRAM 604, which is DDR memory, typically 16G or larger in size, and is used to store data of computing device 601 and / or processing device 603.

[0305] Figure 9 The diagram below shows the internal structure of a single-core computing device. The single-core computing device 700 is used to process input data for computer vision, speech, natural language processing, data mining, etc. The single-core computing device 700 includes three main modules: a control module 701, a processing module 702, and a storage module 703.

[0306] The control module 701 coordinates and controls the operation of the computation module 702 and the storage module 703 to complete the deep learning task. It includes an instruction fetch unit (IFU) 7011 and an instruction decode unit (IDU) 7012. The instruction fetch unit 7011 fetches instructions from the processing device, while the instruction decode unit 7012 decodes the fetched instructions and sends the decoding result as control information to the computation module 702 and the storage module 703.

[0307] The computation module 702 includes a vector operation unit 7021 and a matrix operation unit 7022. The vector operation unit 7021 is used to perform vector operations and can support complex operations such as vector multiplication, addition, and nonlinear transformations. The matrix operation unit 7022 is responsible for the core computations of deep learning algorithms, namely matrix multiplication and convolution.

[0308] Storage module 703 is used to store or move relevant data, including neuron RAM (NRAM) 7031, weight RAM (WRAM) 7032, and direct memory access (DMA) module 7033. NRAM 7031 stores input neurons, output neurons, and intermediate results. WRAM 7032 stores the convolution kernels, i.e., weights, of the deep learning network. DMA 7033 is connected to DRAM 604 via bus 704 and is responsible for data transfer between the single-core computing device 700 and DRAM 604.

[0309] Figure 10 The diagram shows the internal structure of a multi-core computing device. The multi-core computing device 800 adopts a hierarchical design. As a system-on-a-chip, the multi-core computing device 800 includes at least one cluster, and each cluster includes multiple processor cores. In other words, the multi-core computing device 800 is constructed in a hierarchical structure of system-on-a-chip, cluster, and processor cores.

[0310] From the perspective of system-on-a-chip hierarchy, such as Figure 10 As shown, the multi-core computing device 800 includes an external storage controller 801, a peripheral communication module 802, an on-chip interconnect module 803, a synchronization module 804, and multiple clusters 805.

[0311] There can be multiple external storage controllers 801; two are shown as an example in the figure. These controllers are used to respond to access requests from the processor core to access external storage devices, such as… Figure 10The DRAM 604 in the chip allows data to be read from or written to external sources. The peripheral communication module 802 receives control signals from the processing device via the interface device 602 to initiate tasks. The on-chip interconnect module 803 connects the external storage controller 801, the peripheral communication module 802, and multiple clusters 805 to transmit data and control signals between modules. The synchronization module 804 is a global barrier controller (GBC) used to coordinate the working progress of each cluster and ensure information synchronization. The multiple clusters 805 are the computing cores of the multi-core computing device 800. Four are shown exemplary in the figure; however, with hardware development, the multi-core computing device 800 of this application can also include 8, 16, 64, or even more clusters 805. The clusters 805 are used to efficiently execute deep learning algorithms.

[0312] From the perspective of cluster hierarchy, such as Figure 10 As shown, each cluster 805 includes multiple processor cores (IPU cores) 806 and one memory core (MEM core) 807. For example, each cluster 805 includes four processor cores and one memory, which can be DRAM 604.

[0313] Four processor cores 806 are shown in the figure as an example, but this application does not limit the number of processor cores 806. Each processor core 806 is similar to Figure 11 The single-core computing device also includes three main modules: a control module 901, an arithmetic module 902, and a storage module 903. The functions and structures of the control module 901, arithmetic module 902, and storage module 903 are largely the same as those of the control module 701, arithmetic module 702, and storage module 703. The control module 901 includes an instruction fetch unit 9011 and an instruction decode unit 9012. The arithmetic module 902 includes a vector operation unit 9021 and a matrix operation unit 9022. Further details are omitted. It should be noted that the storage module 903 includes an input / output direct memory access (IODMA) module 9033 and a move direct memory access (MVDMA) module 9034. The IODMA 9033 controls memory access between the NRAM 9031 / WRAM 9032 and the DRAM 604 via the broadcast bus 809. The MVDMA 9034 is used to control memory access to the NRAM 9031 / WRAM 9032 and the SRAM 808.

[0314] Back Figure 10The storage core 807 is primarily used for storage and communication, namely storing shared data or intermediate results among processor cores 806, and performing communication between cluster 805 and DRAM 604, communication between clusters 805, and communication between processor cores 806. In other embodiments, the storage core 807 has scalar operation capabilities and is used to perform scalar operations.

[0315] The storage core 807 includes an SRAM 808, a broadcast bus 809, a cluster direct memory access (CDMA) module 810, and a global direct memory access (GDMA) module 811. The SRAM 808 acts as a high-performance data relay station. Data multiplexed between different processor cores 806 within the same cluster 805 does not need to be obtained from the DRAM 604 by each processor core 806. Instead, it is relayed between processor cores 806 via the SRAM 808. The storage core 807 only needs to quickly distribute the multiplexed data from the SRAM 808 to multiple processor cores 806 to improve inter-core communication efficiency and greatly reduce on-chip and off-chip input / output access.

[0316] Broadcast bus 809, CDMA 810, and GDMA 811 are used to perform communication between processor cores 806, communication between clusters 805, and data transfer between cluster 805 and DRAM 604, respectively. These will be explained below.

[0317] The broadcast bus 809 is used to complete high-speed communication between the processor cores 806 within the cluster 805. In this embodiment, the broadcast bus 809 supports inter-core communication methods including unicast, multicast, and broadcast. Unicast refers to point-to-point (e.g., single processor core to single processor core) data transmission. Multicast is a communication method that transmits a piece of data from SRAM 808 to several specific processor cores 806. Broadcast is a communication method that transmits a piece of data from SRAM 808 to all processor cores 806, and is a special case of multicast.

[0318] CDMA 810 is used to control SRAM 808 access between different clusters 805 within the same computing device.

[0319] The GDMA 811 works in conjunction with the external memory controller 801 to control memory access from the SRAM 808 to the DRAM 604 in the cluster 805, or to read data from the DRAM 604 into the SRAM 808. As previously mentioned, communication between the DRAM 604 and the NRAM 9031 or WRAM 9032 can be achieved through two channels. The first channel is a direct connection between the DRAM 604 and the NRAM 9031 or WRAM 9032 via the IODAM 9033. The second channel involves first transferring data between the DRAM 604 and SRAM 808 via the GDMA 811, and then transferring data between the SRAM 808 and the NRAM 9031 or WRAM 9032 via the MVDMA 9034. Although the second channel appears to require more components and has a longer data flow, in some embodiments, the bandwidth of the second channel is actually much greater than that of the first channel. Therefore, communication between the DRAM 604 and the NRAM 9031 or WRAM 9032 may be more efficient through the second channel. Embodiments of this application may select the data transmission channel based on their own hardware capabilities.

[0320] In other embodiments, the functions of GDMA 811 and IODMA 9033 can be integrated into the same component. For ease of description, this application treats GDMA 811 and IODMA 9033 as different components. For those skilled in the art, any component that implements functions and achieves technical effects similar to this application falls within the scope of protection of this application. Furthermore, the functions of GDMA 811, IODMA 9033, CDMA 810, and MVDMA 9034 can also be implemented by the same component.

[0321] This application also provides a computer-readable storage medium storing a computer program, which, when executed by at least one processor, implements the method provided in any of the above embodiments.

[0322] This application also provides a computer program product, including a computer program that, when executed by a processor, implements the method of any of the foregoing embodiments.

[0323] The foregoing may be better understood in view of the following clauses:

[0324] Clause 1. A binary code instrumentation method, comprising:

[0325] Obtain the target code and its corresponding jump function.

[0326] Insert jump instructions pointing to the jump function into the target code to generate the first executable file.

[0327] The first executable file and the preset dynamic library are linked to generate the second executable file, wherein the preset dynamic library includes the implementation information of the inserted functions.

[0328] Based on the implementation information of the insertion function, complete the jumper function in the second executable file and generate the instrumented code.

[0329] Clause 2. Obtain the target code according to the method in Clause 1, including:

[0330] In response to the instrumentation task, the target code is hijacked when the target code's module is loaded, but before relocation.

[0331] Clause 3, the method pursuant to Clause 1, also includes:

[0332] The pre-packaged disassembler interface is called to disassemble the target code and obtain the disassembled result of the target code.

[0333] Clause 4. Obtain the jump table function corresponding to the target code according to the method in Clause 3, including:

[0334] Based on the disassembly results, determine the target instruction corresponding to the instrumentation location specified by the instrumentation task.

[0335] Based on the target instruction, the program counter where the instrumentation location is located, and the insertion function specified by the instrumentation task, determine the identification information of the jumper function. The jumper function includes at least a call to the insertion function.

[0336] Clause 5. Following the method in Clause 3, insert jump instructions pointing to the jump function into the object code to generate the first executable file, including:

[0337] Based on the disassembly results of the target code, determine the target instruction corresponding to the instrumentation position specified by the instrumentation task.

[0338] Determine whether the target instruction is a preset type instruction.

[0339] If the target instruction is not a preset type instruction, the target instruction is replaced with a call instruction to the jump function in the target code to generate the first executable file. The jump function includes at least the relevant instructions for executing the target instruction and a call to the inserted function.

[0340] If the target instruction is a preset type instruction, based on the disassembly result of the target code, a jump instruction pointing to the jump function is inserted into the basic block to which the target instruction belongs, and the first executable file is generated.

[0341] Clause 6. Following the method in Clause 5, based on the disassembly results of the target code, insert jump instructions pointing to the jump function into the basic block to which the target instructions belong, generating the first executable file, including:

[0342] Based on the disassembly results of the target code, determine the mapping relationship between instructions and basic blocks in the target code.

[0343] The target basic block to which the target instruction belongs is determined based on the mapping relationship.

[0344] The target instruction list corresponding to the target basic block is determined based on the mapping relationship.

[0345] Insert jump instructions pointing to the jump function into the target instruction list to obtain the modified instruction list of the target basic block.

[0346] Iterate through the list of basic blocks in the target code, update the program counter corresponding to each basic block based on the list of modified instructions corresponding to each basic block, and generate the first executable file.

[0347] Clause 7, following the method in Clause 6, after traversing the list of basic blocks in the target code and updating the program counter corresponding to each basic block based on the modified instruction list corresponding to each basic block, also includes:

[0348] Iterate through the list of basic blocks in the target code. For basic blocks that end with jump instructions based on instruction offsets, recalculate the corresponding instruction offsets based on the updated program counter.

[0349] Clause 8. According to any one of Clauses 5-7, if the target instruction is a preset type instruction, after inserting a jump instruction pointing to a jump function into the basic block to which the target instruction belongs based on the disassembly result of the target code, and generating the first executable file, the following shall also be included:

[0350] If memory has already been allocated for the target code on the device side, release the allocated memory and reallocate memory according to the instrumented code.

[0351] Clause 9. Following the method described in any one of Clauses 1-8, and based on the implementation information of the inserted function, complete the jumper function in the second executable file to generate the instrumented code, including:

[0352] Based on the implementation information of the insertion function, determine the set of physical registers used by the insertion function.

[0353] Save and restore codes are generated based on the set of physical registers used by the insert function.

[0354] Based on the saved and restored code, the springboard function is completed in the second executable file, generating the instrumented code.

[0355] Clause 10, pursuant to the method of Clause 9, generates save and restore codes based on the set of physical registers used by the insert function, further comprising:

[0356] Perform active analysis on the physical register set.

[0357] Based on the active analysis results, the target physical registers used by the instrumentation location specified by the instrumentation task are determined from the set of physical registers.

[0358] Generate save and restore code roots for the target physical register based on the implementation information of the insertion function.

[0359] Clause 11. A binary code instrumentation framework system, comprising:

[0360] The driver middleware layer is used to obtain the target code.

[0361] The code generator calls a pre-encapsulated instrumentation interface to obtain the jump function corresponding to the target code. It inserts jump instructions pointing to the jump function into the target code, generating a first executable file. The first executable file is linked with a pre-defined dynamic library, generating a second executable file, where the pre-defined dynamic library includes the implementation information of the inserted function. Based on the implementation information of the inserted function, the jump function is completed in the second executable file, generating the instrumented code.

[0362] Clause 12. The system pursuant to Clause 11 also includes:

[0363] The disassembler is used to call a pre-packaged disassembler interface to disassemble the target code and obtain the disassembled result of the target code, so that the code generator can generate instrumented code based on the disassembled result.

[0364] Clause 13. An electronic device comprising:

[0365] At least one processor. And

[0366] A memory that is communicatively connected to at least one processor.

[0367] The memory stores instructions that can be executed by at least one processor to cause the electronic device to perform any one of the methods in Clauses 1-10.

[0368] Clause 14. A computer-readable storage medium storing computer-executable instructions that, when executed by a processor, implement the method of any one of Clauses 1-10.

[0369] Clause 15. A computer program product comprising a computer program that, when executed by a processor, implements the method as described in any one of Clauses 1-10.

[0370] It should be noted that, for the sake of simplicity, the foregoing method embodiments are all described as a series of actions. However, those skilled in the art should understand that this application is not limited to the described order of actions, as some steps may be performed in other orders or simultaneously according to this application. Furthermore, those skilled in the art should also understand that the embodiments described in the specification are all optional embodiments, and the actions and modules involved are not necessarily essential to this application.

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

[0372] It should be understood that the above-described device embodiments are merely illustrative, and the device of this application can also be implemented in other ways. For example, the division of units / modules in the above embodiments is only a logical functional division, and there may be other division methods in actual implementation. For example, multiple units, modules, or components may be combined, or integrated into another system, or some features may be ignored or not executed.

[0373] Furthermore, unless otherwise specified, the functional units / modules in the various embodiments of this application can be integrated into one unit / module, or each unit / module can exist physically separately, or two or more units / modules can be integrated together. The integrated units / modules described above can be implemented in hardware or as software program modules.

[0374] When integrated units / modules are implemented in hardware, the hardware can be digital circuits, analog circuits, etc. The physical implementation of the hardware structure includes, but is not limited to, transistors, memristors, etc. Unless otherwise specified, a neural network processor can be any suitable hardware processor, such as a CPU, GPU, FPGA, DSP, and ASIC, etc. Unless otherwise specified, storage units can be any suitable magnetic or magneto-optical storage medium, such as resistive random access memory (RRAM), dynamic random access memory (DRAM), static random access memory (SRAM), enhanced dynamic random access memory (EDRAM), high-bandwidth memory (HBM), hybrid memory cube (HMC), etc.

[0375] If the integrated unit / module is implemented as a software program module and sold or used as an independent product, it can be stored in a computer-readable storage device (CMD). Based on this understanding, the technical solution of this application, in essence, or the part that contributes to the prior art, or all or part of the technical solution, can be embodied in the form of a software product. This computer software product is stored in a memory and includes several instructions to cause a computer device (which may be a personal computer, server, or network device, etc.) to execute all or part of the steps of the methods of the various embodiments of this application. The aforementioned memory includes various media capable of storing program code, such as a USB flash drive, read-only memory (ROM), random access memory (RAM), portable hard drive, magnetic disk, or optical disk.

[0376] In the above embodiments, the descriptions of each embodiment have their own emphasis. For parts not described in detail in a certain embodiment, please refer to the relevant descriptions of other embodiments. The technical features of the above embodiments can be combined arbitrarily. For the sake of brevity, not all possible combinations of the technical features in the above embodiments are described. However, as long as the combination of these technical features does not contradict each other, it should be considered within the scope of this specification.

[0377] The embodiments of this disclosure have been described in detail above. Specific examples have been used to illustrate the principles and implementation methods of this disclosure. The descriptions of the embodiments above are only for the purpose of helping to understand the methods and core ideas of this disclosure. Furthermore, any changes or modifications made by those skilled in the art based on the ideas of this disclosure, and on the specific implementation methods and application scope of this disclosure, are all within the scope of protection of this disclosure. Therefore, the content of this specification should not be construed as a limitation of this disclosure.

Claims

1. A binary code instrumentation method, characterized in that, include: Obtain the target code and the corresponding jumper function; Insert a jump instruction pointing to the springboard function into the target code to generate the first executable file; The first executable file and the preset dynamic library are linked to generate a second executable file, wherein the preset dynamic library includes the implementation information of the inserted function; Based on the implementation information of the insertion function, the springboard function is completed in the second executable file to generate the instrumented code.

2. The method according to claim 1, characterized in that, The acquisition of the target code includes: In response to the instrumentation task, the target code is hijacked when the module of the target code is loaded, but before relocation.

3. The method according to claim 1, characterized in that, Also includes: The pre-packaged disassembler interface is invoked to disassemble the target code, thereby obtaining the disassembled result of the target code.

4. The method according to claim 3, characterized in that, Obtaining the springboard function corresponding to the target code includes: Based on the disassembly results, determine the target instruction corresponding to the instrumentation position specified by the instrumentation task; Based on the target instruction, the program counter where the instrumentation position is located, and the insertion function specified by the instrumentation task, the identification information of the jumper function is determined, and the jumper function includes at least a call to the insertion function.

5. The method according to claim 3, characterized in that, The step of inserting a jump instruction pointing to the jump function into the target code to generate a first executable file includes: Based on the disassembly results of the target code, determine the target instruction corresponding to the instrumentation position specified by the instrumentation task; Determine whether the target instruction is a preset type instruction; If the target instruction is not the preset type instruction, the target instruction is replaced with a call instruction to the jump function in the target code to generate the first executable file, wherein the jump function includes at least the relevant instructions for executing the target instruction and a call to the inserted function; If the target instruction is the preset type instruction, according to the disassembly result of the target code, a jump instruction pointing to the jump function is inserted into the basic block to which the target instruction belongs, and the first executable file is generated.

6. The method according to claim 5, characterized in that, The step of inserting jump instructions pointing to the jump function into the basic block to which the target instruction belongs, based on the disassembly result of the target code, to generate the first executable file, includes: Based on the disassembly results of the target code, determine the mapping relationship between instructions and basic blocks in the target code; The target basic block to which the target instruction belongs is determined based on the mapping relationship; The target instruction list corresponding to the target basic block is determined based on the mapping relationship; Insert a jump instruction pointing to the springboard function into the target instruction list to obtain the modified instruction list of the target basic block; The list of basic blocks in the target code is traversed, and the program counter corresponding to each basic block is updated according to the modified instruction list corresponding to each basic block to generate the first executable file.

7. The method according to claim 6, characterized in that, After traversing the list of basic blocks in the target code and updating the program counter corresponding to each basic block based on the modified instruction list corresponding to each basic block, the process further includes: Iterate through the list of basic blocks in the target code. For basic blocks that end with jump instructions based on instruction offsets, recalculate the corresponding instruction offsets based on the updated program counter.

8. The method according to any one of claims 5-7, characterized in that, If the target instruction is the preset type instruction, after inserting a jump instruction pointing to the jump function into the basic block to which the target instruction belongs based on the disassembly result of the target code, and generating the first executable file, the method further includes: If the device has already allocated memory for the target code, release the allocated memory and reallocate memory according to the instrumented code.

9. The method according to any one of claims 1-8, characterized in that, The step of completing the springboard function in the second executable file based on the implementation information of the insertion function to generate instrumented code includes: Based on the implementation information of the insertion function, determine the set of physical registers used by the insertion function; Generate save and restore codes based on the set of physical registers used by the insertion function; The springboard function is completed in the second executable file based on the saved code and the restored code, and the instrumented code is generated.

10. The method according to claim 9, characterized in that, The step of generating save code and restore code based on the set of physical registers used by the insertion function further includes: Perform an activity analysis on the physical register set; Based on the active analysis results, determine the target physical register used by the instrumentation location specified by the instrumentation task from the set of physical registers; Based on the implementation information of the insertion function, save code and restore code are generated for the target physical register.

11. A binary code instrumentation framework system, characterized in that, include: The driver middleware layer is used to obtain the target code; A code generator is used to call a pre-encapsulated instrumentation interface to obtain the springboard function corresponding to the target code; Insert a jump instruction pointing to the springboard function into the target code to generate a first executable file; link the first executable file and a preset dynamic library to generate a second executable file, wherein the preset dynamic library includes the implementation information of the inserted function; based on the implementation information of the inserted function, complete the springboard function in the second executable file to generate instrumented code.

12. The system according to claim 11, characterized in that, Also includes: A disassembler is used to call a pre-packaged disassembler interface to disassemble the target code and obtain the disassembled result of the target code, so that the code generator can generate the instrumented code based on the disassembled result.

13. An electronic device, characterized in that, include: At least one processor; as well as A memory that is communicatively connected to the at least one processor; The memory stores instructions executable by the at least one processor, which, when executed by the at least one processor, cause the electronic device to perform the method according to any one of claims 1-10.

14. A computer-readable storage medium, characterized in that, The computer-readable storage medium stores computer-executable instructions, which, when executed by a processor, implement the method as described in any one of claims 1-10.

15. A computer program product, characterized in that, Includes a computer program that, when executed by a processor, implements the method as described in any one of claims 1-10.

Images