A kernel function debugging method and device, and a storage medium
By performing semantic equivalent substitution and simulated execution on the central processing unit side, the problem of low efficiency in kernel function debugging is solved, and an efficient kernel function debugging process is achieved, simplifying debugging complexity and performance overhead.
Patent Information
- Authority / Receiving Office
- CN · China
- Patent Type
- Applications(China)
- Current Assignee / Owner
- SHANGHAI BIREN TECH CO LTD
- Filing Date
- 2026-06-15
- Publication Date
- 2026-07-21
AI Technical Summary
In existing technologies, kernel function debugging is inefficient on artificial intelligence chips. Hardware debuggers synchronously control a large number of threads, resulting in slow single-step debugging. Software emulators are slow and have difficulty supporting all GPU instructions. Printing function outputs masks the true running characteristics, leading to a complex debugging process and high performance overhead.
Semantic equivalent replacement is performed on the central processing unit side to convert the kernel function source code into the target source code. Simulation execution is performed through log printing functions and conditional return functions, avoiding the hardware debugger's synchronous control of threads, and directly simulating execution and debugging on the CPU.
It improves the efficiency of kernel function debugging, avoids the complex steps of hardware debuggers synchronously controlling a large number of threads, realizes the transparency and efficient debugging of thread bundle level instructions, and simplifies the debugging process.
Smart Images
Figure CN122432022A_ABST
Abstract
Description
Technical Field
[0001] This application relates to the field of artificial intelligence chip technology, and in particular to a kernel function debugging method, device and storage medium. Background Technology
[0002] With the continuous evolution of AI chip architectures (such as Graphics Processing Units, or GPUs), modern AI chip designs often improve overall performance by using larger-granularity instruction operations. Examples include Tensor Memory Access (TMA) instructions and TensorCore instructions. The pipelined arrangement of these large-granularity instructions has become a key implementation method for high-performance kernel functions. To ensure that kernel functions execute correctly during the application phase, correctness debugging of kernel functions during the development phase is crucial.
[0003] In related technologies, the Central Processing Unit (CPU) uses a dedicated hardware debugger to perform step-by-step debugging of kernel functions on the GPU. However, during the execution of kernel functions on the GPU, there are large-scale parallel threads. Thus, during debugging, the hardware debugger needs to synchronously control a large number of threads, and stepping through a single instruction often takes several seconds or even longer. While using the printf function to output variable values or execution status in the kernel function code can also achieve kernel function debugging, this method has a huge performance overhead and masks the true characteristics of the operation, resulting in low efficiency for kernel function debugging and development. Summary of the Invention
[0004] This invention provides a kernel function debugging method, device, and storage medium for correctness debugging of kernel functions during development, thereby improving the efficiency of kernel function debugging and development.
[0005] On one hand, embodiments of this application provide a kernel function debugging method applied to a central processing unit, the method comprising: Obtain the kernel function source code; The kernel function source code is semantically equivalently replaced to obtain the target source code, wherein the kernel function source code is the source code that the artificial intelligence chip can run, and the target source code is the source code that the central processing unit can run; The target source code and the configured set of simulated preset parameters are compiled to obtain a first executable file; Execute the first executable file to obtain the simulated execution result; Based on the simulation execution results and the preset reference execution results, the debugging results of the kernel function source code are obtained.
[0006] On one hand, embodiments of this application provide a kernel function debugging device applied to a central processing unit, the device comprising: The acquisition module is used to obtain the source code of kernel functions; The replacement module is used to perform semantically equivalent replacement on the kernel function source code to obtain the target source code, wherein the kernel function source code is the source code that the artificial intelligence chip can run, and the target source code is the source code that the central processing unit can run; The compilation module is used to compile the target source code and the configured simulated preset parameter set to obtain the first executable file; The execution module is used to execute the first executable file and obtain the simulated execution result; The debugging module is used to obtain the debugging results of the kernel function source code based on the simulated execution results and the preset reference execution results.
[0007] Optionally, the replacement module is further configured to: When the compiler macro switch is enabled, the kernel function source code is semantically equivalently replaced to obtain the target source code.
[0008] Optionally, the replacement module is further configured to: When the compiler macro switch is off, the kernel function source code is compiled to obtain a second executable file; The second executable file is sent to the artificial intelligence chip for execution to obtain the actual execution result.
[0009] Optionally, the simulated preset parameter set is passed to the central processing unit in the form of a configuration file or command line; or, the simulated preset parameter set is pre-coded in the debugging script of the central processing unit.
[0010] Optionally, the execution module is specifically used for: The thread indexes contained in the first executable file are adjusted using command-line parameters to obtain the adjusted first executable file; Execute the adjusted first executable file to obtain the simulated execution result.
[0011] Optionally, the semantic equivalent substitution includes at least one of the following operations: The built-in variables and built-in functions in the kernel function source code are encapsulated into a conditional return function; The device modifier in the kernel function source code is encapsulated as an empty flag by macro definition; Replace the architecture instructions in the kernel function source code with log printing functions; Replace the kernel startup instructions in the kernel function source code with direct function call instructions.
[0012] Optionally, the compilation module is specifically used for: The built-in variables and built-in functions in the kernel function source code are encapsulated into a conditional return function controlled by a macro. The conditional return function returns a preset fixed value from the simulated preset parameter set.
[0013] Optionally, the simulation execution result includes: key parameters of the architecture instructions obtained by executing the log printing function.
[0014] On one hand, embodiments of this application provide a computer device, including: a memory, an artificial intelligence chip, a central processing unit, and a computer program stored in the memory and running on the central processing unit, wherein the central processing unit executes the computer program to implement the steps of the above-described kernel function debugging method.
[0015] On one hand, embodiments of this application provide a computer-readable storage medium storing a computer program executable by a computer device, which, when run on the computer device, causes the computer to perform the steps of the above-described kernel function debugging method.
[0016] On one hand, embodiments of this application provide a computer program product, including a computer program stored on a computer-readable storage medium. The computer program includes program instructions, which, when executed by a computer device, cause the computer device to perform the steps of the above-described kernel function debugging method.
[0017] In this embodiment, the central processing unit (CPU) performs semantically equivalent substitution on the acquired kernel function source code to obtain target source code. The kernel function source code is the source code that the AI chip can run, and the target source code is the source code that the CPU can run. The target source code and the configured set of preset simulation parameters are compiled to obtain a first executable file. Then, the first executable file is executed to obtain simulated execution results. This transforms the kernel function debugging process, which originally needed to be performed on the AI chip, to be performed on the CPU, avoiding the complex steps of hardware debuggers simultaneously controlling a large number of threads. Simultaneously, the CPU compiles and directly runs the transformed target source code to obtain simulated execution results. Then, based on the simulated execution results and preset reference execution results, the debugging results of the kernel function source code are obtained, instead of performing single-step debugging of the kernel function. This solves the problem of excessively slow debugging processes and improves the debugging efficiency of kernel functions. Attached Figure Description
[0018] To more clearly illustrate the technical solutions in the embodiments of the present invention, the accompanying drawings used in the description of the embodiments will be briefly introduced below. Obviously, the accompanying drawings described below are only some embodiments of the present invention. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort.
[0019] Figure 1 This is a schematic diagram of a chip architecture provided in an embodiment of this application; Figure 2 A flowchart illustrating a kernel function debugging method provided in an embodiment of this application; Figure 3 A flowchart illustrating another kernel function debugging method provided in this application embodiment; Figure 4 This is a schematic diagram of the structure of a kernel function debugging device provided in an embodiment of this application; Figure 5 This is a schematic diagram of the structure of a computer device provided in an embodiment of this application. Detailed Implementation
[0020] To make the objectives, technical solutions, and beneficial effects of this invention clearer, the invention will be further described in detail below with reference to the accompanying drawings and embodiments. It should be understood that the specific embodiments described herein are merely illustrative and not intended to limit the invention. All other embodiments obtained by those skilled in the art based on the embodiments of this invention without inventive effort are within the scope of protection of this invention.
[0021] It should be noted that in this application, "first" and "second" are used for descriptive purposes only and should not be construed as indicating relative importance or implying the number of technical features indicated.
[0022] refer to Figure 1 This is a structural diagram of an artificial intelligence chip applicable to the embodiments of this application. The artificial intelligence chip 100 includes at least: video memory 101 and computing unit 102. This application does not specifically limit the number of video memory 101 and computing unit 102.
[0023] The computing unit 102 can be a streaming processing cluster (SPC). The video memory 101 can be high-bandwidth memory (HBM) or other types of memory.
[0024] Each computing unit 102 can launch one or more warps, each warp consisting of multiple threads (e.g., 32 threads).
[0025] In the embodiments of this application, in debug mode (with the compiler macro switch on), the central processing unit performs semantic equivalent substitution on the kernel function source code to obtain the target source code, and compiles the target source code into a first executable file; then the central processing unit executes the first executable file.
[0026] In non-debug mode (with the compiler macro switch off), the central processing unit compiles the kernel function source code to obtain a second executable file, and then sends the second executable file to the artificial intelligence chip 100 for execution, so that the kernel function runs correctly on the artificial intelligence chip 100.
[0027] In addition to the structure described above, the artificial intelligence chip 100 in this application may also include other structures, and this application does not specifically limit such structures.
[0028] Artificial intelligence chips 100 can be: Graphics Processing Unit (GPU), General-purpose computing on graphics processing units (GPGPU), Domain Specific Architecture (DSA), etc.
[0029] In related technologies, the host side provides a mechanism for debugging kernel function source code on artificial intelligence chips through a hardware debugger (GDB). This allows setting breakpoints in the kernel function source code, single-stepping through hardware threads, inspecting and modifying the memory and variables of any hardware thread, and switching focus between hardware threads and hardware thread bundles using thread-related commands (e.g., `info cuda threads`) to analyze the execution behavior of specific hardware threads. However, GDB needs to synchronously control thousands of hardware threads. In this case, single-stepping through the instructions corresponding to a single hardware thread takes significantly more time and consumes excessive resources, slowing down the entire debugging process and making it difficult to quickly and efficiently locate logical errors in the kernel function source code. Furthermore, while GDB's compilation options (such as -G -O0) can generate executables with device debugging information for kernel function source code, these options disable optimization by default, disabling all optimization features such as instruction reordering, dead code removal, and constant folding. This results in an excessively bloated executable, with runtime and register pressure far exceeding the expected release version. If the -G -O2 compilation option provided by GDB is used, although optimization features are retained, the generated debugging information has garbled line numbers and unreadable variable values, causing line numbers to jump around during single-step debugging, thus negating the purpose of single-step debugging.
[0030] While related technologies employ software emulators to simulate the execution of kernel function source code in a pure software environment, which can alleviate the aforementioned problems associated with debugging using GDB to some extent, software emulators still have a performance gap compared to real hardware, resulting in slow simulation speeds. Furthermore, the pure software environment simulated by software emulators is difficult to meet the characteristics of all GPU instructions, such as thread bundle level instructions (e.g., TMA).
[0031] Secondly, hardware debuggers struggle to visually demonstrate complex data layouts and tensor transformations when debugging kernel function source code. While related techniques can address this by using printf functions to output variable values in the kernel function source code, each call to printf triggers global memory writes and potential synchronization operations, slowing down kernel function execution. Furthermore, the variable values output by printf are written to a circular buffer, requiring a call to the cudaDeviceSynchronize instruction after each kernel function call to ensure the output variable values are refreshed. Thus, the resources consumed by numerous printf runs obscure the true execution of the kernel function, resulting in significant performance overhead.
[0032] In view of this, this application provides a kernel function debugging method, which enables debugging of kernel functions on the CPU side. achievement Simulation and debugging See Figure 2 This method is performed by a computer device, which includes... Figure 1 The artificial intelligence chip shown includes the following steps: Step 201: Obtain the kernel function source code.
[0033] Specifically, the kernel function debugging method of this application can be applied to various scenarios, such as image processing, speech processing, and text processing. In different application scenarios, the physical meaning of the data processed by the kernel function being debugged can be different.
[0034] For example, in a text processing scenario, the data processed by a kernel function can be text data used in tasks such as text generation and text recognition.
[0035] For example, in a speech processing scenario, the data processed by the kernel function can be speech data used in tasks such as speech enhancement, speech recognition, and speech synthesis.
[0036] For example, in image processing scenarios, the data processed by kernel functions can be image data used in tasks such as image preprocessing, image segmentation, and object detection.
[0037] Step 202: Perform semantic equivalent replacement on the kernel function source code to obtain the target source code; wherein, the kernel function source code is the source code that the artificial intelligence chip can run, and the target source code is the source code that the central processing unit can run.
[0038] Specifically, semantic equivalent substitution refers to replacing specific instructions in the kernel function source code that need to be executed on the artificial intelligence chip with implementation instructions on the central processing unit, while ensuring consistency in logical behavior.
[0039] In some embodiments, when the compiler macro switch is enabled, the kernel function source code is semantically equivalently replaced to obtain the target source code.
[0040] Specifically, the compiler macro switch is a conditional compilation mechanism that utilizes the conditional compilation features of preprocessor macros to selectively enable or disable device-specific code segments during compilation. In the embodiments of this application, when the compiler macro switch is enabled, semantically equivalent substitution is performed on the kernel function source code to obtain the target source code, thereby achieving simulated compilation and execution on the central processing unit, rather than compiling and executing the kernel function source code in the environment of an artificial intelligence chip.
[0041] In some embodiments, semantic equivalent substitution includes at least one of the following operations: encapsulating built-in variables and built-in functions in the kernel function source code into conditional return functions; encapsulating device modifiers in the kernel function source code into null flags through macro definitions; replacing architecture instructions in the kernel function source code with log printing functions; and replacing kernel boot instructions in the kernel function source code with direct function call instructions.
[0042] Specifically, when the compiler macro switch is enabled, semantic equivalent substitution of the kernel function source code is initiated; when the semantic equivalent substitution includes multiple operations as described above, these multiple operations can be processed in parallel or sequentially, and this application does not impose any specific limitations on this.
[0043] In this embodiment, the device modifiers in the kernel function source code include global execution modifier (_global_), device execution modifier (_device_), shared memory modifier (_shared_), etc. These device modifiers are replaced with null flags or inline modifiers through macro definitions. The inline modifier is a standard keyword in C++, which is used to quickly copy the modifier function modified by inline to all call sites of the function during the compilation process, so that the modifier function does not need to be written again at the call site.
[0044] The architecture instructions in the kernel function source code include, but are not limited to: Tensor Memory Access (TMA) instruction, Load Matrix (LDMATRIX) instruction, General Matrix Multiply Accumulate (GMMA) instruction, and Multicast Barrier (MBARRIER) instruction. In semantic equivalence substitution, the above architecture instructions are replaced with log printing functions, which can print the key parameters of the architecture instructions. For different architecture instructions, the key parameters printed by the corresponding log printing functions are different.
[0045] Among them, the TMA instruction is a high-efficiency memory access instruction used to load data directly from global memory to shared memory or to transfer data between clusters of different artificial intelligence chips. When performing equivalent semantic replacement operations, the TMA instruction is replaced by a log printing function to print key parameters such as the TMA instruction descriptor pointer, shared memory pointer (i.e., shared memory address), coordinate parameters, size parameters, and barrier addresses.
[0046] The LDMATRIX instruction is a warp-level matrix loading instruction that can load multiple matrix data into registers simultaneously within a single warp. When performing an equivalent semantic substitution operation, the LDMATRIX instruction is replaced with a log printing function to print key parameters of the LDMATRIX instruction, such as matrix dimensions, transpose flag, barrier address, step size, shared memory pointer (i.e., shared memory address), and commit queue.
[0047] The GMMA instruction is an accumulation operation instruction that combines matrix multiplication. It is used to perform matrix multiplication on operands in one instruction and accumulate the result to the destination register. When performing the equivalent semantic substitution operation, the GMMA instruction is replaced with a log printing function to print the data type (e.g., BF16 / FP16 / F32 type), transpose flag, input / output register pointer, shared memory pointer, step size, and wait barrier of the operands of the GMMA instruction.
[0048] The MBARRIER instruction is a hardware barrier instruction that supports multicast and is used to achieve one-to-many signal synchronization. When performing an equivalent semantic substitution operation, the MBARRIER instruction is replaced with a log printing function to print key parameters of the MBARRIER instruction, such as the barrier address, parity bit, arrival of synchronization instruction point, and waiting for synchronization instruction point.
[0049] While existing technologies use the `printf` function to output variable values or execution status, which can solve the problems of difficulty in visually displaying complex data layouts and tensor transformations, and the difficulty in recognizing highly templated function symbol names when debugging kernel function source code with hardware debuggers, the following issues remain: each thread in the kernel function source code executes `printf`, and with a large number of threads running in parallel, it is difficult to coordinate the output order of `printf` from different threads. The output order of `printf` is highly dependent on the thread scheduling order, and the output time of `printf` is completely different for different threads, resulting in a chaotic and complex debugging process. Furthermore, `printf` has limited functionality, often only able to print basic types (such as numbers and characters), and cannot print complex data structures.
[0050] Therefore, in this embodiment, only the architecture instructions are replaced with log printing functions. The architecture instructions are all thread bundle granular instructions. The same thread bundle contains multiple threads, and multiple threads execute the same task. In this way, it is equivalent to only needing to focus on whether the key parameters printed by one thread in a thread bundle granular instruction meet expectations, avoiding thread-by-thread simulation compilation and execution, and not needing to pay attention to the execution details of each thread, which improves the efficiency of kernel function debugging by orders of magnitude.
[0051] In addition, when the kernel startup instructions in the kernel function source code are replaced with direct function call instructions, startup configuration parameters are passed to the direct function call instructions. These startup configuration parameters include, but are not limited to, grid dimension (grid_dim), thread block dimension (block_dim), and thread cluster dimension (cluster_dim). Only by correctly passing these startup configuration parameters can the target source code called by the direct function call instructions obtain the correct startup information.
[0052] In this embodiment, the kernel function source code is losslessly converted into printable target source code through a semantic equivalent substitution operation. This solves the problem of not being able to intuitively observe the debugging process in related technologies at a lower development cost and achieves transparency of thread bundle level instructions.
[0053] In some embodiments, the built-in variables and built-in functions in the kernel function source code are encapsulated into a conditional return function controlled by a macro. The conditional return function returns a preset fixed value from a simulated preset parameter set.
[0054] Specifically, the built-in variables include, but are not limited to, thread block dimension variables (blockDim.x), thread index variables (threadIdx.y), and cluster dimension variables (cluster_dim.x), etc., and the built-in functions include, but are not limited to, thread synchronization functions and atomic operation functions. These built-in variables and built-in functions are encapsulated into conditional return functions controlled by macros, and a preset fixed value is set for the return of the conditional return function in the simulated preset parameter set. This preset fixed value can be a thread index range, cluster dimension, etc., and this application does not specifically limit it.
[0055] In this embodiment, built-in variables and built-in functions are encapsulated into conditional return functions that return preset fixed values. This allows for the accurate reproduction of the thread bundle query index behavior in the kernel function source code without needing to query real values from real hardware registers in simulation mode. This enables the rapid and efficient reproduction of the kernel function source code's behavioral logic on the central processing unit.
[0056] Step 203: Compile the target source code and the configured simulation preset parameter set to obtain the first executable file.
[0057] In some embodiments, the simulated preset parameter set is passed to the central processing unit in the form of a configuration file or command line; or, the simulated preset parameter set is pre-coded in the central processing unit's debugging script.
[0058] Specifically, the simulated preset parameter set includes not only the preset fixed values returned by the conditional return function, but also other configuration data. For example, grid size, thread block size, cluster dimension, shared memory size, etc., defined by developers or testers.
[0059] The simulation preset parameter set is used to ensure that the debugging environment is consistent with the actual operating environment of the GPU device when debugging kernel functions in simulation mode, thereby ensuring that the context during simulation execution is consistent with the debugging target of the actual device.
[0060] Step 204: Execute the first executable file to obtain the simulated execution results.
[0061] Specifically, the first executable file is a simulated executable file that can be run directly on the central processing unit.
[0062] In some embodiments, the thread index contained in the first executable file is adjusted by command-line parameters to obtain an adjusted first executable file; the adjusted first executable file is executed to obtain a simulated execution result.
[0063] Specifically, the first executable file and the modified first executable file are used to track different thread behaviors or thread bundle behaviors. The threads or thread bundles to be tracked are configured using command-line parameters such as -thr or -blk to achieve flexible tracking of different threads or thread bundles. In practical applications, by adjusting the thread index in the first executable file, a first executable file for tracking another thread or thread bundle can be obtained.
[0064] In this embodiment, the thread index contained in the first executable file is adjusted by command line parameters to obtain the first executable file that tracks other threads. This allows the executable file that tracks different threads / thread bundles to be obtained by compiling the target source code only once on the central processing unit, which facilitates comparative analysis of the behavior of multiple threads / thread bundles and thus improves compilation efficiency.
[0065] In some embodiments, the simulation execution results include key parameters of the architecture instructions obtained from the execution log printing function.
[0066] Specifically, during the execution of the first executable file, the CPU actually executes the logic of the kernel function source code. However, it no longer executes the architecture instructions in the kernel function source code. Instead, it executes the replaced log printing function to obtain the key parameters of the architecture instructions printed line by line. These key parameters include: the descriptor pointer, shared memory pointer, coordinate parameters, size parameters, and barrier address of the TMA instruction; the matrix dimension, transpose flag, barrier address, step size, shared memory pointer, and commit queue of the LDMATRIX instruction; the data type of the operand, transpose flag, input / output register pointer, shared memory pointer, step size, and waiting barrier of the GMMA instruction; and the barrier address, parity bit, reached synchronization instruction point, and waiting synchronization instruction point of the MARRRIER instruction.
[0067] Similarly, during the execution of the first executable file, the CPU no longer executes device modifiers in the kernel function source code, but instead executes the replaced null flags or inline modifiers to eliminate call overhead during compilation in emulation mode. It no longer executes kernel startup instructions in the kernel function source code, but instead executes replaced direct function call instructions; it no longer executes built-in variables or built-in functions in the kernel function source code, but instead executes replaced conditional return functions, returning preset fixed values from the simulated preset parameter set.
[0068] Thus, in addition to printing the key parameters of each architectural instruction, the simulation execution results also include: the startup configuration parameters passed to direct function call instructions and the preset fixed values returned by conditional return functions. Furthermore, the simulation execution results also include the execution status of instructions other than the architectural instructions, such as the shared memory allocation tables of thread bundle-level instructions like TensorCore and SYNCTHREADS. These shared memory allocation tables record information such as the starting position, alignment position, and size of each memory bank.
[0069] In the embodiments of this application, in simulation mode, the complete parameters of each instruction in the kernel function source code and the visualization of the layout type of the implementation data are obtained through simulation execution results. Compared with compiling and executing kernel function source code in a black box environment such as GPU, simulation mode makes various information during the debugging process transparent.
[0070] Step 205: Based on the simulation execution results and the preset reference execution results, obtain the debugging results of the kernel function source code.
[0071] Specifically, the reference execution result is preset by the developers; for example, the reference execution result includes: key parameters of the architecture instruction, startup configuration parameters of the direct function call instruction, and preset fixed values returned by the return function; in addition, the reference execution result may also include: the layout type corresponding to the preset fixed values returned by the conditional return function, and other configuration data customized by the developers for other instructions.
[0072] The reference execution results are compared with the simulated execution results to obtain comparison results. The comparison results indicate logical errors or configuration problems in the kernel function source code, such as whether the data flow is correct or whether the address calculation is as expected. Then, the kernel function source code is optimized based on the comparison results to obtain the debugging results of the kernel function source code.
[0073] In this embodiment, the central processing unit (CPU) performs semantically equivalent substitution on the acquired kernel function source code to obtain target source code. The kernel function source code is the source code that the AI chip can run, and the target source code is the source code that the CPU can run. The target source code and the configured set of preset simulation parameters are compiled to obtain a first executable file. Then, the first executable file is executed to obtain simulated execution results. This transforms the kernel function debugging process, which originally needed to be performed on the AI chip, to be performed on the CPU, avoiding the complex steps of hardware debuggers simultaneously controlling a large number of threads. Simultaneously, the CPU compiles the transformed target source code and runs it directly to obtain simulated execution results. Then, based on the simulated execution results and preset reference execution results, the debugging results of the kernel function source code are obtained, instead of performing single-step debugging of the kernel function. This solves the problem of excessively slow debugging processes and improves the debugging efficiency of kernel functions.
[0074] In some embodiments, when the compiler macro switch is off, the kernel function source code is compiled to obtain a second executable file; the second executable file is sent to the artificial intelligence chip for execution to obtain the actual execution result.
[0075] Specifically, when the compiler macro switch is detected to be in the off state, the built-in variables, built-in functions, architecture instructions, device modifiers, etc. in the kernel function source code remain unchanged. Then, the kernel function source code is directly compiled into an executable file (i.e., the second executable file) that generates a real artificial intelligence chip by the compiler. This executable file can be directly sent to the artificial intelligence chip for execution.
[0076] In some cases, after successfully compiling, executing, and debugging the kernel function source code through simulation mode, the debugged kernel function source code can be compiled into an executable file by disabling the compiler macro switch and sent to the artificial intelligence chip for execution, so as to achieve the final performance test in the development stage or the actual deployment when going online.
[0077] To better explain the embodiments of this application, the following is combined with... Figure 3 This application provides a flowchart of a kernel function debugging method, which includes the following steps: Please see Figure 3 This demonstrates that this application controls different compilation paths through compiler macro switches, specifically including the following steps: Step 301: Obtain the kernel function source code.
[0078] Step 302: Determine whether the compiler macro switch is enabled. If yes, proceed to step 303; otherwise, proceed to step 307.
[0079] Step 303: In simulation mode, the kernel function source code is semantically equivalently replaced by the replacement layer and the preset parameter set in the host simulation debugging framework to obtain the target source code.
[0080] Specifically, the host simulation debugging framework is located in the central processing unit, and the replacement layers include: built-in variable replacement layer, device modifier replacement layer, hardware instruction log replacement layer, and kernel boot replacement layer.
[0081] The built-in variable substitution layer replaces built-in variables in the kernel function source code with conditional return functions, which return preset fixed values from the simulated preset parameter set. The kernel boot substitution layer replaces kernel boot instructions in the kernel function source code with direct function call instructions, and passes boot configuration parameters from the simulated preset parameter set to the direct function call instructions, such as grid dimension (grid_dim), thread block dimension (block_dim), and thread cluster dimension (cluster_dim). The device modifier substitution layer replaces device modifiers such as global execution modifier (_global_), device execution modifier (_device_), and shared memory modifier (_shared_) in the kernel function source code with null flags or inline modifiers. The hardware instruction log substitution layer replaces architecture instructions in the kernel function source code with log printing functions, which print key parameters of the architecture instructions, such as architecture instruction name, shared memory address, coordinate parameters, compensation, size parameters, multicast mask, step size, and other input / output information.
[0082] Specifically, regarding the printing of shared memory addresses in key parameters, since the shared memory in an AI chip is divided into multiple independent banks, different banks can be accessed in parallel, but parallel access to the same bank can lead to serialization conflicts. Therefore, the shared memory in an AI chip is divided into low banks and high banks. This application clearly shows the location of different buffers in the low and high banks by printing the shared memory addresses in the architecture instructions, assisting developers in identifying and avoiding serialization conflicts in the banks.
[0083] Step 304: Compile the target source code to obtain the first executable file.
[0084] Step 305: Execute the first executable file on the CPU.
[0085] Step 306: Output the simulation execution results.
[0086] Specifically, the simulation execution results include: key parameters of each architecture instruction printed, preset fixed values returned by conditional return functions, and startup configuration parameters of direct function call instructions.
[0087] Step 307: Compile the kernel function source code to obtain the second executable file.
[0088] Specifically, when the compiler macro switch is turned off, all replacement layers are disabled, and the architecture instructions, built-in variables, built-in functions, device modifiers, and kernel boot instructions in the kernel function source code remain unchanged. The kernel function source code is then compiled by the compiler to generate a second executable file that can run on a real GPU.
[0089] Step 308: Execute the second executable file on the GPU.
[0090] Specifically, executing the second executable file on the GPU does not produce any log output, and there is no additional overhead in executing the second executable file.
[0091] In this embodiment, a compilation path for simulation mode can be added to the actual GPU compilation path by using a compiler macro switch. Without a hardware debugger or any modification to the kernel function source code, the kernel function source code can be simulated and compiled and executed on the CPU, achieving seamless migration from GPU code to CPU code and significantly accelerating the kernel function development process.
[0092] Based on the same technical concept, this application provides a schematic diagram of the structure of a kernel function debugging device, such as... Figure 4 As shown, the kernel function debugging device 400 includes: Module 401 is used to retrieve the source code of kernel functions; Replacement module 402 is used to perform semantic equivalent replacement on the kernel function source code to obtain target source code, wherein the kernel function source code is source code that can be run by an artificial intelligence chip, and the target source code is source code that can be run by a central processing unit; The compilation module 403 is used to compile the target source code and the configured simulated preset parameter set to obtain the first executable file; Execution module 404 is used to execute the first executable file and obtain the simulated execution result; The debugging module 405 is used to obtain the debugging results of the kernel function source code based on the simulated execution results and the preset reference execution results.
[0093] Optionally, the replacement module 402 is further configured to: When the compiler macro switch is enabled, the kernel function source code is semantically equivalently replaced to obtain the target source code.
[0094] Optionally, the replacement module 402 is further configured to: When the compiler macro switch is off, the kernel function source code is compiled to obtain a second executable file; The second executable file is sent to the artificial intelligence chip for execution to obtain the actual execution result.
[0095] Optionally, the simulated preset parameter set is passed to the central processing unit in the form of a configuration file or command line; or, the simulated preset parameter set is pre-coded in the debugging script of the central processing unit.
[0096] Optionally, the execution module 404 is specifically used for: The thread indexes contained in the first executable file are adjusted using command-line parameters to obtain the adjusted first executable file; Execute the adjusted first executable file to obtain the simulated execution result.
[0097] Optionally, the semantic equivalent substitution includes at least one of the following operations: The built-in variables and built-in functions in the kernel function source code are encapsulated into a conditional return function; The device modifier in the kernel function source code is encapsulated as an empty flag by macro definition; Replace the architecture instructions in the kernel function source code with log printing functions; Replace the kernel startup instructions in the kernel function source code with direct function call instructions.
[0098] Optionally, the compilation module 403 is specifically used for: The built-in variables and built-in functions in the kernel function source code are encapsulated into a conditional return function controlled by a macro. The conditional return function returns a preset fixed value from the simulated preset parameter set.
[0099] Optionally, the simulation execution result includes: key parameters of the architecture instructions obtained by executing the log printing function.
[0100] In this embodiment, the central processing unit (CPU) performs semantically equivalent substitution on the acquired kernel function source code to obtain target source code. The kernel function source code is the source code that the AI chip can run, and the target source code is the source code that the CPU can run. The target source code and the configured set of preset simulation parameters are compiled to obtain a first executable file. Then, the first executable file is executed to obtain simulated execution results. This transforms the kernel function debugging process, which originally needed to be performed on the AI chip, to be performed on the CPU, avoiding the complex steps of hardware debuggers simultaneously controlling a large number of threads. Simultaneously, the CPU compiles and directly runs the transformed target source code to obtain simulated execution results. Then, based on the simulated execution results and preset reference execution results, the debugging results of the kernel function source code are obtained, instead of performing single-step debugging of the kernel function. This solves the problem of excessively slow debugging processes and improves the debugging efficiency of kernel functions.
[0101] Based on the same technical concept, embodiments of this application provide a computer device, such as... Figure 5 As shown, it includes at least one artificial intelligence chip 100, and a memory 501 and a central processing unit 502 connected to the at least one artificial intelligence chip 100. In this embodiment, the specific connection medium between the artificial intelligence chip 100 and the memory 501 and central processing unit 502 is not limited. Figure 5 Taking the AI chip 100, memory 501, and central processing unit 502 as an example, they are connected via a bus. The bus can be divided into address bus, data bus, control bus, etc.
[0102] In this embodiment of the application, the memory 501 stores instructions that can be executed by the central processing unit 502 to perform the steps of the above-described kernel function debugging method.
[0103] The central processing unit 502 is the control center of the computer device. It can connect to various parts of the computer device through various interfaces and lines. It can debug kernel functions by running or executing instructions stored in memory 501 and calling data stored in memory 501.
[0104] In some embodiments, the central processing unit 502 and the memory 501 may be implemented on the same chip, while in other embodiments they may be implemented on separate chips.
[0105] Memory 501, as a non-volatile computer-readable storage medium, can be used to store non-volatile software programs, non-volatile computer-executable programs, and modules. Memory 501 may include at least one type of storage medium, such as flash memory, hard disk, multimedia card, card-type memory, random access memory (RAM), static random access memory (SRAM), programmable read-only memory (PROM), read-only memory (ROM), electrically erasable programmable read-only memory (EEPROM), magnetic storage, magnetic disk, optical disk, etc. Memory 501 can be any other medium capable of carrying or storing desired program code in the form of instructions or data structures and accessible by a computer device, but is not limited thereto. Memory 501 in the embodiments of this application may also be a circuit or any other device capable of implementing storage functions for storing program instructions and / or data.
[0106] Based on the same inventive concept, embodiments of this application provide a computer-readable storage medium storing a computer program executable by a computer device, which, when run on the computer device, causes the computer device to perform the steps of the above-described kernel function debugging method.
[0107] Based on the same inventive concept, this application provides a computer program product, which includes a computer program stored on a computer-readable storage medium. The computer program includes program instructions, which, when executed by a computer device, cause the computer device to perform the steps of the above-described kernel function debugging method.
[0108] Those skilled in the art will understand that embodiments of this application can be provided as methods, systems, or computer program products. Therefore, this application can take the form of a completely hardware embodiment, a completely software embodiment, or an embodiment combining software and hardware aspects. Furthermore, this application can take the form of a computer program product embodied on one or more computer-usable storage media (including but not limited to disk storage, CD-ROM, optical storage, etc.) containing computer-usable program code.
[0109] This application is described with reference to flowchart illustrations and / or block diagrams of methods, apparatus (systems), and computer program products according to this application. It should be understood that each block of the flowchart illustrations and / or block diagrams, and combinations of blocks in the flowchart illustrations and / or block diagrams, can be implemented by computer program instructions. These computer program instructions can be provided to a processor of a general-purpose computer, special-purpose computer, embedded processor, or other programmable data processing apparatus to produce a machine, such that the instructions, which execute via the processor of the computer or other programmable data processing apparatus, generate instructions for implementing the flowchart illustrations and / or block diagrams. Figure 1 One or more processes and / or boxes Figure 1 A device that provides the functions specified in one or more boxes.
[0110] These computer program instructions may also be stored in a computer-readable storage medium that can direct a computer or other programmable data processing device to function in a particular manner, such that the instructions stored in the computer-readable storage medium produce an article of manufacture including instruction means, which are implemented in a process Figure 1 One or more processes and / or boxes Figure 1 The function specified in one or more boxes.
[0111] These computer program instructions may also be loaded onto a computer or other programmable data processing equipment to cause a series of operational steps to be performed on the computer or other programmable equipment to produce a computer-implemented process, thereby providing instructions that execute on the computer or other programmable equipment for implementing the process. Figure 1 One or more processes and / or boxes Figure 1 The steps of the function specified in one or more boxes.
[0112] Obviously, those skilled in the art can make various modifications and variations to this application without departing from the spirit and scope of this application. Therefore, if such modifications and variations fall within the scope of the claims of this application and their equivalents, this application also intends to include such modifications and variations.
Claims
1. A kernel function debugging method, characterized in that, Applied to a central processing unit, the method includes: Obtain the kernel function source code; The kernel function source code is semantically equivalently replaced to obtain the target source code, wherein the kernel function source code is the source code that the artificial intelligence chip can run, and the target source code is the source code that the central processing unit can run; The target source code and the configured set of simulated preset parameters are compiled to obtain a first executable file; Execute the first executable file to obtain the simulated execution result; Based on the simulation execution results and the preset reference execution results, the debugging results of the kernel function source code are obtained.
2. The method as described in claim 1, characterized in that, The step of performing semantic equivalent replacement on the kernel function source code to obtain the target source code includes: When the compiler macro switch is enabled, the kernel function source code is semantically equivalently replaced to obtain the target source code.
3. The method as described in claim 1, characterized in that, Also includes: When the compiler macro switch is off, the kernel function source code is compiled to obtain a second executable file; The second executable file is sent to the artificial intelligence chip for execution to obtain the actual execution result.
4. The method as described in claim 1, characterized in that, The simulated preset parameter set is passed to the central processing unit in the form of a configuration file or command line; or, the simulated preset parameter set is pre-coded in the debugging script of the central processing unit.
5. The method as described in claim 1, characterized in that, The execution of the first executable file to obtain the simulated execution result includes: The thread indexes contained in the first executable file are adjusted using command-line parameters to obtain the adjusted first executable file; Execute the adjusted first executable file to obtain the simulated execution result.
6. The method according to any one of claims 1 to 5, characterized in that, The semantic equivalent substitution includes at least one of the following operations: The built-in variables and built-in functions in the kernel function source code are encapsulated into a conditional return function; The device modifier in the kernel function source code is encapsulated as an empty flag by macro definition; Replace the architecture instructions in the kernel function source code with log printing functions; Replace the kernel startup instructions in the kernel function source code with direct function call instructions.
7. The method as described in claim 6, characterized in that, The step of encapsulating the built-in variables and built-in functions in the kernel function source code into a conditional return function includes: The built-in variables and built-in functions in the kernel function source code are encapsulated into a conditional return function controlled by a macro. The conditional return function returns a preset fixed value from the simulated preset parameter set.
8. The method as described in claim 6, characterized in that, The simulation execution results include: key parameters of the architecture instructions obtained by executing the log printing function.
9. A computer device, comprising a memory, an artificial intelligence chip, a central processing unit (CPU), and a computer program stored in the memory and running on the CPU, characterized in that, When the central processing unit executes the computer program, it implements the steps of the method as described in any one of claims 1 to 8.
10. A computer-readable storage medium, characterized in that, It stores a computer program that is executed by a computer device, which, when run on the computer device, causes the computer device to perform the steps of the method as described in any one of claims 1 to 8.