A method for memory access safety checks in GPU compilers

Abstract

This invention provides a memory access safety check method for GPU compilers. It converts GPU code into high-level IR instructions; extracts the boundaries of each memory allocation based on metadata information in the high-level IR instructions, and converts the high-level IR instructions into low-level IR instructions; based on a transfer function range analysis method, it performs data flow analysis on the low-level IR instructions in the compiler backend to calculate the symbolic address range of each memory access; and checks the safety of each memory access based on the symbolic address range of each memory access and the boundaries of each memory allocation. This invention provides a method for memory access safety checks at the low-level IR language layer in GPU compilers, used to check for potential memory out-of-bounds access or buffer overflow problems that may occur after multiple optimizations in the compiler backend. It is independent of specific hardware architecture and specific IR language, and has a certain degree of universality.

Description

Technical Field

[0001] This invention relates to the field of GPU compilers, and more specifically, to a method for memory access security checks for GPU compilers. Background Technology

[0002] In recent years, with continuous investment in high-performance computing hardware, domestically produced GPU chips have made significant progress. Domestically produced GPU products are now widely used in scenarios such as artificial intelligence training, graphics rendering, and scientific computing. Especially driven by large language models and intelligent computing platforms like DeepSeek, the demand for parallel computing capabilities has surged to unprecedented levels. GPUs, with their thousands of computing cores and superior throughput, have become a core support for the domestic AI computing power ecosystem.

[0003] However, the high parallelism of GPU architecture also brings significant challenges to software development and runtime security. Among these, memory safety issues are becoming increasingly prominent. GPU kernel programs often need to execute concurrently across tens of thousands of thread instances, which frequently access global memory, local shared memory, and private thread memory. Under the complex thread scheduling and address calculation logic, security vulnerabilities such as buffer overflows and illegal address access are highly likely to occur, leading to unstable model training results, system crashes, and even security risks. Summary of the Invention

[0004] This invention addresses the technical problems existing in the prior art by providing a memory access security check method for GPU compilers. It overcomes the problem that under complex thread scheduling and address calculation logic, security vulnerabilities such as buffer overflows and illegal address access are prone to occur, which can lead to system crashes or even security risks.

[0005] This invention provides a method for memory access security checking in a GPU compiler, comprising:

[0006] Step 1: Use a standard toolchain to convert GPU code into high-level IR instructions, which include metadata information;

[0007] Step 2: Based on the metadata information in the high-level IR instructions, extract and symbolize the boundary and runtime information of each memory allocation, and convert the high-level IR instructions into low-level IR instructions;

[0008] Step 3: Based on the range analysis method of the transfer function, in the compiler backend, data flow analysis is performed on the low-level IR instructions to calculate the symbolic address range of each memory access. Based on the boundary of each memory allocation and runtime information in the metadata information, the security of each memory access is checked.

[0009] Step 4: Mark memory accesses that pose a security risk and generate a security analysis report.

[0010] This invention provides a memory access safety check method for GPU compilers. The method converts GPU code into high-level IR instructions; extracts the boundaries of each memory allocation based on metadata information in the high-level IR instructions, and converts the high-level IR instructions into low-level IR instructions; based on a transfer function range analysis method, data flow analysis is performed on the low-level IR instructions in the compiler backend to calculate the symbolic address range of each memory access; and the safety of each memory access is checked based on the symbolic address range of each memory access and the boundaries of each memory allocation. This invention provides a method for memory access safety checks at the low-level IR language layer in GPU compilers, used to check for potential memory out-of-bounds access or buffer overflow problems that may occur after multiple optimizations in the compiler backend. It is independent of specific hardware architecture and specific IR language, and has a certain degree of versatility. Attached Figure Description

[0011] Figure 1 A flowchart of a memory access security check method for a GPU compiler is provided as an embodiment of the present invention;

[0012] Figure 2 This is a flowchart illustrating the register state information update process according to an embodiment of the present invention.

[0013] Figure 3 This is a flowchart illustrating a memory access out-of-bounds check according to an embodiment of the present invention. Detailed Implementation

[0014] To make the objectives, technical solutions, and advantages of the embodiments of the present invention clearer, the technical solutions of the embodiments of the present invention will be clearly and completely described below with reference to the accompanying drawings. Obviously, the described embodiments are only some embodiments of the present invention, not all embodiments. Based on the embodiments of the present invention, all other embodiments obtained by those skilled in the art without creative effort are within the scope of protection of the present invention. In addition, the technical features of the various embodiments or individual embodiments provided by the present invention can be arbitrarily combined with each other to form feasible technical solutions. Such combinations are not constrained by the order of steps and / or structural composition patterns, but must be based on the ability of those skilled in the art to implement them. When the combination of technical solutions is contradictory or cannot be implemented, it should be considered that such a combination of technical solutions does not exist and is not within the scope of protection claimed by the present invention.

[0015] A GPU compiler is a special software program that translates high-level code (such as CUDA C++, OpenCL, HLSL) written by programmers for running on a GPU (graphics processing unit) into low-level instructions (machine code) that the GPU hardware can directly understand and execute.

[0016] In detail, the GPU compiler pipeline is a multi-stage process that translates high-level parallel code (such as CUDA C++ or GLSL) into machine code executable by the specific GPU hardware. This process begins at the front end, which parses the source code and generates a generic, hardware-independent intermediate representation (IR), such as LLVM IR or SPIR-V. These IRs then enter the optimization phase, where a series of hardware-independent optimization passes (such as function inlining and dead code elimination) are executed to improve the general efficiency of the code. Next, at the back end, this generic IR is translated into a vendor-specific IR (such as NVIDIA's PTX). Here, a series of hardware-dependent back-end passes play a crucial role, performing instruction scheduling to hide memory latency, performing precise register allocation to improve GPU utilization, merging memory accesses to maximize bandwidth, and handling branch divergence issues. Finally, the fully optimized vendor IR is compiled into the final, GPU-specific native machine code (such as SASS). This final compilation step is typically performed by the GPU driver at runtime (JIT), ensuring that the code runs at optimal performance on the user's specific hardware.

[0017] The reason why a dedicated GPU compiler is needed stems from the fundamental architectural differences between CPUs and GPUs:

[0018] CPU (Central Processing Unit): Low latency, high single-core performance. Excels at handling complex logic, branching, and sequential tasks. It has a small number of cores (a few to dozens), but each core is very powerful, possessing a complex control unit and a large amount of cache.

[0019] GPU (Graphics Processing Unit): High throughput, massively parallel computing. It excels at handling a large number of simple computational tasks with similar structures that can be performed simultaneously. It has hundreds or thousands of simple computational cores, but its control unit and cache are relatively simple.

[0020] In domestically developed GPU compilation toolchains, numerous optimization passes are performed in the GPU compiler backend to optimize the IR code in order to improve program execution efficiency. Problems can arise during compiler source code parsing, frontend code transformation, and backend optimization passes, leading to errors in the compiled program. Among these, issues caused by memory access problems are particularly serious, easily causing system instability and making it difficult to trace the root cause.

[0021] To analyze and examine problems caused by memory access, this invention proposes a method for performing memory security analysis on low-level memory access (IR) in the GPU compiler backend.

[0022] See Figure 1 The present invention provides a method for memory access security checking for GPU compilers, comprising the following steps:

[0023] Step 1: Use a standard toolchain to convert GPU code into high-level IR instructions, which include metadata information.

[0024] Intermediate representation (IR) is a data structure or code used by the compiler to represent source code. It serves as a program representation between the source and target languages ​​during the compilation process. Almost all compilers require some form of intermediate representation to model the code being analyzed, transformed, and optimized. During compilation, the intermediate representation must be sufficiently expressive to accurately represent the source code without losing information, and it must fully consider the completeness of the compilation from source code to target code, the ease of use of compilation optimization, and performance.

[0025] Step 1, which involves using a standard toolchain to convert GPU code into high-level IR instructions, includes:

[0026] Step 11: Convert the GPU code into high-level IR instructions (such as SPIR-V) using a standard compiler front-end such as Clang. The high-level IR instructions already contain the core metadata required for subsequent processing.

[0027] Step 12: Establish a mapping from high-level IR instructions to the original source code files and line numbers, which will be used as debugging information later.

[0028] Step 2: Based on the metadata information in the high-level IR instructions, extract and symbolize the boundary and runtime information of each memory allocation, and convert the high-level IR instructions into low-level IR instructions.

[0029] Understandably, after converting the GPU source code into high-level IR instructions, metadata is extracted from the high-level IR instructions, including memory object metadata extraction and runtime information metadata extraction.

[0030] Step 2 includes:

[0031] Step 21, Extracting Metadata from Memory Objects: Based on the metadata information in the high-level IR instructions, create a unique identifier for each memory allocation in the code and record its boundary size. Different identifiers are used to record the boundary size for different types of memory spaces; for pointers to global address spaces of unknown size, a symbol (such as size(@buffer_A)) is used to represent their boundaries; for pointers to local address spaces and private address spaces of known size, their exact byte size (such as 1024) is recorded.

[0032] The following section will introduce the different types of memory spaces.

[0033] Global Memory:

[0034] Global memory has a large amount of memory, and its access speed is slower than other types of memory such as shared memory and registers. All running threads can read and write to global memory, and the CPU can also read and write to global memory. Global memory is allocated and released by the host.

[0035] Constant Memory:

[0036] This memory is also part of the GPU's main memory; it has its own cache and is independent of the global L1 and L2 caches. All threads can access the same constant memory, but they can only read it, not write to it. The CPU sets the value in the constant memory before starting the kernel.

[0037] Shared Memory:

[0038] Shared memory is used to enable fast communication between threads within a thread block. Shared memory exists only during the lifetime of the block.

[0039] Local Memory:

[0040] Local memory is also part of the GPU's main memory (the same as global memory), so it is typically slow. Threads automatically use local memory when registers run out or become unavailable. This is called a register overflow. This occurs if each thread has too many variables to use registers or if the kernel uses structures. Additionally, arrays without constant indices use local memory because registers don't have addresses and must use addressable memory space. The scope of local memory is per thread.

[0041] Step 22, Runtime Information Metadata Extraction: Based on the metadata information in the high-level IR instructions, an initial value range is assigned to built-in functions such as get_global_id(). Generally, this range is symbolic, such as the range of get_local_id(0) being [0, get_local_size(0) - 1]. When there are additional constraints, such as (reqd_work_group_size(128, ...)), this information is used to refine the range to a constant, such as the range of get_local_id(0) becoming more precise [0, 127].

[0042] Step 23: After extracting metadata information from the high-level IR instructions, convert the high-level IR instructions into low-level IR instructions.

[0043] Step 3: Based on the range analysis method of the transfer function, data flow analysis is performed on the low-level IR instructions in the compiler backend to calculate the symbolic address range of each memory access. Based on the boundary of each memory allocation and runtime information in the metadata information, the security of each memory access is checked.

[0044] Understandably, in the compiler domain, a Pass is an independent, modular processing unit. It traverses the intermediate representation (IR) of the program and performs a specific task, which can be:

[0045] (1) Analysis Pass: Examine the code, collect information, but do not modify the code. For example, a pass might be used to analyze which variables in a loop are no longer used outside the loop.

[0046] (2) Transform Pass: Modify code based on analysis results or specific rules. For example, a pass might be used to remove calculations of variables that are no longer used outside of loops (dead code elimination).

[0047] In this embodiment of the invention, a data flow analysis pass is added during the compiler backend process. This pass performs data flow analysis on the low-level IR in step S2, and its core idea is range-based pointer analysis.

[0048] In one embodiment of the present invention, step 3 specifically includes the following steps:

[0049] Step 31: Define register status information.

[0050] Registers are the fastest memory on the GPU. Variables declared in the kernel will use registers unless they run out or cannot be stored in registers, in which case local memory will be used.

[0051] To complete the data flow analysis, the register status information is defined as follows:

[0052] State(%reg) = { BasePointerID, AddressSpace, OffsetRange}

[0053] in:

[0054] State(%reg): Represents the status information of register %reg.

[0055] BasePointerID: Represents a unique identifier, indicating the base address of the pointer.

[0056] AddressSpace: Indicates the memory space type, which can be private, global, local, or NonPointer.

[0057] OffsetRange: Represents pointer offset [min_offset, max_offset], which is the minimum and maximum byte offset of the pointer relative to BasePointerID.

[0058] Step 32: Obtain the metadata information, the low-level IR instructions in the current state of the compiler backend, and the corresponding control flow graph (CFG). Perform data flow analysis based on the control flow graph (CFG) and update the register state information.

[0059] In one embodiment, step 32 specifically includes the following steps:

[0060] Step 321: Allocate storage space for each basic block B in the control flow graph CFG to record its entry state IN[B] and exit state OUT[B].

[0061] For each basic block in the control flow graph (CFG), storage space is allocated to record its entry state IN[B] and exit state OUT[B]. For the entry block of the CFG, its IN state is set to an initial condition, which represents prior knowledge before the kernel or function execution. For all other basic blocks, their IN and OUT states are initialized to the "top element," which represents the "least information" or "most uncertain" state.

[0062] Step 322: Create a work list, which is used to store the data structure of the basic blocks to be processed.

[0063] Step 323: After initialization, place the entry basic blocks in the control flow graph (CFG) into the working list.

[0064] Understandably, a worklist is created. The worklist is a data structure used to store basic blocks to be processed. Initially, the entry basic blocks in the control flow graph (CFG) are added to the worklist to start the iteration process.

[0065] Step 324: Take a currently pending basic block from the work list, denoted as B_current.

[0066] Step 325: Calculate the new entry state IN[B_current] for B_current.

[0067] Specifically, the entry state of a basic block is calculated, which is a collection of the exit state information of all its predecessor nodes. For a selected basic block B_current, its entry state IN[B_current] is calculated as follows:

[0068] In the control flow graph (CFG), all basic blocks that directly point to B_current are identified, i.e., their predecessor node set Predecessors(B_current). For the exit state OUT[P] of all predecessor nodes P ∈ Predecessors(B_current), a predefined merge operator (Meet / Join Operator) is applied. This operator merges the data flow information from multiple input paths into a single state. Its mathematical representation is:

[0069] IN[B_current] = Meet(OUT[P_1], OUT[P_2], ..., OUT[P_i], ..., OUT[P_n]);

[0070] Where P_i ∈ Predecessors(B_current), and n is the number of predecessor basic blocks of B_current.

[0071] It's important to note that, for example, when analyzing an if-else branch structure, traditional methods split into two independent paths for exploration. If the branches are nested, the number of paths increases exponentially. However, this invention performs a state merge operation at the branch convergence point. If the range of a variable x after the if branch is [10, 10], and after the else branch it's [100, 100], then the merged unified state is [10, 100]. Subsequent analysis only needs to continue based on this more general range of [10, 100], without needing to track two separate paths. This method of using a single abstract state to cover multiple possibilities fundamentally avoids the exponential growth of the number of paths, ensuring that even when facing complex control flows (such as multiple branches and loops), the analysis can be completed efficiently within a predictable timeframe, avoiding the path explosion of traditional methods and improving operational efficiency.

[0072] Step 326: Based on the new entry state IN[B_current] of B_current, which is B_current, calculate the new exit state OUT[B_current] and update the register state information.

[0073] Understandably, the exit state of a basic block is determined by processing the instructions within that block, based on its entry state. During the execution of instructions within the basic block, the register states are updated.

[0074] Specifically, all instructions in the basic block B_current are traversed and processed, and the exit state OUT[B_current] of B_current is calculated. For each instruction type in the basic block B_current, a transfer function is defined to complete the register state update. The workflow for defining a transfer function to complete the register state update for each instruction type in the basic block can be found in [link to documentation]. Figure 2 .

[0075] Apart from memory access instructions, each instruction requires a corresponding transfer function to update the register state. The pseudo-instructions and state update methods for typical instructions are shown in Table 1.

[0076] Table 1 Typical instructions and status update methods

[0077]

[0078] Table 1 shows various instruction types and their methods for updating register states.

[0079] Step 327: Check whether the calculated new exit state of B_current causes a change in the state of subsequent basic blocks. If so, store the data structure of the subsequent basic blocks of B_current into the working list and remove B_current from the working list.

[0080] Understandably, after calculating the new exit state of the current B_current, it is determined whether the new exit state of B_current causes a change in the state of subsequent basic blocks. If so, the data structure of the subsequent basic blocks of B_current is stored in the working list so that the new entry and exit states of the subsequent basic blocks in the working list can be calculated later. After calculating the new entry and exit states of the current basic block B_current, the B_current basic block is removed from the working list.

[0081] Step 33: When a memory access instruction is encountered, obtain the current register status information.

[0082] Step 34: Obtain the corresponding memory allocation boundary based on the base address and memory space type in the current register status information;

[0083] Step 35: Determine whether the pointer offset is out of bounds based on the pointer offset in the current register status information and the corresponding memory allocation boundary.

[0084] Understandably, during the execution of each instruction in a basic block, a safety check is performed when a memory access instruction (such as load / store) is encountered. The process of the memory access safety check is as follows: Figure 3 As shown.

[0085] Specifically, obtain the current status information of the register to get {BasePointerID, AddressSpace,OffsetRange}, where OffsetRange=[min_offset, max_offset].

[0086] Based on the base address BasePointerID and memory space type, find the corresponding memory allocation boundary information [0, alloc_size] from the metadata information. Perform checks based on the information available at this time:

[0087] (1) Check the lower bound of the pointer offset: whether min_offset >= 0 is satisfied. If it is satisfied, the lower bound has not been exceeded.

[0088] If min_offset < 0 and min_offset is a negative constant (e.g., -1), report an out-of-bounds error; if min_offset depends on a symbol (e.g., sym_id_10), report a potential out-of-bounds warning.

[0089] (2) Check the upper bound of the pointer offset: whether max_offset + access_size <= alloc_size. If it is satisfied, the upper bound has not exceeded the limit.

[0090] For private and local memory, the memory boundary `alloc_size` is a compile-time constant. It can be directly calculated and compared: for example, `max_offset + access_size > alloc_size` will report a clear out-of-bounds error.

[0091] If global memory is used, alloc_size is a symbol, such as size(@global_buffer). The check is a proof of a symbolic inequality. For example, max_offset might be (global_size_0 - 1) × 4. The check is (global_size_0 - 1) × 4 + 4 <= size(@buffer_A), i.e., global_size_0 × 4 <= size(@buffer_A). If the compiler can prove this inequality always holds (which might be possible in some simple cases), then it is safe. Otherwise, a conditional warning is reported: "Out-of-bounds access may have occurred when the startup parameter does not satisfy global_work_size[0] × 4 <= size_of_buffer_A".

[0092] If the lower and upper bounds are not exceeded after the check, the address space is then checked. Memory access instructions typically specify the address space, calculating the address range based on the pointer base address and pointer offset in the register status. For example, the current register status information is...

[0093] {@local_buf_a, local_address, [0, 128]}, where @local_buf_a is the pointer base address and [0, 128] is the pointer offset, so the address range is [local_buf_a , local_buf_a + 128].

[0094] The calculated address range is matched against the specified address space range. If they do not match, it indicates that the pointer is being referenced incorrectly. The address space ranges for different types are determined at compile time (the ranges for private and local address spaces are determined at compile time).

[0095] For example, if the local address space is specified by the compiler as [2048, 4096], then all local type pointers must be within this range. If the calculated address range of the pointer is [2048, 6144], then it may exceed the specified address space range of the local type, resulting in an incomplete match and a security risk.

[0096] During data flow analysis, it's crucial to determine whether the analysis has reached convergence, i.e., a stable state. The specific method for determining this is as follows:

[0097] For each basic block in the control flow graph CFG, the newly calculated exit state OUT_new[B_current] is compared with the previously stored old exit state OUT_old[B_current].

[0098] If OUT_new and OUT_old are different, it indicates that the data flow information has changed and the data flow analysis has not yet converged.

[0099] Update the exit state of B_current to OUT_new, and add the data structures of all successor nodes of B_current in the control flow graph CFG to the working list;

[0100] If OUT_new is the same as OUT_old, no operation is performed;

[0101] When the work list is empty, it means that the state of all basic blocks has stabilized and no longer changes, the system has reached a fixed point, and the iteration process terminates.

[0102] Step 4: Mark memory accesses that pose a security risk and generate a security analysis report.

[0103] Based on the security check results of memory access obtained through the above steps, a security analysis report is generated. Specifically, this includes the following steps:

[0104] Step 41: Extract all instances marked as risk and their associated context data from the analyzed code.

[0105] Step 42: Assign a stable and unique ID to each individual risk instance for continuous tracking. Classify the risk into high (deterministic error), medium (potential risk), or low (informative) levels based on its certainty.

[0106] Step 43 precisely maps the risk points in the low-level IR back to their filenames and line numbers in the high-level source code. Then, a concise and clear text is generated programmatically to explain the nature of the risk and its potential triggering conditions.

[0107] Step 44: Compile all generated risk records into a structured file (such as JSON or SARIF format) and output it.

[0108] The memory access security check method for GPU compilers provided by this invention has the following advantages:

[0109] (1) The analysis process of this invention directly acts on the private low-level IR after the back-end processing of the domestic GPU compiler. The object of analysis is the final code that is closer to the actual execution of the hardware. Compared with all existing technologies that perform analysis at the high-level IR (such as LLVM IR), this invention can discover and report memory security issues introduced or exposed at the back-end optimization stage of the domestic GPU manufacturer's compiler, thus eliminating the blind spots of existing technologies.

[0110] (2) Before analysis, this invention captures the symbolic boundaries of memory objects and key variables of the execution environment through a systematic metadata extraction and symbolization mechanism. During security checks, an inequality composed of symbols is constructed based on the precise context for judgment, which improves the accuracy of the judgment.

[0111] (3) The core algorithm of this invention uses symbolic range analysis instead of heavyweight, path-by-path symbolic execution. When encountering branches or loops, the states of multiple paths are merged into a unified abstract range, avoiding the exponential growth of the number of paths. This enables the analysis of complex control flows to be completed in a predictable and reasonable time, resulting in high operating efficiency.

[0112] It should be noted that the descriptions of each embodiment in the above embodiments have different focuses. For parts that are not described in detail in a certain embodiment, please refer to the relevant descriptions in other embodiments.

[0113] Those skilled in the art will understand that embodiments of the present invention can be provided as methods, systems, or computer program products. Therefore, the present invention can take the form of a completely hardware embodiment, a completely software embodiment, or an embodiment combining software and hardware aspects. Furthermore, the present invention 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.

[0114] This invention is described with reference to flowchart illustrations and / or block diagrams of methods, apparatus (systems), and computer program products according to embodiments of the invention. It will 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 computer, 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. Figure 1 One or more processes and / or boxes Figure 1 A device that provides the functions specified in one or more boxes.

[0115] 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.

[0116] 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.

[0117] Although preferred embodiments of the invention have been described, those skilled in the art, upon learning the basic inventive concept, can make other changes and modifications to these embodiments. Therefore, the appended claims are intended to be interpreted as including both the preferred embodiments and all changes and modifications falling within the scope of the invention.

[0118] Obviously, those skilled in the art can make various modifications and variations to this invention without departing from its spirit and scope. Therefore, if these modifications and variations fall within the scope of the claims of this invention and their equivalents, this invention also intends to include these modifications and variations.

Claims

1. A method for memory access security checking in a GPU compiler, characterized in that, include: Step 1: Use a standard toolchain to convert GPU code into high-level IR instructions, which include metadata information; Step 2: Based on the metadata information in the high-level IR instructions, extract and symbolize the boundary and runtime information of each memory allocation, and convert the high-level IR instructions into low-level IR instructions; Step 3: Based on the range analysis method of the transfer function, in the compiler backend, data flow analysis is performed on the low-level IR instructions to calculate the symbolic address range of each memory access. Based on the boundary of each memory allocation and runtime information in the metadata information, the security of each memory access is checked. Step 4: Mark memory accesses that pose a security risk and generate a security analysis report.

2. The memory access security check method according to claim 1, characterized in that, Step 1, which involves using a standard toolchain to convert GPU code into high-level IR instructions, includes: Step 11: Based on the standard compiler front-end, convert the GPU code into high-level IR instructions, which include core metadata; Step 12: Establish the mapping between the high-level IR instructions and the GPU source code files and line numbers.

3. The memory access security check method according to claim 1, characterized in that, Step 2 involves extracting and symbolizing the boundaries and runtime information of each memory allocation based on the core metadata in the high-level IR instructions, including: Step 21, Memory object metadata extraction: Based on the metadata information in the high-level IR instructions, create a unique identifier for each memory allocation in the code. The unique identifier is used to record the boundary size of each memory allocation. For pointers to the global address space with unknown size, a symbol is used to represent their boundaries; for pointers to the local address space and the private address space with known size, the exact byte size is used to represent their boundaries. Step 22, runtime information metadata extraction: Based on the metadata information in the high-level IR instructions, assign an initial value range to each built-in function. The value range is symbolic. When there are additional constraints, the value range is refined into a constant based on the constraint information. Step 23: After extracting the metadata information, convert the high-level IR instructions into low-level IR instructions.

4. The memory access security check method according to claim 1, characterized in that, Step 3, based on the range analysis method of the transfer function, performs data flow analysis on the low-level IR instructions in the compiler backend, calculates the symbolic address range of each memory access, and checks the security of each memory access based on the boundaries and runtime information of each memory allocation in the metadata information, including: Step 31: Define register status information, which is represented as State(%reg) = {BasePointerID, AddressSpace, OffsetRange}, where State(%reg) represents the status information of register %reg; BasePointerID represents a unique identifier, indicating the pointer base address; AddressSpace represents the memory space type, which can be private, global, local, or NonPointer; OffsetRange represents the pointer offset [min_offset, max_offset], which are the minimum and maximum byte offsets of the pointer relative to BasePointerID. Step 32: Obtain the metadata information, the low-level IR instructions in the current state of the compiler backend, and the corresponding control flow graph (CFG). Perform data flow analysis based on the control flow graph (CFG) and update the register state information. Step 33: When a memory access instruction is encountered, obtain the current register status information; Step 34: Obtain the corresponding memory allocation boundary based on the base address and memory space type in the current register status information; Step 35: Determine whether the pointer offset is out of bounds based on the pointer offset in the current register status information and the corresponding memory allocation boundary.

5. The memory access security check method according to claim 4, characterized in that, Step 32 involves acquiring the metadata information, the low-level IR instructions in the compiler backend at their current state, and the corresponding control flow graph (CFG). Based on the CFG, data flow analysis is performed, and the register state information is updated, including: Step 321: Allocate storage space for each basic block B in the control flow graph CFG to record its entry state IN[B] and exit state OUT[B]; Step 322: Create a work list, which is used to store the data structure of the basic blocks to be processed. Step 323: After initialization, place the entry basic blocks in the control flow graph (CFG) into the working list; Step 324: Take a currently pending basic block from the work list, denoted as B_current; Step 325: Calculate the new entry state IN[B_current] for B_current; Step 326: Based on the new entry state IN[B_current] of B_current, which is B_current, calculate the new exit state OUT[B_current] and update the register state information; Step 327: Check whether the calculated new exit state of B_current causes a change in the state of subsequent basic blocks. If so, store the data structure of the subsequent basic blocks of B_current into the working list and remove B_current from the working list.

6. The memory access security check method according to claim 5, characterized in that, Step 325, calculating the new entry state IN[B_current] for B_current, includes: In the control flow graph CFG, all basic blocks that directly point to B_current are identified, namely the set of predecessor basic blocks of B_current, Predecessors(B_current). For the exit state OUT[P] of all predecessor basic blocks P∈Predecessors(B_current), a predefined merge operator Meet / Join Operator is applied. This operator merges the data flow information from multiple input paths into a single state, which is mathematically represented as: IN[B_current] = Meet(OUT[P_1], OUT[P_2], ..., OUT[P_i], ..., OUT[P_n]); Where P_i ∈ Predecessors(B_current), and n is the number of predecessor basic blocks of B_current.

7. The memory access security check method according to claim 5, characterized in that, Step 326: Based on the new entry state IN[B_current] of B_current, which is B_current, calculate the new exit state OUT[B_current], and update the register state information, including: Iterate through and process all instructions in B_current, calculate the new exit state OUT[B_current] of B_current, and update the register state for each instruction type in B_current by defining a transfer function.

8. The memory access security check method according to claim 4, characterized in that, Step 35, based on the pointer offset in the current register status information and the corresponding memory allocation boundary, determines whether the pointer offset is out of bounds, including: Get the pointer offset [min_offset, max_offset] from the current register status information, where the boundary of memory allocation obtained from the metadata information is denoted as [0, alloc_size]. The lower bound min_offset and the upper bound max_offset of the pointer offset are checked respectively. The check on the lower bound min_offset includes: If min_offset >= 0, then the lower bound has not been exceeded; If min_offset < 0, and min_offset is a negative constant, then it is an explicit lower bound out-of-bounds error; If it cannot be completely determined that min_offset >= 0, and min_offset depends on the sign, then it is a potential lower bound out-of-bounds error; The upper bound max_offset is checked, including: If max_offset + access_size <= alloc_size, then perform an address space check; If max_offset + access_size > alloc_size, and max_offset + access_size is a constant, then it is a clear upper bound out-of-bounds error; If max_offset + access_size > alloc_size, and max_offset + access_size depends on the symbol, then it is a potential upper bound out-of-bounds error; The address space check includes: The address range is calculated based on the pointer base address and pointer offset in the current register status information. The address range is then matched with the specified address space range. If they do not match, it indicates that the pointer is being referenced incorrectly. The specified address space range for different types of memory spaces is determined at compile time.

9. The memory access security check method according to claim 5, characterized in that, Step 327 involves checking whether the calculated new exit state of B_current causes a change in the state of subsequent basic blocks. If so, the data structure of the subsequent basic blocks of B_current is stored in the working list, and B_current is removed from the working list. The step also includes determining whether the data flow analysis has reached convergence. If it has, the data flow analysis ends. The determination of whether the data stream analysis has converged includes: For each basic block in the control flow graph CFG, the newly calculated exit state OUT_new[B_current] is compared with the previously stored old exit state OUT_old[B_current]. If OUT_new and OUT_old are different, it indicates that the data flow information has changed and the data flow analysis has not yet converged. Update the exit state of B_current to OUT_new, and add the data structures of all successor nodes of B_current in the control flow graph CFG to the working list; If OUT_new is the same as OUT_old, no operation is performed; When the work list is empty, the state of all basic blocks has stabilized, and the data flow analysis has converged.

10. The memory access security check method according to claim 5, characterized in that, Step 4 involves marking memory accesses that pose a security risk and generating a security analysis report, including: Step 41: Extract all memory access instances marked as risk and their associated context data from the analyzed low-level IR instructions; Step 42: Assign a stable and unique ID to each individual instance of memory access with risk, and classify it into high-risk, medium-risk, or low-risk levels based on the certainty of the risk. Step 43: Map the risk points in the low-level IR instructions back to their file names and line numbers in the high-level GPU source code, and programmatically generate descriptive text to explain the nature of the risk and potential triggering conditions. Step 44: Compile all generated risk records into a structured file and output it.

Images