An address mapping method and system for GPU convolution acceleration

Abstract

The application relates to the field of computer architecture and parallel computing technology, and discloses an address mapping method and system for GPU convolution acceleration, which is applied to the process of performing convolution operation based on general matrix multiplication by a GPU, dynamically remaps the memory access address of a workspace matrix, and comprises the following steps: identifying a data repetition mode in the workspace matrix; based on the data repetition mode, remapping the original address of a current memory access request to a target address by using an address mapping strategy; performing a data access operation by using the target address; if the target data is cached in a cache level of the GPU, directly returning the same value in the original access address; otherwise, loading data from a global memory and storage. The application avoids frequent global memory and storage access caused by scattered storage of repeated data in a traditional method, significantly improves the cache hit rate, and further breaks through the bottleneck of low cache utilization of an existing strategy.

Description

Technical Field

[0001] This application relates to the fields of computer architecture and parallel computing technology, such as an address mapping method and system for GPU convolution acceleration. Background Technology

[0002] With the widespread application of deep neural networks in fields such as computer vision and natural language processing, their model structures are becoming increasingly complex, with numerous convolutional layers accounting for the majority of the overall computation time. Therefore, improving the speed of convolution operations has become crucial for accelerating neural network training and inference. Among various acceleration techniques, transforming convolution operations into Generalized Matrix Multiplication (GEMM) has become the mainstream implementation method on parallel processors such as GPUs. This method, by reorganizing the input data and convolutional kernels into matrix form, can fully utilize the thread-level parallelism of GPUs to achieve a higher speedup.

[0003] To further optimize the performance of GEMM-based convolutional computation on GPUs, researchers have proposed various strategies primarily from the perspective of memory and storage access optimization, which can be summarized into three categories: First, matrix recombination methods, such as the MEC method that reduces memory and storage usage through specific matrix transformations, or the im2col method that optimizes data locality through implicit channel preference. However, these methods are usually accompanied by significant data recombination overhead or require substantial modifications to existing systems. Second, precision and computation optimization methods, such as the msGeMM method that avoids low-precision data multiplication and addition operations by pre-building lookup tables. However, it faces limitations in storage overhead and construction time in high-precision scenarios. Third, access mode optimization methods, such as the Duplo method that achieves register access reuse through runtime history buffers. Although this method is transparent, it introduces additional hardware overhead. These existing optimization methods are either limited by specific scenarios or introduce significant overhead, failing to systematically and cost-effectively solve the core bottleneck problem of low cache hit rate caused by data duplication in the GEMM workspace matrix. Therefore, there is an urgent need for a deeply optimized cache access method in GPU convolutional computation to address these issues.

[0004] It should be noted that the information disclosed in the background section above is only used to enhance the understanding of the background of this application, and therefore may include information that does not constitute prior art known to those skilled in the art. Summary of the Invention

[0005] To provide a basic understanding of some aspects of the disclosed embodiments, a brief summary is given below. This summary is not intended as a general commentary, nor is it intended to identify key / important components or describe the scope of protection of these embodiments, but rather as a prelude to the detailed description that follows.

[0006] This disclosure provides an address mapping method and system for GPU convolution acceleration, which solves the technical problems of existing GEMM-based convolution optimization strategies being limited by scenarios, increasing additional overhead, and having low cache hit rates.

[0007] In some embodiments, the address mapping method for GPU convolution acceleration is applied during the GPU's execution of convolution operations based on general matrix multiplication to dynamically remap the memory access addresses of the workspace matrix, including the following steps:

[0008] Identify data repetition patterns in the workspace matrix, where data repetition patterns include intra-block element repetition caused by horizontal sliding of the convolution kernel and block repetition caused by vertical sliding of the convolution kernel;

[0009] Based on the data duplication pattern, an address mapping strategy is adopted to remap the original address of the current memory access request to the target address. The address mapping strategy includes at least one of the upper right mapping strategy, the lower left mapping strategy, and the extra workspace mapping strategy.

[0010] Data access operations are performed using the target address. If the target data is already cached in the GPU's cache level, the same value at the original access address is returned directly; otherwise, the data is loaded from global memory and storage.

[0011] In some embodiments, employing an address mapping strategy includes selecting an execution strategy from the address mapping strategies based on the configuration parameters of the current convolutional layer through an adaptive mapping strategy, wherein the execution of the adaptive mapping strategy includes:

[0012] Determine the relationship between the kernel height and the stride;

[0013] Address remapping is not enabled when the kernel height is less than or equal to the stride.

[0014] When the kernel height is greater than the stride, address remapping is enabled, and the address mapping strategy is selected according to the predetermined decision rules, taking into account the global memory and storage resources, the numerical relationship between the kernel height and the stride.

[0015] In some embodiments, the predetermined decision rules include:

[0016] If the difference between the kernel height and the stride is 1, then when there is enough global memory and storage to allocate extra workspace, the extra workspace mapping strategy is selected; when there is insufficient global memory and storage to allocate extra workspace, the top-right mapping strategy is selected.

[0017] If the difference between the kernel height and the stride is greater than or equal to 2, and the stride is 1, then when global memory and storage are sufficient to allocate additional workspace, the additional workspace mapping strategy or the lower left mapping strategy is selected; when global memory and storage are insufficient to allocate additional workspace, the lower left mapping strategy is selected.

[0018] If the stride is greater than or equal to 2, the extra workspace mapping strategy is selected when global memory and storage are sufficient to allocate extra workspace; if global memory and storage are insufficient to allocate extra workspace, the upper right mapping strategy is selected if the kernel height is divisible by the stride, and the lower left mapping strategy is selected if the kernel height is not divisible by the stride.

[0019] In some embodiments, the top-right mapping strategy remaps access requests for duplicate elements to the top-right corner where they first appear in the workspace matrix; the execution process of the top-right mapping strategy includes:

[0020] Convert the original address corresponding to the original memory access request into the global row and column coordinates of the original element;

[0021] Convert global row and column coordinates to block and element coordinates in the original workspace;

[0022] Map block coordinates and element coordinates to their corresponding target block coordinates and target element coordinates;

[0023] A one-dimensional index of the target element is synthesized based on the target block coordinates and the target element coordinates.

[0024] In some embodiments, block coordinates and element coordinates are mapped to their corresponding target block coordinates and target element coordinates, specifically through the following formula:

[0025]

[0026] In the formula, Represents the row coordinates of the target block or the row coordinates of the target element within the block. Represents the column coordinates of the target block or the column coordinates of the target element within the block. and Represents the row and column coordinate indices of the original block or element. Indicates the step size.

[0027] In some embodiments, the lower-left mapping strategy remaps access requests for duplicate elements in the lower-right region of the workspace matrix to the corresponding duplicate element position at the bottom of the lower-right region. The lower-right region refers to the area in the workspace matrix consisting of all elements that can be mapped to the right column, where the right column is the rightmost column of elements in the original workspace. The lower-left mapping strategy execution process includes:

[0028] Convert the original address corresponding to the original memory access request into the global row and column coordinates of the original element;

[0029] Convert global row and column coordinates to block and element coordinates in the original workspace;

[0030] If the original element belongs to the bottom right region, then the original coordinates of the element in the bottom right region are mapped to the corresponding coordinates at the bottom of the bottom right region;

[0031] If the original element does not belong to the lower right region, the block coordinates and element coordinates are mapped to their corresponding target block coordinates and target element coordinates according to the upper right mapping strategy.

[0032] A one-dimensional index of the target element is synthesized based on the target block coordinates and the target element coordinates.

[0033] In some embodiments, the original coordinates of elements in the lower right region are mapped to the corresponding coordinates at the bottom of the lower right region, specifically through the following formula:

[0034]

[0035] In the formula, Represents the row coordinates of the target block or the row coordinates of the target element within the block. Represents the column coordinates of the target block or the column coordinates of the target element within the block. and Represents the row and column coordinate indices of the original block or element. Indicates the step size.

[0036] In some embodiments, the additional workspace mapping strategy involves storing the elements in the rightmost column of all target blocks contiguously in additionally allocated storage space, and remapping accesses to the rightmost column elements to the additionally allocated storage space; the execution process of the additional workspace mapping strategy includes:

[0037] Based on the ratio of CUDA thread block size, output matrix width, and stride, the arrangement of the right column elements in the extra workspace is determined to be one of column-first, row-first, or N-shaped arrangement. Column-first arrangement means that all right column elements are arranged in column-first order, so that elements in the same right column are consecutive in the storage space. Row-first arrangement means that elements in the same row of all right columns are consecutive in the storage space. N-shaped arrangement means that elements in the same right column are consecutive in the vertical direction, right columns between different blocks are consecutive in the horizontal direction, and all right column elements are expanded into the extra workspace in N-shaped order.

[0038] In some embodiments, the arrangement of the right column elements in the additional workspace is determined based on the ratio of the CUDA thread block size, the output matrix width, and the stride, including:

[0039] pass Calculate the ratio, where, For CUDA thread block size, To output the matrix width, Step size;

[0040] If the ratio is less than or equal to the first preset value, then the sorting method is column-major sorting;

[0041] If the ratio is greater than the first preset value and less than or equal to the second preset value, then the arrangement method is selected as N-shaped arrangement, wherein the first preset value is less than the second preset value;

[0042] If the ratio is greater than the second preset value, then the sorting method is row-major.

[0043] In some embodiments, an address mapping system for GPU convolution acceleration is used to execute any of the address mapping methods for GPU convolution acceleration described above.

[0044] The address mapping method and system for GPU convolution acceleration provided in this disclosure can achieve the following technical effects:

[0045] This application is applied to scenarios where GPUs perform convolution operations based on general matrix multiplication. First, it identifies data repetition patterns in the workspace matrix caused by horizontal sliding of the convolution kernel (internal element repetition) and vertical sliding (block repetition). This identification process does not rely on specific convolution kernel sizes, input scales, or data precision, breaking the scenario limitations of existing optimization strategies and adapting to various convolution configurations based on general matrix multiplication. Second, based on the aforementioned data repetition patterns, it dynamically remaps the original memory access addresses using upper-right mapping, lower-left mapping, and additional workspace mapping strategies. This unifies access requests that originally pointed to different addresses but contained duplicate data, directing them to the target address. This remapping avoids the overhead of modifying matrix recombination methods and hardware maintenance overhead, achieving optimization only through address logic translation, thus reducing the additional overhead of existing strategies. Finally, when performing data access using the remapped target address, since the data corresponding to the target address has most likely been cached in the GPU cache level due to previous repeated accesses, a cache hit can be directly achieved. This avoids the frequent global memory and storage access caused by the scattered storage of duplicate data in traditional methods, significantly improving the cache hit rate and thus breaking through the bottleneck of low cache utilization in existing strategies.

[0046] The above general description and the description below are exemplary and illustrative only and are not intended to limit this application. Attached Figure Description

[0047] One or more embodiments are illustrated by way of example with reference to the accompanying drawings. These illustrations and drawings do not constitute a limitation on the embodiments. Elements having the same reference numerals in the drawings are shown as similar elements. The drawings are not to be scaled. And wherein:

[0048] Figure 1 This is a flowchart of an address mapping method for GPU convolution acceleration provided in an embodiment of this disclosure;

[0049] Figure 2 This is a schematic diagram of a data repetition pattern provided in an embodiment of this disclosure;

[0050] Figure 3 This is a schematic diagram of an upper-right mapping strategy provided in an embodiment of this disclosure;

[0051] Figure 4 This is a schematic diagram of a lower left mapping strategy provided in an embodiment of this disclosure;

[0052] Figure 5 This is a schematic diagram of an additional workspace mapping strategy provided in an embodiment of this disclosure;

[0053] Figure 6 This is a flowchart of an adaptive mapping strategy selection provided in an embodiment of this disclosure;

[0054] Figure 7 This is an implementation architecture diagram of an address mapping system for GPU convolution acceleration provided in this disclosure. Detailed Implementation

[0055] To provide a more detailed understanding of the features and technical content of the embodiments of this disclosure, the implementation of the embodiments of this disclosure will be described in detail below with reference to the accompanying drawings. The accompanying drawings are for illustrative purposes only and are not intended to limit the embodiments of this disclosure. In the following technical description, for ease of explanation, several details are used to provide a full understanding of the disclosed embodiments. However, one or more embodiments may still be implemented without these details. In other cases, well-known structures and devices may be simplified in their depiction to simplify the drawings.

[0056] The terms "first," "second," etc., used in the embodiments of this disclosure are used to distinguish similar objects and are not necessarily used to describe a specific order or sequence. It should be understood that such data can be interchanged where appropriate for the embodiments of this disclosure described herein. Furthermore, the terms "comprising" and "having," and any variations thereof, are intended to cover non-exclusive inclusion.

[0057] Unless otherwise stated, the term "multiple" means two or more.

[0058] In this embodiment of the disclosure, the character " / " indicates that the objects before and after it are in an "or" relationship. For example, A / B means: A or B.

[0059] The term "and / or" describes an association between objects, indicating that three relationships can exist. For example, A and / or B means: A or B, or A and B.

[0060] The term "correspondence" can refer to an association or binding relationship. The correspondence between A and B means that there is an association or binding relationship between A and B.

[0061] Modern deep neural networks contain numerous stacked convolutional layers, whose execution time accounts for a significant portion of the overall process. Therefore, accelerating convolutional operations is crucial for improving the overall network performance. The logic of direct convolution involves the convolutional kernel sliding from the top left corner to the bottom right corner of the input matrix, calculating the inner product between the kernel and the corresponding region of the input matrix at each sliding position, and then generating the elements of the output matrix. To accelerate this process, researchers have proposed traditional optimization methods such as FFT and Winograd. However, FFT is not suitable for handling small-scale input matrices and convolutional kernels, and Winograd is only effective for small convolutional kernels such as 3×3 and 5×5. Furthermore, neither method can handle convolutions with non-unit strides and requires additional storage space for data transformation, resulting in significant space overhead.

[0062] In contrast, methods based on Generalized Matrix Multiplication (GEMM) can adapt to all convolution configurations and fully utilize the thread-level parallelism of GPUs to achieve higher speedup ratios, becoming an important direction for GPU convolution acceleration. To further optimize the performance of GEMM-based convolution operations on GPUs, researchers have proposed three types of memory and storage access optimization strategies from different perspectives, but these strategies still have shortcomings: Matrix recombination methods (such as MEC, implicit channel-first im2col, etc.) reduce memory and storage overhead through matrix transformation or data reorganization, but some require redesigning the entire data flow, which involves significant modifications to existing systems, and some have matrix recombination overhead, limiting their applicability and generality; Precision optimization methods (such as msGeMM) accelerate computation by pre-compiling lookup tables for low-precision data, which only adapts to low-precision neural networks, and the storage and construction of the lookup tables incur additional overhead, failing to cover high-precision scenarios; Access pattern optimization methods (such as Duplo) dynamically track data access redirection register access, which does not require modification of application code, but maintaining the loading history buffer and pattern matching incurs additional hardware overhead.

[0063] As mentioned above, current convolution optimization strategies based on GEMM are either limited by the scenario (such as adapting to accuracy or limited convolution configuration) or require additional overhead (such as data reorganization, lookup table storage, and hardware maintenance overhead). Furthermore, none of them have effectively solved the core problem of low cache hit rate, becoming a key bottleneck restricting further improvement in GPU convolution operation performance.

[0064] To address the aforementioned problems, this disclosure provides an address mapping method and system for GPU convolution acceleration. Applied to the GPU's execution of convolution operations based on general matrix multiplication, this application identifies repetitive distribution patterns in the workspace matrix generated by convolution using general matrix multiplication. It remaps the original address requested by the thread to the address of an element in a non-repeating region already in the cache, thereby bypassing underlying storage access and significantly reducing memory access latency. Compared to the original path's "fetch data by address, access global memory and storage if a miss occurs" process, the address mapping strategy allows duplicate elements to be redirected to the target element in the cache for direct return, reducing unnecessary subsequent memory processes and significantly accelerating convolution calculations based on general matrix multiplication. The address mapping method provided by this disclosure is described in detail below with reference to the accompanying drawings.

[0065] Figure 1 This is a flowchart of an address mapping method for GPU convolution acceleration provided in this disclosure embodiment, combined with... Figure 1 As shown, the address mapping method includes the following steps:

[0066] S101: Identify data repetition patterns in the workspace matrix, wherein the data repetition patterns include intra-block element repetition caused by horizontal sliding of the convolution kernel, and block repetition caused by vertical sliding of the convolution kernel.

[0067] In some embodiments, in convolution operations based on general matrix multiplication, the convolution kernel expands the data covered at each position as it traverses the original input matrix and connects these expanded fragments to form a working space matrix. Figure 2 This is a schematic diagram of a data repetition pattern provided in an embodiment of this disclosure. The colors and numbers in the diagram are corresponding; identical numbers and colors indicate repetition. (Combined with...) Figure 2 As shown, the movement of the convolution kernel includes two methods: horizontal sliding and vertical sliding. Horizontal sliding refers to the convolution kernel gradually moving from the left side of the input matrix to the right according to a predetermined stride until it reaches the end of the row, as shown below. Figure 2 The process from HStep0 to HStep2 is presented in the diagram; while vertical sliding involves moving the convolution kernel a specific step downwards after each complete horizontal traversal, such as... Figure 2The process from VStep0 to VStep2 is shown in the diagram. During horizontal sliding, elements in the same row of the input matrix converge into a continuous region in the workspace. This application refers to such a region as a "row block" (or simply "block"), and its dimension corresponds to the product of the output width and the height of the convolution kernel. For example, in the case of HStep, the 3×3 region with index 0 in the upper left corner of the workspace forms the block with index 0, where the elements in the 0th row have indices from 0 to 4. When sliding vertically, although the convolution kernel captures the same elements in the same row of the input matrix, these elements are scattered into different blocks in the workspace. For example, elements 10, 11, and 12 appear in blocks 2, 4, and 6 respectively during VStep0, VStep1, and VStep2. These two different sliding operations generate two distinct data duplication patterns in the workspace: one is intra-block element duplication, where some elements within a block are repeated due to horizontal sliding, such as element 2 appearing multiple times in block 0; the other is block duplication, where the entire block contains identical elements due to vertical sliding, such as blocks 1 and 3 having completely identical content. Therefore, in the workspace matrix, the distribution pattern of duplicate elements with the same value but different memory addresses consistently exhibits a directional characteristic extending from the upper right to the lower left. This directionality provides a clear characteristic basis for subsequent address mapping strategies based on data duplication patterns.

[0068] S102: Based on the data duplication pattern, an address mapping strategy is adopted to remap the original address of the current memory access request to the target address. The address mapping strategy includes at least one of the upper right mapping strategy, the lower left mapping strategy, and the extra workspace mapping strategy.

[0069] The address mapping strategy disclosed in the embodiments of this application will now be described in conjunction with the accompanying drawings.

[0070] In some embodiments, the top-right mapping strategy remaps access requests for duplicate elements to the top-right corner where the element first appears in the workspace matrix. This strategy leverages the element loading characteristics of the workspace matrix, where elements are loaded into the GPU cache in a left-to-right, top-to-bottom order. Therefore, remapping access requests for duplicate elements to the top-right corner where the element first appears in the workspace matrix utilizes the temporal locality of the cache, significantly improving cache hit rates during repeated accesses. Figure 3This is a schematic diagram of a top-right mapping strategy provided in an embodiment of this disclosure. To accurately identify the target address, both block repetition and element repetition modes need to be utilized simultaneously; therefore, two levels of mapping—block mapping and element mapping—must be implemented. First, the corresponding target block is located, which is usually located in the top-right direction of the source block. Then, the target element is found within this block; the target element is always the top-rightmost element among all repeating elements. All repeating elements are mapped to the top-right, as shown in the white area in Figure 3. This area can be divided into "upper row" and "right column." Combined with... Figure 3 The execution process of the top-right mapping strategy is as follows:

[0071] First, the original address corresponding to the original memory access request is converted into the global row and column coordinates of the original element. This application uses IDs in the range of 0 to (K×M-1) to distinguish elements in the workspace, converting the one-dimensional index of the original element in the workspace into the global row and column coordinates of that element in the workspace matrix, realizing the mapping from "one-dimensional address to two-dimensional spatial position", laying the foundation for subsequent splitting of block coordinates and element coordinates. The index of the global row and column coordinates can be calculated using the following formula:

[0072]

[0073] In the formula, The one-dimensional index of each element in the original workspace matrix, ranging from [0, K×M-1] (where K is the cumulative dimension length and M is the M-dimensional workspace dimension length), is used to uniquely identify the initial position of the original element in global memory and storage. This represents the global row index of the element in the workspace matrix. This represents the global column index of the element in the workspace matrix.

[0074] Then, the global row and column coordinates are converted into block and element coordinates in the original workspace. The original one-dimensional matrix in the workspace can be converted into the original block and element row and column coordinate indices using the global row and column coordinate indices. The calculation formula is as follows:

[0075]

[0076]

[0077] In the formula, Represents the row coordinates of the original block (identifying the row position of the element's row block within all blocks). Represents the column coordinates of the original block (identifying the column position of the row block to which the element belongs within all blocks). Indicates the row index of the element within the original block. This indicates the column index of the element within the original block. Indicates the height of the convolution kernel. Indicates the height of the output matrix. Indicates the kernel width. This indicates the width of the output matrix.

[0078] Secondly, the block coordinates and element coordinates are mapped to their corresponding target block coordinates and target element coordinates, which is achieved through the following formula:

[0079]

[0080] In the formula, Represents the row coordinates of the target block or the row coordinates of the target element within the block. Represents the column coordinates of the target block or the column coordinates of the target element within the block. and This represents the row and column coordinate indices of the original block or element. The asterisk (*) is a wildcard that can represent either block coordinates or element coordinates within a block; this formula is used for both block mappings and element mappings. Indicates the step size.

[0081] Finally, based on the target block coordinates and the target element coordinates, a one-dimensional index of the target element is synthesized, realizing the conversion from "two-dimensional target location to one-dimensional memory address," providing an executable address for the final cache / memory access. This is specifically achieved through the following formula:

[0082]

[0083] In the formula, This represents the global one-dimensional index of the target element in the workspace matrix. and Indicates the row / column coordinates of the target block. and This indicates the row / column coordinates of the target element within the target block.

[0084] The top-right mapping strategy can significantly improve computation speed because the target element is likely to be repeatedly accessed within a short period and is therefore highly likely to remain in the cache. When the requested duplicate element is not in the cache, but the target element is, the corresponding data will be read directly from the cache. However, the top-right mapping strategy also has its problems, such as... Figure 3 In the row with row index 1, although elements 1 and 2 can be mapped to their target elements in the row with row index 0, requesting 3 in row 1 will still result in a cache miss, and that row will still be loaded into the cache. In this case, the first three rows in the example all need to be loaded into the cache.

[0085] To address the issue of optimizing the top-right mapping strategy, this application proposes a bottom-left mapping strategy.

[0086] In some embodiments, the lower-left mapping strategy is to remap access requests for duplicate elements in the lower-right region of the workspace matrix to the corresponding duplicate element position at the bottom of the lower-right region. The lower-right region refers to the region consisting of all elements in the workspace matrix that can be mapped to the right column, and the right column is the rightmost column of elements in the original workspace. Figure 4 This is a schematic diagram of a lower left mapping strategy provided in an embodiment of this disclosure, combined with... Figure 4 All elements 13 can be mapped to the white 13, instead of the 13 in the row with row index 1. The same applies to 3 and 8. This way, the row with row index 1 doesn't need to be loaded into the cache, thus improving cache utilization compared to the top-right mapping strategy. Specifically, the "bottom-right region" is defined as the region containing all elements that can be mapped to the "right column". For example, for a block whose "right column" contains 13 and 14, its bottom-right region, as shown by the double-lined box in Figure 4, contains two 13s and one 14. In this block, element 12 belongs to the "upper row" element. The bottom-left mapping strategy region of the entire workspace matrix contains all blocks that can be mapped to the right column from a block perspective. For example, blocks 5, 7, and 8 constitute the bottom-left mapping strategy region of the workspace matrix.

[0087] In some embodiments, the positioning logic of the lower right region is as follows: First, the workspace matrix is ​​divided into independent row blocks according to the dimension of "output width × kernel height"; then, the rightmost column of each row block is determined, which is the "right column" of the block, excluding the "upper row" elements; second, elements within the row block whose values ​​are the same as the right column elements and can be associated with the right column in the repeating direction from "upper right to lower left" are selected; finally, the continuous region formed by these selected elements is the "lower right region" of the block. Figure 4 (Use double-lined boxes to annotate).

[0088] In some embodiments, the lower-left mapping strategy execution process includes:

[0089] First, the original address corresponding to the original memory access request is converted into the global row and column coordinates of the original element.

[0090] Then, the global row and column coordinates are converted into block coordinates and element coordinates in the original workspace.

[0091] Secondly, determine whether the original element belongs to the bottom right region. If the original element belongs to the bottom right region, map the original coordinates of the element in the bottom right region to the corresponding coordinates at the bottom of the bottom right region; if the original element does not belong to the bottom right region, map the block coordinates and element coordinates to their corresponding target block coordinates and target element coordinates according to the top right mapping strategy.

[0092] The mapping of the original coordinates of elements in the lower right region to the corresponding coordinates at the bottom of the lower right region is achieved through the following formula:

[0093]

[0094] For the top-right mapping formula (Formula 4) and the bottom-left mapping formula (Formula 6), from the perspective of the mapping direction, the bottom-left mapping formula maps to the bottom left, which is manifested as the row coordinate shifting downwards and the column coordinate shifting to the left; while the top-right mapping formula maps to the top right, which is manifested as the row coordinate shifting upwards and the column coordinate shifting to the right. In calculating the offset multiplier, the offset multiplier K of the bottom-left mapping is obtained through... Calculate (numerator is the distance from the original row coordinates to the bottom of the block); the offset factor K of the top-right mapping is obtained through... Calculations are performed (directly based on the original row coordinates). For row coordinate adjustment, the bottom-left mapping is the original row coordinate plus an offset multiplier multiplied by the step size, shifting the row coordinate downwards; the top-right mapping is the original row coordinate minus an offset multiplier multiplied by the step size, shifting the row coordinate upwards. For column coordinate adjustment, the bottom-left mapping is the original column coordinate minus an offset multiplier, shifting the column coordinate to the left; the top-right mapping is the original column coordinate plus an offset multiplier, shifting the column coordinate to the right.

[0095] Finally, a one-dimensional index of the target element is synthesized based on the target block coordinates and the target element coordinates.

[0096] In summary, compared to the top-right mapping strategy, the new mapping scheme in the bottom-right region maps previously non-adjacent elements to adjacent positions, further improving data locality. However, the cache utilization of the bottom-left mapping strategy is still not optimal, for example, in... Figure 4 In the cache, elements 2 through 14 in row 2 will not be accessed because elements 2, 7, and 12 have already been mapped to the target element in row 0. Therefore, even after this row is loaded into the cache, some space will still be wasted.

[0097] To further improve cache utilization and completely resolve the continuity issue, this application stores the "right column" elements in a contiguous storage space. This makes the requested data more contiguous in the cache, i.e., it employs an extra workspace mapping strategy. The extra workspace mapping strategy involves storing the elements in the rightmost column of all target blocks contiguously in an extra-allocated storage space and remapping accesses to the right column elements to the extra-allocated storage space.

[0098] In some embodiments, since the CUDA (Compute Unified Device Architecture) thread block size, output matrix width, and stride affect the loading order of extra workspace elements, the extra workspace mapping strategy determines the arrangement of the right column elements in the extra workspace as either column-first, row-first, or N-shaped, based on the ratio of the CUDA thread block size, output matrix width, and stride.

[0099] Figure 5 This is a schematic diagram of an additional workspace mapping strategy provided in an embodiment of this disclosure. Column-first arrangement means that all elements in the right column are arranged in column-first order, so that elements in the same right column are contiguous in the storage space, emphasizing locality in the vertical direction. For example... Figure 5 As shown, the "right columns" of all target blocks are arranged in column-major order. For example, 10, 11, and 12 are arranged consecutively first, followed by 23, 24, and 25 in the same column, until all "right columns" of the current column are sorted. After completing the same column, the adjacent right-hand block columns 36, 37, and 38 are processed. All elements in the "right columns" are expanded into the extra workspace in a top-to-bottom, left-to-right order, so that the "right columns" are adjacent in the vertical direction but not in the horizontal direction. When the convolution kernel and stride are large, they may not be in the same cache sector. If the CUDA thread block is large, one block may cover the "right columns" of multiple blocks. In this case, the same rows in the covered "right columns" may be executed in the same cycle, resulting in lower cache utilization for column-major sorting. Therefore, column-major sorting is suitable for cases with small CUDA thread blocks and strides. It should be noted that... Figure 5 Taking a convolution stride of 2 as an example, the "up row" elements include the elements in rows 0 and 1. Taking block 0 as an example, all the "right column" elements include 8 and 9, and 8 and 9 belong to the "up row" elements.

[0100] The row-first arrangement refers to the contiguous storage of elements in the same row within all elements in the right column, emphasizing locality in the horizontal direction. For example... Figure 5 As shown, when using row-major ordering, all elements in the same row of the "right column" are arranged consecutively in the extra workspace, such as 10, 36, and 62 in the second row. All elements are expanded in order from left to right and from top to bottom. When the CUDA thread block is small, the same row in the "right column" of two adjacent blocks may not be executed in the same SM (Streaming Multiprocessor), leading to low utilization. Therefore, row-major ordering is suitable for situations where the CUDA thread block is large and the output size is small.

[0101] The N-shaped arrangement refers to elements in the same right column being vertically continuous, right columns between different blocks being horizontally continuous, and all right column elements being expanded into the additional workspace in an N-shaped order. The N-shaped arrangement combines the advantages of column-first and row-first order, while also considering vertical and horizontal locality. Figure 5 Elements within the same "right column" are vertically contiguous, while "right columns" between different blocks are horizontally contiguous. For example, first 10, 11, 12, then adjacent 36, 37, 38, expanding in an N-shape into additional workspace. When a CUDA thread block is large enough to cover the "right columns" of multiple blocks, although identical rows may not be contiguous in memory and storage, they are likely to reside in the same cache line or even the same sector due to proximity. When a CUDA thread block is small, the "right columns" within a block are contiguous, which effectively utilizes locality when loading the current block.

[0102] In some embodiments, when selecting an additional workspace mapping strategy, the arrangement order needs to be determined based on the CUDA thread block size, output width (Q), and stride (S). The output width determines the distance between two "right columns" in the horizontal direction, the CUDA thread block size determines how many "right columns" can be accommodated within the thread block, and the stride determines the number of "right columns" in the vertical direction. Therefore, through... Calculate the ratio, where, For CUDA thread block size, To output the matrix width, The step size.

[0103] In some embodiments, based on statistics of dozens of commonly used convolutional layers, a first preset value of 0.8 and a second preset value of 2 are set. Based on these two preset values, the results can be divided into the following three intervals:

[0104] When the ratio is ≤0.8, a CUDA thread block typically only processes one "right column" of the block. As the step size increases, more elements in the same column will be processed, so column-first sorting is chosen.

[0105] When 0.8 < ratio ≤ 2, a CUDA thread block may span multiple blocks. In order to take into account both horizontal and vertical locality, an N-shaped arrangement is preferred.

[0106] When the ratio is greater than 2, horizontal locality takes precedence over vertical locality, so row-major ordering should be chosen.

[0107] In some embodiments, the additional workspace mapping strategy uses additional GPU memory to store the elements in the "right column," as described in this application. The global index representing the additional workspace has the following range: ,in, Indicates the batch size. Indicates the number of channels in the input matrix. Indicates the height of the output matrix. This indicates the height of the convolution kernel.

[0108] Since different CUDA thread blocks, strides, and output matrix widths affect the optimal arrangement order of the extra workspace matrix, and different arrangement orders affect data locality and final performance, the formulas for three different arrangements of the extra workspace mapping strategy proposed in this application are as follows:

[0109]

[0110]

[0111]

[0112]

[0113] In this context, formula 7 represents the target index for column-major ordering; formulas 8 and 9 represent the mapping of the non-bottom-right regions for row-major and N-shaped orderings, respectively; and formula 10 represents the mapping of the bottom-right regions for row-major and N-shaped orderings. Figure 5 Taking row 49 with row index 9 as an example, it is first mapped to row 49 in the "right column" of row with row index 7. When using column-major sorting, there is... The index 9 is used to locate 49 in the additional workspace.

[0114] In summary, this application systematically identifies and utilizes the element and block repetition patterns in the GEMM workspace matrix in the cache. By dynamically remapping memory access requests to addresses of high-probability cached duplicate data, it overcomes the technical bottleneck of low workspace matrix cache hit rate in traditional methods. This innovative design not only maintains the functional integrity of the original GEMM operation but also introduces three efficient address mapping strategies: the upper-right mapping strategy implements mapping based on the first-occurrence principle, the lower-left mapping strategy optimizes data locality layout, and the additional workspace mapping strategy achieves optimal cache utilization by reorganizing key data.

[0115] In some embodiments, an adaptive mapping strategy is proposed based on the features of different convolutional layer configurations. This strategy can dynamically select the most suitable mapping strategy according to the configuration parameters of the current convolutional layer, which greatly improves the computational efficiency of the GPU when performing convolution operations. Figure 6 This is a flowchart of an adaptive mapping strategy selection provided in an embodiment of this disclosure, combined with... Figure 6 The execution process of the adaptive mapping strategy is as follows:

[0116] Determine the relationship between the kernel height F and the stride S. When there are no duplicate elements in the workspace, address remapping is not enabled. When the kernel height is greater than the stride, address remapping is enabled, and the address mapping strategy is selected according to a predetermined decision rule, taking into account global memory and storage resources, the numerical relationship between kernel height and stride.

[0117] In some embodiments, the predetermined decision rule is specifically as follows:

[0118] If the difference between the kernel height and the stride is 1, that is... (As in ResNet's C3 layer), only one element is accessed per sector, resulting in low space utilization. In this case, an additional workspace mapping strategy can be used to achieve the highest locality. Therefore, when global memory and storage are sufficient to allocate additional workspace, the additional workspace mapping strategy is selected; when global memory and storage are insufficient to allocate additional workspace, the top-right mapping strategy is selected.

[0119] If the difference between the kernel height and the stride is greater than or equal to 2 and the stride is 1, then... 2 and All elements in the "right column" can be placed at the bottom of the lower right area, achieving optimal results without additional storage space. If an additional workspace mapping strategy is selected, comparable performance can also be obtained. Therefore, when global memory and storage are sufficient to allocate additional workspace, the aforementioned additional workspace mapping strategy or the lower left mapping strategy should be selected; when global memory and storage are insufficient to allocate additional workspace, the lower left mapping strategy should be selected.

[0120] If the step size is greater than or equal to 2, that is If global memory and storage are sufficient to allocate additional workspace, the additional workspace mapping strategy is selected; if global memory and storage are insufficient to allocate additional workspace, if the kernel height is divisible by the stride, i.e., F%S == 0 (such as the C4 layer of DCGAN), the upper-right mapping strategy is selected; if the kernel height is not divisible by the stride, the lower-left mapping strategy is selected.

[0121] S103: Perform data access operation using the target address. If the target data is already cached in the GPU's cache level, return the same value at the original access address directly; otherwise, load the data from global memory and storage.

[0122] In some embodiments, if the target data is frequently and repeatedly mapped and accessed, it is highly likely that it has already resided in the GPU's cache level. If a hit occurs, that is, if the target data in the data access operation is already cached in the GPU's cache level, the same value in the original access address is returned directly; otherwise, the data is loaded from global memory and storage.

[0123] The address mapping method in this application focuses on the caching level, improving caching behavior and accelerating GEMM operations by discovering and utilizing data repetition patterns in the workspace matrix. This method achieves significant acceleration without modifying the matrix storage format or limiting network and data precision, exhibiting good versatility and can be combined with traditional caching optimization techniques such as prefetching and cache locking for even greater performance improvements. Therefore, this application analyzes the repetition patterns in GEMM-based convolution computation at the caching level and proposes various address mapping strategies and their variants based on these patterns, improving cache hit rate. Furthermore, an adaptive mapping strategy is proposed to address the configuration characteristics of different convolutional layers, automatically selecting the optimal mapping strategy based on the configuration of different convolutional layers, greatly improving computational efficiency.

[0124] Based on the same inventive concept as the address mapping method for GPU convolution acceleration, this application also discloses an address mapping system for GPU convolution acceleration in some embodiments. This system is used to execute the address mapping method for GPU convolution acceleration disclosed in any of the above embodiments, and the process will not be described in detail here.

[0125] Figure 7 This is an implementation architecture diagram of an address mapping system for GPU convolution acceleration provided in this disclosure. Combined with... Figure 7 To implement the address mapping method described above, the address mapping matrix can be specified during programming. For example, the newly defined function `cudaMalloc_map()` can be used instead of `cudaMalloc()`. The compiler then inherits the mapping tag and generates corresponding instructions, such as `ld_map()` and `st_map()`, for the corresponding memory accesses. When these instructions are executed in the GPU, instructions marked with "_map" are sent to the proposed address mapping module. In this module, an appropriate mapping strategy is selected and enabled according to the adaptive mapping strategy, converting the original memory addresses into remapped addresses. After this process, the `ld_map()` and `st_map()` instructions are converted into `ld()` and `st()` instructions with remapped address parameters. These instructions, along with other `ld()` and `st()` instructions, are sent to the memory merging unit module for subsequent memory merging processing.

[0126] In addition, to verify the effectiveness of the address mapping method for GPU convolution acceleration in this application, the following experimental configuration was used and the performance results were obtained:

[0127] (1) Regarding the experimental setup

[0128] Simulation environment: Based on the GPGPU-Sim 4.0.0 simulator, with configurations comparable to Titan V (Volta architecture); Software environment: CUDA SDK 9.1, GEMM calls based on cuBLAS, and column-first storage in the workspace; Test benchmark: ResNet, DCGAN, and ConvNeXt models (covering most convolutional configurations) were selected, with traditional GPU convolutional operations without enabling the address mapping strategy of this invention serving as the comparison benchmark.

[0129] (2) Regarding the experimental results

[0130] Under the above configuration, the address mapping strategy proposed in this application achieves performance improvements on different computing units of the GPU. On CUDA Cores: the top-right mapping strategy achieves an average speedup of 17.7%, the bottom-left mapping strategy achieves an average speedup of 24.6%, the extra workspace mapping strategy achieves an average speedup of 24.9%, and the combined adaptive mapping strategy achieves an average speedup of 25.7%. On Tensor Cores, the average improvement reaches 21.0%.

[0131] The above results demonstrate that this application effectively improves the cache utilization and computational efficiency of GPU convolution operations by optimizing the cache access mode of the workspace matrix.

[0132] The technical solutions of this disclosure can be embodied in the form of a software product. This computer software product is stored in a storage medium and includes one or more instructions executed by a computer device (which may be a personal computer, server, or network device, etc.) to perform all or part of the steps of the method described in this disclosure. The aforementioned storage medium can be a non-transitory storage medium, including: USB flash drive, portable hard drive, read-only memory (ROM), random access memory (RAM), magnetic disk, or optical disk, and other media capable of storing program code; it can also be a transient storage medium.

[0133] The foregoing description and accompanying drawings fully illustrate embodiments of this disclosure to enable those skilled in the art to practice them. Other embodiments may include structural, logical, electrical, procedural, and other changes. The embodiments represent only possible variations. Individual components and functions are optional unless explicitly required, and the order of operation may vary. Parts and features of some embodiments may be included in or replace parts and features of other embodiments. Moreover, the terminology used in this application is for descriptive purposes only and is not intended to limit the scope of protection. As used in the description herein, the singular forms “a,” “an,” and “the” are intended to equally include the plural forms unless the context clearly indicates otherwise. Similarly, the term “and / or” as used herein means including one or more of the associated listed items and all possible combinations thereof. Additionally, when used in this application, the term "comprise" and its variations "comprises" and / or "comprising" refer to the presence of stated features, integrals, steps, operations, elements, and / or components, but do not exclude the presence or addition of one or more other features, integrals, steps, operations, elements, components, and / or groups thereof. Without further limitations, an element defined by the phrase "comprises a..." does not exclude the presence of other identical elements in the process, method, or apparatus that includes said element. In this document, each embodiment may focus on the differences from other embodiments, and similar or identical parts between embodiments can be referred to mutually. For methods, products, etc., disclosed in the embodiments, if they correspond to the method section disclosed in the embodiments, the relevant parts can be referred to the description of the method section.

[0134] Those skilled in the art will recognize that the units and algorithm steps of the various examples described in conjunction with the embodiments disclosed herein can be implemented in electronic hardware, or a combination of computer software and electronic hardware. Whether these functions are implemented in hardware or software depends on the specific application and design constraints of the technical solution. Those skilled in the art can use different methods to implement the described functions for each specific application, but such implementation should not be considered beyond the scope of the embodiments of this disclosure. Those skilled in the art will clearly understand that, for the sake of convenience and brevity, the specific working processes of the systems, devices, and units described above can be referred to the corresponding processes in the foregoing method embodiments, and will not be repeated here.

[0135] The methods and products (including but not limited to devices and equipment) disclosed in the embodiments herein can be implemented in other ways. For example, the device embodiments described above are merely illustrative. For instance, the division of units may be merely a logical functional division, and in actual implementation, there may be other division methods. For example, multiple units or components may be combined or integrated into another system, or some features may be ignored or not executed. In addition, the coupling or direct coupling or communication connection shown or discussed between each other may be through some interfaces, and the indirect coupling or communication connection between devices or units may be electrical, mechanical, or other forms. The units described as separate components may or may not be physically separate. The components shown as units may or may not be physical units, that is, they may be located in one place or distributed across multiple network units. Some or all of the units may be selected to implement this embodiment according to actual needs. In addition, the functional units in the embodiments of this disclosure may be integrated into one processing unit, or each unit may exist physically separately, or two or more units may be integrated into one unit.

Claims

1. An address mapping method for GPU convolution acceleration, characterized in that, The method for dynamically remapping the memory access addresses of the workspace matrix during convolution operations based on general matrix multiplication performed by a GPU includes the following steps: Identify data repetition patterns in the workspace matrix, wherein the data repetition patterns include intra-block element repetition caused by horizontal sliding of the convolution kernel, and block repetition caused by vertical sliding of the convolution kernel; Based on the data duplication pattern, an address mapping strategy is adopted to remap the original address of the current memory access request to the target address. This address mapping strategy includes at least one of the following: an upper-right mapping strategy, a lower-left mapping strategy, and an additional workspace mapping strategy. The upper-right mapping strategy remaps access requests for duplicate elements to the upper-right corner where they first appear in the workspace matrix. The lower-left mapping strategy remaps access requests for duplicate elements in the lower-right region of the workspace matrix to the corresponding duplicate element position at the bottom of that region. The lower-right region refers to the area in the workspace matrix that can be mapped to the right column, where the right column is the rightmost column of elements in the original workspace. The additional workspace mapping strategy stores all elements in the rightmost column of all target blocks contiguously in additional allocated storage space and remaps accesses to the right column elements to the additional allocated storage space. Data access operations are performed using the target address. If the target data is already cached in the GPU's cache level, the same value at the original access address is returned directly; otherwise, the data is loaded from global memory and storage.

2. The address mapping method for GPU convolution acceleration according to claim 1, characterized in that, The adoption of the address mapping strategy includes selecting an execution strategy from the address mapping strategy based on the configuration parameters of the current convolutional layer through an adaptive mapping strategy, wherein the execution of the adaptive mapping strategy includes: Determine the relationship between the kernel height and the stride; Address remapping is not enabled when the kernel height is less than or equal to the stride. When the kernel height is greater than the stride, address remapping is enabled, and the address mapping strategy is selected according to the predetermined decision rules, taking into account the global memory and storage resources, the numerical relationship between the kernel height and the stride.

3. The address mapping method for GPU convolution acceleration according to claim 2, characterized in that, The predetermined decision rules include: If the difference between the kernel height and the stride is 1, then when global memory and storage are sufficient to allocate additional workspace, the additional workspace mapping strategy is selected; when global memory and storage are insufficient to allocate additional workspace, the upper right mapping strategy is selected. If the difference between the kernel height and the stride is greater than or equal to 2, and the stride is 1, then when global memory and storage are sufficient to allocate additional workspace, the additional workspace mapping strategy or the lower left mapping strategy is selected; when global memory and storage are insufficient to allocate additional workspace, the lower left mapping strategy is selected. If the stride is greater than or equal to 2, the extra workspace mapping strategy is selected when global memory and storage are sufficient to allocate extra workspace; if global memory and storage are insufficient to allocate extra workspace, the upper-right mapping strategy is selected if the kernel height is divisible by the stride, and the lower-left mapping strategy is selected if the kernel height is not divisible by the stride.

4. The address mapping method for GPU convolution acceleration according to claim 1, characterized in that, The execution process of the top-right mapping strategy includes: Convert the original address corresponding to the original memory access request into the global row and column coordinates of the original element; Convert global row and column coordinates to block and element coordinates in the original workspace; Map block coordinates and element coordinates to their corresponding target block coordinates and target element coordinates; A one-dimensional index of the target element is synthesized based on the target block coordinates and the target element coordinates.

5. The address mapping method for GPU convolution acceleration according to claim 4, characterized in that, The mapping of block coordinates and element coordinates to their corresponding target block coordinates and target element coordinates is specifically achieved through the following formula: In the formula, Represents the row coordinates of the target block or the row coordinates of the target element within the block. Represents the column coordinates of the target block or the column coordinates of the target element within the block. and Represents the row and column coordinate indices of the original block or element. Indicates the step size.

6. The address mapping method for GPU convolution acceleration according to claim 1, characterized in that, The execution process of the lower left mapping strategy includes: Convert the original address corresponding to the original memory access request into the global row and column coordinates of the original element; Convert global row and column coordinates to block and element coordinates in the original workspace; If the original element belongs to the lower right region, then the original coordinates of the element in the lower right region are mapped to the corresponding coordinates at the bottom of the lower right region; If the original element does not belong to the lower right region, the block coordinates and element coordinates are mapped to their corresponding target block coordinates and target element coordinates according to the upper right mapping strategy. A one-dimensional index of the target element is synthesized based on the target block coordinates and the target element coordinates.

7. The address mapping method for GPU convolution acceleration according to claim 6, characterized in that, The mapping of the original coordinates of elements in the lower right region to the corresponding coordinates at the bottom of the lower right region is achieved through the following formula: In the formula, Represents the row coordinates of the target block or the row coordinates of the target element within the block. Represents the column coordinates of the target block or the column coordinates of the target element within the block. and Represents the row and column coordinate indices of the original block or element. Indicates step size, This indicates the width of the convolution kernel.

8. The address mapping method for GPU convolution acceleration according to claim 1, characterized in that, The execution process of the additional workspace mapping strategy includes: Based on the ratio of CUDA thread block size, output matrix width, and stride, the arrangement of the right column elements in the extra workspace is determined to be one of column-first, row-first, or N-shaped arrangement. Column-first arrangement means that all right column elements are arranged in column-first order, so that elements in the same right column are consecutive in the storage space. Row-first arrangement means that elements in the same row of all right columns are consecutive in the storage space. N-shaped arrangement means that elements in the same right column are consecutive in the vertical direction, right columns between different blocks are consecutive in the horizontal direction, and all right column elements are expanded into the extra workspace in an N-shaped order.

9. The address mapping method for GPU convolution acceleration according to claim 8, characterized in that, The method of determining the arrangement of the right column elements in the extra workspace based on the ratio of CUDA thread block size, output matrix width, and stride includes: pass Calculate the ratio, where, For CUDA thread block size, To output the matrix width, Step size; If the ratio is less than or equal to the first preset value, then the sorting method is selected as column priority sorting; If the ratio is greater than the first preset value and less than or equal to the second preset value, then the arrangement method is selected as the N-shaped arrangement, wherein the first preset value is less than the second preset value; If the ratio is greater than the second preset value, then the sorting method is selected as row priority sorting.

10. An address mapping system for GPU convolution acceleration, characterized in that, The system is used to execute the address mapping method for GPU convolution acceleration as described in any one of claims 1 to 9.

Images