A reasoning request processing method and an electronic device

Abstract

The application discloses a reasoning request processing method and an electronic device, relates to the technical field of model reasoning, and comprises the following steps: generating a topology execution graph based on a pre-training language model initialization or a captured calculation process after deployment, combining the distribution characteristics of historical sequence lengths, constructing a graph pool containing different sequence length exclusive topology execution graphs, and the construction logic of the graph pool can cover a plurality of sequence length scenes frequently appearing in actual applications; meanwhile, based on the construction mode of the execution graph captured by the pre-training language model calculation process, the structural consistency and calculation correctness of the topology execution graphs of different lengths are ensured, so that the reasoning framework can flexibly respond to reasoning requests of different lengths; when the actual sequence length of the reasoning request is greater than a preset length threshold, the reasoning request is split into a plurality of independent sub-blocks, then for each sub-block, a corresponding target topology execution graph is matched from the graph pool and execution is started, and the situation that a dynamic kernel function is returned to due to an ultra-long sequence exceeding a fixed range is avoided.

Description

Technical Field

[0001] This application relates to the field of model reasoning technology, and in particular to a method for processing reasoning requests and an electronic device. Background Technology

[0002] In the rapid development of large-scale language model technology, computational optimization during the inference phase has become crucial for improving system performance. To reduce the startup overhead of graphics processor kernel functions and improve instruction execution efficiency, mainstream inference frameworks widely adopt the Compute Unified Device Architecture Graph (CUDA Graph) technology. This technology captures the computational process as a fixed-topology execution graph during warm-up, allowing direct API calls during the inference phase, effectively reducing scheduling and synchronization overhead.

[0003] However, in practical applications, large language models have a very wide range of dynamic changes, but CUDA Graph requires tensor shapes to be fixed. Mainstream CUDA Graph inference frameworks need to pre-allocate GPU memory for each fixed length to cache the execution graph, and short sequences also need to be filled to the maximum length, which exacerbates the GPU memory pressure. Furthermore, they can only capture topological execution graphs for a few fixed sequence lengths (such as 128, 512, etc.), and select the closest execution during inference. If the length exceeds the predefined maximum length, they fall back to the dynamic kernel function mode, which leads to a decrease in inference performance. Summary of the Invention

[0004] This application provides a method for processing inference requests and an electronic device to at least solve the technical problems of mainstream CUDA Graph inference frameworks in the related art, such as wasted video memory, limited coverage of fixed lengths, and performance degradation of ultra-long sequences due to the fixed requirements of CUDA Graph on tensor shape.

[0005] This application provides a method for processing inference requests, including: obtaining the actual sequence length of the inference request; if the actual sequence length is greater than a preset length threshold, dividing the actual sequence of the inference request into multiple sub-blocks; for each sub-block, determining a target topology execution graph matching the sub-block length from a pre-constructed graph pool; wherein, the graph pool is constructed based on the execution graph obtained after initialization or deployment of a pre-trained language model and the captured computation process, combined with the distribution characteristics of historical sequence lengths, and the graph pool includes topology execution graphs with different sequence lengths; and starting the target topology execution graph to obtain the inference result of the inference request.

[0006] This application also provides a processing apparatus for reasoning requests, comprising:

[0007] The acquisition module is used to obtain the actual sequence length of the inference request;

[0008] The partitioning module is used to divide the actual sequence of the inference request into multiple sub-blocks if the actual sequence length is greater than a preset length threshold.

[0009] The topology execution graph matching module is used to determine the target topology execution graph that matches the length of each sub-block from a pre-built graph pool. The graph pool is constructed based on the execution graphs obtained after the initialization or deployment of the pre-trained language model and the captured computation process, combined with the distribution characteristics of the length of historical sequences. The graph pool includes topology execution graphs with different sequence lengths.

[0010] The inference module is used to initiate the target topology execution graph to obtain the inference result of the inference request.

[0011] This application also provides an electronic device, including: a memory for storing a computer program; and a processor for implementing the steps of the above-described reasoning request processing method when executing the computer program.

[0012] This application also provides a computer-readable storage medium storing a computer program, wherein the computer program, when executed by a processor, implements the steps of the above-described reasoning request processing method.

[0013] This application also provides a computer program product, including a computer program that, when executed by a processor, implements the steps of the above-described reasoning request processing method.

[0014] This application generates a basic execution graph based on the computational flow captured after the initialization or deployment of a pre-trained language model. It then constructs a graph pool containing dedicated topological execution graphs for different sequence lengths, taking into account the distribution characteristics of historical sequence lengths. The construction logic of the graph pool fully considers the distribution characteristics of historical sequence lengths, covering a variety of sequence length scenarios that frequently occur in practical applications, rather than being limited to a few fixed preset lengths. Simultaneously, the execution graph construction method based on the pre-trained language model computational flow ensures the structural consistency and computational correctness of topological execution graphs of different lengths. This allows the inference framework to flexibly respond to inference requests of different lengths, expanding the supported sequence length range and meeting diverse business needs. When the actual sequence length of the inference request exceeds a preset length threshold, the excessively long sequence is split into multiple independent sub-blocks, each with a length within the effective range covered by the graph pool. Subsequently, for each sub-block, the corresponding target topological execution graph is matched from the graph pool and executed in parallel, avoiding the situation where the excessively long sequence exceeds the fixed CUDA Graph support range and falls back to the dynamic kernel function. Attached Figure Description

[0015] To more clearly illustrate the embodiments of this application, the accompanying drawings used in the embodiments will be briefly introduced below. Obviously, the drawings described below are only some embodiments of this application. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort.

[0016] Figure 1 A schematic diagram of the specific hardware architecture on which the execution of a reasoning request processing method provided in this application embodiment depends;

[0017] Figure 2 A method for processing inference requests is provided as an embodiment of this application;

[0018] Figure 3 A flowchart illustrating the graph pool construction process provided in this application embodiment;

[0019] Figure 4 A schematic diagram of the structure of a reasoning request processing device provided in an embodiment of this application;

[0020] Figure 5 This is a schematic diagram of the structure of an electronic device provided in an embodiment of this application. Detailed Implementation

[0021] The technical solutions of the embodiments of this application will be clearly and completely described below with reference to the accompanying drawings. Obviously, the described embodiments are only some embodiments of this application, and not all embodiments. Based on the embodiments of this application, all other embodiments obtained by those of ordinary skill in the art without creative effort are within the protection scope of this application.

[0022] It should be noted that, in the description of this application, the terms "comprising," "including," or any other variations thereof are intended to cover non-exclusive inclusion, such that a process, method, article, or apparatus that comprises a list of elements includes not only those elements but also other elements not expressly listed, or elements inherent to such a process, method, article, or apparatus. The terms "first," "second," etc., in this application are used to distinguish similar objects and are not used to describe a specific order or sequence.

[0023] In the rapid development of Large Language Model (LLM) technology, computational optimization during the inference phase has become a key direction for improving overall system performance. To effectively reduce the kernel launch overhead on Graphics Processing Units (GPUs) while improving instruction pipeline consistency and execution efficiency, current mainstream modern inference frameworks widely adopt CUDA Graph technology.

[0024] The core working principle of CUDA Graph technology is as follows: During the system warm-up phase, a specific computation process (typically the operation process of a multi-layer Transformer module) is captured once as an execution graph with a fixed topology. In the subsequent inference phase, there is no need to repeat the complex construction of the computation process. The pre-captured execution graph can be executed directly by simply calling the "cudaGraphLaunch()" interface. In this way, the scheduling overhead and synchronization overhead in the computation process can be significantly reduced, effectively improving inference efficiency.

[0025] However, in practical applications of large-scale language models for natural language processing tasks, the length of the input sequence (seq_len) is highly dynamic, often ranging from tens of tokens to tens of thousands of tokens. Because CUDA Graph technology imposes strict requirements on tensor shapes during execution graph capture—that is, the shapes of all tensors involved in the computation must remain unchanged—existing inference frameworks typically resort to capturing multiple execution graphs for several fixed sequence lengths (e.g., 128, 512, 2048) when dealing with dynamically long input sequences. Then, during inference, the framework selects the captured execution graph closest to the actual input sequence length for execution.

[0026] The aforementioned static multi-graph processing mechanism based on fixed sequence length has revealed several problems in practical applications, specifically as follows:

[0027] First, there is a serious problem of wasted video memory. For each execution graph captured with a fixed sequence length, a video memory buffer corresponding to that sequence length needs to be pre-allocated. Even if the actual input sequence length is short, it must still be filled to the fixed maximum length corresponding to the selected execution graph. This directly leads to a large amount of video memory resources being ineffectively occupied, and the video memory utilization rate is greatly reduced.

[0028] Secondly, it is difficult to achieve a balance between the number of execution graphs and the sequence length coverage. If too many execution graphs are captured in order to more comprehensively cover different sequence lengths, it will cause a sharp increase in video memory usage and trigger a video memory explosion problem. Conversely, if the number of execution graphs captured is reduced in order to control video memory usage, it will result in a lot of video memory waste when short sequence inputs are used, and it will be unable to support long sequence inputs that exceed the length of the captured execution graphs.

[0029] Third, there is a lack of support for ultra-long sequences. When the length of the input sequence exceeds the maximum sequence length predefined by the system (e.g., 16k, 32k tokens), existing technical solutions can usually only fall back to the dynamic kernel function running mode and can no longer reuse the pre-captured CUDA Graph execution graph. This situation will lead to a significant decrease in inference performance and make it difficult to meet the needs of ultra-long sequence inference scenarios.

[0030] Fourth, execution graph capture is too costly. For large execution graphs corresponding to large language models, each capture operation not only takes a long time, but also occupies a large amount of resident video memory on the GPU during the capture process. This not only affects the efficiency of the execution graph capture stage, but also interferes with the concurrent execution of other tasks on the device, reducing the overall resource utilization of the device.

[0031] It is evident that existing inference frameworks based on CUDA Graph technology have significant shortcomings in efficiently supporting inference scenarios with dynamic and ultra-long sequence lengths. These shortcomings are particularly pronounced in long context inference tasks involving large language models, and have become a key factor restricting the improvement of system performance.

[0032] Furthermore, mainstream deep learning frameworks, such as PyTorch, face similar technical bottlenecks when enabling CUDA Graph technology. These frameworks typically capture and cache a separate CUDA Graph execution graph for each set of computational tasks with a unique input shape. When a new inference request arrives, the system first checks the cache for an execution graph that matches the current input shape; if no matching graph is found, it can only execute the non-static, regular execution mode, resulting in a significant decrease in inference speed.

[0033] Furthermore, each independent CUDA Graph execution graph needs to be bound to its own input and output tensors. When the system needs to support a large range of input sequence lengths (e.g., seq_len∈[1,1024]), hundreds of independent execution graphs may need to be cached to cover different sequence lengths within this range. This undoubtedly puts enormous pressure on GPU memory resources, exceeding the capacity of GPU memory. At the same time, caching a large number of execution graphs also significantly increases the complexity of system management, requiring the design of complex execution graph cache eviction and hit judgment mechanisms, further increasing the difficulty of system design and implementation.

[0034] Most importantly, in application scenarios involving extremely long input sequences (such as 16k or 32k tokens) or frequent changes in input sequence length, the existing technical solution of capturing and caching the execution graph separately according to each input shape has become almost impractical in practice.

[0035] It is evident that the relevant technologies still face insurmountable performance and resource bottlenecks in reasoning scenarios where ultra-long sequence inputs and a wide range of dynamic sequence lengths coexist, failing to meet the actual needs of efficient reasoning for large-scale language models.

[0036] To address all or part of the aforementioned technical problems, this application provides a method for processing inference requests. First, the actual sequence length of the inference request is obtained, avoiding the static memory allocation model based on a maximum fixed length in traditional frameworks. When the actual sequence length exceeds a preset length threshold, the long sequence is divided into multiple sub-blocks, ensuring that the length of each sub-block fits within the efficient range of memory allocation. This avoids the need for excessively large memory for a single long sequence and achieves on-demand allocation of memory resources through sub-blocking, further improving memory utilization efficiency. A pre-constructed version containing multiple adaptations of different lengths is built based on the distribution characteristics of historical sequence lengths. This topology execution graph pool ensures that there are topology execution graphs in the pool that precisely match the lengths of various common sequences (and sub-block lengths), rather than being limited to a few fixed-length options. When processing inference requests, for the entire sequence or the split sub-blocks, the target topology execution graph with the matching length can be directly determined from the graph pool, without having to use an approximate matching method that selects the execution graph with the closest length, thus avoiding the execution efficiency loss caused by length mismatch. At the same time, for long sequences that exceed the predefined maximum length of traditional frameworks, the topology execution graph in the graph pool that is adapted to the sub-block length is reused after splitting the sub-blocks, without having to fall back to the dynamic kernel function mode, maintaining the high performance advantage of the CUDA Graph static execution graph and solving the performance degradation problem during long sequence inference.

[0037] To enable those skilled in the art to better understand the present application, the present application will be further described in detail below with reference to the accompanying drawings and specific embodiments.

[0038] like Figure 1 The diagram shows a specific hardware architecture on which the inference request processing method depends, including a graphics processor, video memory, host memory, data transfer bus, and auxiliary computing unit.

[0039] The GPU's built-in Streaming Multiprocessors (SMs) are responsible for parallel execution of computational logic fused within the CUDA Kernel. This includes fusion operations of operators such as matrix multiplication (MatMul), bias addition (Bias), activation functions (GeLU), and layer normalization (LayerNorm), as well as critical operations such as dynamic boundary judgment and attention calculation. The GPU's multi-stream parallel capabilities support the creation of dedicated CUDA Streams for each tensor shape group, ensuring that the graph capture process does not interfere with each other. Furthermore, the pre-captured execution graph can be directly scheduled through the "cudaGraphLaunch()" interface, aligning with the static topology execution characteristics of CUDAGraph.

[0040] The video memory is used to pre-allocate fixed-shape input / output buffers (input_buf / output_buf), providing dedicated storage space for each Graph instance grouped by tensor shape to ensure the tensor shape is fixed during the capture phase; and to maintain a global Key / Value Cache (KVCache) buffer to store the intermediate context state during block inference of ultra-long sequences, achieving cross-block data continuity through attention offsets; it is also used to cache reusable "cudaGraphExec_t" objects to avoid repeated capture or reconstruction of the Graph at runtime, reducing video memory usage and scheduling latency.

[0041] The host memory is responsible for storing the input sequence data to be inferred, the sequence length distribution statistics of historical requests, and performance metrics (such as graph hit rate, GPU memory usage, latency, etc.) during the offline adaptive optimization phase. Before inference starts, the input data in the host memory is copied to the corresponding buffer in the GPU memory; after inference is completed, the output results in the GPU memory are returned to the host memory for subsequent processing or returned to the user.

[0042] The data transfer bus is responsible for data transfer between the host and the GPU. Its bandwidth directly affects the efficiency of input data copying, output result return, and cross-block data transfer in the KVCache. In ultra-long sequence block inference, the data transfer bus needs to efficiently complete the rapid migration of data from each sub-block between the host memory and the GPU memory to ensure the continuity and low latency of block execution and avoid impacting overall inference performance due to data transfer bottlenecks.

[0043] The auxiliary computing unit includes the GPU's instruction scheduling unit and cache hierarchy: the instruction scheduling unit is responsible for distributing the execution logic of the fused kernel to each stream processor to achieve efficient scheduling of parallel computing; the cache hierarchy is used to temporarily store intermediate data during kernel execution, reducing the frequency of access to video memory and improving computing efficiency, especially playing a key role in high-frequency operations such as attention mechanism calculation and dynamic boundary judgment.

[0044] Embodiments of this application provide a method for processing inference requests, such as... Figure 2 As shown, the method includes the following steps:

[0045] S201. Obtain the actual sequence length of the inference request.

[0046] Inference requests refer to requests initiated by the client to the pre-trained language model, demanding that the model perform calculations on the input data and return results. These requests are the triggering conditions for the inference process in the topological execution graph. The input data is a sequence stored in tensor form in a natural language task, with the core feature being the sequence length (seq_len). For example, requests from users to the pre-trained language model such as "continuing text based on 1800 tokens" or "parsing the semantics of a document using 300 tokens" are both inference requests.

[0047] In some embodiments, after obtaining the actual sequence length of the inference request, it is first determined whether the actual sequence length is less than or equal to a preset length threshold. If the actual sequence length is less than or equal to the preset length threshold, a target topology execution graph that matches the actual sequence length and has the smallest tensor shape is determined from a pre-built graph pool; then, the target topology execution graph is started to obtain the inference result of the inference request.

[0048] The preset length threshold (L_max) is 1024. For example, if L_real=75, the length range is [65, 128], and the corresponding tensor shape (Shape) of 128 is selected for the topology execution graph. Shape=128 means that the tensor corresponds to a sequence length of 128 tokens.

[0049] The input data is then copied to the input buffer corresponding to Shape=128, occupying the storage space of the first 75 tokens. The "actual_seq_len" parameter is set to L_real=75, and the corresponding target topology execution graph is executed through "cudaGraphLaunch()" to complete the inference calculation.

[0050] The above embodiments compare the actual sequence length with a preset length threshold to select suitable processing paths, avoiding redundant operations for inference requests with short sequence lengths and reducing unnecessary resource consumption. Matching the target topology execution graph with the actual sequence length ensures adaptation to inference requirements, avoiding computational waste caused by topology mismatch. At the same time, the tensor shape of the target topology execution graph is minimal, which can minimize memory consumption and reduce data transfer overhead, making the inference process more lightweight.

[0051] S202. If the actual sequence length is greater than the preset length threshold, the actual sequence of the inference request will be divided into multiple sub-blocks.

[0052] For example, assuming the actual sequence length L_real = 1800, which is greater than L_max = 1024, the actual sequence length is divided into sub-blocks (Chunk) 0 and Chunk 1. The length of Chunk 0 is 1024, and the length of Chunk 1 is 776.

[0053] S203. For each sub-block, determine the target topology execution graph that matches the length of the sub-block from the pre-built graph pool.

[0054] The graph pool comprises execution graphs captured during the initialization phase or deployment pre-defined phase of the pre-trained language model, constructed by combining the distribution characteristics of historical sequence lengths. It includes topological execution graphs (CUDA Graphs) of different sequence lengths. This allows for rapid matching of the optimal execution graph to the actual sequence length, eliminating the need for real-time construction and shortening inference preparation time. All instances in the Graph pool are static graphs pre-captured during the offline warm-up phase; dynamic capture or graph stitching is not required during inference, and they can be directly invoked via "cudaGraphLaunch()".

[0055] The topology execution graph is the execution graph obtained by capturing the computation process after the pre-trained language model is initialized or deployed. It extracts and solidifies the computation process into a static topology structure, which can be directly reused during model inference to avoid redundant scheduling overhead.

[0056] In some embodiments, the length of the tail block among multiple sub-blocks is less than or equal to a preset length threshold, and the lengths of the other sub-blocks are equal to the length threshold. When the actual sequence length (L_real) exceeds the maximum pre-capture length (L_max) of the captured graph, the input sequence is divided into multiple consecutive sub-blocks (chunks), with the first k blocks having a length of L_max, and the tail sub-blocks having less than one block using a graph with a smaller shape.

[0057] When performing step S203, the attention offset is determined for each sub-block by reading the historical context of the global buffer. Then, the target topology execution graph is determined from the pre-built graph pool based on the sub-block length and the attention offset.

[0058] The global buffer, serving as a unified context storage carrier, continuously records the previous sequence processing records, user interaction states, and data association information of the inference task, ensuring that the sub-block can obtain complete historical dependencies.

[0059] Specifically, for each sub-block after division, the corresponding historical context fragment is quickly retrieved through the buffer index, and key information related to attention calculation is extracted, including the length of the historical sequence, token position mapping, attention weight distribution, etc. Then, the attention offset is calculated through a preset algorithm: the offset is used to correct the attention direction of the sub-block when it is processed independently and to compensate for the context breakage caused by the block division. The specific calculation will combine the sequence length of the historical context, the starting position of the sub-block in the overall sequence and the decay coefficient of the attention mechanism to obtain a quantitative value that represents the strength of the association between the sub-block and the historical context.

[0060] Subsequently, the calculated attention offset is combined with the actual length of the current sub-block. The sub-block length determines the basic computational scale required to support the execution graph, while the attention offset limits the contextual association capabilities that the execution graph must be compatible with. For the pre-built graph pool, the system traverses the topological execution graphs with different sequence lengths and attention configurations contained in the pool, and filters out candidate execution graphs that simultaneously satisfy "sub-block length adaptation" and "attention offset compatibility". The optimal solution is determined from the candidate execution graphs as the target topological execution graph, prioritizing the execution graph with the highest matching degree with sub-block length and attention offset and the smallest tensor shape, thereby maximizing resource utilization efficiency while ensuring inference accuracy.

[0061] Following the previous example, the length of Chunk0 is 1024, which corresponds to the target topology execution graph G6 with Shape=1024 in the graph pool. The length of Chunk1 is 776, which falls into the length range [513-1024] of the target topology execution graph G6. Therefore, the target topology execution graph G6 is reused to reduce graph switching overhead.

[0062] In the above embodiments, the first k blocks use a uniform maximum length L_max to ensure processing efficiency, while the last few sub-blocks (less than one block) are adapted to a smaller shape graph, avoiding resource waste caused by uniform block division and ensuring that sub-blocks of different lengths receive precisely matched computational resources. By reading the historical context of the global buffer to determine the attention offset, it is ensured that the sub-blocks after block division can still maintain complete contextual association, avoiding semantic breakage problems that may be caused by block processing, and ensuring the accuracy and coherence of the inference results. Block division reduces the computational pressure of single-round inference. Combined with the efficient matching and context association mechanism of the graph pool, when processing long sequence inference requests, it can ensure that the inference accuracy is not lost, while improving the utilization of computational resources and the inference response speed.

[0063] S204. Start the target topology execution graph to obtain the inference result of the inference request.

[0064] For each sub-block, the graph of the corresponding shape is selected sequentially for inference. After each execution, the intermediate state (such as the key / value cache) is written to the global memory buffer. When the next sub-block is executed, the historical context in the global buffer is read and the correct attention offset (such as kv_cache_offset) is passed in to ensure the continuity of the attention mechanism across blocks. For sub-blocks that are not long enough to be considered a block, a smaller shape graph that matches its length is selected for processing.

[0065] Taking L_real=1800 (L_max=1024) as an example, the system will divide it into two sub-blocks with lengths of 1024 and 776, both of which will be executed using a Graph with Shape=1024. The actual computation lengths of 1024 and 776 are specified by the actual_seq_len parameter, respectively. At the same time, kv_cache_offset=1024 ensures that the second sub-block can read the KVCache of the previous sub-block, thus achieving logically continuous inference.

[0066] During the execution of each sub-block, the pre-captured graph of the corresponding shape is invoked. At the same time, a global key / value (KV) cache buffer and attention offset are maintained through GPU memory to ensure cross-block context continuity and achieve semantically consistent inference. The efficient pre-captured CUDA graph is reused, avoiding the scheduling and synchronization overhead of dynamic kernel mode. Even with block execution, it is still significantly better than dynamic mode. There is no need to recapture a huge graph. Only a limited number of pre-captured graphs are needed to support sequences of arbitrary length, reducing memory usage and graph capture costs.

[0067] Following the previous example, Chunk0 copies 1024 tokens to the G6 input buffer (input_buf[6]), sets actual_seq_len=1024, and calls "cudaGraphLaunch()". The core intermediate state (KV Cache) of the pre-trained language model attention calculation is written to the global cache to preserve the historical context for subsequent blocks. When Chunk1 is executed, the next 776 tokens are copied to the first 776 bits of the G6 input buffer, and actual_seq_len=776 is set to skip invalid calculations. At the same time, when the attention calculation is performed, kv_cache_offset=1024 is set so that the current block can read the KV Cache of Chunk0 from the 1024th position of the global KV cache to ensure the continuity of the cross-block context. Then, cudaGraphLaunch() is called to update the KV Cache to the 1800th position.

[0068] like Figure 3 As shown, the graph pool construction process provided in this application embodiment includes the following steps:

[0069] S301, Obtain the historical sequence of the historical reasoning request.

[0070] In some embodiments, after executing step S301, the operator sequence of the pre-trained language model during the inference process is first obtained, and static program analysis is performed on the operator sequence to identify the dynamic dimension in the input tensor that is related to the length of the historical sequence, and to determine the change characteristics of the dynamic dimension in different historical inference requests; a fusion kernel function template is generated according to the change characteristics, and the fusion kernel function template is edited into an executable device function; wherein, the fusion kernel function template is adapted to the sequence processing requirements of different length ranges.

[0071] Specifically, static program analysis is performed on the operator sequences of the pre-trained language model to identify the lengths of historical sequences with dynamic changing characteristics in the input tensor. A fusion kernel function template is generated, which integrates the logic of multiple related operators such as matrix multiplication (MatMul), bias addition (Bias), activation function (GeLU), and layer normalization (LayerNorm) into a single GPU function, reducing scheduling overhead and data transmission latency between operators and improving computational efficiency.

[0072] The key is that runtime shape boundary judgment logic is inserted into the template, introducing the actual sequence length (actual_seq_len) as a parameter, allowing the kernel to dynamically adjust the calculation range based on the actual input length at runtime. For example, the kernel explicitly receives the actual sequence length as an input parameter, which is dynamically passed in during inference based on the specific requested sequence length. In the kernel's thread execution logic, the code snippet "int seq_idx = tid / hidden_size; if (seq_idx >= actual_seq_len) return;" calculates the sequence index (seq_idx) corresponding to the current thread and checks whether this index exceeds the actual sequence length. If it does, the invalid calculation of the current thread is terminated directly, achieving dynamic pruning of the filling region.

[0073] The generated kernel template will be compiled into executable device functions using the nvcc compiler or offline compilation tools, with fixed function pointers. It's worth noting that all fixed-shape groups will share the same kernel logic. There's no need to compile a separate kernel for each shape; instead, during the subsequent graph capture and execution phases, the corresponding startup parameters (such as actual_seq_len, hidden_size, etc.) are passed in according to the needs of different shape groups. This adapts to the sequence processing requirements of different length ranges, ensuring the uniformity of the computational logic while improving kernel reusability and system flexibility.

[0074] In an online deployment scenario, the system will call the nvcc compiler, set the compilation options according to the target GPU architecture, compile the .cu format Kernel template into PTX intermediate code, and then further assemble it into binary device code (cubin). In an offline deployment scenario, the cubin file can be compiled in advance to avoid the time-consuming online compilation.

[0075] The pseudocode is as follows:

[0076]

[0077] In some embodiments, when performing step S301, the length of the historical sequence can be counted and presented in the form of a histogram, with the horizontal axis of the histogram representing the sequence length interval and the vertical axis representing the number of requests for the corresponding interval.

[0078] The historical sequence length (seq_len) of the pre-trained language model during inference is statistically analyzed, dividing the historical sequence length into multiple sequence length intervals. This is presented as a histogram, with the horizontal axis representing the sequence length interval and the vertical axis representing the number of requests within the corresponding interval. The histogram visually reflects which length type of requests are more frequent.

[0079] Alternatively, during step S301, the sequence length distribution of historical requests can be updated using a sliding window. The sliding window always retains the latest request data and removes expired data, ensuring that the distribution pattern keeps up with load changes and avoiding inaccurate subsequent segmentation due to outdated early request distributions. For example, the size of the sliding window is a preset number of historical inference requests, and the sliding window size is set to 1000 requests.

[0080] For example, the seq_len distribution of 1000 requests is: [1,32] (150 times), [33,64] (300 times), [65,128] (350 times), [129,256] (150 times), [257,1024] (50 times).

[0081] S302. Divide the historical sequence into multiple non-overlapping continuous intervals to obtain subsequences.

[0082] In some embodiments, an executable device function is invoked to segment the historical sequence. Specifically, if the length of the historical sequence exceeds a preset length threshold, it indicates that the historical sequence is an excessively long sequence. The device function then divides the historical sequence into multiple non-overlapping, continuous subsequences. Independent computation is then performed on each subsequence. Contextual continuity between the multiple subsequences can be maintained by reading the attention offset from the global buffer, thus preserving semantic consistency when the pre-trained language model performs inference based on the historical sequence.

[0083] In some embodiments, heuristic algorithms (such as greedy segmentation), clustering algorithms (such as K-means), or optimization strategies based on a joint memory-delay objective function are used to divide the historical sequence length into multiple non-overlapping continuous intervals as subsequences. Each interval corresponds to a fixed upper limit length.

[0084] Optionally, the length of the historical sequence can be divided into multiple non-overlapping continuous intervals based on the frequency of historical inference requests. The heuristic algorithm groups sequences with high request frequencies into the same interval. For example, if requests of length [65, 128] account for 35%, it is divided into a separate interval (Shape=128) to avoid increasing the number of graphs due to splitting this interval. By prioritizing the grouping of frequently requested seq_len sequences into the same interval, the number of graphs is reduced, thus lowering memory usage. Simultaneously, the length span of a single interval is controlled to avoid excessive padding caused by short inputs matching graphs with excessively large shapes, reducing memory waste and improving execution efficiency. For example, the interval [65, 128] (350 times / 1000 times) with a high request rate is divided into a separate interval (Shape=128), reducing the number of graphs (no need to build a separate graph for each length within 65-128) and avoiding padding waste caused by short inputs (such as 65) matching a 256 Shape Graph, achieving a balance between memory usage and execution efficiency.

[0085] Based on the clustering algorithm, the historical sequence length is clustered to obtain multiple non-overlapping continuous intervals of subsequences corresponding to multiple clusters. The clustering algorithm clusters historical inference requests with similar sequence lengths into clusters, with each cluster corresponding to an interval. For example, the two high-frequency clusters [1,32] (150 times) and [33,64] (300 times) are divided into separate intervals to reduce the filling waste caused by cross-cluster matching.

[0086] Based on a joint memory-latency objective function, the strategy first calculates the memory usage and execution latency corresponding to different partitioning methods, and then finds the optimal balance between the two through function solving. Next, the historical sequence length is divided according to the partitioning method corresponding to the optimal balance point, resulting in multiple non-overlapping continuous interval subsequences. Specifically, a larger number of graphs leads to higher memory usage; a larger interval span requires more padding with short inputs, resulting in higher latency. For example, when memory is limited, smaller intervals are appropriately merged to reduce the number of graphs; when latency is sensitive, larger regions are split to reduce padding, ensuring that the partitioning result does not exceed memory limits while meeting inference speed requirements.

[0087] The resulting non-overlapping continuous intervals are shown in Table 1:

[0088] Table 1

[0089]

[0090] S303. Based on the upper limit length of multiple non-overlapping continuous intervals, the computation process of the pre-trained language model is converted into a topological execution graph with a fixed tensor shape.

[0091] The upper limit length of non-overlapping continuous intervals is the basis for grouping tensor shapes. The Graph pool uses a smart shape grouping strategy to capture only one Graph instance for each non-overlapping continuous interval's upper limit length, eliminating the need to construct a separate Graph for each possible sequence length and reducing the number of Graphs. For example, covering the length range [1~1024] requires only 4~6 Graph instances, keeping total GPU memory usage manageable and completely avoiding the memory explosion problem caused by capturing hundreds of Graphs in traditional solutions.

[0092] Based on the aforementioned embodiments, after compilation, the executable device function is called to allocate input / output buffers (in / output_buf) for non-overlapping continuous intervals, and the processor memory addresses of the input / output buffers are recorded. Then, placeholder data is copied to the input buffer, and the device function is started and the input / output buffers are bound. After the pre-trained language model executes the computation process, the computation process is converted into a topological execution graph with a fixed tensor shape.

[0093] Specifically, each Shape group Gi is first allocated an independent input / output buffer, with the format [1, Shape, hidden_size], and the GPU memory address of the buffer is recorded. Here, "1" represents a fixed batch_size, "Shape_i" is the upper limit length of the interval of the group, and "hidden_size" is the fixed hidden layer dimension of the model. This fixed-dimensional design perfectly matches the capture requirements of CUDA Graph, avoiding the capture failure problem caused by dynamic Shapes from the root.

[0094] For each group of Gi, input and output buffers are pre-allocated in GPU memory according to fixed dimensions, and a fixed-shape computing environment is customized to ensure that the Graph can be repeatedly called after capture.

[0095] Then, the CUDA Graph capture interface (cudaStreamBeginCapture / cudaStreamEndCapture) is called to start the executable device function (i.e., the compiled kernel function template) in the dedicated stream (CUDA stream). At this time, the actual_seq_len parameter is passed in as the upper limit length of the current Shape group (e.g., 64 when Shape=64), and the buffer pointer of the group is bound. After the capture is completed, the generated cudaGraph_t object is instantiated into an executable cudaGraphExec_t object, and the mapping relationship between the instance and the kernel function address and the buffer address is recorded.

[0096] Copying all-zero dummy data to input_buf[i] is intended to provide placeholder data for Graph capture. When capturing a CUDA Graph, a computation process needs to be actually executed once to record the instructions. Dummy data does not need to have real semantics, but only needs to meet the tensor shape requirements.

[0097] After starting the stream capture mode via "cudaStreamBeginCapture", the pre-generated fused kernel (fused_kernel_ptr) is launched. At this time, the parameters passed in strictly match the characteristics of the group: the input / output pointers are bound to the buf[i] of the group, actual_seq_len is set to Shape_i (the upper limit length of the group), and the appropriate grid (number of thread blocks) and block (number of threads per block) dimensions are calculated according to Shape_i and hidden_size (e.g., when hidden_size=768, block is set to 256 to adapt to the GPU warp size); after the capture is completed, the Graph[i] object that records the complete computation process is obtained through "cudaStreamEndCapture".

[0098] Calling "cudaGraphInstantiate" transforms Graph[i] into an executable Graph_exec[i] object. This step is similar to compiling the graph structure into a runnable program. During subsequent inference, there is no need to re-parse the graph; execution can be started directly through Graph_exec[i].

[0099] Since all Shape groups share the same compiled Kernel, different instances are distinguished only by startup parameters (such as actual_seq_len and buffer pointers), which reduces code redundancy and ensures that each instance can accurately adapt to the computational needs of the corresponding Shape, ultimately forming a reusable Graph instance pool, laying the foundation for fast runtime invocation.

[0100] Ultimately, the tuples (Graph_exec[i], input_buf[i], output_buf[i]) of all groups are stored in the Graph pool. The value of this design lies in two aspects: efficient reuse and controllable resources. On the one hand, each Graph_exec[i] is a pre-built static execution unit, which can be directly called when the corresponding group is matched during subsequent inference, without repeated capture (avoiding capture time and a surge in GPU memory usage). On the other hand, the pre-allocated buffer and Graph instance are completed during the initialization phase, and the total GPU memory usage can be calculated in advance (e.g., 6 groups only occupy <2GB), which will not cause GPU memory explosion due to dynamic capture like in traditional multi-graph mechanisms.

[0101] The above embodiments pre-build efficient execution templates for each length interval, which not only meets the static requirements of CUDA Graph, but also provides a low-overhead and highly available execution foundation for subsequent dynamic and ultra-long sequence inference.

[0102] To adapt to dynamic changes in different workloads, the system also introduces an offline adaptive optimization mechanism. This process is completely separate from the online inference path, avoiding impact on service stability. The system periodically (e.g., every 5 minutes) collects various performance metrics, including the call hit rate of each graph, total memory usage, block inference frequency, and end-to-end latency. During service idle periods (e.g., within 10 seconds without requests) or in a preset reconfiguration window, the system analyzes the collected metrics. If it detects that some graphs have not been used for a long time (e.g., no calls for nearly 1 hour), the proportion of excessively long block requests has increased significantly (e.g., from 5% to 30%), or the current grouping is causing memory waste and performance degradation, a regrouping process will be triggered. At this point, the system will resample the sequence length distribution of the latest 1000 requests, recalculate the optimal Shape grouping (e.g., adjust to [1–64], [65–128], [129–256], [257–512], [513–2048]), release the old Graph pool, re-execute the Graph capture process based on the new grouping, build a new Graph pool, and gradually switch to the new pool in subsequent inference requests to achieve continuous self-optimization of system performance.

[0103] In the process of constructing the graph pool described above, firstly, during the inference initialization phase or deployment warm-up phase, the distribution of the sequence length (seq_len) of historical inference requests is automatically sampled. Then, through a greedy algorithm, K-means clustering, or joint optimization of the objective function based on GPU memory and latency, the sequence length range is divided into multiple non-overlapping continuous intervals, each bound to a fixed topology execution graph. Next, static program analysis is performed on the model operator sequence to identify the dynamic seq_len dimension of the input tensor, automatically generating a fused CUDAKernel template. This template integrates the core operator logics such as MatMul, Bias, GeLU, and LayerNorm into a single GPU function and inserts runtime Shape boundary judgment logic. Through the actual_seq_len parameter and the thread exit mechanism of "seq_idx>=actual_seq_len", dynamic adaptation to inputs of arbitrary length within the interval is achieved. This kernel is compiled into an executable device function by nvcc or offline. All Shape groups share the same kernel logic, requiring only different parameters to be passed at startup.

[0104] Subsequently, for each fixed-shape group, the system performs CUDA Graph pre-capture: Input buffers (input_buf) and output buffers (output_buf) of corresponding shapes are pre-allocated in GPU memory for each group, with the buffer shapes following the format [1, Shape_i, hidden_size]. A dedicated CUDA stream is created for each group to avoid mutual interference during the capture process. On the dedicated stream, all-zero dummy data is copied to the input buffer. "cudaStreamBeginCapture()" is called to start the capture, initiating the fusion kernel and binding the current group's buffer. After executing the computation process, "cudaStreamEndCapture()" completes the capture, obtaining a Graph object. The Graph object is then instantiated into an executable cudaGraphExec_t object. Finally, the Graph_exec instance and the input / output buffers are stored as tuples in the graph pool, completing the construction of the topology execution graph for all groups.

[0105] The graph pool uses an intelligent shape grouping strategy to capture only one graph instance for the upper limit of each length interval, eliminating the need to build a separate topology execution graph for each possible sequence length, thus significantly reducing the number of topology execution graphs. At the same time, the dynamic pruning mechanism of the kernel allows each topology execution graph to adapt to inputs of any length within the interval, eliminating the need for redundant padding for short inputs and further reducing memory waste.

[0106] In summary, this application addresses the issue of wasted GPU memory in mainstream CUDA Graph inference frameworks by sampling the sequence length distribution of historical inference requests during the warm-up phase. Through greedy algorithms, K-means clustering, or joint memory-latency optimization, it divides the data into non-overlapping continuous intervals covering the entire dynamic range. Only a fixed-tensor-shaped topological execution graph is captured for the upper limit of each interval, eliminating the need to construct a separate graph for each possible length. Furthermore, it integrates the explicitly introduced `actual_seq_len` runtime parameter from the CUDA Kernel, dynamically skipping invalid computation units exceeding the actual sequence length in the thread execution logic through boundary checks using `if(seq_idx>=actual_seq_len)return;`.

[0107] To address the issue that mainstream CUDA Graph inference frameworks only support a few preset lengths, resulting in limited coverage for fixed lengths, this system automatically divides intervals based on historical input distribution during the grouping phase. This ensures maximum length coverage with the fewest possible graphs. At runtime, topology execution graphs of the same shape can adapt to inputs of any length within an interval via the `actual_seq_len` parameter, eliminating the need to capture a separate topology execution graph for each length and expanding the effective coverage range. Simultaneously, the system periodically collects metrics such as call hit rate and memory usage for each topology execution graph. During service idle periods, it resamples the latest input distribution, dynamically adjusts shape grouping boundaries, and updates the topology execution graph pool, ensuring that coverage continuously adapts to load changes and completely resolving the issue of insufficient coverage for fixed preset lengths.

[0108] To address the performance degradation issue of ultra-long sequences in mainstream CUDA Graph inference frameworks, this method, when the input sequence length L_real exceeds the maximum length L_max of the pre-captured topology execution graph, avoids backing down the dynamic kernel function. Instead, it automatically divides the ultra-long sequence into multiple contiguous sub-blocks (the first k blocks are of length L_max, and the last block is ≤ L_max). Each sub-block reuses an existing efficient Graph instance in the graph pool for inference. By maintaining a global Key / Value Cache (KVCache) buffer in GPU memory, combined with attention offsets (e.g., kv_cache_offset=1024), the continuity of cross-block context is ensured. The KVCache of the previous sub-block is written to the global buffer, and the historical context is read through the offset when the next sub-block is executed, achieving semantically consistent continuous inference. The actual_seq_len parameter is used to adapt to different sub-block lengths, eliminating the need to build a giant topology execution graph or back down the dynamic mode. This maintains the low scheduling overhead advantage of CUDA Graph while solving the performance degradation problem of ultra-long sequences.

[0109] Through the above description of the embodiments, those skilled in the art can clearly understand that the methods according to the above embodiments can be implemented by means of software plus necessary general-purpose hardware platforms. Of course, they can also be implemented by hardware, but in many cases the former is a better implementation method.

[0110] like Figure 4 As shown, embodiments of this application also provide a processing apparatus for inference requests, the apparatus comprising:

[0111] Module 401 is used to obtain the actual sequence length of the inference request;

[0112] The partitioning module 402 is used to divide the actual sequence of the inference request into multiple sub-blocks if the actual sequence length is greater than a preset length threshold.

[0113] The topology execution graph matching module 403 is used to determine the target topology execution graph that matches the length of each sub-block from a pre-built graph pool. The graph pool is constructed based on the execution graph obtained by capturing the computation process after the initialization or deployment of the pre-trained language model, combined with the distribution characteristics of the length of historical sequences. The graph pool includes topology execution graphs with different sequence lengths.

[0114] The inference module 404 is used to initiate the target topology execution graph to obtain the inference result of the inference request.

[0115] As an optional implementation provided in this application, the topology execution graph matching module 403 is used to: read the historical context in the global buffer for each sub-block and determine the attention offset; determine the target topology execution graph from the pre-built graph pool according to the sub-block length and the attention offset; wherein, the length of the tail block among multiple sub-blocks is less than or equal to the length threshold, and the length of other sub-blocks other than the tail block is equal to the preset length threshold.

[0116] As an optional implementation provided in this application, the topology execution graph matching module 403 is further configured to: if the actual sequence length is less than or equal to a preset length threshold, select a target topology execution graph from the pre-constructed graph pool that matches the actual sequence length and has the smallest tensor shape.

[0117] As an optional implementation provided in this application, the device further includes a graph pool construction module, used for: obtaining a historical sequence of historical inference requests; dividing the historical sequence into multiple non-overlapping continuous intervals as subsequences; and converting the computation process of the pre-trained language model into a topological execution graph with a fixed tensor shape according to the upper limit length of the multiple non-overlapping continuous intervals.

[0118] As an optional implementation provided in this application, after obtaining the historical sequence of historical inference requests, and before dividing the historical sequence into multiple non-overlapping continuous interval subsequences, the graph pool construction module is further configured to: obtain the operator sequence of the pre-trained language model during the inference process; perform static program analysis on the operator sequence to identify the dynamic dimension in the input tensor related to the length of the historical sequence, and determine the variation characteristics of the dynamic dimension in different historical inference requests; generate a fusion kernel function template based on the variation characteristics, and compile the fusion kernel function template into an executable device function; wherein, the fusion kernel function template adapts to the sequence processing requirements of different length intervals; correspondingly, when the length of the historical sequence is greater than a preset length threshold, the historical sequence is divided into multiple non-overlapping continuous interval subsequences.

[0119] As an optional implementation provided in this application, when the graph pool construction module obtains the historical sequence of historical inference requests, it is used to: count the length of the historical sequence and present it in the form of a histogram; the horizontal axis of the histogram is the sequence length interval, and the vertical axis is the number of requests in the corresponding interval; or, update the distribution of the historical sequence length through a sliding window, the size of which is a preset number of historical inference requests.

[0120] As an optional implementation provided in this application, when the graph pool construction module divides the historical sequence into multiple non-overlapping continuous interval subsequences, it is used to: divide the historical sequence according to the frequency of historical inference requests to obtain multiple non-overlapping continuous interval subsequences; or, cluster the length of the historical sequence based on a clustering algorithm to obtain multiple non-overlapping continuous interval subsequences corresponding to multiple clusters.

[0121] As an optional implementation method provided in this application, when the graph pool construction module divides the historical sequence into multiple non-overlapping continuous interval subsequences, it is used to: calculate the video memory usage and execution latency corresponding to different division methods, and determine the optimal balance point between video memory usage and execution latency; divide the historical sequence according to the division method corresponding to the optimal balance point to obtain multiple non-overlapping continuous interval subsequences.

[0122] As an optional implementation provided in this application, the graph pool construction module, when converting the computation process of the pre-trained language model into a topological execution graph with a fixed tensor shape based on the upper limit length of multiple non-overlapping continuous intervals, is used to: allocate input / output buffers for the non-overlapping continuous intervals and record the processor memory addresses of the input / output buffers; the dimensions of the input / output buffers include the upper limit length of the non-overlapping continuous intervals; copy placeholder data to the input buffers; start executable device functions and bind the input / output buffers, and convert the computation process into a topological execution graph with a fixed tensor shape after the pre-trained language model executes the computation process.

[0123] For a description of the features in the embodiment corresponding to the inference request processing device, please refer to the relevant description in the embodiment corresponding to the inference request processing method, which will not be repeated here.

[0124] like Figure 5 As shown, embodiments of this application also provide an electronic device, including a memory 501 and a processor 502. The memory 501 stores a computer program, and the processor 502 is configured to run the computer program to perform the steps in any of the above-described inference request processing method embodiments.

[0125] Embodiments of this application also provide a computer-readable storage medium storing a computer program configured to execute the steps in any of the above-described inference request processing method embodiments when running.

[0126] In one exemplary embodiment, the aforementioned computer-readable storage medium may include, but is not limited to, various media capable of storing computer programs, such as a USB flash drive, read-only memory (ROM), random access memory (RAM), portable hard disk, magnetic disk, or optical disk.

[0127] Embodiments of this application also provide a computer program product, which includes a computer program that, when executed by a processor, implements the steps in any of the above-described inference request processing method embodiments.

[0128] Embodiments of this application also provide another computer program product, including a non-volatile computer-readable storage medium storing a computer program, which, when executed by a processor, implements the steps in any of the above-described inference request processing method embodiments.

[0129] Those skilled in the art will further recognize that the units and algorithm steps of the various examples described in conjunction with the embodiments disclosed herein can be implemented in electronic hardware, computer software, or a combination of both. To clearly illustrate the interchangeability of hardware and software, the components and steps of the various examples have been generally described in terms of functionality in the foregoing description. Whether these functions are implemented in hardware or software depends on the specific application and design constraints of the technical solution. Those skilled in the art can use different methods to implement the described functions for each specific application, but such implementation should not be considered beyond the scope of this application.

[0130] The foregoing has provided a detailed description of a reasoning request processing method and electronic device. Specific examples have been used to illustrate the principles and implementation methods of this application. The descriptions of the embodiments above are only intended to help understand the method and core ideas of this application. It should be noted that those skilled in the art can make various improvements and modifications to this application without departing from its principles, and these improvements and modifications also fall within the protection scope of the claims of this application.

Claims

1. A method for processing inference requests, characterized in that, The method comprises: obtaining an actual sequence length of an inference request; if the actual sequence length is greater than a preset length threshold, dividing the actual sequence of the inference request into a plurality of sub-blocks; for each sub-block, determining a target topology execution graph matching the sub-block length from a pre-constructed graph pool; wherein the graph pool is constructed based on a pre-trained language model initialization or an execution graph obtained by capturing a computing process after deployment, in combination with the distribution characteristics of historical sequence lengths; the graph pool comprises topology execution graphs of different sequence lengths; starting the target topology execution graph to obtain an inference result of the inference request; the method further comprises:

2. The method of claim 1, wherein, if the actual sequence length is less than or equal to the preset length threshold, selecting a target topology execution graph matching the actual sequence length and having the smallest tensor shape from the pre-constructed graph pool. The construction process of the graph pool comprises:

3. The method of claim 1, wherein, obtaining a historical sequence of a historical inference request; dividing the historical sequence to obtain a plurality of non-overlapping continuous interval subsequences; converting the computing process of the pre-trained language model into a topology execution graph with fixed tensor shape according to the upper limit length of the plurality of non-overlapping continuous intervals. The method further comprises:

4. The method of claim 3, wherein, obtaining an operator sequence of the pre-trained language model in the inference process; performing static program analysis on the operator sequence, identifying dynamic dimensions in the input tensor related to the historical sequence length, and determining the variation characteristics of the dynamic dimensions in different historical inference requests; generating a fusion kernel function template according to the variation characteristics, and compiling the fusion kernel function template into an executable device function; wherein the fusion kernel function template adapts to the sequence processing requirements of different length intervals; correspondingly, the method further comprises: calling the device function and executing, and in the case that the historical sequence length is greater than the preset length threshold, dividing the historical sequence into a plurality of non-overlapping continuous interval subsequences. The method further comprises:

5. The method of claim 3, wherein, statistically analyzing the historical sequence length and presenting it in the form of a histogram; the horizontal axis of the histogram is the sequence length interval, and the vertical axis is the number of requests in the corresponding interval; or, updating the distribution of the historical sequence length through a sliding window; the size of the sliding window is a preset number of historical inference requests. ​ 6. The method of claim 3, wherein, The process of dividing the historical sequence into multiple non-overlapping continuous intervals to obtain subsequences includes: The historical sequence is divided according to the frequency of the historical reasoning requests to obtain subsequences of the multiple non-overlapping continuous intervals; Alternatively, clustering algorithms can be used to cluster the length of historical sequences, resulting in multiple non-overlapping continuous intervals corresponding to multiple clusters.

7. The method of claim 3, wherein, The process of dividing the historical sequence into multiple non-overlapping continuous intervals to obtain subsequences includes: Calculate the video memory usage and execution latency corresponding to different partitioning methods, and determine the optimal balance point between the video memory usage and the execution latency; The historical sequence is divided according to the partitioning method corresponding to the optimal equilibrium point to obtain the subsequences of the multiple non-overlapping continuous intervals.

8. The method of claim 4, wherein, The step of converting the computation process of the pre-trained language model into a topological execution graph with a fixed tensor shape based on the upper limit length of the plurality of non-overlapping continuous intervals includes: An input / output buffer is allocated for the non-overlapping continuous interval, and the processor memory address of the input / output buffer is recorded; the dimension of the input / output buffer includes the upper limit length of the non-overlapping continuous interval. Copy the placeholder data to the input buffer; The device function is started and the input / output buffer is bound. After the pre-trained language model executes the computation process, the computation process is converted into a topological execution graph with a fixed tensor shape.

9. An electronic device, comprising: include: Memory, used to store computer programs; A processor, configured to implement the steps of the processing method for the inference request as described in any one of claims 1 to 8 when executing the computer program.

Images