A GPU intermediate representation program testing method, device and equipment

Abstract

This invention relates to a method, apparatus, and device for testing GPU intermediate representation programs. The method includes: acquiring the LLVM IR of the GPU kernel to be tested; parsing the LLVM IR to extract kernel function signatures, parameter data types, memory layout, and access mode information; generating GPU host-side driver code and CPU reference implementation code based on the parsed kernel function signatures, parameter data types, memory layout, and access mode information; compiling the LLVM IR into GPU executable code and executing the host-side driver code; and comparing the GPU execution results with the CPU reference implementation results to verify the correctness of the kernel. This method provides efficient, universal, and accurate GPU compiler correctness verification capabilities.

Description

Technical Field

[0001] This invention relates to the field of GPU compiler verification technology, and in particular to a GPU intermediate representation program testing method, apparatus and device. Background Technology

[0002] The GPU compiler is a crucial bridge connecting high-level programming languages ​​and GPU hardware; the correctness of the intermediate representation code it generates directly determines the accuracy of the final execution result. During compiler development, efficiently and accurately verifying the correctness of the intermediate representation generated by the compiler is a core issue in ensuring the reliability of GPU computing.

[0003] Currently, the correctness testing of intermediate representations in GPU compilers mainly faces the following technical challenges: Test case construction is challenging: Existing compiler testing methods typically employ end-to-end testing, unit testing, or regression testing. End-to-end testing verifies correctness by running the complete program and comparing results, but struggles to pinpoint the source of errors. Unit testing targets individual operators or compiler passes, but lacks a realistic execution environment, making it impossible to detect runtime errors. Regression testing uses historical test cases to prevent the introduction of new errors, but its coverage is limited and it lacks sufficient support for new operators. More critically, a significant semantic gap exists between the bare intermediate representation output by the compiler and the complete executable test cases. Traditional methods cannot automatically construct executable test cases that include complete logic such as device initialization, memory allocation, data transfer, and kernel startup, leading to a high reliance on manual intervention in the testing process, resulting in inefficiency and a high risk of errors.

[0004] Input / output shapes are difficult to infer automatically: GPU kernel input / output parameters are typically represented as raw pointers in intermediate representations. Key information such as data type, tensor shape, memory layout, and alignment is implicit in the compiler-inserted memory attributes, load / store instructions, and thread index calculation logic. For example, in LLVM IR (LLVM Intermediate Representation), the `dereferenceable` attribute indicates the number of bytes the pointer can dereference, the `align` attribute indicates memory alignment requirements, and the `noalias` attribute indicates whether parameters overlap. However, this information requires manual parsing and inference. Current technology cannot automatically reconstruct the complete tensor description from this low-level information, forcing testers to manually analyze intermediate representation code and construct test data, which is almost impractical for complex operators or batch testing scenarios.

[0005] Lack of Reference Implementations: Verifying the correctness of a GPU kernel requires comparing it with expected results, but CPU reference implementations for complex operators require extensive manual coding. Especially during compiler development iterations, operator semantics are frequently adjusted, making manual maintenance of reference implementations extremely costly and prone to human error. Current technologies lack the ability to automatically generate semantically equivalent CPU reference implementations from intermediate representations, causing the verification process to rely on limited benchmark sets or manually coded results, making it difficult to guarantee coverage and accuracy.

[0006] Strong environment dependency and poor portability: Traditional testing methods are typically tightly coupled to specific compiler ecosystems, making it difficult to migrate test cases to different GPU architectures or compiler backends. For example, test cases developed for the XLA compiler cannot be directly used on the TVM compiler, and tests designed for the NVIDIA CUDA architecture cannot be directly used to verify the MooreThreads MUSA platform. When cross-platform consistency verification is required, a complete compilation environment and test framework must be rebuilt on each target platform, resulting in long testing cycles, high costs, and difficult maintenance. Furthermore, the testing process depends on the runtime environment of the original compiler, making it difficult to distinguish between front-end lowering semantic deviations and errors introduced by back-end optimizations, thus hindering defect localization.

[0007] In summary, existing GPU compiler intermediate representation correctness testing methods suffer from technical problems such as difficulty in constructing test cases, inability to automatically infer input and output shapes, lack of reference implementations, and strong environment dependence. There is an urgent need for a general, automated testing solution that is independent of any specific compiler ecosystem. Summary of the Invention

[0008] Therefore, it is necessary to provide a GPU intermediate representation program testing method, apparatus, and device that can achieve efficient, universal, and accurate GPU compiler testing to address the aforementioned technical problems.

[0009] A method for testing GPU intermediate representation programs, the method comprising: Obtain the LLVM IR of the GPU kernel to be tested.

[0010] Parse LLVM IR to extract kernel function signatures, parameter data types, memory layout, and access mode information.

[0011] Based on the parsed kernel function signatures, parameter data types, memory layout, and access mode information, GPU host-side driver code and CPU reference implementation code are generated.

[0012] Compile LLVM IR into GPU executable code and execute host-side driver code.

[0013] The GPU execution results are compared with the CPU reference implementation results to verify the correctness of the kernel.

[0014] A GPU intermediate representation program testing device, the device comprising: The input module is used to obtain the LLVM IR of the GPU kernel to be tested.

[0015] The parsing module is used to parse LLVM IR and extract kernel function signatures, parameter data types, memory layout, and access mode information.

[0016] The code generation module is used to generate GPU host-side driver code and CPU reference implementation code based on the parsed kernel function signatures, parameter data types, memory layout, and access mode information.

[0017] The compilation and execution module is used to compile LLVM IR into GPU executable code and execute host-side driver code.

[0018] The testing and verification module is used to compare the GPU execution results with the CPU reference implementation results to verify the correctness of the kernel.

[0019] A computer device includes a memory and a processor, the memory storing a computer program, and the processor executing the computer program performing the following steps: Obtain the LLVM IR of the GPU kernel to be tested.

[0020] Parse LLVM IR to extract kernel function signatures, parameter data types, memory layout, and access mode information.

[0021] Based on the parsed kernel function signatures, parameter data types, memory layout, and access mode information, GPU host-side driver code and CPU reference implementation code are generated.

[0022] Compile LLVM IR into GPU executable code and execute host-side driver code.

[0023] The GPU execution results are compared with the CPU reference implementation results to verify the correctness of the kernel.

[0024] The aforementioned GPU intermediate representation program testing method, apparatus, and device adopt a technical solution that takes LLVM IR (low-level virtual machine compiler intermediate representation) as the entry point, extracts kernel function signatures, parameter data types, memory layout, and access mode information through automated parsing technology, generates GPU host-side driver code and CPU reference implementation code based on the parsing results, compiles and executes them into GPU executable code, and then compares the execution results to verify the correctness of the kernel. This system automates test case construction, transforming raw intermediate representations into complete test cases executable on real GPU hardware. It systematically analyzes function parameter memory attributes, data flow, and thread mapping relationships to automatically reconstruct tensor shapes and memory layouts. Based on the parsed control flow and memory access patterns, it automatically generates semantically equivalent CPU serial reference implementations. Using LLVM IR as a unified interface, it adapts to multiple compiler backends, generating executable code adapted to different GPU architectures within the same host environment by switching physical assembler toolchains, significantly improving the portability and reusability of test cases. By decoupling from the original compiler runtime environment, it narrows the defect localization scope from the entire compilation chain to a single degradation stage, effectively distinguishing between front-end semantic deviations and back-end optimization errors, greatly shortening the cross-platform verification cycle, and achieving efficient, universal, and accurate GPU compiler correctness verification capabilities. Attached Figure Description

[0025] Figure 1 This is a flowchart illustrating a GPU intermediate representation program testing method in one embodiment; Figure 2 This is a pseudocode diagram of the GPU Kernel correctness verification execution flow with LLVM IR as input in one embodiment; Figure 3 This is a schematic diagram illustrating how a complete host-side driver framework is automatically generated based on a driver code template and kernel feature information in one embodiment. Figure 4 This is a structural block diagram of a GPU intermediate representation program testing device in one embodiment; Figure 5 This is an internal structural diagram of a computer device in one embodiment. Detailed Implementation

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

[0027] In one embodiment, such as Figure 1 As shown, a method for testing GPU intermediate representation programs is provided, including the following steps: Step 102: Obtain the LLVM IR of the GPU kernel to be tested.

[0028] Specifically, obtain the LLVM IR format file of the GPU Kernel to be verified. This LLVM IR can originate from any GPU compiler backend, including but not limited to the XLA (Accelerated Linear Algebra) compiler, MLIR-AIE (MLIR for AI Engine, a domain-specific compiler infrastructure built on the LLVM / MLIR architecture), and TVM (TensorVirtual Machine, an end-to-end deep learning compiler framework). This intermediate representation file contains the complete computational logic of the kernel functions, memory access patterns, thread mapping relationships, and various memory attribute information inserted by the compiler.

[0029] Step 104: Parse LLVM IR to extract kernel function signatures, parameter data types, memory layout, and access mode information.

[0030] Specifically, the process of parsing LLVM IR includes: First, analyzing the memory attributes of function parameters, including the number of dereferenceable bytes, alignment, and alias information; calculating the tensor size using the `dereferenceable` attribute; determining memory alignment requirements using the `align` attribute; and checking for parameter overlap using the `noalias` attribute. Second, identifying the read / write direction of data through load and store instructions to determine the role of each parameter in the computation, where parameters pointed to by load instructions are input parameters and parameters pointed to by store instructions are output parameters. Then, parsing the parallel mapping relationship of data based on thread index instructions, including reading LLVM. Built-in thread index functions such as .read.ptx.sreg.tid.x are used to analyze conditional judgment instructions to filter active threads. The correspondence between threads and data is inferred through division and modulo operations, and finally a complete mapping model between threads and data is constructed.

[0031] Step 106: Based on the parsed kernel function signature, parameter data type, memory layout, and access mode information, generate GPU host-side driver code and CPU reference implementation code.

[0032] Specifically, generating GPU host-side driver code includes: generating GPU host-side driver code based on the parsing results using a large language model. This driver code contains complete logic such as device initialization, memory allocation, input data transmission, kernel function startup, synchronous waiting, and output data return. Generating CPU reference implementation code includes: extracting control flow and memory access patterns from LLVM IR, constructing a loop structure based on thread mapping relationships to simulate the parallel execution semantics of the GPU kernel, converting the loop structure into a serially executed CPU function, and ensuring that the computational logic of the CPU reference implementation is consistent with the GPU kernel. When test cases come from the XLA compiler, input data also needs to be extracted from the constant tensors generated by the XLA compiler. Specifically, this includes: locating constant tensor definition fragments from the XLA compiler compilation artifacts, using a lexical analyzer to perform lexical segmentation on the tensor text to generate a token stream, parsing the token stream based on syntax rules to construct an abstract syntax tree, traversing the syntax tree to extract numerical sequences and converting them into a space-separated format readable by the host driver code, ensuring that the input data during independent testing is consistent with that when the kernel is executed through the XLA compiler.

[0033] Step 108: Compile LLVM IR into GPU executable code and execute the host-side driver code.

[0034] Specifically, LLVM IR is compiled into executable code for the target GPU architecture using a physical assembler, which may include, but is not limited to, ptxas or a platform-specific assembler; host-side driver code is generated and linked with the GPU executable code; the host-side driver code is executed on the target GPU platform, and the driver code automatically completes operations such as device initialization, memory allocation, input data transfer, kernel startup, synchronization waiting, and output data return; the kernel is executed on the real GPU hardware and the execution results at the physical addresses of the video memory are collected.

[0035] Step 110: Compare the GPU execution results with the CPU reference implementation results to verify the correctness of the kernel.

[0036] Specifically, after the GPU finishes execution, the output buffer data is read from the video memory and compared with the output data generated by the CPU reference implementation. If the comparison results are consistent, the kernel is determined to be correct, indicating that the conversion process from LLVM IR to GPU executable code has not introduced semantic errors. If the comparison results are inconsistent, it is determined that there is a semantic deviation or optimization error, which can narrow the scope of defect localization from the entire compilation chain to a single degradation stage, making it easier to troubleshoot and fix subsequent problems. By executing the compilation artifacts on different GPU architectures, the semantic consistency of the compiler on different platforms can be quickly verified without rebuilding a complete compilation environment on each GPU platform, significantly shortening the compiler's testing cycle.

[0037] The aforementioned GPU intermediate representation program testing method adopts an LLVM IR as the entry point, and uses automated parsing technology to extract kernel function signatures, parameter data types, memory layout and access mode information. Based on the parsing results, GPU host-side driver code and CPU reference implementation code are generated, compiled and executed into GPU executable code, and the execution results are compared to verify the correctness of the kernel. This system automates test case construction, transforming raw intermediate representations into complete test cases executable on real GPU hardware. It systematically analyzes function parameter memory attributes, data flow, and thread mapping relationships to automatically reconstruct tensor shapes and memory layouts. Based on the parsed control flow and memory access patterns, it automatically generates semantically equivalent CPU serial reference implementations. Using LLVM IR as a unified interface, it adapts to multiple compiler backends such as XLA and MLIR-AIE. By switching physical assembler toolchains, it generates executable code adapted to different GPU architectures within the same host environment, significantly improving the portability and reusability of test cases. By decoupling from the original compiler runtime environment, it narrows the defect localization scope from the entire compilation chain to a single degradation stage, effectively distinguishing between front-end semantic deviations and back-end optimization errors, greatly shortening the cross-platform verification cycle, and achieving efficient, universal, and accurate GPU compiler correctness verification capabilities.

[0038] In one embodiment, the LLVM IR is parsed, and the memory attributes of the function parameters are analyzed, including the number of dereferenceable bytes, alignment, and alias information, to infer the shape of the input and output tensors. The read and write directions of data are identified through load and store instructions to determine the role of each function parameter in the computation. The parallel mapping relationship of data is parsed based on thread index instructions to construct the correspondence between threads and data.

[0039] In one embodiment, control flow and memory access patterns are extracted based on the parsed kernel function signature, parameter data types, memory layout, and access mode information. A loop structure is constructed based on the thread mapping relationship to simulate the parallel execution semantics of the GPU kernel. The loop structure is then converted into a serially executed CPU function to make the computational logic consistent with the GPU kernel.

[0040] In one embodiment, based on the extracted kernel function signature, parameter data type, memory layout, and access mode information, a large language model is invoked to generate a host driver code template. Thread mapping logic, memory allocation, and data transfer code are embedded into the host driver code template to generate a complete host-side driver program.

[0041] In one embodiment, in the host-side driver, the input data comes from constant tensors generated by the XLA compiler. The extraction process specifically involves: locating the constant tensor definition fragment from the XLA compiler's output, using a lexical analyzer to segment the text of the constant tensor to generate a token stream. The token stream is then parsed according to syntax rules to construct an abstract syntax tree (AST), which is traversed to extract numerical sequences and converted into a format readable by the driver code.

[0042] In one embodiment, after locating the constant tensor definition fragment from the compilation output of the XLA compiler, a lexical analyzer is used to segment the text of the constant tensor and extract its numerical value. ValueList=Flatten(Parse(Tokenize(TensorLiteral))); In this context, ValueList represents the numerical values ​​of constant tensors, Tokenize is the lexical segmentation function, Parse is the syntax parsing function, Flatten is the multidimensional numerical expansion function, and TensorLiteral is a literal data structure used to represent constant tensor data in the Accelerated Linear Algebra Compiler (XLA) Intermediate Representation (HLO).

[0043] In one embodiment, after the GPU executes, it reads output buffer data from video memory and compares the buffer data with the output data generated by the CPU reference implementation. If the comparison is consistent, the GPU kernel program is determined to be correct; otherwise, the GPU kernel is identified as having a semantic deviation or optimization error.

[0044] In one embodiment, such as Figure 2 The pseudocode shown illustrates the GPU Kernel correctness verification execution flow using LLVM IR as input. Taking temp_2061bd0fb3902b38.ll as an example, this file is an LLVM IR that processes the Concatenate / Fusion operator, merging three input tensors and writing them into a single output buffer. The analysis follows the steps outlined in Table 1 below: Table 1 Analysis Steps

[0045] Step 1: Analyze the memory properties of the function prototype LLVM IR function parameters carry compiler-inserted memory attributes: `ptr noalias align 16dereferenceable(80)%0`, as shown in Table 2. Table 2 Memory Attributes

[0046] Comparative analysis of the four parameters: %0: dereferenceable(80) → 20 floats → candidate input A; %1:dereferenceable(320) → 80 floats → Candidate input B; %2:dereferenceable(24) → 6 floats → Candidate input C; %3:dereferenceable(640) → 160 floats → Candidate output (largest size, usually the output); align 256 (Special alignment, implying performance optimization for writing to the buffer).

[0047] Step 2: Determine the data flow direction through memory access instructions. In SSA-style IR, the load and store instructions are the sole source of the data flow: Recognize input: %13=getelementptr inbounds [20 x float], ptr%0, i32 0, i32 %12; %14=load float, ptr%13, align 4, !invariant.load !2.

[0048] Analyzing the two instructions, we can see that `load` reads data from `%13`, and `%13` is calculated from the base address `%0`. Also, `[20x float]` matches the previously calculated 80 bytes. Furthermore, the `!invariant.load` metadata indicates that this memory location remains unchanged during kernel execution, a typical marker of read-only input.

[0049] The output is: %16=getelementptr inbounds [160 x float], ptr%3, i32 0, i32%15; store float %14, ptr%16, align 4.

[0050] The two instructions show that `store` writes `%14` (the value just read from `%0`) to `%16`, and `%16` is calculated from the base address of `%3`. `[160xfloat]` corresponds to 640 bytes, consistent with the dereferenceable (640) of `%3`. Since the instructions performed a write operation on the memory of `%3`, `%3` is determined to be an output parameter.

[0051] Step 3: Inferring data parallel semantics through thread mapping The GPU Kernel's thread indexing logic reflects the mapping relationship between input elements and threads: %5=call i32 @llvm.musa.read.ptx.sreg.tid.x(); Get threadIdx.x; %6=urem i32 %5, 4; lane=tid % 4; %7=icmp eq i32 %6, 0; lane==0; %8=icmp sle i32 %5, 76; tid<=76; %9=and i1 %7, %8; Only threads in lane 0 and tid<=76 are active.

[0052] Semantic reasoning: Active thread calculation: The condition tid%4==0 is used to filter out the 0th thread in each warp.

[0053] The condition tid<=76 limits the total number of threads; Active thread IDs: 0, 4, 8, …, 76 → a total of 20 threads.

[0054] Input index calculation: %11=udiv i32 %5, 4; idx=tid / 4 → range [0, 19] %13=getelementptr ..., i32 %11; Accessing in0[idx] 20 threads access in0[0] to in0

[19] respectively. Mapping relationship: 1 thread 1. Input elements.

[0055] Output index calculation: %15 = mul i32 %11, 8; out_idx = idx 8%16 = getelementptr ..., i32%15 outputs values ​​at positions 0, 8, 16, …, 152 (step size 8). Data reorganization mode: Scatters the compact input in0 to the stride-8 position of the output.

[0056] The aforementioned semantics are used to implement the CPU Reference (a serial implementation of the CPU version used for comparison and verification). In practice, an IR file containing the aforementioned shape reasoning elements and semantic reasoning elements, along with corresponding host-side driver code, is provided as the prompt word for the single-round prompt word project. The host-side driver code contains the CPU reference implementation corresponding to the IR semantics and the call template for the target platform's driver API. The prompt word and the IR to be analyzed are input into a large language model, which then generates new host-side driver code according to the analysis method described above. Next, the IR and the generated driver are compiled, and finally, the IR is run to verify the consistency between the GPU and CPU results.

[0057] When test cases originate from the XLA compiler, the constant tensor input data generated by the XLA compiler for the kernel can be extracted to ensure that the input data is consistent with that when the kernel is executed through the XLA compiler during subsequent independent kernel testing.

[0058] HLO constant tensor localization: dump the input data generated by the XLA compiler, identify and extract constant tensor definition fragments (such as multidimensional array literals in the format f32[128,256]{1,0}), and separate the shape metadata (data type, dimension) from the numerical content; Lexical Analysis: The XlaTensorLexer generated by ANTLR4 is used to perform lexical segmentation on the extracted tensor text, identify tokens such as numeric values, type identifiers, dimension separators, and parentheses, and generate a normalized token stream. Syntax Analysis: Based on custom XlaTensor.g4 syntax rules, XlaTensorParser is used to parse the Token stream into an Abstract Syntax Tree (AST), establishing a hierarchical structure of root node (tensor) → shape node (shape) → value list node (valueList); AST Traversal and Numerical Extraction: ParseTreeWalker is used to traverse the AST, locate the semantic nodes of valueList, and recursively extract all numerical literals in dimensional order, maintaining the original data precision and dimensional arrangement; Standardized format conversion: Converts the extracted numerical sequence into a space-separated format that can be directly read sequentially by the host driver code.

[0059] In one embodiment, such as Figure 3As shown, this paper presents a method that utilizes a large language model to automatically generate complete host-side driver pseudocode based on driver code templates and kernel feature information. Taking the file temp_2061bd0fb3902b38.ll as an example, this file is an LLVM IR that processes the Concatenate / Fusion operator, merging three input tensors and writing them into an output buffer. The prompt words containing the contents of this IR file, the driver code template, and the CPU reference implementation framework are then input into the large language model.

[0060] The large language model first parses the function signature `void @input_concatenate_fusion(ptr%0,ptr%1, ptr%2, ptr%3)` in the IR, identifying four pointer parameters. Through the attributes `dereferenceable(80),` `dereferenceable(320),` `dereferenceable(24),` and `dereferenceable(640)`, combined with the 4-byte size of the float type, it infers that `%0` contains 20 floats, `%1` contains 80 floats, `%2` contains 6 floats, and `%3` contains 160 floats. The `align 16` and `align 256` attributes determine that the input parameters are 16-byte aligned and the output parameters are 256-byte aligned. By analyzing the instructions `loadfloat, ptr%13` and `store float %14, ptr%16`, it determines that `%0`, `%1`, and `%2` are input parameters, and `%3` is the output parameter. Finally, by analyzing the thread index function `llvm.musa.read.ptx.sreg.tid.x` and the conditional instruction `%6 = icmpule i32 %5`, 76 and %7 = icmp eq i32 %8, 0, division instruction %11 = udiv i32 %5, 4, parsing out 20 active thread IDs of 0, 4, 8... respectively, accessing in0[0] to in0

[19] , outputting positions of 0, 8, 16...152 with a step size of 8.

[0061] Based on the above analysis results, the host-side driver code generated by the large language model includes: allocating input buffers of 20 floats, 80 floats, and 6 floats, and an output buffer of 160 floats; generating test input data or loading it from a constant tensor; configuring the kernel startup parameters as a one-dimensional thread block of 1 grid and 76 threads; reading the results from the output buffer after the kernel starts; and generating the CPU reference implementation code as for (int idx=0; idx<20; idx++) { out[idx 8]=in0[idx];} and similar in1, in2 processing logic; finally, compare the consistency between GPU output and CPU output.

[0062] It is worth noting that the typical host-side driver code framework has good scalability. It can be adapted to different GPU platforms by replacing the underlying driver API. For example, by replacing the CUDA driver API with the MUSA driver API, a smooth migration from the NVIDIA platform to the Moore Threads platform can be achieved. Different thread mapping modes can be adapted by modifying the kernel startup configuration, such as extending from one-dimensional thread blocks to two-dimensional or three-dimensional thread blocks. More complex operator semantics can be adapted by expanding the reference implementation module, such as supporting advanced features such as atomic operations, cooperative groups, and dynamic parallelism.

[0063] Furthermore, the initialization module is responsible for selecting and initializing the GPU device. This module first calls the device query function of the driver API to obtain a list of available GPU devices in the system and their attributes, including device name, computing power, memory size, maximum number of threads, etc.; then, it selects the target device according to a user-specified or automatic selection strategy; next, it calls the device setup function to create the GPU context and set runtime attributes such as working directory and cache configuration; finally, it verifies whether the device status is ready to ensure that subsequent operations can be executed normally. This module also includes error handling logic, which logs error information and terminates the test process when the device is unavailable or initialization fails.

[0064] It should be understood that, although Figure 1 The steps in the flowchart are shown sequentially as indicated by the arrows, but these steps are not necessarily executed in the order indicated by the arrows. Unless otherwise specified herein, there is no strict order in which these steps are executed, and they can be performed in other orders. Figure 1 At least some of the steps in the process 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 executed in turn or alternately with other steps or at least some of the sub-steps or stages of other steps.

[0065] In one embodiment, such as Figure 4 As shown, a GPU intermediate representation program testing device is provided, including: an input module 402, a parsing module 404, a code generation module 406, a compilation and execution module 408, and a testing and verification module 410, wherein: Input module 402 is used to obtain the LLVM IR of the GPU kernel to be tested.

[0066] The parsing module 404 is used to parse LLVM IR and extract kernel function signatures, parameter data types, memory layout, and access mode information.

[0067] The code generation module 406 is used to generate GPU host-side driver code and CPU reference implementation code based on the parsed kernel function signature, parameter data type, memory layout and access mode information.

[0068] The compilation and execution module 408 is used to compile LLVM IR into GPU executable code and execute host-side driver code.

[0069] The test and verification module 410 is used to compare the GPU execution results with the CPU reference implementation results to verify the correctness of the kernel.

[0070] Specific limitations regarding the GPU intermediate representation program testing device can be found in the limitations of the GPU intermediate representation program testing method described above, and will not be repeated here. Each module in the aforementioned GPU intermediate representation program testing device can be implemented entirely or partially through software, hardware, or a combination thereof. These modules can be embedded in hardware or independently of the processor in the computer device, or stored in software in the memory of the computer device, so that the processor can call and execute the operations corresponding to each module.

[0071] In one embodiment, a computer device is provided, which may be a terminal, and its internal structure diagram may be as follows: Figure 5 As shown, the computer device includes a processor, memory, network interface, display screen, and input devices connected via a system bus. The processor provides computing and control capabilities. The memory includes non-volatile storage media and internal memory. The non-volatile storage media stores the operating system and computer programs. The internal memory provides an environment for the operation of the operating system and computer programs stored in the non-volatile storage media. The network interface is used to communicate with external terminals via a network connection. When the computer program is executed by the processor, it implements a GPU intermediate representation program testing method. The display screen can be an LCD screen or an e-ink screen. The input devices can be a touch layer covering the display screen, buttons, a trackball, or a touchpad mounted on the computer device casing, or an external keyboard, touchpad, or mouse.

[0072] Those skilled in the art will understand that Figures 4-5The structure shown is merely a block diagram of a portion of the structure related to the present invention and does not constitute a limitation on the computer device to which the present invention is applied. A specific computer device may include more or fewer components than those shown in the figure, or combine certain components, or have different component arrangements.

[0073] In one embodiment, a computer device is provided, including a memory and a processor, the memory storing a computer program, the processor executing the computer program to perform the following steps: Obtain the LLVM IR of the GPU kernel to be tested.

[0074] Parse LLVM IR to extract kernel function signatures, parameter data types, memory layout, and access mode information.

[0075] Based on the parsed kernel function signatures, parameter data types, memory layout, and access mode information, GPU host-side driver code and CPU reference implementation code are generated.

[0076] Compile LLVM IR into GPU executable code and execute host-side driver code.

[0077] The GPU execution results are compared with the CPU reference implementation results to verify the correctness of the kernel.

[0078] Those skilled in the art will understand that all or part of the processes in the methods of the above embodiments can be implemented by a computer program instructing related hardware. The computer program can be stored in a non-volatile computer-readable storage medium. When executed, the computer program can include the processes of the embodiments of the above methods. Any references to memory, storage, databases, or other media used in the embodiments provided by this invention can include non-volatile and / or volatile memory. Non-volatile memory can include read-only memory (ROM), programmable ROM (PROM), electrically programmable ROM (EPROM), electrically erasable programmable ROM (EEPROM), or flash memory. Volatile memory can include random access memory (RAM) or external cache memory. By way of illustration and not limitation, RAM is available in various forms, such as static RAM (SRAM), dynamic RAM (DRAM), synchronous DRAM (SDRAM), dual data rate SDRAM (DDRSDRAM), enhanced SDRAM (ESDRAM), Synchlink, DRAM (SLDRAM), RAMbus direct RAM (RDRAM), direct memory bus dynamic RAM (DRDRAM), and RAMbus dynamic RAM (RDRAM), etc.

[0079] The technical features of the above embodiments can be combined in any way. For the sake of brevity, not all possible combinations of the technical features in the above embodiments are described. However, as long as there is no contradiction in the combination of these technical features, they should be considered to be within the scope of this specification.

[0080] The embodiments described above are merely illustrative of several implementations of the present invention, and while the descriptions are specific and detailed, they should not be construed as limiting the scope of the invention. It should be noted that those skilled in the art can make various modifications and improvements without departing from the concept of the present invention, and these modifications and improvements all fall within the scope of protection of the present invention. Therefore, the scope of protection of the present invention should be determined by the appended claims.

Claims

1. A GPU intermediate representation program testing method, characterized by, The method includes: Obtain the LLVM IR of the GPU kernel to be tested; Parse the LLVM IR to extract kernel function signatures, parameter data types, memory layout, and access mode information; Based on the parsed kernel function signature, parameter data type, memory layout, and access mode information, GPU host-side driver code and CPU reference implementation code are generated. The LLVM IR is compiled into GPU executable code, and the host-side driver code is executed. The GPU execution results are compared with the CPU reference implementation results to verify the correctness of the kernel.

2. The method of claim 1, wherein, Parse the LLVM IR to extract kernel function signatures, parameter data types, memory layout, and access mode information, including: The LLVM IR is parsed, and the memory properties of the function parameters are analyzed, including the number of dereferenceable bytes, alignment, and alias information, in order to infer the shape of the input and output tensors; The read / write direction of data is identified by loading and storing instructions, and the role of each function parameter in the calculation is determined. The parallel mapping relationship of data is parsed based on the thread index instruction, and the correspondence between threads and data is constructed.

3. The method of claim 1, wherein, Based on the parsed kernel function signature, parameter data type, memory layout, and access mode information, GPU host-side driver code and CPU reference implementation code are generated, including: Based on the parsed kernel function signature, parameter data type, memory layout, and access mode information, extract the control flow and memory access mode; A loop structure is constructed based on the thread mapping relationship to simulate the parallel execution semantics of the GPU kernel; The loop structure is converted into a CPU function that is executed serially, so that the computational logic is consistent with the GPU kernel.

4. The method according to any one of claims 1 to 3, characterized in that, The LLVM IR is compiled into GPU executable code, and the host-side driver code is executed. Based on the extracted kernel function signature, parameter data type, memory layout, and access mode information, the large language model is invoked to generate a host driver code template; The thread mapping logic, memory allocation, and data transfer code are embedded into the host driver code template to generate a complete host-side driver program.

5. The method of claim 4, wherein, In the host-side driver, the input data comes from constant tensors generated by the XLA compiler. The extraction process is as follows: after locating the constant tensor definition fragment from the compilation output of the XLA compiler, the lexical analyzer is used to segment the text of the constant tensor to generate a token stream. The token stream is parsed according to the syntax rules, an abstract syntax tree is constructed, the abstract syntax tree is traversed to extract the numerical sequence, and it is converted into a format that can be read by the driver code.

6. The method according to claim 5, characterized in that, After locating the constant tensor definition fragment from the compilation output of the XLA compiler, the text of the constant tensor is segmented using a lexical analyzer to generate a token stream, including: After locating the constant tensor definition fragment from the compilation output of the XLA compiler, the text of the constant tensor is segmented using a lexical analyzer to extract the numerical value of the constant tensor: ValueList=Flatten(Parse(Tokenize(TensorLiteral))) Among them, ValueList is the numerical value of constant tensor, Tokenize is the lexical segmentation function, Parse is the syntax parsing function, Flatten is the multidimensional numerical expansion function, and TensorLiteral is the literal data structure.

7. The method according to claim 6, characterized in that, The kernel's correctness is verified by comparing the GPU execution results with the CPU reference implementation results, including: After the GPU executes, it reads the output buffer data from the video memory and compares the buffer data with the output data generated by the CPU reference implementation. If the comparison is consistent, the GPU kernel program is determined to be correct; otherwise, the GPU kernel is identified as having a semantic deviation or optimization error.

8. A GPU intermediate representation program testing device, characterized in that, The device includes: The input module is used to obtain the LLVM IR of the GPU kernel to be tested; The parsing module is used to parse the LLVM IR and extract kernel function signatures, parameter data types, memory layout, and access mode information. The code generation module is used to generate GPU host-side driver code and CPU reference implementation code based on the parsed kernel function signature, parameter data type, memory layout and access mode information. The compilation and execution module is used to compile the LLVM IR into GPU executable code and execute the host-side driver code; The testing and verification module is used to compare the GPU execution results with the CPU reference implementation results to verify the correctness of the kernel.

9. A computer device comprising a memory and a processor, wherein the memory stores a computer program, characterized in that, When the processor executes the computer program, it implements the steps of the method according to any one of claims 1 to 7.

Images