Calculation method of matrix vector multiplication, model reasoning method and graphics processor

Abstract

The invention provides a matrix vector multiplication calculation method, a model reasoning method and a graphics processor, and the method comprises the steps: obtaining a plurality of target tasks; respectively distributing the target tasks to thread bundles of a graphics processor; and loading the first vector and the second vector corresponding to the target task to a thread register of each thread in a thread bundle, so that the thread bundle executes the target task. According to the scheme, the calculation task obtained by multiplying the first matrix and the first vector is decomposed into the multiple target tasks, the target tasks are distributed to the thread bundles of the graphics processor, so that the thread bundles execute the target tasks in parallel, the parallel calculation capability of the graphics processor can be fully utilized, the calculation speed of the result vector is increased, and the calculation efficiency is improved. And the calculation efficiency of matrix vector multiplication is improved.

Description

TECHNICAL FIELD

[0001] The present application relates to the technical field of artificial intelligence, in particular to a matrix-vector multiplication calculation method, a model inference method and a graphics processing unit. BACKGROUND

[0002] In the field of artificial intelligence today, the application scope of model inference technology continues to expand, covering multiple important directions such as computer vision, natural language processing, and recommendation systems. Model inference refers to using a trained model to process new input data to achieve prediction or decision-making tasks. As the model size continues to grow, its demand for computing resources is also rising sharply, and higher requirements are placed on inference efficiency. Especially when dealing with complex tasks, the model often needs to perform a large number of matrix operations, and the efficiency of these operations directly affects the overall performance and response speed of the model. Therefore, how to efficiently perform model inference tasks has become a problem of concern in the field of artificial intelligence.

[0003] Existing model inference methods usually rely on the General Matrix Multiplication (GEMM) optimization paradigm. However, in long sequence, single batch processing or small batch processing scenarios, these methods have obvious shortcomings, which are as follows: when facing long sequence, single batch processing or small batch processing scenarios, the existing scheme still starts a complex kernel designed for large-scale GEMM, resulting in low GPU core utilization and increased inference delay. SUMMARY

[0004] The purpose of the embodiments of the present application is to provide a matrix-vector multiplication calculation method, a model inference method and a graphics processing unit to solve the above technical problems.

[0005] In a first aspect, the embodiments of the present application provide a matrix-vector multiplication calculation method for calculating the product of a first matrix and a first vector to obtain a result vector, the method comprising: obtaining a plurality of target tasks; wherein the target task is used to represent the calculation task of multiplying a second vector in the first matrix with the first vector to obtain an element in the result vector; allocating the target tasks to thread bundles of a graphics processing unit respectively; wherein each thread bundle is configured to execute at least one target task to obtain at least one element in the result vector; loading the first vector and the second vector corresponding to the target task into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task.

[0006] In the implementation of the above scheme, by decomposing the computation task of multiplying the first matrix and the first vector into multiple target tasks and allocating them to the thread bundles of the graphics processor, the target tasks can be executed in parallel by each thread bundle. This fully utilizes the parallel computing capabilities of the graphics processor, thereby accelerating the computation speed of the result vector and improving the computation efficiency of matrix-vector multiplication. On the other hand, by loading the vector data corresponding to the target tasks into the thread registers of each thread in the thread bundle, the frequent access to global memory by the threads during execution is reduced, effectively reducing memory access latency and thus improving the computation efficiency of matrix-vector multiplication. Furthermore, by decomposing the computation task into multiple target tasks for parallel processing, the above matrix-vector multiplication computation method can dynamically adapt to the computational needs of matrix-vector multiplication at different scales. It can fully utilize the computing resources of the graphics processor in different scenarios, which is beneficial to improving the adaptability of the above matrix-vector multiplication computation method.

[0007] In one implementation of the first aspect, loading the first vector and the second vector corresponding to the target task into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task, includes: The third vector and the second vector corresponding to the target task are loaded into the thread registers of each thread in the thread bundle, so that the thread bundle executes the target task; wherein, the third vector is a vector in the second matrix; the second matrix is ​​a matrix after quantization of the first matrix; the execution of the target task includes: performing dequantization on the third vector.

[0008] In the implementation of the above scheme, by loading the second vector corresponding to the target task and the quantized third vector into the thread registers of each thread in the thread bundle, and performing dequantization on the third vector during execution, the parallel computing capabilities of the graphics processor can be fully utilized, thereby accelerating the calculation speed of multiplying the quantized matrix and vector, and effectively improving the computational efficiency of quantized matrix-vector multiplication. On the other hand, by directly loading the quantized third vector into the thread register so that the thread can perform dequantization, the thread's access requirements to the original matrix data in global memory are reduced, the access frequency and latency of the thread to global memory are reduced, and the computational performance of quantized matrix-vector multiplication is improved. Furthermore, by combining the processing of the quantized matrix and the dequantization operation, the above matrix-vector multiplication calculation method can flexibly adapt to the requirements of different quantization bit rates and matrix sizes, and can efficiently utilize graphics processor resources in various computing scenarios, which is conducive to improving the adaptability of the above matrix-vector multiplication.

[0009] In one implementation of the first aspect, the method further includes: The second matrix is ​​loaded into the global memory of the graphics processor in the form of quantized data blocks; wherein, the quantized data blocks include quantized data; The thread bundle is configured to read the quantized data from the global memory and perform dequantization on the quantization process.

[0010] In the implementation of the above scheme, the storage format of quantized data blocks allows the quantized data to be arranged compactly. By loading the second matrix into the global memory of the graphics processor in the form of quantized data blocks, the storage space occupation can be effectively reduced and the storage efficiency can be improved. On the other hand, the thread bundle can read the quantized data from the global memory and perform inverse quantization processing, which reduces the bandwidth requirements during data transmission and is conducive to improving data transmission efficiency. Furthermore, using quantized data blocks as the storage format allows the above matrix-vector multiplication calculation method to better adapt to the requirements of different quantization precisions, which is conducive to improving the flexibility and adaptability of the above matrix-vector multiplication calculation method.

[0011] In one implementation of the first aspect, the quantized data block further includes shared metadata; The target thread in the thread bundle is configured to read the shared metadata from the global memory and broadcast the shared metadata to other threads in the thread bundle via an intra-bundle data exchange instruction.

[0012] In the implementation of the above scheme, after the target thread in the thread bundle reads the shared metadata, it broadcasts it to other threads in the same thread bundle using the intra-bundle data exchange instruction and resides in the registers of each thread. This reduces the number of repeated accesses to global memory by threads, lowers the memory access frequency and access latency, thereby improving the computational efficiency of the matrix-vector multiplication method. On the other hand, the rapid broadcasting of shared metadata through the intra-bundle data exchange instruction significantly improves the data transmission efficiency within the thread bundle, which is beneficial for improving the efficiency of data synchronization between threads, thereby reducing the waiting time caused by thread loading data, and further improving the computational efficiency of the matrix-vector multiplication method. Furthermore, threads within the thread bundle can quickly obtain the key parameters required for performing dequantization processing, thereby further improving the computational efficiency of the matrix-vector multiplication method.

[0013] In one implementation of the first aspect, the target thread is configured as follows: The local calculation results obtained by each wire bundle within the thread bundle are loaded into the thread register of the target thread by the intra-bundle data exchange instruction, and the local calculation results are accumulated in the thread register to obtain the global accumulated result.

[0014] In the implementation of the above scheme, by using intra-bundle data exchange instructions to load the local calculation results obtained by each thread within the thread bundle into the thread register of the target thread, and performing accumulation processing in the register to obtain the global accumulation result, the number of accesses to global memory can be effectively reduced, thereby reducing the memory access latency of the thread and improving the computational efficiency of the above matrix-vector multiplication calculation method. On the other hand, performing accumulation processing in the thread register can make full use of the low latency and high bandwidth characteristics of the register, improve the efficiency of data reduction and summarization, and thus improve the overall computational performance of the above matrix-vector multiplication calculation method. Furthermore, the register-level reduction method adopted can simplify the data synchronization process, help maintain data consistency within the thread bundle, and avoid the competition and data inconsistency problems caused by multiple threads accessing the same memory location at the same time, thereby improving the accuracy of the calculation results.

[0015] In one implementation of the first aspect, each thread in the thread bundle is configured to execute at least one subtask in the target task; wherein the subtask is used to characterize a local multiplication and accumulation operation between the second vector and the first vector; The thread bundle is configured such that, for each thread in the thread bundle, when the thread executes the current subtask, the partial data of the second vector corresponding to the next subtask to be executed by the thread and the partial data of the first vector are loaded into the thread register of the thread.

[0016] In the implementation of the above scheme, by preloading the local data corresponding to the next subtask to be executed by the thread into the thread register while executing the current subtask, the idle time caused by the thread waiting for data loading during execution is effectively reduced, thereby improving the thread utilization and computational efficiency. On the other hand, by preloading the data of the next subtask, the high-speed read and write performance of the thread register can be fully utilized, reducing the latency caused by data transmission and further improving the calculation speed of matrix-vector multiplication. Furthermore, by overlapping the data loading process of the next subtask with the calculation process of the current subtask, the thread bundle can execute multiple subtasks more smoothly, reducing the overhead caused by thread idling time, improving the overall throughput of the thread bundle, and thus improving the computational performance of the above matrix-vector multiplication calculation method.

[0017] In one implementation of the first aspect, the allocation of the target tasks to the thread bundles of the graphics processor includes: When the number of target tasks is less than a preset task number threshold, a graphics processor thread bundle is allocated to each target task; wherein each thread bundle executes one target task independently. When the number of target tasks is not less than the preset task number threshold, a thread bundle of the graphics processor is allocated to at least one target task; wherein each thread bundle executes at least one target task.

[0018] In the implementation of the above scheme, when the number of target tasks is small, each task is assigned a thread bundle for fast processing; when the number of target tasks is large, multiple tasks are allowed to share a thread bundle to avoid resource waste and improve overall throughput. Flexibly adjusting the thread bundle allocation strategy based on the number of target tasks helps improve the resource utilization of the graphics processor. Furthermore, the flexible task allocation mechanism can dynamically adapt to different computational loads, ensuring that the matrix-vector multiplication calculation method maintains high computational performance even when the number of target tasks changes, thus improving the adaptability of the matrix-vector multiplication calculation method. Moreover, by reasonably allocating the number of tasks executed by the thread bundle, the competition and synchronization overhead between thread bundles can be effectively reduced, thereby improving computational stability and reliability.

[0019] Secondly, embodiments of this application provide a model reasoning method, the method comprising: Multiple target tasks are obtained; wherein, the target task is used to characterize the computation task of multiplying the second vector in the first matrix with the first vector to obtain an element in the result vector; the first matrix is ​​the weight matrix of the model to be inferred; the first vector is the input activation vector of the model to be inferred; The target tasks are assigned to the thread bundles of the graphics processor; wherein each thread bundle is configured to execute at least one of the target tasks to obtain at least one element in the result vector; The first vector and the second vector corresponding to the target task are loaded into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task and completes the model inference of the model to be inferred.

[0020] Thirdly, embodiments of this application provide a graphics processor for executing the method provided by the first aspect or any possible implementation of the first aspect or the second aspect or any possible implementation of the second aspect.

[0021] Fourthly, embodiments of this application provide an electronic device, including: a processor, a memory, and a communication bus, wherein the processor and the memory communicate with each other through the communication bus; the memory stores program instructions that can be executed by the processor, and the processor can execute the method provided by the first aspect or any possible implementation of the first aspect or the second aspect or any possible implementation of the second aspect by calling the program instructions.

[0022] Fifthly, embodiments of this application provide a computer-readable storage medium storing computer program instructions, which, when read and executed by a processor, perform the method provided by the first aspect or any possible implementation thereof, or the second aspect or any possible implementation thereof.

[0023] In a sixth aspect, embodiments of this application provide a computer program product, the computer program product including a computer program, which, when executed by a processor, implements the method provided by the first aspect or any possible implementation of the first aspect or the second aspect or any possible implementation of the second aspect.

[0024] Other features and advantages of this application will be set forth in the following description and will be apparent in part from the description or may be learned by practicing embodiments of this application. The objectives and other advantages of this application may be realized and obtained by means of the structures particularly pointed out in the written description, claims and drawings. Attached Figure Description

[0025] To more clearly illustrate the technical solutions of the embodiments of this application, the accompanying drawings used in the embodiments of this application will be briefly introduced below. It should be understood that the following drawings only show some embodiments of this application and should not be regarded as a limitation of the scope. For those skilled in the art, other related drawings can be obtained based on these drawings without creative effort.

[0026] Figure 1 A flowchart illustrating the matrix-vector multiplication calculation method provided in this application embodiment; Figure 2 A schematic diagram of the memory structure of the quantized data block provided in the embodiments of this application; Figure 3 A flowchart illustrating the model reasoning method provided in this application embodiment in a certain application scenario; Figure 4 This is a schematic diagram of the structure of an electronic device provided in an embodiment of this application. Detailed Implementation

[0027] The technical solutions of the embodiments of this application will now be described with reference to the accompanying drawings. The following embodiments are only used to more clearly illustrate the technical solutions of this application, and are therefore merely examples and should not be used to limit the scope of protection of this application.

[0028] Most existing model inference methods are based on the Large-Scale Matrix Multiplication (GEMM) optimization paradigm, which demonstrates high efficiency when handling large-scale matrix operations. However, its limitations in certain scenarios cannot be ignored. The GEMM optimization paradigm is primarily designed for parallel computation of large-scale matrices, and its optimization strategy typically assumes that the matrix has a large dimension and can fully utilize the GPU's computing resources. However, in practice, especially when dealing with long sequences, single-batch processing, or mini-batch inference tasks, the matrix size is often small, failing to fully utilize the GPU's parallel computing capabilities. This leads to partial idle computing resources and reduced inference efficiency.

[0029] In view of this, embodiments of this application provide a matrix-vector multiplication calculation method. By decomposing the calculation task of multiplying a first matrix and a first vector into multiple target tasks and allocating them to thread bundles of a graphics processor, the target tasks can be executed in parallel by each thread bundle. This fully utilizes the parallel computing capabilities of the graphics processor, thereby accelerating the calculation speed of the result vector and improving the calculation efficiency of matrix-vector multiplication. On the other hand, by loading the vector data corresponding to the target tasks into the thread registers of each thread in the thread bundle, the frequent access to global memory by the threads during execution is reduced, effectively reducing memory access latency and thus improving the calculation efficiency of matrix-vector multiplication. Furthermore, by decomposing the calculation task into multiple target tasks for parallel processing, the above matrix-vector multiplication calculation method can dynamically adapt to matrix-vector multiplication calculation needs of different scales. It can fully utilize the computing resources of the graphics processor in different scenarios, which is beneficial to improving the adaptability of the above matrix-vector multiplication calculation method.

[0030] Before introducing the calculation method for matrix-vector multiplication, let's first introduce the application scenarios of matrix-vector multiplication: The matrix-vector multiplication method described above is applicable to scenarios requiring efficient matrix-vector operations, particularly during inference in artificial intelligence models (such as deep learning models and large language models). This method effectively utilizes the parallel computing capabilities of the graphics processing unit (GPU) by decomposing the computational task into multiple target tasks and allocating them to the GPU's thread bundles, thereby accelerating the calculation of the result vector.

[0031] Furthermore, it is understandable that the above matrix-vector multiplication calculation method can also be integrated into mainstream inference frameworks (TensorRT, vLLM, etc.) as a low-level high-efficiency operator to improve the inference performance of large models in specific scenarios.

[0032] The technical solutions in the embodiments of this application will now be described with reference to the accompanying drawings.

[0033] Please see Figure 1This application provides a method for calculating matrix-vector multiplication, used to calculate the product of a first matrix and a first vector to obtain a result vector, including: Step S110: Obtain multiple target tasks; wherein, the target task is used to characterize the calculation task of multiplying the second vector in the first matrix with the first vector to obtain an element in the result vector.

[0034] General Matrix-Vector Multiplication (GEMV) is one of the fundamental operations in linear algebra, describing the product relationship between matrices and vectors. Mathematically, given an M×N matrix... (i.e., Matrix A) and a dimensional vector B (i.e., Vector) B The result of matrix-vector multiplication is a dimensional column vector (i.e., Vector) C ), result vector Each element From the matrix The Rows and Vectors The dot product is calculated as follows:

[0035] The first matrix and the first vector mentioned above refer to matrix A and vector B in the above formula, and the target task refers to the calculated result vector. a certain element The task.

[0036] The aforementioned objective task refers to the decomposition of matrix-vector multiplication into multiple independent computational units to obtain the result vector by multiplying the first matrix and the first vector. Specifically, the objective task can be set as performing a dot product operation between a second vector in the first matrix and the first vector to calculate an element in the result vector. This decomposition method breaks down large-scale matrix operations into multiple fine-grained vector dot product calculation tasks, each task being independent of the others, facilitating parallel processing.

[0037] The above scheme can obtain multiple target tasks by decomposing the matrix-vector multiplication operation. Specifically, when performing the multiplication operation between the first matrix and the first vector, the operation is decomposed into a corresponding number of independent target tasks based on the dimension of the result vector. For example, if the size of the first matrix is ​​M×N and the size of the first vector is N×1, then the size of the result vector is M×1, meaning the result vector contains M elements. Therefore, M target tasks can be decomposed, each responsible for calculating one element in the result vector.

[0038] Step S120: Assign the target tasks to the thread bundles of the graphics processor respectively; wherein each thread bundle is configured to execute at least one target task to obtain at least one element in the result vector.

[0039] The aforementioned Graphics Processing Unit (GPU), also known as a display core, display chip, or video processor, is a hardware device specifically designed to perform massively parallel computing tasks. Compared to traditional central processing units (CPUs), GPUs have a large number of simple computational cores, enabling them to handle multiple tasks simultaneously and making them suitable for graphics rendering and high-performance computing. In fields such as deep learning and scientific computing, the parallel computing capabilities of GPUs are widely used to accelerate computationally intensive tasks such as matrix operations and convolution operations.

[0040] The aforementioned warp is a form of thread organization in a GPU. In the parallel computing architecture of a GPU, a thread is the basic unit of execution, and a warp typically consists of multiple threads. For example, in NVIDIA's GPU architecture, a warp typically contains 32 threads. These warps can execute the same or different instructions simultaneously, improving computational efficiency through parallel execution. Threads within a warp share resources such as the instruction pointer and register stack, enabling efficient synchronization and communication within the warp.

[0041] The multiple target tasks obtained in step S110 are assigned to GPU thread bundles for execution. Each thread bundle is configured to execute at least one target task, accelerating the computation of the result vector by processing multiple tasks in parallel. This allocation method fully utilizes the parallel architecture of the GPU, enabling each thread bundle to independently compute one or more elements in the result vector, ultimately efficiently completing the entire matrix-vector multiplication operation.

[0042] Step S130: Load the first vector and the second vector corresponding to the target task into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task.

[0043] The aforementioned thread registers are high-speed storage units dedicated to each thread in the graphics processing unit (GPU). Located within each GPU core, they are used to temporarily store data required by the thread during execution, such as variables and intermediate calculation results. Thread registers have extremely high access speeds, typically several orders of magnitude faster than global memory, providing low-latency data read / write services to threads and thus accelerating thread execution.

[0044] Loading the data of the first and second vectors corresponding to the target task into the thread registers of each thread in the thread bundle is to fully utilize the high-speed access characteristics of registers and improve computational efficiency. When a thread executes the target task, frequent access to global memory will cause access latency because global memory access speed is relatively slow. By pre-loading the required data into the thread registers, the number of times the thread accesses global memory during execution can be reduced, thereby reducing memory access latency and improving overall computational speed. Furthermore, the data in the thread registers can be quickly reused, which is particularly important for computational tasks that require multiple accesses to the same data, further optimizing computational performance.

[0045] Optionally, step S120 above may include: loading the third vector and the second vector corresponding to the target task into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task; wherein, the third vector is a vector in the second matrix; the second matrix is ​​a matrix after quantization of the first matrix; executing the target task includes: performing inverse quantization on the third vector.

[0046] Quantization refers to the process of converting data in a matrix from a high-precision representation to a low-precision representation. In model inference scenarios, quantizing model parameters can reduce the model's storage requirements and computational complexity. By converting model parameters (such as weight matrices) from a high-precision format (such as 32-bit floating-point numbers) to a low-precision format (such as 8-bit, 4-bit, or even lower integer formats), quantization reduces the model's storage space requirements, while also reducing the computing power and memory bandwidth required during computation. In model inference applications, quantized models are suitable for resource-constrained environments, such as mobile devices and edge computing devices. These devices typically have limited computing power and storage capacity, making it impossible to efficiently run large, unquantized models. Through quantization, models can run more efficiently on these devices, resulting in faster inference speeds and lower power consumption. Commonly used low-bit quantization forms for large language models include Q2_K, Q3_K, Q4_K, Q5_K, and Q6_K, where the number after Q represents the number of bits used to represent each parameter, and K represents the number of bits in the integer part. Specifically: Q2_K indicates that each parameter is quantized using 2 bits, Q3_K uses 3 bits, Q4_K uses 4 bits, and Q5_K and Q6_K use 5 bits and 6 bits respectively. The integer part of Q4_2 is 2 bits, and its quantization step size can be set to 1 / (2^2) = 0.25.

[0047] The main advantage of these quantization methods is that they can reduce the storage requirements of the model, while also reducing the computational resources required during inference.

[0048] When a thread in a thread bundle executes a target task, the quantized third vector can be dequantized. This is because quantization reduces data precision, and when performing specific matrix-vector multiplication calculations, it is necessary to restore the data precision to ensure the accuracy of the calculation results. Dequantization is the process of converting low-precision quantized data back to a high-precision representation, typically achieved through linear transformations or other correction algorithms.

[0049] Furthermore, the dequantization process described above can be performed in the thread registers, thereby quickly restoring data precision without increasing the burden of global memory access, thus improving computational efficiency and result accuracy. Specifically: when the dequantization process does not require shared metadata, threads can operate entirely based on locally stored data. Each thread independently performs the dequantization operation, converting low-precision data into high-precision data according to predetermined dequantization rules. Each thread can independently complete the data conversion, thereby quickly restoring data precision and providing high-precision data support for subsequent calculations. When the dequantization process requires shared metadata, the threads in the thread bundle rely on the shared metadata to perform dequantization. Each thread can read the required shared metadata from global memory and then use the shared metadata to convert the quantized low-precision data into high-precision data.

[0050] The aforementioned shared metadata refers to auxiliary information associated with quantized data, which is crucial for quantization processing. Shared metadata primarily includes key parameters related to the quantization process, such as the quantization step size (Scale) and zero point (ZeroPoint). These parameters are used to restore low-precision quantized data to its original high-precision form during dequantization. Shared metadata is shared, meaning it can be used by multiple processing units (such as multiple threads within a thread bundle). In matrix-vector multiplication scenarios, shared metadata is typically stored in global memory, and each thread bundle can read its required shared metadata from global memory for dequantization operations.

[0051] The above scheme, by loading the second vector corresponding to the target task and the quantized third vector into the thread registers of each thread in the thread bundle, and performing dequantization on the third vector during execution, can fully utilize the parallel computing capabilities of the graphics processor, thereby accelerating the calculation speed of multiplying the quantized matrix and vector, and effectively improving the computational efficiency of quantized matrix-vector multiplication. On the other hand, by directly loading the quantized third vector into the thread register so that the thread can perform the dequantization operation, the thread's access requirements to the original matrix data in global memory are reduced, lowering the thread's access frequency and latency to global memory, thus improving the computational performance of quantized matrix-vector multiplication. Furthermore, by combining the processing of the quantized matrix and the dequantization operation, the above matrix-vector multiplication calculation method can flexibly adapt to the requirements of different quantization bit rates and matrix sizes, and can efficiently utilize graphics processor resources in various computing scenarios, which is conducive to improving the adaptability of the above matrix-vector multiplication.

[0052] Optionally, the above matrix-vector multiplication calculation method may further include: loading the second matrix into the global memory of the graphics processor in the form of a quantized data block; wherein the quantized data block includes quantized data; and the thread bundle is configured to: read the quantized data from the global memory and perform inverse quantization processing on the quantization process.

[0053] The aforementioned quantized data block refers to a unit that organizes and stores quantized matrix data using a specific data structure. Each quantized data block can include quantized data and related metadata, which typically includes information such as quantization step size and zero points, used for subsequent inverse quantization processing. The size and structure of the quantized data block can be adjusted according to specific computational requirements and hardware characteristics. Taking a scenario where a large language model is quantized using Q4_K quantization as an example, the memory format of the quantized data block is as follows: Figure 2As shown, quants is used to store quantization values, scales is used to store scaling factors, and dm is used to store bias or other auxiliary parameters. These auxiliary parameters can be used to further refine the dequantization process to improve the accuracy of the dequantized data.

[0054] The aforementioned method of storing the second matrix in the global memory of the graphics processor (GPU) in the form of quantized data blocks primarily aims to improve storage efficiency and data access convenience. Organizing the quantized data into a block structure makes data storage in global memory more compact, reducing wasted storage space. Furthermore, quantized data blocks help optimize the data loading and transmission process. When performing matrix-vector multiplication, a thread bundle can read quantized data blocks in batches from global memory and then perform dequantization processing in its local thread register. This approach reduces the number of global memory accesses, lowers memory access latency, and improves overall computational efficiency. In addition, the use of quantized data blocks simplifies data management and task allocation. During matrix operations, tasks can be assigned to different thread bundles in blocks, with each thread within a bundle responsible for processing one or more quantized data blocks. This block-based processing approach helps achieve load balancing of computational tasks and facilitates the utilization of the GPU's massively parallel computing capabilities.

[0055] The storage format of the quantized data blocks in the above scheme allows the quantized data to be arranged compactly. By loading the second matrix into the global memory of the graphics processor in the form of quantized data blocks, the storage space occupied can be effectively reduced and the storage efficiency can be improved. On the other hand, the thread beam can read the quantized data from the global memory and perform inverse quantization processing, which reduces the bandwidth requirements during data transmission and is conducive to improving data transmission efficiency. Furthermore, using quantized data blocks as the storage format allows the above matrix-vector multiplication calculation method to better adapt to the requirements of different quantization precisions, which is conducive to improving the flexibility and adaptability of the above matrix-vector multiplication calculation method.

[0056] Optionally, the quantized data block also includes shared metadata; the target thread in the thread bundle is configured to read the shared metadata in global memory and broadcast the shared metadata to other threads in the thread bundle via intra-bundle data exchange instructions.

[0057] The target thread mentioned above refers to a thread specifically designated to perform a particular task. In the above scheme, the target thread is mainly responsible for reading and broadcasting shared metadata. Within a thread bundle, a thread is typically selected as the target thread. This target thread is responsible for reading shared metadata from global memory and broadcasting this data to other threads in the bundle via intra-bundle data exchange instructions. This role allocation ensures the efficiency of data reading and broadcasting and reduces redundant global memory accesses.

[0058] Shared metadata includes key parameters such as quantization step size and zero point, which are common when dequantizing the quantized data handled by the thread bundle. Therefore, a thread (the target thread) can be configured to read this data from global memory and then broadcast it to other threads in the bundle via intra-bundle data exchange instructions. This broadcasting method not only improves data transmission efficiency but also ensures data consistency and accuracy.

[0059] The aforementioned intra-beam data exchange instructions are a low-latency communication mechanism provided by the graphics processing unit (GPU), allowing threads within a beam to directly exchange data. Intra-beam data exchange instructions typically feature high execution efficiency, enabling rapid data transfer within the beam. Common intra-beam data exchange instructions include `__shfl_sync` and `__shfl_down_sync`. In NVIDIA's GPU architecture, these instructions allow threads to share data with extremely low latency, achieving efficient broadcasting and data synchronization. Through these instructions, the target thread can quickly broadcast the shared metadata it has read to other threads within the beam, ensuring that each thread can promptly obtain the necessary parameters for dequantization.

[0060] The above scheme, by having the target thread in the thread bundle read the shared metadata and then broadcast it to other threads in the same thread bundle using intra-bundle data exchange instructions, and having it reside in each thread's register, reduces the number of repeated accesses to global memory by threads, lowers the memory access frequency and latency, thereby improving the computational efficiency of the matrix-vector multiplication method. Furthermore, the rapid broadcasting of shared metadata through intra-bundle data exchange instructions significantly improves the data transmission efficiency within the thread bundle, which is beneficial for improving the efficiency of data synchronization between threads, thereby reducing the waiting time caused by data loading and further improving the computational efficiency of the matrix-vector multiplication method. Additionally, threads within the thread bundle can quickly obtain the key parameters required for dequantization processing, further improving the computational efficiency of the matrix-vector multiplication method.

[0061] Optionally, the target thread can also be configured to: load the local calculation results obtained by each wire in the thread bundle into the thread register of the target thread through the intra-bundle data exchange instruction, and perform accumulation processing on the local calculation results in the thread register to obtain the global accumulation result.

[0062] In the above scheme, the target thread can also be responsible for the global reduction operation within the thread bundle to obtain the global accumulation result. Specifically, after each thread within the thread bundle completes its local calculation, the target thread can use the data exchange instruction within the bundle to collect the local results obtained by each thread into its own thread register. Within the target thread's thread register, the target thread performs an accumulation operation on these local results, ultimately obtaining the global accumulation result. This process efficiently utilizes the fast data processing capability of the thread register, reduces the need for global memory access, and the low-latency characteristics of the data exchange instruction within the bundle ensure the high efficiency of data collection and accumulation.

[0063] Understandably, traditional methods often rely on shared memory in global memory for reduction operations. This not only leads to frequent global memory accesses but also introduces synchronization complexity and potential data conflicts. The proposed solution, however, reduces the number of times threads access global memory by performing global reduction within thread registers, effectively reducing memory access latency and improving overall computational efficiency. Furthermore, operations within thread registers offer low latency and high bandwidth, making the accumulation process faster. In addition, this reduction method simplifies the synchronization process, leveraging the efficient communication capabilities of intra-binary data exchange instructions to ensure low latency and high efficiency in data transmission between threads, thereby enhancing the stability and reliability of the computation process.

[0064] The above scheme utilizes intra-thread data exchange instructions to load the local computation results obtained by each thread within the thread bundle into the target thread's thread register, and performs accumulation processing within the register to obtain the global accumulation result. This effectively reduces the number of accesses to global memory, thereby reducing thread memory access latency and improving the computational efficiency of the matrix-vector multiplication method. On the other hand, performing accumulation processing within the thread register fully utilizes the low latency and high bandwidth characteristics of the register, improving the efficiency of data reduction and summarization, thus improving the overall computational performance of the matrix-vector multiplication method. Furthermore, the register-level reduction method simplifies the data synchronization process, helps maintain data consistency within the thread bundle, and avoids contention and data inconsistency problems caused by multiple threads accessing the same memory location simultaneously, thereby improving the accuracy of the computation results.

[0065] Optionally, each thread in the above thread bundle can be configured to execute at least one subtask in the target task; wherein the subtask is used to characterize the local multiplication and accumulation operation between the second vector and the first vector.

[0066] The thread bundle is configured such that, for each thread in the thread bundle, when the thread is executing the current subtask, the local data of the second vector and the local data of the first vector corresponding to the next subtask to be executed by the thread are loaded into the thread's thread register.

[0067] The subtasks of the aforementioned objective task refer to finer-grained computational units formed by further decomposing the local multiplication and summation operation between the second and first vectors. Each subtask is responsible for calculating the multiplication and summation result within a local range, that is, performing element-wise multiplication and summation on the corresponding local data of the second and first vectors to obtain the local multiplication and summation. These subtasks are the basic computational units constituting the entire objective task. Through the collaborative completion of multiple subtasks, the computational objective of the entire objective task is ultimately achieved.

[0068] Each thread in the aforementioned thread bundle is configured to preload the local data of the second vector and the local data of the first vector required for the next subtask into its thread register while executing the current subtask. Once a thread completes the computation of its current subtask, the data required for the next subtask is already in its thread register, allowing the thread to immediately begin executing the next subtask, thus reducing the time required for loading data from global memory into its thread register.

[0069] Furthermore, threads within a thread bundle can perform implicit or explicit synchronization after completing a subtask. Implicit or explicit synchronization is a concept in parallel computing used to ensure that threads within a thread bundle execute in a coordinated manner at specific points, guaranteeing data consistency and computational correctness. Implicit synchronization is automatically implemented through features of programming languages ​​or frameworks, eliminating the need for programmers to manually insert synchronization instructions. For example, when using intra-bundle data exchange instructions (such as NVIDIA GPU's __shfl_sync), the hardware automatically ensures that related threads execute synchronously, thus implicitly guaranteeing data consistency. Explicit synchronization, on the other hand, is implemented through specific synchronization instructions. When a thread bundle is executing a computational task, explicit synchronization instructions force all threads to pause at that point, waiting for all threads to reach that point before continuing execution. This ensures that data exchange and computational steps between threads proceed in the expected order, preventing data races and inconsistencies.

[0070] Furthermore, a pipelined architecture can be implemented through careful design. For example, a multi-buffer mechanism can be used, where while computation uses data in one buffer, other buffers preload data for the next subtask in parallel, ensuring seamless integration of computation and data loading operations. Additionally, instruction-level parallelism optimization can be used to allow threads to preload data for the next subtask using idle execution units while executing the current computation. Through these designs, threads prepare for the next subtask while executing the current one, thus achieving a highly efficient pipelined architecture. This architecture not only reduces data loading latency but also improves the utilization of computational resources, making matrix-vector multiplication more efficient.

[0071] The above scheme effectively reduces the idle time of the thread during execution due to waiting for data loading by preloading the local data corresponding to the next subtask to be executed into the thread register while executing the current subtask, thereby improving the thread utilization and computational efficiency. On the other hand, by preloading the data of the next subtask, the high-speed read and write performance of the thread register can be fully utilized, reducing the latency caused by data transmission and further improving the calculation speed of matrix-vector multiplication. Furthermore, by overlapping the data loading process of the next subtask with the calculation process of the current subtask, the thread bundle can execute multiple subtasks more smoothly, reducing the overhead caused by thread idling time and improving the overall throughput of the thread bundle, thereby improving the computational performance of the above matrix-vector multiplication calculation method.

[0072] Optionally, the above step S120, which assigns the target tasks to the thread bundles of the graphics processor, may include: when the number of target tasks is less than a preset task number threshold, assigning a thread bundle of the graphics processor to each target task; wherein each thread bundle executes one target task independently; when the number of target tasks is not less than the preset task number threshold, assigning a thread bundle of the graphics processor to at least one target task; wherein each thread bundle executes at least one target task.

[0073] In the above scheme, the mechanism for flexibly allocating target tasks based on their quantity dynamically adjusts the thread bundle allocation strategy to optimize the resource utilization and computational efficiency of the graphics processor. When the number of target tasks is small, each target task is assigned to an independent thread bundle, allowing each thread bundle to focus on executing a single task and thus complete the computation quickly. This one-to-one task allocation method reduces competition between thread bundles, ensuring that each task can efficiently utilize the computational resources of the thread bundle. When the number of target tasks reaches or exceeds a preset task quantity threshold, the allocation strategy is adjusted, assigning multiple target tasks to the same thread bundle. This adjustment aims to fully utilize the parallel processing capabilities of the thread bundle, improving overall throughput by having one thread bundle handle multiple tasks.

[0074] The preset task quantity threshold can be determined based on the hardware characteristics of the graphics processor (such as the number of thread bundles and their parallel processing capabilities) and the computational complexity of the target task. When the number of target tasks is large, a multi-task allocation strategy can reduce the overhead of thread bundle creation and destruction, improve resource utilization, and avoid performance degradation caused by too many thread bundles competing for resources. Through the above task allocation method, resource allocation can be dynamically adjusted under different task loads to ensure efficient computing performance.

[0075] In the above scheme, when the number of target tasks is small, each task is assigned a thread bundle for fast processing; while when the number of target tasks is large, multiple tasks are allowed to share a thread bundle, avoiding resource waste and improving overall throughput. Flexibly adjusting the thread bundle allocation strategy based on the number of target tasks helps improve the resource utilization of the graphics processor. Furthermore, the flexible task allocation mechanism can dynamically adapt to different computational loads, ensuring that the matrix-vector multiplication calculation method maintains high computational performance even when the number of target tasks changes, thus improving its adaptability. Moreover, by reasonably allocating the number of tasks executed by each thread bundle, competition and synchronization overhead between thread bundles can be effectively reduced, thereby improving computational stability and reliability.

[0076] Based on the same inventive concept, this application also provides a model inference method, including: acquiring multiple target tasks; wherein, the target task is used to characterize the computation task of multiplying a second vector in a first matrix with a first vector to obtain an element in the result vector; the first matrix is ​​the weight matrix of the model to be inferred; the first vector is the input activation vector of the model to be inferred; assigning the target tasks to thread bundles of a graphics processor respectively; wherein, each thread bundle is configured to execute at least one target task to obtain at least one element in the result vector; loading the first vector and the second vector corresponding to the target task into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task and completes the model inference of the model to be inferred.

[0077] It is understood that the above model reasoning method can realize any function of the matrix-vector multiplication calculation method provided in the embodiments of this application. For the sake of simplicity, the way the model reasoning method realizes each function will not be described in detail here. For specific implementation methods, please refer to the matrix-vector multiplication calculation method provided in the embodiments of this application.

[0078] To facilitate understanding of the working principle of the above-described model inference method, this application also provides a specific application scenario. This application scenario mainly utilizes the above-described model inference method to perform model inference on a quantized large language model, with the quantization method chosen being Q4_K. Please refer to... Figure 3 The relevant aspects of the above model inference method in this application scenario include: Task mapping section: Quantizes the low-bit weight matrix to be calculated (i.e., Figure 3 In the matrix A) and the input activation vector (i.e. Figure 3 The multiplication operation of Vector B in the GPU is performed by a Warp thread bundle. Each Warp is designed to: collaboratively multiply a row vector in the low-bit quantization weight matrix with the input activation vector to obtain the result vector C (i.e., ...). Figure 3An element in a Vector (C).

[0079] In addition, it should be noted that, Figure 3 Both Grid and Block are concepts in GPU architecture, specifically: In GPU computing architecture, Grid and Block are key concepts for organizing and managing parallel computing tasks. A Grid is a two-dimensional or three-dimensional array of Blocks, each containing several Warp threads. This hierarchical structure enables GPUs to efficiently handle large-scale parallel computing tasks. The Grid is the top-level logical unit of the entire computing task, used to decompose large-scale computing tasks into multiple relatively independent subtasks. Each subtask is executed by a Block. A Block is a computing unit within a Grid, containing a certain number of Warp threads, each capable of independently executing the same instructions. The design of Blocks allows tasks to be further subdivided and parallelized, thereby improving computational efficiency. Each Block consists of several Warp threads, each containing a certain number of threads. In NVIDIA's GPU architecture, a Warp typically consists of 32 threads. These Warp threads can execute the same instructions simultaneously, achieving SIMD (Single Instruction Multiple Data) parallelism. When performing matrix-vector multiplication, each Block (i.e., Warp thread) can be responsible for calculating one or more elements of the result vector, accelerating the entire computation process by processing multiple Warp threads in parallel.

[0080] In this application scenario, the Grid consists of multiple Blocks, and each Block corresponds to a Warp thread bundle. Each Warp thread bundle executes its computational tasks independently, and the Warp threads can run in parallel.

[0081] Phased data operations by threads: Taking the example of each thread bundle executing one target task to obtain one element from the result vector, multiple threads within a thread bundle collaborate to compute and obtain one element from the result vector. The quantized model weight matrix is ​​shown as follows: Figure 2 The quantized data blocks shown are stored in the GPU's global memory. Each thread is responsible for processing the dequantization operation of one or more quantized data points and the multiplication operation between the quantized data and the corresponding data in the input activation vector. In this application scenario, the processing of a single thread's dequantization operation of a quantized data point and the multiplication operation between that quantized data and the corresponding data in the input activation vector is called a Stage, which can also be understood as a loop (or a subtask in the target task). Threads can process multiple quantized data points through multiple stages, and threads can perform the following overlap operation: Calculation of stage i: Using the weights of the current stage (from Matrix A) and the input activation vector (from Vector B), perform dequantization and multiply-accumulate calculations to generate a partial cumulative sum.

[0082] Data loading for stage i+1: While stage i is being calculated, the next stage, i.e., stage i+1, is started asynchronously and loaded from the GPU global memory into the thread register.

[0083] The data loading in stage i+1 overlaps with the computation in stage i, and the memory access delay caused by the data loading is hidden by the subsequent computation process.

[0084] The dequantized data processing in thread registers: The dequantization process of the low-bit quantization weight matrix and the multiplication and addition operations of the dequantized data and the input vector can be fused at the register level, thus forming an efficient computational pipeline.

[0085] Shared metadata is acquired using Warp Shuffle instructions: For shared metadata required during dequantization (such as the scaling factor for Q4_K blocks), a predefined target thread within the Warp can load it from global memory and then broadcast it to the registers of all other threads within the Warp using Warp Shuffle instructions (such as __shfl_sync, __shfl_up_sync, etc.). Warp Shuffle is an efficient data exchange mechanism in GPU architecture that allows threads within a Warp bundle to directly exchange data. This mechanism enables threads within a Warp bundle to quickly share and transfer data without going through global or shared memory, thereby reducing data transfer latency.

[0086] Efficient reduction based on Warp Shuffle: Multiple threads within a Warp compute their own local sums in their respective thread registers. The target thread in the thread bundle can use Warp Shuffle instructions (such as __shfl_xor_sync) to perform efficient reduction operations at the register level, quickly accumulating all partial sums to generate the final output of the GEMV operation.

[0087] The implementation steps of the above model inference method in this application scenario mainly include: Environment preparation: Use MR100, SDK 4.1.0 or higher development environment.

[0088] Kernel definition: Define a CUDA kernel function __global__voidwarp_gemv_q4k_kernel(...).

[0089] Warp mapping: Assuming the output vector has M elements, start M warps (i.e., M blocks), each warp is responsible for computing one output element. .

[0090] Initialization: Thread 0 within Warp initializes the final sum final_sum=0.0f.

[0091] The thread bundle performs pipelined loops, that is, iterates through each Q4_K block of the weight matrix to achieve model inference. The steps of the thread bundle to implement pipelined loops include: Step 1: Data and computation overlap: Warp's 64 threads collaborate to asynchronously load the next stage's quantized weight data block and corresponding vector B element fragments from global memory into registers. This loading operation overlaps with the computation in Step 2.

[0092] Step 2: Thread computation: Calculations are performed using the quantized weights and fragments of input activation vector data loaded into the thread register in the previous stage (or initialization process). Rapid Broadcast: Thread 0 reads the scaling factor 'scale' from global memory, executes reg_scale=__shfl_sync(0xFFFFFFFF,scale,0), and broadcasts the scaling factor to other threads in the thread bundle.

[0093] Computational fusion: Each thread i is based on the quantization weights stored in the thread register. and Calculate the local results:

[0094] Efficient reduction: All threads within Warp use the __shfl_xor_sync instruction to perform reduction operations, accumulating partial_sum into reg_sum.

[0095] Accumulate to the final result: Thread 0 accumulates reg_sum to final_sum.

[0096] Step 3: Perform data synchronization: Wait for the memory loading in Step 1 to complete (implicit or explicit synchronization).

[0097] Loop: After step three is completed, jump back to step one to process the next group.

[0098] Write-back result: After the loop ends, thread 0 writes final_sum back to the output buffer in global memory.

[0099] Each Warp completes the calculation of one element in the result, and multiple Blocks work together to complete the calculation of the entire result, ultimately completing the inference process for the model.

[0100] It is understandable that the above model inference method can be applied to various types of models, including but not limited to large language models and deep learning models, among which: Large language models, such as generative pre-trained models (GPT) and Transformer models, typically have a large parameter scale and are used to handle complex natural language tasks, such as text generation, machine translation, and question answering systems. Model inference methods accelerate the execution of language processing tasks and improve response speed and processing efficiency by efficiently handling the computational requirements of these models.

[0101] Deep learning models encompass convolutional neural networks (CNNs), recurrent neural networks (RNNs), and their variants, and are widely used in computer vision, speech recognition, and other tasks requiring complex pattern recognition. Model inference methods can optimize the computational flow of deep learning models, reduce inference time, and improve the usability and performance of models in practical applications.

[0102] Figure 4 This is a schematic diagram of an electronic device provided in an embodiment of this application. (Refer to...) Figure 4 The electronic device 200 includes a graphics processor 210, a memory 220, and a communication interface 230. These components are interconnected and communicate with each other via a communication bus 240 and / or other forms of connection mechanism (not shown).

[0103] The memory 220 includes one or more (only one is shown in the figure), which may be, but is not limited to, Random Access Memory (RAM), Read Only Memory (ROM), Programmable Read-Only Memory (PROM), Erasable Programmable Read-Only Memory (EPROM), Electrically Erasable Programmable Read-Only Memory (EEPROM), etc. The graphics processor 210 and other possible components can access the memory 220 to read and / or write data therein. The graphics processor 210 includes one or more (only one is shown in the figure), and the communication interface 230 includes one or more (only one is shown in the figure), which can be used to communicate directly or indirectly with other devices to exchange data. For example, the communication interface 230 may be an Ethernet interface; a mobile communication network interface, such as an interface for 3G, 4G, or 5G networks; or other types of interfaces with data transmission and reception capabilities.

[0104] One or more computer program instructions may be stored in memory 220. Graphics processor 210 may read and run these computer program instructions to implement the matrix-vector multiplication calculation method, model reasoning method and other desired functions provided in the embodiments of this application.

[0105] Understandable. Figure 4 The structure shown is for illustrative purposes only; the electronic device 200 may also include more than [other components]. Figure 4 The more or fewer components shown, or having the same Figure 4 The different configurations shown. Figure 4 The components shown can be implemented using hardware, software, or a combination thereof. For example, electronic device 200 can be a single server (or other device with computing power), a combination of multiple servers, a cluster of a large number of servers, etc., and can be either a physical device or a virtual device.

[0106] This application also provides a computer-readable storage medium storing computer program instructions. These instructions are read and executed by a processor to perform the matrix-vector multiplication calculation method and model reasoning method provided in this application. For example, the computer-readable storage medium can be implemented as follows: Figure 4The memory 220 in the electronic device 200, or a separate storage product (such as a USB flash drive, portable hard drive, etc.).

[0107] This application also provides a computer program product, which includes computer program instructions. These computer program instructions are read and executed by a processor to perform the matrix-vector multiplication calculation method and model reasoning method provided in this application. For example, these computer program instructions can be stored in... Figure 4 The memory 220 in the electronic device 200 is located inside the memory, or it is stored in a separate storage product (such as a USB flash drive, portable hard drive, etc.).

[0108] In this document, relational terms such as first and second are used only to distinguish one entity or operation from another entity or operation, without necessarily requiring or implying any such actual relationship or order between these entities or operations.

[0109] The above description is merely an embodiment of this application and is not intended to limit the scope of protection of this application. Various modifications and variations can be made to this application by those skilled in the art. Any modifications, equivalent substitutions, improvements, etc., made within the spirit and principles of this application should be included within the scope of protection of this application.

Claims

1. A method for calculating matrix-vector multiplication, characterized in that, The method for calculating the product of a first matrix and a first vector to obtain a result vector includes: Multiple target tasks are obtained; wherein, the target task is used to characterize the computation task of multiplying the second vector in the first matrix with the first vector to obtain an element in the result vector; The target tasks are respectively assigned to the thread bundles of the graphics processor; wherein each thread bundle is configured to execute at least one of the target tasks to obtain at least one element in the result vector; The first vector and the second vector corresponding to the target task are loaded into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task.

2. The method for calculating matrix-vector multiplication according to claim 1, characterized in that, The step of loading the first vector and the second vector corresponding to the target task into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task, includes: The third vector and the second vector corresponding to the target task are loaded into the thread registers of each thread in the thread bundle, so that the thread bundle executes the target task; wherein, the third vector is a vector in the second matrix; the second matrix is ​​a matrix after quantization of the first matrix; the execution of the target task includes: performing dequantization on the third vector.

3. The method for calculating matrix-vector multiplication according to claim 2, characterized in that, The method further includes: The second matrix is ​​loaded into the global memory of the graphics processor in the form of quantized data blocks; wherein, the quantized data blocks include quantized data; The thread bundle is configured to read the quantized data from the global memory and perform dequantization on the quantization process.

4. The method for calculating matrix-vector multiplication according to claim 3, characterized in that, The quantized data block also includes shared metadata; The target thread in the thread bundle is configured to read the shared metadata from the global memory and broadcast the shared metadata to other threads in the thread bundle via an intra-bundle data exchange instruction.

5. The method for calculating matrix-vector multiplication according to claim 4, characterized in that, The target thread is configured as follows: The local calculation results obtained by each wire bundle within the thread bundle are loaded into the thread register of the target thread by the intra-bundle data exchange instruction, and the local calculation results are accumulated in the thread register to obtain the global accumulated result.

6. The method for calculating matrix-vector multiplication according to any one of claims 1 to 5, characterized in that, Each thread in the thread bundle is configured to execute at least one subtask in the target task; wherein the subtask is used to characterize a local multiplication and accumulation operation between the second vector and the first vector; The thread bundle is configured such that, for each thread in the thread bundle, when the thread executes the current subtask, the partial data of the second vector corresponding to the next subtask to be executed by the thread and the partial data of the first vector are loaded into the thread register of the thread.

7. The method for calculating matrix-vector multiplication according to any one of claims 1 to 5, characterized in that, The process of assigning the target tasks to the thread bundles of the graphics processor includes: When the number of target tasks is less than a preset task number threshold, a graphics processor thread bundle is allocated to each target task; wherein each thread bundle executes one target task independently. When the number of target tasks is not less than the preset task number threshold, a thread bundle of the graphics processor is allocated to at least one target task; wherein each thread bundle executes at least one target task.

8. A model reasoning method, characterized in that, The method includes: Multiple target tasks are obtained; wherein, the target task is used to characterize the computation task of multiplying the second vector in the first matrix with the first vector to obtain an element in the result vector; the first matrix is ​​the weight matrix of the model to be inferred; the first vector is the input activation vector of the model to be inferred; The target tasks are respectively assigned to the thread bundles of the graphics processor; wherein each thread bundle is configured to execute at least one of the target tasks to obtain at least one element in the result vector; The first vector and the second vector corresponding to the target task are loaded into the thread register of each thread in the thread bundle, so that the thread bundle executes the target task and completes the model inference of the model to be inferred.

9. A graphics processor, characterized in that, Used to perform the method as described in any one of claims 1 to 8.

10. An electronic device, characterized in that, include: A graphics processor, a memory, and a communication bus, wherein the graphics processor and the memory communicate with each other via the communication bus; The memory stores program instructions that can be executed by the graphics processor, and the graphics processor can execute the method as described in any one of claims 1 to 9 by calling the program instructions.

11. A computer-readable storage medium, characterized in that, The computer-readable storage medium stores computer instructions that, when executed by a computer, cause the computer to perform the method as described in any one of claims 1 to 9.

12. A computer program product, characterized in that, The computer program product includes a computer program that, when executed by a processor, implements the method as described in any one of claims 1 to 9.

Images