Contention Analysis Method and System for GPU L1 Data Cache
Patent Information
- Application Number
- CN202611244105.0
- Authority / Receiving Office
- CN · China
- Patent Type
- Applications(China)
- Current Assignee / Owner
- Filing Date
- 2026-08-17
- Publication Date
- 2026-10-23
AI Technical Summary
这种次序的随机性直接导致缓存争用与替换行为难以断定，由此极大地增加了分析复杂度
[0040]本申请可以在核函数的每条访存指令之间插入同步语句，分别分析各次循环访问中的访存地址，将条件分支访问中所有分支按程序顺序合并为统一分析路径；对每条访存指令构建对应的访存集合；其中，访存集合包括访存指令所有候选访问的缓存行条目；根据访存集合计算每个缓存行条目对应的干扰集合；根据干扰集合逐指令分析访存指令对 L1数据缓存的最坏争用情况，得到各条访存指令对应的总未命中次数。本申请先预处理核函数，再以访存指令为分析粒度，构建访存集合和干扰集合，进而计算得到各条访存指令的总未命中次数，能准确地反映缓存未命中的最坏情况数量。
Smart Images
Figure CN122963982A_ABST
Abstract
Description
Technical Field
[0001] This application relates to the field of computer system technology, and in particular to a contention analysis method and system for GPU L1 data cache. Background Technology
[0002] In the complex memory hierarchy of GPUs, the L1 cache has the most direct and significant impact on the accuracy of WCET estimation. As the high-speed cache closest to the streaming multiprocessor core, the L1 cache exhibits access latency orders of magnitude different from back-end caches (such as L2 cache) and even global memory. Specifically, an L1 cache hit takes approximately 30 clock cycles, an L2 cache hit takes approximately 200 cycles, while accessing global memory takes as much as 290 cycles. This significant latency difference means that the accuracy of predicting L1 cache hits and misses directly determines the tightness of the final upper bound of WCET. Therefore, constructing an accurate L1 cache contention analysis model is a crucial step in reducing the pessimism of WCET estimation and improving the end-to-end predictability of the entire system.
[0003] Achieving accurate L1D cache contention analysis on GPU platforms faces two major challenges. First, there is the uncertainty of memory addresses. Many applications' memory access addresses depend on dynamic input data, making it difficult to determine the exact memory addresses in advance during static analysis, thus hindering accurate judgment of whether access by different threads will conflict. Second, there is the opacity of the GPU hardware scheduling mechanism. NVIDIA has not disclosed the decision logic of the Warp scheduler, making the order in which memory requests arrive at the L1D cache unpredictable. This randomness directly makes it difficult to determine cache contention and replacement behavior, significantly increasing the complexity of the analysis. Summary of the Invention
[0004] In view of this, embodiments of this application provide a contention analysis method and system for GPU L1 data cache to accurately analyze the worst-case contention situation of L1 data cache.
[0005] One aspect of this application provides a contention analysis method for GPU L1 data cache, the method comprising the following steps:
[0006] Insert synchronization statements between each memory access instruction in the kernel function, analyze the memory access addresses in each loop access separately, and merge all branches in the conditional branch access into a unified analysis path according to the program order.
[0007] For each memory access instruction, a corresponding memory access set is constructed; wherein, the memory access set includes cache line entries for all candidate accesses of the memory access instruction;
[0008] Calculate the interference set corresponding to each cache line entry based on the memory access set;
[0009] Based on the interference set, analyze the worst-case contention situation of the memory access instruction on the L1 data cache, and obtain the total number of misses for each memory access instruction.
[0010] In some embodiments, inserting synchronization statements between each memory access instruction of the kernel function includes the following steps:
[0011] For a single thread block, a block synchronization statement is inserted between each memory access instruction in the kernel function, so that all threads in the thread block execute the block synchronization statement before continuing to execute the subsequent memory access instructions.
[0012] For multiple thread blocks that are executed in parallel, inter-block synchronization statements are inserted between each memory access instruction in the kernel function, so that all threads in each thread block execute the inter-block synchronization statements before continuing to execute the subsequent memory access instructions.
[0013] In some embodiments, constructing a corresponding memory access set for each memory access instruction includes the following steps:
[0014] Allocate a memory block of the first target byte for the array based on the array's original size, cache line size, and number of cache groups;
[0015] The first address in the memory block that satisfies the second target byte alignment is located using pointer arithmetic and used as the starting pointer for the actual use of the array;
[0016] Construct corresponding memory access sets for regular memory access instructions and irregular memory access instructions respectively.
[0017] In some embodiments, constructing the corresponding memory access sets for regular memory access instructions and irregular memory access instructions respectively includes the following steps:
[0018] The memory access set is defined to include multiple cache line entries; wherein each cache line entry includes a cache group index, a cache line label, the number of times the thread block accesses the cache line, the age of the cache line, a first boolean variable to identify whether the current access is within a conditional branch, and a second boolean variable to identify whether the current memory access is an irregular memory access.
[0019] If the first bool variable is determined to be true through static analysis, it is determined that the access corresponding to the current cache line entry is within the conditional branch. When updating the age value of the cache line entry, the size of the interference set is added to the original age value. If the first bool variable is false, it is determined that the access corresponding to the current cache line entry is not within the conditional branch, and the age value of the cache line entry is set to the size of the interference set.
[0020] If the specific address cannot be determined through static analysis, the second bool variable is set to true, thus determining that the current memory access instruction is the irregular memory access instruction, and the mapping range of the interference set of the irregular memory access instruction is set to any cache group.
[0021] In some embodiments, the method further includes the following steps:
[0022] If, in the memory access set corresponding to any of the aforementioned rule memory access instructions, the cache group indexes of each cache line entry are concentrated in a set number of cache groups, then redundant space is inserted at the end of the rows of the array to change the step size between adjacent data rows, so that the cache line entries are mapped to each cache group as evenly as possible, thereby reducing the upper bound of the analysis of the interference set size and the number of cache misses.
[0023] In some embodiments, calculating the interference set corresponding to each cache line entry based on the memory access set includes the following steps:
[0024] For the first cache line entry of the rule memory access instruction, calculate the corresponding cache group index and cache line label according to the index, and determine whether different first cache line entries constitute interference items, and then construct the corresponding interference set according to each interference item;
[0025] For the second cache line entry of the irregular memory access instruction, define the interference between the second cache line entry and any other second cache line entry, and then construct the corresponding interference set.
[0026] In some embodiments, the step of analyzing the worst-case contention of the memory access instructions for the L1 data cache on an instruction-by-instruction basis according to the interference set, and obtaining the total number of misses for each memory access instruction, includes the following steps:
[0027] For each access request of the cache line entry, the following steps are repeated until all accesses to the current cache line entry are analyzed or the remaining interference items in the interference set are insufficient to cause a cache miss: For the first access, if the sum of the initial age value of the cache line and the size of the interference set is greater than or equal to the correlation of the L1 data cache, it is determined to be a miss; wherein, for the first memory access instruction of the kernel function, the age value of all cache lines is set to be consistent with the value of the correlation; for subsequent accesses, a greedy strategy is used to select the interference item with the most remaining accesses from the interference set for analysis; if the number of remaining interference items is still greater than or equal to the correlation, it is determined to be a miss.
[0028] After the analysis of a single access instruction is completed, if the cache line entry is not in a conditional branch, the age value of the cache line entry is set to the size of the interference set; if the cache line entry is in a conditional branch, the size of the interference set is added to the original age value.
[0029] After calculating the worst-case execution of each access instruction in the order of instructions, the total number of misses corresponding to all cache line entries is accumulated to obtain the total number of misses.
[0030] Another aspect of this application embodiment provides a contention analysis system for GPU L1 data cache, the system comprising:
[0031] The kernel function preprocessing unit is used to insert synchronization statements between each memory access instruction of the kernel function, analyze the memory access addresses in each loop access, and merge all branches in the conditional branch access into a unified analysis path according to the program order.
[0032] The memory access set construction unit is used to construct a corresponding memory access set for each memory access instruction; wherein, the memory access set includes cache line entries for all candidate accesses of the memory access instruction;
[0033] An interference set construction unit is used to calculate the interference set corresponding to each cache line entry based on the memory access set.
[0034] The worst-case contention analysis unit is used to analyze the worst-case contention situation of the memory access instruction on the L1 data cache according to the interference set, and to obtain the total number of misses for each memory access instruction.
[0035] Another aspect of this application embodiment provides an electronic device, including a processor and a memory;
[0036] The memory is used to store programs;
[0037] The processor executes the program to implement any of the methods described above.
[0038] Another aspect of this application provides a computer-readable storage medium storing a program that is executed by a processor to implement the method described in any of the above embodiments.
[0039] This application includes at least the following beneficial effects:
[0040] This application inserts synchronization statements between each memory access instruction in the kernel function, analyzes the memory access addresses in each loop access separately, and merges all branches in the conditional branch access into a unified analysis path according to the program order. It constructs a corresponding memory access set for each memory access instruction, where the memory access set includes cache line entries for all candidate accesses of the memory access instruction. Based on the memory access set, it calculates the interference set corresponding to each cache line entry. Based on the interference set, it analyzes the worst-case contention situation of the memory access instruction on the L1 data cache instruction by instruction, obtaining the total number of cache misses for each memory access instruction. This application first preprocesses the kernel function, then constructs the memory access set and interference set at the memory access instruction level, and then calculates the total number of cache misses for each memory access instruction, accurately reflecting the worst-case number of cache misses. Attached Figure Description
[0041] To more clearly illustrate the technical solutions in the embodiments of this application, 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 this application. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort.
[0042] Figure 1 A flowchart illustrating a contention analysis method for GPU L1 data cache provided in an embodiment of this application;
[0043] Figure 2 An example flowchart of a contention analysis method for GPU L1 data cache provided in this application embodiment;
[0044] Figure 3 Example code diagrams for alignment operations provided in embodiments of this application;
[0045] Figure 4 Example code diagrams for rule access provided in embodiments of this application;
[0046] Figure 5 Example code diagrams for fill operations provided in embodiments of this application;
[0047] Figure 6 Example code diagrams for inserting synchronization statements provided in embodiments of this application;
[0048] Figure 7 Algorithm code diagram for calculating worst-case contention in a single-threaded block provided in this application embodiment;
[0049] Figure 8 This is a block diagram of a contention analysis system for GPU L1 data cache provided in an embodiment of this application. Detailed Implementation
[0050] To make the objectives, technical solutions, and advantages of this application clearer, the following detailed description is provided in conjunction with the accompanying drawings and embodiments. It should be understood that the specific embodiments described herein are merely illustrative and not intended to limit the scope of this application.
[0051] Before providing a detailed description of the embodiments of this application, some related technologies involved in the embodiments of this application will be described first, as follows:
[0052] Terminology Explanation:
[0053] GPU Thread Hierarchy: The GPU thread hierarchy is a hierarchical structure defined by the NVIDIA CUDA programming model for organizing parallel threads. This structure divides parallel computing tasks into three logical levels: Thread, Thread Block, and Thread Grid. A thread is the most basic unit of execution; multiple threads form a thread block; and multiple thread blocks form a thread grid. Each kernel function (i.e., GPU-executed code) initiates a thread grid. This hierarchy effectively maps to the physical hardware of the GPU and is the foundation for realizing massively parallel computing on GPUs.
[0054] Streaming Multiprocessor (SM): A streaming multiprocessor is the core computing unit in a GPU architecture, responsible for the creation, management, scheduling, and execution of threads. Each SM contains a large number of CUDA cores, registers, L1 data cache, and other hardware resources. In a GPU, thread blocks are bound to a specific SM for execution by the hardware scheduler, and their concurrency is limited by the hardware resources of the SM. The SM schedules all resident threads in units of warp (32 consecutive threads), which is the foundation for achieving large-scale data parallel processing.
[0055] Warp: A warp is a basic execution unit in the GPU execution model. A warp typically consists of 32 consecutive threads. Streaming multiprocessors create, manage, schedule, and execute threads in units of warps. All threads within the same warp execute the same instructions within the same cycle but operate on different data, exhibiting a single-instruction multithreaded (SIMT) execution mode.
[0056] Warp scheduling refers to the process by which a streaming multiprocessor selects one warp from the set of currently ready warps and issues it to an execution unit for execution within each clock cycle. Its main goal is to hide instruction and memory access latency to maintain the saturation of computational units. When a warp stalls while waiting for data, the scheduler quickly switches to another ready warp for execution, thus achieving high throughput. Common scheduling strategies include Greedy Then Oldest (GTO) and Loose Round Robin (LRR), but NVIDIA has not yet disclosed its actual warp scheduling strategy.
[0057] L1 Data Cache: The L1 data cache is a high-speed, first-level data cache located within the streaming multiprocessor (SM) in the GPU's memory hierarchy. Its primary function is to cache frequently accessed data from Global Memory (GNMemory) to reduce latency for subsequent accesses. Each SM typically has its own independent L1 data cache. The L1 data cache uses a set-associative architecture, consisting of several cache sets, each containing multiple cache lines; the mapping from memory addresses to cache sets is usually determined using modulo operations.
[0058] Cache contention: Cache contention refers to resource access conflicts that occur when multiple threads or warps access data mapped to the same cache set in the L1 data cache, due to competition for cache lines. This contention leads to frequent cache line replacements, significantly reducing cache hit rates, thereby increasing effective memory access latency and ultimately impacting the overall computing performance of the GPU. Accurate quantitative analysis of cache contention is a prerequisite for optimizing memory access patterns and improving system performance.
[0059] With the increasing demand for intelligent computing, graphics processing units (GPUs) are being increasingly used in safety-critical systems that require stringent timing requirements, such as autonomous driving, aerospace, and industrial control. These systems must verify end-to-end timing behavior to ensure functional correctness and reliability. Worst-case execution time (WCET) analysis is a core technology for verifying the safety of system timing behavior. Its goal is to determine the upper bound of execution time that a task can achieve under any possible input and a defined execution environment. A safe and compact WCET upper bound is crucial for the rational allocation of system resources and the satisfaction of timing constraints.
[0060] Among the closest related technologies to this application, two key assumptions are defined to simplify the analysis: first, it is assumed that the addresses of all memory accesses are known in advance during the static analysis phase to avoid the challenge of address uncertainty; second, to handle the opacity of hardware scheduling, this related technology presupposes a fixed and idealized warp scheduling order. Although these strong assumptions enable cache behavior analysis in theoretical modeling, their analytical premise deviates from the dynamic and input-dependent actual environment of GPU runtime, which limits the practical application value of this related technology in security-critical scenarios.
[0061] Existing contention analysis techniques for GPU L1 data caches, especially the aforementioned techniques, have the following limitations:
[0062] This approach assumes that all memory addresses are statically known. To simplify the analysis, it assumes that all memory access addresses are known and fixed in advance during the analysis phase. However, in practical applications, a program's memory access patterns are usually determined by dynamically input data, resulting in significant address uncertainty. This strong assumption makes this analysis technique unsuitable for general input-dependent scenarios, greatly limiting its application scope and the practicality of the analysis results.
[0063] An idealized and fixed warp scheduling order is adopted. To avoid the uncertainty of warp scheduling, existing technologies presuppose a fixed warp execution sequence. However, a fixed scheduling model cannot reflect the actual cache access interleaving and contention, leading to a deviation between the analysis results and the actual runtime behavior of the GPU, thus affecting its reliability and trustworthiness in security-critical scenarios.
[0064] The purpose of this application is to overcome the shortcomings of the prior art and provide a safe and compact worst-case contention analysis method for GPU L1 data cache. This technology possesses the following key technical solutions:
[0065] Using a single memory access as the analysis granularity, a corresponding interference set is constructed. By combining factors such as the size of the interference set, its age, and cache associativity, it is determined whether the access may cause a cache invalidation in the worst case, thereby obtaining a safe upper bound on the number of cache invalidations.
[0066] For regular memory access patterns, the NVIDIA API is used to apply memory address alignment constraints to the starting address of data objects to eliminate the uncertainty of the base address in cache set mapping. Simultaneously, when there is an imbalance in cache set mapping, a uniform cache set mapping technique is used to optimize the size of the interference set, thereby tightening the upper bound of the cache miss count analysis.
[0067] For irregular memory accesses where addresses are difficult to determine statically, such as pointer access and indirect addressing, a conservative interference set construction strategy is adopted. These are treated as potentially mapping to any cache set and interfering with other memory accesses, in order to ensure the safety of worst-case analysis.
[0068] NVIDIA-provided synchronization instructions are inserted into the kernel function to transform dynamic Warp scheduling into a deterministic, phased synchronous execution model. Based on this, loop structures and conditional branches in the kernel function are preprocessed. For loop structures, the loop body is statically expanded according to the loop's upper bound, and the memory access addresses in each loop iteration are analyzed. For conditional branch structures, all possible branches are merged into a unified analysis path according to the program sequence, and a monotonic age update rule is applied to memory accesses within each branch.
[0069] After the above processing in this application, the maximum cache contention can be analyzed according to the memory access instruction order, and the number of worst-case cache misses can be calculated.
[0070] Reference Figure 1 This application provides a contention analysis method for GPU L1 data cache, specifically including the following steps S100~S130:
[0071] S100: Insert synchronization statements between each memory access instruction in the kernel function, analyze the memory access addresses in each loop access separately, and merge all branches in the conditional branch access into a unified analysis path according to the program order;
[0072] S110: Construct a corresponding memory access set for each memory access instruction; wherein, the memory access set includes cache line entries for all candidate accesses of the memory access instruction;
[0073] S120: Calculate the interference set corresponding to each cache line entry based on the memory access set;
[0074] S130: Analyze the worst-case contention situation of the memory access instruction on the L1 data cache according to the interference set, and obtain the total number of misses for each memory access instruction.
[0075] Optionally, inserting synchronization statements between each memory access instruction in the kernel function includes the following steps:
[0076] For a single thread block, a block synchronization statement is inserted between each memory access instruction in the kernel function, so that all threads in the thread block execute the block synchronization statement before continuing to execute the subsequent memory access instructions.
[0077] For multiple thread blocks that are executed in parallel, inter-block synchronization statements are inserted between each memory access instruction in the kernel function, so that all threads in each thread block execute the inter-block synchronization statements before continuing to execute the subsequent memory access instructions.
[0078] Optionally, constructing a corresponding memory access set for each memory access instruction includes the following steps:
[0079] Allocate a memory block of the first target byte for the array based on the array's original size, cache line size, and number of cache groups;
[0080] The first address in the memory block that satisfies the second target byte alignment is located using pointer arithmetic and used as the starting pointer for the actual use of the array;
[0081] Construct corresponding memory access sets for regular memory access instructions and irregular memory access instructions respectively.
[0082] Optionally, constructing the corresponding memory access sets for regular and irregular memory access instructions respectively includes the following steps:
[0083] The memory access set is defined to include multiple cache line entries; wherein each cache line entry includes a cache group index, a cache line label, the number of times the thread block accesses the cache line, the age of the cache line, a first boolean variable to identify whether the current access is within a conditional branch, and a second boolean variable to identify whether the current memory access is an irregular memory access.
[0084] If the first bool variable is determined to be true through static analysis, it is determined that the access corresponding to the current cache line entry is within the conditional branch. When updating the age value of the cache line entry, the size of the interference set is added to the original age value. If the first bool variable is false, it is determined that the access corresponding to the current cache line entry is not within the conditional branch, and the age value of the cache line entry is set to the size of the interference set.
[0085] If the specific address cannot be determined through static analysis, the second bool variable is set to true, thus determining that the current memory access instruction is the irregular memory access instruction, and the mapping range of the interference set of the irregular memory access instruction is set to any cache group.
[0086] Optionally, the method further includes the following steps:
[0087] If, in the memory access set corresponding to any of the aforementioned rule memory access instructions, the cache group indexes of each cache line entry are concentrated in a set number of cache groups, then redundant space is inserted at the end of the rows of the array to change the step size between adjacent data rows, so that the cache line entries are mapped to each cache group as evenly as possible, thereby reducing the upper bound of the analysis of the interference set size and the number of cache misses.
[0088] Optionally, calculating the interference set corresponding to each cache line entry based on the memory access set includes the following steps:
[0089] For the first cache line entry of the rule memory access instruction, calculate the corresponding cache group index and cache line label according to the index, and determine whether different first cache line entries constitute interference items, and then construct the corresponding interference set according to each interference item;
[0090] For the second cache line entry of the irregular memory access instruction, define the interference between the second cache line entry and any other second cache line entry, and then construct the corresponding interference set.
[0091] Optionally, the step of analyzing the worst-case contention situation of the memory access instructions on the L1 data cache according to the interference set, and obtaining the total number of misses for each memory access instruction, includes the following steps:
[0092] For each access request of the cache line entry, the following steps are repeated until all accesses to the current cache line entry are analyzed or the remaining interference items in the interference set are insufficient to cause a cache miss: For the first access, if the sum of the initial age value of the cache line and the size of the interference set is greater than or equal to the correlation of the L1 data cache, it is determined to be a miss; wherein, for the first memory access instruction of the kernel function, the age value of all cache lines is set to be consistent with the value of the correlation; for subsequent accesses, a greedy strategy is used to select the interference item with the most remaining accesses from the interference set for analysis; if the number of remaining interference items is still greater than or equal to the correlation, it is determined to be a miss.
[0093] After the analysis of a single access instruction is completed, if the cache line entry is not in a conditional branch, the age value of the cache line entry is set to the size of the interference set; if the cache line entry is in a conditional branch, the size of the interference set is added to the original age value.
[0094] After calculating the worst-case execution of each access instruction in the order of instructions, the total number of misses corresponding to all cache line entries is accumulated to obtain the total number of misses.
[0095] The following section will provide a detailed introduction and explanation of the solutions in the embodiments of this application, using specific application examples.
[0096] The problem modeling in this embodiment is described below:
[0097] This embodiment focuses on a single Streaming Multiprocessor (SM) and considers the cache access behavior of thread blocks within the target GPU kernel function when executed on that SM. The thread block configuration of this kernel function is denoted as... , and These represent the number of threads in the thread block along the x and y dimensions, respectively; that is, each thread block contains... One thread.
[0098] In this embodiment, the regular memory access mode includes continuous memory access and step-by-step memory access, which can be uniformly represented as: In the form of. Among them, This represents the index of the thread in the corresponding dimension. Indicates the array size or stride parameter. Representation and thread index An irrelevant fixed offset. For this type of regular memory access, the memory access address can be represented as... ,in The base address of the array. The size of the element being accessed is denoted as . Other memory accesses that are difficult to represent in the above-described regular form, such as pointer access and indirect addressing, typically have memory addresses that depend on runtime data and are difficult to determine precisely during static analysis. This embodiment treats these as irregular memory accesses and handles them conservatively in subsequent cache contention analysis.
[0099] For each memory access instruction, its corresponding memory access set needs to be constructed. This set consists of cache line entries (also called entries or memory access entries) that the memory access instruction may access, where each entry... It can be represented as a six-tuple .in, Indicates the cache group index. Indicates the cache line label. This indicates the number of times the thread block accesses this cache line. This indicates the age of the cached line. Used to identify whether the access is within a conditional branch. Used to identify whether the memory access is an irregular memory access.
[0100] Furthermore, the hardware parameters of the L1 data cache are determined by the number of cache sets, the size of cache lines, and the correlation, and its total capacity can be expressed as follows: ,in For the number of cache groups, The cache line size (in bytes). For relevance. Cache set mapping uses a linear modulo rule, meaning that given a memory access address... Its corresponding cache group index With cache line labels From the formula and The calculation yielded the result.
[0101] For example, Figure 2 This is an example flowchart of a contention analysis method for GPU L1 data cache. The following describes the specific steps of this embodiment. It should be noted that the step numbers are not limited in execution order, and the specific order of execution can be adjusted appropriately according to the actual situation.
[0102] Step 1: Memory access address analysis and memory access set construction.
[0103] Step 1.1: First, to eliminate the uncertainty of the array's starting address in the cache set mapping, the array needs to be... Byte alignment ( Figure 3 (Example code diagram for alignment operation), mapping its starting address to the first position of the cache group. The specific implementation is as follows: first, through... The method allocates a size of [size missing] to the array. A memory block of bytes. Among them, This is the original size of the array. Then, pointer arithmetic is used to locate the first element in this memory block that satisfies this value. A byte-aligned address, used as the starting pointer for the actual array.
[0104] Step 1.2: Construct the corresponding memory access set for the regular memory access instructions. Memory access set Each cache line entry The attribute is calculated as follows: cache group index Cache line tags and the number of cached lines accessed Obtained through cache mapping rules; and As a boolean variable, it can be determined through simple static analysis; Instead of being given as a static parameter in advance, it is calculated iteratively during contention analysis, as detailed in Step 3.5 regarding the age update process. Figure 4Example code diagram for rule-based access (matrix multiplication).
[0105] Step 1.3: If the cache set index of each cache line entry in the memory access set corresponding to a certain rule memory access instruction is... If the interference set is concentrated in a few cache groups, resulting in an excessively large size, padding techniques can be introduced for optimization. This involves inserting redundant space at the end of array rows to change the stride between adjacent data rows, ensuring that cache line entries are mapped as evenly as possible across cache groups. This tightens the upper bound on the size of the interference set and the number of cache misses. Specifically, each row of the array is padded to a specific value, which can be represented as a coefficient. The product of an odd number, Figure 5 The image shows a sample code diagram for the fill operation. This step does not affect the safety of cache contention analysis; it is only used to tighten the upper bound of the cache miss count under specific conditions.
[0106] Step 1.4: For irregular memory accesses where the specific address is difficult to determine statically, instead of enumerating the specific memory access address and cache set index, the irregular memory access flag of the entry is set. Set it to true, and conservatively treat it as potentially mapping to any cache set in the subsequent interference set construction. Thus, the memory access set corresponding to the current memory access instruction is obtained. .
[0107] Step 2: Kernel function performs constraint and control flow preprocessing.
[0108] Step 2.1 (Loop Preprocessing): Since the address accessed in each loop in the kernel function may be different, it is necessary to statically unroll the loop according to the upper bound of the loop and analyze the operation address of each loop in turn.
[0109] Step 2.2 (Branch Preprocessing): For conditional branches whose execution path cannot be determined through static analysis, all possible branches need to be executed. Since this introduces redundant accesses, to avoid interfering with the L1 data cache, it is required that when executing conditional branches, It can only increase, not decrease. Specifically, update. At that time, if If true, then Reductions are not allowed. Detailed update procedures are provided in the algorithm section of Step 3.6.
[0110] Step 2.3 (Synchronization Preprocessing): Insert intra-block synchronization statements between each memory access instruction in the kernel function, such as those provided by NVIDIA. method, Figure 6The image shows sample code for inserting a synchronization statement. This operation ensures that all threads within a thread block must reach this barrier point before continuing to execute subsequent instructions. This allows for safe analysis of the worst-case cache contention that may occur in the L1 data cache for a single thread block, at the instruction level.
[0111] Step 3: Worst-case cache contention analysis for single-threaded blocks.
[0112] Step 3.1: After preprocessing the kernel function in Step 2, under single-threaded block conditions, we can analyze the worst-case contention of memory access instructions on the L1 data cache instruction by instruction. For each memory access instruction, we first obtain its corresponding memory access set through Step 1. .
[0113] Step 3.2: Determine the memory access set Then, it is necessary to calculate each cache line entry. The corresponding set of interference. For regular memory access, since Step 1.1 has eliminated the uncertainty of the array starting address in the cache set mapping, its cache set index can be calculated based on the subscript. and cache line labels And determine whether different cache line entries interfere with each other. Specifically, for any cache line entry... Its interference set It can be represented as:
[0114] .
[0115] in, This indicates that two cache line entries are mapped to the same cache group. This indicates that the two correspond to different cache lines, which may cause cache replacement issues. Also, for For truly irregular memory accesses, the system conservatively assumes that they may interfere with any other cache line entries.
[0116] Step 3.3: For irregular memory access instructions, their corresponding memory access sets Irregular memory access flags for each cache line entry If all values are true, the interference set can be directly calculated using the interference set formula in Step 3.2.
[0117] Step 3.4: After obtaining the interference set, the worst-case execution situation for a single instruction can be calculated. This is done for each cache line entry. The access request can be analyzed as follows, and the specific steps are given in the algorithm in Step 3.6.
[0118] For the first visit, if the initial age value The sum of the sizes of the interference sets is greater than or equal to If the cache access instruction is not found, it is considered a cache miss. For the first memory access instruction of the kernel function (i.e., the initial phase), all cache lines... Set all to .
[0119] For subsequent visits, a greedy strategy is used from... Select the distractor with the most remaining access counts for analysis. If the number of remaining distractors is still greater than or equal to... If the result is not met, it is considered a miss.
[0120] Repeat this process until All accesses have been analyzed or there are not enough remaining distractors to cause a cache miss.
[0121] Step 3.5: After analyzing a single access instruction, the cache lines in each cache group need to be updated. To reflect The worst-case state after execution, and the updated This serves as the initial age value for analyzing subsequent memory access instructions. The specific update strategy is: if... If it is not in a conditional branch, then... Let this be the size of the interference set; if If it is in a conditional branch, then in the original... The size of the interference set is added on top of the existing set to ensure that the cache line age within the conditional branch can only increase monotonically.
[0122] Step 3.6: Calculate the worst-case execution scenario for each instruction in sequence, then sum all the results. The corresponding number of misses gives the total number of misses. Figure 7 The algorithm code diagram for calculating the worst-case contention in a single-threaded block.
[0123] Step 4: Worst-case cache contention analysis for multi-threaded blocks.
[0124] Step 4.1: For cases where multiple thread blocks execute in parallel, the intra-block synchronization in Step 2.3 is extended to inter-thread block synchronization. Specifically, the `grid.sync()` method provided by CUDA Cooperative Groups can be called between memory access phases. This synchronization introduces additional synchronization wait and scheduling overhead.
[0125] Step 4.2: Under inter-block synchronization constraints, cache line entries The scope of analysis is expanded from a single thread block to all thread blocks involved in synchronization. For the same memory access instruction, the memory access set... It includes all memory accesses generated by all thread blocks during this memory access phase.
[0126] Step 4.3: Obtain the expanded memory access set Then, following the steps in step 3, including constructing the interference set, determining cache misses, and... The update rules are analyzed to obtain the upper bound of the worst-case cache miss in a multi-threaded block scenario.
[0127] In addition, the following alternative solutions can also be provided in this embodiment:
[0128] 1. Memory address alignment can be achieved by using pointer offset after cudaMalloc, or by using other memory allocation methods that can guarantee that the array starting address meets the specified byte alignment.
[0129] 2. Intra-block synchronization can be implemented using NVIDIA's __syncthreads() function, or other GPU programming interfaces with equivalent intra-block barrier synchronization semantics.
[0130] In summary, this embodiment may include the following key technical solutions:
[0131] 1. A worst-case contention analysis method for GPU L1 data cache is proposed. This method uses memory access instructions as the analysis granularity, constructs a memory access set and a disturbance set, and calculates the number of worst-case cache misses based on cache mapping relationships and cache state update rules.
[0132] 2. A cache set mapping constraint and optimization method oriented towards rule-based memory access is proposed. This method eliminates the uncertainty of array start address in cache set mapping by memory address alignment, and introduces array padding technology when the cache set index distribution is uneven, so that cache line entries are mapped to each cache set as evenly as possible, thereby tightening the upper bound of cache miss count analysis without affecting the analysis security.
[0133] 3. A synchronization analysis mechanism for GPU Warp scheduling uncertainty is proposed. This method transforms the opaque dynamic Warp scheduling into a deterministic phased execution model by inserting synchronization instructions between kernel function memory access operations, making cache contention analysis independent of specific Warp scheduling strategies. At the same time, it combines loop static unrolling and conditional branch merging to ensure the safety of worst-case analysis performed in the order of memory access instructions.
[0134] 4. Address Mapping Controllability: By analyzing the memory access instructions in the kernel function and combining custom memory alignment and array filling techniques, the mapping relationship of data in the L1 data cache is actively controlled to eliminate the uncertainty of array base address mapping and reduce the conservative nature of analysis caused by cache set mapping imbalance under certain conditions, thus providing a foundation for building a deterministic cache contention analysis model.
[0135] 5. Analyzability of scheduling behavior: By inserting synchronization barriers between memory access instructions in kernel functions, constraints are imposed on the scheduling and execution order of Warp, thereby transforming the originally opaque parallel execution process into an analyzable phased execution scenario. This is the key to safely quantifying cache interference between Warp instances.
[0136] 6. Contention upper bound calculation mechanism: Based on the memory access set and the interference set, combined with the cache line age update rule, calculate the worst-case number of cache misses within and between thread blocks.
[0137] Beneficial effects:
[0138] Compared to existing technologies that assume statically known memory access addresses, this embodiment can handle input-related uncertainties in memory access addresses. For regular memory accesses, this embodiment reduces cache set mapping uncertainty and analysis conservatism by using memory address alignment and array filling constraints to constrain the mapping relationship of data in the L1 data cache. For irregular memory accesses where addresses are difficult to determine statically, this embodiment incorporates a conservative disturbance modeling approach into the worst-case analysis, thereby expanding the applicability of the method.
[0139] Compared to existing technologies that employ a fixed warp scheduling order, this embodiment does not rely on idealized scheduling assumptions. Instead, it transforms dynamic warp scheduling into a deterministic, phased synchronous execution model by inserting synchronization instructions between kernel function memory access instructions. Based on this, this embodiment can analyze the maximum cache contention between warps according to the execution order of memory access instructions, thereby obtaining a safer and more compact upper bound for the worst-case cache miss.
[0140] Reference Figure 8 This application provides a contention analysis system for GPU L1 data cache, including:
[0141] The kernel function preprocessing unit is used to insert synchronization statements between each memory access instruction of the kernel function, analyze the memory access addresses in each loop access, and merge all branches in the conditional branch access into a unified analysis path according to the program order.
[0142] The memory access set construction unit is used to construct a corresponding memory access set for each memory access instruction; wherein, the memory access set includes cache line entries for all candidate accesses of the memory access instruction;
[0143] An interference set construction unit is used to calculate the interference set corresponding to each cache line entry based on the memory access set.
[0144] The worst-case contention analysis unit is used to analyze the worst-case contention situation of the memory access instruction on the L1 data cache according to the interference set, and to obtain the total number of misses for each memory access instruction.
[0145] It is understood that the content of the above method embodiments is applicable to this embodiment. The specific functions implemented in this embodiment are the same as those in the above method embodiments, and the beneficial effects achieved are also the same as those achieved in the above method embodiments.
[0146] In some alternative embodiments, the functions / operations mentioned in the block diagrams may not occur in the order shown in the operation diagrams. For example, depending on the functions / operations involved, two consecutively shown blocks may actually be executed substantially simultaneously, or the blocks may sometimes be executed in reverse order. Furthermore, the embodiments presented and described in the flowcharts of this application are provided by way of example to provide a more comprehensive understanding of the technology. The disclosed methods are not limited to the operations and logic flows presented herein. Alternative embodiments are contemplated in which the order of various operations is changed and sub-operations described as part of a larger operation are executed independently.
[0147] Furthermore, although this application is described in the context of functional modules, it should be understood that, unless otherwise stated to the contrary, one or more of the described functions and / or features may be integrated into a single physical system and / or software module, or one or more functions and / or features may be implemented in a separate physical system or software module. It is also understood that a detailed discussion of the actual implementation of each module is unnecessary for understanding this application. Rather, given the properties, functions, and internal relationships of the various functional modules in the system disclosed herein, the actual implementation of the module will be understood within the scope of conventional technology for an engineer. Therefore, those skilled in the art can implement the application set forth in the claims using ordinary techniques without excessive experimentation. It is also understood that the specific concepts disclosed are merely illustrative and not intended to limit the scope of this application, which is determined by the full scope of the appended claims and their equivalents.
[0148] If the aforementioned functions are implemented as software functional units and sold or used as independent products, they can be stored in a computer-readable storage medium. Based on this understanding, the technical solution of this application, in essence, or the part that contributes to the prior art, or a portion of the technical solution, can be embodied in the form of a software product. This computer software product is stored in a storage medium and includes several instructions to cause a computer device (which may be a personal computer, server, or network device, etc.) to execute all or part of the steps of the methods described in the various embodiments of this application. The aforementioned storage medium includes various media capable of storing program code, such as USB flash drives, portable hard drives, read-only memory (ROM), random access memory (RAM), magnetic disks, or optical disks.
[0149] The logic and / or steps represented in the flowchart or otherwise described herein, for example, can be considered as a ordered list of executable instructions for implementing logical functions, and can be embodied in any computer-readable medium for use by, or in conjunction with, an instruction execution system, system, or device (such as a computer-based system, a processor-including system, or other system that can fetch and execute instructions from, an instruction execution system, system, or device). For the purposes of this specification, "computer-readable medium" can mean any system that can contain, store, communicate, propagate, or transmit programs for use by, or in conjunction with, an instruction execution system, system, or device.
[0150] More specific examples of computer-readable media (a non-exhaustive list) include: electrical connections (electronic systems) with one or more wires, portable computer disk drives (magnetic systems), random access memory (RAM), read-only memory (ROM), erasable and editable read-only memory (EPROM or flash memory), fiber optic systems, and portable optical disc read-only memory (CDROM). Furthermore, computer-readable media can even be paper or other suitable media on which the program can be printed, because the program can be obtained electronically, for example, by optically scanning the paper or other medium, followed by editing, interpreting, or otherwise processing as necessary, and then stored in computer memory.
[0151] It should be understood that various parts of this application can be implemented using hardware, software, firmware, or a combination thereof. In the above embodiments, multiple steps or methods can be implemented using software or firmware stored in memory and executed by a suitable instruction execution system. For example, if implemented in hardware, as in another embodiment, it can be implemented using any one or a combination of the following techniques known in the art: discrete logic circuits having logic gates for implementing logical functions on data signals, application-specific integrated circuits (ASICs) having suitable combinational logic gates, programmable gate arrays (PGAs), field-programmable gate arrays (FPGAs), etc.
[0152] In the description of this specification, the references to terms such as "one embodiment," "some embodiments," "example," "specific example," or "some examples," etc., indicate that a specific feature, structure, material, or characteristic described in connection with that embodiment or example is included in at least one embodiment or example of this application. In this specification, the illustrative expressions of the above terms do not necessarily refer to the same embodiment or example. Furthermore, the specific features, structures, materials, or characteristics described may be combined in any suitable manner in one or more embodiments or examples.
[0153] Although embodiments of this application have been shown and described, those skilled in the art will understand that various changes, modifications, substitutions and alterations can be made to these embodiments without departing from the principles and spirit of this application, the scope of which is defined by the claims and their equivalents.
[0154] The above is a detailed description of the preferred embodiments of this application, but this application is not limited to the embodiments described. Those skilled in the art can make various equivalent modifications or substitutions without departing from the spirit of this application, and these equivalent modifications or substitutions are all included within the scope defined by the claims of this application.
Claims
1. A contention analysis method for GPU L1 data cache, characterized in that, The method includes the following steps: Insert synchronization statements between each memory access instruction in the kernel function, analyze the memory access addresses in each loop access separately, and merge all branches in the conditional branch access into a unified analysis path according to the program order. For each memory access instruction, a corresponding memory access set is constructed; wherein, the memory access set includes cache line entries for all candidate accesses of the memory access instruction; Calculate the interference set corresponding to each cache line entry based on the memory access set; Based on the interference set, analyze the worst-case contention situation of the memory access instruction on the L1 data cache, and obtain the total number of misses for each memory access instruction.
2. The contention analysis method for GPU L1 data cache according to claim 1, characterized in that, Inserting synchronization statements between each memory access instruction in the kernel function includes the following steps: For a single thread block, a block synchronization statement is inserted between each memory access instruction in the kernel function, so that all threads in the thread block execute the block synchronization statement before continuing to execute the subsequent memory access instructions. For multiple thread blocks that are executed in parallel, inter-block synchronization statements are inserted between each memory access instruction in the kernel function, so that all threads in each thread block execute the inter-block synchronization statements before continuing to execute the subsequent memory access instructions.
3. The contention analysis method for GPU L1 data cache according to claim 1, characterized in that, The step of constructing a corresponding memory access set for each memory access instruction includes the following steps: Allocate a memory block of the first target byte for the array based on the array's original size, cache line size, and number of cache groups; The first address in the memory block that satisfies the second target byte alignment is located using pointer arithmetic and used as the starting pointer for the actual use of the array; Construct corresponding memory access sets for regular memory access instructions and irregular memory access instructions respectively.
4. The contention analysis method for GPU L1 data cache according to claim 3, characterized in that, The step of constructing the corresponding memory access sets for regular and irregular memory access instructions includes the following steps: The memory access set is defined to include multiple cache line entries; wherein each cache line entry includes a cache group index, a cache line label, the number of times the thread block accesses the cache line, the age of the cache line, a first boolean variable to identify whether the current access is within a conditional branch, and a second boolean variable to identify whether the current memory access is an irregular memory access. If the first bool variable is determined to be true through static analysis, it is determined that the access corresponding to the current cache line entry is within the conditional branch. When updating the age value of the cache line entry, the size of the interference set is added to the original age value. If the first bool variable is false, it is determined that the access corresponding to the current cache line entry is not within the conditional branch, and the age value of the cache line entry is set to the size of the interference set. If the specific address cannot be determined through static analysis, the second bool variable is set to true, thus determining that the current memory access instruction is the irregular memory access instruction, and the mapping range of the interference set of the irregular memory access instruction is set to any cache group.
5. The contention analysis method for GPU L1 data cache according to claim 4, characterized in that, The method further includes the following steps: If, in the memory access set corresponding to any of the aforementioned rule memory access instructions, the cache group indexes of each cache line entry are concentrated in a set number of cache groups, then redundant space is inserted at the end of the rows of the array to change the step size between adjacent data rows, so that the cache line entries are mapped to each cache group as evenly as possible, thereby reducing the upper bound of the analysis of the interference set size and the number of cache misses.
6. The contention analysis method for GPU L1 data cache according to claim 4, characterized in that, The step of calculating the interference set corresponding to each cache line entry based on the memory access set includes the following steps: For the first cache line entry of the rule memory access instruction, calculate the corresponding cache group index and cache line label according to the index, and determine whether different first cache line entries constitute interference items, and then construct the corresponding interference set according to each interference item; For the second cache line entry of the irregular memory access instruction, define the interference between the second cache line entry and any other second cache line entry, and then construct the corresponding interference set.
7. The contention analysis method for GPU L1 data cache according to any one of claims 1 to 6, characterized in that, The step of analyzing the worst-case contention of memory access instructions on the L1 cache based on the interference set, and obtaining the total number of misses for each memory access instruction, includes the following steps: For each access request of the cache line entry, the following steps are repeated until all accesses to the current cache line entry are analyzed or the remaining interference items in the interference set are insufficient to cause a cache miss: For the first access, if the sum of the initial age value of the cache line and the size of the interference set is greater than or equal to the correlation of the L1 data cache, it is determined to be a miss; wherein, for the first memory access instruction of the kernel function, the age value of all cache lines is set to be consistent with the value of the correlation; for subsequent accesses, a greedy strategy is used to select the interference item with the most remaining accesses from the interference set for analysis; if the number of remaining interference items is still greater than or equal to the correlation, it is determined to be a miss. After the analysis of a single access instruction is completed, if the cache line entry is not in a conditional branch, the age value of the cache line entry is set to the size of the interference set; if the cache line entry is in a conditional branch, the size of the interference set is added to the original age value. After calculating the worst-case execution of each access instruction in the order of instructions, the total number of misses corresponding to all cache line entries is accumulated to obtain the total number of misses.
8. A contention analysis system for GPU L1 data cache, characterized in that, The system includes: The kernel function preprocessing unit is used to insert synchronization statements between each memory access instruction of the kernel function, analyze the memory access addresses in each loop access, and merge all branches in the conditional branch access into a unified analysis path according to the program order. The memory access set construction unit is used to construct a corresponding memory access set for each memory access instruction; wherein, the memory access set includes cache line entries for all candidate accesses of the memory access instruction; An interference set construction unit is used to calculate the interference set corresponding to each cache line entry based on the memory access set. The worst-case contention analysis unit is used to analyze the worst-case contention situation of the memory access instruction on the L1 data cache according to the interference set, and to obtain the total number of misses for each memory access instruction.
9. An electronic device, characterized in that, The electronic device includes a processor and a memory; The memory is used to store programs; The processor executes the program to implement the method as described in any one of claims 1 to 7.
10. A computer-readable storage medium, characterized in that, The storage medium stores a program that is executed by a processor to implement the method as described in any one of claims 1 to 7.