Data processing method and device based on GPU, electronic equipment and medium

Abstract

The invention provides a GPU-based data processing method and device, electronic equipment and a medium, and relates to the technical field of chips, in particular to the technical field of GPU memories. According to the implementation scheme, a first matrix to be processed and a second matrix to be processed are obtained from a global memory of a GPU; storing the first matrix into a first sub-memory of a shared memory corresponding to a target thread block in a GPC (Graphics Processing Cluster); storing the second matrix into a second sub-memory of the shared memory corresponding to the target thread block; performing multiplication operation on the first matrix and the second matrix in the shared memory; the third matrix obtained through multiplication operation is stored in a shared GPC memory of a plurality of thread blocks in a GPC of the GPU, the plurality of thread blocks comprise target thread blocks, and the shared GPC memory is addressed uniformly for the plurality of thread blocks; and performing reduction operation on the third matrix in the shared GPC memory.

Description

Technical Field

[0001] The present disclosure relates to the field of chip technology, in particular to the field of GPU memory technology, and specifically to a GPU-based data processing method, device, electronic device, computer-readable storage medium, and computer program product. Background Art

[0002] With the development of large models and neural network models, the bottleneck of data access in GPUs has become increasingly prominent. Improving the efficiency of data reuse and synchronization will effectively increase the calculation speed of the model. Therefore, how to improve the efficiency of data reuse and synchronization in GPUs has become an urgent problem to be solved.

[0003] The methods described in this section are not necessarily methods that have been previously conceived or employed. Unless otherwise indicated, it should not be assumed that any method described in this section is considered to be prior art simply because it is included in this section. Similarly, unless otherwise indicated, the issues mentioned in this section should not be considered to have been recognized in any prior art. Summary of the invention

[0004] The present disclosure provides a GPU-based data processing method, device, electronic device, computer-readable storage medium, and computer program product.

[0005] According to one aspect of the present disclosure, a GPU-based data processing method is provided, the method comprising: obtaining a first matrix and a second matrix to be calculated from a global memory of a GPU; storing the first matrix in a first sub-memory of a shared memory corresponding to a target thread block within a graphics processing cluster GPC; storing the second matrix in a second sub-memory of a shared memory corresponding to the target thread block; performing a multiplication operation on the first matrix and the second matrix in the shared memory; storing a third matrix obtained by the multiplication operation in a shared GPC memory of multiple thread blocks within a graphics processing cluster GPC of the GPU, wherein the multiple thread blocks include the target thread block, and wherein the shared GPC memory is uniformly addressed for the multiple thread blocks; and performing a reduction operation on the third matrix in the shared GPC memory.

[0006] According to another aspect of the present disclosure, a GPU-based data processing device is provided, comprising: an acquisition module configured to acquire a first matrix and a second matrix to be calculated from a global memory of a GPU; a first storage module configured to store the first matrix in a first sub-memory of a shared memory corresponding to a target thread block within a graphics processing cluster (GPC); a second storage module configured to store the second matrix in a second sub-memory of a shared memory corresponding to the target thread block; a first computing module configured to perform a multiplication operation on the first matrix and the second matrix in the shared memory; a third storage module configured to store a third matrix obtained by the multiplication operation in a shared GPC memory of multiple thread blocks within a graphics processing cluster (GPC) of the GPU, wherein the multiple thread blocks include the target thread block, and wherein the shared GPC memory is uniformly addressed for the multiple thread blocks; and a second computing module configured to perform a reduction operation on the third matrix in the shared GPC memory.

[0007] According to another aspect of the present disclosure, an electronic device is provided, comprising: at least one processor; and a memory communicatively connected to the at least one processor; wherein the memory stores instructions executable by the at least one processor, and the instructions are executed by the at least one processor so that the at least one processor can execute the above method.

[0008] According to another aspect of the present disclosure, a non-transitory computer-readable storage medium storing computer instructions is provided, wherein the computer instructions are used to enable the computer to execute the above method.

[0009] According to another aspect of the present disclosure, a computer program product is provided, comprising a computer program, wherein the computer program implements the above method when executed by a processor.

[0010] According to one or more embodiments of the present disclosure, a GPU-based data processing method is provided, which implements data sharing across thread blocks by introducing a uniformly addressed shared GPC memory inside a graphics processing cluster (GPC), thereby enabling direct reduction operations on the obtained product matrix after matrix multiplication, thereby reducing dependence on global memory and speeding up the processing of matrix data.

[0011] It should be understood that the content described in this section is not intended to identify the key or important features of the embodiments of the present disclosure, nor is it intended to limit the scope of the present disclosure. Other features of the present disclosure will become easily understood through the following description. BRIEF DESCRIPTION OF THE DRAWINGS

[0012] The accompanying drawings exemplarily illustrate the embodiments and constitute a part of the specification, and together with the text description of the specification, are used to explain the exemplary implementation of the embodiments. The embodiments shown are for illustrative purposes only and do not limit the scope of the claims. In all drawings, the same reference numerals refer to similar but not necessarily identical elements.

[0013] Figure 1 is a schematic diagram illustrating an example system in which the various methods described herein may be implemented according to an exemplary embodiment.

[0014] Figure 2 A flowchart of a GPU-based data processing method according to an embodiment of the present disclosure is shown;

[0015] Figure 3 A schematic diagram showing a hardware structure of a shared GPC memory according to an embodiment of the present disclosure;

[0016] Figure 4 A structural block diagram of a GPU-based data processing device according to an embodiment of the present disclosure is shown; and

[0017] Figure 5 A structural block diagram of an exemplary electronic device that can be used to implement the embodiments of the present disclosure is shown. DETAILED DESCRIPTION

[0018] The following is a description of exemplary embodiments of the present disclosure in conjunction with the accompanying drawings, including various details of the embodiments of the present disclosure to facilitate understanding, which should be considered as merely exemplary. Therefore, it should be recognized by those of ordinary skill in the art that various changes and modifications may be made to the embodiments described herein without departing from the scope of the present disclosure. Similarly, for the sake of clarity and conciseness, the description of well-known functions and structures is omitted in the following description.

[0019] In the present disclosure, unless otherwise specified, the use of the terms "first", "second", etc. to describe various elements is not intended to limit the positional relationship, temporal relationship, or importance relationship of these elements, and such terms are only used to distinguish one element from another element. In some examples, the first element and the second element may refer to the same instance of the element, and in some cases, based on the description of the context, they may also refer to different instances.

[0020] The terms used in the description of various examples in this disclosure are only for the purpose of describing specific examples and are not intended to be limiting. Unless the context clearly indicates otherwise, if the number of elements is not specifically limited, the element can be one or more. In addition, the term "and / or" used in this disclosure covers any one of the listed items and all possible combinations.

[0021] In the related art, data is mainly stored in shared memory (SMEM) and shared among multiple threads in the same thread block. However, this method limits data communication to calculations within the same thread block. For calculations across thread blocks, data needs to be transferred through global memory, resulting in performance loss and bandwidth bottlenecks.

[0022] To solve the above problems, the present disclosure provides a GPU-based data processing method, which realizes cross-block data sharing by introducing a shared GPC memory with unified addressing inside the GPC, so that after matrix multiplication, the obtained product matrix can be directly reduced, reducing the dependence on global memory and improving the processing speed of matrix data.

[0023] The embodiments of the present disclosure will be described in detail below with reference to the accompanying drawings.

[0024] Figure 1 FIG. 1 is a schematic diagram of an exemplary system 100 in which various methods and apparatuses described herein may be implemented according to an embodiment of the present disclosure. Figure 1 , the system 100 includes one or more client devices 101, 102, 103, 104, 105, and 106, a server 120, and one or more communication networks 110 coupling the one or more client devices to the server 120. The client devices 101, 102, 103, 104, 105, and 106 may be configured to execute one or more applications.

[0025] In an embodiment of the present disclosure, the server 120 may run one or more services or software applications that enable execution of a GPU-based data processing method.

[0026] In some embodiments, server 120 may also provide other services or software applications that may include non-virtualized environments and virtualized environments. In some embodiments, these services may be provided as web-based services or cloud services, such as provided to users of client devices 101, 102, 103, 104, 105, and / or 106 under a software as a service (SaaS) model.

[0027] exist Figure 1In the configuration shown, the server 120 may include one or more components that implement the functions performed by the server 120. These components may include software components, hardware components, or a combination thereof that can be executed by one or more processors. Users operating client devices 101, 102, 103, 104, 105, and / or 106 may in turn utilize one or more client applications to interact with the server 120 to utilize the services provided by these components. It should be understood that a variety of different system configurations are possible, which may differ from the system 100. Therefore, Figure 1 is an example of a system for implementing the methods described herein and is not intended to be limiting.

[0028] The user may use client devices 101, 102, 103, 104, 105, and / or 106 to perform the GPU-based data processing method. The client device may provide an interface that enables a user of the client device to interact with the client device. The client device may also output information to the user via the interface. Figure 1 Only six client devices are depicted, but one skilled in the art will appreciate that the present disclosure may support any number of client devices.

[0029] Client devices 101, 102, 103, 104, 105 and / or 106 may include various types of computer devices, such as portable handheld devices, general-purpose computers (such as personal computers and laptop computers), workstation computers, wearable devices, smart screen devices, self-service terminal devices, service robots, game systems, thin clients, various messaging devices, sensors or other sensing devices, etc. These computer devices may run various types and versions of software applications and operating systems, such as MICROSOFT Windows, APPLE iOS, UNIX-like operating systems, Linux or Linux-like operating systems (such as GOOGLE Chrome OS); or include various mobile operating systems, such as MICROSOFT Windows Mobile OS, iOS, Windows Phone, Android. Portable handheld devices may include cellular phones, smart phones, tablet computers, personal digital assistants (PDAs), etc. Wearable devices may include head-mounted displays (such as smart glasses) and other devices. Game systems may include various handheld game devices, Internet-enabled game devices, etc. Client devices are capable of executing various different applications, such as various Internet-related applications, communication applications (such as email applications), short message service (SMS) applications, and may use various communication protocols.

[0030] The network 110 may be any type of network known to those skilled in the art that may support data communications using any of a variety of available protocols, including but not limited to TCP / IP, SNA, IPX, etc. By way of example only, the one or more networks 110 may be a local area network (LAN), an Ethernet-based network, a token ring, a wide area network (WAN), the Internet, a virtual network, a virtual private network (VPN), an intranet, an extranet, a public switched telephone network (PSTN), an infrared network, a wireless network (e.g., Bluetooth, WIFI), and / or any combination of these and / or other networks.

[0031] Server 120 may include one or more general purpose computers, dedicated server computers (e.g., PC (personal computer) servers, UNIX servers, mid-range servers), blade servers, mainframe computers, server clusters, or any other suitable arrangement and / or combination. Server 120 may include one or more virtual machines running virtual operating systems, or other computing architectures involving virtualization (e.g., one or more flexible pools of logical storage devices that may be virtualized to maintain a server's virtual storage device). In various embodiments, server 120 may run one or more services or software applications that provide the functionality described below.

[0032] The computing units in the server 120 may run one or more operating systems including any of the above operating systems and any commercially available server operating systems. The server 120 may also run any of a variety of additional server applications and / or middle-tier applications, including HTTP servers, FTP servers, CGI servers, JAVA servers, database servers, etc.

[0033] In some implementations, server 120 may include one or more applications to analyze and consolidate data feeds and / or event updates received from users of client devices 101, 102, 103, 104, 105, and 106. Server 120 may also include one or more applications to display data feeds and / or real-time events via one or more display devices of client devices 101, 102, 103, 104, 105, and 106.

[0034] In some embodiments, the server 120 may be a server of a distributed system, or a server combined with a blockchain. The server 120 may also be a cloud server, or an intelligent cloud computing server or intelligent cloud host with artificial intelligence technology. A cloud server is a host product in a cloud computing service system to solve the defects of difficult management and weak business scalability in traditional physical hosts and virtual private servers (VPS) services.

[0035] The system 100 may also include one or more databases 130. In some embodiments, these databases may be used to store data and other information. For example, one or more of the databases 130 may be used to store information such as audio files and video files. The databases 130 may reside in various locations. For example, the database used by the server 120 may be local to the server 120, or may be remote from the server 120 and may communicate with the server 120 via a network-based or dedicated connection. The databases 130 may be of different types. In some embodiments, the databases used by the server 120 may be, for example, relational databases. One or more of these databases may store, update, and retrieve data to and from the databases in response to commands.

[0036] In some embodiments, one or more of the databases 130 may also be used by applications to store application data. The databases used by the applications may be different types of databases, such as a key-value store, an object store, or a conventional store backed by a file system.

[0037] Figure 1 The system 100 may be configured and operated in various ways to enable the application of various methods and apparatuses described in the present disclosure.

[0038] Figure 2 A flowchart of a GPU-based data processing method according to an embodiment of the present disclosure is shown.

[0039] like Figure 2 As shown, the GPU-based data processing method 200 includes:

[0040] Step S201, obtaining a first matrix and a second matrix to be processed from the global memory of the GPU;

[0041] Step S202: storing the first matrix in a first sub-memory of a shared memory corresponding to a target thread block in a graphics processing cluster GPC;

[0042] Step S203, storing the second matrix in the second sub-memory of the shared memory corresponding to the target thread block;

[0043] Step S204, performing a multiplication operation on the first matrix and the second matrix in the shared memory;

[0044] Step S205: storing the third matrix obtained by the multiplication operation into a shared GPC memory of a plurality of thread blocks within a graphics processing cluster GPC of the GPU, wherein the plurality of thread blocks include the target thread block, and wherein the shared GPC memory is uniformly addressed for the plurality of thread blocks;

[0045] Step S206: performing a reduction operation on the third matrix in the shared GPC memory.

[0046] It is understandable that in the prior art, the matrix multiplication method based on GPU mainly stores data in the shared memory SMEM corresponding to the block, and each block has its own corresponding SMEM. The multiplied product result matrix needs to be written to the external global memory first for block division and reduction operations, and then the reduction result is stored in the corresponding shared memory inside the GPC. The blocks inside the GPC can only synchronize data through the external global memory, and only the threads inside the block can share the shared memory SMEM corresponding to the block. In the calculation process of large models and / or neural network models, a large number of matrix operations are involved, and the additional data writing and reading operations between the block and the global memory will seriously affect the computing performance of the model.

[0047] To this end, the present application introduces a shared GPC memory with unified addressing inside the GPC. This memory can be shared by all blocks inside the GPC, thereby realizing data sharing across blocks. Therefore, after matrix multiplication, the result matrix of the product can be directly reduced without the need for additional data communication through the global memory to achieve data synchronization, thereby reducing the dependence on the global memory and improving the processing speed of the matrix data.

[0048] Matrix multiplication and reduction operations are very basic operations in physical computing and play an important role in many complex scientific simulations. In recent years, they have been widely used in large models, machine learning and neural network models. Through parallel computing and efficient matrix operations, the performance of large models, machine learning and neural network models can be improved.

[0049] According to some embodiments, method 200 further includes: allocating an address space for the multiple thread blocks within the graphics processing cluster GPC as the shared GPC memory of the multiple thread blocks, wherein the size of the shared GPC memory is associated with the size of the shared memory corresponding to the multiple thread blocks.

[0050] Exemplarily, the size of the shared GPC memory can be configured, and is a fixed value with the total space size of the shared memory SMEM. Therefore, the size of the shared GPC memory can be flexibly configured according to different computing tasks to adapt to different application scenarios. For smaller matrices or data sets, a smaller shared GPC memory can be configured to reduce memory waste; for large-scale calculations, a larger shared GPC memory can be configured to ensure sufficient memory for sharing and calculation across blocks. This flexibility can improve computing efficiency and avoid performance degradation caused by insufficient or wasted memory.

[0051] Figure 3 A schematic diagram of a hardware structure of a shared GPC memory according to an embodiment of the present disclosure is shown.

[0052] like Figure 3 As shown, four streaming processors (Streaming Multiprocessor, SM) located inside the GPC and the computing unit CU therein and the corresponding shared memory SMEM are shown. It can be seen that each shared memory SMEM includes multiple banks. Based on the shared memory SMEM, the shared GPC memory is distributed inside the GPC and is a continuous encoding space. This part of the address space is shared between blocks belonging to the same grid and distributed inside the same GPC. The shared GPC memory is uniformly addressed for all thread blocks inside the GPC. This means that all thread blocks in the same GPC use the same address space to read and write data. This unified addressing method allows different blocks to work together efficiently and avoids repeated memory loading operations. Since the shared GPC memory is shared between thread blocks and uses a continuous address space, the memory access efficiency is high, the memory access overhead is reduced, and the communication speed between thread blocks is improved.

[0053] According to some embodiments, step S202 includes: determining the width of the first matrix and the size of the target thread block; dividing the first matrix into at least one block according to the width of the first matrix and the size of the target thread block; and writing each block in the at least one block in sequence into the first sub-memory of the shared memory corresponding to the target thread block.

[0054] According to some embodiments, step S203 includes: determining the width of the second matrix and the size of the target thread block; dividing the second matrix into at least one block according to the width of the second matrix and the size of the target thread block; and writing each block of the at least one block of the second matrix into the second sub-memory of the shared memory corresponding to the target thread block in sequence.

[0055] The following shows example code to illustrate the process of loading the first matrix and the second matrix from global memory into shared memory.

[0056]

[0057] As shown in the above code, width is used to represent the widths of the first matrix and the second matrix, and BLOCKSIZE represents the size of the target thread block. Then, through the loop statement for(int m = 0; m < width / BLOCKSIZE; m++), the first matrix and the second matrix can be divided into at least one block of size BLOCKSIZE respectively, and the blocks are written into the shared memory sequentially using the loop, thereby implementing the storage of the first matrix and the second matrix from global memory to shared memory, so that all threads in the target thread block can share this part of the data and perform calculation operations on it, avoiding repeatedly reading data from global memory.

[0058] According to some embodiments, step S206 includes: performing an addition operation on every two adjacent rows of data of the third matrix in the shared GPC memory.

[0059] In many matrix calculation scenarios, multiple thread blocks perform calculations independently. After each completes its calculation, it is necessary to summarize their calculation results (for example, summing the partial matrices calculated by each block, or adding the results calculated by multiple blocks). Using a reduction operation can efficiently integrate the calculation results of multiple thread blocks and improve the calculation efficiency of the final result. Through addition reduction, multiple accesses to global memory can be reduced, and data can be quickly integrated at the same time.

[0060] According to some embodiments, the data in the shared GPC memory is indexed based on the position of the thread block corresponding to the data in the grid and the index value of the thread.

[0061] By using unified addressing based on the thread block and the thread index, it can be ensured that data access in different thread blocks is consistent and predictable. The unified addressing method can effectively manage the data in the shared memory and ensure that data access between multiple thread blocks does not conflict.

[0062] The following gives example code for the shared GPC memory for the multiplication operation of two matrices and the reduction operation of the product matrix.

[0063] Exemplarily, the declaration of variables can be represented by the following code:

[0064] __gpc_shared__ float M_gpc_Shared[width][width]

[0065] Int gpc_x=blockIdx.x*blockDim.x+threadIdx.x

[0066] Int gpc_y=blockIdx.y*blockDim.y+threadIdx.y

[0067] It is understandable that shared GPC memory is declared and accessed in a similar way to shared memory. __gpc_shared_ is used to declare a specific memory space that can be shared among all thread blocks in the same GPC.

[0068] float M_gpc_Shared[width][width] declares a two-dimensional array M_gpc_Shared, whose size is width x width and type is float. This is the matrix data storage space allocated to the shared GPC memory, and multiple thread blocks can share access to the array. This helps to exchange data between different thread blocks and reduce redundant calculations and memory transfers.

[0069] The thread index in the GPC is calculated by gpc_x and gpc_y, and is indexed based on the position of the thread block corresponding to the current thread (or current data) in the grid and the thread index value. _gpc_shared__ data storage involves data communication between blocks, so the blockIdx parameter is added to the declaration of shared memory to locate the specific block, and then threadIdx.x and threadIdx.y are used to locate the current thread.

[0070] The multiplication of the first matrix and the second matrix and the reduction operation of the resulting product, the third matrix, can be represented by the following code:

[0071]

[0072]

[0073] Among them, M_device points to the storage location of the first matrix M in the global memory, N_device points to the storage location of the second matrix N in the global memory, and P_device points to the storage location of the third matrix P of the multiplication result.

[0074] M_deviceShared and N_deviceShared declare local matrices for each thread block in shared memory to store the block size (BLOCKSIZE * BLOCKSIZE) and temporarily store the matrix data of each block in the blocked first matrix and the blocked second matrix.

[0075] A_deviceShared declares shared GPC memory through the __gpc_shared__ keyword, allowing multiple thread blocks in the same GPC to share a section of memory with a unified address space to store the third matrix and perform reduction operations.

[0076] During the calculation, M_deviceShared[ty][tx] = M_device[y * width + (m * BLOCKSIZE + tx)] is used to load a part of the data of the first matrix M_device in global memory into the shared memory M_deviceShared. Each thread is responsible for loading a block of data. Similarly, N_deviceShared[ty][tx] = N_device[(m * BLOCKSIZE + ty) * width + x] is used to load the corresponding part of the second matrix N_device into the shared memory N_deviceShared. __syncthreads() is used to synchronize the threads to ensure that all threads complete the data loading before continuing to execute.

[0077] Thus, by blocking the matrix and using the loop for(int m = 0; m < width / BLOCKSIZE; m++) to successively load each block of the first matrix into the shared memory, and then using the loop for(int k = 0; k < BLOCKSIZE; k++) to traverse the elements of each row and each column of the first matrix and the second matrix in the shared memory, the first matrix and the second matrix are multiplied to obtain P_element, which is the partial matrix product result calculated by the current thread. After the loop ends, the product result of the first matrix and the second matrix, the third matrix, is stored in the shared GPC memory through P_device[gpc_y][gpc_x] = P_element.

[0078] Furthermore, the following code segment performs a reduction operation on the third matrix P in the shared GPC memory by the dichotomy method, that is, adding the data of every two adjacent rows of the P matrix.

[0079] If(y / 2 == 0){

[0080] A_deviceShared[y] = A_deviceShared[y][x] + A_deviceShared[y + 1][x];

[0081] }

[0082] __syncthreads();

[0083] If (y / 2 == 0) {

[0084] P_deviceShared[y / 2 + x] = A_deviceShared[y][x];

[0085] }

[0086] According to some embodiments, the data reuse efficiency can be further improved by using shared GPC memory. As shown in the following exemplary code, the statements M_deviceShared = M_device and N_deviceShared = N_device can be used to directly load the first matrix M and the second matrix N from the global memory into the shared GPC memory, and the multiplication operation on the first matrix M and the second matrix N can be performed in the shared GPC memory using for (int k = 0; k < BLOCKSIZE; k++) { P_element += M_deviceShared[gpc_y][k] * N_deviceShared[k][gpc_x]}, and the reduction operation on the result of the product, the third matrix P, can be performed. Thus, the data stored in the shared GPC memory can be reused between different blocks without repeatedly reading from the global memory.

[0087]

[0088]

[0089] According to another aspect of the present disclosure, a GPU-based data processing device is provided. As Figure 4As shown, the GPU-based data processing device 400 includes: an acquisition module 401, configured to acquire a first matrix and a second matrix to be processed from a global memory of a GPU; a first storage module 402, configured to store the first matrix in a first sub-memory of a shared memory corresponding to a target thread block within a graphics processing cluster GPC; a second storage module 403, configured to store the second matrix in a second sub-memory of a shared memory corresponding to the target thread block; a first calculation module 404, configured to perform a multiplication operation on the first matrix and the second matrix in the shared memory; a third storage module 405, configured to store a third matrix obtained by the multiplication operation in a shared GPC memory of a plurality of thread blocks within a graphics processing cluster GPC of the GPU, wherein the plurality of thread blocks include the target thread block, and wherein the shared GPC memory is uniformly addressed for the plurality of thread blocks; and a second calculation module 406, configured to perform a reduction operation on the third matrix in the shared GPC memory.

[0090] The GPU-based data processing device 400 introduces a shared GPC memory with unified addressing inside the GPC. This memory can be shared by all blocks inside the GPC, thereby realizing data sharing across blocks. Therefore, after matrix multiplication, the result matrix of the product can be directly reduced without the need for additional data communication through the global memory to achieve data synchronization, thereby reducing the dependence on the global memory and improving the calculation speed of the matrix.

[0091] According to some embodiments, the GPU-based data processing device 400 also includes: an allocation module, configured to allocate an address space for the multiple thread blocks inside the graphics processing cluster GPC as the shared GPC memory of the multiple thread blocks, wherein the size of the shared GPC memory is associated with the size of the shared memory corresponding to the multiple thread blocks.

[0092] Exemplarily, the size of the shared GPC memory can be configured, and is a fixed value with the total space size of the shared memory SMEM. Therefore, the size of the shared GPC memory can be flexibly configured by the allocation module according to different computing tasks to adapt to different application scenarios. For smaller matrices or data sets, a smaller shared GPC memory can be configured to reduce memory waste; for large-scale calculations, a larger shared GPC memory can be configured to ensure sufficient memory for sharing and calculation across blocks. This flexibility can improve computing efficiency and avoid performance degradation caused by insufficient or wasted memory.

[0093] According to some embodiments, the first storage module 402 includes: a first determination unit, configured to determine the width of the first matrix and the size of the target thread block; a first division unit, configured to divide the first matrix into at least one block according to the width of the first matrix and the size of the target thread block; and a first writing unit, configured to write each block of the at least one block in sequence into the first sub-memory of the shared memory corresponding to the target thread block.

[0094] According to some embodiments, the second storage module 403 includes: a second determination unit, configured to determine the width of the second matrix and the size of the target thread block; a second division unit, configured to divide the second matrix into at least one block according to the width of the second matrix and the size of the target thread block; and a second writing unit, configured to write each block of the at least one block of the second matrix into the second sub-memory of the shared memory corresponding to the target thread block in sequence.

[0095] By dividing the first matrix and the second matrix into at least one block respectively, and writing the blocks into the shared memory in sequence by using a loop, the first matrix and the second matrix are stored from the global memory to the shared memory, so that all threads in the target thread block can share this part of data and perform calculation operations on it, thereby avoiding repeated reading of data from the global memory.

[0096] According to some embodiments, the second calculation module 406 is further configured to: perform an addition operation on every two adjacent rows of data of the third matrix in the shared GPC memory.

[0097] In many matrix calculation scenarios, multiple thread blocks perform calculations independently. After each block completes its calculation, its calculation results need to be summarized (for example, summing the partial matrix calculated by each block, or summing the results calculated by multiple blocks). Using reduction operations can efficiently integrate the calculation results of multiple thread blocks and improve the calculation efficiency of the final result. Additive reduction can reduce multiple accesses to global memory and quickly integrate data.

[0098] According to some embodiments, data in the shared GPC memory is indexed based on the position of the thread block corresponding to the data in the grid and the index value of the thread.

[0099] Through unified addressing based on thread blocks and thread indexes, data access in different thread blocks can be guaranteed to be consistent and predictable. The unified addressing method can effectively manage data in shared memory and ensure that data access between multiple thread blocks does not conflict.

[0100] According to another aspect of the present disclosure, an electronic device is also provided, comprising: at least one processor; and a memory communicatively connected to the at least one processor; wherein the memory stores instructions executable by the at least one processor, and the instructions are executed by the at least one processor so that the at least one processor can execute the aforementioned method.

[0101] According to another aspect of the present disclosure, a non-transitory computer-readable storage medium storing computer instructions is further provided, wherein the computer instructions are used to enable the computer to execute the aforementioned method.

[0102] According to another aspect of the present disclosure, a computer program product is also provided, including a computer program, wherein the computer program implements the aforementioned method when executed by a processor.

[0103] like Figure 5 As shown, the electronic device 500 includes a computing unit 501, which can perform various appropriate actions and processes according to a computer program stored in a read-only memory (ROM) 502 or a computer program loaded from a storage unit 508 into a random access memory (RAM) 503. In the RAM 503, various programs and data required for the operation of the electronic device 500 can also be stored. The computing unit 501, the ROM 502, and the RAM 503 are connected to each other via a bus 504. An input / output (I / O) interface 505 is also connected to the bus 504.

[0104] Multiple components in the electronic device 500 are connected to the I / O interface 505, including: an input unit 506, an output unit 507, a storage unit 508 and a communication unit 509. The input unit 506 can be any type of device that can input information to the electronic device 500. The input unit 506 can receive input digital or character information, and generate key signal input related to user settings and / or function control of the electronic device, and can include but is not limited to a mouse, a keyboard, a touch screen, a trackpad, a trackball, a joystick, a microphone and / or a remote control. The output unit 507 can be any type of device that can present information, and can include but is not limited to a display, a speaker, a video / audio output terminal, a vibrator and / or a printer. The storage unit 508 can include but is not limited to a disk, an optical disk. The communication unit 509 allows the electronic device 500 to exchange information / data with other devices through a computer network such as the Internet and / or various telecommunication networks, and can include but is not limited to a modem, a network card, an infrared communication device, a wireless communication transceiver and / or a chipset, such as Bluetooth TM devices, 802.11 devices, WiFi devices, WiMax devices, cellular communication devices, and / or the like.

[0105] The computing unit 501 may be a variety of general and / or special processing components with processing and computing capabilities. Some examples of the computing unit 501 include, but are not limited to, a central processing unit (CPU), a graphics processing unit (GPU), various dedicated artificial intelligence (AI) computing chips, various computing units running machine learning model algorithms, digital signal processors (DSPs), and any appropriate processors, controllers, microcontrollers, etc. The computing unit 501 performs the various methods and processes described above, such as a GPU-based data processing method. For example, in some embodiments, a GPU-based data processing method may be implemented as a computer software program, which is tangibly contained in a machine-readable medium, such as a storage unit 508. In some embodiments, part or all of the computer program may be loaded and / or installed on the electronic device 500 via ROM 502 and / or communication unit 509. When the computer program is loaded into RAM 503 and executed by the computing unit 501, one or more steps of the GPU-based data processing method described above may be performed. Alternatively, in other embodiments, the computing unit 501 may be configured to perform a GPU-based data processing method in any other appropriate manner (e.g., by means of firmware).

[0106] Various implementations of the systems and techniques described above herein can be implemented in digital electronic circuit systems, integrated circuit systems, field programmable gate arrays (FPGAs), application specific integrated circuits (ASICs), application specific standard products (ASSPs), systems on chips (SOCs), complex programmable logic devices (CPLDs), computer hardware, firmware, software, and / or combinations thereof. These various implementations can include: being implemented in one or more computer programs that can be executed and / or interpreted on a programmable system including at least one programmable processor, which can be a special purpose or general purpose programmable processor that can receive data and instructions from a storage system, at least one input device, and at least one output device, and transmit data and instructions to the storage system, the at least one input device, and the at least one output device.

[0107] The program code for implementing the method of the present disclosure may be written in any combination of one or more programming languages. These program codes may be provided to a processor or controller of a general-purpose computer, a special-purpose computer, or other programmable data processing device, so that the program code, when executed by the processor or controller, enables the functions / operations specified in the flow chart and / or block diagram to be implemented. The program code may be executed entirely on the machine, partially on the machine, partially on the machine and partially on a remote machine as a stand-alone software package, or entirely on a remote machine or server.

[0108] In the context of the present disclosure, a machine-readable medium may be a tangible medium that may contain or store a program for use by or in conjunction with an instruction execution system, device, or equipment. A machine-readable medium may be a machine-readable signal medium or a machine-readable storage medium. A machine-readable medium may include, but is not limited to, an electronic, magnetic, optical, electromagnetic, infrared, or semiconductor system, device, or equipment, or any suitable combination of the foregoing. A more specific example of a machine-readable storage medium may include an electrical connection based on one or more lines, a portable computer disk, a hard disk, a random access memory (RAM), a read-only memory (ROM), an erasable programmable read-only memory (EPROM or flash memory), an optical fiber, a portable compact disk read-only memory (CD-ROM), an optical storage device, a magnetic storage device, or any suitable combination of the foregoing.

[0109] To provide interaction with a user, the systems and techniques described herein can be implemented on a computer having: a display device (e.g., a CRT (cathode ray tube) or LCD (liquid crystal display) monitor) for displaying information to the user; and a keyboard and pointing device (e.g., a mouse or trackball) through which the user can provide input to the computer. Other types of devices can also be used to provide interaction with the user; for example, the feedback provided to the user can be any form of sensory feedback (e.g., visual feedback, auditory feedback, or tactile feedback); and input from the user can be received in any form (including acoustic input, voice input, or tactile input).

[0110] The systems and techniques described herein may be implemented in a computing system that includes back-end components (e.g., as a data server), or a computing system that includes middleware components (e.g., an application server), or a computing system that includes front-end components (e.g., a user computer with a graphical user interface or a web browser through which a user can interact with implementations of the systems and techniques described herein), or a computing system that includes any combination of such back-end components, middleware components, or front-end components. The components of the system may be interconnected by any form or medium of digital data communication (e.g., a communication network). Examples of communication networks include: a local area network (LAN), a wide area network (WAN), and the Internet.

[0111] A computer system may include a client and a server. The client and the server are generally remote from each other and usually interact through a communication network. The relationship of client and server is generated by computer programs running on respective computers and having a client-server relationship with each other. The server may be a cloud server, a server of a distributed system, or a server combined with a blockchain.

[0112] It should be understood that the various forms of processes shown above can be used to reorder, add or delete steps. For example, the steps described in this disclosure can be performed in parallel, sequentially or in different orders, as long as the desired results of the technical solutions disclosed in this disclosure can be achieved, and this document does not limit this.

[0113] Although the embodiments or examples of the present disclosure have been described with reference to the accompanying drawings, it should be understood that the above-mentioned methods, systems and devices are merely exemplary embodiments or examples, and the scope of the present invention is not limited by these embodiments or examples, but only by the claims after authorization and their equivalent scope. Various elements in the embodiments or examples can be omitted or replaced by their equivalent elements. In addition, each step can be performed in an order different from that described in the present disclosure. Further, the various elements in the embodiments or examples can be combined in various ways. It is important that with the evolution of technology, many elements described herein can be replaced by equivalent elements that appear after the present disclosure.

Claims

1. A data processing method based on GPU, comprising: Get the first matrix and the second matrix to be processed from the global memory of the GPU; storing the first matrix in a first sub-memory of a shared memory corresponding to a target thread block within a graphics processing cluster GPC; storing the second matrix in a second sub-memory of the shared memory corresponding to the target thread block; Performing a multiplication operation on the first matrix and the second matrix in the shared memory; storing the third matrix obtained by the multiplication operation into a shared GPC memory of a plurality of thread blocks within a graphics processing cluster (GPC) of the GPU, wherein the plurality of thread blocks include the target thread block, and wherein the shared GPC memory is uniformly addressed for the plurality of thread blocks; and A reduction operation is performed on the third matrix in the shared GPC memory.

2. The method according to claim 1, further comprising: An address space is allocated to the multiple thread blocks inside the graphics processing cluster GPC as the shared GPC memory of the multiple thread blocks, wherein the size of the shared GPC memory is associated with the size of the shared memory corresponding to the multiple thread blocks.

3. The method according to claim 1 or 2, wherein: The step of storing the first matrix in a first sub-memory of a shared memory corresponding to a target thread block within a graphics processing cluster GPC comprises: Determining a width of the first matrix and a size of the target thread block; Dividing the first matrix into at least one block according to the width of the first matrix and the size of the target thread block; and Each block of the at least one block is sequentially written into the first sub-memory of the shared memory corresponding to the target thread block.

4. The method according to any one of claims 1 to 3, wherein: The storing the second matrix into the second sub-memory of the shared memory corresponding to the target thread block comprises: Determining a width of the second matrix and a size of the target thread block; dividing the second matrix into at least one block according to the width of the second matrix and the size of the target thread block; and Each block of at least one block of the second matrix is ​​sequentially written into the second sub-memory of the shared memory corresponding to the target thread block.

5. The method according to any one of claims 1 to 4, wherein: The reducing operation on the third matrix in the shared GPC memory includes: An addition operation is performed on every two adjacent rows of data of the third matrix in the shared GPC memory.

6. The method according to any one of claims 1 to 5, wherein: The data in the shared GPC memory is indexed based on the position of the thread block corresponding to the data in the grid and the index value of the thread.

7. A data processing device based on a GPU, comprising: An acquisition module is configured to acquire a first matrix and a second matrix to be processed from a global memory of a GPU; A first storage module is configured to store the first matrix in a first sub-memory of a shared memory corresponding to a target thread block in a graphics processing cluster GPC; A second storage module is configured to store the second matrix in a second sub-memory of the shared memory corresponding to the target thread block; A first computing module is configured to perform a multiplication operation on the first matrix and the second matrix in the shared memory; a third storage module configured to store the third matrix obtained by the multiplication operation into a shared GPC memory of a plurality of thread blocks within a graphics processing cluster GPC of the GPU, wherein the plurality of thread blocks include the target thread block, and wherein the shared GPC memory is uniformly addressed for the plurality of thread blocks; and The second computing module is configured to perform a reduction operation on the third matrix in the shared GPC memory.

8. The apparatus according to claim 7, further comprising: The allocation module is configured to allocate an address space for the multiple thread blocks inside the graphics processing cluster GPC as the shared GPC memory of the multiple thread blocks, wherein the size of the shared GPC memory is associated with the size of the shared memory corresponding to the multiple thread blocks.

9. The device according to claim 7 or 8, wherein: The first storage module comprises: A first determining unit is configured to determine a width of the first matrix and a size of the target thread block; a first dividing unit configured to divide the first matrix into at least one block according to the width of the first matrix and the size of the target thread block; and The first writing unit is configured to write each block of the at least one block in sequence into the first sub-memory of the shared memory corresponding to the target thread block.

10. The device according to any one of claims 7 to 9, wherein: The second storage module comprises: a second determining unit configured to determine a width of the second matrix and a size of the target thread block; a second dividing unit configured to divide the second matrix into at least one block according to the width of the second matrix and the size of the target thread block; and The second writing unit is configured to write each block of the at least one block of the second matrix into the second sub-memory of the shared memory corresponding to the target thread block in sequence.

11. The device according to any one of claims 7 to 10, wherein: The second computing module is further configured to: An addition operation is performed on every two adjacent rows of data of the third matrix in the shared GPC memory.

12. The device according to any one of claims 7 to 11, wherein: The data in the shared GPC memory is indexed based on the position of the thread block corresponding to the data in the grid and the index value of the thread.

13. An electronic device comprising: at least one processor; as well as a memory communicatively coupled to the at least one processor; in The memory stores instructions that can be executed by the at least one processor, and the instructions are executed by the at least one processor to enable the at least one processor to perform the method according to any one of claims 1 to 6.

14. A non-transitory computer-readable storage medium storing computer instructions, wherein: The computer instructions are used to cause the computer to execute the method according to any one of claims 1-6.

15. A computer program product comprising a computer program, wherein: When the computer program is executed by a processor, the method according to any one of claims 1 to 6 is implemented.

Images