Large language model weight compression method, inference operation method, system and medium

Abstract

The application discloses a large language model weight compression method, an inference operation method, a system and a medium, and belongs to the technical field of large language model weight compression and inference operation optimization. The method comprises the following steps: acquiring a weight matrix of a large language model and constructing an exponential high-frequency window and a benchmark value, classifying and encoding each weight data into a fixed-length code word and splitting the fixed-length code word into an independent bitmap, synchronously generating a compact value stream and a rollback value stream, and packing the compact value stream and the rollback value stream into weight compression data; loading the weight compression data, generating a state mask and a channel prefix mask in a register based on the bitmap, calculating a reading offset by using a population count instruction, reading a compressed data segment from a corresponding data stream, reconstructing weight data in the register in combination with the code word, and finally directly inputting the weight data into a matrix multiplication unit for operation; therefore, by implementing the application, the problem of efficiency reduction caused by control flow divergence and redundant memory access in GPU inference of existing variable-length encoding can be solved, and the model inference efficiency under lossless compression can be improved.

Description

Technical Field

[0001] This application relates to the field of large language model compression and inference operation optimization technology, and in particular to a large language model weight compression method, inference operation method, system and medium. Background Technology

[0002] With the rapid development of artificial intelligence technology, large language models (LMs) have become a core technology in fields such as natural language processing. However, the improvement of model performance depends on the exponential growth of parameter size, which puts storage and bandwidth pressure on inference deployment on graphics processing units (GPUs).

[0003] Most current mainstream models use the BFloat16 format to store weights, which has significant redundancy in its exponent portion, and its information entropy is far lower than its physical storage width, theoretically possessing compression potential. Traditional general-purpose compression algorithms (such as Gzip), due to their inherent serial decoding characteristics, cannot adapt to the single-instruction multithreaded (SIMT) parallel computing architecture of GPUs, making it difficult to meet the high-throughput, low-latency real-time inference requirements. To adapt to GPU parallel computing, existing technical solutions (such as DietGPU, nvCOMP, etc.) mostly adopt variable-length compression methods based on entropy coding. However, variable-length coding leads to inconsistent decoding step sizes for each data codeword, which fundamentally conflicts with the GPU's SIMT architecture. In the GPU's single-instruction multithreaded execution model, threads within the same thread bundle need to maintain instruction synchronization. Variable-length decoding disrupts this synchronization, causing inconsistent decoding step sizes among threads within a thread bundle, resulting in frequent pipeline pauses and significantly reducing the utilization of hardware computing units. To circumvent the aforementioned decoding synchronization problem, existing solutions generally adopt a two-stage decoupling process of "global decompression first, then computation." The process involves first reading the compressed data from the video memory, then decompressing it on the chip, writing the decompressed data back to the video memory, and finally rereading it for matrix operations. This inserts redundant video memory round-trip operations between decoding and computation, which not only fails to effectively alleviate the video memory bandwidth pressure but also generates additional data transfer during the decoding stage, exacerbating the "memory wall" bottleneck and ultimately leading to a decrease in end-to-end inference throughput.

[0004] Therefore, existing technologies struggle to achieve efficient lossless compression inference on GPUs due to control flow divergence in variable-length encoding and redundant data transport issues in decoupled architectures. There is an urgent need for a compression and inference solution that integrates decoding and computation, eliminates intermediate data write-back, and deeply collaborates with the parallel capabilities of GPUs. Summary of the Invention

[0005] This application provides a method for compressing the weights of large language models, an inference computation method, a system, and a medium, which can solve the problem that the efficiency of model inference computation is reduced in GPU inference computation due to control flow divergence and redundant memory access in existing variable-length encoding compression methods, and improve the efficiency of model inference under lossless compression.

[0006] In a first aspect, the present invention provides a weight compression method for a large language model, comprising: Obtain the weight matrix of the large language model, construct a high-frequency window based on the exponential information distribution of the weight matrix, and determine the benchmark value; The weight matrix is ​​divided into weight sub-matrices adapted for hardware parallel processing. Each first weight data in each weight sub-matrix is ​​classified and judged: when the first weight data is high-frequency weight data, a fixed-length codeword corresponding to the first weight data is generated based on the exponent information of the high-frequency weight data and the benchmark value; when the first weight data is low-frequency weight data, the exponent information of the low-frequency weight data is mapped to the backoff codeword corresponding to the first weight data. Based on the fixed-length codeword and backoff codeword corresponding to each of the first weight data in each of the weight submatrices, initial compressed data corresponding to each of the weight submatrices is generated; wherein, the initial compressed data includes: bitmap, compact value stream and backoff value stream; After aligning the initial compressed data corresponding to each of the weight sub-matrices according to the hardware access requirements, the first weighted compressed data for storage in the global memory is obtained.

[0007] This application leverages the highly concentrated statistical characteristics of weight exponents to construct high-frequency windows and employs an encoding strategy of "baseline value + fixed-length codewords," abandoning the variable-length encoding used in existing technologies. This fundamentally eliminates the possibility of control flow divergence due to inconsistent code lengths during subsequent GPU parallel decoding. By using fixed-length codewords and backoff codewords, the uncertainty of the decoding path is eliminated from a data structure perspective. The generation of bitmaps, compact value streams, and backoff value streams reduces the storage footprint of high-frequency weights. The final alignment and packing process ensures that the compressed data meets the GPU's efficient merging and memory access requirements, providing a directly usable, hardware-aware compressed data format for subsequent high-performance parallel decoding computation.

[0008] Further, the step of generating initial compressed data corresponding to each weight submatrix based on the fixed-length codeword and backoff codeword corresponding to each of the first weight data in each weight submatrix specifically involves: The fixed-length codewords and backoff codewords corresponding to the first weight data in each weight submatrix are split bitwise to obtain each bitmap corresponding to each weight submatrix. When the first weight data in each of the weight submatrices is high-frequency weight data, the sign bit and mantissa bit corresponding to the high-frequency weight data are stored to obtain the compact value stream corresponding to each of the weight submatrices. When the first weight data in each of the weight submatrices is low-frequency weight data, the floating-point value corresponding to the low-frequency data is stored to obtain the backoff value stream corresponding to each of the weight submatrices.

[0009] By separating bit planes and storing the same bit in all codewords within the weight submatrix, GPU threads can read the same bit information of a large number of weights in a single memory transaction, improving memory access efficiency. By storing only the sign and mantissa bits of high-frequency weights, data storage is reduced, achieving compression. Simultaneously, storing the original values ​​of low-frequency weights ensures no precision loss is introduced during encoding, guaranteeing subsequent lossless reconstruction. This differentiated storage strategy achieves an optimal balance between compression ratio and data integrity.

[0010] Furthermore, the construction of a high-frequency window and determination of a benchmark value based on the exponential information distribution of the weight matrix specifically involves: The frequency distribution of the exponent is obtained by counting the frequency of each exponent value in the weight matrix. Select continuous exponential intervals from the exponential frequency distribution, calculate the cumulative frequency of each continuous exponential interval, and determine the continuous exponential interval with the highest cumulative frequency as the high-frequency window; The baseline value is obtained by subtracting the preset offset from the minimum value of the high-frequency window.

[0011] By statistically analyzing and selecting the continuous exponential interval with the highest cumulative frequency as the high-frequency window, the coverage and compression potential of the weighted data are maximized. By using the minimum value of the window minus one as the benchmark value, all exponents falling into the high-frequency window can be encoded as a small, non-zero positive integer offset (fixed-length codeword), while exponents outside the window are uniformly marked as zero (backoff codeword). The encoding logic is simple and efficient, and facilitates rapid decoding and numerical reconstruction by subsequent hardware.

[0012] Secondly, this application also provides a weighted inference operation method for a large language model, including: The second weight compressed data is loaded from the global memory to the shared memory; wherein the second weight compressed data is obtained by compressing the second weight matrix of the second large language model according to the weight compression method of the large language model as described in the first aspect above, and the second weight compressed data includes a second bitmap, a second compact value stream and a second backoff value stream; Based on the second bitmap, a status mask is generated in the register, and a channel prefix mask is generated based on the processor's thread index; Based on the state mask and the channel prefix mask, the read offset corresponding to each second weight data is calculated using the population counting instruction; Based on each read offset, the corresponding compressed data segment is read from the second compact value stream or the second backoff value stream, and based on each compressed data segment and the second bitmap, each second weight data is obtained; Each of the second weight data is input into the matrix multiplication unit to perform matrix multiplication.

[0013] This application enables inference computations to be performed directly in the compressed state by generating data in a specific compressed format. By utilizing state masks and channel prefix masks, combined with population counting instructions, branchless dynamic addressing is achieved, resolving the control flow divergence problem in traditional variable-length decoding and ensuring efficient parallel execution within GPU thread bundles. Most importantly, by fusing the reading, decoding, reconstruction, and matrix multiplication operations of compressed data at the register level, the redundant step of "decompressing data, writing it back to GPU memory, and then reading it again" is eliminated, minimizing GPU memory access and significantly alleviating the "memory wall" bottleneck, achieving the dual benefits of storage savings and computational acceleration.

[0014] Furthermore, the processor loads the second weighted compressed data from the global memory to the shared memory, specifically as follows: The second weighted compressed data is loaded into the shared memory in batches using an asynchronous copy method. When processing the current batch of compressed data, the next batch of compressed data is asynchronously prefetched into the shared memory.

[0015] This asynchronous copying approach overlaps data loading with core computation, effectively reducing data transfer latency from global memory to shared memory. Combined with double buffering or pipelined prefetching strategies, it ensures a continuous data supply to computing units, preventing idle units due to data waiting and further improving throughput and hardware utilization throughout the inference process.

[0016] Furthermore, the step of generating a state mask in the register based on the second bitmap specifically involves: The corresponding bitmap slice is read from the second bitmap into the register, and bit logic operations are performed on the bitmap slice to obtain the state mask; wherein, the value of each bit in the state mask is used to indicate whether the corresponding second weight data is in a compressed state or a back-up state.

[0017] This method efficiently synthesizes a mask identifying the state of each weighted data point by performing intra-register bitwise operations on bitmap slices stored separately in bit planes. This process is completed entirely within registers, without the need for branching judgments, providing crucial input for subsequent branchless parallel data routing and decoding based on the mask.

[0018] Further, the step of calculating the read offset corresponding to each second weight data point using the population counting instruction based on the state mask and the channel prefix mask specifically involves: Perform a bitwise AND operation between the state mask and the channel prefix mask to obtain the local channel prefix mask; The number of bits set in the local channel prefix mask is counted, and the resulting count is used as the read offset corresponding to the second weighted data processed by the current thread.

[0019] This approach leverages natively supported population counting instructions in conjunction with masking operations to achieve parallel prefix sum computation. This step dynamically calculates the position of the required data in the compact or backtracked value stream for each thread without introducing any conditional branches. This branchless parallel decoding ensures complete consistency of thread execution paths under the SIMT architecture.

[0020] Further, the step of reading the corresponding compressed data segment from the second compact value stream or the second backoff value stream according to the read offset specifically includes: When the second weighted data is in a compressed state, the corresponding sign bit and mantissa bit data segments are read from the second compact value stream using the read offset as an index. When the state of the second weight data is in the rollback state, the corresponding floating-point value of the second weight data is read from the second rollback value stream using the read offset as the index.

[0021] Based on the indication from the state mask, and using the calculated precise offset, corresponding data segments are read from different data streams. For compressed data, only the sign and mantissa portions are read; for backflip data, the complete original value is read. This step achieves branchless, parallel, and precise routing of data.

[0022] Furthermore, this application also proposes that before loading the second weighted compressed data from global memory to shared memory, the following steps are also included: Based on the length characteristics of the input sequence, determine whether the current inference stage is the pre-filling stage or the decoding stage; If it is the pre-filling stage, the second weight compressed data is decompressed into a third weight matrix, and the third weight matrix is ​​stored in the global memory. Then, the standard matrix multiplication library is called to perform matrix multiplication operations. If it is the decoding stage, the second weighted compressed data is loaded from the global memory to the shared memory.

[0023] This approach dynamically switches execution strategies based on the computational characteristics of the inference phase. During the computationally intensive pre-filling phase, a strategy of decompressing once and then calling a highly optimized standard computing library is employed to maximize computational throughput. During the memory-intensive decoding phase, this application is enabled to maximize bandwidth savings and increase generation speed, achieving global optimization of end-to-end inference performance.

[0024] Thirdly, this application also provides a weighted compression system for a large language model, including: a baseline determination module, an encoding processing module, a compressed data construction module, and a data storage module; The benchmark determination module is used to obtain the weight matrix of the large language model, construct a high-frequency window based on the exponential information distribution of the weight matrix, and determine the benchmark value. The encoding processing module is used to divide the weight matrix into weight sub-matrices adapted for hardware parallel processing, and to classify and judge each first weight data in each weight sub-matrix: when the first weight data is high-frequency weight data, a fixed-length codeword corresponding to the first weight data is generated based on the exponent information of the high-frequency weight data and the benchmark value; when the first weight data is low-frequency weight data, the exponent information of the low-frequency weight data is mapped to the backoff codeword corresponding to the first weight data. The compressed data construction module is used to generate initial compressed data corresponding to each weight submatrix based on the fixed-length codeword and backoff codeword corresponding to each first weight data in each weight submatrix; wherein, the initial compressed data includes: bitmap, compact value stream and backoff value stream; The data storage module is used to align the multiple independent bitmaps, the compact value stream, and the fallback value stream according to the hardware access requirements, and then obtain first weighted compressed data and store it in the global memory. This application employs a modular design, dividing the weight compression process into four stages: statistical analysis and benchmark determination, parallel encoding, compressed data generation, and aligned storage. The system can automatically and efficiently convert raw weights into a hardware-friendly compressed format, providing a standardized compressed data source for subsequent high-performance inference.

[0025] Fourthly, this application also provides a weighted inference operation system for a large language model, including: a data loading module, a decoding and reconstruction module, and a matrix operation module; The data loading module is used to load the second weight compressed data from the global memory to the shared memory; wherein the second weight compressed data is obtained by compressing the second weight matrix of the second large language model according to the weight compression method of the large language model as described in the first aspect, and the second weight compressed data includes a second bitmap, a second compact value stream and a second backoff value stream; The decoding and reconstruction module is used to generate a state mask in a register based on the second bitmap, and generate a channel prefix mask based on the processor's thread index; calculate the read offset corresponding to each second weight data using a population counting instruction based on the state mask and the channel prefix mask; read the corresponding compressed data segment from the second compact value stream or the second backoff value stream according to each read offset, and obtain each second weight data based on each compressed data segment and the second bitmap; The matrix operation module is used to directly input the second weight data into the matrix multiplication unit and perform matrix multiplication operations.

[0026] This application achieves the loading of compressed weights, branchless parallel decoding and reconstruction, and efficient integration with computing units through an integrated system design. This system serves as the specific hardware and software implementation of the weight inference calculation method for large language models described in the second aspect above. It can be deployed and run in actual GPU inference engines, improving inference efficiency and bandwidth utilization.

[0027] Fifthly, this application also provides a computer-readable storage medium, comprising: a stored computer program, wherein, when the computer program is executed, the device on which the computer-readable storage medium is located is controlled to perform the weight compression method for a large language model as described in the first aspect or the weight inference calculation method for a large language model as described in the second aspect. Attached Figure Description

[0028] To more clearly illustrate the technical solution of this application, the 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 from these drawings without creative effort.

[0029] Figure 1 This is a flowchart illustrating an embodiment of the weight compression method for large language models provided in this application; Figure 2 This is a schematic diagram of the compression and storage process of one embodiment of the weighted compression method for large language models provided in this application; Figure 3 This is a flowchart illustrating one embodiment of the weighted inference calculation method for a large language model provided in this application; Figure 4 This is a schematic diagram of the state mask generation process of one embodiment of the weighted inference operation method for the large language model provided in this application; Figure 5 This is a schematic diagram of an embodiment of the weighted compression system for a large language model provided in this application.

[0030] Figure 6 This is a schematic diagram of the structure of one embodiment of the weighted inference computation system for a large language model provided in this application.

[0031] Labeling Explanation: 301, Baseline Determination Module; 302, Encoding Processing Module; 303, Compressed Data Construction Module; 304, Data Storage Module; 401, Data Loading Module; 402, Decoding and Reconstruction Module; 403, Matrix Operation Module. Detailed Implementation

[0032] To make the objectives, technical solutions, and advantages of this application clearer, the technical solutions of this application will be clearly and completely described below with reference to the accompanying drawings of the embodiments. Obviously, the described embodiments are only some embodiments of this application, not all embodiments. Based on the embodiments of this application, all other embodiments obtained by those skilled in the art without creative effort are within the scope of protection of this application.

[0033] In the description of the embodiments of this application, technical terms such as "first" and "second" are used only to distinguish different objects and should not be construed as indicating or implying relative importance or implicitly specifying the number, specific order, or primary and secondary relationship of the indicated technical features. In the description of the embodiments of this application, "multiple" means two or more, unless otherwise explicitly defined.

[0034] In this document, the term "embodiment" means that a particular feature, structure, or characteristic described in connection with an embodiment may be included in at least one embodiment of this application. The appearance of this phrase in various places throughout the specification does not necessarily refer to the same embodiment, nor is it a separate or alternative embodiment mutually exclusive with other embodiments. It will be explicitly and implicitly understood by those skilled in the art that the embodiments described herein can be combined with other embodiments.

[0035] When traditional large language models are used for inference on graphics processing units (GPUs), the use of variable-length encoding can cause thread control flow divergence and redundant GPU memory reads and writes, which limits end-to-end inference performance and reduces overall computational efficiency.

[0036] Please refer to Figure 1 , Figure 1 This is a flowchart illustrating an embodiment of the weight compression method for large language models provided in this application. To address the problem of reduced model inference efficiency caused by thread control flow divergence due to the use of variable-length encoding in existing technologies, an embodiment of this application provides a weight compression method for large language models, including steps S101 to S104, each step as follows: Step S101: Obtain the weight matrix of the large language model, construct a high-frequency window based on the exponential information distribution of the weight matrix, and determine the benchmark value; Step S102: Divide the weight matrix into weight sub-matrices adapted for hardware parallel processing, and classify and judge each first weight data in each weight sub-matrix: when the first weight data is high-frequency weight data, generate a fixed-length codeword corresponding to the first weight data based on the exponent information and benchmark value of the high-frequency weight data; when the first weight data is low-frequency weight data, map the exponent information of the low-frequency weight data to the backoff codeword corresponding to the first weight data. Step S103: Generate initial compressed data corresponding to each weight submatrix based on the fixed-length codeword and backoff codeword corresponding to each first weight data in each weight submatrix; wherein, the initial compressed data includes: bitmap, compact value stream and backoff value stream; Step S104: After aligning the initial compressed data corresponding to each weight submatrix according to the hardware access requirements, the first weight compressed data is obtained for storage in the global memory.

[0037] In large language models, the weight matrix values ​​are usually floating-point numbers in BFloat16 format, representing the features and correlations learned by the model.

[0038] A high-frequency window refers to a range of index values ​​that appear relatively frequently and continuously in the distribution of index information. The index values ​​within this range represent the exponential characteristics of most of the data in the weight matrix.

[0039] A benchmark value is a reference value determined based on the characteristics of a high-frequency window. It is used to normalize or encode the exponential information of high-frequency weighted data in order to achieve fixed-length compression.

[0040] Fixed-length codewords refer to binary codewords with a fixed bit width generated after encoding the exponential information of high-frequency weighted data. Their length does not change with the original exponential value, which is beneficial for parallel decoding.

[0041] The backoff codeword refers to the codeword generated after mapping the exponential information of low-frequency weighted data. It is used to indicate that the data is low-frequency data and to guide the decoder to obtain the complete data from the backoff value stream.

[0042] In some preferred embodiments, all the bitmaps, compact value streams, and fallback value streams are packaged according to the cache line alignment requirements of the target hardware (such as a GPU), and the starting address of the compressed data packet in global memory is aligned to a 128-byte boundary. This alignment specification matches the access granularity of modern GPU memory controllers, ensuring that subsequent loading operations achieve the highest bandwidth utilization.

[0043] Specifically, in some embodiments, please refer to Figure 2 , Figure 2This is a schematic diagram of the compression and storage process of an embodiment of the weighted compression method for large language models provided in this application. In step S103, initial compressed data corresponding to each weighted submatrix is ​​generated based on the fixed-length codeword and backoff codeword corresponding to each first weight data in each weighted submatrix, specifically as follows: The fixed-length codewords and backoff codewords corresponding to the first weight data in each weight submatrix are split bitwise to obtain the bit diagrams corresponding to each weight submatrix. When the first weight data in each weight submatrix is ​​high-frequency weight data, store the sign bit and mantissa bit corresponding to the high-frequency weight data to obtain the compact value stream corresponding to each weight submatrix. When the first weight data in each weight submatrix is ​​low-frequency weight data, store the floating-point value corresponding to the low-frequency data to obtain the backoff value stream corresponding to each weight submatrix.

[0044] The following is a specific example illustrating the compression process of an 8×8 BFloat16 format weight submatrix. First, a high-frequency window and a baseline value are determined. Then, 3-bit fixed-length codewords are generated for each of the 64 weight data points in the 8×8 weight submatrix. Codewords 001 to 111 indicate that the weight data is "high-frequency hit" data, and its exponential information can be reconstructed using "baseline value + codeword"; codeword 000 indicates that the weight is "backoff" data.

[0045] To adapt to the parallel memory access characteristics of GPUs, a bit-plane separation method is used to store these codewords. The 64 3-bit codewords are split into three independent bit planes according to their bits: the 0th bit of the 64 codewords is arranged sequentially to form a 64-bit binary sequence, which is stored as the first bit plane; the 1st bit of the 64 codewords is arranged sequentially to form a second 64-bit binary sequence, which is stored as the second bit plane; and the 2nd bit of the 64 codewords is arranged sequentially to form a third 64-bit binary sequence, which is stored as the third bit plane.

[0046] With the aforementioned bit-plane separation storage, when a GPU warp (typically containing 32 threads) needs to read codewords for multiple weights, it can efficiently read the 0th, 1st, and 2nd bits of all 32 target weights at once through three aligned and merged memory accesses. This design, which fixes the bitmap length to the word length natively supported by the GPU (e.g., 64 bits), allows the entire bitmap to be read with a single memory load instruction, thereby reducing the number of memory transactions and improving data loading efficiency.

[0047] For the weight data identified as "high-frequency hits" in the weight submatrix, the sign bit (1 bit) and mantissa bit (7 bits) are extracted from its BFloat16 format data. These extracted bits are then tightly packed and stored contiguously according to their original weight order within the submatrix, forming a compact value stream. Since the exponent information of the "high-frequency hit" weight data is deterministically represented through a mapping relationship of "base value + fixed-length codeword," its original 8-bit exponent can be safely discarded during storage. By storing only the sign bit and mantissa bit, the storage overhead for each "high-frequency hit" weight data is reduced from the original 16 bits to 8 bits, achieving effective compression of high-frequency data.

[0048] For weight data marked as "back off" (i.e., codeword 000) in the weight submatrix, its complete original BFloat16 format data is preserved and stored sequentially in the order it appears in the submatrix as another independent data stream, namely the back off value stream. The back off value stream ensures that all weight data that does not conform to high-frequency statistical characteristics are completely preserved. During decoding, the "back off" data will be read directly from the original value in this stream without any approximation or calculation.

[0049] Through the above steps, for an 8×8 BFloat16 weighted submatrix, the generated initial compressed data will be stored in memory with the data components aligned according to predetermined rules.

[0050] Specifically, in some embodiments, in step S101, a high-frequency window is constructed based on the exponential information distribution of the weight matrix and a benchmark value is determined, specifically as follows: The frequency distribution of the exponent is obtained by counting the frequency of each exponent value in the weight matrix. Select continuous exponential intervals from the exponential frequency distribution, calculate the cumulative frequency of each continuous exponential interval, and determine the continuous exponential interval with the highest cumulative frequency as the high-frequency window; The baseline value is obtained by subtracting the preset offset from the minimum value of the high-frequency window.

[0051] The frequency of each index value in the statistical weight matrix is ​​determined by traversing all weight data in the matrix and extracting an 8-bit index value (range 0-255) for each weight data. The number of times each possible index value appears in the entire matrix is ​​counted, thus constructing an index frequency distribution histogram. This histogram visually reflects the probability of different index values ​​appearing. After obtaining the index frequency distribution, instead of simply selecting the most frequent discrete index values, a selection strategy based on continuous intervals is adopted to simplify the subsequent encoding and decoding logic.

[0052] In some preferred embodiments, to achieve the best balance between compression ratio and encoding / decoding complexity, the number k of consecutive exponential values ​​covered by the high-frequency window is set to 7. Statistics show that this setting can cover the vast majority of exponential distributions for the weight matrix of a typical large language model in BFloat16 format, achieving a high compression ratio. A window with a preset window length K=7 slides across the exponential value range (e.g., 0-255). For each possible starting position i, the sum of the frequencies corresponding to the seven consecutive exponential values ​​i, i+1, ..., i+6 is calculated, which is the cumulative coverage frequency of the window. The cumulative coverage frequencies of all sliding windows are compared, and the consecutive exponential interval with the highest cumulative coverage frequency is selected as the "high-frequency window". After determining the high-frequency window, a baseline value is calculated according to the formula: Baseline value = Minimum value of high-frequency window − 1.

[0053] The obtained high-frequency window range and benchmark value will be used as global parameters to guide the subsequent encoding operation of each weight data in the weight matrix, that is, to determine whether it falls into the high-frequency window and generate the corresponding fixed-length codeword.

[0054] refer to Figure 3 , Figure 3 This is a flowchart illustrating an embodiment of the weighted inference operation method for a large language model provided in this application. To address the problem of reduced model inference efficiency caused by redundant memory read / write operations due to variable-length encoding in existing technologies, an embodiment of this application provides a weighted inference operation method for a large language model, including steps S201 to S205, each step as follows: Step S201: Load the second weight compressed data from the global memory to the shared memory; wherein, the second weight compressed data is obtained by compressing the second weight matrix of the second large language model according to the weight compression method of the large language model described above, and the second weight compressed data includes the second bitmap, the second compact value stream and the second backoff value stream; Step S202: Based on the second bitmap, generate a status mask in the register and generate a channel prefix mask based on the processor's thread index; Step S203: Calculate the read offset corresponding to each second weight data using the population counting instruction based on the state mask and the channel prefix mask; Step S204: Read the corresponding compressed data segments from the second compact value stream or the second backoff value stream according to each read offset, and obtain each second weight data based on each compressed data segment and the second bitmap; Step S205: Input each second weight data into the matrix multiplication unit and perform matrix multiplication operation.

[0055] Among them, from step S202 to step S204, that is, from state mask generation, offset calculation, to data segment reading and floating-point reconstruction, the entire decoding process is completed entirely within the registers of the GPU thread. The reconstructed weight values directly flow from the registers into the input registers of the matrix multiplication unit (such as Tensor Core), and are never written back to the shared memory or global memory. This design eliminates the redundant data movement in the traditional "decompression - write-back - reread" architecture.

[0056] Loading the second weight compressed data into the shared memory reduces global memory access and saves video memory bandwidth. The generation of the state mask and channel prefix mask enables each thread to independently calculate the offset, avoiding thread synchronization waiting. The population count instruction ensures fast and synchronous offset calculation, supporting parallel decoding. The data reconstruction is based on the second bitmap, ensuring losslessness. Finally, it is directly input into the matrix multiplication unit for operation, improving the GPU inference efficiency and bandwidth utilization of the large language model.

[0057] The following is illustrated by a specific example. In this embodiment, taking the processing of an 8×8 BFloat16 weight sub-matrix as an example, after the operator is started, the compressed data block corresponding to the 8×8 weight sub-matrix to be processed is loaded from the global memory to the shared memory. Taking the thread numbered T0 in the warp as an example, it is responsible for processing 8 weight elements from the 0th to the 7th in this weight sub-matrix. T0 reads the bitmap slices related to the elements it is responsible for in the three bitmaps from the shared memory into its private register. By performing a bitwise "or" operation on these slices, T0 synthesizes an 8-bit state mask in its register. For example, the synthesis result may indicate that among the 8 elements it is responsible for, the state bits of the 2nd and 5th elements are '1' (fallback state), and the state bits of the remaining elements are '0' (compressed state).

[0058] T0 generates a channel prefix mask in the register by calculating (1<<lane_id) - 1 according to its relative index (lane_id) within the warp.

[0059] T0 then performs branchless dynamic addressing calculation. For the 2nd element (fallback state), T0 performs a bitwise AND operation on the state mask and a local prefix mask that is '1' only at the position before the 2nd element, and then uses the hardware population count instruction to count the number of '1's in the result. Assuming the count result is 0, the read offset is 0, which means the 2nd element is the first fallback state element that T0 is responsible for and should be read from the starting position of the fallback value stream. For the 5th element (fallback state), the offset is calculated to be 1 in the same way, indicating that it is the second fallback state element that T0 is responsible for. Similarly, the offsets for reading compressed state elements from the compact value stream can be calculated.

[0060] Based on the calculated precise offsets, T0 reads the required segments from the data stream in shared memory into the register in parallel: for the 2nd and 5th backflip-state elements, two complete raw BFloat16 values ​​are read directly from the second backflip-value stream, indexed by their offsets (0 and 1). For the remaining 6 compressed-state elements, 6 8-bit compressed data segments (sign + mantissa) are read from the second compact-value stream, indexed by their offsets in the compressed stream.

[0061] Subsequently, T0 performs arithmetic and bit operations in the register: for compressed elements, the 3-bit codeword decoded from the bitmap is added to the base value to reconstruct the original exponent; then the reconstructed exponent is concatenated with the sign bit and mantissa bit read from the compact stream to recover 6 complete BFloat16 values ​​without loss.

[0062] At this point, T0's registers already hold the eight fully reconstructed BF16 weight values ​​it is responsible for. These values ​​are directly used as operands to fill the input register segment of the matrix multiplication unit (such as the Tensor Core). The thread bundle containing T0 then executes the Tensor Core instructions to perform high-speed matrix multiplication and accumulation operations on these weights and their corresponding activation values. Finally, only the accumulated result of the matrix multiplication is written back to global memory. The intermediate weight data recovered from the compressed state resides within the chip core's registers or computation units throughout the entire process and is never written back to any level of memory as independent data.

[0063] In some other embodiments, in step S201, the processor loads the second weighted compressed data from the global memory to the shared memory, specifically as follows: The second-weighted compressed data is loaded into the shared memory in batches using an asynchronous copy method. While processing the current batch of compressed data, the next batch of compressed data is asynchronously prefetched into the shared memory.

[0064] In this process, asynchronous copy instructions (such as cp.async) are used to move the weighted compressed data from global memory to on-chip shared memory. Unlike traditional synchronous loading, asynchronous copying allows the computing core to continue executing subsequent computation instructions that do not depend on the data immediately after issuing the data load instruction.

[0065] To completely mask the latency of data transfer from global memory to shared memory, the system employs a double-buffering strategy: dividing shared memory into two logical buffers (Buffer A and Buffer B). While the computation core is processing a compressed data block in Buffer A, the `cp.async` instruction simultaneously prefetches the next data block into Buffer B. Once the current computation is complete, the buffers switch roles, thus achieving continuous overlap between computation and memory access.

[0066] Specifically, in some embodiments, in step S202, a state mask is generated in the register based on the second bitmap, specifically as follows: The corresponding bitmap slice is read from the second bitmap into the register, and bit logic operations are performed on the bitmap slice to obtain the state mask; wherein, the value of each bit in the state mask is used to indicate whether the corresponding second weight data is in the compressed state or the back-off state.

[0067] Each thread performs a bitwise OR operation on the three bitmap slices it reads in its register to generate a state mask.

[0068] The following is a specific example for illustration; please refer to it. Figure 4 , Figure 4 This is a schematic diagram of the state mask generation process of this application. This embodiment takes processing a weight submatrix as an example to describe how multiple threads in a GPU thread bundle collaborate to generate their respective state masks from a second bitmap array. Assume the second bitmap array consists of three bitmaps, each a 64-bit integer. A thread bundle containing 32 threads is assigned to process this weight submatrix. Each thread is responsible for processing a continuous data segment within a micro-block. For example, thread T0 is responsible for weight data from 0 to 7, thread T1 for weight data from 8 to 15, and so on. Each thread reads the corresponding bitmap slices from the three bitmaps in shared memory into its private register according to the index range of the weight data it is responsible for. Each thread performs a bitwise OR operation on the three bitmap slices read in its register; the result of this operation generates a state mask, the number of bits of which is equal to the number of weights the thread is responsible for.

[0069] In the generated state mask, each bit directly corresponds to a weight that the thread is responsible for. A bit value of "1" indicates that the weight data is in a "compressed state," and its complete exponent information can be restored through "base value + codeword," while the sign and mantissa are stored in a compact value stream. If a bit is "0," it indicates that the exponent of the weight data does not fall into the high-frequency window, and its complete original value is stored in the backtracking value stream.

[0070] Specifically, in some embodiments, in step S203, the read offset corresponding to each second weight data is calculated using a population counting instruction based on the state mask and the channel prefix mask, specifically as follows: Perform a bitwise AND operation between the state mask and the channel prefix mask to obtain the local channel prefix mask; The number of bits set in the local channel prefix mask is counted, and the resulting count is used as the read offset corresponding to the second weight data processed by the current thread.

[0071] Among them, each thread generates a channel prefix mask according to its thread index (lane_id) within the warp. The generation rule of this mask is: (1<<lane_id)-1. Its function is to create a binary mask. In this mask, all bits lower than the current thread lane_id are set to 1, and the current lane_id and higher bits are set to 0.

[0072] Specifically, in some embodiments, in step S204, corresponding compressed data segments are read from the second compact value stream or the second fallback value stream according to each read offset respectively, specifically: When the state of the second weight data is the compressed state, taking the read offset as an index, read the corresponding symbol bit and tail bit data segments from the second compact value stream; When the state of the second weight data is the fallback state, taking the read offset as an index, read the corresponding floating-point value of the second weight data from the second fallback value stream.

[0073] Among them, for compressed state data, the thread performs the following bit operations in the register to accurately reconstruct the original BFloat16 value: The reconstructed value = (symbol bit << 15) | ((reference value + code word) << 7) | tail bit Where << represents the left shift operation, and | represents the bitwise OR operation. This formula ensures a bit-exact lossless reconstruction from the compressed representation to the original binary data.

[0074] In some other embodiments, before loading the second weight compressed data from the global memory to the shared memory, it further includes: Judging whether the current inference stage is the prefill stage or the decoding stage according to the length characteristics of the input sequence; If it is the prefill stage, decompress the second weight compressed data into a third weight matrix consistent with the second weight matrix of the second large language model, store the third weight matrix, and call the standard matrix multiplication library to perform matrix multiplication operations; If it is the decoding stage, load the second weight compressed data from the global memory to the shared memory.

[0075] Among them, the prefill stage refers to processing the entire prompt sequence of the user input, and it is necessary to calculate the attention scores and mappings of all input tokens and all weights. This stage is computationally intensive and has a large amount of computation, but the number of accesses to each weight data is relatively small. Therefore, adopting the strategy of decompressing once and calling a highly optimized standard matrix multiplication library (such as cuBLAS) can maximize the utilization of computing resources.

[0076] The decoding stage refers to the autoregressive generation of each output token (word-by-word generation). This stage calculates only one new token at a time, resulting in relatively low computational cost, but it requires repeatedly and frequently reading the entire weight matrix from GPU memory. The performance of this stage is severely limited by GPU memory bandwidth. Therefore, by switching to the ZipGEMM fusion operator of this application, which eliminates redundant data transport through "load-and-compute," bandwidth pressure can be effectively alleviated and generation speed improved.

[0077] Please refer to Figure 5 , Figure 5 This is a schematic diagram of the structure of a weighted compression system for a large language model in some embodiments of this application. The system includes: a baseline determination module 301, an encoding processing module 302, a compressed data construction module 303, and a data storage module 304. The benchmark determination module 301 is used to obtain the weight matrix of the large language model, construct a high-frequency window based on the exponential information distribution of the weight matrix, and determine the benchmark value. The encoding processing module 302 is used to divide the weight matrix into weight sub-matrices adapted for hardware parallel processing, and to classify and judge each first weight data in each weight sub-matrix: when the first weight data is high-frequency weight data, a fixed-length codeword corresponding to the first weight data is generated based on the exponent information and benchmark value of the high-frequency weight data; when the first weight data is low-frequency weight data, the exponent information of the low-frequency weight data is mapped to the backoff codeword corresponding to the first weight data. The compressed data construction module 303 is used to generate the initial compressed data corresponding to each weight submatrix based on the fixed-length codeword and backoff codeword corresponding to each first weight data in each weight submatrix; wherein, the initial compressed data includes: bitmap, compact value stream and backoff value stream; The data storage module 304 is used to align multiple independent bitmaps, compact value streams and fallback value streams according to the hardware access requirements, obtain first weighted compressed data, and store it in the global memory.

[0078] It is understood that the above system item embodiments correspond to the large language model weight compression method item embodiments of this application, and can implement the large language model weight compression method provided by any of the above method item embodiments of this application.

[0079] Please refer to Figure 6 , Figure 6 This is a schematic diagram of the structure of a large language model reasoning system in some embodiments of this application. The system includes: a data loading module 401, a decoding and reconstruction module 402, and a matrix operation module 403. The data loading module 401 is used to load the second weight compressed data from the global memory to the shared memory; wherein, the second weight compressed data is obtained by compressing the second weight matrix of the second large language model according to the weight compression method of any large language model in this application, and the second weight compressed data includes a second bitmap, a second compact value stream and a second backoff value stream; The decoding and reconstruction module 402 is used to generate a state mask in a register based on the second bitmap and generate a channel prefix mask based on the thread index of the processor; calculate the read offset corresponding to each second weight data using a population counting instruction based on the state mask and the channel prefix mask; read the corresponding compressed data segment from the second compact value stream or the second backoff value stream according to each read offset; and obtain each second weight data based on each compressed data segment and the second bitmap. The matrix operation module 403 is used to directly input the second weight data into the matrix multiplication unit and perform matrix multiplication operations.

[0080] It is understood that the above system item embodiments correspond to the reasoning operation method item embodiments of the large language model of this application, and can implement the weighted reasoning operation method of the large language model provided by any of the above method item embodiments of this application.

[0081] It should be noted that the system embodiments described above are merely illustrative, and some or all of the modules can be selected to achieve the purpose of this embodiment according to actual needs. Furthermore, in the accompanying drawings of the system embodiments provided in this application, the connection relationships between modules indicate that they have communication connections, which can be specifically implemented as one or more communication buses or signal lines. Those skilled in the art can understand and implement this without any creative effort.

[0082] Based on the above-described method embodiments, another embodiment of this application provides a computer-readable storage medium including a stored computer program, wherein, when the computer program is executed, it controls the device where the computer-readable storage medium is located to execute the weight compression method for a large language model or the weight inference operation method for a large language model as described in any of the above-described method embodiments of this application.

[0083] The above description is the preferred 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 are also considered to be within the scope of protection of this application.

Claims

1. A weight compression method for a large language model, characterized in that, include: Obtain the weight matrix of the large language model, construct a high-frequency window based on the exponential information distribution of the weight matrix, and determine the benchmark value; The weight matrix is ​​divided into weight sub-matrices adapted for hardware parallel processing. Each first weight data in each weight sub-matrix is ​​classified and judged: when the first weight data is high-frequency weight data, a fixed-length codeword corresponding to the first weight data is generated based on the exponent information of the high-frequency weight data and the benchmark value; when the first weight data is low-frequency weight data, the exponent information of the low-frequency weight data is mapped to the backoff codeword corresponding to the first weight data. Based on the fixed-length codeword and backoff codeword corresponding to each of the first weight data in each of the weight submatrices, initial compressed data corresponding to each of the weight submatrices is generated; wherein, the initial compressed data includes: bitmap, compact value stream and backoff value stream; After aligning the initial compressed data corresponding to each of the weight sub-matrices according to the hardware access requirements, the first weighted compressed data for storage in the global memory is obtained.

2. The weight compression method for large language models according to claim 1, characterized in that, The step of generating initial compressed data corresponding to each weight submatrix based on the fixed-length codeword and backoff codeword corresponding to each first weight data in each weight submatrix is ​​specifically as follows: The fixed-length codewords and backoff codewords corresponding to the first weight data in each weight submatrix are split bitwise to obtain each bitmap corresponding to each weight submatrix. When the first weight data in each of the weight submatrices is high-frequency weight data, the sign bit and mantissa bit corresponding to the high-frequency weight data are stored to obtain the compact value stream corresponding to each of the weight submatrices. When the first weight data in each of the weight submatrices is low-frequency weight data, the floating-point value corresponding to the low-frequency data is stored to obtain the backoff value stream corresponding to each of the weight submatrices.

3. The weight compression method for large language models according to claim 1, characterized in that, The process of constructing a high-frequency window and determining a benchmark value based on the exponential information distribution of the weight matrix is ​​as follows: The frequency distribution of the exponent is obtained by counting the frequency of each exponent value in the weight matrix. Select continuous exponential intervals from the exponential frequency distribution, calculate the cumulative frequency of each continuous exponential interval, and determine the continuous exponential interval with the highest cumulative frequency as the high-frequency window; The baseline value is obtained by subtracting the preset offset from the minimum value of the high-frequency window.

4. A weighted inference calculation method for a large language model, characterized in that, include: The second weight compressed data is loaded from the global memory to the shared memory; wherein the second weight compressed data is obtained by compressing the second weight matrix of the second large language model according to the weight compression method of the large language model as described in any one of claims 1-3, and the second weight compressed data includes a second bitmap, a second compact value stream, and a second backoff value stream; Based on the second bitmap, a status mask is generated in the register, and a channel prefix mask is generated based on the processor's thread index; Based on the state mask and the channel prefix mask, the read offset corresponding to each second weight data is calculated using the population counting instruction; Based on each read offset, the corresponding compressed data segment is read from the second compact value stream or the second backoff value stream, and based on each compressed data segment and the second bitmap, each second weight data is obtained; Each of the second weight data is input into the matrix multiplication unit to perform matrix multiplication.

5. The weighted inference operation method for a large language model according to claim 4, characterized in that, The processor loads the second-weighted compressed data from the global memory to the shared memory, specifically as follows: The second weighted compressed data is loaded into the shared memory in batches using an asynchronous copy method. When processing the current batch of compressed data, the next batch of compressed data is asynchronously prefetched into the shared memory.

6. The weighted inference operation method for a large language model according to claim 4, characterized in that, The step of generating a state mask in the register based on the second bitmap specifically involves: The corresponding bitmap slice is read from the second bitmap into the register, and bit logic operations are performed on the bitmap slice to obtain the state mask; wherein, the value of each bit in the state mask is used to indicate whether the corresponding second weight data is in a compressed state or a back-up state.

7. The weighted inference operation method for a large language model according to claim 4, characterized in that, The step of calculating the read offset corresponding to each second weight data point based on the state mask and channel prefix mask using a population counting instruction is as follows: Perform a bitwise AND operation between the state mask and the channel prefix mask to obtain the local channel prefix mask; The number of bits set in the local channel prefix mask is counted, and the resulting count is used as the read offset corresponding to the second weighted data processed by the current thread.

8. The weighted inference operation method for a large language model according to claim 4, characterized in that, The step of reading the corresponding compressed data segment from the second compact value stream or the second backoff value stream according to each read offset is specifically as follows: When the second weighted data is in a compressed state, the corresponding sign bit and mantissa bit data segments are read from the second compact value stream using the read offset as an index. When the state of the second weight data is in the rollback state, the corresponding floating-point value of the second weight data is read from the second rollback value stream using the read offset as the index.

9. The weighted inference operation method for a large language model according to claim 4, characterized in that, Before loading the second weighted compressed data from global memory to shared memory, the method further includes: Based on the length characteristics of the input sequence, determine whether the current inference stage is the pre-filling stage or the decoding stage; If it is the pre-filling stage, the second weight compressed data is decompressed into a third weight matrix, and the third weight matrix is ​​stored in the global memory. Then, the standard matrix multiplication library is called to perform matrix multiplication operations. If it is the decoding stage, the second weighted compressed data is loaded from the global memory to the shared memory.

10. A weighted compression system for a large language model, characterized in that, include: The module includes a baseline determination module, an encoding processing module, a compressed data construction module, and a data storage module. The benchmark determination module is used to obtain the weight matrix of the large language model, construct a high-frequency window based on the exponential information distribution of the weight matrix, and determine the benchmark value. The encoding processing module is used to divide the weight matrix into weight sub-matrices adapted for hardware parallel processing, and to classify and judge each first weight data in each weight sub-matrix: when the first weight data is high-frequency weight data, a fixed-length codeword corresponding to the first weight data is generated based on the exponent information of the high-frequency weight data and the benchmark value; when the first weight data is low-frequency weight data, the exponent information of the low-frequency weight data is mapped to the backoff codeword corresponding to the first weight data. The compressed data construction module is used to generate initial compressed data corresponding to each weight submatrix based on the fixed-length codeword and backoff codeword corresponding to each first weight data in each weight submatrix; wherein, the initial compressed data includes: bitmap, compact value stream and backoff value stream; The data storage module is used to align the multiple independent bitmaps, the compact value stream, and the fallback value stream according to the hardware access requirements, and then obtain first weighted compressed data and store it in the global memory.

11. A weighted inference computation system for a large language model, characterized in that, include: Data loading module, decoding and reconstruction module, and matrix operation module; The data loading module is used to load the second weight compressed data from the global memory to the shared memory; wherein, the second weight compressed data is obtained by compressing the second weight matrix of the second large language model according to the weight compression method of the large language model as described in any one of claims 1-3, and the second weight compressed data includes a second bitmap, a second compact value stream, and a second backoff value stream; The decoding and reconstruction module is used to generate a state mask in a register based on the second bitmap, and generate a channel prefix mask based on the processor's thread index; calculate the read offset corresponding to each second weight data using a population counting instruction based on the state mask and the channel prefix mask; read the corresponding compressed data segment from the second compact value stream or the second backoff value stream according to each read offset, and obtain each second weight data based on each compressed data segment and the second bitmap; The matrix operation module is used to directly input the second weight data into the matrix multiplication unit and perform matrix multiplication operations.

12. A computer-readable storage medium, characterized in that, include: A stored computer program, wherein, when the computer program is executed, it controls the device containing the computer-readable storage medium to perform the weight compression method of the large language model as described in any one of claims 1 to 3 or the weight inference operation method of the large language model as described in any one of claims 4 to 9.

Images