Matrix operation system supporting mask processing, operation method, electronic equipment and readable storage medium

Abstract

The invention relates to a matrix operation system supporting mask processing, an operation method, electronic equipment and a readable storage medium, and the matrix operation system comprises a matrix operation module which is used for executing matrix operation on an input matrix and outputting a matrix operation result; the mask processing module is arranged at the output end of the matrix operation module, and the mask processing module is used for executing mask processing on the matrix operation result; the configuration module is used for configuring parameters of the mask processing module; the result output module is used for outputting a result after mask processing; wherein the mask processing module comprises: a vector register to store a mask count; a scalar register for storing a specified value; and the mask execution module is used for performing mask processing on the matrix operation result according to the mask count in the vector register and the specified value in the scalar register. According to the matrix operation system provided by the invention, the mask processing module is introduced, so that matrix operation and mask operation can be completed at the same time, and the calculation overhead and the occupation of hardware resources are reduced.

Description

Technical Field

[0001] The present invention relates to the field of artificial intelligence technology, and in particular to a matrix operation system supporting mask processing, an operation method, an electronic device, and a readable storage medium. Background Art

[0002] With breakthroughs in existing large models, AI application scenarios have further expanded. Currently, how AI accelerators can accelerate AI models is a key research direction. AI accelerators play a vital role in improving computing efficiency and reducing energy consumption. AI accelerators are primarily used for the training and inference of deep learning models, especially for accelerating large-scale matrix operations. The Transformer architecture is the core technical principle of existing large models, and the key operator is attention. The calculation formula for attention is:

[0003] attention(Q,K,V)=softmax(QK T )V

[0004] Where Q is the query matrix, K is the key matrix, and V is the value matrix. Each of Q, K, and V is a matrix with N rows, corresponding to the number of tokens in the input sequence. Each column is a vector of length D, corresponding to the features of the token embedded in the high-dimensional space. By calculating the similarity between the query matrix (Q), the key matrix (K), and the value matrix (V), attention weights are generated to extract features from the input sequence.

[0005] The calculation of attention is divided into 3 steps: the first step QK T Calculate the orthogonal similarity between each row in Q and each row in K to obtain a similarity matrix of size N*N. The second step is softmax(QK T ) normalizes the similarity to ensure that each row is a probability vector with a value of 0 to 1 and a sum of 1. The third step is softmax(QK T )V uses the previous step probability as the weight and performs weighted summation on the vectors of each row in V to obtain the final result.

[0006] In time-step-based models such as text generation, the tokens corresponding to each row in Q, K, and V are generated or given in chronological order, and all historical data are collected for calculation at the current moment. During calculation, the calculation result of historical data at a certain moment may be affected by the data generated later. Therefore, in the first step, QK T After the calculation, the model will introduce a causal mask (causal-mask). The causal mask will be QK TAll elements above the main diagonal of the calculation result are set to negative infinity. This ensures that the output of the second step still meets the requirement of a probability vector for each row, and also ensures that the similarity weight of each row of data with the data generated after it is 0.

[0007] Existing AI accelerators are already relatively mature in matrix operations (first and third steps) and vector calculations (second step).

[0008] The existing technology sets the upper triangular part of the result of matrix multiplication to negative infinity. The existing technology includes:

[0009] 1. Implementation based on scalar instructions. Using scalar instructions, loop through the upper triangular part of the matrix multiplication result and set it to negative infinity, while the rest retain the original value. The following is the C language code to implement this process:

[0010]

[0011] 2. Implementation based on SIMD vector instructions. Using vector instructions, a single instruction sets multiple locations to negative infinity. The following is the C language code based on the riscv-vector vector instruction set that implements this process:

[0012]

[0013]

[0014] 3. Implementation based on SIMT instructions. Using SIMT instructions, multiple threads operate on different matrix regions, and the thread ID determines whether the region should retain its original value or be set to negative infinity. The following is the CUDA code that implements this process:

[0015]

[0016] To accelerate computation, AI accelerators must maximize hardware parallelization and vectorization capabilities. Each execution component must also fully utilize possible concurrency opportunities. In attention scenarios, the aforementioned causal mask calculation solutions are insufficient.

[0017] 1. Scalar calculations do not utilize any parallelization or vectorization acceleration capabilities. For an N*N matrix multiplication result, causal masking with scalar calculations requires approximately 1 / 2N*N scalar writes.

[0018] 2. SIMD / SIMT implementations occupy vector / multi-threaded execution units. Accelerating causal mask calculations using SIMD / SIMT utilizes hardware acceleration capabilities, but execution time remains proportional to N*N. Furthermore, causal mask calculations occupy the AI ​​accelerator's vector / multi-threaded execution units, blocking subsequent softmax calculations and increasing overall computation time.

[0019] Existing patents (such as CN115859011 B) propose using masks to control the range of matrix operations, but they focus on attention window partitioning and do not address the hardware acceleration of causal masks. CN118871902A also addresses sparse mask optimization but does not incorporate pipeline design for matrix multiplication. Therefore, an efficient, flexible, and hardware-friendly causal mask processing solution is urgently needed. Summary of the Invention

[0020] The present invention aims to provide a matrix operation system supporting mask processing to solve the problems of high computational overhead and high hardware resource occupation in the prior art causal mask implementation scheme.

[0021] To achieve the above object, the present invention first provides a matrix operation system supporting mask processing, comprising:

[0022] A matrix operation module, which is used to perform matrix operations on input matrices and output matrix operation results;

[0023] a mask processing module, which is provided at the output end of the matrix operation module, and is used to perform mask processing on the matrix operation result;

[0024] a configuration module, configured to configure parameters of the mask processing module; and

[0025] A result output module, which is used to output the result after mask processing;

[0026] The mask processing module includes:

[0027] A vector register, which is used to store the mask count;

[0028] scalar registers, which are used to store specified values; and

[0029] A mask execution module is used to perform mask processing on the matrix operation result according to the mask count in the vector register and the specified value in the scalar register.

[0030] Furthermore, the mask count is a vector, and the length of the mask count is the same as the number of columns of the matrix operation result.

[0031] Furthermore, when mask_count[i]≥0, the mask_count[i] rows before the i-th column of the matrix operation result are set to x;

[0032] When mask_count[i]<0, the first -mask_count[i]-1 rows of the i-th column of the matrix operation result are retained, and the results of the remaining rows of the matrix operation result are set to x;

[0033] Wherein, mask_count[i] is the value of the i-th position of the vector register, and x is the specified value in the scalar register.

[0034] Furthermore, the configuration module includes:

[0035] A mask count configuration module, configured to configure the mask count;

[0036] The specified value configuration module is used to configure the specified value.

[0037] Furthermore, the matrix operation and the mask processing are performed in parallel through a three-stage pipeline, including:

[0038] A calculation module, which is used to generate matrix operation results row by row;

[0039] a masking module for applying a column mask to the current row of data; and

[0040] The write module is used to write the masked data into the memory.

[0041] Furthermore, the matrix operation module, the mask processing module and the configuration module are connected via a bus.

[0042] Furthermore, the matrix operation system is a matrix operation system based on CPU, GPU or ASIC.

[0043] The present invention also provides a method for operating the matrix operation system supporting mask processing, the method comprising the following steps:

[0044] Configure the parameters of the mask processing module, including mask count and specified value;

[0045] Input the input matrix to the matrix operation module;

[0046] The matrix operation module performs matrix operations and outputs matrix operation results to the mask processing module;

[0047] The mask processing module performs mask processing on the matrix operation result according to the mask count and the specified value; and

[0048] Output the masked result.

[0049] The present invention further provides an electronic device, comprising:

[0050] a processor configured to execute machine-readable instructions;

[0051] A graphics card with an artificial intelligence chip, configured to train the method for operating the matrix operation system supporting mask processing; and

[0052] The memory is configured to store machine-readable instructions, which, when executed by a processor and / or a graphics card, perform the steps of the method for operating the matrix operation system supporting mask processing.

[0053] The present invention also provides a computer-readable storage medium having machine-readable instructions stored thereon. When the machine-readable instructions are executed by a processor, the steps of the operating method of the matrix operation system supporting mask processing are executed.

[0054] The technical solution provided by the present invention has the following advantages:

[0055] 1. The matrix operation system supporting mask processing provided by the present invention introduces a mask processing module, so that the matrix operation system can simultaneously complete matrix operations and mask operations, reducing computing overhead and hardware resource usage.

[0056] 2. For a matrix of size NxN, the additional overhead introduced by the mask operation is O(N 2 / P), where P is the degree of hardware parallelism. The matrix operation system supporting mask processing provided by the present invention introduces a new hardware structure and pipeline design, so that the additional overhead is only one clock cycle, thereby greatly reducing the additional overhead introduced by the matrix operation system for mask operations.

[0057] 3. The matrix operation system supporting mask processing provided by the present invention has a flexibly configurable mask processing module. It is not limited to causal masking, but can also perform other operations such as lower triangle masking and custom masking, thereby improving computing efficiency and enhancing hardware flexibility and adaptability. BRIEF DESCRIPTION OF THE DRAWINGS

[0058] To further illustrate the above and other advantages and features of various embodiments of the present invention, a more detailed description of various embodiments of the present invention will be presented with reference to the accompanying drawings. It will be understood that these drawings depict only typical embodiments of the present invention and are not to be considered as limiting the scope thereof. In the drawings, for clarity, identical or corresponding components will be represented by the same or similar reference numerals.

[0059] Figure 1 A schematic structural diagram of a matrix operation system supporting mask processing according to an embodiment of the present invention is shown;

[0060] Figure 2 A schematic diagram of pipeline execution of an existing matrix engine is shown;

[0061] Figure 3 A schematic diagram showing pipeline parallel execution of a matrix operation module according to an embodiment of the present invention is shown;

[0062] Figure 4 A schematic diagram showing mask processing according to an embodiment of the present invention is shown;

[0063] Figure 5 A schematic diagram illustrating mask processing according to another embodiment of the present invention; and

[0064] Figure 6 The flowchart of the operation method of the matrix operation system supporting mask processing according to one embodiment of the present invention is shown. DETAILED DESCRIPTION

[0065] In the following description, the present invention is described with reference to various embodiments. However, those skilled in the art will recognize that the embodiments may be implemented without one or more of the specific details or with other alternative and / or additional methods or components. In other cases, well-known structures or operations are not shown or described in detail to avoid obscuring the inventive aspects of the present invention. Similarly, specific numbers and configurations are set forth for illustrative purposes in order to provide a comprehensive understanding of the embodiments of the present invention. However, the present invention is not limited to these specific details.

[0066] In this specification, reference to "one embodiment" or "the embodiment" means that a particular feature, structure, or characteristic described in connection with the embodiment is included in at least one embodiment of the present invention. The appearances of the phrase "in one embodiment" in various places in this specification are not necessarily all referring to the same embodiment.

[0067] It should be noted that the embodiments of the present invention describe the method steps in a specific order, but this is only for the purpose of illustrating the specific embodiment and does not limit the order of the steps. On the contrary, in different embodiments of the present invention, the order of the steps can be adjusted according to actual needs.

[0068] In the present invention, each module of the system according to the present invention can be implemented using software, hardware, firmware or a combination thereof. When the module is implemented using software, the function of the module can be implemented by a computer program flow, for example, the module can be implemented by a code segment (such as a code segment in a language such as C, C++) stored in a storage device (such as a hard disk, memory, etc.), wherein when the code segment is executed by a processor, the corresponding function of the module can be implemented. When the module is implemented using hardware, the function of the module can be implemented by setting a corresponding hardware structure, for example, by hardware programming a programmable device such as a field programmable gate array (FPGA) to implement the function of the module, or by designing an application-specific integrated circuit (ASIC) including electronic devices such as a plurality of transistors, resistors and capacitors to implement the function of the module. When the module is implemented using firmware, the function of the module can be written into a read-only memory such as an EPROM or EEPROM of the device in the form of program code, and when the program code is executed by the processor, the corresponding function of the module can be implemented. In addition, certain functions of the module may need to be implemented by separate hardware or through collaboration with the hardware, for example, the detection function is implemented by corresponding sensors (such as proximity sensors, accelerometers, gyroscopes, etc.), the signal transmission function is implemented by corresponding communication devices (such as Bluetooth devices, infrared communication devices, baseband communication devices, Wi-Fi communication devices, etc.), the output function is implemented by corresponding output devices (such as displays, speakers, etc.), and so on.

[0069] The present invention provides a matrix operation system supporting mask processing. Software and hardware design is performed to implement causal masks, enabling efficient matrix operations and setting the upper triangular portion of the result to negative infinity. The matrix operation system supporting mask processing, provided by the present invention, introduces a mask processing module, enabling the system to simultaneously perform matrix operations and mask operations, reducing computational overhead and hardware resource usage.

[0070] The solution of the present invention is further described below in conjunction with the accompanying drawings of the embodiments.

[0071] Figure 1 FIG. 1 shows a schematic diagram of a matrix operation system supporting mask processing according to an embodiment of the present invention. Figure 1 As shown, a matrix operation system supporting mask processing includes a matrix operation module, a mask processing module, a configuration module, and a result output module. The matrix operation module is used to perform matrix operations on an input matrix, which may include a left matrix and a right matrix, and output the matrix operation results. The mask processing module is provided at the output end of the matrix operation module and is used to perform mask processing on the matrix operation results. The configuration module is used to configure the parameters of the mask processing module. The result output module is used to output the masked results.

[0072] The mask processing module includes a vector register, a scalar register, and a mask execution module. The vector register (mask_count) is used to store a mask count, which is a vector. The length of the mask count is the same as the number of columns of the matrix operation result. In one embodiment of the present invention, the matrix operation can be a matrix multiplication operation. The scalar register is used to store a specified value. The mask execution module is used to perform mask processing on the matrix operation result according to the mask count in the vector register and the specified value in the scalar register. mask_count[i] is the value of the i-th position of the vector register, and x is the specified value in the scalar register. When mask_count[i]≥0, the mask_count[i] rows before the i-th column of the matrix operation result are set to x; when mask_count[i]<0, the rows before the i-th column of the matrix operation result - mask_count[i]-1 are retained, and the results of the remaining rows of the matrix operation result are set to x.

[0073] In one embodiment of the present invention, the configuration module includes: a mask count configuration module, configured to configure the mask count; and a specified value configuration module, configured to configure the specified value.

[0074] Figure 2 FIG. 1 shows a schematic diagram of the pipeline execution of an existing matrix engine. Figure 2 As shown, the computation unit in a conventional matrix engine completes the calculation of a row of matrix results in one clock cycle, and then the write unit completes writing the results in one clock cycle. The computation unit and the write unit execute in parallel. If the output matrix has m rows, the calculation is completed in a total of m+1 clock cycles.

[0075] Figure 3 FIG. 1 shows a schematic diagram of pipeline parallel execution of a matrix operation module according to an embodiment of the present invention. Figure 3 As shown, matrix operations and mask processing are performed in parallel through a three-stage pipeline. The calculation module is used to generate matrix operation results row by row, the mask module is used to apply column masks to the current row data, and the write module is used to write the masked data into the memory. Figure 3 As shown, the introduction of the mask processing unit adds one level of pipeline execution to the matrix operation module. After each calculation, the generated rows are first masked according to the values ​​in the mask_count vector register and the x scalar register, and then the write module performs the write operation. The improved matrix operation module's internal calculation module, mask module, and write module are pipelined in parallel. If the output matrix has m rows, the matrix multiplication and mask calculation can be completed in a total of m+2 clock cycles. Compared to matrix operations without masks, the additional overhead introduced by masking is only one cycle.

[0076] In one embodiment of the present invention, the matrix operation module, the mask processing module and the configuration module are connected via a bus. In one embodiment of the present invention, the matrix operation system is a matrix operation system based on a CPU, a GPU or an ASIC.

[0077] Figure 4 Shown Figure 4 FIG. 1 is a schematic diagram showing mask processing according to an embodiment of the present invention. Figure 4 As shown, for causal masking, the first i rows of the i-th column (starting from 0) of the matrix multiplication result need to be set to negative infinity. Figure 4 As shown, for Q and K T The result of matrix multiplication is QK, which is of size NxN. T , the positions that need to be set to negative infinity through masking are shaded. The mask_count vector corresponding to this masking operation is Figure 4 The C language code for the overall calculation process of matrix multiplication and upper triangular masking operation is as follows:

[0078]

[0079] To set the upper triangular matrix of the matrix product c (the product of matrix a and matrix b) to negative infinity, set the first i rows of the i-th column of matrix c to negative infinity. That is, the i-th element of the mask_count vector is i. Such a vector can be generated using __riscv_vid_v_u16m8. Then, using __me_masked_matmul to call MatrixEngine for matrix multiplication and masking, the final result is obtained. This example describes the design of a matrix operation system that supports masking operations and how to use the corresponding software.

[0080] Figure 5 FIG1 shows a schematic diagram of mask processing of another embodiment of the present invention. In some cases, due to the memory layout of the input data or the layout requirements of subsequent operators, the computational overhead may be greater than computing its transposed version. At this time, for the correctness of the operation, we need to set the lower triangular part of the result to negative infinity to be equivalent to the previous causal mask. That is, the first i+1 rows of the i-th column (subscript starts from 0) of the output matrix are retained, and the rest are set to negative infinity. The specific mask part is as follows Figure 5The shaded area is shown in the middle. When the elements in mask_count[i] are designed to be negative, the first -mask_count[i]-1 rows are not masked, and the remaining rows are masked. Taking advantage of the fact that the inverse of mask_count[i] is equal to -mask_count[i]-1, the mask_count vector corresponding to the lower triangular part of the matrix can be masked to [~1,~2,~3,…,~n]. Using this feature, the following C language code can implement the overall operation of matrix multiplication and lower triangular masking:

[0081] / / causal_mask(matmul(a^T,b^T))^T=tril(matmul(b,a))

[0082] void matmul_transposed_and_causal_mask(

[0083] float*c,const float*a,const float*b,size_t m,size_t k,size_t n){

[0084] vuint16m8_t mask_count=__riscv_vid_v_u16m8(m);

[0085] mask_count=__riscv_vadd_vx_u16m8(mask_count,1,m);

[0086] mask_count=__riscv_vnot_v_u16m8(mask_count,m);

[0087] __me_masked_matmul(c,b,a,n,k,m,mask_count,-INFINITY);

[0088] }

[0089] As you can see, by configuring different values ​​of the mask_count register, the mask processing unit is not limited to the standard causal mask, but can also perform other operations such as lower triangle masking and custom masking.

[0090] The matrix operation system supporting mask processing provided by the present invention introduces a mask processing module, so that the matrix operation system can simultaneously complete matrix operations and mask operations, thereby reducing calculation overhead and occupation of hardware resources.

[0091] The existing mask processing schemes all perform a uniform mask operation on the matrix multiplication result after completing the matrix multiplication operation. For a matrix of size NxN, the additional overhead introduced by the mask operation is O(N 2 / P), where P is the degree of hardware parallelism. The matrix operation system supporting mask processing provided by the present invention introduces a novel hardware structure and pipeline design, reducing overhead to only one clock cycle. This significantly reduces the overhead introduced by mask operations in the matrix operation system. Furthermore, this overhead is independent of matrix size; the larger the matrix, the more significant the advantages of the present invention over existing techniques.

[0092] The matrix operation system supporting mask processing provided by the present invention has a flexibly configurable mask processing module that is not limited to causal masking but can also perform other operations such as lower triangle masking and custom masking, thereby improving computational efficiency and enhancing hardware flexibility and adaptability.

[0093] Figure 6 FIG. 1 is a flow chart showing an operating method of a matrix operation system supporting mask processing according to an embodiment of the present invention. Figure 6 As shown, the method includes the following steps:

[0094] Configure the parameters of the mask processing module, including mask count and specified value;

[0095] Input the input matrix to the matrix operation module;

[0096] The matrix operation module performs matrix operations and outputs matrix operation results to the mask processing module;

[0097] The mask processing module performs mask processing on the matrix operation result according to the mask count and the specified value;

[0098] Output the masked result.

[0099] In one embodiment of the present invention, a computer system is provided, comprising a processor, a graphics card having an artificial intelligence chip, and a memory, wherein the memory is configured to store machine-readable instructions, the graphics card is configured to train the method for operating the matrix operation system supporting mask processing, and the processor is configured to execute the machine-readable instructions. When the processor and / or the graphics card execute the machine-readable instructions, the following processing steps are implemented: configuring parameters of a mask processing module, including a mask count and a specified value; inputting an input matrix into the matrix operation module; the matrix operation module performing a matrix operation and outputting a matrix operation result to the mask processing module; the mask processing module performing a mask operation on the matrix operation result according to the mask count and the specified value; and outputting the masked result.

[0100] The graphics card can preferably have a GPU computing power higher than 5.0. Since the amount of data required for training is large, providing a graphics card configuration can significantly improve the training speed.

[0101] The memory includes various media that can store machine-readable instructions, such as a USB flash drive, a read-only memory (ROM), a random access memory (RAM), a mobile hard disk, a magnetic disk or an optical disk.

[0102] It can be understood that in addition to the memory and processor mentioned above, the above-mentioned computer system also includes other software and hardware components not listed in this specification. The specific components can be determined according to the model of the specific data processing equipment in different application scenarios. This specification will not list them one by one in detail.

[0103] In one embodiment of the present invention, the present invention also provides a computer-readable storage medium having machine-readable instructions stored thereon, which implement the following processing steps when executed by a processor: configuring parameters of a mask processing module, including a mask count and a specified value; inputting an input matrix into a matrix operation module; the matrix operation module performing a matrix operation and outputting a matrix operation result to the mask processing module; the mask processing module performing a mask process on the matrix operation result according to the mask count and the specified value; and outputting the masked result.

[0104] Although various embodiments of the present invention have been described above, it should be understood that they are presented as examples only and not as limitations. It will be apparent to those skilled in the relevant art that various combinations, modifications, and variations may be made thereto without departing from the spirit and scope of the present invention. Therefore, the breadth and scope of the present invention disclosed herein should not be limited by the exemplary embodiments disclosed above, but should be defined solely in terms of the technical solutions and their equivalents.

Claims

1. A matrix operation system supporting mask processing, characterized in that: include: A matrix operation module, which is used to perform matrix operations on input matrices and output matrix operation results; a mask processing module, which is provided at the output end of the matrix operation module, and is used to perform mask processing on the matrix operation result; A configuration module, used to configure parameters of the mask processing module; as well as A result output module, which is used to output the result after mask processing; The mask processing module includes: A vector register, which is used to store the mask count; scalar registers, which are used to store specified values; and A mask execution module is used to perform mask processing on the matrix operation result according to the mask count in the vector register and the specified value in the scalar register.

2. The matrix operation system supporting mask processing according to claim 1, characterized in that: The mask count is a vector, and the length of the mask count is the same as the number of columns of the matrix operation result.

3. The matrix operation system supporting mask processing according to claim 1, characterized in that: When mask_count[i]≥0, set the mask_count[i] rows before the i-th column of the matrix operation result to x; When mask_count[i]<0, the first -mask_count[i]-1 rows of the i-th column of the matrix operation result are retained, and the results of the remaining rows of the matrix operation result are set to x; Wherein, mask_count[i] is the value of the i-th position of the vector register, and x is the specified value in the scalar register.

4. The matrix operation system supporting mask processing according to claim 1, characterized in that: The configuration module includes: A mask count configuration module, configured to configure the mask count; The specified value configuration module is used to configure the specified value.

5. The matrix operation system supporting mask processing according to claim 1, characterized in that: The matrix operation and the mask processing are performed in parallel through a three-stage pipeline, including: A calculation module, which is used to generate matrix operation results row by row; a masking module for applying a column mask to the current row of data; and The write module is used to write the masked data into the memory.

6. The matrix operation system supporting mask processing according to claim 1, characterized in that: The matrix operation module, the mask processing module and the configuration module are connected via a bus.

7. The matrix operation system supporting mask processing according to claim 1, characterized in that: The matrix operation system is a matrix operation system based on CPU, GPU or ASIC.

8. A method for operating a matrix operation system supporting mask processing according to any one of claims 1 to 7, characterized in that: The following steps are involved: Configure the parameters of the mask processing module, including mask count and specified value; Input the input matrix to the matrix operation module; The matrix operation module performs matrix operations and outputs matrix operation results to the mask processing module; The mask processing module performs mask processing on the matrix operation result according to the mask count and the specified value; as well as Output the masked result.

9. An electronic device, characterized in that: include: a processor configured to execute machine-readable instructions; A graphics card having an artificial intelligence chip configured to train the method according to claim 8; as well as A memory configured to store machine-readable instructions, which, when executed by a processor and / or a graphics card, perform the steps of the method according to claim 8.

10. A computer-readable storage medium, characterized in that Machine-readable instructions are stored thereon, which, when executed by a processor, perform the steps of the method according to claim 8.

Images