Asynchronous prefetching system based on 3D stacking

Abstract

The invention discloses an asynchronous prefetching system based on 3D stacking. The asynchronous prefetching system comprises an on-chip structure and at least one layer of off-chip structure interconnected with the on-chip structure. The on-chip structure comprises at least one computing core, at least one asynchronous prefetching module corresponding to the computing core and an on-chip prefetching queue, the off-chip structure comprises at least one off-chip prefetching queue, and the depth of the off-chip prefetching queue is greater than that of the on-chip prefetching queue. The on-chip prefetching queue and the off-chip prefetching queue are connected in series to form an asynchronous prefetching assembly line; the asynchronous prefetching module is used for prefetching data from a global memory to the tail of an off-chip prefetching queue based on an asynchronous prefetching assembly line, transmitting the data from the head of the off-chip prefetching queue to the tail of an on-chip prefetching queue, and providing required data for a computing core of the on-chip structure. The overhead of the computing core during data access processing is reduced, and the memory access delay is reduced.

Description

Technical Field

[0001] The present application relates to the field of computer technology, and in particular to an asynchronous pre-fetching system based on 3D stacking. Background Art

[0002] In application scenarios such as deep learning and scientific computing, the amount of data that the graphics processing unit (GPU) needs to process is often very large. The parallel computing capability of the GPU enables it to process multiple blocks of data at the same time. This parallelism requires a sufficient amount of data to fill the GPU's computing units to avoid wasting computing resources. Therefore, a large amount of data needs to be quickly pre-fetched into the GPU for calculation.

[0003] Different types of memory in the GPU, such as global memory, shared memory, SRAM, etc., have different access speeds. Although global memory has a large capacity, it has a high access latency; while SRAM has a fast access speed, but its capacity is limited. At present, in order to reduce the access latency of global memory, GPUs usually use SRAM as a cache to store high-frequency access data. However, due to the limited capacity of SRAM, it is impossible to cache all data, resulting in the need to frequently read data from global memory, which increases access latency. Therefore, the traditional GPU architecture is limited by the capacity of on-chip SRAM and is difficult to support the needs of large-scale data prefetching. Summary of the invention

[0004] Based on the above problems, the present application provides an asynchronous pre-fetching system based on 3D stacking.

[0005] To solve the above problems, the technical solutions provided in the embodiments of the present application are as follows:

[0006] The first aspect of the present application provides an asynchronous prefetching system based on 3D stacking, comprising:

[0007] An on-chip structure and at least one layer of off-chip structure interconnected with the on-chip structure through 3D stacking; the on-chip structure includes at least one computing core, at least one asynchronous prefetch module corresponding to the computing core, and an on-chip prefetch queue; the off-chip structure includes at least one off-chip prefetch queue, the depth of the off-chip prefetch queue is greater than the depth of the on-chip prefetch queue, the on-chip prefetch queue and the off-chip prefetch queue are connected in series to form an asynchronous prefetch pipeline, and the asynchronous prefetch pipeline is used to instruct the off-chip structure farthest from the on-chip structure to sequentially transfer data in the global memory to the on-chip structure in descending order of distance from the on-chip structure;

[0008] The asynchronous prefetch module is used to prefetch data from the global memory to the end of the off-chip prefetch queue based on the asynchronous prefetch pipeline, and transfer the data from the head of the off-chip prefetch queue to the end of the on-chip prefetch queue to provide the required data for the computing core of the on-chip structure.

[0009] In one possible implementation, when the number of layers of the off-chip structure is 1, the asynchronous prefetch module is used to prefetch data from the global memory to the end of the off-chip prefetch queue when the depth of the off-chip prefetch queue is less than a preset threshold, or the queue length of the prefetch task is not less than a preset length and there is at least one free position in the off-chip prefetch queue.

[0010] In one possible implementation, when the number of layers of the off-chip structure is 1, the asynchronous prefetch module is used to transfer data from the head of the off-chip prefetch queue to the end of the on-chip prefetch queue when there is vacant space in the on-chip prefetch queue and there is data in the off-chip prefetch queue.

[0011] In one possible implementation, when the asynchronous prefetch module is used to execute a prefetch task, it performs address translation on each prefetch task in the prefetch task list to obtain a physical address corresponding to each prefetch task; when the physical addresses of multiple prefetch tasks in the prefetch task list are continuous, the multiple prefetch tasks are aggregated into an overall task so that the aggregated overall task is executed when the prefetch condition is met.

[0012] In one possible implementation, the on-chip structure includes multiple computing cores, multiple asynchronous prefetch modules corresponding to each computing core, an on-chip prefetch queue and an arbiter. The on-chip prefetch queue and the off-chip prefetch queue are connected in series to form multiple asynchronous prefetch pipelines. The arbiter is used to set the scheduling priority of each asynchronous prefetch pipeline based on the amount of data scheduled for each asynchronous prefetch pipeline.

[0013] In a possible implementation, when the number of layers of the off-chip structure is m, the off-chip structures are interconnected, and m is a positive integer greater than 1;

[0014] The asynchronous prefetch module is used to prefetch data from the global memory to the end of the off-chip prefetch queue of the m-th layer off-chip structure when the depth of the prefetch queue is less than a preset threshold, or the queue length of the prefetch task is not less than a preset length and there is at least one item of free position in the off-chip prefetch queue of the m-th layer off-chip structure, wherein the m-th layer off-chip structure is the off-chip structure farthest from the computing core among the multi-layer off-chip structures;

[0015] When the depth of the off-chip prefetch queue of a layer of off-chip structures among the m off-chip structures is less than a preset threshold, data is transferred from the head of the off-chip prefetch queue of the next layer of off-chip structures farther from the computing core to the tail of the off-chip structure;

[0016] When there is a vacant position in the on-chip prefetch queue and there is data in the off-chip prefetch queue of the first layer off-chip structure, the data is transferred from the head of the off-chip prefetch queue of the first layer off-chip structure to the end of the on-chip prefetch queue.

[0017] In one possible implementation, when the depth of the off-chip prefetch queue of multiple layers of off-chip structures between m off-chip structures is less than a preset threshold, data is preferentially transferred from the head of the off-chip prefetch queue of the off-chip structure closer to the computing core in the multiple layers of off-chip structures to the tail of the off-chip prefetch queue of the previous layer of off-chip structure.

[0018] In a possible implementation, from the off-chip structure farthest from the on-chip structure to the on-chip structure, the bandwidth of each layer of storage structure increases and the capacity of each layer of storage structure decreases.

[0019] In a possible implementation, the physical medium of the off-chip structure is a static random access memory SRAM or a dynamic random access memory DRAM, and the physical medium of the on-chip structure is SRAM.

[0020] In a possible implementation, the on-chip structure and the off-chip structure, as well as each layer of the off-chip structure, are interconnected via a 3D packaging interface.

[0021] Compared with the prior art, this application has the following beneficial effects:

[0022] By dividing the storage structure into on-chip and off-chip structures, and using the asynchronous prefetch pipeline mechanism, the system can prefetch data from the global memory to the off-chip structure in advance, and then quickly transfer the data to the on-chip structure when needed. This prefetch and transfer mechanism reduces the time the computing core waits for data, thereby improving the efficiency of data access and the performance of the overall system. The traditional GPU architecture is limited by the capacity of on-chip SRAM and is difficult to support the demand for large-scale data prefetching. The 3D stacking technology not only greatly increases the storage capacity by stacking on-chip and off-chip structures, but also maintains the high-speed access characteristics of SRAM. This design enables the GPU to have a large-capacity cache and fast memory access capabilities at the same time, and the asynchronous prefetch module can automatically perform data prefetch and transfer operations based on the asynchronous prefetch pipeline without the direct intervention of the computing core. Through the asynchronous prefetch mechanism, data can be prepared before the computing core needs it, reducing waiting time and effectively reducing system latency. As a result, the computing core's overhead in processing data access is reduced, memory access latency is reduced, and the computing core can be more focused on performing computing tasks. BRIEF DESCRIPTION OF THE DRAWINGS

[0023] In order to more clearly illustrate the technical solutions in this embodiment or the prior art, the drawings required for use in the embodiments or the description of the prior art will be briefly introduced below. Obviously, the drawings described below are only some embodiments of the present application. For ordinary technicians in this field, other drawings can be obtained based on these drawings without paying creative work.

[0024] Figure 1 The memory hierarchical architecture of the AI ​​accelerator provided for this application;

[0025] Figure 2A NVIDIA Amper architecture diagram provided for this application;

[0026] Figure 2B A schematic diagram of the Hopper architecture provided for this application;

[0027] Figure 3 A schematic diagram of the structure of a first asynchronous pre-fetching system based on 3D stacking provided in an embodiment of the present application;

[0028] Figure 4 An asynchronous prefetch system structure based on 3D stacking from a hardware perspective provided in an embodiment of the present application;

[0029] Figure 5 A schematic diagram of the structure of a second asynchronous prefetching system based on 3D stacking provided in an embodiment of the present application;

[0030] Figure 6 A schematic diagram of the structure of a third 3D stacking-based asynchronous prefetching system provided in an embodiment of the present application;

[0031] Figure 7 A schematic structural diagram of a fourth 3D stacking-based asynchronous prefetching system provided in an embodiment of the present application. DETAILED DESCRIPTION

[0032] In order to enable those skilled in the art to better understand the solution of the present application, the technical solution in the embodiments of the present application will be clearly and completely described below in conjunction with the drawings in the embodiments of the present application. Obviously, the described embodiments are only part of the embodiments of the present application, not all of the embodiments. Based on the embodiments in the present application, all other embodiments obtained by ordinary technicians in this field without creative work are within the scope of protection of this application.

[0033] To facilitate understanding of the technical solutions provided by the embodiments of the present application, the professional terms involved in the embodiments of the present application will be explained below.

[0034] GPGPU (General-purpose computing on graphics processing units): A general-purpose graphics processing unit, a general-purpose computing chip for artificial intelligence and high-performance computing. Representative products include NVIDIA's A100.

[0035] DSA (Domain-Specific Architectures): Domain-specific architecture, in this article it mainly refers to specific chip architecture for artificial intelligence scenarios.

[0036] GPU (Graphics Processing Unit): Graphics processing unit is a microprocessor that performs image and graphics related operations on personal computers, workstations, game consoles and some mobile devices (such as tablets, smartphones, etc.). GPU reduces the graphics card's dependence on the CPU and performs some of the processing work originally handled by the CPU, especially accelerating 3D graphics processing, thereby greatly improving the computer's graphics processing capabilities.

[0037] SRAM (Static Random-Access Memory): Static random access memory. Its "static" characteristic means that as long as the power remains on, the data stored in it can be kept constant, without the need for periodic refresh like DRAM.

[0038] DRAM (Dynamic Random Access Memory): Dynamic random access memory, whose data can only be kept for a short time and needs to be refreshed periodically to maintain data preservation.

[0039] SMEM (Shared MEMory): Shared memory. Refers to large-capacity memory that can be accessed by different central processing units (CPUs) in a multi-processor computer system.

[0040] LLC (Last Level Cache): last level cache.

[0041] TMA (TensorMemoryAccelerator): A tensor memory accelerator is a hardware unit specifically designed to accelerate the transfer of tensors (multidimensional arrays) between memories.

[0042] DMA (Direct Memory Access): Direct memory access is a technology that allows certain hardware subsystems (such as disk drives, network interface cards, etc.) to directly access main memory without CPU intervention.

[0043] FIFO (First-In-First-Out): First-in, first-out. A first-in, first-out data cache structure where data is stored and accessed in the order in which it enters.

[0044] TSV, Hybrid Bonding: some advanced technologies of 3D packaging. TSV stands for Through-Silicon Via, which can be translated as through-silicon via technology. It is a vertical interconnection technology realized by through-silicon via technology, which allows high-density and high-speed vertical interconnection within the chip or between chips. Hybrid Bonding can be translated as hybrid bonding technology. It is an advanced packaging technology that combines multiple interconnection technologies (such as copper-copper bonding, oxide-oxide bonding, etc.) to achieve high-density and high-reliability interconnection between chips. Hybrid Bonding technology can significantly improve the performance and reliability of 3D packaging.

[0045] To facilitate understanding of the technical solutions provided by the embodiments of the present application, the background technology involved in the embodiments of the present application will be described below.

[0046] As mentioned above, AI accelerators essentially move data to the computing core, which then moves the data out after computing. As the computing core capabilities of GPUs increase, the overhead of data movement (i.e., IO) becomes more and more significant. Figure 1 The memory hierarchical architecture of the AI ​​accelerator provided for this application includes SRAM, HBM and DRAM, and their corresponding speeds are 19TB / s, 1.5TB / s and 12.8GB / s. It can be seen that the closer to the computing core, the greater the bandwidth of data access and the smaller the latency, but at the same time the capacity is also smaller. Therefore, a multi-level cache structure is generally adopted to use the locality of data to reduce the average data transfer time. However, cache misses are still inevitable, resulting in significant IO memory access latency.

[0047] Although GPGPU can prevent the computing core from being blocked by IO by switching thread warps, data is still transported and calculated synchronously, increasing the overall latency. Moreover, the cost of zero-overhead thread warp switching is a large amount of register resource overhead.

[0048] Asynchronous Copy can achieve asynchronous pre-fetching of data in advance. Ideally, the computing core can always obtain data from shared memory for calculation without switching thread bundles, thus hiding the memory access latency to the greatest extent. The earliest asynchronous copy feature in the prior art is cudaMemcpyAsync, which supports the transfer of data from CPU memory to GPU global memory and the overlap of CUDA kernel execution operations. Since then, CudaDMA supports asynchronous copy from GPU global memory to shared memory, but this asynchronous copy is guided by binding a dedicated thread bundle.

[0049] like Figure 2A As shown, Figure 2A The schematic diagram of the NVIDIA Amper architecture provided for this application shows that the NVIDIA Amper architecture can use cuda::memcpy_async to asynchronously copy data between GPU global memory and shared memory without binding dedicated thread bundles. Specifically, before the emergence of cuda::memcpy_async, the data flow moved to shared memory was GMEM→L2 Cache→L1 Cache→thread registers→shared memory. After the emergence of cuda::memcpy_async, it was simplified to: GMEM→L2 Cache→L1 Cache→shared memory. Obviously, the efficiency of the latter has been significantly improved.

[0050] On this basis, the Hopper architecture provides a new dedicated hardware TMA for asynchronous prefetching, which is specifically used to handle asynchronous transfers between globalmemory and sharedmemory, such as Figure 2B As shown, Figure 2B A schematic diagram of the Hopper architecture provided for this application. The core advantage of TMA is that it does not impose any burden on threads when performing asynchronous copies. In particular, a single thread creates a descriptor for global memory data before starting TMA, which is then processed by the TMA hardware, while the A100 thread needs to calculate the tedious global memory address, occupying valuable register resources.

[0051] First, the latency of loading data from L2 cache, GPU local memory, host CPU memory, and even remote memory to on-chip SMEM increases exponentially. To avoid warp switching, the solution is to use a sufficiently deep asynchronous prefetch pipeline, that is, to prefetch as much data as possible to on-chip SMEM in advance, which also means that on-chip SMEM needs to retain multi-stage pipeline data buffers. At the same time, as the computing cores (such as Figure 2B As the computing power of Tensor Core in the MCU increases, the amount of data that needs to be pre-fetched at each pipeline level also increases, further exacerbating the shortage of on-chip SMEM resources. Therefore, the biggest disadvantage of the existing technology, such as H100 TMA, is that the SMEM capacity is insufficient and cannot support the pre-fetching of large quantities of data.

[0052] In order to solve this problem, an asynchronous prefetch system based on 3D stacking is provided in the embodiment of the present application. By connecting the on-chip prefetch queue and the off-chip prefetch queue in series, an efficient asynchronous prefetch pipeline is formed, which can prefetch data from the global memory to the on-chip structure in advance, reducing the time the computing core waits for data. The depth of the off-chip prefetch queue is greater than the depth of the on-chip prefetch queue, so that the system can prefetch more data, effectively hiding the access delay of DRAM and improving the overall performance.

[0053] By prefetching data from global memory to off-chip structures and then to on-chip structures, the locality of data access can be fully utilized to reduce unnecessary data transfer. This structure helps keep frequently used data closer to the computing core, reduces frequent access to global memory, and reduces the overall latency of the system.

[0054] The asynchronous prefetch module can prefetch most of the data required by the computing core into the on-chip structure, so that the computing core can continuously and efficiently perform computing tasks. The data can be prepared before the computing core needs it, reducing waiting time and effectively reducing system latency. As a result, the utilization rate of the computing core is maximized, thereby improving the throughput of the entire system.

[0055] It should be noted that the asynchronous prefetch system based on 3D stacking provided by the present application can be applied to the field of computer technology. The above is only an example and does not limit the application field of the asynchronous prefetch system based on 3D stacking provided by the present application. In addition, the embodiments of the present application may not limit the execution subject of the asynchronous prefetch system based on 3D stacking. For example, the asynchronous prefetch system based on 3D stacking of the embodiments of the present application can be applied to data processing devices such as terminal devices or servers. Among them, the terminal device can be an electronic device such as a computer, a personal digital assistant (PDA). The server can be an independent server, a cloud server, or a cluster server composed of multiple servers.

[0056] In order to make the purpose, technical solution and advantages of the embodiments of the present application clearer, the technical solution in the embodiments of the present application will be clearly and completely described below in conjunction with the drawings in the embodiments of the present application. Obviously, the described embodiments are only part of the embodiments of the present application, not all of the embodiments. Based on the embodiments in the present application, all other embodiments obtained by ordinary technicians in this field without creative work are within the scope of protection of this application.

[0057] Embodiment 1

[0058] The following describes, through Example 1, the system architecture, basic principles, functions and hardware forms of an asynchronous prefetch system based on 3D stacking provided by the present application.

[0059] See also Figure 3 ,Should Figure 3 A schematic diagram of the structure of the first asynchronous prefetch system based on 3D stacking provided in an embodiment of the present application. The asynchronous prefetch system structure provided in an embodiment of the present application is composed of a prefetch queue of an on-chip SMEM, several prefetch queues of an off-chip SMEM, and a hardware asynchronous prefetcher. Among them, the prefetch queue of the on-chip SMEM usually has a faster access speed, but limited capacity. Due to the capacity limitation of the on-chip SMEM, a large amount of prefetched data needs to be stored in an off-chip SMEM. These off-chip SMEMs may use 3D stacking technology to increase capacity and be able to retain enough prefetched data for subsequent use.

[0060] Among them, the hardware asynchronous prefetcher is responsible for performing specific prefetch operations. Unlike traditional software prefetching methods, the hardware asynchronous prefetcher can run independently of the computing core, thereby reducing the occupation of computing resources. It is responsible for prefetching data from global memory (such as GPU DRAM) to off-chip SMEM, and transferring data from off-chip SMEM to on-chip SMEM when appropriate to ensure that the computing core can obtain the required data in a timely manner.

[0061] In addition, an embodiment of the present application provides at least two levels of pipelines on the chip to store the data required by the computing core in the current and next pipeline time periods, respectively saving the data required by the computing core in the current and next pipeline time periods. The remaining levels of the pipeline are located in the 3D stacked off-chip SMEM, which has a large capacity and can retain enough pre-fetched data. In this way, the data that will be needed can be prepared in advance without blocking the computing core. The pre-fetch queue of the on-chip SMEM and the pre-fetch queue of one (or more) off-chip SMEM are connected in series to form a complete asynchronous pre-fetch pipeline. This design allows data to continuously pass through the off-chip SMEM from the global memory (such as DRAM) and finally reach the on-chip SMEM for use by the computing core.

[0062] In the 3D stack-based asynchronous prefetching system provided in the embodiment of the present application, each queue can be organized in the form of FIFO (first in, first out), thereby ensuring that data can be processed and transmitted in order. The circular queue is an efficient implementation method, which uses a continuous physical memory space to simulate the head and tail of the queue through the circular movement of the pointer, thereby avoiding frequent movement and copying of data.

[0063] In the asynchronous prefetch system setting process provided in the embodiment of the present application, the depth of the on-chip prefetch queue can be set shallower, because the capacity of the on-chip SMEM is limited, and the computing core has high real-time requirements for data. Ideally, the depth of two data blocks can ensure that when the computing core needs the next data block, the data block is ready and located in the on-chip SMEM. The premise of this design is that the calculation time (including the time for data to be loaded / stored from the on-chip SMEM to the register) is greater than the time for data to be transferred from the off-chip SMEM to the on-chip SMEM. Under 3D stacking technology, the physical distance between the off-chip SMEM and the on-chip SMEM is greatly shortened, and the transmission bandwidth is improved.

[0064] In a possible implementation, the asynchronous prefetcher uniformly manages the on-chip and off-chip SMEM resources required by the pipeline, and can configure a prefetch task queue and a queue status table for each asynchronous prefetch pipeline. The prefetch task queue stores the future prefetch data information issued by the software at runtime, and the queue status table records the status information of all prefetch queues on the pipeline (such as queue number, base address, length, head and tail pointer positions, etc.). This design enables the hardware asynchronous prefetcher to efficiently manage and schedule prefetch tasks to ensure the correctness and timeliness of the data. Therefore, in the actual application of the asynchronous prefetch system, the hardware asynchronous prefetcher cascades data from the off-chip SMEM to the on-chip SMEM according to the status of each FIFO queue, and updates the context information in real time. This dynamic data handling and queue status table management mechanism enables the system to flexibly respond to different computing requirements and data access modes.

[0065] It should be noted that the above-mentioned queue status table can also be expressed as the term in RDMA programming: context. In each embodiment of the present application, the queue status table and the context have the same meaning and protection scope, and do not limit the protection scope of the embodiments of the present application.

[0066] Compared with the above description of the asynchronous prefetch system structure from a logical perspective, the following further describes the embodiment of the present application from a hardware perspective. Figure 4 , Figure 4 The asynchronous prefetch system structure based on 3D stacking from the hardware perspective provided in the embodiment of the present application includes an on-chip SMEM at the bottom layer including computing core 1 and computing core 2. It should be noted that the meanings and protection scopes expressed by the stream multiprocessor, stream processor and computing core in each embodiment of the present application are the same, and do not limit the protection scope of the embodiment of the present application.

[0067] like Figure 4 As shown, each on-chip SMEM has a private off-chip SMEM, which can be organized in a single layer or multiple layers, namely L2 SMEM and L3 SMEM in the figure. Among them, stream processor 1 and stream processor 2 are responsible for processing data flow and computing tasks respectively. The stream processor is the core processing unit in the GPU, which is used to perform parallel computing tasks. Each stream multiprocessor is connected to an L2SMEM block 1 and an L2 SMEM block 2. In addition, according to the problem description, each on-chip SMEM (which can be assumed to be part of or closely related to the stream processor) has a private off-chip SMEM, namely the L2SMEM block here.

[0068] Among them, each SMEM can be interconnected through a high-speed 3D packaging interface, such as TSV and Hybrid Bonding. In addition, in order to improve flexibility, the method of fusing L1 cache with on-chip SMEM in GPGPU can be adopted to fuse the off-chip SMEM (L3SMEM block) with the off-chip LLC (Last Level Cache). LLC, as the last level cache in the GPU memory hierarchy, is responsible for caching data from slower external memory (such as GDDR memory). By integrating the off-chip SMEM and LLC more tightly or optimizing the data flow between them, the number of accesses to the external memory can be reduced, further reducing data access latency.

[0069] In summary, the structure provided by this application realizes efficient data prefetching and transmission through the two-stage pipeline design of on-chip SMEM, the large-capacity 3D stacked storage of off-chip SMEM, and the intelligent scheduling and management of the hardware asynchronous prefetcher. This design not only improves the data access speed, but also reduces the waiting time of the computing core, thereby significantly improving the overall performance and efficiency of the system.

[0070] Embodiment 2

[0071] The structure provided by the embodiment of the present application is described below through the second embodiment. The asynchronous prefetch system structure provided by the present embodiment includes a layer of on-chip structure and a layer of off-chip structure.

[0072] See also Figure 5 , Figure 5 The second 3D stacking-based asynchronous prefetching system structure diagram provided in the embodiment of the present application adopts a 2-level SMEM packaged in 2-layer 3D stacking, including two main levels: the first level is the logic die, and the second level is the memory die.

[0073] The logic die is used to process data and convert it into useful information. It is mainly responsible for processing data and performing logical operations. In the CPU, the logic die is the most important component of the CPU, also known as the "kernel" or "core". It contains all the computing units, control units, caches and other key components of the CPU, and is responsible for executing instructions and data processing tasks in computer programs. Therefore, the Chinese name of the logic die can be translated as "logic chip" or "kernel chip". In the embodiment of the present application, it includes a computing core and an asynchronous computing module. In the first layer of logic die, each computing core is equipped with its own private on-chip SMEM, referred to as L1 SMEM. The second layer of memory die consists of distributed off-chip SMEM blocks, referred to as L2SMEM, and each L1 SMEM is directly connected to its corresponding L2 SMEM by a physical channel (such as TSV and Hybrid Bounding).

[0074] The memory die, which may also be called the storage die in the embodiment of the present application, is mainly responsible for the data storage function. It may be a tiny memory chip or flash memory chip in a storage device such as a memory stick or a solid-state drive. These chips contain actual circuits for storing data and are one of the most basic storage units in a storage device. Therefore, the Chinese name of the storage die or memory die can be translated as "storage chip" or "storage grain". It is composed of distributed off-chip SMEM blocks (L2 SMEM), which can be based on SRAM or DRAM. L2 SMEM is used to store larger capacity data to support the needs of the computing core. The data queue created in L2SMEM is used to temporarily store data from the global memory DMA (direct memory access).

[0075] The asynchronous computing module supports asynchronous prefetch and data handling operations to improve the efficiency of data processing. Each computing core in the logic die is equipped with an asynchronous prefetcher to manage the prefetch task queue and context, and maintain a prefetch task queue and context for each software-defined asynchronous prefetch pipeline. Each element in the queue contains specific prefetch task information, such as physical address, length, and L2 data queue number. When the software defines the asynchronous prefetch pipeline, data queues are created in L1 and L2SMEM, respectively, referred to as L1 and L2 data queues. The status information of the two queues is saved in the context of the pipeline.

[0076] When the software is running, the asynchronous prefetcher can first translate the address of the received prefetch task to obtain the corresponding physical address. Before adding the prefetch task to the prefetch task queue, the asynchronous prefetcher can try to aggregate it with the tail task of the queue, provided that the two can be merged into a continuous physical address.

[0077] The following describes the data transfer process of the asynchronous prefetch system with a two-layer structure. In actual application, the asynchronous prefetcher can perform two types of data transfer operations, namely, transferring data from the global memory DMA to the tail of the L2 SMEM queue and from the head of the L2 SMEM to the free block of the L1 SMEM.

[0078] For the first type of data handling in this embodiment, the asynchronous prefetcher can perform prefetch operations on the global memory when one of the following two conditions is met. Condition one is that the length of the prefetch task queue is greater than or equal to 2 and the L2 data queue is not full, that is, when the current prefetch task can no longer be aggregated with the previous task, the previous data prefetch task can be executed directly. Condition two is that the queue depth of L2 SMEM is less than a preset threshold, and the preset threshold is determined based on the latency from L2 SMEM to L1 SMEM and the latency of global memory DMA. For example, the preset threshold can ensure that the L2 data queue is not emptied during the time period of obtaining data from the global memory. In addition, if the data length of the prefetch task to be executed will cause the L2 data queue to overflow, the task needs to be split.

[0079] For the second type of data transfer in this embodiment, as long as a certain asynchronous prefetch pipeline satisfies that there is remaining space in the L1 data queue and the L2 data queue is not empty, the data transfer can be executed.

[0080] Embodiment 3

[0081] The following is an explanation of an asynchronous prefetch system based on 3D stacking provided by the present application through Example 3. Compared with the asynchronous prefetch system structure provided in the aforementioned Example 2, which includes a layer of on-chip structure and a layer of off-chip structure, the asynchronous prefetch system structure provided in Example 3 of the present application includes a layer of on-chip structure and multiple layers of off-chip structure.

[0082] See also Figure 6 ,Should Figure 6 A schematic structural diagram of a third 3D stacking-based asynchronous prefetching system provided in an embodiment of the present application, which adopts a multi-level SMEM structure (face-to-back packaging) of multi-layer 3D stacking packaging.

[0083] The first layer of logic die is the same as the previous embodiment. The remaining dies implement multi-level SMEM, which can use SRAM or DRAM. Different from the previous embodiment, this embodiment logically divides the multi-layer die into data queues organized in FIFO, and each data queue can span multiple physical dies. Exemplarily, a typical configuration is a 3D package with a layer of logic die, a layer of SRAM die, and multiple layers of DRAM die, corresponding to three levels of data queues.

[0084] The logic die is the core of the system, which contains an asynchronous computing core that is connected to the streaming multiprocessor 1 through the L1 queue. In addition, the logic die also contains a prefetch module, which is responsible for quickly prefetching data from the storage die to the logic die for computing processing.

[0085] The following is an explanation of the data transport process of the asynchronous prefetch system with a multi-layer structure. In actual application, the asynchronous prefetcher contained in the logic die also needs to support the transport of class data, that is, from the global memory to the last level data queue, from L2 to L1 data queue, and between other adjacent data queues. The implementation logic of the first two categories is basically the same as that of the first embodiment, except that the condition 2 of the first category must be based on the last level queue L M .L M The queue depth of SMEM is less than the preset threshold, which is based on L M The latency from SMEM to L1 SMEM is determined by the latency of global memory DMA. For example, the preset threshold can ensure that L1 SMEM is within the time period of obtaining data from global memory. M The data queue is not emptied. In addition, if the data length of the prefetch task to be executed will cause L M When the data queue overflows, the task needs to be split.

[0086] For the transfer between intermediate data queues, a strategy similar to the first type of condition 2 can be adopted to prevent the sudden occurrence of a large number of DMAs. That is, only when the depth of a certain level of intermediate data queue is less than the preset threshold, the data is moved from the head of the next level queue to the tail of the queue.

[0087] In one possible implementation, in order to avoid idle computing cores as much as possible, when multiple levels of prefetch queues meet the transfer conditions at the same time, the asynchronous prefetcher traverses according to the strategy from front to back, and supports cross-level transfer of data to improve efficiency. Here, "front" refers to the storage level closer to the computing core), and prioritizes those queues that are closer to the computing core or have higher priorities. That is, data can be directly transferred from a farther storage level to a storage level near the computing core as needed, thereby further improving the efficiency of data access. This cross-level transfer capability increases the flexibility and response speed of data transfer.

[0088] The following is an explanation of the detailed technical features that can be adaptively set in Example 2 and Example 3. It should be noted that the following technical features can be selectively applied to the technical scenarios of Example 2 and / or Example 3 according to actual needs. Figure 7 , Figure 7 A schematic structural diagram of a fourth 3D stacking-based asynchronous prefetching system provided in an embodiment of the present application.

[0089] In one possible implementation, the total depth of the pipeline and the number of stages allocated to each SMEM on-chip and off-chip are software-definable. This means that when the pipeline is initialized, the structure of the pipeline can be customized according to the specific application scenario and performance requirements. For example, for compute-intensive applications, the number of off-chip SMEM stages can be increased to expand the capacity of the prefetch queue.

[0090] In one possible implementation, when the memory access address can be accurately predicted and the physical address is continuous, the pre-fetching of multiple consecutive blocks of data can reduce latency through memory access aggregation. Memory access aggregation combines multiple small memory access requests into a large memory access request, thereby reducing the number of memory accesses and overhead. The better the locality of the data in the memory, the more significant the effect of memory access aggregation. Through a reasonable pre-fetch strategy and memory layout design, the locality of the data can be fully utilized to hide the DRAM memory access latency, thereby improving the overall performance of the system.

[0091] The traditional prefetch mechanism performs DRAM prefetching in each pipeline time slot, which may cause unnecessary DRAM access and power consumption. However, the embodiment of the present application caches prefetch instructions and attempts to aggregate multiple prefetch tasks with consecutive physical addresses into one, and then executes it when the prefetch conditions are met, which significantly reduces the number of DRAM accesses and improves memory access efficiency.

[0092] In one possible implementation, multiple pipelines managed by the asynchronous prefetcher may meet the data transfer conditions at the same time. Therefore, the structure provided in the embodiment of the present application may also include an arbitrator for allocating prefetch resources fairly and efficiently according to a predetermined scheduling strategy.

[0093] Exemplarily, the pipelines can be sorted according to the amount of data scheduled, and pre-fetching tasks can be arranged preferentially for pipelines with less data. This helps to balance the processing speed of each pipeline and prevent some pipelines from being idle due to data starvation. The number of tasks scheduled for each pipeline can also be used, and pipelines with fewer tasks can obtain higher scheduling priorities. This is applicable to situations where the amount of tasks is unbalanced and helps to achieve load balancing. For example, the amount of data or the number of tasks that have been scheduled for the pipeline is recorded in the context, and the arbitrator schedules them in ascending order according to the amount of data or the number of tasks. The embodiments of the present application also support setting different scheduling priorities for pipelines to meet different quality of service (QoS) requirements.

[0094] In actual application scenarios, in addition to fixed scheduling strategies, the scheduling priority of each pipeline can be dynamically adjusted according to the current system status (such as CPU utilization, memory occupancy, etc.) or the quality of service (QoS) requirements set by the user. This increases the flexibility and adaptability of the system.

[0095] Compared with the prior art, the asynchronous prefetching system provided by the embodiment of the present application has the following beneficial effects:

[0096] The multi-level SMEM structure implemented by the 3D stacking technology in the embodiment of the present application greatly expands the on-chip storage capacity of the GPU, allowing the GPU to store more data pre-fetching redundancy, thereby more effectively hiding the memory access time. This design overcomes the limitations of limited on-chip SRAM resources in traditional GPU architectures. Unlike extended memory that requires explicit software control, the off-chip SMEM data queue in the embodiment of the present application is accessed and controlled by a hardware asynchronous prefetcher, which greatly improves the efficiency of data pre-fetching and management and reduces software-level intervention and overhead.

[0097] The traditional GPU architecture is limited by the capacity of on-chip SRAM and cannot support the needs of large-scale data pre-fetching. The 3D stacking technology not only greatly increases the storage capacity by vertically stacking SRAM layers, but also maintains the high-speed access characteristics of SRAM. This design enables the GPU to have both large-capacity cache and fast memory access capabilities, which is crucial for processing large-scale data sets and complex computing tasks. The large-capacity and high-speed SRAM cache reduces dependence on DRAM and reduces memory access latency, thereby improving overall computing efficiency. Because AI models usually involve frequent access and calculation of large amounts of data, the computing efficiency advantage is particularly obvious in AI applications.

[0098] In the embodiment of the present application, the execution condition of the pre-fetch task is based on the depth of the last level pre-fetch queue and the aggregated physical address. When the queue depth is below the waterline or the aggregated physical address cannot be used, the pre-fetch task in the waiting state is executed. This mechanism effectively relieves the pressure of the hardware pre-fetcher, avoids unnecessary pre-fetch operations, and further improves the memory access efficiency.

[0099] Unlike technologies such as TMA (Tile-based Memory Access) that passively rely on data locality, the embodiments of the present application utilize the predictable characteristics of AI scene memory access, analyze the current queue status and memory access mode through intelligent algorithms, and actively pre-fetch and move data. This active data management strategy can more effectively reduce the number of DRAM accesses and improve the hit rate of data access.

[0100] Traditional prefetch tasks usually have only one level of prefetch data queues, and the embodiments of the present application support multi-level queues, so that data can be flexibly transferred between queues of different levels. When the depth of a certain level of queue is less than a threshold, the asynchronous prefetcher can try to move data from the subsequent queue to the current queue to maintain the continuity and availability of the data.

[0101] Furthermore, the embodiment of the present application can carry out data transfer by priority from front to back and allow cross-level transfer, which further improves the efficiency and flexibility of data transfer. By optimizing the data transfer management between queues of different levels, the embodiment of the present application can more effectively utilize the storage resources of the GPU and improve the overall performance.

[0102] In summary, the embodiments of the present application improve the memory access performance of the GPU by introducing a multi-level SMEM structure of 3D stacked memory dies, an aggregated prefetch address mechanism, prefetch task execution conditions, and a multi-level prefetch queue management mechanism. As a result, not only the data processing capability and efficiency of the GPU are improved, but also strong support is provided for more complex and larger-scale computing tasks in the future.

[0103] It should be noted that the on-chip structure in the embodiments of the present application may refer to the logic die in the aforementioned embodiments, and there is a structural layer of a stream multiprocessor. The off-chip structure may refer to the storage die, multi-layer storage die, and the layer structure where the L2 SMEM and L3 SMEM are located in the aforementioned embodiments. The asynchronous prefetch module may refer to the asynchronous prefetcher in the aforementioned embodiments. For details about the system structure composed of the on-chip structure and the off-chip structure, please refer to the previous embodiments, which will not be repeated here.

[0104] The embodiments of the present application also provide corresponding devices and computer storage media for implementing the 3D stacking-based asynchronous prefetching system provided in the embodiments of the present application.

[0105] The device includes a memory and a processor, the memory is used to store instructions or codes, and the processor is used to execute the instructions or codes, so that the device executes the asynchronous prefetch system based on 3D stacking described in any embodiment of the present application.

[0106] The computer storage medium stores codes, and when the codes are executed, a device executing the codes implements the asynchronous pre-fetching system based on 3D stacking described in any one of the embodiments of the present application.

[0107] It should be noted that the various embodiments in this specification are described in a progressive manner, and each embodiment focuses on the differences from other embodiments, and the same or similar parts between the various embodiments can be referred to each other. For the system or device disclosed in the embodiment, since it corresponds to the method disclosed in the embodiment, the description is relatively simple, and the relevant parts can be referred to the method part description.

[0108] It should be understood that in the present application, "at least one (item)" means one or more, and "plurality" means two or more. "And / or" is used to describe the association relationship of associated objects, indicating that three relationships may exist. For example, "A and / or B" can mean: only A exists, only B exists, and A and B exist at the same time, where A and B can be singular or plural. The character " / " generally indicates that the objects associated before and after are in an "or" relationship. "At least one of the following" or similar expressions refers to any combination of these items, including any combination of single or plural items. For example, at least one of a, b or c can mean: a, b, c, "a and b", "a and c", "b and c", or "a and b and c", where a, b, c can be single or multiple.

[0109] It should be understood that the terms "center", "longitudinal", "lateral", "up", "down", "front", "back", "left", "right", "vertical", "horizontal", "top", "bottom", "inside", "outside" and the like indicate orientations or positional relationships based on the orientations or positional relationships shown in the accompanying drawings, and are only for the convenience of describing the present invention and simplifying the description, and do not indicate or imply that the referred device or element must have a specific orientation, be constructed and operated in a specific orientation, and therefore should not be understood as a limitation on the present invention.

[0110] It should be noted that, unless otherwise clearly specified and limited, the terms "installed", "connected" and "connected" should be understood in a broad sense, for example, it can be a fixed connection, a detachable connection, or an integral connection; it can be a mechanical connection or an electrical connection; it can be a direct connection or an indirect connection through an intermediate medium, or it can be the internal communication of two components. For ordinary technicians in this field, the specific meanings of the above terms in the present invention can be understood according to specific circumstances.

[0111] It should also be noted that, in this article, relational terms such as first and second, etc. are only used to distinguish one entity or operation from another entity or operation, and do not necessarily require or imply any such actual relationship or order between these entities or operations. Moreover, the terms "include", "comprise" or any other variants thereof are intended to cover non-exclusive inclusion, so that a process, method, article or device including a series of elements includes not only those elements, but also other elements not explicitly listed, or also includes elements inherent to such process, method, article or device. In the absence of further restrictions, the elements defined by the sentence "comprise a ..." do not exclude the presence of other identical elements in the process, method, article or device including the elements.

[0112] The steps of the method or algorithm described in conjunction with the embodiments disclosed herein may be implemented directly using hardware, a software module executed by a processor, or a combination of the two. The software module may be placed in a random access memory (RAM), a memory, a read-only memory (ROM), an electrically programmable ROM, an electrically erasable programmable ROM, a register, a hard disk, a removable disk, a CD-ROM, or any other form of storage medium known in the art.

[0113] The above description of the disclosed embodiments enables those skilled in the art to implement or use the present application. Various modifications to these embodiments will be apparent to those skilled in the art, and the general principles defined herein may be implemented in other embodiments without departing from the spirit or scope of the present application. Therefore, the present application will not be limited to the embodiments shown herein, but will conform to the widest scope consistent with the principles and novel features disclosed herein.

Claims

1. An asynchronous prefetching system based on 3D stacking, characterized in that: include: An on-chip structure and at least one layer of off-chip structure interconnected with the on-chip structure through 3D stacking; the on-chip structure includes at least one computing core, at least one asynchronous prefetch module corresponding to the computing core, and an on-chip prefetch queue; the off-chip structure includes at least one off-chip prefetch queue, the depth of the off-chip prefetch queue is greater than the depth of the on-chip prefetch queue, the on-chip prefetch queue and the off-chip prefetch queue are connected in series to form an asynchronous prefetch pipeline, and the asynchronous prefetch pipeline is used to instruct the off-chip structure farthest from the on-chip structure to sequentially transfer data in the global memory to the on-chip structure in descending order of distance from the on-chip structure; The asynchronous prefetch module is used to prefetch data from the global memory to the end of the off-chip prefetch queue based on the asynchronous prefetch pipeline, and transfer the data from the head of the off-chip prefetch queue to the end of the on-chip prefetch queue to provide the required data for the computing core of the on-chip structure.

2. The system according to claim 1, characterized in that When the number of layers of the off-chip structure is 1, the asynchronous prefetch module is used to prefetch data from the global memory to the end of the off-chip prefetch queue when the depth of the off-chip prefetch queue is less than a preset threshold, or the queue length of the prefetch task is not less than a preset length and there is at least one empty position in the off-chip prefetch queue.

3. The system according to claim 1, characterized in that When the number of layers of the off-chip structure is 1, the asynchronous prefetch module is used to transfer data from the head of the off-chip prefetch queue to the end of the on-chip prefetch queue when there is a vacant position in the on-chip prefetch queue and there is data in the off-chip prefetch queue.

4. The system according to claim 1, characterized in that When the asynchronous prefetch module is used to execute a prefetch task, it performs address translation on each prefetch task in the prefetch task list to obtain a physical address corresponding to each prefetch task; when the physical addresses of multiple prefetch tasks in the prefetch task list are continuous, the multiple prefetch tasks are aggregated into an overall task so that the aggregated overall task is executed when the prefetch condition is met.

5. The system according to claim 1, characterized in that The on-chip structure includes multiple computing cores, multiple asynchronous prefetch modules corresponding to each computing core, an on-chip prefetch queue and an arbiter. The on-chip prefetch queue and the off-chip prefetch queue are connected in series to form multiple asynchronous prefetch pipelines. The arbiter is used to set the scheduling priority of each asynchronous prefetch pipeline based on the amount of data scheduled by each asynchronous prefetch pipeline.

6. The system according to claim 1, characterized in that When the number of layers of the off-chip structure is m, the off-chip structures are interconnected, and m is a positive integer greater than 1; The asynchronous prefetch module is used to prefetch data from the global memory to the end of the off-chip prefetch queue of the m-th layer off-chip structure when the depth of the prefetch queue is less than a preset threshold, or the queue length of the prefetch task is not less than a preset length and there is at least one item of free position in the off-chip prefetch queue of the m-th layer off-chip structure, wherein the m-th layer off-chip structure is the off-chip structure farthest from the computing core among the multi-layer off-chip structures; When the depth of the off-chip prefetch queue of a layer of off-chip structures among the m off-chip structures is less than a preset threshold, data is transferred from the head of the off-chip prefetch queue of the next layer of off-chip structures farther from the computing core to the tail of the off-chip structure; When there is a vacant position in the on-chip prefetch queue and there is data in the off-chip prefetch queue of the first layer off-chip structure, the data is transferred from the head of the off-chip prefetch queue of the first layer off-chip structure to the end of the on-chip prefetch queue.

7. The system according to claim 6, characterized in that When the depth of the off-chip prefetch queue of multiple layers of off-chip structures between m off-chip structures is less than a preset threshold, data is preferentially transferred from the head of the off-chip prefetch queue of the off-chip structure closer to the computing core in the multiple layers of off-chip structures to the tail of the off-chip prefetch queue of the previous layer of off-chip structures.

8. The system according to any one of claims 1 to 7, characterized in that: From the off-chip structure farthest from the on-chip structure to the on-chip structure, the bandwidth of each layer of storage structure increases, and the capacity of each layer of storage structure decreases.

9. The system according to any one of claims 1 to 7, characterized in that: The physical medium of the off-chip structure is a static random access memory SRAM or a dynamic random access memory DRAM, and the physical medium of the on-chip structure is SRAM.

10. The system according to any one of claims 1 to 7, characterized in that: The on-chip structure and the off-chip structure, as well as the off-chip structures of each layer, are interconnected via a 3D packaging interface.

Images