Inference acceleration method and device, electronic equipment and storage medium

Abstract

The invention relates to the technical field of computers, in particular to a reasoning acceleration method and device, electronic equipment and a storage medium, and is used for improving the reasoning speed when a model executes a task. The method comprises the steps that when a target model is adopted to execute a to-be-reasoned task, a matched reasoning mode is activated according to hardware state information of computing resources occupied by the target model; when the reasoning mode is an instruction acceleration mode, at least one round of reasoning operation is executed, and each round of reasoning operation comprises the steps that storage continuity check is conducted in computing resources for reasoning dependency data related to the input sequence of the round; the reasoning dependency data is generated and stored through the reasoning operation of the previous round; when the storage position of the reasoning dependency data is in a non-continuous state in the computing resources, the storage position is adjusted to be in a continuous state by rearranging the reasoning dependency data; and calling a specified hardware acceleration instruction to load the adjusted reasoning dependency data, and determining the reasoning result of the round based on the loading result.

Description

Technical Field

[0001] The present application relates to the field of computer technology, and in particular to an inference acceleration method, device, electronic device, and storage medium. Background Art

[0002] In large-model inference scenarios, variable-length sequence reasoning involves processing input tasks of dynamically changing lengths (such as text generation and time series prediction). The number of tokens in different sequences varies significantly, causing the storage requirements of the Graphics Processing Unit (GPU) to grow nonlinearly with the sequence length, especially when storing reusable intermediate results, such as the key-value (KV) vectors in the attention mechanism. Therefore, how to effectively manage the storage of this data is crucial to supporting efficient variable-length sequence reasoning.

[0003] Related technologies typically use a fixed-length storage pre-allocation strategy, pre-allocating a maximum-length contiguous storage block regardless of the actual sequence length. For shorter sequences, this not only wastes more storage space and leads to low storage utilization, but also fills the remaining storage space with zeros, introducing redundant computations and wasting computing power.

[0004] In summary, in inference scenarios, how to improve the efficiency of hardware resource utilization and achieve efficient inference is an urgent problem to be solved. Summary of the Invention

[0005] Embodiments of the present application provide an inference acceleration method, device, electronic device, and storage medium for increasing the inference speed when a model performs a task.

[0006] An inference acceleration method provided in an embodiment of the present application includes:

[0007] When the target model is used to execute the task to be inferred, based on the hardware status information of the computing resources occupied by the target model, an inference mode that matches the current hardware status of the computing resources is activated; and when the inference mode is the instruction acceleration mode, at least one round of inference operation is performed, each round of inference operation including:

[0008] Performing a storage continuity check in the computing resource for the inference dependency data associated with the input sequence of the current round, wherein the inference dependency data is generated and stored by the inference operation of the previous round;

[0009] When the storage location of the inference-dependent data is in a discontinuous state in the computing resource, adjusting the storage location to a continuous state by rearranging the inference-dependent data;

[0010] The specified hardware acceleration instruction is called to load the adjusted inference dependency data, and after determining the result of the current round of inference based on the loading result, the input sequence of the next round is generated according to the result of the current round of inference, until a preset stop condition is reached.

[0011] An inference acceleration device provided in an embodiment of the present application includes:

[0012] A decision unit, configured to activate, when executing a task to be inferred using a target model, an inference mode that matches the current hardware state of the computing resources occupied by the target model based on the hardware state information of the computing resources;

[0013] An inference unit, configured to, when the inference mode is the instruction acceleration mode, perform at least one round of inference operations, each round of inference operations including:

[0014] Performing a storage continuity check in the computing resource for the inference dependency data associated with the input sequence of the current round, wherein the inference dependency data is generated and stored by the inference operation of the previous round;

[0015] When the storage location of the inference-dependent data is in a discontinuous state in the computing resource, adjusting the storage location to a continuous state by rearranging the inference-dependent data;

[0016] The specified hardware acceleration instruction is called to load the adjusted inference dependency data, and after determining the result of the current round of inference based on the loading result, the input sequence of the next round is generated according to the result of the current round of inference, until a preset stop condition is reached.

[0017] Optionally, after performing a storage continuity check in the computing resource for the inference dependency data related to the input sequence of this round, the inference unit is further configured to:

[0018] When the storage location of the reasoning-dependent data is in a continuous state in the computing resource, the specified hardware acceleration instruction is directly called to load the reasoning-dependent data, and after determining the result of this round of reasoning based on the loading result, the input sequence of the next round is generated according to the result of this round of reasoning until the preset stop condition is reached.

[0019] Optionally, before calling a specified hardware acceleration instruction to load the inference-dependent data and determining the result of the current round of inference based on the loading result, the inference unit is further configured to:

[0020] Allocating a first buffer and a second buffer in a shared memory of the computing resource;

[0021] The reasoning unit is specifically used for:

[0022] Call the hardware acceleration instruction to alternately load different batches of inference-dependent sub-data from the first buffer and the second buffer for calculation and processing, respectively, to obtain the result of this round of inference; wherein the first buffer and the second buffer are used to alternately load different batches of inference-dependent sub-data from the inference-dependent data stored in the global memory of the computing resource.

[0023] Optionally, the reasoning unit is specifically configured to:

[0024] Based on the first buffer and the second buffer, the following steps are executed in a loop until the inference-dependent data is processed:

[0025] Calling the hardware acceleration instruction to perform computation on the inference-dependent sub-data of this batch in the current computation buffer;

[0026] During the computation process, a next batch of inference-dependent sub-data is loaded from the global memory to the current preload buffer via an asynchronous transmission interface; wherein the computation buffer is one of the first buffer and the second buffer, and the preload buffer is the other of the first buffer and the second buffer;

[0027] When it is detected that the calculation of the calculation buffer is completed and the data loading of the preload buffer is completed, the inference dependent sub-data and the calculation results of this batch in the calculation buffer are written back to the global memory; and the preload buffer is switched to a new calculation buffer, and the calculation buffer is switched to a new preload buffer.

[0028] Optionally, the reasoning dependency data includes: the latest reasoning dependency data related to the latest reasoning result in the input sequence of the current round, and the historical reasoning dependency data related to other information in the input sequence of the current round;

[0029] The reasoning unit is specifically used for:

[0030] Checking storage continuity of the latest reasoning dependency data and the historical reasoning dependency data in the computing resource;

[0031] When the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource, the latest reasoning dependency data and the historical reasoning dependency data are adjusted to a continuous storage state by rearranging the latest reasoning dependency data.

[0032] Optionally, the inference unit is further configured to determine the storage locations of the latest inference dependency data and the historical inference dependency data in a discontinuous state in the computing resource by any of the following methods:

[0033] When the storage location of the latest reasoning dependency data is in a discontinuous state in the computing resource, determining that the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource;

[0034] When the storage location of the latest reasoning dependency data and the storage location of the historical reasoning dependency data are not adjacent in the computing resource, it is determined that the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource.

[0035] Optionally, the reasoning unit is specifically configured to:

[0036] For each memory block storing the inference-dependent data, decompose each memory block into a plurality of sub-blocks according to a preset data block size; wherein the size of each sub-block is an integer multiple of the preset data block size; the preset data block size is determined based on the memory access rule of the computing resource;

[0037] The obtained sub-blocks are rearranged according to the target continuous storage order to adjust the storage positions to a continuous state.

[0038] Optionally, the reasoning unit is specifically configured to:

[0039] Determine the effective data length of the loading result;

[0040] According to the effective data length, using the first reasoning precision, performing a main reasoning operation on the loaded result to obtain an intermediate result;

[0041] Using the second reasoning precision, a secondary reasoning operation is performed on the intermediate result to obtain the current round reasoning result; wherein the first reasoning precision is lower than the second reasoning precision.

[0042] Optionally, the decision unit is specifically configured to:

[0043] Determining a mode decision factor corresponding to the computing resource by quantifying the hardware status information;

[0044] If the mode decision factor is greater than a first threshold, activating the instruction acceleration mode;

[0045] If the mode decision factor is less than a second threshold, a paging cache mode is activated; wherein the paging cache mode is used to divide the input sequence of this round into multiple short sequences according to a preset length, and the inference dependency data corresponding to each short sequence is continuously stored using different paging units;

[0046] If the mode decision factor is not less than the second threshold and not greater than the first threshold, the current reasoning mode is maintained.

[0047] Optionally, the hardware status information includes: computing unit utilization, memory access efficiency, and bandwidth utilization; and the decision unit is specifically configured to:

[0048] Determining the utilization rate of the computing unit according to the number of idle cycles of the target computing unit in the computing resource within a preset time period;

[0049] Determining the memory access efficiency according to the number of accesses to the target memory layer in the computing resource within a preset time period;

[0050] determining the bandwidth utilization rate according to an amount of bandwidth used by a target storage unit in the computing resource within a preset time period;

[0051] A weighted sum is performed based on the computing unit utilization, the memory access efficiency, and the bandwidth utilization to obtain the mode decision factor.

[0052] Optionally, before performing at least one round of reasoning operations, the reasoning unit is further configured to:

[0053] For the task to be inferred, obtaining an initial input sequence corresponding to the task to be inferred;

[0054] The length of the initial input sequence is compared with a preset length threshold, and it is determined that the length of the initial input sequence is less than the preset length threshold.

[0055] An electronic device provided by an embodiment of the present application includes a processor and a memory, wherein the memory stores a computer program, and when the computer program is executed by the processor, the processor performs the steps of any one of the above-mentioned inference acceleration methods.

[0056] An embodiment of the present application provides a computer-readable storage medium, which includes a computer program. When the computer program is run on an electronic device, the computer program is used to enable the electronic device to perform the steps of any one of the above-mentioned inference acceleration methods.

[0057] An embodiment of the present application provides a computer program product, which includes a computer program stored in a computer-readable storage medium; when a processor of an electronic device reads the computer program from the computer-readable storage medium, the processor executes the computer program, causing the electronic device to perform the steps of any one of the above-mentioned inference acceleration methods.

[0058] The beneficial effects of this application are as follows:

[0059] Embodiments of the present application provide an inference acceleration method, apparatus, electronic device, and storage medium. By dynamically monitoring the hardware status of computing resources, the inference mode that best matches the current hardware state is selected. This dynamic adjustment mechanism optimizes resource utilization and ensures efficient inference operations under varying hardware conditions.

[0060] In instruction acceleration mode, the input sequence-related reasoning-dependent data usually includes intermediate results (such as KV vectors) generated by the previous round. By checking the storage continuity of this data in the storage, this inspection step can identify performance issues caused by non-continuous storage, lay the foundation for subsequent optimization, and ensure that the reasoning process can run under efficient data layout. For example, when it is detected that the reasoning-dependent data is stored non-continuously, the data is adjusted to a continuous state. This rearrangement operation reduces storage fragmentation, improves data loading efficiency, and avoids redundant calculations caused by zero filling. Compared with static pre-allocation filling technology, this method significantly improves storage utilization and reduces invalid overhead. After the data storage continuity is optimized, the specified hardware acceleration instruction is called to load continuous reasoning-dependent data and perform efficient operations. This process fully utilizes the hardware acceleration capabilities and further improves the reasoning efficiency.

[0061] Overall, this technical solution addresses the issues of low storage utilization and latency fluctuations in related technologies through dynamic hardware adaptation, storage optimization, and efficient instruction calls. Compared to static pre-allocated padding and dynamic storage requests, this approach offers significant advantages in storage management and computational efficiency. It is particularly well-suited for inference scenarios involving large models with variable-length sequences, effectively improving inference performance and hardware resource utilization.

[0062] Other features and advantages of the present application will be described in the following description, and in part will become apparent from the description, or will be understood by practicing the present application. The objectives and other advantages of the present application can be realized and obtained by the structures particularly pointed out in the written description, claims, and drawings. BRIEF DESCRIPTION OF THE DRAWINGS

[0063] The drawings described herein are used to provide a further understanding of the present application and constitute a part of the present application. The illustrative embodiments of the present application and their descriptions are used to explain the present application and do not constitute an improper limitation on the present application. In the drawings:

[0064] Figure 1 This is an optional schematic diagram of an application scenario in an embodiment of the present application;

[0065] Figure 2 A flowchart of an implementation of an inference acceleration method provided in an embodiment of the present application;

[0066] Figure 3 A schematic diagram of the calculation logic of a mode decision factor in an embodiment of the present application;

[0067] Figure 4 A schematic diagram of the design concept of a mode decision maker in an embodiment of the present application;

[0068] Figure 5 A schematic diagram of memory reorganization in an embodiment of the present application;

[0069] Figure 6 A schematic diagram of the calculation logic of an attention mechanism in an embodiment of the present application;

[0070] Figure 7 A schematic diagram of a transformation of an input sequence for a text generation scenario in an embodiment of the present application;

[0071] Figure 8 This is a schematic diagram of a discontinuous state in an embodiment of the present application;

[0072] Figure 9A This is a schematic diagram of the first TMA asynchronous characteristic in an embodiment of the present application;

[0073] Figure 9B This is a schematic diagram of the second TMA asynchronous characteristic in an embodiment of the present application;

[0074] Figure 10 A schematic diagram of a buffer switching logic in an embodiment of the present application;

[0075] Figure 11A A schematic diagram of processing logic of an instruction acceleration mode in an embodiment of the present application;

[0076] Figure 11B This is a processing logic diagram of another instruction acceleration mode in an embodiment of the present application;

[0077] Figure 12 A schematic diagram of a GPU according to an embodiment of the present application;

[0078] Figure 13This is a schematic diagram of an optional interaction implementation timing logic in an embodiment of the present application;

[0079] Figure 14 A schematic diagram of the structure of an inference acceleration device in an embodiment of the present application;

[0080] Figure 15 A schematic diagram of the hardware structure of an electronic device to which an embodiment of the present application is applied. DETAILED DESCRIPTION

[0081] To make the purpose, technical solutions, and advantages of the embodiments of this application more clear, the technical solutions of this application will be clearly and completely described below in conjunction with the drawings in the embodiments of this application. Obviously, the described embodiments are part of the embodiments of the technical solutions of this application, but not all of them. Based on the embodiments described in this application document, all other embodiments obtained by ordinary technicians in this field without making creative efforts shall fall within the scope of protection of the technical solutions of this application.

[0082] The following is an introduction to some concepts involved in the embodiments of this application.

[0083] 1. Target model: This refers to a pre-trained or customized machine learning model designed to perform a specific task, such as a transformer or bidirectional encoder representations from transformers (BERT). It typically contains multiple layers and parameters, capable of processing complex input data and generating corresponding outputs. The design and optimization of the target model directly impacts inference efficiency and accuracy, and is applicable to a variety of application scenarios, including text generation and image recognition.

[0084] 2. Tasks to be inferred: These are specific tasks that require analysis and processing by the target model, such as text classification in natural language processing and future value prediction in time series forecasting. These tasks typically involve input data of dynamically varying lengths and require efficient use of computing resources to obtain accurate results. Effective task design can significantly improve inference performance and user experience.

[0085] 3. Input sequence: This refers to the ordered set of data elements provided to the target model during an inference cycle, such as tokens in text or numerical points in a time series. The length of the input sequence can vary depending on the task, requiring flexible handling. Effective management and optimization of the input sequence can significantly improve the speed and accuracy of the inference process, meeting the needs of diverse application scenarios.

[0086] 4. Inference-dependent data: This is the intermediate result data generated and stored by the previous round of inference operations, such as the key-value vector in the attention mechanism. This data facilitates subsequent reasoning because it contains the results of previous calculations and contextual information. Ensuring efficient storage and fast access to inference-dependent data is key to achieving efficient inference, especially when processing long sequences or complex models.

[0087] 5. Hardware status information: This includes, but is not limited to, key metrics such as the current utilization of computing resources (such as GPUs), memory access efficiency, and bandwidth utilization. This information reflects the workload of hardware resources and helps the system dynamically adjust to the optimal inference mode. By monitoring hardware status information, resource allocation can be optimized, waiting time and resource waste can be reduced, and overall computing efficiency can be improved.

[0088] 6. Inference Mode: This refers to selecting the most appropriate execution strategy based on the hardware status of the computing resources occupied by the target model. In the embodiments of this application, inference modes include, but are not limited to, instruction acceleration mode and page cache mode. Each mode is designed to maximize computing resource utilization and inference speed based on different hardware conditions and task requirements. Selecting the appropriate inference mode facilitates efficient inference operations.

[0089] 7. Instruction Acceleration Mode: This is a specially optimized inference mode that utilizes hardware acceleration instructions (such as GMMA instructions) to directly load and process contiguously stored data blocks, significantly improving computational efficiency. In this mode, inference-dependent data is rearranged to ensure contiguous storage, reducing storage access latency and fragmentation, making large-scale matrix operations more efficient. It is particularly suitable for variable-length sequence inference scenarios.

[0090] 8. Paging cache mode: It is another specially optimized reasoning mode. It achieves efficient management by dividing the input sequence into multiple short sequences according to a preset length (for example, 64 tokens / block) and assigning an independent paging unit to each short sequence. The reasoning within each paging unit relies on continuous storage of data, and can be combined with the asynchronous prefetch strategy driven by the Tensor Memory Accelerator (TMA) to load the data of the next paging unit into the shared memory in advance, reducing storage fragmentation and data handling overhead. This mode is particularly suitable for variable-length sequence reasoning scenarios, which can effectively reduce storage requirements and increase reasoning speed while maintaining low latency and high hardware resource utilization. Through the memory-storage asynchronous transfer pipeline mechanism, the data handling process is further optimized to achieve overlapping execution of computing and data transmission.

[0091] 9. Paging Unit: In paging cache mode, a paging unit is a storage unit for each independent short sequence after the input sequence is divided into preset lengths (such as 64 tokens / block). Inference-dependent data within each paging unit is stored continuously, supporting TMA-driven asynchronous prefetching to pre-load the next unit's data into shared memory, reducing storage fragmentation and data transfer overhead, thereby improving inference speed, reducing latency, and optimizing hardware resource utilization.

[0092] 10. Compute Buffer: This refers to an area allocated in the shared memory of a computing resource (such as a GPU) specifically for storing data currently being processed. It is often paired with another preload buffer to form a double buffering mechanism. When a buffer serves as a compute buffer, the data in it is directly used to execute inference steps such as matrix multiplication and accumulation operations. This design maximizes GPU utilization, reduces latency, and improves overall computing efficiency.

[0093] 11. Preload Buffer: This is an area allocated in the shared memory of a compute resource (such as a GPU) for asynchronously loading data that will be processed soon. While one buffer is processing data as a compute buffer, the preload buffer loads the next batch of data from global memory. Once the data in the current compute buffer is processed and the data in the preload buffer is loaded, the two roles are swapped, ensuring a seamless data processing and loading process, significantly improving compute resource utilization and throughput.

[0094] 12. Mode Decision Factor: This is a value or set of values derived from a comprehensive assessment of the hardware status of computing resources. It determines which inference mode to activate at a given moment to optimize task execution efficiency. It is calculated based on multiple key hardware metrics, such as compute unit utilization, memory access efficiency, and bandwidth utilization, and aggregates these metrics through weighted summation or other mathematical methods. The primary function of the Mode Decision Factor is to dynamically adjust the inference strategy to ensure optimal performance under varying hardware conditions.

[0095] The following is a brief summary of the design ideas of the embodiments of this application:

[0096] Continuous advancements in computer technology have led to the emergence of large language models (LLMs) with powerful reasoning capabilities. These models enable computers to perform reasoning tasks, such as classifying articles, by invoking a vast array of model parameters to obtain inference results for category labels. Inference in large language models is an autoregressive process, where the current prediction depends on previously generated tokens. Key intermediate results (such as key-value vectors) generated during this process can be cached for use in subsequent reasoning steps. This mechanism is known as key-value caching.

[0097] Variable-length sequence reasoning is particularly important in tasks such as natural language processing (NLP) and time series analysis, as the input data length varies. For example, in text generation tasks, one sentence may contain 50 words, while another may contain 200 words. In time series forecasting, some time periods may have more data points than others. Therefore, when performing KV caching, the required storage size varies for sequences of different lengths.

[0098] In related technologies, static pre-allocation and padding techniques are currently commonly used for storage management. For example, mainstream frameworks such as FasterTransformer typically adopt a fixed-length storage pre-allocation strategy. This method pre-allocates continuous storage blocks of a maximum length (e.g., 4096 tokens) regardless of the actual sequence length. For shorter sequences, the remaining space needs to be filled with zeros, which not only leads to low storage utilization (storage fragmentation rate as high as 27%), but also introduces redundant calculations, wasting about 15%-30% of computing power.

[0099] In summary, the relevant technical solutions have the following two core contradictions:

[0100] (1) Waste of computing power in short sequence scenarios: Traditional solutions cannot use the GPU architecture to achieve zero-fill free computing, which limits the throughput of short sequence reasoning.

[0101] (2) Inefficient storage management: Current technical solutions have problems such as severe storage fragmentation and high data transfer overhead, which affect the deployment efficiency of large-scale model real-time inference services.

[0102] In view of this, embodiments of the present application provide an inference acceleration method, apparatus, electronic device, and storage medium. By dynamically monitoring the hardware status information of computing resources, the inference mode that best matches the current hardware status is selected. This dynamic adjustment mechanism optimizes resource utilization and ensures efficient inference operations under different hardware conditions.

[0103] In instruction acceleration mode, the input sequence-related reasoning-dependent data usually includes intermediate results generated by the previous round. By checking the storage continuity of this data in the storage, this inspection step can identify performance issues caused by non-contiguous storage, laying the foundation for subsequent optimization, and ensuring that the reasoning process can run under efficient data layout. For example, when it is detected that the reasoning-dependent data is stored non-contiguously, the data is adjusted to a continuous state. This rearrangement operation reduces storage fragmentation, improves data loading efficiency, and avoids redundant calculations caused by zero padding. Compared with the static pre-allocation padding technology, this method significantly improves storage utilization and reduces invalid overhead. After the data storage continuity is optimized, the specified hardware acceleration instruction is called to load the continuous reasoning-dependent data and perform efficient operations. This process fully utilizes the hardware acceleration capabilities and further improves the reasoning efficiency.

[0104] Overall, this technical solution addresses the issues of low storage utilization and latency fluctuations in related technologies through dynamic hardware adaptation, storage optimization, and efficient instruction calls. Compared to static pre-allocated padding and dynamic storage requests, this approach offers significant advantages in storage management and computational efficiency. It is particularly well-suited for inference scenarios involving large models with variable-length sequences, effectively improving inference performance and hardware resource utilization.

[0105] The preferred embodiments of the present application are described below in conjunction with the drawings in the specification. It should be understood that the preferred embodiments described herein are only used to illustrate and explain the present application and are not used to limit the present application. In addition, the embodiments and features in the embodiments of the present application can be combined with each other if there is no conflict.

[0106] like Figure 1 , which is a schematic diagram of an application scenario of an embodiment of the present application. The application scenario diagram includes a terminal device 110 and a server 120.

[0107] In the embodiment of the present application, the terminal device 110 includes but is not limited to mobile phones, tablet computers, laptop computers, desktop computers, e-book readers, intelligent voice interaction devices, smart home appliances, car terminals and other devices; a task reasoning-related client can be installed on the terminal device, which can be software, such as a browser, artificial intelligence (AI) dialogue software, image processing software, voice assistant, etc., or a web page, applet, etc. The server 120 is a background server corresponding to the software or web page, applet, etc., or a server specifically used for task reasoning acceleration, or a server for deploying target models, etc., which is not specifically limited in this application. The server 120 can be an independent physical server, or a server cluster or distributed system composed of multiple physical servers, or a cloud server that provides cloud services, cloud databases, cloud computing, cloud functions, cloud storage, network services, cloud communications, middleware services, domain name services, security services, content delivery networks (CDNs), and basic cloud computing services such as big data and artificial intelligence platforms.

[0108] It should be noted that the inference acceleration method in each embodiment of the present application can be executed by an electronic device, which can be a terminal device 110 or a server 120, that is, the method can be executed by the terminal device 110 or the server 120 alone, or can be executed jointly by the terminal device 110 and the server 120.

[0109] For example, when executing independently on server 120, the target model can be deployed on server 120 and computing resources can be provided for the target model. The task to be inferred can be uploaded by the object through terminal device 110, or obtained by the server from a database or other end, which is not specifically limited in this document. In this way, when using the target model deployed on server 120 to execute the task to be inferred, server 120 can first activate the inference mode that matches the current hardware status of the computing resources occupied by the target model based on the hardware status information of the computing resources.

[0110] When the inference mode is the instruction acceleration mode, the server 120 may perform at least one round of inference operations, each round of inference operations including the following steps:

[0111] For the inference dependency data related to the input sequence of this round, storage continuity check is performed in the computing resources; the inference dependency data is generated and stored through the inference operation of the previous round; when the storage location of the inference dependency data is in a discontinuous state in the computing resources, the storage location is adjusted to a continuous state by reordering the inference dependency data; calling the specified hardware acceleration instruction to load the adjusted inference dependency data, and after determining the result of this round of inference based on the loading result, generating the input sequence of the next round according to the result of this round of inference, until the preset stop condition is reached.

[0112] In an optional implementation, the terminal device 110 and the server 120 may communicate via a communication network.

[0113] In an optional implementation, the communication network is a wired network or a wireless network.

[0114] It should be noted that Figure 1 The examples shown are just for illustration. In fact, the number of terminal devices and servers is not limited and is not specifically limited in the embodiments of this application.

[0115] In an embodiment of the present application, when there are multiple servers, the multiple servers can be combined into a blockchain, and the servers are nodes on the blockchain; as disclosed in the embodiment of the present application, the relevant data involved can be saved on the blockchain, such as hardware status information, input sequence, task execution results, reasoning dependency data, etc.

[0116] In addition, the embodiments of the present application can be applied to various scenarios, including but not limited to cloud technology, artificial intelligence, smart transportation, assisted driving, etc. The following are a few examples:

[0117] Scenario 1: Model training phase in cloud technology.

[0118] In a cloud computing environment, the training of large-scale deep learning models usually requires a large amount of computing resources and storage space. In the training scenario, taking natural language processing tasks as an example, the input sequence of the model is usually a text fragment or a complete sentence after word segmentation. The trained model can be applied to application scenarios such as machine translation and text generation. The reasoning acceleration method proposed in this application can significantly improve the efficiency of the model training process. For example, when training a language model based on the Transformer architecture in the cloud, the hardware status information will be monitored in real time to dynamically adjust the reasoning mode. When it is detected that the hardware status information meets certain conditions, the system will activate the instruction acceleration mode, check and reorder the storage continuity of the reasoning-dependent data, and ensure the efficiency of data loading. In this way, not only is the fragmentation of the video memory reduced, but also redundant calculations are reduced, so that large-scale parallel computing tasks can be executed efficiently. In addition, by utilizing the elastic resource allocation mechanism provided by the cloud platform, this method can automatically expand or reduce computing resources according to actual needs, thereby optimizing cost and performance.

[0119] Scenario 2: Application stage in the field of artificial intelligence.

[0120] The method of this application is also of great value in practical applications of artificial intelligence, such as text generation in natural language processing (NLP) tasks. Suppose a BERT-based question-answering system is being deployed that needs to quickly respond to user queries. In this process, the continuity and efficient loading of inference-dependent data (such as the key-value vectors in the attention mechanism) are crucial. Using this method, the storage continuity of this data can be checked before inference begins, and any discontinuous data can be rearranged to ensure continuous storage of data blocks in shared memory. After invoking hardware acceleration instructions, the system can quickly load this data and perform efficient computational operations (such as matrix multiplication), significantly reducing inference time. This not only improves the user experience but also enables the system to operate stably in high-concurrency environments, meeting real-time requirements. For example, in a news scenario, the above method can be used to automatically analyze and generate news summaries, categorize news articles, and provide personalized content recommendations to users. It also supports intelligent search and question-answering systems, making information acquisition more efficient and accurate, thereby improving content management and user experience. Furthermore, it can assist content creators in their creation, check and optimize article quality, thereby increasing the overall appeal and interactivity of the content.

[0121] Scenario 3: Assisted driving stage in smart transportation.

[0122] In the field of smart transportation, especially in assisted driving systems, real-time processing of large amounts of sensor data is one of the key challenges, that is, the input sequence of the model can be environmental data obtained by sensors (such as cameras, radars), and the feature vectors or image representations formed after processing. For example, autonomous driving vehicles need to continuously analyze input data from devices such as cameras and radars to make immediate decisions. The method of the present application can improve the response speed of the system by optimizing the management and loading process of reasoning-dependent data. Specifically, in each round of reasoning operation, the system performs a storage continuity check on the reasoning-dependent data related to the current input sequence and rearranges it in discontinuous cases. This method effectively reduces data handling overhead and improves bandwidth utilization. In addition, it can also be combined with the asynchronous prefetch strategy driven by TMA to further improve data transmission efficiency, so that the assisted driving system can complete complex perception and decision-making tasks within milliseconds, thereby enhancing driving safety.

[0123] The above three application scenarios simply demonstrate the applicability and advantages of the present method in different fields, and can significantly improve the performance and efficiency of the system in both the training and application stages. In addition, the above application scenarios are just simple examples, and other application scenarios are also applicable to the embodiments of the present application, which will not be detailed here.

[0124] It is also important to emphasize that in the specific implementations of this application, object-related data is involved, such as the object input-related data to be inferred as listed above. When the above embodiments of this application are applied to specific products or technologies, the object's permission or consent must be obtained, and the collection, use, and processing of relevant data must comply with the relevant laws, regulations, and standards of the relevant countries and regions.

[0125] The following describes the reasoning acceleration method provided by the exemplary embodiment of the present application in combination with the application scenarios described above and with reference to the accompanying drawings. It should be noted that the above application scenarios are only shown to facilitate understanding of the spirit and principles of the present application, and the implementation of the present application is not limited in this respect.

[0126] See Figure 2 The figure is a flowchart of an implementation of an inference acceleration method provided in an embodiment of the present application. Taking the server as the execution subject as an example, the specific implementation process of the method is as follows:

[0127] S21: When the target model is used to execute the task to be inferred, based on the hardware status information of the computing resources occupied by the target model, an inference mode that matches the current hardware status of the computing resources is activated.

[0128] In the embodiments of this application, the target model refers to a pre-trained or customized machine learning model used to perform a specific task, such as a deep learning model, a reinforcement learning model, a large language model, etc., which is used to obtain the task result through step-by-step reasoning when performing the task to be reasoned. This document does not specifically limit this. The following are a few examples:

[0129] (1) Transformer, suitable for tasks such as text generation, translation, and question-answering systems.

[0130] In simple terms, the Transformer model, especially its variants such as the Generative Pre-trained Transformer (GPT), generates output token by token through an attention mechanism. For example, in a text generation task, the model predicts the next token based on the previous token until a predetermined stopping condition is reached.

[0131] (2) BERT, suitable for tasks such as text classification, question answering systems, and named entity recognition.

[0132] In short, BERT is mainly used to understand the context of the entire sentence, and in some tasks, it also needs to predict token by token. For example, in a question-answering system, the model may need to generate answers token by token.

[0133] (3) Long Short-Term Memory (LSTM) or its simplified variant, Gated Recurrent Unit (GRU), is suitable for tasks such as time series prediction, speech recognition, and natural language processing.

[0134] Simply put, LSTM is a variant of the Recurrent Neural Network (RNN), particularly well-suited for processing sequential data and capturing dependencies within long time series. For example, in speech recognition, speech synthesis, or text generation tasks, LSTM typically generates outputs token by token, updating the current state based on the previous state and predicting the next token. Similarly, GRU is also applicable to these tasks, so I won't elaborate on them here.

[0135] (4) Pointer Networks, suitable for tasks such as question answering systems and text summarization.

[0136] Simply put, this type of network is often used for generative tasks, where the model not only generates new words but can also copy words directly from the input sequence. During the generative process, the model decides whether to generate a new word or copy an existing word on a token-by-token basis. For example, in text summarization, a pointer network can decide whether to generate a new word or copy a word directly from the input.

[0137] It's important to note that all of the above models require step-by-step reasoning (e.g., token-by-token prediction), especially when processing sequential data. These models rely on previously generated partial results as they generate output, gradually building up the final output sequence. This makes them particularly well-suited for tasks that require sequential output generation, such as natural language processing, speech recognition, and time series forecasting.

[0138] In addition, it should be noted that the target models and tasks to be inferred listed above are just simple examples. In addition, other models or tasks are also applicable to the embodiments of this application and will not be described one by one here.

[0139] Among them, the computing resources occupied by the target model refer to the hardware facilities and supporting technologies required to run and reason about the model, including but not limited to GPU, Central Processing Unit (CPU), Tensor Processing Unit (TPU), etc., which are not specifically defined in this article. These computing resources provide the necessary computing power and memory space to support the model for efficient numerical calculations, data processing and storage. Specifically, GPUs excel at parallel computing and are suitable for large-scale matrix operations; CPUs provide strong general computing power and are suitable for processing logically complex tasks; and TPUs are accelerators specially designed for machine learning and deep learning, which can efficiently perform tensor operations and increase the speed of model training and reasoning. Choosing appropriate computing resources is conducive to optimizing model performance, reducing latency and improving energy efficiency.

[0140] Computing resource hardware status information refers to a series of indicators that reflect the current operating status of computing resources. Specifically, these indicators can be hardware-related metrics used to evaluate and optimize system performance during the execution of inference tasks. These metrics include, but are not limited to, key indicators such as compute unit utilization, memory access efficiency, and bandwidth utilization. This information can help identify system bottlenecks, optimize resource utilization, and ensure optimal task execution. By continuously monitoring and analyzing these indicators, problems can be promptly identified and appropriate optimization measures can be implemented to ensure that the system is always performing at its best.

[0141] Specifically, in step S21, when the target model is used to execute the task to be inferred, the hardware status information of the computing resources occupied by the target model must be monitored in real time. This hardware status information reflects the workload of the current computing resources. Based on this information, the system will dynamically activate the inference mode that matches the current hardware status to optimize the task execution efficiency. The following is a specific example:

[0142] For example, high compute unit utilization (e.g., GPU cores nearing full capacity) indicates efficient use of computing resources; low memory access efficiency (e.g., low cache hit rate) indicates significant latency during data loading. In such cases, instruction acceleration mode can be activated. This allows hardware-accelerated instructions to directly load contiguously stored data blocks and leverage efficient matrix operations to reduce the impact of memory access latency on performance.

[0143] For example, low compute unit utilization (e.g., high GPU core idle cycles) indicates underutilized compute resources; high memory access efficiency (e.g., high cache hit rate) indicates smooth data loading without significant bottlenecks. In this case, you can maintain the current inference mode or switch to a lighter-weight mode (e.g., low-precision inference mode). By reducing computational precision (e.g., switching from FP32 to BF16), computational throughput can be further improved, leveraging the benefits of memory access efficiency.

[0144] For example, a moderate level of compute unit utilization (e.g., a busy GPU core) indicates that the compute resources are under load but still have sufficient capacity. Low bandwidth utilization (e.g., underutilized video memory bandwidth) indicates that data transfer channels are idle. In this case, the paging cache mode can be activated. This divides the input sequence into multiple short sequences, assigning each short sequence to a different paging unit for continuous storage, and combines it with a TMA-driven asynchronous prefetching strategy to fully utilize video memory bandwidth resources.

[0145] It should be noted that the several mode selection methods listed above are just simple examples. Other methods of setting the inference mode based on hardware status information are also applicable to the embodiments of this application and will not be repeated here. This dynamic adjustment mechanism can flexibly optimize resource utilization based on the actual hardware status, significantly improve the execution efficiency of inference tasks, and adapt to the performance requirements in different scenarios.

[0146] To further refine this process, this application provides a specific implementation method that quantifies hardware status information to determine the mode decision factor and selects a more appropriate inference mode based on this factor. This method not only improves the accuracy of inference mode selection but also ensures that the system can maintain optimal performance under different hardware conditions. The specific implementation method is as follows:

[0147] An optional implementation is to implement S21 according to the following process:

[0148] First, the mode decision factor corresponding to the computing resources is determined by quantifying the hardware status information.

[0149] In an embodiment of the present application, the hardware status information of the computing resources occupied by the target model can be obtained through real-time monitoring. For example, the system can use built-in or external monitoring tools to regularly collect key performance indicators of computing resources, including but not limited to computing unit utilization, memory access efficiency, and bandwidth utilization. These monitoring tools can be standard interfaces provided by the operating system or advanced monitoring frameworks designed specifically for deep learning tasks, such as the nvidia-smi tool or other third-party performance monitoring software.

[0150] nvidia-smi is a command-line tool for managing and monitoring GPU devices. It displays GPU usage (such as video memory usage and GPU utilization), operating status, temperature, power consumption, and other information. It also supports basic management operations such as adjusting performance mode, limiting power consumption, or resetting the GPU.

[0151] For example, if the computing resource is a GPU, you can use tools such as nvidia-smi or similar application programming interfaces (APIs) to obtain real-time GPU hardware status information. These tools can provide hardware status information such as GPU utilization, memory / video memory usage, bandwidth usage, temperature, power consumption, etc.

[0152] For example, the operating system's performance monitoring interface can be used to obtain hardware status information such as GPU utilization, memory / video memory usage, bandwidth usage, temperature, and power consumption.

[0153] Optionally, hardware status information can also be obtained through a custom monitoring module integrated into the target model. This can be done by developing a dedicated monitoring module or integrating an existing monitoring framework, such as Prometheus (an open source system monitoring and alarm toolkit) or Grafana (an open source visualization and analysis platform), to periodically collect hardware status information about the computing resources used by the target model. This module can more finely control the time interval and frequency of data collection, ensuring that changes in hardware status can be captured in a timely manner without affecting system performance. For example, key indicators can be recorded at the beginning and end of each inference cycle, or data collection can be triggered based on specific events (such as task completion, anomaly detection, etc.).

[0154] Of course, in addition to the real-time collection method mentioned above, periodic collection can also be performed, such as setting a reasonable collection period (for example, every second or every few milliseconds) to ensure that changes in hardware status are captured in a timely manner. Furthermore, event-driven collection can be used, such as triggering data collection when specific events occur (such as task start, end, or exception) to more accurately reflect actual workload changes. This article does not specifically limit this.

[0155] After collecting hardware status information, the next step is to quantify it. This step involves converting the raw monitoring data into a numerical form that can be used for subsequent analysis. For example, compute unit utilization can be determined by calculating the proportion of idle cycles per unit time; memory access efficiency can be measured by counting the ratio of cache hits to total accesses; bandwidth utilization is the ratio of used bandwidth to total available bandwidth; and so on.

[0156] Based on this quantified hardware status information, the system comprehensively determines a mode decision factor, which guides the system in dynamically selecting the most appropriate inference mode, ensuring optimal performance under varying load conditions. This approach not only enables the system to efficiently utilize existing resources but also flexibly adapts to various complex operating environments, providing stable and efficient inference services.

[0157] Based on the above, to more specifically illustrate the quantification method of hardware status information and its impact on the mode decision factor, the following further explains how the mode decision factor is determined based on computing unit utilization, memory access efficiency, and bandwidth utilization:

[0158] In an optional implementation, the hardware status information includes computing unit utilization, memory access efficiency, and bandwidth utilization. In this case, determining the mode decision factor corresponding to the computing resource by quantifying the hardware status information includes the following steps:

[0159] First, it is necessary to determine the computing unit utilization based on the number of idle cycles of the target computing unit in the computing resource within the preset time period; determine the memory access efficiency based on the number of accesses to the target memory layer in the computing resource within the preset time period; and determine the bandwidth utilization based on the amount of bandwidth used by the target storage unit in the computing resource within the preset time period.

[0160] The preset time period can be a unit of time or any length of time set based on specific application scenarios, such as one second, one minute, or even longer, to more accurately reflect system performance under different load conditions. By flexibly adjusting the preset time period, it is possible to effectively monitor and analyze trends in computing unit utilization, memory access efficiency, and bandwidth utilization, providing a basis for optimization decisions.

[0161] This flexibility allows the monitoring period to be customized for specific application scenarios or test requirements, ensuring that the system's transient state or long-term performance characteristics are captured. For example, for rapidly changing tasks, a shorter period can be selected to capture real-time performance fluctuations; while for more stable workloads, a longer period can be used for evaluation.

[0162] The target compute unit refers to the processor core that performs specific computational tasks within the computing resources occupied by the target model. The target memory layer typically refers to the storage layer between the fast L1 cache and the slower video memory, such as L2 cache. The target storage unit refers to the video memory used for long-term data storage and exchange, such as DRAM. These components work together to ensure efficient data processing.

[0163] The target computing unit refers to the core component that performs specific computing tasks in the computing resources occupied by the target model, such as the streaming multiprocessors (SM) in the GPU, or the core or dedicated accelerator in the CPU. This article does not make specific restrictions on this.

[0164] The target memory tier is a specific layer in the computing resource's storage architecture, used to temporarily store data or instructions being processed for fast access. It typically includes the L1 cache, L2 cache, and closely coordinated main memory, such as dynamic random access memory (DRAM), though this article does not specifically define this. For example, in data-intensive applications, such as machine learning inference, frequently used model parameters are loaded into the L2 cache, while more intermediate data may reside in DRAM, ensuring high and stable access rates and supporting efficient computing operations.

[0165] The target storage unit refers to a storage unit in the computing resource used for long-term data storage, such as DRAM, Video Random Access Memory (VRAM), etc., which is not specifically limited in this article.

[0166] After calculating the computing unit utilization, memory access efficiency, and bandwidth utilization, a weighted sum is performed based on the computing unit utilization, memory access efficiency, and bandwidth utilization to obtain a mode decision factor.

[0167] The following first explains how to calculate computing unit utilization, memory access efficiency, and bandwidth utilization:

[0168] (1) The computing unit utilization is determined based on the number of idle cycles of the target computing unit in the computing resources within a preset time period. An optional calculation method is as follows:

[0169] C idle =T total / T idle (Formula 1);

[0170] Among them, C idle represents the computing unit utilization, T total Indicates the total number of cycles (such as time periods) within a preset time period, and T idle It is the number of cycles that the target computing unit is idle during a preset time period.

[0171] In the embodiment of the present application, this ratio C idle Reflects the utilization of the target computing unit. When T idle The smaller the C idle The larger it is, the less idle time the target computing unit has and the higher the utilization rate is, which means the target computing unit is busier. idle A high value may indicate a high computational load. On the contrary, when T idle The larger the C idle The smaller it is, the more idle time the target computing unit has and the lower the utilization rate is, that is, the target computing unit is more idle. Therefore, C idle Low probably indicates a low computational load.

[0172] For example, the total time is 100ms, the idle time is 10ms, C idle =100 / 10=10; if the idle time increases to 50ms, C idle =100 / 50=2.

[0173] Taking a GPU as an example of a computing resource, the target computing unit may refer to a streaming multiprocessor (SM) in the GPU.

[0174] Specifically, an SM is the core computing unit in the GPU architecture, responsible for executing the various parallel computing tasks assigned to it. In a GPU (such as the Hopper GPU), a GPU chip typically contains multiple SMs. Each SM contains multiple computing cores, shared memory, register files, and other specialized hardware units such as Tensor Cores and Ray Tracing Cores (RT Cores). These SMs are responsible for performing actual computational operations such as matrix multiplication and vector operations. The specific implementation depends on the architecture and is not specifically defined in this article.

[0175] SM idle cycles refer to the proportion of time these processing units are unused over a period of time. Specifically, they are expressed as the number of cycles in which the computing units are idle. Each SM in the GPU executes instructions synchronously according to an internal clock signal. This internal clock signal determines when each operation (such as addition, multiplication, etc.) starts and ends. One clock cycle is the time required to complete such a basic operation. High SM idle cycles mean that computing resources are not fully utilized, indicating that it may be necessary to switch to a more efficient computing mode to improve resource utilization. Therefore, by monitoring and optimizing SM usage, the overall performance and efficiency of the GPU can be significantly improved.

[0176] Of course, the computing resources and target computing units can also be other, not limited to GPU and SM. Similarly, the target memory layer is not limited to L2 Cache below, and the target storage unit is not limited to DRAM below. This article does not make specific restrictions on this. The following still uses the GPU as an example to explain other hardware status information:

[0177] (2) Memory access efficiency is determined based on the number of accesses to the target memory layer in the computing resource within a preset time period. Taking the target memory layer as the L2 cache in the GPU as an example, an optional calculation method is as follows:

[0178] H L2 =N hit / N total (Formula 2);

[0179] Among them, H L2 Indicates memory access efficiency, N total It is the total number of memory accesses within a preset time period, which can be referred to as the total number of accesses. N hit It is the number of times the required data is found in the L2 Cache during these accesses, which can be referred to as the number of L2 access hits.

[0180] In the embodiment of the present application, this ratio represents the memory access efficiency, H L2 When high, it means that a higher proportion of memory access requests can directly obtain the required data from the L2 Cache, which means that the memory access efficiency is good. At this time, 1-H L2 A lower value of H indicates that fewer requests need to go to slower video memory (such as DRAM), resulting in improved overall system performance. This usually occurs when the load is moderate or low, because efficient cache usage means that the computing unit can quickly obtain the required data, reducing waiting time; conversely, H L2 Low, 1-H L2 A higher value of indicates that more requests need to be directed to slower video memory (such as DRAM), which usually occurs under high load conditions.

[0181] Specifically, a GPU is a processor designed for fast graphics processing and massively parallel computing tasks. The GPU's memory architecture follows a hierarchy from fastest to slowest, from smallest to largest, typically going from registers to L1 cache to L2 cache to video memory. Data is first attempted from the fastest cache level; if a miss occurs, the search continues to the next level, ultimately reaching video memory. To optimize data access speed, GPUs utilize a multi-level cache system (such as L1 and L2 caches) to reduce the need for direct access to video memory, thereby improving efficiency.

[0182] (3) Bandwidth utilization is determined based on the amount of bandwidth used by the target storage unit in the computing resource within a preset time period. An optional calculation method is as follows:

[0183] U DRAM =B used / B max (Formula 3);

[0184] Among them, U DRAM Indicates bandwidth utilization, B max is the maximum bandwidth of the video memory within the preset time period, and B used It is the actual usage bandwidth of video memory during the preset time period.

[0185] In the embodiment of the present application, this ratio is used to measure the pressure on the video memory. A value close to 1 indicates that the video memory bandwidth is almost fully utilized and may become a bottleneck. This situation usually occurs under high load conditions; while a lower value indicates that there is sufficient bandwidth that is not fully utilized. This situation usually occurs under low load conditions.

[0186] Based on the hardware status information obtained by the above calculation, the mode decision factor can be further calculated by weighted summation. The weight reflects the influence of different indicators on the overall performance.

[0187] In an optional implementation, the weights of various hardware status indicators (i.e., hardware status information / signals) can be set based on the needs of specific application scenarios. Different application scenarios may require assigning different weights to various hardware status indicators to reflect their impact on overall performance. For example, in compute-intensive tasks, the weight of computing unit utilization may be higher; while in data-intensive tasks, the weights of memory access efficiency and bandwidth utilization may be more important; and so on.

[0188] In another optional implementation, the weights of various hardware status indicators can be dynamically adjusted according to the characteristics of the application scenario to more accurately reflect the impact of these indicators on the overall performance. For example, for real-time data analysis scenarios, this scenario requires not only fast data processing capabilities, but also efficient memory management and data transmission rates. Therefore, the weights of memory access efficiency (such as L2 Cache hit rate) and bandwidth utilization (such as DRAM usage) can be increased, which means that the system pays more attention to how to reduce data read latency and improve data throughput to ensure that it can quickly respond to changing data streams. For graphics-intensive applications or high-performance computing tasks, when it comes to complex image rendering or large-scale scientific computing, the importance of computing unit utilization is once again highlighted. At this time, a higher weight can be given to computing unit utilization, while appropriately considering bandwidth utilization, because such tasks usually require a large number of parallel computing resources and efficient data exchange to maintain high computing throughput; and so on.

[0189] By flexibly setting the weights of various hardware status indicators based on specific application requirements, we can better meet performance requirements in different scenarios and achieve optimal resource allocation and utilization. This approach emphasizes the importance of tailoring monitoring and optimization strategies to different task types, helping to improve system adaptability and efficiency.

[0190] In addition, it should be noted that the weight setting methods listed above are just simple examples. Other weight setting methods can also be used, such as 1 / 3 (weighted average), and this article does not make specific restrictions on this.

[0191] When performing weighted summation, an optional calculation method is as shown in the following formula 4:

[0192] D=α·C idle +β·(1-H L2 )+γ·U DRAM (Formula 4);

[0193] Among them, D is the mode decision factor, C idle The weight is α, (1-H L2 ) is weighted by β, U DRAM The weights are γ, which are used to adjust the influence of each indicator on the final decision. Their values can be determined through offline training and calibration to ensure that the model accurately reflects the priorities and influences in real-world application scenarios.

[0194] Among them, θ1 is the first threshold and θ2 is the second threshold. These two thresholds are used to decide when to switch to different operation modes. These thresholds can also be set through offline training to adapt to specific workloads and hardware configurations.

[0195] In formula 4, H is not directly L2 Perform weighted summation, but (1-H L2 ) is weighted summed to convert the L2 Cache hit rate H L2 Converted into an indicator reflecting the miss rate to reflect the low efficiency of memory access.

[0196] Specifically, H L2 Indicates the L2 Cache hit rate, that is, the proportion of memory access requests that can directly obtain data from the L2 Cache. L2 A higher value means that more data accesses can be obtained directly from the cache, reducing reliance on slower video memory (such as DRAM), thereby improving performance.

[0197] And 1-H L2 Indicates the miss rate of L2 Cache, that is, the proportion of memory access requests that fail to obtain data from L2 Cache and must turn to slower video memory (such as DRAM). L2 A higher value means that more memory accesses need to rely on video memory, which increases data access latency and video memory pressure, thus affecting system performance.

[0198] In the above dynamic decision model, we hope to measure the impact of negative factors on system performance. For example: C idle It measures the idle time of computing units and reflects the inefficient use of computing resources. DRAM It measures the memory bandwidth utilization and reflects the pressure of the memory. Similarly, (1-H L2 ) is used to reflect the low efficiency of memory access because it represents the proportion of cache misses, which lead to higher latency and lower efficiency.

[0199] If you use H directly L2 , then a high hit rate will appear as a lower value, which is consistent with C idle and U DRAM The logic is inconsistent. By using (1-H L2 ), which can be converted into a load-related negative impact indicator to maintain consistency with other terms.

[0200] Therefore, this conversion makes the various indicators C in formula 4 idle 、(1-H L2 ) and U DRAM They can all point to the negative impact on system performance, facilitating comprehensive evaluation and decision-making.

[0201] Optional, taking into account (1-H L2 ) and U DRAMThey are all floating point numbers between 0 and 1. Before calculating with formula 4, you can also first idle Normalization processing is performed, and the processing method is not specifically limited in this article.

[0202] In the embodiment of the present application, the dynamic decision model uses the weighted threshold comparison method to make the decision of the reasoning mode, that is, when D>θ1, C idle High (high computing utilization), H L2 Low (poor memory efficiency), U DRAM High (high video memory pressure), the weighted sum of these three factors leads to high D, so the overall load is high, and you can switch to instruction acceleration mode (such as GMMA mode, also known as GMMA high-speed mode). In this mode, it is mainly used to process short sequences; when D<θ2, C idl e is low (computing is more idle), H L2 High (good memory efficiency), U DRAM Low (low video memory pressure), low weighted sum, so the overall load is low, and you can switch to page cache mode, which is mainly used to process long sequences.

[0203] like Figure 3 As shown, it is a calculation logic diagram of a mode decision factor in an embodiment of the present application. Assuming that a text generation task is being executed, and it is desired to dynamically adjust the inference mode to optimize performance, nvidia-smi can be used to collect the hardware status information of the GPU every second, wherein the calculation of the computing unit utilization needs to be determined in combination with the idle cycle and the total cycle, and the calculation method is as shown in Formula 1; the calculation of the memory access efficiency needs to be determined in combination with the number of L2 access hits and the total number of accesses, and the calculation method is as shown in Formula 2; the calculation of the bandwidth utilization needs to be determined in combination with the actual bandwidth used and the maximum bandwidth, and the calculation method is as shown in Formula 3. On this basis, the mode decision factor can be determined by referring to Formula 4. The specific calculation process can be found in the above embodiment, and will not be repeated here.

[0204] Assuming the compute unit utilization is 0.8, the memory access efficiency is 0.9, and the bandwidth utilization is 0.7, with weights of 1 / 3, the mode decision factor is 0.8. Based on this mode decision factor, the system can select the appropriate inference mode (such as instruction acceleration mode or page cache mode) based on the preset first and second thresholds, thereby achieving dynamic adjustment to optimize task execution efficiency.

[0205] The hardware status information quantification method in the above embodiment makes the calculation of the mode decision factor more accurate and reliable. By considering multiple key indicators such as computing unit utilization, memory access efficiency, and bandwidth utilization, the system can comprehensively evaluate the current hardware status and make the optimal inference mode selection accordingly. This method not only improves inference efficiency, but also enhances system stability and responsiveness, especially in high-concurrency and resource-constrained environments.

[0206] It should be noted that the calculation method of the mode decision factor listed above is just a simple example. For example, only part of the computing unit utilization, memory access efficiency and bandwidth utilization can be considered for weighted summation. Alternatively, the mode decision factor can be evaluated in combination with other types of hardware status information. The specific evaluation method is not specifically limited in this article, such as weighted summation after quantification, and the hierarchical analysis method, fuzzy logic method, machine learning method, decision tree method, etc. mentioned below. In addition to the key indicators mentioned above, the following hardware status signals can also be considered as the basis for optimization decisions in this scenario:

[0207] (1) Power Efficiency: Power efficiency refers to the ratio of the energy consumed by a GPU when performing a specific task to the speed at which it completes that task. In embodiments of the present application, the power efficiency or average power efficiency of the GPU during historical inference tasks can be monitored to inform the inference mode for subsequent tasks.

[0208] Specifically, the actual power consumption of the GPU (in watts) can usually be measured through dedicated sensors, and the power efficiency can be evaluated in combination with its current workload. Specifically, the power efficiency can be evaluated by comparing the power consumption and performance (such as computing power per watt) under different workloads. For example, the energy required to complete a specific task (power consumption multiplied by time) can be recorded, and then the power efficiency can be evaluated based on the quality or speed of task completion.

[0209] (2) Temperature: Temperature monitoring involves real-time detection of the GPU chip's operating temperature to ensure it remains within a safe operating range. Excessive temperatures may cause throttling (i.e., automatically lowering the frequency to reduce heat), thereby affecting performance.

[0210] Specifically, most GPUs have built-in temperature sensors that can provide real-time temperature readings. This data can be obtained through APIs provided by the operating system or dedicated monitoring software.

[0211] It should be noted that these additional monitoring indicators provide a more comprehensive understanding of the system's operating status, which helps to make more refined optimization adjustments. By comprehensively considering these factors, resource utilization can be better optimized, potential problems can be avoided, and overall work efficiency can be improved. Of course, in addition to these, other hardware status information is also applicable to the embodiments of this application, such as the number of active threads, throughput, etc., which will not be detailed here.

[0212] Optionally, in addition to the above-mentioned method of determining the mode decision factor by quantifying the hardware status information, in the embodiment of the present application, the hardware status information may be processed in other ways to determine the mode decision factor. The following briefly lists several methods:

[0213] In an alternative implementation, hardware status information can be decomposed into a target layer, a criterion layer, and a solution layer using the Analytic Hierarchy Process (AHP). Each pair of hardware status indicators is then compared to construct a judgment matrix. Subsequently, the weight of each hardware status indicator is calculated and combined into the final mode decision factor. This approach takes into account the relative importance of each indicator, making it suitable for multi-dimensional evaluation and ensuring balanced and reasonable decisions in complex environments.

[0214] In another alternative implementation, fuzzy logic can be used to map hardware status information into fuzzy sets and define corresponding fuzzy rule sets. These fuzzy rules are used to evaluate the hardware status and generate fuzzy outputs, which are then converted into specific pattern decision factors through defuzzification techniques. This approach has the advantage of being able to handle uncertainty and ambiguity, making it well-suited for decision-making in complex environments and offering greater flexibility and adaptability.

[0215] In another optional implementation, a machine learning model can be used. By collecting historical hardware state information and its corresponding optimal inference patterns as training data, a supervised learning algorithm (such as a random forest or support vector machine) can be used to train the model. When the current hardware state information is input, the model can predict the optimal mode decision factor. The advantage of this approach is that it can automatically learn complex nonlinear relationships and provide highly personalized optimization strategies, which is particularly suitable for dynamically changing workloads.

[0216] In another alternative implementation, a decision tree approach can be used, using hardware status information as input features. Based on a pre-defined splitting criterion (such as information gain or the Gini coefficient), nodes are gradually partitioned until a leaf node is reached. The value corresponding to the leaf node is the model decision factor. This approach is highly interpretable, easy to understand, and implement, making it particularly suitable for applications requiring rapid response and transparent decision-making.

[0217] The several implementation methods listed above are just simple examples. In actual applications, the most suitable method can be flexibly selected to determine the mode decision factor according to the needs of the specific application scenario, so as to achieve more accurate reasoning mode adjustment and maximize system performance. This article does not make specific restrictions on this.

[0218] After determining the mode decision factor by any of the above methods:

[0219] Next, the mode decision factor is compared with the pre-set factor threshold to select the appropriate inference mode, which is divided into the following cases:

[0220] Case 1: If the mode decision factor is greater than the first threshold, the instruction acceleration mode is activated.

[0221] Specifically, under high load conditions, computing resources are near full capacity, memory access efficiency is relatively low, and bandwidth utilization is relatively high. In these situations, improving computing efficiency and reducing latency and data handling overhead are prioritized. Therefore, by activating instruction acceleration mode, hardware-accelerated instructions (such as the GMMA instruction) are used to directly load contiguously stored data blocks and perform efficient matrix operations, maximizing computing throughput.

[0222] That is to say, in this case, activating instruction acceleration mode can increase the processing speed of compute-intensive tasks, reduce memory access latency, and thus improve overall performance. It is particularly suitable for tasks that require fast response, such as real-time speech recognition, image processing and other scenarios.

[0223] In an embodiment of the present application, an instruction acceleration mode such as the GMMA high-speed mode is suitable for short sequence processing because it can perform efficient matrix multiplication operations by utilizing continuous memory blocks, reducing data handling overhead. For example, for short sequences (less than 512 tokens), the GMMA instruction stream in the Hopper GPU architecture is used for calculations. This method uses continuous memory blocks to perform calculations without zero padding (Zero-Padding Free), which can improve computing efficiency and reduce unnecessary memory usage. For example, the instruction stream throughput efficiency of short sequence reasoning is increased by 2.3 times, completely releasing the hardware potential of the Matrix Multiply Accelerator in the Hopper architecture.

[0224] Case 2: If the mode decision factor is less than the second threshold, the paging cache mode is activated; wherein, the paging cache mode is used to divide the input sequence of this round into multiple short sequences according to a preset length, and the inference dependency data corresponding to each short sequence is continuously stored using different paging units.

[0225] Specifically, under low load conditions, computing resource utilization is low, and there may be a large amount of idle resources. In these situations, optimizing data management and transmission to improve overall efficiency is a priority. The paged cache mode divides the input sequence into multiple short sequences and allocates a different paging unit for continuous storage of each short sequence, effectively reducing memory fragmentation and data transfer overhead.

[0226] In other words, activating paged cache mode in this case can improve the efficiency of memory and video memory usage, reduce waiting time during data transfer, and is particularly suitable for long-sequence inference tasks. Furthermore, the asynchronous transmission mechanism enables overlapping execution of data transfer and computation, further improving system throughput.

[0227] In an embodiment of the present application, the paged cache mode is more suitable for long sequence processing because it can effectively manage memory when processing large amounts of data, avoiding video memory from becoming a bottleneck. Therefore, when it is necessary to process longer sequences or when the GMMA mode is not suitable, the paged cache mode can be used. This method calls the Hopper architecture TMA module to implement asynchronous data transmission between paged memory and shared memory, which can significantly reduce data handling overhead and reduce data transmission costs by 74%, thereby improving overall performance. Simply put, by dividing the KV cache into dynamic memory units of preset lengths, such as 64 tokens / blocks, according to the division of the input sequence, combined with the TMA-driven asynchronous prefetch strategy, the video memory fragmentation rate is reduced to 1.2%. At the same time, the data handling overhead is reduced by 74% through the memory-video memory asynchronous transmission pipeline, and an end-to-end efficient long sequence processing system is jointly constructed.

[0228] Specifically, the management of inference-dependent data in the paging cache mode is achieved through partitioning. During each round of inference, the input sequence of this round is divided into multiple short sequences according to a preset length (such as 64 tokens / block). The inference-dependent data corresponding to each short sequence is assigned to different paging units to ensure that the data in each paging unit is stored continuously. Among them, the inference-dependent data in each paging unit is stored continuously, which can maximize the use of memory bandwidth and reduce memory access latency. Each paging unit can be stored continuously or discontinuously, and this article does not make specific restrictions. In this way, even if the overall input sequence is long, the problem of video memory fragmentation can be effectively avoided and the data handling overhead can be reduced.

[0229] Optionally, in the case of variable-length input sequences, the input sequence for each round of reasoning is obtained by concatenating the latest reasoning results based on the previous round, as shown below: Figure 7As shown, when dividing the input sequence, an incremental approach can also be used. That is, compared to the previous round, only the inference dependency data corresponding to the latest inference result needs to be divided into an existing paging unit or a new paging unit. For example, assuming that the previous round processed 128 tokens and generated a new result of 32 tokens, only the data of these 32 tokens needs to be appended to the existing paging unit or formed into a new unit to avoid re-dividing the entire sequence, thereby improving efficiency and reducing overhead. This approach is particularly suitable for dynamic adjustment and optimization of long sequence processing tasks.

[0230] When loading inference-dependent data, the TMA-driven asynchronous prefetching strategy allows the system to pre-load the next paging unit's data into shared memory without waiting for the current computation to complete. This asynchronous operation allows data movement and computation to overlap, reducing GPU wait time and improving overall efficiency.

[0231] Optionally, this application also supports dynamic adjustment of the KV cache to memory units of a preset length. Combined with the asynchronous prefetching strategy of TMA, the system can flexibly adjust the data management method according to actual needs to ensure efficient data access and processing.

[0232] Through the above approach, paged cache mode effectively manages memory when processing long-sequence tasks, preventing video memory from becoming a bottleneck. Especially when processing large amounts of data or when GMMA mode is unsuitable, using paged cache mode combined with TMA's asynchronous prefetching strategy can significantly improve overall system performance. This end-to-end, efficient long-sequence processing system not only improves the efficiency of memory and video memory usage, but also significantly reduces latency during data transmission, creating a more flexible and efficient inference framework.

[0233] Case 3: If the mode decision factor is not less than the second threshold and not greater than the first threshold, the current reasoning mode is maintained.

[0234] Specifically, when the mode decision factor is between two thresholds, it indicates that the system is under moderate load, with neither obvious performance bottlenecks nor excessive idle resources. Therefore, maintaining the current inference mode avoids the additional overhead caused by unnecessary mode switching and maintains system stability and consistency.

[0235] That is to say, in this case, maintaining the current inference mode can ensure that the system runs smoothly under medium load, reduce performance fluctuations caused by frequent mode switching, and provide a balanced option that neither over-optimizes nor wastes resources, which is suitable for most common application scenarios.

[0236] In general, by reasonably setting the first threshold and the second threshold and dynamically adjusting the inference mode according to the mode decision factor, the system performance and resource utilization efficiency can be optimized under different load conditions.

[0237] In the embodiment of the present application, the second threshold is smaller than the first threshold to ensure that there is a sufficient gap between the two to avoid frequent switching of the inference mode. The following briefly lists several feasible threshold setting methods:

[0238] As for the first threshold, the threshold is used to distinguish whether the system is in a high-load state, that is, a situation where the computing resource utilization is high and may become a performance bottleneck.

[0239] In an optional embodiment, the first threshold may be set based on historical data or benchmark test results. For example, under a typical workload, a value of the mode decision factor is found through experiments as the first threshold, above which the system performance begins to significantly degrade.

[0240] In another optional embodiment, the first threshold can be set according to different application scenarios and hardware configurations. For example, a higher first threshold can be set for a high-performance computing environment, while a lower first threshold may be required for a resource-constrained edge device.

[0241] As for the second threshold, the threshold is used to identify whether the system is in a low-load state or resource underutilization state, so as to activate a more efficient reasoning mode.

[0242] In an alternative embodiment, the second threshold can also be set based on historical data and benchmark results, but with a focus on the system's performance under light load. For example, a mode decision factor value is typically selected as the second threshold. A value below this value indicates low system resource utilization, suggesting room for optimization.

[0243] In another optional implementation, the second threshold may also be set based on different application scenarios and hardware configurations. For example, a lower second threshold may be set for a high-performance computing environment, while a higher second threshold may be required for a resource-constrained edge device.

[0244] Specifically, in a high-performance computing environment, due to the abundance of resources and high task complexity, the demand for computing and memory is relatively high. Therefore, a higher first threshold can be set to activate the instruction acceleration mode only when the system is close to full load, ensuring that the mode switch is only performed when there is a real need to improve computing efficiency. For example:

[0245] In a high-performance computing environment, for large model training tasks in data centers or cloud services, the first threshold can be set to 0.85. This means that when the mode decision factor exceeds 0.85, the system will consider that it is currently in a high-load state and activate the instruction acceleration mode to increase the processing speed.

[0246] For resource-constrained edge devices, to prevent premature high load, the first threshold can be set to 0.70. This allows computing efficiency to be optimized before the system is fully saturated, avoiding performance bottlenecks.

[0247] Similarly, the second threshold value should also be different in different application scenarios and hardware configurations. For example:

[0248] In a high-performance computing environment with ample resources, even under relatively low load, some idle resources may exist. To fully utilize these resources, the second threshold can be set lower, such as 0.30. This means that only when the mode decision factor is below 0.30 will the system be considered under low load and the paging cache inference mode or other lightweight modes activated.

[0249] For resource-constrained edge devices, any underutilized resources can lead to inefficiency. Therefore, the second threshold can be set slightly higher, such as 0.45. This allows the corresponding optimization mode to be activated when the system load decreases slightly, maximizing the utilization of existing resources.

[0250] By flexibly adjusting these two thresholds based on different application scenarios and hardware configurations, the system can better adapt to changes in actual workloads, thereby achieving better performance and resource utilization. This dynamic adjustment mechanism not only improves the system's flexibility but also enhances its ability to handle a variety of complex tasks.

[0251] Specifically, in actual applications, the first threshold or the second threshold can be automatically adjusted through automation tools according to changes in actual workload, and continuous optimization can be carried out based on real-time monitoring data to ensure that the system is always in the optimal operating state.

[0252] Optionally, these two thresholds can also be calibrated through offline training. For example, based on a large amount of historical operating data, the system performance under different workloads is simulated, and the optimal threshold range is found by analyzing the relationship between the mode decision factor and the inference mode switching. Specifically, simulation tools or experimental platforms can be used to test key indicators such as the system's response speed and resource utilization in a variety of typical task scenarios, and the thresholds can be continuously adjusted in combination with optimization algorithms (such as grid search or genetic algorithms) until the optimal performance balance point is reached. This offline calibration method can ensure that the threshold settings are more in line with actual application scenarios and improve the dynamic adaptability of the system.

[0253] Optionally, the above-described method for switching the inference mode can be implemented by a mode decider. Taking the inference mode including instruction acceleration mode and paging cache mode as an example, the mode decider can dynamically select whether to use instruction acceleration mode or paging cache mode based on real-time hardware signals, such as SM idle cycles, L2 Cache hit rate, and DRAM bandwidth utilization.

[0254] like Figure 4 As shown in FIG, it is a schematic diagram of the design concept of a mode decision maker in an embodiment of the present application. The mode decision maker collects hardware status information of the Hopper GPU in real time, establishes a dynamic decision model, and realizes automatic switching between instruction acceleration mode and page cache mode. Let the mode decision factor be D, the first threshold be θ1, and the second threshold be θ2. The decision basis includes:

[0255] Case 1: D>θ1, activate instruction acceleration mode to improve computing efficiency and reduce memory access latency.

[0256] Case 2: D<θ2, activate the page cache mode to optimize data management and transmission, reduce fragmentation and data handling overhead.

[0257] Case 3: θ2≤D≤θ1, maintain the current inference mode to maintain system stability and avoid unnecessary switching overhead.

[0258] The method proposed in the above implementation uses hardware status information to determine the mode decision factor and dynamically adjusts the inference mode based on the mode decision factor. This not only improves the flexibility and accuracy of inference mode selection, but also ensures that the system can maintain optimal performance under different hardware conditions. Compared with static configuration, this approach can significantly reduce waiting time and resource waste, improving overall computing efficiency.

[0259] S22: When the inference mode is the instruction acceleration mode, at least one round of inference operation is performed, and each round of inference operation includes S221 to S223.

[0260] In real-world applications, model input sequence lengths often vary significantly, and the number of tokens required for different tasks can vary by several or even dozens of times. Using a unified processing path can introduce performance bottlenecks or waste resources in some scenarios.

[0261] Therefore, before performing at least one round of reasoning operations, in order to achieve dynamic adaptation to different input sequence lengths, in an optional implementation, the length of the initial input sequence corresponding to the task to be reasoned can be obtained for judgment.

[0262] Optionally, before performing at least one round of reasoning operations, an initial input sequence corresponding to the task to be reasoned is obtained for the task to be reasoned; the length of the initial input sequence is compared with a preset length threshold to determine whether the length of the initial input sequence is lower than the preset length threshold.

[0263] In an embodiment of the present application, it is first necessary to obtain the initial input sequence corresponding to the task to be inferred, and compare its length with a preset length threshold to determine whether the length of the initial input sequence is lower than the preset length threshold, that is, to complete short sequence recognition.

[0264] The preset length threshold can be flexibly set based on factors such as the model deployment scenario, hardware platform performance, memory resource constraints, and statistical characteristics of typical input data. In an optional embodiment, if the deployment platform is a resource-constrained edge device, the threshold can be set to 256; if deployed in a high-performance GPU server environment and the model input is generally long, it can be set to 512 or higher. For example, the threshold is set to 512 to adapt to the token number distribution characteristics of the generated class model in a typical NLP task.

[0265] It should be noted that the above threshold settings are only examples and the specific values are not limited. In actual applications, they can also be customized or dynamically adjusted according to actual business scenarios.

[0266] After the threshold setting is completed, the initial input sequence of the task to be inferred is obtained, and the length of the initial input sequence can be counted. The statistical result is then compared with the preset length threshold. The input sequence length is branched and judged by the preset threshold (L=512) to determine whether the current input sequence belongs to the "below the threshold" situation. If it is lower than the preset length threshold (assuming it is 512), it indicates that the input sequence is a short sequence. Short sequences (<512 tokens) are prone to memory discontinuity due to insufficient sequence length. Invalid data (Zero-Padding) needs to be filled to meet the matrix alignment requirements (such as filling to a multiple of 512), resulting in 35%-50% computing power waste, that is, the memory access mode has a significant impact on performance and requires targeted optimization. Therefore, when determining that the input sequence is a short sequence, the application performs the next step of storage continuity check and subsequent processing, which can effectively improve the reasoning speed. Otherwise, a conventional processing path is adopted, such as a block scheduling strategy or other strategy, which is not specifically limited in this article.

[0267] In addition, the judgment operation in this embodiment can be performed on the initial input sequence, that is, the first set of inputs before the inference task is started. However, in other optional expansion schemes, the judgment logic can also be configured to have a certain degree of dynamicity. For example, it can be designed to perform a judgment once for each round of inference, which is particularly suitable for autoregressive generation tasks where the length of the input sequence may change. It can also be set to perform a judgment every several rounds, such as every 5 or 10 rounds, reducing the judgment frequency to save system overhead. It can also be combined with historical inference data to perform sliding window statistics to improve the accuracy and stability of the judgment. This application does not specifically limit the timing of input sequence comparison.

[0268] It's also worth noting that a multi-threshold strategy based on model type can be introduced. Different preset length thresholds can be configured for different models (such as text generation models, machine translation models, and question-answering systems) to accommodate the natural differences in input token size for different types of tasks. A personalized inference path scheduling mechanism can also be implemented based on user preferences or scenario configurations.

[0269] Through the above implementation, a one-time, lightweight judgment logic can be completed based on the length of the initial input sequence before the inference task is executed. This judgment not only simplifies the process of obtaining path selection conditions, but also serves as an upstream control signal for modules such as memory scheduling, cache strategy, and inference accuracy setting, thereby building a more adaptive and refined inference scheduling mechanism. Especially in scenarios where multiple models are deployed concurrently and the input structure varies greatly, it can effectively improve the system's adaptability and computing resource utilization efficiency, and enhance the performance stability and response sensitivity of the overall inference service.

[0270] S221: Perform storage continuity check in computing resources for the inference dependency data related to the input sequence of this round; the inference dependency data is generated and stored through the inference operation of the previous round.

[0271] In an embodiment of the present application, reasoning-dependent data refers to the intermediate result data generated and stored by the reasoning operation of the previous round during the reasoning process. In the first round of reasoning, these data can be preset values or initialization parameters to ensure that the algorithm can start execution from a certain state. This document does not make specific restrictions on this. These data typically include the context information and calculation results required by the model when processing the current input sequence. For example, in a model based on the attention mechanism, the reasoning-dependent data can be specifically expressed as a KV vector. These KV vectors are generated by encoding the input sequence and used for subsequent attention calculations to capture the correlation between each position in the input sequence.

[0272] The importance of data-dependent reasoning is particularly prominent in multi-round reasoning tasks. For example, in natural language processing tasks, reasoning over long texts or in dialogue systems often requires the step-by-step generation of output sequences. Each round of reasoning requires computing the output of the current step based on the previously generated key-value vectors. Therefore, these key-value vectors effectively constitute the context of the reasoning process, directly influencing the model's understanding of the current input and the accuracy of its output.

[0273] Of course, in addition to KV vectors, inference-dependent data can also be other types of intermediate result data, such as the embedding vectors of each token in the input sequence, the probability distribution of historical tokens, intermediate layer activation values, hidden layer states, encoder outputs, decoder state information, model parameters, etc. Among them, the probability distribution of historical tokens reflects the probability distribution of previously generated tokens in the vocabulary, which helps the model predict the next most likely token and enhance the coherence and accuracy of the output sequence; the intermediate layer activation values represent the response of the network's intermediate layer to the input data. These values contain feature information after multi-layer conversion, which helps the model perform in-depth feature extraction and understanding; the hidden layer state contains the feature representation after processing at each layer, which helps capture deeper semantic information; the encoder output provides the overall contextual information of the input sequence, and the decoder state information plays an important role in generating a coherent output sequence; model parameters include but are not limited to the weights and biases of each network layer in the target model.

[0274] By effectively managing and utilizing this rich set of intermediate data, the efficiency and accuracy of model reasoning can be significantly improved. This article does not specifically limit this approach, but rather emphasizes the diversity and importance of data-driven reasoning. This flexible design makes this method applicable to a variety of application scenarios and model architectures.

[0275] Ensuring efficient storage and fast access to inference-dependent data is key to achieving efficient inference, especially when processing long sequences or complex models. If the storage of inference-dependent data is discontinuous, it may lead to inefficient memory access, increased latency, and reduced overall performance. For example, in high-performance computing resources such as GPUs, non-continuously stored data will significantly increase video memory access time, thereby slowing down inference speed. Therefore, the storage continuity check in step S221 is intended to optimize the storage layout of inference-dependent data to make it more suitable for efficient loading and processing of hardware-accelerated instructions.

[0276] Therefore, performing storage continuity checks in S221, especially for inference-dependent data in the KV cache, can effectively improve inference efficiency. This optimization is particularly important when processing long sequence tasks.

[0277] The following is a simple example of how to check storage continuity:

[0278] First, you need to determine the scope and type of inference-dependent data to be checked. This includes identifying which data is critical inference-dependent data, such as key-value vectors, embedding vectors, and the probability distribution of historical tokens, and recording their starting addresses and sizes in computing resources such as the GPU's global memory (GMEM). Global memory is the largest but slowest-access memory type on the GPU, primarily used to store input data, output data, and intermediate results. It boasts a large storage capacity that can accommodate entire datasets, making it suitable for large-scale data processing tasks.

[0279] Next, specific methods or tools are used to check the storage continuity of these inference-dependent data. Specifically, this check can be performed by writing specialized Compute Unified Device Architecture (CUDA) kernel functions. These kernel functions traverse the data structure and verify whether the addresses of adjacent elements in GMEM are continuous. This is to assess whether the layout of the inference-dependent data in GMEM is row-major or column-major continuous.

[0280] For example, an optional inspection process is: first, traverse each memory block related to the inference dependency data and check its physical storage location in GMEM. Next, confirm whether these memory blocks exist in the form of continuous blocks. If the inference dependency data spans multiple discontinuous memory areas, it is marked as a discontinuous state. The specific operation can be achieved by comparing the memory address difference of adjacent elements. If the difference does not meet expectations (that is, it is not a fixed step size), it is judged to be discontinuous. Then, the continuity status of each data block is marked, and the relevant information is recorded for subsequent processing.

[0281] Of course, in addition to using CUDA kernel functions, you can also use the GPU's built-in memory inspection tools or APIs to evaluate storage continuity. For example, Nsight Compute and Nsight Systems provide detailed memory analysis capabilities that can automatically detect data layout and access patterns. In addition, developers can use CUDA's memory management API (such as cudaMemGetInfo) in combination with custom scripts to count and analyze the allocation of data blocks in GMEM and their continuity, thereby simplifying the inspection process and improving accuracy; these methods provide multiple ways to optimize data loading and processing.

[0282] It should be noted that the storage continuity checking methods listed above are just simple examples. In addition, other checking methods are also applicable to the embodiments of the present application and will not be described in detail here.

[0283] After the storage continuity check is completed, the next step can be determined based on the result of the storage continuity check, such as S222 to S223 (non-continuous) or S224 (continuous). S222 to S223 are first described below:

[0284] S222: When the storage location of the inference-dependent data is in a discontinuous state in the computing resources, the storage location is adjusted to a continuous state by rearranging the inference-dependent data.

[0285] In an embodiment of the present application, when it is detected that the inference-dependent data is stored in a non-contiguous location, memory reorganization technology can be used to rearrange the data so that the data is stored contiguously, thereby reducing subsequent read time. During the reorganization, the data can be reorganized according to the actual usage order of the data, and memory copy or move operations can be used to concentrate the scattered data into a contiguous memory area.

[0286] Specifically, memory reorganization technology is an optimization strategy that improves data access efficiency by breaking down data blocks scattered across discontinuous memory addresses into multiple sub-blocks, rearranging them in a specific order, and concentrating them in contiguous memory areas. This technology reduces the high-latency memory access caused by discontinuous storage, fully utilizing the memory bandwidth and cache mechanisms of computing resources, thereby accelerating data reading and processing. This approach not only reduces unnecessary memory access overhead but also enhances system responsiveness and overall efficiency.

[0287] To further optimize the storage layout of inference-dependent data and ensure its continuity across computing resources, thereby reducing subsequent read time and improving access efficiency, this application proposes a specific rearrangement method. This method not only reorganizes and moves non-contiguously stored data, but also defines in detail how to achieve continuous storage by decomposing memory blocks into multiple sub-blocks and rearranging these sub-blocks according to specific rules. The specific implementation is as follows:

[0288] When the storage location of the inference-dependent data is in a discontinuous state in the computing resources, an optional implementation method is to implement S222 according to the following flowchart, including the following steps S2221 to S2222 ( Figure 2 not shown):

[0289] S2221: For each memory block storing inference-dependent data, decompose each memory block into multiple sub-blocks according to a preset data block size.

[0290] The size of each sub-block is an integer multiple of a preset data block size; the preset data block size is determined based on the memory access rules of the computing resource. In other words, the selection of the preset data block size in the embodiments of the present application takes into account the memory access characteristics and rules of the computing resource (such as a GPU) to optimize performance. Selecting an appropriate preset data block size can help better utilize hardware characteristics, reduce memory access latency, and improve bandwidth utilization.

[0291] In one optional implementation, the default data block size is 64 bytes, based on the memory alignment mechanism. For example, if the GPU's memory system requires data to be aligned in 64-byte units, selecting 64 bytes as the default data block size ensures that each memory access is aligned, avoiding the additional overhead caused by unaligned accesses. With this configuration, each memory request can efficiently transmit a complete cache line, reducing partial load and store operations and improving memory access efficiency.

[0292] In another optional embodiment, according to the memory alignment mechanism, the preset data block size is 128 bytes or larger. For example, some high-performance GPUs may support larger alignment units, such as 128 bytes or 256 bytes. Selecting these larger preset data block sizes can further optimize the memory access pattern in specific application scenarios. Under such a configuration, for applications that need to process large amounts of data (such as large-scale matrix operations or deep learning model inference), using larger data blocks can reduce the number of memory accesses and increase the amount of data transmitted each time, thereby more effectively utilizing the high bandwidth characteristics of the GPU.

[0293] In another optional implementation, hardware adaptation can also be considered, such as based on the L1 / L2 cache line size. Assuming the GPU's L1 or L2 cache line size is 128 bytes, selecting 128 bytes as the default data block size ensures that the data block fits completely within the cache line, maximizing the cache hit rate. This reduces cache misses, speeds up data access, and ultimately improves overall computing performance.

[0294] In another optional implementation, memory transaction granularity can also be considered. For example, some GPU architectures may have a fixed memory transaction granularity, for example, each memory transaction processes a fixed amount of data (such as 256 bytes). In this case, selecting 256 bytes as the preset data block size allows each memory transaction to fully utilize its capacity. By matching the memory transaction granularity, unnecessary small transaction overhead can be avoided, memory bandwidth usage can be optimized, and data transmission efficiency can be improved.

[0295] In summary, to achieve optimal performance, an appropriate data block size can be selected based on the specific memory access characteristics and rules of the target hardware, thereby optimizing data loading, cache utilization, and memory bandwidth usage. The above is just a simple example of the specific value of the preset data block size, and this article does not make any specific restrictions on this.

[0296] When implementing S2221, you first need to identify the memory blocks corresponding to the inference-dependent data involved in this round and their storage locations in GMEM. Then, each memory block is decomposed into multiple sub-blocks according to the preset data block size. The size of each sub-block should be an integer multiple of the preset data block size. This design is designed to align with the memory access rules of the computing resources to optimize performance. This design ensures that each sub-block can fully match the hardware's memory transaction granularity or cache line size, reducing the additional overhead caused by unaligned access, maximizing memory bandwidth utilization, and improving data loading and processing efficiency.

[0297] S2222: Rearrange the obtained sub-blocks according to the target continuous storage order to adjust the storage positions to a continuous state.

[0298] In S2222, it is necessary to rearrange the obtained sub-blocks according to the target continuous storage order. This arrangement ensures the connection between each two adjacent sub-blocks, that is, the end position of the previous sub-block is continuous with the starting address of the next sub-block, thereby forming a continuous storage structure. Among them, the target continuous storage order can be determined based on the actual use order of the data and / or the memory access characteristics of the computing resources. Usually, the sub-blocks are arranged according to the time sequence or logical order in which the data is accessed during the reasoning process to ensure that the data with high frequency access is aligned first. At the same time, the order is adjusted in combination with the hardware characteristics (such as cache line size, memory transaction granularity) so that adjacent sub-blocks are continuous in physical address, thereby reducing access latency and improving memory utilization, laying the foundation for subsequent efficient computing.

[0299] Optionally, when performing the rearrangement, memory management APIs (such as CUDA's cudaMemcpy or cudaMemset) or custom memory copy / move functions can be used to perform the actual data rearrangement operation. For example, a new contiguous memory area is allocated and then each sub-block is copied into the area in the target order. During this process, asynchronous memory copying and stream mechanisms can be used to further improve efficiency, allowing computation and data relocation to proceed in parallel, reducing overall processing time, ensuring that data is reorganized in an optimal manner, and improving subsequent access speed and performance.

[0300] In another optional embodiment, S222 (including S2221-S2222) can be implemented using a TMA memory reordering mechanism. In this case, the TMA module is used to optimize the non-contiguous memory layout, thereby reducing redundant memory operations and improving performance. The TMA can help reorganize data structures to make them more suitable for subsequent calculations, avoiding inefficient memory access caused by non-contiguous memory access during direct calculations.

[0301] Taking a preset data block size of 64 bytes as an example, if data is not aligned on 64-byte boundaries, each access must span multiple cache lines, resulting in some data being wasted. 64-byte memory reorganization technology is designed to fully utilize the GPU architecture's L2 cache and memory bandwidth, avoiding performance bottlenecks caused by non-contiguous memory access. In the Hopper architecture, the L2 cache line size is 64 bytes, so the data layout in memory needs to be optimized based on this cache line size to achieve more efficient memory access.

[0302] In this embodiment, 64-byte alignment is performed on non-contiguous memory layouts to ensure that each memory access perfectly matches the GPU's memory access pattern, thereby maximizing cache utilization and computing efficiency. Rearranging data in memory using the TMA module avoids the fragmentation problem of traditional memory allocation and improves the locality of memory access.

[0303] Specifically, when it detects that data storage is non-contiguous, the TMA module automatically rearranges the related data blocks to meet the 64-byte alignment requirement, thereby improving the efficiency of subsequent accesses. The detailed steps are described as follows:

[0304] First, analyze the memory layout of the inference-dependent data for this round to determine whether it conforms to the optimal GPU memory access pattern. For example, by writing a dedicated CUDA kernel or using existing tool libraries, check whether the inference-dependent data (such as an n×m matrix, denoted as matrix D) is aligned on 64-byte boundaries. If the data is not aligned to 64 bytes, it needs to be re-arranged.

[0305] When rearranging, the TMA module is called, and the TMA module reorganizes non-contiguous memory blocks into 64-byte aligned blocks through hardware instructions (that is, in this article, each memory block is decomposed into multiple sub-blocks according to the preset data block size).

[0306] For example, first calculate the minimum alignment unit size for each element in matrix D, where each element occupies b bytes of data, so that it is an integer multiple of 64 bytes. If b = 8 bytes, the alignment unit size is 64 / gcd(64,b) = 8, meaning each alignment unit contains 8 elements. Next, split the original matrix D into multiple small blocks (i.e., sub-blocks) according to the calculated alignment unit size, and use the TMA module to rearrange these small blocks to ensure that they comply with the 64-byte alignment rule. The TMA packages the data blocks into 64-byte units and places them in new contiguous memory locations.

[0307] After completing the above steps, a new matrix, Aligned(D), is generated. Each data block (i.e., sub-block) in this matrix is arranged in a 64-byte alignment. This process not only improves data transmission efficiency but also reduces memory bandwidth usage while ensuring that each data block can be efficiently accessed within a single GPU cache line, significantly improving subsequent computational performance.

[0308] The above process can be expressed by the following algorithm:

[0309] Assume that the memory data D (i.e., the data to be inferred that needs to be rearranged) is an n×m matrix, where each element occupies b bytes. If b is not a multiple of 64, memory rearrangement is required. The algorithm steps can be expressed as the following formula 5:

[0310] Aligned(D)=TMA(D,64-byte block alignment)(Formula 5);

[0311] Aligned(D) represents the new matrix after 64-byte memory rearrangement of matrix D. TMA (64-byteblock alignment) indicates that the algorithm splits and rearranges the input matrix D according to the 64-byte alignment rule through TMA.

[0312] For example, the original memory block is split into multiple sub-blocks, each with a size of 64 bytes; if a sub-block is less than 64 bytes, data is extracted from subsequent memory blocks to fill it; then, the split sub-blocks are rearranged in sequence to ensure that the new memory addresses are continuous, and the TMA module is used to place the reorganized data into the new continuous memory area to form the final continuous memory block.

[0313] It should be noted that the remaining portion of less than 64 bytes can be processed as a separate sub-block or combined with the previous adjacent sub-block to form a new sub-block that is an integer multiple of 64 bytes. This approach optimizes memory layout, reduces fragmentation, and improves data access efficiency.

[0314] like Figure 5 As shown in FIG, it is a schematic diagram of a memory reorganization in an embodiment of the present application. Assume that the reasoning-dependent data of this round passes through 3 memory blocks, such as Figure 5 Memory blocks 1, 2, and 3 in the memory block are stored in a distributed manner in global memory. Using the above method, sub-blocks 1, 2, and 3 can be reorganized. Each sub-block can be considered a new memory block. These sub-blocks are stored contiguously, with no restrictions on their storage locations. For example, memory reorganization 1 is stored in one contiguous manner, while memory reorganization 2 is stored in another manner.

[0315] After data reorganization, each memory block is accessed in a manner consistent with the GPU's cache line alignment. Each cache line (64 bytes) is fully utilized, effectively avoiding conflicts and wasteful memory accesses. Since 64 bytes is the size of a GPU's L2 cache line, memory accesses can fully utilize the L2 cache, improving cache hit rates. Furthermore, the reordered data access pattern conforms to the Hopper architecture's optimal memory access strategy, maximizing the utilization of parallel computing units, reducing memory bottlenecks, and significantly improving performance.

[0316] Optionally, TMA can dynamically adjust the data storage method based on actual needs to ensure that data is distributed in memory in an optimal manner, avoiding additional overhead caused by fragmentation.

[0317] In the above implementation, a sophisticated memory management strategy is used to decompose non-contiguously stored inference-dependent data into multiple sub-blocks and rearrange them in a specific order, achieving continuous data storage. This method significantly reduces latency and memory access frequency during subsequent data reads, improving data loading speed and processing efficiency. This can effectively reduce memory access bottlenecks and enhance overall system performance, particularly when processing large-scale datasets or multi-round inference tasks, significantly improving the speed and accuracy of model inference.

[0318] S223: Calling the specified hardware acceleration instruction to load the adjusted inference dependency data, and after determining the result of this round of inference based on the loading result, generating the input sequence of the next round according to the result of this round of inference, until a preset stop condition is reached.

[0319] In an embodiment of the present application, hardware acceleration instructions refer to a specific instruction set provided by a computing resource (such as a GPU, TPU, or other dedicated accelerator) for efficiently loading and processing optimized reasoning-dependent data. These instructions can make full use of the parallel computing capabilities and high-bandwidth memory access characteristics of the hardware, thereby significantly improving data loading speed and computing efficiency. By calling these instructions, the matrix operations, attention mechanism calculations, or other key operations required in the reasoning process can be quickly completed, and then the current round of reasoning output is determined based on the loading results, and the next round of input sequence is generated until the preset stop condition is met.

[0320] Taking the computing resource GPU as an example, the GPU can contain specialized hardware acceleration units, such as TensorCores or Matrix Multiply Accelerator (MMA), which can significantly accelerate specific types of calculations, such as matrix multiplication, matrix addition, etc.

[0321] In the GPU, hardware acceleration instructions include but are not limited to: Grouped Matrix Multiply Accumulate (GMMA) instructions, Warpgroup Matrix Multiply Accumulate (WGMMA) instructions, Tensor Core instructions, Matrix-Fused Multiply-Add (MFMA) instructions, etc.

[0322] Specifically, GMMA and WGMMA are specialized instruction sets for matrix multiplication and accumulation operations in GPU architectures (such as the Hopper GPU). They are particularly suitable for large-scale matrix operations, such as the attention mechanism in convolutional neural networks (CNNs), RNNs, and Transformer models in deep learning. These instructions can efficiently process matrices of different sizes and shapes, making them particularly suitable for the complex computational requirements of deep learning tasks.

[0323] Tensor Core is a hardware unit in the GPU specifically designed to accelerate matrix operations. It supports mixed-precision calculations (such as FP16 / FP32) and is mainly used for matrix multiplication and convolution operations in deep learning tasks. Compared with traditional computing units, tensor cores provide higher computing density and lower energy consumption, making them very suitable for tasks that require a large number of matrix operations. For example, some Tensor Core technologies provide efficient support for tensor operations, including processing matrix operations with automatic zero padding, so that non-aligned or different-sized input data can be directly processed when performing matrix multiplication without manual padding.

[0324] The MFMA instruction is optimized for matrix operations. It can perform multiplication and addition in a single operation, significantly improving the speed and efficiency of matrix processing in fields such as deep learning and scientific computing.

[0325] Optionally, instruction parameters can be automatically adapted to hardware characteristics. Specifically, when executing the aforementioned hardware-accelerated instructions, parameters such as matrix size, data type, degree of parallelism, and memory layout can be intelligently adjusted to match the hardware characteristics of the specific GPU. For example, the optimal thread block size, optimal data loading mode, and computational pipeline configuration for different GPUs can all be automatically optimized; this document does not specify these requirements.

[0326] In the embodiments of the present application, the designated hardware accelerated instruction stream can be directly used to perform efficient zero-fill free calculations. For example, using Hopper's WGMMA instruction, the input sequence is divided into M×K matrix blocks and directly loaded into the Tensor Core for calculation, eliminating the waste of computing power caused by traditional fill.

[0327] Taking the text generation scenario as an example, in this scenario, the specific details of calculating the reasoning results of this round in S223 are as follows:

[0328] In text generation tasks, the input sequence and / or related inference dependency data are first divided into multiple M×K matrix blocks. For example, M is set to 512 tokens (it can also be other values, which are not specified in this article), which means that each matrix block contains information about 512 tokens. The K here (not the K in the KV vector, but the matrix dimension K) depends on the specific configuration of the model and the need for contextual information, which is not specified in this article. This partitioning method allows the model to process longer sequences or more contextual information while maintaining efficient data loading and computation.

[0329] In traditional implementations, in order to adapt the input sequence to a fixed-size matrix operation, it is usually necessary to pad the shorter sequences with zeros. This not only increases memory usage but also leads to unnecessary computational overhead because these padded zero values also need to be involved in the calculation.

[0330] In the embodiments of this application, the WGMMA instruction can be used to directly load these M×K matrix blocks into the Tensor Core for calculation without the need for traditional zero padding. This means that even if the actual length of the input sequence is less than the maximum supported length, there is no need to pad with zeros to achieve a fixed matrix size.

[0331] Among them, loading onto the tensor core for calculation can specifically refer to attention calculation. Taking the target model as LLM as an example, LLM uses the attention mechanism when reasoning token by token. Figure 6 As shown in FIG, it is a calculation logic diagram of an attention mechanism in an embodiment of the present application. In the attention mechanism, for each word in the input sequence, the model calculates a query vector (query), a key vector (key), and a value vector (value). These vectors are obtained by learning the weight matrix (such as Figure 6 The Wq, Wk, and Wc in the image are obtained by linearly transforming the input word. In the case of KV caching, the key vector (key) and the value vector (value) can be directly loaded as inference dependency data.

[0332] Assume that the input sequence of the large language model is X = [x1, x2, ..., xn]. For word units xi and xj, the large language model can calculate an attention score, which is used to represent the degree of attention of word unit xi to word unit xj, that is, the degree of attention. The attention score is obtained by performing a dot product operation on the query vector and the key vector, and then dividing it by a scaling factor (usually the square root of the key vector dimension). The calculation formula of the attention score is as follows:

[0333]

[0334] Among them, score(Qi, Kj) is the attention score of xi to xj, Qi is the query vector corresponding to xi, Kj is the key vector corresponding to xj, T represents the transpose operation, dk is the dimension of the key vector Kj, i and j are both positive integers less than or equal to n, n is a positive integer greater than 1, and i can be equal to j.

[0335] Next, the large language model can call the softmax function to convert the attention score to a value between 0 and 1, and the sum of all attention scores is 1 (that is, the attention score is normalized), and then the attention weight is obtained. The calculation formula of the attention weight is as follows:

[0336] wij=softmax(score(Qi,Kj))(Formula 7);

[0337] Finally, the large language model can multiply the value vector of each word unit with its corresponding attention weight, and then sum up the multiplication results to get the final output:

[0338]

[0339] Among them, Vj is the value vector corresponding to xj, and the final output Zi is the attention calculation result. The attention calculation result is also a vector representation, which covers the context information. Based on the output attention calculation result, the next output token of the large language model can be predicted.

[0340] like Figure 6 As shown in , the dimensions of Q, K, and V vectors are all 2×3, and the dimension of KT vector is 3×2. The result calculated by formula 6 is a 2×2 vector, and the result of Z calculated by formula 8 is a 2×3 vector. Among them, Figure 6The calculation logic is simply illustrated. For the specific calculation method, please refer to the above embodiment, and the repeated parts will not be repeated.

[0341] The WGMMA instruction avoids this waste of computational power by dynamically adjusting the matrix size and automatically handling irregular inputs. This allows the model to focus only on the actual useful data, improving computational efficiency and resource utilization.

[0342] In the above process, reasoning-dependent data (such as KV vectors, hidden layer states, etc.) play a key role. They contain the reasoning results of the previous round and provide the context information required for the current reasoning step. Before using the WGMMA instruction, you need to ensure that these reasoning-dependent data have been properly organized and arranged (for example, 64-byte alignment through the TMA module, or the memory itself is continuous) so that they can be efficiently loaded into the tensor core. Good data layout can reduce memory access latency and increase computing speed. In addition, this continuity also helps to reduce bandwidth consumption during data transmission, thereby improving the overall reasoning efficiency.

[0343] In summary, by utilizing the WGMMA instructions of the Hopper architecture and combining them with an optimized inference-dependent data management strategy, an efficient, zero-padding-free computational process can be achieved in text generation tasks, significantly improving the model's inference performance and resource utilization.

[0344] It should be noted that the hardware acceleration instructions listed above are just simple examples. In addition, other similar hardware acceleration instructions are also applicable to the embodiments of the present application and will not be described in detail here.

[0345] By calling these hardware-accelerated instructions, particularly the GMMA or WGMMA instructions, the complex computational tasks required during inference can be quickly completed. The output of the current round of inference is then determined based on the loaded results. The next round of input sequences is then generated based on the results of the current round of inference until a preset stopping condition is met. This approach not only significantly improves inference efficiency but also provides strong support for processing large datasets and complex models.

[0346] Preset stopping conditions include, but are not limited to, reaching a maximum generation length or generating a specific end marker. The maximum generation length limits the maximum size of the generated sequence, preventing infinite loops. The end marker (such as the EOS marker) indicates a natural stopping point, indicating that the output is complete. When either condition is met, the inference process stops, ensuring that the generated sequence is both complete and meets the expected length requirements, thereby improving output quality and computational efficiency.

[0347] In the embodiment of the present application, in the reasoning process of the current model, the reasoning task is usually analyzed step by step, that is, multiple reasoning operations are performed, and each time the next reasoning is performed based on the token analyzed and reasoned first, and tokens are generated one by one in sequence, and each new token generated depends on the token sequence generated previously.

[0348] Therefore, when generating the next round of input sequence based on the result of this round of reasoning, the result of this round of reasoning can be spliced with the input sequence of this round to obtain the next round of input sequence.

[0349] like Figure 7 As shown, it is a schematic diagram of the transformation of a text generation scenario input sequence in an embodiment of the present application. Assuming that the initial input sequence (i.e., the first inference) is "What does this picture contain?", LLM outputs "contains" when performing the first inference; then, the input sequence during the second inference is "What does this picture contain? Contains", LLM outputs "dog" when performing the second inference; then, the input sequence during the third inference is "What does this picture contain? Contains dog", LLM outputs "and" when performing the third inference; then, the input sequence during the fourth inference is "What does this picture contain? Contains dog and", ..., and so on.

[0350] Based on the above, considering that each round of reasoning generates a corresponding inference result, which serves as input for the next round, the inference dependency data required for subsequent rounds, with the exception of the first (non-initial) round of reasoning, will include both the latest inference dependency data related to the previous round's inference result (i.e., the current latest inference result) and the historical inference dependency data related to historical data. Furthermore, in each round of reasoning, discontinuous inference dependency data is rearranged and made continuous. Therefore, when checking the continuity of inference dependency data, the concept of incrementation can be used as a reference.

[0351] That is, when implementing S221, an optional implementation method is as follows:

[0352] Optionally, the reasoning dependency data includes: the latest reasoning dependency data related to the latest reasoning result in the input sequence of this round, and the historical reasoning dependency data related to other information in the input sequence of this round; therefore, when performing a storage continuity check on the reasoning dependency data related to the input sequence of this round in the computing resources, it is necessary to check the storage continuity of the latest reasoning dependency data and the historical reasoning dependency data in the computing resources.

[0353] Among them, the latest reasoning dependency data contains all information related to the new results generated by the previous round of reasoning operations.

[0354] Taking the inference dependency data as a KV vector as an example, for example, in a text generation task, assuming the current round is N+1, if the previous round (round N) generates a new sequence segment S_N containing 50 tokens, then the KV vectors corresponding to these newly generated tokens are the latest inference dependency data, that is, the KV vector corresponding to S_N is the latest inference dependency data.

[0355] Historical reasoning dependency data refers to all relevant data generated during previous reasoning rounds. This data provides the necessary context for the current round of reasoning. If, before round N of reasoning, sequence segments S_1, S_2, ..., S_{N-1} have been generated, the key-value vectors corresponding to these sequence segments are the historical reasoning dependency data.

[0356] In each round of inference, discontinuous inference dependency data is rearranged to a continuous state to ensure efficient data loading and computation. The specific implementation steps are as follows:

[0357] First, data division and identification are performed. That is, at the beginning of each round of reasoning, the reasoning dependency data needed for this round is first identified, including the latest reasoning dependency data and historical reasoning dependency data.

[0358] For example, during the N+1 round of reasoning, the latest reasoning dependency data generated in the Nth round and all previous historical reasoning dependency data need to be used.

[0359] Since each inference generates new inference-dependent data and adds it to the existing dataset, an incremental approach can be used to perform storage continuity checks. Specifically, only the continuity between the newly added data (i.e., the latest inference-dependent data) and the existing data (i.e., the historical inference-dependent data) needs to be checked. The specific check method can be found in the above-mentioned embodiments and will not be repeated here.

[0360] It should be noted that if the latest reasoning dependency data itself is also discontinuous, then it can also be regarded as a discontinuity between the latest reasoning dependency data and the historical reasoning dependency data.

[0361] An optional implementation is to determine the storage locations of the latest reasoning dependency data and the historical reasoning dependency data in a discontinuous state in the computing resources by any of the following methods:

[0362] (1) When the storage location of the latest reasoning dependency data is in a discontinuous state in the computing resource, it is determined that the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource.

[0363] Specifically, when checking the storage location of the latest reasoning-dependent data, if it is found that this data is not stored continuously in the computing resources (such as GPU global memory), it can be determined that the overall storage status of the latest reasoning-dependent data and the historical reasoning-dependent data is discontinuous. In other words, even if the historical reasoning-dependent data is stored continuously, as long as the latest data block itself is discontinuous, it will affect the overall continuity. In this case, the system needs to perform additional data rearrangement operations to optimize subsequent processing steps to ensure efficient data access and computing performance.

[0364] Assume that the historical reasoning dependency data is located at addresses 0x1000 to 0x103F, and the latest dependency data itself is stored in multiple non-contiguous memory blocks, such as: the first block is 0x1080 to 0x109F, and the second block is 0x10C0 to 0x10DF, which indicates that the latest reasoning dependency data itself is discontinuous.

[0365] (2) When the storage location of the latest reasoning dependency data and the storage location of the historical reasoning dependency data are not adjacent in the computing resources, it is determined that the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resources.

[0366] Specifically, when comparing the storage locations of the latest inference-dependent data with the historical inference-dependent data, if the two are not stored adjacently in the computing resources, that is, there are gaps or other data blocks between them, then it can be determined that the storage locations of these data are non-contiguous. This means that even if the data within each is continuous, due to the non-adjacent storage locations between them, efficient continuous access as a whole cannot be achieved. In this case, data rearrangement is also required to optimize the storage layout to reduce memory fragmentation and improve data loading efficiency.

[0367] Assuming that the historical inference dependency data is located at addresses 0x1000 to 0x103F, and the latest inference dependency data is located at 0x1080 to 0x10BF, there is a gap between the two.

[0368] like Figure 8 As shown, it is a schematic diagram of a discontinuous state in an embodiment of the present application; Figure 8 This article briefly lists three discontinuity states between historical and latest reasoning dependency data. For example, if the historical reasoning dependency data is continuous, discontinuity state 1 indicates that the latest reasoning dependency data is also continuous, but discontinuous with the historical reasoning dependency data. Discontinuity state 2 indicates that the latest reasoning dependency data itself is discontinuous, and discontinuous with the historical reasoning dependency data. Discontinuity state 3 indicates that the latest reasoning dependency data itself is discontinuous, and some of the latest reasoning dependency data is discontinuous with the historical reasoning dependency data. These states are merely examples and are not specifically defined in this article.

[0369] If discontinuity is found, it is necessary to rearrange it through the TMA module or other memory management technology. Thus, when implementing S222, an optional implementation method is as follows:

[0370] When the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resources, the storage locations are adjusted to a continuous state by rearranging the latest reasoning dependency data.

[0371] Specifically, when the latest reasoning dependency data and the historical reasoning dependency data are stored discontinuously within the computing resources, the gap between them and the discontinuous portion within the latest reasoning dependency data are first identified. Next, the latest reasoning dependency data is rearranged using a TMA module or other memory management tool. For example, the data is decomposed into pre-set data blocks and the decomposed sub-blocks are moved from their original locations and merged into the continuous space immediately following the historical reasoning dependency data. For a detailed implementation, please refer to the above embodiment, and any repetitions will not be repeated here.

[0372] It should be noted that when decomposing the latest inference-dependent data, if there is a sub-block in the historical inference-dependent data that does not meet the requirements of the preset data block size (the remaining small portion), the latest inference-dependent data can be decomposed to a portion that can be integrated with the sub-block to form a sub-block that meets the requirements of the preset data block size. This can better optimize memory layout, reduce fragmentation, ensure data storage continuity in computing resources, and thus improve data loading and processing efficiency, maximizing hardware resource utilization. This strategy helps improve overall inference performance and system stability.

[0373] For example, the storage address range of the last sub-block in the historical reasoning dependency data is 0x1100 to 0x110F (only 16 bytes), which does not meet the requirement of an integer multiple of 64 bytes, and 48 bytes are required to reach an integer multiple of 64 bytes. However, the original storage address range of the latest reasoning dependency data is discontinuous, so 48 bytes can be taken from the latest reasoning dependency data and merged with the insufficient sub-blocks in the historical reasoning dependency data. In this way, the sub-blocks in the historical reasoning dependency data that originally did not meet the integer multiple of 64 bytes are completed, ensuring that all sub-blocks meet the preset data block size requirements.

[0374] For example, the 48 bytes of the latest inference dependency data are moved to the 0x1100 to 0x110F area of the historical inference dependency data, expanding it to 64 bytes (from 0x1100 to 0x113F). The remaining latest inference dependency data can be reallocated as a sub-block to the continuous area after 0x113F, that is, 0x1140 to 0x117F (a total of 64 bytes).

[0375] If there are four consecutive 64-byte sub-blocks before the last sub-block in the historical reasoning dependency data, then the total storage range of the historical reasoning dependency data is 0x1000 to 0x110F (a total of 272 bytes, including four consecutive 64-byte sub-blocks and the last 16-byte sub-block). The total storage range of the latest reasoning dependency data after adjustment is 0x1110 to 0x117F. In this way, the historical reasoning dependency data and the latest reasoning dependency data total 384 bytes, and the total storage range is continuous, from 0x1000 to 0x117F.

[0376] After this, the pointers or indexes pointing to these data blocks can be updated to ensure that they point to the new continuous address range, thereby optimizing data access efficiency and computing performance. This process ensures the storage continuity of all inference-dependent data in the computing resources.

[0377] The above implementation reduces memory fragmentation and optimizes data layout by ensuring that the latest and most recent inference-dependent data are stored contiguously within computing resources. This approach not only improves memory access efficiency and computing resource utilization, but also significantly enhances inference performance, reduces computing power waste, and supports processing longer sequences. Ultimately, it enhances system stability and flexibility, enabling the model to more efficiently execute complex inference tasks and improving overall application performance and reliability.

[0378] Based on the above, and taking into account the efficient use of hardware resources, to further optimize the data loading and processing efficiency during the inference process, this application proposes a double buffering mechanism. Double buffers are allocated in shared memory. In this way, when executing S223 or S224, hardware acceleration instructions can be called to alternately load data, ensuring maximum computational density. An optional implementation is as follows:

[0379] Before executing S223 or S224 , a first buffer and a second buffer may be allocated in a shared memory of computing resources occupied by the target model.

[0380] Among them, shared memory (SMEM) is a small-capacity cache located on each SM that can be shared by all threads in the same thread block. Compared with global memory, shared memory has a smaller capacity, usually between tens of KB and hundreds of KB, depending on the hardware model. However, due to its location advantage close to the computing unit, the access speed of shared memory is significantly faster than that of global memory. In addition, developers can flexibly control which data is loaded into shared memory through programming and optimize data access patterns to improve performance. It is particularly suitable for application scenarios that require frequent access to the same data.

[0381] In the process of GPU performing data loading and calculation, global memory and shared memory each play an important role. In the initial stage, data is loaded from global memory to the computing unit for processing, but due to the high latency of global memory access, frequent global memory access may become a performance bottleneck. In order to optimize this process, this application loads data that needs to be used multiple times from global memory to shared memory, so that multiple threads in the same thread block can share these data copies, reducing repeated access to global memory. This method not only reduces the number of global memory accesses, but also takes advantage of the low latency of shared memory. It is particularly suitable for application scenarios such as matrix operations and convolutional neural networks, thereby significantly improving computing efficiency and data processing speed.

[0382] In this embodiment, double buffering requires allocating two memory buffers with identical data layout and size. This pre-allocation of the double buffers provides an efficient temporary data storage area for subsequent data processing, ensuring that different batches of inference-dependent sub-data can be quickly accessed and processed.

[0383] Next, when executing S223 or S224, the hardware acceleration instruction can be called to alternately load different batches of inference-dependent sub-data from the first buffer and the second buffer for calculation and processing to obtain the result of this round of inference.

[0384] The first and second buffers are used to alternately load different batches of inference-dependent sub-data from the inference-dependent data stored in the computing resource's global memory. The dual-buffering mechanism proposed in this application supports cyclically loading different batches of data from the two buffers while simultaneously performing computations on the data in the other buffer.

[0385] Optionally, the two buffers can be organized in a 64-byte (preset data block size) aligned manner to adapt to the memory access requirements of the GMMA instruction; and it is also necessary to ensure that the data layout in the buffer is continuous to avoid performance degradation caused by non-contiguous memory.

[0386] That is, the inference-dependent sub-data of the same batch is also stored contiguously in memory. This is achieved by partitioning the contiguous inference-dependent data required for the current round according to the buffer size of the computing resource. For example, if the current inference round requires processing 128 bytes of data and the buffer size is 64 bytes, the inference-dependent data is divided into two batches and loaded into the first and second buffers respectively. This continuity and partitioning strategy ensures efficient data loading and processing, reduces memory fragmentation, and optimizes access speed.

[0387] Specifically, when these hardware acceleration instructions are called, the system alternately loads different batches of inference-dependent sub-data from the first and second buffers for computational processing. This dual-buffering mechanism allows data to be loaded in one buffer while computation is performed in the other. The two buffers are used alternately, avoiding conflicts during data loading and computation, effectively reducing wait time and increasing data throughput. Ultimately, this efficient data loading and processing method enables accurate inference results for the current round, significantly improving inference efficiency and performance.

[0388] Optionally, the above implementation can be implemented using the TMA module mentioned above. The asynchronous nature of TMA allows for the overlap of computation and data transfer, further reducing redundant memory operations and increasing the effective computation ratio.

[0389] The following first briefly describes the asynchronous characteristics of TMA.

[0390] like Figure 9A As shown, it is a schematic diagram of the first TMA asynchronous characteristic in the embodiment of the present application. Figure 9A The figure shows the parallel and collaborative relationship between TMA threads (TMA Warps), Math Warps 0, and Math Warps 1. TMA threads are primarily responsible for data loading and task distribution (TMA issue), providing data support for subsequent calculations. Math Warps 0 and Math Warps 1 are responsible for specific computational tasks, including matrix operations (WGMMA) and result optimization (Promotion). Specifically, they alternate between WGMMA calculations and Promotion, demonstrating a parallel computing strategy in which data loading works in conjunction with matrix calculations and optimization operations to improve efficiency and performance.

[0391] This architecture uses TMA threads to build a data pipeline, while math threads maintain computational throughput with dual peaks. The two are synchronized at the nanosecond level via a hardware scheduler, forming a software-defined memory hierarchy. After WGMMA completes matrix multiplication and addition, Promotion improves the accuracy of the results, ultimately achieving full computation-communication parallelism, making it suitable for intensive tasks such as deep learning.

[0392] like Figure 9B As shown, it is a schematic diagram of the second TMA asynchronous characteristic in an embodiment of the present application, which illustrates a parallel computing strategy for optimizing the reasoning process of the deep learning model. Figure 9BIn

[15] , computational tasks are divided into multiple warp groups (e.g., Warp Group 0 and Warp Group 1). Each warp group is responsible for performing a series of computational operations, including GEMM0, GEMM1, and Softmax. GEMM0 and GEMM1 represent matrix operations, which are common, computationally intensive operations in deep learning models. They are typically used to multiply weights by input data in forward propagation, known as General Matrix Multiply (GEMM). Softmax is an activation function commonly used in the output layer, converting the model's raw output into a probability distribution to facilitate decision-making in classification tasks.

[0393] from Figure 9B As can be seen in the figure, there is some overlap in computation between different warp groups. For example, both warp group 0 and warp group 1 include GEMM0 and Softmax operations. This design aims to improve computational efficiency and reduce overall inference time through parallel processing. Specifically, while one warp group is executing GEMM0, another warp group can simultaneously perform Softmax calculations, achieving efficient pipeline processing.

[0394] above Figure 9A and Figure 9B This simple example illustrates the asynchronous nature of TMA. Based on this, a double-buffering mechanism can be introduced. This mechanism enables efficient loading and processing of inference-dependent data. By alternating between two buffers for data processing and preloading, computing resources can be maximized, latency can be reduced, and throughput can be increased. This approach is particularly suitable for scenarios requiring large amounts of data processing, significantly improving system responsiveness and processing capabilities.

[0395] Based on the above, in order to further optimize the efficiency of data loading and processing, avoid idle computing resources, and ensure the continuity and efficiency of data processing, this application proposes detailed execution steps to implement the buffer mechanism mentioned above. An optional implementation is as follows:

[0396] Based on the first and second buffers, the following steps are executed in a loop until all inference-dependent data is processed:

[0397] S01: Call the hardware acceleration instruction to calculate and process the inference-dependent sub-data of this batch in the current calculation buffer;

[0398] S02: During the computation process, a next batch of inference-dependent sub-data is loaded from the global memory to the current preload buffer via the asynchronous transmission interface; wherein the computation buffer is one of the first buffer and the second buffer, and the preload buffer is the other of the first buffer and the second buffer;

[0399] S03: When it is detected that the calculation of the calculation buffer is completed and the data loading of the preload buffer is completed, the inference dependent sub-data in the calculation buffer and the calculation results of this batch are written back to the global memory; and the preload buffer is switched to a new calculation buffer, and the calculation buffer is switched to a new preload buffer.

[0400] Specifically, the implementation is refined by looping through data loading and processing tasks and utilizing an asynchronous transmission interface to preload the next set of data while the data in the current buffer is being processed. This approach not only ensures a seamless data processing flow but also maximizes the use of computing resources and reduces waiting time, making it a key technical means for achieving efficient inference operations.

[0401] It details how to loop through data loading and processing tasks, preloading the next set of data (i.e., the next batch of inference-dependent sub-data) through an asynchronous transmission interface while the data in the current buffer is being processed. This approach not only ensures the continuity and efficiency of data processing, but also effectively utilizes computing resources and avoids waiting time.

[0402] Assume that two buffers are allocated in shared memory, named buffer 1 (also referred to as buffer 1) and buffer 2 (also referred to as buffer 2). These two buffers have the same data layout and size, and each buffer is large enough to accommodate a batch of inference-dependent sub-data to be processed.

[0403] The following combination Figure 10 , the cycle process of S01 to S03 is described in detail:

[0404] like Figure 10 As shown, it is a buffer switching logic diagram in an embodiment of the present application; during initialization, buffer 1 can be loaded with the first batch of inference dependency sub-data, and buffer 2 remains idle.

[0405] Afterwards, buffer 1 serves as the calculation buffer during the first round of loops, and buffer 2 serves as the pre-load buffer during the first round of loops. During calculation processing, buffer 1 contains the currently processed data, such as the first batch of inference-dependent sub-data. The GMMA instruction can be called to perform relevant inference calculations on the first batch of inference-dependent sub-data in buffer 1, such as the matrix multiplication and addition calculations involved in the attention mechanism. During this process, the GMMA instruction will load data directly from buffer 1, perform calculations, and write the results back to the target location (such as global memory or shared memory). At the same time, the data loading operation of buffer 2 is started to load the next batch of data, such as the second batch of inference-dependent sub-data, from global memory to buffer 2. Data loading can be completed through CUDA kernel functions or other asynchronous loading mechanisms.

[0406] When the calculation in buffer 1 is complete, the system checks whether buffer 2 has finished loading data. If buffer 2 is loaded, the system can seamlessly switch to buffer 2 to continue the calculation. The GMMA instruction will directly load data from buffer 2, such as the second batch of inference dependent sub-data, perform the calculation, and write the results back to the target location. At this time, buffer 1 begins loading the next batch of data, such as the third batch of inference dependent sub-data.

[0407] When the calculation in buffer 2 is complete, the system checks whether buffer 1 has finished loading data. If buffer 1 is loaded, the system can seamlessly switch to buffer 1 to continue the calculation. The GMMA instruction will directly load data from buffer 1, such as the third batch of inference dependent sub-data, perform the calculation, and write the results back to the target location. At this point, buffer 2 begins loading the next batch of data, such as the fourth batch of inference dependent sub-data. And so on.

[0408] Optionally, this application supports data communication between buffers. For example, when subsequent calculations require the calculation results of the previous batch, they can be obtained through communication and further calculated.

[0409] When switching buffers, CUDA's synchronization mechanism (such as __syncthreads()) can be used to ensure the safety of switching between buffers. Specifically, when buffer 1 completes the calculation, the thread block checks whether buffer 2 has completed data loading. If buffer 2 is loaded, synchronization operations are used to ensure that all threads are ready to switch to buffer 2 for the next round of calculations. At this point, buffer 1 can be released immediately and start loading the next batch of inference-dependent sub-data. This double-buffering mechanism takes advantage of the characteristics of asynchronous loading and calculation to avoid waiting time, while preventing data competition or uninitialized access problems through synchronization mechanisms, thereby ensuring the efficiency and correctness of the inference process.

[0410] Of course, in addition to CUDA's __syncthreads() synchronization mechanism, the following similar buffer switching mechanism can also be used to ensure the security and efficiency of data loading and calculation:

[0411] (1) Event synchronization (CUDAEvents): Use CUDA events (cudaEvent_t) to mark the completion status of buffer loading and calculation, and use cudaStreamWaitEvent to achieve synchronization between streams to ensure that data loading and calculation do not conflict when switching buffers. (2) Dual-stream asynchronous execution (CUDA Streams): Distribute data loading and calculation to different CUDA streams, and use the asynchronous characteristics of streams to achieve seamless switching of buffers. For example, one stream is responsible for the calculation of buffer 1, and the other stream is responsible for the data loading of buffer 2. The switching is coordinated through the dependency relationship between streams. (3) Semaphores: In a multi-threaded or distributed environment, the semaphore mechanism can be used to control the access rights of the buffer to ensure that a buffer will not be accessed by the computing thread before the loading is completed. (4) Hardware-supported double buffering mechanism: Some GPU architectures provide hardware-level double buffering support, which automatically manages buffer switching through dedicated instructions to reduce software synchronization overhead.

[0412] The switching mechanisms listed above are just simple examples and are not specifically limited in this article. In actual applications, they can be selected and used according to specific application scenarios to optimize buffer switching efficiency and ensure the correctness of the inference process. We will not go into details here.

[0413] Follow the above logic loop until all data is processed.

[0414] Care must be taken to ensure consistency of state between buffers at each switch (e.g., to avoid data races or incomplete loads).

[0415] The mathematical expression of the above method can be expressed as the following formula 9:

[0416]

[0417] That is, buffer 1 is used for computing tasks, and buffer 2 is used for data loading, and the two are performed alternately.

[0418] In practical applications, GPU hardware can implement data exchange between two buffers through a scheduler. After each computation, the data in the current buffer is immediately written back to memory, while the data in the other buffer is loaded into the computational unit. To avoid data dependency conflicts, the hardware provides asynchronous access mechanisms and data synchronization strategies. This process involves hardware pipeline scheduling, ensuring that computation and data loading always maintain maximum parallelism. Specifically, utilizing the GPU's asynchronous data transfer channels, computation tasks and data loading tasks can be executed alternately within the same clock cycle.

[0419] In summary, the double buffering mechanism ensures maximum utilization of the GPU computing unit and avoids waiting time between data loading and calculation. The mathematical formula is expressed as follows:

[0420]

[0421] Among them, Throughput refers to the number of tasks that the system can complete per unit time; Compute Time refers to the time required to complete a specific computing task, specifically the time required for the computing buffer to complete the current batch of inference-dependent sub-data; Data Load Time specifically refers to the time required for the preload buffer to load the next batch of inference-dependent sub-data from global memory to shared memory; Latency refers to the total time from starting a request to obtaining the result.

[0422] This formula is used to calculate the amount of data a system can process or the number of tasks it can complete under given latency conditions. Specifically, it measures the amount of work a system can complete per unit time. The numerator, Compute Time + Data Load Time, represents the total time required to complete a complete operation (including data loading and calculation). The denominator, Latency, represents the total delay from initiating a request to obtaining the result. By dividing the total operation time by the latency time, we get the number of operations that the system can complete per unit delay time, i.e., throughput.

[0423] The double-buffering mechanism improves overall system efficiency by enabling computation and data loading to proceed simultaneously. Specifically, in parallel mode, computation and data loading can proceed simultaneously, avoiding idle waiting. Furthermore, the double-buffering mechanism allows for more efficient utilization of both computing units and memory bandwidth, bringing hardware performance close to its theoretical peak. Furthermore, because the processing time for each batch is shortened, the system can complete more tasks within the same timeframe, improving throughput.

[0424] In other words, since data loading and calculation can be performed in parallel, computing throughput and overall efficiency are significantly improved.

[0425] The above implementation details how to perform specific computation and data transfer tasks within a dual-buffer mechanism, preloading data via an asynchronous transmission interface and immediately switching buffer roles upon detecting the completion of the current batch of data. This approach not only ensures seamless data processing but also minimizes idle time in the computing unit, thereby improving overall system efficiency and resource utilization. It is particularly suitable for applications requiring sustained high-load computing.

[0426] Optionally, to further optimize the performance and adaptability of the dual-buffering mechanism, a dynamic expansion function is introduced in the embodiments of the present application. Specifically, in actual applications, the system continuously monitors the load status of the first buffer and the second buffer. This monitoring includes analyzing the ratio of data loading time to calculation time to ensure efficient operation of the system.

[0427] Specifically, the load status of the double buffer can be evaluated by the following indicators:

[0428] (1) Single buffer loading time: Monitors the data loading time of each buffer, that is, the time required to read the inference-dependent sub-data from the global memory and load it into the shared memory.

[0429] (2) Computation time: monitors the time spent on executing computational tasks in the buffer, such as matrix multiplication and addition operations.

[0430] Dynamic expansion of the buffer is triggered when one of the following conditions occurs:

[0431] (1) The loading time of a single buffer exceeds 150% of the calculation time: If the data loading time of a buffer is significantly longer than its calculation time (for example, the loading time is more than 1.5 times the calculation time), this indicates that data loading has become a bottleneck, affecting the overall efficiency.

[0432] (2) Imbalance in the ratio of data loading time to computing time (>1.5:1): Even if the loading time of a single buffer does not exceed 150% of the computing time, if the overall ratio of data loading time to computing time remains above 1.5:1 for a long time, it also means that more buffers are needed to balance the load.

[0433] If any of the above conditions are detected, the number of buffers will be dynamically expanded to three. This third buffer (Buffer 3) is primarily used for preemptive loading of high-priority data with high real-time requirements. This means that when the system identifies data that needs to be processed as soon as possible, it can be loaded into Buffer 3 first, ensuring that this data is not delayed due to busy other buffers.

[0434] This allows the system to better handle sudden data flows or high-priority tasks, improving overall responsiveness and stability. For example, in applications such as video streaming and real-time speech recognition, Buffer 3 ensures that key frames or important speech segments are processed promptly, avoiding quality degradation caused by delays.

[0435] In summary, 64-byte memory reordering and double-buffered write technology optimize memory access patterns while ensuring parallelism between computation and data transfer. This efficient memory and compute resource scheduling strategy significantly improves GPU performance when processing large-scale model inference, especially in scenarios involving short and long sequences. The strategy can be dynamically adjusted based on hardware status and data characteristics to achieve optimal computational efficiency and memory utilization.

[0436] Optionally, after S221, there is another situation (parallel to S222 to S223), that is, if it is determined in S221 that the inference dependent data itself is stored continuously, then there is no need to reorganize the memory, and the following S224 is directly executed ( Figure 2 not shown):

[0437] S224: When the storage location of the inference-dependent data is in a continuous state in the computing resources, directly call the specified hardware acceleration instruction to load the inference-dependent data, and after determining the result of this round of inference based on the loading result, generate the input sequence of the next round according to the result of this round of inference, until the preset stop condition is reached.

[0438] That is, during the execution of the reasoning operation, after the storage continuity check is performed on the reasoning-dependent data related to the input sequence, if it is found that these data are already in a continuous state in the computing resources, there is no need to perform additional data rearrangement steps. At this time, the specified hardware acceleration instruction can be directly called to load these data, and the result of this round of reasoning can be determined based on the loading result, and then the input sequence of the next round is generated based on the result. This process follows the logical framework of steps S221 to S223, but simplifies the process under certain circumstances. It is equivalent to skipping the data rearrangement step of S222 if the S221 check result shows that the data is continuous, directly calling the hardware acceleration instruction to load the data, and starting the reasoning process, further improving efficiency.

[0439] like Figure 11A As shown, it is a processing logic diagram of an instruction acceleration mode in an embodiment of the present application. By judging the length of the input sequence and determining that it is less than 512 tokens, which is a short sequence, the storage continuity check of the inference dependency data related to the input sequence is performed (i.e., S222). If the state is continuous, the GMMA instruction is directly called (i.e., S224 is executed). If the state is discontinuous, the TMA reordering is triggered (i.e., S222 to S223 are executed).

[0440] The data loading and reasoning process can be referred to in the above embodiment. A simple example is given below:

[0441] For example, if the inference relies on data stored contiguously, the Hopper architecture's GMMA instructions can be directly used for efficient computation. These instructions are particularly well-suited for processing data in contiguous memory layouts, significantly reducing data transfer overhead and improving computational efficiency.

[0442] Optionally, in the process of directly calling the GMMA instruction, the double buffering mechanism can also be used to further improve efficiency, that is, when executing S223 or S224, the double buffering mechanism can be combined to achieve, such as Figure 11B This is a processing logic diagram of another instruction acceleration mode in an embodiment of the present application; wherein, Figure 11B This example illustrates the logic of performing inference calculations in both discontinuous and continuous states using a dual-buffering mechanism. In the discontinuous state, memory reorganization is required to adjust the inference-dependent data to a continuous state. Buffers 1 and 2 are then used to alternately load the inference-dependent sub-data from global memory for calculation. In the continuous state, memory reorganization is unnecessary; Buffers 1 and 2 can be used to alternately load the inference-dependent sub-data from global memory for calculation. For detailed implementation details, please refer to the previous examples, and any repetitions will not be repeated here.

[0443] like Figure 12 As shown, it is a schematic diagram of a GPU in an embodiment of the present application, which simply illustrates two GPUs of reasoning precision, wherein the throughput of the GPU on the left is increased by 6 times by adopting the method in this application. In the GPU reasoning process, especially when processing short sequences or continuous memory scenarios, traditional computing methods often become performance bottlenecks due to redundant memory access operations. To avoid these problems, the Hopper GPU architecture introduces the GMMA instruction, which specifically optimizes matrix multiplication and addition operations (such as GEMM), which is particularly suitable for efficient use of tensor cores.

[0444] Designed for contiguous memory layouts, the GMMA instruction supports mixed-precision computations (such as BF16 primary computations and FP32 accumulations), significantly improving throughput. By directly accessing contiguous blocks of memory, it bypasses unnecessary memory accesses and data padding steps, significantly improving computational efficiency. This contiguous memory layout ensures that data is stored in order, enabling a single instruction to perform multiple multiply-add operations, maximizing the parallel computing capabilities of the GPU.

[0445] In order to use the GMMA instruction, the data must first be stored according to the requirements of the GPU memory system, ensuring that each matrix data block is arranged continuously in the global memory. For example, matrices A and B should be arranged in adjacent memory locations, as shown in the following mathematical expression:

[0446] A = contiguous memory block, B = contiguous memory block (Formula 11);

[0447] In this way, each matrix element has a fixed, adjacent storage location, without the need for additional memory rearrangement or zero padding, reducing computing power waste.

[0448] Once the data is prepared in a contiguous memory layout, the GMMA instruction can be called to perform matrix multiplication and accumulation operations. The GMMA instruction directly loads the data from this contiguous memory and completes the multiplication and accumulation of the entire matrix in a single operation cycle, avoiding the need for piecemeal calculations and frequent data loading. The mathematical representation is as follows:

[0449] C = A·B + C (Formula 12);

[0450] Here, A and B are input matrices, and C is the accumulated result matrix, achieving incremental updates. Because the memory is continuous, the GPU can perform matrix multiplication and addition operations in parallel while loading data, significantly reducing waiting time.

[0451] The GMMA instruction also optimizes memory access patterns, particularly in contiguous memory scenarios. Due to the contiguous nature of data, GPU memory accesses are no longer subject to fragmentation, allowing each memory access to fully utilize bandwidth. For example, the memory access stride matches the matrix dimensions and memory block size, ensuring that each load operation maximizes memory bandwidth.

[0452] In actual implementation, the Hopper architecture's Tensor Cores use GMMA instructions to directly perform multiplication and addition operations on consecutive matrix blocks in memory, avoiding the overhead of memory reordering and significantly improving computational efficiency and throughput. This approach not only reduces data transfer time but also improves the overall system response speed, making it particularly suitable for applications requiring real-time and high performance.

[0453] The GMMA instruction significantly improves GPU inference performance in contiguous memory scenarios by optimizing memory access and computational processes. Its core advantages lie in efficient computational throughput, reduced computational bottlenecks, lower computational overhead, and significant performance improvements.

[0454] The GMMA instruction leverages memory continuity to directly access contiguous data blocks in global memory, eliminating redundant memory accesses and data padding operations (such as zero padding), significantly improving computational throughput. This design significantly improves throughput, especially when processing short sequences (e.g., less than 512 tokens), unleashing the GPU's computational potential.

[0455] In addition, the GMMA instruction can complete multiple matrix multiplication and addition operations in a single clock cycle, reducing memory access latency and computational wait time. This feature effectively avoids performance bottlenecks caused by memory fragmentation or non-contiguous data layout, enabling the GPU to execute tasks more efficiently. At the same time, the GMMA instruction supports mixed-precision calculations (such as BF16 main calculations and FP32 accumulation), which not only reduces computational complexity but also improves computational efficiency, making it particularly suitable for large-scale matrix operation tasks.

[0456] In compute-intensive tasks, the GMMA instruction can increase the throughput of each GPU core, delivering several-fold performance improvements. Especially in the Hopper architecture, combined with the hardware acceleration capabilities of the Tensor Cores, the GMMA instruction fully leverages the advantages of contiguous memory layout. By optimizing memory layout, reducing redundant operations, and leveraging mixed-precision computations, the GMMA instruction significantly reduces memory bandwidth usage, improving computational density and throughput.

[0457] Overall, the GMMA instruction demonstrates exceptional efficiency when processing large-scale matrix operations, reducing memory access latency and improving overall inference performance. This optimization enables the GPU to more efficiently execute deep learning inference tasks, meeting high-performance computing requirements, particularly in applications with high real-time requirements.

[0458] The above implementation provides a simplified data processing path. When the inference-dependent data is already stored in a continuous format, hardware-accelerated instructions are directly invoked for loading and processing, eliminating the need for additional data reordering steps. This approach reduces unnecessary computational overhead and improves overall inference efficiency, especially when the data itself is already optimized.

[0459] Optionally, on the basis of simplifying the data loading process, in order to further improve the overall performance of the reasoning stage, this application also introduces a precision control mechanism to adapt to the differentiated requirements of different tasks for reasoning quality and execution efficiency. That is, in order to effectively reduce the computing load and improve the overall execution efficiency while ensuring the reasoning accuracy, this application proposes a hierarchical precision reasoning strategy to further optimize the process of generating reasoning results. In an optional implementation, precision control can be performed through the following steps to determine the current round of reasoning results based on the loading results:

[0460] First, each time a load is performed and a calculation is performed, the effective data length of the load results of this batch is determined.

[0461] The loading result refers to the data content loaded for this round of reasoning after calling the hardware acceleration instruction from the computing resource, such as the reasoning-dependent sub-data loaded in this batch.

[0462] In summary, since the actual loading process may read data from global memory according to a fixed step size, block size, or alignment boundary, the load result is generally a portion of the inference-dependent data. Of course, when the inference-dependent data is small, the load result can also be the entire inference-dependent data. Because the inference-dependent data is guaranteed to be continuous before loading, no zero padding is required. Therefore, when performing calculations, the effective data length of the data loaded in this batch is used, that is, the data range actually involved in the calculation, which does not include zero values.

[0463] For example, when the inference-dependent data is a KV cache in the attention mechanism, the loading result may include a continuous range of KV values. In this case, the effective data length can be determined based on the index interval of the loaded KV data. In the embodiment of the present application, determining the effective data length helps to rationally configure the calculation accuracy and loading strategy in the subsequent inference stage, avoid redundant processing of invalid data, and improve inference efficiency and resource utilization.

[0464] Secondly, according to the effective data length, the first reasoning precision is adopted to perform the main reasoning operation on the loaded result to obtain the intermediate result.

[0465] The main inference operation refers to the main mathematical calculation step performed on the data in the loaded result during the inference process, which can be a multiplication operation between matrices. For example, in the matrix multiply-accumulate (MMA) calculation process, the main inference operation is the multiplication operation between matrix A and matrix B.

[0466] During the main inference process, in order to improve execution efficiency, the first inference accuracy can be used for the main calculation operation to reduce the memory bandwidth usage and improve the parallel execution capability.

[0467] In an optional embodiment, the first inference precision can be BF16 (bfloat16), or other low-bit-width floating-point or fixed-point formats with fast computational capabilities. BF16 is a 16-bit floating-point format that retains the same number of exponent bits as 32-bit floating-point (FP32) and reduces only the mantissa bits. This significantly compresses the storage space required for data representation while retaining a large dynamic range. It is suitable for inference tasks that require relatively moderate accuracy but require increased speed.

[0468] In the embodiment of the present application, in the multiplication-addition calculation process of the GMMA instruction, the main reasoning operation corresponds to the matrix multiplication part thereof, such as A·B in the above formula 12.

[0469] The following will take the loading result as a set of key vectors (Key) as an example to illustrate the actual execution process of the main inference operation.

[0470] In inference tasks, the input sequence is typically converted to vector form to participate in the attention mechanism calculation. When the loading result is a set of key vectors from the historical cache, the main inference operation can be expressed as a matrix multiplication calculation between the vector corresponding to the current input and the set of key vectors to obtain the attention level of different historical positions, such as the above formula 6. In this calculation process, Q can represent the query vector corresponding to a word in the current input sequence, and K represents the key vector corresponding to a word (i.e., the inference dependency data).

[0471] This process actually performs a set of multiplications between vectors and matrices, and the output is the matching score of the current input at different historical positions. It serves as the intermediate result of this round and provides a basis for subsequent attention weighted calculations or generation of control logic.

[0472] This type of matrix multiplication operation is usually computationally intensive in actual deployment. By adopting the first inference precision for execution, it not only reduces the load of a single calculation but also improves the overall inference throughput. It is particularly suitable for use in high-concurrency, low-latency model inference environments.

[0473] Finally, the second reasoning accuracy is used to perform a secondary reasoning operation on the intermediate result to obtain the result of this round of reasoning; wherein, the first reasoning accuracy is lower than the second reasoning accuracy.

[0474] Secondary inference operations are minor mathematical calculations performed on the data in the loaded results during the inference process. These operations can primarily take the form of vector addition or weighted summation. These operations are used to further integrate the loaded data with the intermediate results obtained during the primary inference phase to generate the final inference output.

[0475] In the GMMA instruction's multiplication-addition formula C = A·B + C, the secondary inference operation corresponds to the addition part + C, which is used to fuse the current round's calculation results with the historical accumulated state to ensure the stability and accuracy of the final inference output.

[0476] Taking the above-mentioned attention mechanism as an example, the primary reasoning operation mainly involves performing a multiplication operation on the Key (K) vector in the loaded result to obtain an attention score, while the secondary reasoning operation is based on the attention score. The weighted addition process of the Value (V) vector in the loaded result is performed to generate the reasoning output.

[0477] For example, the attention score obtained for the current input can be used as a weight to correspond to multiple Value vectors in the loaded result, and the output result is generated by weighted summation, as shown in the above formula 8. The weighted addition operation in this calculation process is a typical vector fusion operation and can be executed with the second inference accuracy to ensure numerical stability and the accuracy of the final result.

[0478] In an optional implementation, the second inference precision can be FP32 (32-bit floating point) or other floating point formats with high numerical representation capabilities, used for calculations such as addition and weighted summation during the secondary inference operation. Different inference scenarios can flexibly select the second inference precision within the supported precision range based on the combined requirements for accuracy and performance.

[0479] It should be noted that the specific types of first and second reasoning precisions are not limited to the examples above; the specific combination can be flexibly configured based on the model structure, hardware capabilities, and application scenarios. However, regardless of the combination, the basic relationship between the first and second reasoning precisions must be maintained. This ensures that the primary reasoning phase prioritizes performance, while the secondary reasoning phase prioritizes accuracy, thus achieving a balance between efficiency and accuracy in the reasoning process.

[0480] In addition to the first and second inference precisions, a third inference precision can be introduced for tasks such as intermediate computational processing, error correction, lightweight fusion, or post-processing optimization at specific stages. In actual implementation, the third inference precision can be between the first and second. For example, when the first inference precision is INT8 and the second inference precision is FP32, the third inference precision can be set to BF16 to further refine the transition between accuracy and performance. It can also be higher than the second inference precision, such as using FP64 to process some highly sensitive data to meet the numerical stability requirements in specific scenarios.

[0481] By introducing a third level of inference precision, a more flexible hierarchical precision path can be constructed, improving the model's adaptability to diverse task requirements and operating environment changes during inference. This application does not limit the number of levels of inference precision and their combination.

[0482] Through the above implementation method, the valid data range can be dynamically identified according to the actual usage of the loaded results during the reasoning process, avoiding repeated processing of invalid or redundant data, and improving the utilization of computing resources. At the same time, a precision grading strategy is introduced, and the main reasoning operation and the secondary reasoning operation are processed with the first reasoning precision and the second reasoning precision respectively, thereby achieving the coordinated optimization of computing efficiency and reasoning accuracy. Among them, the main reasoning stage adopts lower precision to speed up the execution of high computing load operations such as matrix multiplication, while the secondary reasoning stage adopts higher precision to ensure the numerical stability and output accuracy of precision-sensitive operations such as weighted addition. The overall solution effectively improves the throughput and execution flexibility of the reasoning process, and is suitable for large-scale model reasoning scenarios that have requirements for both response latency and prediction accuracy.

[0483] See Figure 13 As shown, it is an optional interactive implementation timing logic diagram in the embodiment of the present application. Figure 13Taking the target model as LLM as an example, this paper briefly illustrates the task processing process in the intelligent AI assistant scenario. In the intelligent AI assistant scenario, users can enter the detailed information of the task to be inferred in the intelligent AI assistant interface on the terminal device side ( Figure 13 The terminal device sends this information to the server. On the server side, this information is used as the input of LLM. After the process shown in S21 to S22, multiple rounds of reasoning are performed to obtain the final task reasoning result ( Figure 13 The specific reasoning process is not introduced in detail, and it depends on the actual situation) and returned to the terminal device side, which is displayed by the terminal device through the intelligent AI assistant interface. The specific reasoning process can be found in the above embodiment, and the repeated parts will not be repeated.

[0484] Optionally, when the inference acceleration method in the embodiment of the present application is applied to the model training stage, there may be multiple tasks to be inferred. When the target model batch processes requests for multiple tasks to be inferred with different sequence lengths, there will be storage gaps between different requests, resulting in storage fragmentation, thereby affecting the inference speed. The inference acceleration method in the embodiment of the present application can effectively improve the above problems and improve the inference speed.

[0485] In summary, this application aims to dynamically optimize the computing paradigm based on hardware status and sequence characteristics, which is particularly suitable for scenarios with variable-length input sequences. By introducing an intelligent mode decision mechanism, the most suitable processing mode (such as GMMA high-speed mode or paged cache mode) can be automatically selected under different workloads, and performance can be further improved through hardware instruction stream optimization. This approach not only improves computing efficiency, but also optimizes video memory management, and is particularly suitable for large language model (LLM) reasoning tasks.

[0486] Specifically, this application uses a hardware-aware dual-mode dynamic switching engine and a page cache co-scheduling mechanism to optimize the GPU accelerated inference process. This design enables hardware-aware dynamic mode switching, reduces computing power waste in short sequence scenarios, and optimizes video memory management, thereby improving overall system performance. Especially when processing variable-length sequence inference tasks, this method can effectively manage video memory resources while maintaining high performance, achieving a global optimal balance between video memory efficiency and computing performance.

[0487] This approach effectively improves the deployment efficiency and performance of large-scale model real-time inference services, providing a more flexible and efficient solution. Whether dealing with short or long sequence inputs, this approach ensures optimal utilization of computing resources and significantly improves the performance of large language models in real-world applications.

[0488] To more intuitively demonstrate the performance of the method provided by the embodiment of this application on an actual hardware platform, this application also provides experimental data based on a Hopper architecture GPU. The experiment compares the computing performance of the method of this application with the method of related technologies under multiple typical matrix configurations. The test results are shown in Table 1:

[0489] Table 1

[0490] M N K Computation (throughput) Speedup 4096 7168 16384 1358 TFLOPS 1.2x 4096 4096 7168 1304 TFLOPS 1.2x 4096 7168 2048 1025 TFLOPS 1.3x

[0491] Among them, M, N, and K represent the three dimensions in matrix multiplication, specifically: M represents the number of rows of the first input matrix (A), N represents the number of columns of the second input matrix (B), which is also the number of columns of the result matrix (C), and K represents the number of columns of the first matrix and the number of rows of the second matrix, which is the shared dimension when the two matrices are multiplied; Computation represents the computational throughput measured under the corresponding matrix configuration, in TFLOPS (trillion floating-point operations per second); Speedup represents the performance improvement multiple compared with related technologies (such as the traditional cuBLAS library) under the same matrix configuration, which is used to measure the acceleration effect of the method of this application in actual hardware execution.

[0492] The results shown in Table 1 show that the method provided by the embodiment of the present application achieves significant computing performance improvements under various typical matrix dimension configurations, with the computing throughput reaching up to 1358 TFLOPS and the overall acceleration ratio being stable between 1.2 and 1.3 times, demonstrating the efficiency of this method in large-scale matrix multiplication calculations. Specifically, in scenarios with larger dimensions (such as K = 16384), a high throughput capacity is still maintained, indicating that the solution of this application has good scalability and stability; and in smaller-scale configurations (such as K = 2048), a speedup ratio of up to 1.3 times is still achieved, reflecting its wide applicability under different computing scales.

[0493] The above results verify that the technical solution of this application, when executed in low-precision computing mode (such as FP8, BF16) combined with video memory continuity optimization and hardware instruction acceleration path (such as GMMA), can effectively release hardware potential, reduce redundant computing and data handling overhead, and improve the inference throughput in actual operation. It is particularly suitable for large model inference, training acceleration and other high-density parallel computing scenarios.

[0494] Based on the same inventive concept, the present application also provides an inference acceleration device. Figure 14 As shown, it is a schematic diagram of the structure of the inference acceleration device 1400, which may include:

[0495] The decision unit 1401 is configured to activate a reasoning mode that matches the current hardware status of the computing resources occupied by the target model when the target model is used to execute the task to be inferred.

[0496] The inference unit 1402 is configured to perform at least one round of inference operations when the inference mode is the instruction acceleration mode, where each round of inference operations includes:

[0497] Performing a storage continuity check in the computing resource for the inference dependency data associated with the input sequence of the current round, wherein the inference dependency data is generated and stored by the inference operation of the previous round;

[0498] When the storage location of the inference-dependent data is in a discontinuous state in the computing resource, adjusting the storage location to a continuous state by rearranging the inference-dependent data;

[0499] The specified hardware acceleration instruction is called to load the adjusted inference dependency data, and after determining the result of the current round of inference based on the loading result, the input sequence of the next round is generated according to the result of the current round of inference, until a preset stop condition is reached.

[0500] Optionally, after performing a storage continuity check on the inference dependency data related to the input sequence of this round in the computing resource, the inference unit 1402 is further configured to:

[0501] When the storage location of the reasoning-dependent data is in a continuous state in the computing resource, the specified hardware acceleration instruction is directly called to load the reasoning-dependent data, and after determining the result of this round of reasoning based on the loading result, the input sequence of the next round is generated according to the result of this round of reasoning until the preset stop condition is reached.

[0502] Optionally, before calling a specified hardware acceleration instruction to load the inference dependency data and determining the result of the current round of inference based on the loading result, the inference unit 1402 is further configured to:

[0503] Allocating a first buffer and a second buffer in a shared memory of the computing resource;

[0504] The reasoning unit 1402 is specifically configured to:

[0505] Call the hardware acceleration instruction to alternately load different batches of inference-dependent sub-data from the first buffer and the second buffer for calculation and processing, respectively, to obtain the result of this round of inference; wherein the first buffer and the second buffer are used to alternately load different batches of inference-dependent sub-data from the inference-dependent data stored in the global memory of the computing resource.

[0506] Optionally, the reasoning unit 1402 is specifically configured to:

[0507] Based on the first buffer and the second buffer, the following steps are executed in a loop until the inference-dependent data is processed:

[0508] Calling the hardware acceleration instruction to perform computation on the inference-dependent sub-data of this batch in the current computation buffer;

[0509] During the computation process, a next batch of inference-dependent sub-data is loaded from the global memory to the current preload buffer via an asynchronous transmission interface; wherein the computation buffer is one of the first buffer and the second buffer, and the preload buffer is the other of the first buffer and the second buffer;

[0510] When it is detected that the calculation of the calculation buffer is completed and the data loading of the preload buffer is completed, the inference dependent sub-data and the calculation results of this batch in the calculation buffer are written back to the global memory; and the preload buffer is switched to a new calculation buffer, and the calculation buffer is switched to a new preload buffer.

[0511] Optionally, the reasoning dependency data includes: the latest reasoning dependency data related to the latest reasoning result in the input sequence of the current round, and the historical reasoning dependency data related to other information in the input sequence of the current round;

[0512] The reasoning unit 1402 is specifically configured to:

[0513] Checking storage continuity of the latest reasoning dependency data and the historical reasoning dependency data in the computing resource;

[0514] When the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource, the latest reasoning dependency data and the historical reasoning dependency data are adjusted to a continuous storage state by rearranging the latest reasoning dependency data.

[0515] Optionally, the reasoning unit 1402 is further configured to determine the storage locations of the latest reasoning dependency data and the historical reasoning dependency data in a discontinuous state in the computing resource by any of the following methods:

[0516] When the storage location of the latest reasoning dependency data is in a discontinuous state in the computing resource, determining that the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource;

[0517] When the storage location of the latest reasoning dependency data and the storage location of the historical reasoning dependency data are not adjacent in the computing resource, it is determined that the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource.

[0518] Optionally, the reasoning unit 1402 is specifically configured to:

[0519] For each memory block storing the inference-dependent data, decompose each memory block into a plurality of sub-blocks according to a preset data block size; wherein the size of each sub-block is an integer multiple of the preset data block size; the preset data block size is determined based on the memory access rule of the computing resource;

[0520] The obtained sub-blocks are rearranged according to the target continuous storage order to adjust the storage positions to a continuous state.

[0521] Optionally, the reasoning unit 1402 is specifically configured to:

[0522] Determine the effective data length of the loading result;

[0523] According to the effective data length, using the first reasoning precision, performing a main reasoning operation on the loaded result to obtain an intermediate result;

[0524] Using the second reasoning precision, a secondary reasoning operation is performed on the intermediate result to obtain the current round reasoning result; wherein the first reasoning precision is lower than the second reasoning precision.

[0525] Optionally, the decision unit 1401 is specifically configured to:

[0526] Determining a mode decision factor corresponding to the computing resource by quantifying the hardware status information;

[0527] If the mode decision factor is greater than a first threshold, activating the instruction acceleration mode;

[0528] If the mode decision factor is less than a second threshold, a paging cache mode is activated; wherein the paging cache mode is used to divide the input sequence of this round into multiple short sequences according to a preset length, and the inference dependency data corresponding to each short sequence is continuously stored using different paging units;

[0529] If the mode decision factor is not less than the second threshold and not greater than the first threshold, the current reasoning mode is maintained.

[0530] Optionally, the hardware status information includes: computing unit utilization, memory access efficiency, and bandwidth utilization; the decision unit 1401 is specifically configured to:

[0531] Determining the utilization rate of the computing unit according to the number of idle cycles of the target computing unit in the computing resource within a preset time period;

[0532] Determining the memory access efficiency according to the number of accesses to the target memory layer in the computing resource within a preset time period;

[0533] determining the bandwidth utilization rate according to an amount of bandwidth used by a target storage unit in the computing resource within a preset time period;

[0534] A weighted sum is performed based on the computing unit utilization, the memory access efficiency, and the bandwidth utilization to obtain the mode decision factor.

[0535] Optionally, before performing at least one round of reasoning operations, the reasoning unit 1402 is further configured to:

[0536] For the task to be inferred, obtaining an initial input sequence corresponding to the task to be inferred;

[0537] The length of the initial input sequence is compared with a preset length threshold, and it is determined that the length of the initial input sequence is less than the preset length threshold.

[0538] This application dynamically monitors the hardware status of computing resources and selects the inference mode that best matches the current hardware state. This dynamic adjustment mechanism optimizes resource utilization and ensures efficient inference operations under different hardware conditions.

[0539] In instruction acceleration mode, the input sequence-related reasoning-dependent data usually includes intermediate results generated by the previous round. By checking the storage continuity of this data in the storage, this inspection step can identify performance issues caused by non-contiguous storage, laying the foundation for subsequent optimization, and ensuring that the reasoning process can run under efficient data layout. For example, when it is detected that the reasoning-dependent data is stored non-contiguously, the data is adjusted to a continuous state. This rearrangement operation reduces storage fragmentation, improves data loading efficiency, and avoids redundant calculations caused by zero padding. Compared with the static pre-allocation padding technology, this method significantly improves storage utilization and reduces invalid overhead. After the data storage continuity is optimized, the specified hardware acceleration instruction is called to load the continuous reasoning-dependent data and perform efficient operations. This process fully utilizes the hardware acceleration capabilities and further improves the reasoning efficiency.

[0540] Overall, this technical solution addresses the issues of low storage utilization and latency fluctuations in related technologies through dynamic hardware adaptation, storage optimization, and efficient instruction calls. Compared to static pre-allocated padding and dynamic storage requests, this approach offers significant advantages in storage management and computational efficiency. It is particularly well-suited for inference scenarios involving large models with variable-length sequences, effectively improving inference performance and hardware resource utilization.

[0541] For the convenience of description, the above parts are divided into modules (or units) according to their functions and described separately. Of course, when implementing this application, the functions of each module (or unit) can be implemented in the same or multiple software or hardware.

[0542] In the embodiments of the present application, the term "module" or "unit" refers to a computer program or a part of a computer program that has a predetermined function and works together with other related parts to achieve a predetermined goal, and can be implemented in whole or in part by using software, hardware (such as processing circuits or memories) or a combination thereof. Similarly, a processor (or multiple processors or memories) can be used to implement one or more modules or units. In addition, each module or unit can be part of an overall module or unit that includes the function of the module or unit.

[0543] After introducing the inference acceleration method and apparatus according to an exemplary embodiment of the present application, an electronic device according to another exemplary embodiment of the present application is introduced next.

[0544] Those skilled in the art will appreciate that various aspects of the present application can be implemented as systems, methods, or program products. Therefore, various aspects of the present application can be specifically implemented in the following forms: a complete hardware implementation, a complete software implementation (including firmware, microcode, etc.), or an implementation that combines hardware and software aspects, which may be collectively referred to herein as a "circuit," "module," or "system."

[0545] Based on the same inventive concept as the above method embodiment, an electronic device is also provided in the embodiment of the present application. In one embodiment, the electronic device may be a server, such as Figure 1 In this embodiment, the structure of the electronic device can be as follows: Figure 15 As shown, it includes a memory 1501 , a communication module 1503 and one or more processors 1502 .

[0546] Memory 1501 is used to store computer programs executed by processor 1502. Memory 1501 may primarily include a program storage area and a data storage area. The program storage area may store an operating system and programs required for running instant messaging functions, while the data storage area may store various instant messaging messages and operating instruction sets.

[0547] Memory 1501 may be a volatile memory, such as random-access memory (RAM); a non-volatile memory, such as read-only memory, flash memory, a hard disk drive (HDD), or a solid-state drive (SSD); or any other medium capable of carrying or storing a desired computer program in the form of instructions or data structures and accessible by a computer, but is not limited thereto. Memory 1501 may be a combination of the aforementioned memories.

[0548] The processor 1502 may include one or more central processing units (CPUs) or digital processing units, etc. The processor 1502 is configured to implement the above-mentioned inference acceleration method when calling the computer program stored in the memory 1501 .

[0549] The communication module 1503 is used to communicate with terminal devices and other servers.

[0550] The specific connection medium between the memory 1501, the communication module 1503 and the processor 1502 is not limited in the embodiment of the present application. Figure 15 In the embodiment, the memory 1501 and the processor 1502 are connected via a bus 1504. The bus 1504 is connected to the processor 1502 via a bus 1504. Figure 15 The connections between the other components are shown in bold lines for illustration only and are not intended to be limiting. The bus 1504 can be divided into an address bus, a data bus, a control bus, etc. For ease of description, Figure 15 The diagram shows a single thick line, but this does not indicate that there is only one bus or one type of bus.

[0551] The memory 1501 stores a computer storage medium, which stores computer executable instructions. The computer executable instructions are used to implement the reasoning acceleration method of the embodiment of the present application. The processor 1502 is used to execute the above-mentioned reasoning acceleration method, such as Figure 2 shown.

[0552] In some possible implementations, various aspects of the inference acceleration method provided in this application may also be implemented in the form of a program product, which includes a computer program. When the program product is run on an electronic device, the computer program is used to enable the electronic device to execute the steps of the inference acceleration method according to various exemplary embodiments of the present application described above in this specification. For example, the electronic device may execute the following steps: Figure 2 Follow the steps shown in .

[0553] The program product may employ any combination of one or more readable media. The readable medium may be a readable signal medium or a readable storage medium. The readable storage medium may be, for example, but not limited to, an electrical, magnetic, optical, electromagnetic, infrared, or semiconductor system, device, or component, or any combination thereof. More specific examples of readable storage media (a non-exhaustive list) include: an electrical connection with one or more wires, a portable disk, a hard disk, a random access memory (RAM), a read-only memory (ROM), an erasable programmable read-only memory (EPROM or flash memory), an optical fiber, a portable compact disk read-only memory (CD-ROM), an optical storage device, a magnetic storage device, or any suitable combination thereof.

[0554] The program product of the embodiment of the present application may be a portable compact disc read-only memory (CD-ROM) and include a computer program, and can be run on an electronic device. However, the program product of the present application is not limited thereto. In this document, a readable storage medium may be any tangible medium containing or storing a program that can be used by or in conjunction with a command execution system, apparatus, or device.

[0555] A readable signal medium may include a data signal propagated in baseband or as part of a carrier wave, which carries a readable computer program. Such a propagated data signal may take a variety of forms, including but not limited to electromagnetic signals, optical signals, or any suitable combination thereof. A readable signal medium may also be any readable medium other than a readable storage medium that can transmit, propagate, or transfer a program for use by or in conjunction with a command execution system, apparatus, or device.

[0556] The computer program embodied on the readable medium may be transmitted using any appropriate medium, including but not limited to wireless, wireline, optical fiber cable, RF, etc., or any suitable combination of the foregoing.

[0557] The computer program for performing the operations of the present application can be written in any combination of one or more programming languages, including object-oriented programming languages such as Java, C++, etc., and conventional procedural programming languages such as "C" or similar programming languages. The computer program can be executed entirely on the user electronic device, partially on the user electronic device, as a separate software package, partially on the user electronic device and partially on a remote electronic device, or entirely on a remote electronic device or server. In cases involving remote electronic devices, the remote electronic device can be connected to the user electronic device through any type of network, including a local area network (LAN) or a wide area network (WAN), or can be connected to an external electronic device (for example, using an Internet service provider to connect through the Internet).

[0558] It should be noted that although several units or subunits of the device are mentioned in the detailed description above, this division is merely exemplary and not mandatory. In fact, depending on the embodiment of the application, the features and functions of two or more units described above can be embodied in a single unit. Conversely, the features and functions of a single unit described above can be further divided and embodied by multiple units.

[0559] Furthermore, although the operations of the method of the present application are described in a particular order in the accompanying drawings, this does not require or imply that the operations must be performed in this particular order, or that all illustrated operations must be performed to achieve the desired results. Additionally or alternatively, some steps may be omitted, multiple steps may be combined into one step, and / or one step may be decomposed into multiple steps.

[0560] Those skilled in the art will appreciate that the embodiments of the present application may be provided as methods, systems, or computer program products. Therefore, the present application may adopt the form of a complete hardware embodiment, a complete software embodiment, or an embodiment combining software and hardware. Furthermore, the present application may adopt the form of a computer program product implemented on one or more computer-usable storage media (including but not limited to magnetic disk storage, CD-ROM, optical storage, etc.) that contain a computer-usable computer program.

[0561] The present application is described with reference to the flowcharts and / or block diagrams of the methods, devices (systems), and computer program products according to the embodiments of the present application. It should be understood that each process and / or box in the flowchart and / or block diagram, as well as the combination of the processes and / or boxes in the flowchart and / or block diagram, can be implemented by computer program commands. These computer program commands can be provided to a processor of a general-purpose computer, a special-purpose computer, an embedded processor, or other programmable data processing device to produce a machine, so that the commands executed by the processor of the computer or other programmable data processing device generate commands for implementing the processes in the flowchart and / or block diagram. Figure 1 a process or multiple processes and / or boxes Figure 1 A device that provides the functions specified in a block or multiple blocks.

[0562] These computer program instructions may also be stored in a computer readable memory capable of directing a computer or other programmable data processing device to operate in a specific manner, so that the instructions stored in the computer readable memory produce an article of manufacture comprising a command device, which implements the process Figure 1 a process or multiple processes and / or boxes Figure 1 The function specified in one or more boxes.

[0563] These computer program instructions can also be loaded onto a computer or other programmable data processing device so that a series of operating steps are executed on the computer or other programmable device to produce a computer-implemented process, thereby providing instructions for implementing the process. Figure 1 a process or multiple processes and / or boxes Figure 1 A step that specifies a function in one or more boxes.

[0564] Although the preferred embodiments of the present application have been described, those skilled in the art may make additional changes and modifications to these embodiments once they have learned the basic creative concept. Therefore, the appended claims are intended to be interpreted as including the preferred embodiments and all changes and modifications that fall within the scope of the present application.

[0565] Obviously, those skilled in the art may make various changes and modifications to this application without departing from the spirit and scope of this application. Thus, if these modifications and variations of this application fall within the scope of the claims of this application and their equivalents, this application is intended to include these modifications and variations.

Claims

1. A method for accelerating inference, characterized in that: The method comprises: When the target model is used to execute the task to be inferred, based on the hardware status information of the computing resources occupied by the target model, an inference mode that matches the current hardware status of the computing resources is activated; and when the inference mode is the instruction acceleration mode, at least one round of inference operation is performed, each round of inference operation including: Performing a storage continuity check in the computing resource for the inference dependency data associated with the input sequence of the current round, wherein the inference dependency data is generated and stored by the inference operation of the previous round; When the storage location of the inference-dependent data is in a discontinuous state in the computing resource, adjusting the storage location to a continuous state by rearranging the inference-dependent data; The specified hardware acceleration instruction is called to load the adjusted inference dependency data, and after determining the result of the current round of inference based on the loading result, the input sequence of the next round is generated according to the result of the current round of inference, until a preset stop condition is reached.

2. The method according to claim 1, wherein After performing a storage continuity check on the inference dependency data related to the input sequence of the current round in the computing resource, the method further includes: When the storage location of the reasoning-dependent data is in a continuous state in the computing resource, the specified hardware acceleration instruction is directly called to load the reasoning-dependent data, and after determining the result of this round of reasoning based on the loading result, the input sequence of the next round is generated according to the result of this round of reasoning until the preset stop condition is reached.

3. The method according to claim 1 or 2, wherein: Before calling the specified hardware acceleration instruction to load the inference-dependent data and determining the result of the current round of inference based on the loading result, the method further includes: Allocating a first buffer and a second buffer in a shared memory of the computing resource; The specified hardware acceleration instruction is called to load the inference-dependent data, and the result of this round of inference is determined based on the loaded result, including: Call the hardware acceleration instruction to alternately load different batches of inference-dependent sub-data from the first buffer and the second buffer for calculation and processing, respectively, to obtain the result of this round of inference; wherein the first buffer and the second buffer are used to alternately load different batches of inference-dependent sub-data from the inference-dependent data stored in the global memory of the computing resource.

4. The method according to claim 3, wherein The calling of the hardware acceleration instruction to alternately load different batches of inference-dependent sub-data from the first buffer and the second buffer for calculation processing includes: Based on the first buffer and the second buffer, the following steps are executed in a loop until the inference-dependent data is processed: Calling the hardware acceleration instruction to perform computation on the inference-dependent sub-data of this batch in the current computation buffer; During the computation process, a next batch of inference-dependent sub-data is loaded from the global memory to the current preload buffer via an asynchronous transmission interface; wherein the computation buffer is one of the first buffer and the second buffer, and the preload buffer is the other of the first buffer and the second buffer; When it is detected that the calculation of the calculation buffer is completed and the data loading of the preload buffer is completed, the inference dependent sub-data and the calculation results of this batch in the calculation buffer are written back to the global memory; and the preload buffer is switched to a new calculation buffer, and the calculation buffer is switched to a new preload buffer.

5. The method according to claim 1, wherein The reasoning dependency data includes: the latest reasoning dependency data related to the latest reasoning result in the input sequence of this round, and the historical reasoning dependency data related to other information in the input sequence of this round; Then, the storage continuity check is performed in the computing resources for the reasoning dependency data related to the input sequence of this round, including: Checking storage continuity of the latest reasoning dependency data and the historical reasoning dependency data in the computing resource; When the storage location of the inference-dependent data is in a discontinuous state in the computing resource, adjusting the storage location to a continuous state by rearranging the inference-dependent data includes: When the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource, the latest reasoning dependency data and the historical reasoning dependency data are adjusted to a continuous storage state by rearranging the latest reasoning dependency data.

6. The method according to claim 5, wherein The storage locations of the latest reasoning dependency data and the historical reasoning dependency data are determined by any of the following methods, which are in a discontinuous state in the computing resource: When the storage location of the latest reasoning dependency data is in a discontinuous state in the computing resource, determining that the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource; When the storage location of the latest reasoning dependency data and the storage location of the historical reasoning dependency data are not adjacent in the computing resource, it is determined that the storage locations of the latest reasoning dependency data and the historical reasoning dependency data are in a discontinuous state in the computing resource.

7. The method according to claim 1, wherein The adjusting the storage locations to a continuous state by rearranging the inference-dependent data includes: For each memory block storing the inference-dependent data, decompose each memory block into a plurality of sub-blocks according to a preset data block size; wherein the size of each sub-block is an integer multiple of the preset data block size; the preset data block size is determined based on the memory access rule of the computing resource; The obtained sub-blocks are rearranged according to the target continuous storage order to adjust the storage positions to a continuous state.

8. The method according to claim 1, wherein Determining the current round of reasoning results based on the loading results includes: Determine the effective data length of the loading result; According to the effective data length, using the first reasoning precision, performing a main reasoning operation on the loaded result to obtain an intermediate result; Using the second reasoning precision, a secondary reasoning operation is performed on the intermediate result to obtain the current round reasoning result; wherein the first reasoning precision is lower than the second reasoning precision.

9. The method according to any one of claims 1 to 2, 4 to 8, characterized in that The activating, according to the hardware status information of the computing resources occupied by the target model, an inference mode that matches the current hardware status of the computing resources includes: Determining a mode decision factor corresponding to the computing resource by quantifying the hardware status information; If the mode decision factor is greater than a first threshold, activating the instruction acceleration mode; If the mode decision factor is less than a second threshold, a paging cache mode is activated; wherein the paging cache mode is used to divide the input sequence of this round into multiple short sequences according to a preset length, and the inference dependency data corresponding to each short sequence is continuously stored using different paging units; If the mode decision factor is not less than the second threshold and not greater than the first threshold, the current reasoning mode is maintained.

10. The method according to claim 9, wherein The hardware status information includes: computing unit utilization, memory access efficiency, and bandwidth utilization; then determining the mode decision factor corresponding to the computing resource by quantifying the hardware status information includes: Determining the utilization rate of the computing unit according to the number of idle cycles of the target computing unit in the computing resource within a preset time period; Determining the memory access efficiency according to the number of accesses to the target memory layer in the computing resource within a preset time period; determining the bandwidth utilization rate according to an amount of bandwidth used by a target storage unit in the computing resource within a preset time period; A weighted sum is performed based on the computing unit utilization, the memory access efficiency, and the bandwidth utilization to obtain the mode decision factor.

11. The method according to any one of claims 1 to 2 and 4 to 8, wherein: Before performing at least one round of reasoning operations, the method further includes: For the task to be inferred, obtaining an initial input sequence corresponding to the task to be inferred; The length of the initial input sequence is compared with a preset length threshold, and it is determined that the length of the initial input sequence is less than the preset length threshold.

12. An inference acceleration device, characterized in that: include: A decision unit, configured to activate, when executing a task to be inferred using a target model, an inference mode that matches the current hardware state of the computing resources occupied by the target model based on the hardware state information of the computing resources; An inference unit, configured to, when the inference mode is the instruction acceleration mode, perform at least one round of inference operations, each round of inference operations including: Performing a storage continuity check in the computing resource for the inference dependency data associated with the input sequence of the current round, wherein the inference dependency data is generated and stored by the inference operation of the previous round; When the storage location of the inference-dependent data is in a discontinuous state in the computing resource, adjusting the storage location to a continuous state by rearranging the inference-dependent data; The specified hardware acceleration instruction is called to load the adjusted inference dependency data, and after determining the result of the current round of inference based on the loading result, the input sequence of the next round is generated according to the result of the current round of inference, until a preset stop condition is reached.

13. An electronic device, characterized in that: The method comprises a processor and a memory, wherein the memory stores a computer program, and when the computer program is executed by the processor, the processor is enabled to perform the steps of any one of the methods of claims 1 to 11.

14. A computer-readable storage medium, characterized in that The method comprises a computer program. When the computer program is run on an electronic device, the computer program is used to enable the electronic device to execute the steps of any one of the methods of claims 1 to 11.

15. A computer program product, characterized in that The method comprises a computer program stored in a computer-readable storage medium; when a processor of an electronic device reads the computer program from the computer-readable storage medium, the processor executes the computer program, so that the electronic device performs the steps of any one of the methods described in claims 1 to 11.

Images