Large model key-value pair fusion processing method, device and system for non-continuous video memory architecture

Abstract

This invention discloses a method, apparatus, and system for key-value pair fusion processing of large models with non-contiguous memory architectures. The method includes reading raw data, including raw key vectors and raw value vectors, from the linear layer of the model and storing them in a register; using a floating-point arithmetic logic unit (Floating-point arithmetics) to perform position encoding calculations on the raw key vectors in the registers to obtain new key vectors, and performing a silent residency operation on the raw value vectors; simultaneously using an integer arithmetics calculation unit (Integer arithmetics) to calculate the target storage address of the raw data in global memory; after the floating-point arithmetics calculation unit completes the position encoding calculation and the integer arithmetics calculation unit completes the target storage address calculation, the computation core sends a storage instruction to a loading storage unit, which reads the new key vector and the raw value vector from the registers and stores them in the target storage address in global memory. This invention can reduce cache latency and bandwidth overhead during large model inference, improve computational efficiency, and increase cache resource utilization.

Description

Technical Field

[0001] This invention relates to the fields of artificial intelligence and computer technology, and in particular to a method, apparatus and system for large model key-value pair fusion processing for non-contiguous memory architecture. Background Technology

[0002] With the widespread application of Large Language Models (LLMs) based on the Transformer architecture across various fields, the inference efficiency of long-sequence text generation has become a core challenge. During the inference phase, as the length of the input sequence increases, the GPU memory occupied by the key-value cache expands rapidly. To address the GPU memory fragmentation problem caused by long-sequence generation, mainstream inference engines, represented by vLLM, generally introduce a paging attention mechanism similar to operating system virtual memory management. This mechanism divides the key-value cache into fixed-size physical blocks and uses page tables for non-contiguous memory management, mapping logically contiguous tokens to physically discontiguous GPU memory spaces.

[0003] Although modern GPU hardware architectures generally support dual-issue and concurrent execution of instructions for integer and floating-point units, existing mainstream inference engines are limited by the serialization inertia of code logic and the limitations of general-purpose compilers in scheduling high-latency instructions. This makes it difficult to fully unleash the hardware's potential in specific complex scenarios involving position-encoded computation and non-contiguous memory addressing. Specifically, the integer division and modulo operations required for paging management have high instruction latency. Without targeted low-level operator orchestration, general-purpose software cannot effectively hide these operations within intensive floating-point computation cycles. This results in frequent pipeline idleness due to waiting for address resolution, thus limiting the upper limit of inference throughput in long-sequence, high-concurrency scenarios.

[0004] In the Transformer model decoding process based on the above architecture, after the linear layer of the model outputs the original query (Q), key (K), and value (V) vectors, existing standard data processing pipelines typically employ a serial, step-by-step processing method. See also... Figure 1 The specific process is as follows:

[0005] Step 1: Location Encoding (Executed by RoPE Kernel). The system first starts a dedicated computing kernel (such as...) Figure 1The computational core 1 shown reads the raw key and value data generated by the linear layer (the linear layer generates consecutive K and V) from the video memory. Based on the logical position of the current token in the sequence, it applies Rotary Positional Embedding (RoPE) to rotate the consecutive K read and obtain K'. Then, it writes the computation result K' back to a pre-allocated temporary consecutive buffer in the video memory.

[0006] Step Two: Cache Migration (Executed by Cache Manager / Copy Kernel). The compute unit then restarts another independent cache management core (such as...). Figure 1 The computational core 2 shown reads the rotated Key vector K' from the temporary buffer, reads the page table information in parallel, calculates the non-contiguous physical address V corresponding to the data, and finally moves the data to each discrete physical block of the physical memory, i.e., writes it to the non-contiguous memory (generating memory read and write).

[0007] Although the paging attention mechanism mentioned above effectively solves the problem of memory fragmentation, this serial processing flow that separates position encoding calculation from non-contiguous cache writing has significant performance bottlenecks in high-performance inference scenarios, mainly in the following three aspects.

[0008] (1) Severe waste of video memory bandwidth: In the existing process, the intermediate data generated in step one above needs to be written back to the pre-allocated temporary contiguous buffer in video memory, and then the data needs to be read out of video memory in step two. This write-back-read process constitutes a redundant global video memory round-trip operation. In the decoding stage of large model inference, the computational density is low, and memory access bandwidth is often the core bottleneck restricting performance. This redundant read and write directly squeezes out the valuable video memory bandwidth resources and reduces the inference throughput.

[0009] (2) Multi-core startup and scheduling overhead: Separating position encoding and cache management into two independent computing cores for serial execution increases the instruction issuance and scheduling overhead from CPU to GPU. Especially in real-time interactive scenarios that handle small batches or have extremely high requirements for instant response latency, the fixed latency caused by multiple computing core startups accounts for a significant proportion of the total time consumption, seriously affecting the first character generation time (TTFT) and token generation speed.

[0010] (3) Temporary storage resource usage: In order to connect the two independent steps mentioned above, the system usually needs to pre-allocate additional contiguous video memory space as an intermediate buffer to temporarily store the position-encoded data. This not only increases the peak video memory usage of the system, but may also lead to video memory overflow in high-concurrency scenarios where video memory resources are scarce, thus limiting the maximum number of concurrent users that the model can support.

[0011] The disclosure of the above background technical content is only for the purpose of assisting in understanding the concept and technical solution of this application, and does not necessarily provide technical instruction. Summary of the Invention

[0012] The purpose of this invention is to provide a method, apparatus, and system for large model key-value pair fusion processing for non-contiguous memory architectures, which can reduce computational latency, reduce bandwidth overhead, improve computational efficiency, and improve the utilization of storage resources.

[0013] To achieve the above objectives, the technical solution adopted by the present invention is as follows:

[0014] A method for key-value pair fusion processing of large models for non-contiguous memory architectures includes the following steps:

[0015] The output data of the linear layer of the model is read to obtain the raw data, which includes the raw key vector and the raw value vector, and the raw data is directly loaded into the private register of the computing core.

[0016] A floating-point arithmetic logic unit is used to perform position encoding calculation on the original key vector in the register to obtain a new key vector, and a silent residency operation is performed on the original value vector to maintain the original value vector in a physical residency state in the register;

[0017] During the location encoding calculation performed by the floating-point arithmetic logic unit, the target storage address of the original data in the global video memory is calculated simultaneously using the integer arithmetic logic unit.

[0018] After the floating-point arithmetic logic unit completes the position encoding calculation and the integer arithmetic logic unit completes the target storage address calculation, the computing core sends a storage instruction to the loading storage unit. The loading storage unit reads the new key vector and the original value vector from the register and stores them in the target storage address of the global video memory.

[0019] Furthermore, following any one or a combination of the aforementioned technical solutions, the target storage address of the original data in global video memory is calculated using an integer arithmetic logic unit, including the following steps:

[0020] The computing core synthesizes a global logical index based on the current thread index and the thread block index. The global logical index is then decomposed into a logical page number and an offset within the page to obtain the address calculation instruction sequence.

[0021] The address calculation instruction sequence is concurrently injected into the integer pipeline so that the integer arithmetic logic unit calculates the target memory address;

[0022] The integer arithmetic logic unit accesses the page table entry stored in the preset cache according to the logical page number, obtains the base address of the target key-value cache physical block, and calculates the non-contiguous physical pointer of the original data target to be written to the global video memory as the target storage address by combining the page offset.

[0023] Furthermore, following any one or a combination of the aforementioned technical solutions, calculating the target storage address using an integer arithmetic logic unit preferably involves replacing division and modulo operations with constant multiplication inverses or bitwise operations.

[0024] Furthermore, following any one or a combination of the aforementioned technical solutions, the preset cache includes a page table address translation cache configured in the following manner: a preset page table address translation cache that can be accessed by each computing core is provided, wherein the page table address translation cache stores logical page numbers and base addresses of physical blocks that have a mapping relationship.

[0025] And / or,

[0026] The preset cache includes one or more of L1 cache and shared memory.

[0027] Furthermore, following any one or a combination of the aforementioned technical solutions, a read operation on the output data of the linear layer of the model is initiated to read the original data in the following manner:

[0028] Depending on the model accuracy type, the loading memory unit inside the computing core is configured to execute loading instructions with the maximum bus width;

[0029] At the hardware scheduling level, the addresses accessed by all active threads within the same thread bundle are aligned continuously in the global video memory physical space. The hardware memory controller aggregates multiple independent read requests into a single or fewer video memory bus transactions.

[0030] Furthermore, following any one or a combination of the aforementioned technical solutions, the method further includes implementing the strategy of keeping the original data residing in the register throughout its entire lifecycle in the following manner:

[0031] After the original data is directly loaded into the private register file of the computing core, the original data in the register file is directly used as the input for subsequent position encoding calculation and target storage address calculation, blocking intermediate state video memory write-back, including not releasing the original data to the on-chip shared memory, and not writing the original data to any temporary buffer in the global video memory.

[0032] Furthermore, following any one or a combination of the aforementioned technical solutions, a new key vector is obtained by performing position encoding calculations on the original key vector within a register using a floating-point arithmetic logic unit, including the following steps:

[0033] The new key vector is used to overwrite the original key vector so that the new key vector is stored at the same memory address in the register where the original key vector was stored.

[0034] Furthermore, following any one or a combination of the aforementioned technical solutions, a new key vector is obtained by performing position encoding calculations on the original key vector within a register using a floating-point arithmetic logic unit, including the following steps:

[0035] The new key vector is reassigned a storage address in the register, and the storage address of the new key vector is different from the storage address of the original key vector.

[0036] Furthermore, following any one or a combination of the aforementioned technical solutions, the method further includes the following steps: based on a periodic iterative covering strategy, covering the original key vector with the new key vector after a certain duration or a certain number of iterations.

[0037] Furthermore, following any or a combination of the aforementioned technical solutions, reading the new key vector and the original value vector from the register and storing them in the target storage address of the global video memory includes performing a parallel write strategy on multiple pairs of new key vectors and original value vectors in the following manner:

[0038] The available bandwidth is predetermined, and multiple new key vectors and original value vector pairs are read in parallel based on the available bandwidth, wherein the bandwidth occupied by the multiple new key vectors and original value vector pairs is not higher than the available bandwidth;

[0039] The read new key vectors and original value vectors are stored in the global video memory according to their corresponding target storage addresses.

[0040] Furthermore, following any one or a combination of the aforementioned technical solutions, the location coding calculation process also includes the following steps:

[0041] The format type of the original key vector is detected in real time. If the original key vector is in FP8 format, multiple original key vectors are packaged into a first key vector with a preset bit width. Then, a floating-point arithmetic logic unit is used to perform position encoding calculation on the first key vector with rotation parameters of FP16 or BF16 precision.

[0042] Furthermore, following any one or a combination of the aforementioned technical solutions, if the original key vector is in FP16 or BF16 format, then multiple original key vectors are packaged into a second key vector with a preset bit width, and then a floating-point arithmetic logic unit is used to perform position encoding calculation on the first key vector.

[0043] According to another aspect of the present invention, the present invention provides a key-value pair fusion processing apparatus for large models with non-contiguous memory architecture, which performs key-value pair fusion processing on large model key-value pairs in the inference process of large models with non-contiguous memory architecture based on any one or a combination of the above technical solutions.

[0044] According to another aspect of the present invention, the present invention provides a computer system suitable for key-value pair fusion processing of large models, wherein the key-value pair fusion processing method for large model inference process oriented to non-contiguous memory architecture, as described in any one or a combination of the above technical solutions, performs key-value pair fusion processing on large model components.

[0045] The beneficial effects of the technical solution provided by this invention are as follows:

[0046] a. The key-value pair fusion processing method for large models with non-contiguous video memory architecture proposed in this invention deeply integrates computation and memory access, breaks the boundary between position encoding computation and cache management in traditional large model inference engines, and merges these two originally independent steps into an atomic operation completed within the life cycle of a single computing core. It uses GPU registers as a high-speed transfer channel and embeds position encoding mathematical operations during the process of storing data into non-contiguous video memory, thereby eliminating redundant intermediate video memory read and write operations.

[0047] b. This invention achieves extreme optimization of bandwidth utilization through raw data extraction and register in-situ residence strategy. By merging memory accesses, it aggregates the originally discrete memory read requests into the minimum number of global video memory transactions, minimizing bus idle time and locking data in the register file. This eliminates the redundant bandwidth consumption caused by the repeated movement of data between global video memory and computing core in traditional solutions, and provides a basis for subsequent pipeline delay masking.

[0048] c. This invention enables full-load parallel operation of integer and floating-point resources within the computing core, effectively reducing computational latency and achieving dual-pipeline load balancing. Through a parallel mechanism, the memory address calculation time, which is usually a pure overhead, is effectively hidden / overlapped within the clock cycle of floating-point operations. This avoids the idle floating-point arithmetic logic unit caused by waiting for address calculation to complete in the traditional serial mode, or the idle integer arithmetic logic unit caused by performing floating-point calculations. Thus, without increasing the total instruction cycle, it completes the complex mapping from logical space to non-contiguous physical space and solves the problem of additional latency caused by non-contiguous video memory management.

[0049] d. After the position encoding and address calculation are completed, the present invention performs zero-copy writing, writing the new key vector and the original value vector in the register into the global video memory based on the target storage address. This not only halves the bandwidth consumption, but also further improves the system operating efficiency through parallel writing. Attached Figure Description

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

[0051] Figure 1 This is a flowchart illustrating a method for processing key-value pairs in large models in the prior art;

[0052] Figure 2 A schematic diagram of the large model key-value pair fusion processing provided as an exemplary embodiment of the present invention;

[0053] Figure 3 A schematic diagram illustrating the entire process principle of instruction parallel processing and fusion based on GPU streaming multiprocessors, provided as an exemplary embodiment of the present invention;

[0054] Figure 4 A flowchart of a large model key-value pair fusion processing method for non-contiguous memory architecture provided as an exemplary embodiment of the present invention. Detailed Implementation

[0055] To enable those skilled in the art to better understand the present invention, the technical solutions of the present invention will be clearly and completely described below with reference to the accompanying drawings of the embodiments of the present invention. Obviously, the described embodiments are only some embodiments of the present invention, and not all embodiments. Based on the embodiments of the present invention, all other embodiments obtained by those skilled in the art without creative effort should fall within the scope of protection of the present invention.

[0056] It should be noted that the terms "first," "second," etc., in the specification, claims, and accompanying drawings of this invention are used to distinguish similar objects and are not necessarily used to describe a specific order or sequence. It should be understood that such data can be interchanged where appropriate so that the embodiments of the invention described herein can be implemented in orders other than those illustrated or described herein. Furthermore, the terms "comprising" and "having," and any variations thereof, are intended to cover a non-exclusive inclusion; for example, a process, method, apparatus, product, or device that comprises a series of steps or units is not necessarily limited to those steps or units explicitly listed, but may include other steps or units not explicitly listed or inherent to such processes, methods, products, or devices.

[0057] To address the aforementioned technical challenges in the cutting-edge fields of artificial intelligence and high-performance computing, particularly concerning the performance of large language model inference on GPUs or dedicated AI accelerator architectures, this invention focuses on the decoding stage, which experiences the greatest computational and memory access pressure during large model inference. Specifically, it proposes a method that deeply integrates rotation position encoding computation with non-contiguous key-value cache write operations in complex scenarios involving the generation of extremely long input sequences and the use of non-contiguous memory management mechanisms (such as paging attention) to cope with memory fragmentation. This invention aims to significantly reduce memory access latency and global memory read / write operations in high-concurrency, long-sequence scenarios through operator-level computation-memory fusion, thereby overcoming the memory bandwidth bottleneck and greatly improving the real-time inference speed and system throughput of large models.

[0058] In one embodiment of the present invention, a method for large model key-value pair fusion processing oriented towards non-contiguous memory architecture is provided, see [link to relevant documentation]. Figures 2 to 4 This includes the following steps:

[0059] The output data of the linear layer of the model is read to obtain the raw data, which includes the raw key vector and the raw value vector, and the raw data is directly loaded into the private register of the computing core.

[0060] A floating-point arithmetic logic unit is used to perform position encoding calculation on the original key vector in the register to obtain a new key vector, and a silent residency operation is performed on the original value vector to maintain the original value vector in a physical residency state in the register;

[0061] During the location encoding calculation performed by the floating-point arithmetic logic unit, the target storage address of the original data in the global video memory is calculated simultaneously using the integer arithmetic logic unit.

[0062] After the floating-point arithmetic logic unit completes the position encoding calculation and the integer arithmetic logic unit completes the target storage address calculation, the computing core sends a storage instruction to the loading storage unit. The loading storage unit reads the new key vector and the original value vector from the register and stores them in the target storage address of the global video memory.

[0063] Based on the key-value pair fusion processing method for large models with non-contiguous memory architecture proposed in this invention, this invention provides a computation-memory access deep fusion operator architecture. This architecture breaks the boundary between positional encoding computation and cache management in traditional large model inference engines, merging these two originally independent steps into an atomic operation completed within the lifecycle of a single computation core. Its core lies in utilizing GPU registers as a high-speed transfer channel, embedding positional encoding mathematical operations during the data storage process into non-contiguous memory, thereby eliminating redundant intermediate memory reads and writes. The working process of the above-mentioned computation-memory access deep fusion operator architecture is further illustrated through the following steps.

[0064] Step S1: Extract raw data and retain it in the register.

[0065] First, read operations are performed based on a high-performance loading mechanism. The Load Memory Unit (LSU) within the computation core initiates read operations for the output data of the model's linear layers. To alleviate the high latency bottleneck of global memory, the following standard optimization strategies are adopted.

[0066] Vectorized loading: Based on the model accuracy type, the loading storage unit is configured to execute loading instructions with the maximum bus width, so that a single thread can extract multiple consecutive data elements at once within one instruction cycle, significantly reducing instruction issue density.

[0067] Merged memory accesses: At the hardware scheduling level, this ensures that the addresses accessed by all active threads within the same thread bundle are contiguously aligned in the global video memory physical space. This allows the hardware memory controller to aggregate these independent read requests into a single or very few video memory bus transactions, thereby maximizing bandwidth utilization.

[0068] Secondly, a register residency strategy is implemented throughout its entire lifecycle. The key inventive point of this application lies in the lifecycle management of data after it is read into the computing core. Unlike existing technologies, this invention constructs the following "register in-situ anchoring" mechanism.

[0069] Forced residency: The raw data read is directly loaded into the compute core's private register file. It is strictly forbidden to release or evict this read data to on-chip shared memory, or write it back to any temporary buffer in global video memory, until subsequent steps are completed.

[0070] Blocking intermediate state memory write-back: This design physically blocks the inefficient path of traditional operators having to exchange data through global memory (i.e., "reading the original data, writing back to global memory, and then rereading global memory").

[0071] By employing the original data extraction and register in-situ resident strategy proposed in step S1, this application achieves extreme optimization of bandwidth utilization: by merging memory accesses, the originally discrete memory read requests are aggregated into the minimum number of global video memory transactions, minimizing bus idle time; this application constructs a zero-copy fusion base: by locking data in the register file, redundant bandwidth consumption caused by repeated data transfer between global video memory and computing cores in traditional schemes is eliminated, providing a foundation for subsequent pipeline delay masking.

[0072] The innovation of step S1 in this application lies in the management of the destination and lifecycle of the original data. In the prior art, data loading into the register is temporary, and it must then be written back to the temporary area of ​​the video memory via the bus, resulting in wasted bandwidth. In this application, the data is locked in the register, and the parallel issue mechanism of the hardware is used to allow the data to become the input of the subsequent calculation logic (FP16) and addressing logic (INT32) without leaving the register.

[0073] Step S2: Perform in-situ rotation position encoding calculation within the register.

[0074] First, the immediate scheduling of arithmetic logic units is performed in the following manner.

[0075] In step S1, after the original key and value data are successfully locked into the register file, the computing core immediately activates its internal floating-point arithmetic logic unit. At this time, the computing core implements a "separate data path control" strategy (the core logic of this patent), executing the following two completely different instruction sequences on the two types of resident vectors.

[0076] (1) Perform in-situ rotation on the key vector.

[0077] Parameter acquisition: The calculation core loads pre-calculated geometric rotation parameters (Cos / Sin table) from constant cache or global video memory based on the channel index and Token logical position of the current thread.

[0078] Register-level updates: For the resident key-value vector components, the floating-point arithmetic logic unit (FP ALU) performs complex multiply-accumulate fusion (FMA) operations. The key lies in controlling the data flow: the new key-value vector generated by the operation (i.e., the result with positional encoding) is directly overwritten into the target register in the register file, or written to a newly allocated register address. The entire computation process is strictly constrained within a closed-loop system of "register read to FP computation to register write," and the system strictly prohibits writing this intermediate result back to global memory.

[0079] Preferably, the following steps are employed in the position encoding calculation process to further improve system computational efficiency: The format type of the original key vector is detected in real time. If the original key vector is in FP8 format, multiple original key vectors are packaged into a first key vector with a preset bit width. Then, a floating-point arithmetic logic unit (Floating-point Logic Unit) is used to perform position encoding calculation on the first key vector using rotation parameters with FP16 or BF16 precision. Bit width refers to the physical storage bit width occupied in the physical registers, cache, or data bus of the computational core after data packaging. For example, it can be a fixed bit width of the underlying hardware vector registers such as 128 bit, 256 bit, or 512 bit. More preferably, if the original key vector is in FP16 or BF16 format, multiple original key vectors are packaged into a second key vector with a preset bit width. Then, a floating-point Logic Unit (Floating-point Logic Unit) is used to perform position encoding calculation on the first key vector. This reduces the number of position encoding calculations while ensuring continuous address alignment.

[0080] (2) Perform a silent dwell operation on the value vector.

[0081] For value vectors resident in the same batch, the computation core performs instruction masking, preventing any floating-point arithmetic or logical operations from being performed on the value vectors within the current instruction cycle. The physical storage properties of the register file are utilized to maintain the value vector data in its original location.

[0082] The value vector data is configured as a "silent operand," meaning that while the floating-point arithmetic logic unit is processing the key vector rotation at full speed, the value vector remains physically resident in the register file and is not released. The system utilizes the in-situ retention characteristic of the register file to artificially extend the lifetime of the value vector within the chip, ensuring that it remains operand-ready during the high latency period of the key vector rotation calculation, until the synchronous write condition of the subsequent S4 stage is met.

[0083] This invention constructs a zero-overflow closed loop within the chip through step S2, and through differentiated path control, builds an ultra-high-speed data flow closed loop strictly limited to the "register file to ALU". Compared with the prior art, this step completely eliminates the redundant operation in the traditional solution that requires writing the intermediate calculation result (i.e., the rotated key) back to the global video memory temporary buffer, and then having the kernel reread it for subsequent transport.

[0084] Furthermore, this step enables the transformation and utilization of arithmetic intensity: by intensively performing multiply-accumulate fusion (FMA) operations within registers, the operator is temporarily transformed into a "computation-intensive" state. This high load on the FP ALU creates the necessary "time window" for subsequent step S3 to utilize a dual-issue mechanism to perform address calculations in parallel, thereby masking integer pipeline delays.

[0085] It should be noted that the inventiveness of this application lies not in "what to calculate," but in "where to calculate" and "where the data goes after the calculation." Existing technologies require consuming video memory bandwidth to write the results after computation; however, this step is part of the fusion operator, and its calculation results are forcibly latched into the register file, directly serving as input for subsequent write steps. This application utilizes register-level data transfer instead of traditional video memory-level data transfer, achieving the core means of low-latency inference.

[0086] Step S3: Perform physical address calculation based on instruction-level parallelism.

[0087] (1) Efficient destructuring of the execution logical index. The computational core first synthesizes a global logical index based on the current thread index and the thread block index. The global logical index is decomposed into a "logical page number" and a "page offset" to obtain an address calculation instruction sequence, which is used to map the global logical index to a non-contiguous video memory space.

[0088] Then, concurrent issue operations of the heterogeneous pipeline are executed. The key to this step lies in the timing control of instruction scheduling. The instruction scheduler inside the compute core leverages the hardware feature of modern GPU architectures, where floating-point pipelines and integer pipelines are independent, to execute the following parallel strategies.

[0089] (2) Concurrent detection and emission. During the full-speed execution of the in-situ rotation operation (complex number rotation multiplication and addition instruction) described in step S2 by the floating-point arithmetic logic unit (FP ALU), the instruction scheduler detects that the integer arithmetic logic unit (INT ALU) is in an idle state.

[0090] Instruction Injection: The system then concurrently injects the above-mentioned optimized address calculation instruction sequence (index destructuring, base address lookup) into the integer pipeline for execution.

[0091] Parallel execution: At this time, the chip is in a "dual-track operation" state, that is, the floating-point arithmetic logic unit (FP ALU) is processing the numerical transformation of the data (i.e., the in-situ rotation encoding operation in step S2), while the integer arithmetic logic unit (INT ALU) is simultaneously calculating the target storage location of the data, i.e., the target storage address (step S3).

[0092] Existing technologies typically treat Location Encoding (RoPE) and cache writing as two separate steps, either executed on two separate computing cores or serially on the same computing core (RoPE is calculated first, then the address is calculated and written to memory). This can lead to "bubbles" in the pipeline, where the floating-point arithmetic logic unit is busy while the integer arithmetic logic unit is idle, and vice versa.

[0093] In this application, position encoding calculation is performed purely by floating-point arithmetic logic units, while address calculation is performed by pure integer arithmetic logic units. By overlapping the two processes in time, instruction-level parallelism is manually orchestrated for specific operator characteristics.

[0094] (3) Instant resolution of non-contiguous physical addresses.

[0095] The integer arithmetic logic unit (ALU) accesses the page table entry preloaded in the L1 cache or shared memory based on the calculated logical page number, obtains the base address of the target key-value cache (KV Cache) physical block, and calculates the non-contiguous physical pointer to be written to global video memory, i.e., the target memory address, in conjunction with the page offset. This physical pointer is temporarily stored in a register, awaiting use in the subsequent step S4.

[0096] More preferably, the preset cache further includes a page table address translation cache configured as follows: a preset page table address translation cache accessible to each computing core is provided, wherein the page table address translation cache stores logical page numbers and base addresses of physical blocks with a mapping relationship. The integer arithmetic logic unit first accesses the page table entry in the page table address translation cache based on the calculated logical page number. If a hit occurs, the base address of the target key-value cache physical block is obtained; if a hit does not occur, an access to the L1 cache or shared memory is triggered to obtain the base address of the target key-value cache physical block. This further reduces the frequency of page table lookup cycles, making the INT ALU's computing pipeline almost uninterrupted.

[0097] Considering the high latency of standard integer division instructions in hardware, it is preferable to use constant multiplication inverses or bitwise operations to replace expensive division and modulo instructions. This ensures that the execution cycle of integer arithmetic instruction sequences is strictly controlled, making it easier to fit within the time window of floating-point operations.

[0098] This application enables the full-load parallel operation of integer and floating-point resources within the computational core, effectively reducing computational latency and achieving dual-pipeline load balancing. This application utilizes the execution cycle of the position-encoded intensive floating-point operation in step S2 as a "shield." Through a dual-issue mechanism, the time spent calculating memory addresses, which is typically pure overhead, is effectively hidden / overlapped within the clock cycle of the floating-point operation. From the perspective of total pipeline time, this achieves "zero additional time overhead" for complex non-contiguous address calculation logic. Furthermore, this design avoids the idle floating-point units caused by waiting for address calculations to complete in traditional serial mode, or the idle integer units caused by executing floating-point calculations. Thus, without increasing the total instruction cycle, it completes the complex mapping from logical space to non-contiguous physical space. In simpler terms, the complex addresses that would normally require extra time to calculate are now completed incidentally during "others' busy periods and your own free time," so from the perspective of total time, it seems as if no time was spent at all.

[0099] It should be noted that although GPU hardware supports dual-issue, in existing technologies, position encoding calculation and non-contiguous address resolution often belong to different kernels or serial code segments. The hardware itself cannot automatically couple these two specific sets of logic together. This application discovers and utilizes the high-density floating-point operation cycles of the position encoding operator as a time window, creatively embedding complex non-contiguous memory addressing logic within this window. This is not an automatic hardware behavior, but a proactive instruction stream reconfiguration tailored to specific algorithm characteristics, thereby solving the additional latency problem caused by non-contiguous memory management.

[0100] Step S4: Perform zero-copy write.

[0101] Pipeline convergence and instruction submission. Once the instruction scheduler confirms that the floating-point operation (the RoPE-transformed Key vector) in step S2 and the integer operation (non-contiguous physical address pointers) in step S3 have been completed, the system enters the pipeline convergence state and performs the following operations.

[0102] Data source: The data to be written is directly taken from the position encoding calculation result (new key value vector) residing in the register file and the original value vector that has been silently residing there.

[0103] Address source: The target write address is directly taken from the non-contiguous physical address pointer temporarily stored in the register file, which is the target storage address (the calculation result of step S3).

[0104] Instruction issuance: The compute core sends a store instruction to the Load Memory Unit (LSU). The Load Memory Unit then takes over bus control and performs a direct write from registers to global memory.

[0105] No intermediate steps: The entire writing process strictly follows a single path from register to global video memory, skipping all intermediate steps that must be taken in traditional technology, such as "temporary buffer allocation", "intermediate data writing to disk" and "re-reading and moving".

[0106] Preferably, the system operating efficiency is further improved by performing a parallel write strategy on multiple new key vectors and original value vector pairs in the following manner: the current available bandwidth is determined in advance, multiple new key vectors and original value vector pairs are read in parallel based on the available bandwidth, and the bandwidth occupied by multiple new key vectors and original value vector pairs is not higher than the available bandwidth; the read multiple new key vectors and original value vector pairs are stored in the global video memory according to their corresponding target storage addresses.

[0107] The key-value pair fusion processing method for large models with non-contiguous memory architecture provided in this application can basically halve bandwidth consumption. From a system perspective, this process reduces the number of global memory accesses in the key-value cache preparation stage from 4N times (reading original data + writing intermediate data + reading intermediate data + writing final data) in traditional technology to 2N times (reading original data + writing final data), where N is a positive integer.

[0108] This application merges the "location encoding calculation" and "cache movement" that are traditionally separated into two independent computing cores into a single computing core's lifecycle computational operation at the physical level, eliminating the additional overhead of kernel switching and memory scheduling.

[0109] In contrast to existing technologies where Location-Based Execution (RoPE) and Paging are typically decoupled into different software modules, requiring data to be transferred through video memory, this invention achieves zero-latency synchronization between computation results and physical addresses at the instruction issuance level through steps S1-S4, thereby physically eliminating the need for an intermediate buffer. This architectural simplification cannot be automatically achieved by existing general-purpose hardware instructions and must rely on the specific operator fusion design described in this patent.

[0110] The following specific embodiment further illustrates the large model key-value pair fusion processing method for non-contiguous memory architectures provided by the present invention.

[0111] Model architecture: A large language model based on the Transformer architecture (such as LLaMA-3 or the DeepSeek series), using the standard Grouped Query Attention (GQA) mechanism.

[0112] Current task: In the decoding phase of reasoning, generating the 1025th token in the sequence.

[0113] Hardware environment: GPU computing cores that support dual-issue technology (such as the NVIDIA Ampere / Hopper architecture), with independent integer and floating-point pipelines, supporting parallel execution of INT32 and FP16 instructions.

[0114] Video memory status: Paging memory management is used, the physical space of the key-value cache is not contiguous, and the page block size is set to 16.

[0115] Compute configuration: Execution is performed in units of warp, with each warp containing 32 threads.

[0116] Data parameters: The attention head dimension is 128, and the data precision is FP16 (2 bytes per element). That is, each token (such as the word "The") will generate a 128-dimensional key vector (containing 128 FP16 components) and a 128-dimensional value vector (containing 128 FP16 components).

[0117] Step S1: Extract raw data and retain it in the register.

[0118] Action: The computation core receives instructions from the output of the linear layer of the model.

[0119] Vectorized Loading: The LSU within the compute core, based on a merged memory access mechanism, issues a 128-bit wide load instruction (such as ld.global.v4.b32 or equivalent instructions). Through this instruction, a single thread continuously reads a 128-bit data packet—the raw data—from the linear layer output region of global memory in a single operation. This data packet logically contains four FP16 format key vector components and four FP16 format value vector components.

[0120] Register anchoring: After these 8 raw data (i.e., 4 key components + 4 value components) are read, they are directly stored in the private register file of the computing core (e.g., registers R0-R7).

[0121] Key constraint: At this point, the original data is locked in R0-R7, and the system does not write this data to the on-chip shared memory, nor does it write it back to any temporary area of ​​the global video memory.

[0122] Step S2: Perform in-situ rotation position encoding calculation within the register.

[0123] Action: Raw data is ready in the registers, and the computation core starts the floating-point arithmetic logic unit (FP ALU). The separate path performs the following operations.

[0124] Key vector components (R0-R3): The FP ALU performs RoPE rotation calculations on the key vector components in registers R0-R3. The calculation core pairs the data in R0-R3 (e.g., R0 with R1, R2 with R3), reads the pre-stored Cos / Sin parameters, and performs complex multiplication-addition fusion (FMA) operations. The calculation result either directly overwrites R0-R3 in place or writes it to the corresponding address in a newly allocated register. At this point, the register stores a new key vector with positional information.

[0125] Value vector components (R4-R7): The computation core performs arithmetic isolation on the value vector components in R4-R7. During the dozens of clock cycles that the FPALU is busy calculating the new key vector, the value vector remains silently resident in the register, does not participate in the operation, and only extends its lifespan within the chip to wait for the convergence of step S4.

[0126] Step S3: Perform physical address calculation based on instruction-level parallelism.

[0127] Action: While the FP ALU is calculating the key vector rotation from start to finish (step S2), the instruction scheduler detects that the integer unit (INT ALU) is idle, and then performs the following parallel computation.

[0128] Logical breakdown: The INT ALU receives the logical index (Index 1024) of the current token in parallel. It calculates that the token belongs to logical page number 64 (assuming a page size of 16), with an offset of 0 within the page.

[0129] Physical mapping: The INT ALU looks up the table based on the logical page number and finds that the physical base address corresponding to logical page number 64 is located in physical block 208 of global video memory.

[0130] Output: The INT ALU calculates the final target physical pointer address (Ptr Final) as the target storage address and stores it in the integer register.

[0131] Since the Location Encoding (RoPE) calculation in step S2 is a dense floating-point operation and takes a long time, the time overhead of the address calculation process in step S3 above is completely overlapped / masked within the floating-point calculation window, thus achieving zero additional overhead for address resolution.

[0132] Step S4: Perform zero-copy write.

[0133] Action: The position encoding calculation is completed in step S2, and the target storage address calculation is completed in step S3.

[0134] Data source: Directly extract registers R0-R3 (which store the new Key vector components after RoPE transformation) and registers R4-R7 (which store the original value vector components that remain silently throughout the entire lifecycle).

[0135] Address source: The target storage address (pointing to a specific offset position of the target physical block 208) is directly calculated using step S3 and temporarily stored in the integer register.

[0136] The pipeline convergence is performed according to the following process: (1) The compute core issues a storage instruction to the LSU; (2) The LSU performs a 128-bit wide vectorized storage operation to pass-through the key-value pairs into the global video memory. Preferably, the MemoryCoalescing mechanism is used to directly write the 128-bit data packet (containing 4 key vectors + 4 value vectors) into the corresponding non-contiguous physical block in the global video memory.

[0137] Final state: The data comes out of the linear layer, goes through a register (the key changes, but the value remains the same), and then falls directly into the non-contiguous key-value cache. There are no read or write operations on the temporary video memory buffer in between.

[0138] In one embodiment of the present invention, a large model key-value pair fusion processing apparatus for non-contiguous memory architecture is provided, which performs fusion processing on large model key-value pairs based on the key-value pair fusion processing method in the large model inference process for non-contiguous memory architecture as described in the above embodiment.

[0139] In one embodiment of the present invention, a computer system suitable for key-value pair fusion processing of large models is provided, which performs key-value pair fusion processing on large model inference process based on the key-value pair fusion processing method for non-contiguous memory architecture as described in the above embodiment.

[0140] It should be noted that the above-described embodiments of the large model key-value pair fusion processing apparatus for non-contiguous memory architecture and the computer system embodiments applicable to large model key-value pair fusion processing are based on the same inventive concept as the embodiments of the large model key-value pair fusion processing method for non-contiguous memory architecture. All contents of the embodiments of the large model key-value pair fusion processing method for non-contiguous memory architecture are incorporated into the embodiments of the large model key-value pair fusion processing apparatus for non-contiguous memory architecture and the computer system embodiments applicable to large model key-value pair fusion processing by reference.

[0141] It should be noted that, in this document, relational terms such as "first" and "second" are used only to distinguish one entity or operation from another, and do not necessarily require or imply any such actual relationship or order between these entities or operations. Furthermore, 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. Without further limitations, an element defined by the phrase "comprising one..." does not exclude the presence of other identical elements in the process, method, article, or apparatus that includes said element.

[0142] The above description is only a specific embodiment of this application. It should be noted that for those skilled in the art, several improvements and modifications can be made without departing from the principle of this application, and these improvements and modifications should also be considered within the scope of protection of this application.

Claims

1. A large model key-value pair fusion processing method for a non-continuous video memory architecture, characterized in that, Includes the following steps: The output data of the linear layer of the model is read to obtain the raw data, which includes the raw key vector and the raw value vector, and the raw data is directly loaded into the private register of the computing core. A floating-point arithmetic logic unit is used to perform position encoding calculation on the original key vector in the register to obtain a new key vector, and a silent residency operation is performed on the original value vector to maintain the original value vector in a physical residency state in the register; During the location encoding calculation performed by the floating-point arithmetic logic unit, the target storage address of the original data in the global video memory is calculated simultaneously using the integer arithmetic logic unit. After the floating-point arithmetic logic unit completes the position encoding calculation and the integer arithmetic logic unit completes the target storage address calculation, the computing core sends a storage instruction to the loading storage unit. The loading storage unit reads the new key vector and the original value vector from the register and stores them in the target storage address of the global video memory.

2. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 1, characterized in that, The process of calculating the target storage address of the original data in global video memory using an integer arithmetic logic unit includes the following steps: The computing core synthesizes a global logical index based on the current thread index and the thread block index. The global logical index is then decomposed into a logical page number and an offset within the page to obtain the address calculation instruction sequence. The address calculation instruction sequence is concurrently injected into the integer pipeline so that the integer arithmetic logic unit calculates the target memory address; The integer arithmetic logic unit accesses the page table entry stored in the preset cache according to the logical page number, obtains the base address of the target key-value cache physical block, and calculates the non-contiguous physical pointer of the original data target to be written to the global video memory as the target storage address by combining the page offset.

3. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 2, characterized in that, The calculation of the target storage address using an integer arithmetic logic unit includes replacing division and modulo operations with constant multiplication inverses or bitwise operations.

4. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 2, characterized in that, The preset cache includes a page table address translation cache configured in the following way: a preset page table address translation cache that can be accessed by each computing core is provided, wherein the page table address translation cache stores logical page numbers and base addresses of physical blocks that have a mapping relationship; And / or, The preset cache includes one or more of L1 cache and shared memory.

5. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 1, characterized in that, The original data is read by initiating a read operation on the output data of the linear layer of the model in the following manner: Depending on the model accuracy type, the loading memory unit inside the computation core is configured to execute loading instructions with the maximum bus width; At the hardware scheduling level, the addresses accessed by all active threads within the same thread bundle are aligned continuously in the global video memory physical space. The hardware memory controller aggregates multiple independent read requests into a single or fewer video memory bus transactions.

6. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 1, characterized in that, This also includes implementing the policy of keeping the raw data resident in the register for its entire lifecycle in the following ways: After the original data is directly loaded into the private register file of the computing core, the original data in the register file is directly used as the input for subsequent position encoding calculation and target storage address calculation, blocking intermediate state video memory write-back, including not releasing the original data to the on-chip shared memory, and not writing the original data to any temporary buffer in the global video memory.

7. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 1, characterized in that, The new key vector is obtained by performing position encoding calculations on the original key vector in a register using a floating-point arithmetic logic unit, including the following steps: The new key vector is used to overwrite the original key vector so that the new key vector is stored at the same memory address in the register where the original key vector was stored.

8. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 1, characterized in that, The new key vector is obtained by performing position encoding calculations on the original key vector in a register using a floating-point arithmetic logic unit, including the following steps: The new key vector is reassigned a storage address in the register, and the storage address of the new key vector is different from the storage address of the original key vector.

9. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 8, characterized in that, It also includes the following steps: Based on a periodic iterative covering strategy, the original key vector is covered by the new key vector after a certain duration or a certain number of iterations.

10. The method of claim 1, wherein the method further comprises: Reading the new key vector and the original value vector from the register and storing them in the target memory address of the global video memory includes performing a parallel write strategy on multiple pairs of new key vectors and original value vectors in the following manner: The available bandwidth is predetermined, and multiple new key vectors and original value vector pairs are read in parallel based on the available bandwidth, wherein the bandwidth occupied by the multiple new key vectors and original value vector pairs is not higher than the available bandwidth; The read new key vectors and original value vectors are stored in the global video memory according to their corresponding target storage addresses.

11. The method of claim 1, wherein the method further comprises: The location encoding calculation process also includes the following steps: The format type of the original key vector is detected in real time. If the original key vector is in FP8 format, multiple original key vectors are packaged into a first key vector with a preset bit width. Then, a floating-point arithmetic logic unit is used to perform position encoding calculation on the first key vector with rotation parameters of FP16 or BF16 precision.

12. The large model key-value pair fusion processing method for a non-continuous video memory architecture according to claim 11, characterized in that, If the original key vector is in FP16 or BF16 format, then multiple original key vectors are packaged into a second key vector with a preset bit width, and then a floating-point arithmetic logic unit is used to perform position encoding calculation on the second key vector.

13. A large model key-value pair fusion processing apparatus for a non-continuous video memory architecture, characterized in that, The key-value pair fusion processing method for large model inference process based on the non-contiguous memory architecture described in claim 1 is used to perform fusion processing on large model key-value pairs.

14. A computer system suitable for large model key-value pair fusion processing, comprising: The key-value pair fusion processing method for large model inference process based on the non-contiguous memory architecture described in claim 1 is used to perform fusion processing on large model key-value pairs.

Images