Artificial intelligence chip, method for synchronizing matrix multiplication instructions, computing device, medium and program product

Abstract

The present application relates to an artificial intelligence chip, a method for synchronizing matrix multiplication instructions, a computing device, a medium and a program product. The chip comprises: a shared memory, including a data synchronization management unit configured to send information indicating that input data in the shared memory is ready to a matrix multiplication instruction management unit; a tensor core including the matrix multiplication instruction management unit and a matrix multiplication operation unit, the matrix multiplication instruction management unit is configured to send the address of the data block required for each loop calculation to the shared memory after receiving the information indicating that the input data in the shared memory is ready, and send information indicating that the shared memory space has been released to the data synchronization management unit when the address of the data block of the last loop calculation is sent; the matrix multiplication operation unit is configured to perform matrix multiplication operation for each loop calculation. The present application can significantly improve the execution efficiency of matrix multiplication instructions and the utilization rate of shared memory space.

Description

TECHNICAL FIELD

[0001] Embodiments of the present application generally relate to the field of artificial intelligence, and more particularly, to an artificial intelligence chip, a method for synchronizing a matrix multiplication instruction, a computing device, a computer readable storage medium, and a computer program product. BACKGROUND

[0002] In the existing method for synchronizing a matrix multiplication instruction, initialization synchronization is first performed, for example, by using an instruction "wgmma.fence.sync.aligned" to ensure that a memory operation (such as a global memory load) has been completed, and then the matrix multiplication instruction is acquired. And by using the instruction "wgmma.wait_group", the execution thread is caused to wait until the completion of the waiting matrix multiplication instruction. It should be understood that the memory space release is earlier than the completion of the matrix multiplication instruction, and therefore, the above-mentioned existing method for synchronizing a matrix multiplication instruction is prone to cause the acquisition of the matrix multiplication instruction and the use of the released memory space to lag, thereby reducing the execution efficiency of the matrix multiplication instruction and wasting the space in the shared memory.

[0003] In summary, the existing method for synchronizing a matrix multiplication instruction has the disadvantages of reducing the execution efficiency of the matrix multiplication instruction and wasting the space in the shared memory. SUMMARY

[0004] The present application provides an artificial intelligence chip, a method for synchronizing a matrix multiplication instruction, a computing device, a computer readable storage medium, and a computer program product, which can significantly improve the execution efficiency of the matrix multiplication instruction and the utilization rate of the shared memory space.

[0005] According to a first aspect of the present application, an artificial intelligence chip is provided, which comprises: a shared memory comprising a data synchronization management unit, the data synchronization management unit being configured to send information indicating that input data in the shared memory is ready to a matrix multiplication instruction management unit; and a tensor core comprising the matrix multiplication instruction management unit and a matrix multiplication operation unit, the matrix multiplication instruction management unit being configured to send an address of a data block required for each loop calculation to the shared memory after receiving the information indicating that the input data in the shared memory is ready, and to send information indicating that the shared memory space has been released to the data synchronization management unit when the address of the data block required for the last loop calculation is sent; and the matrix multiplication operation unit being configured to perform each loop calculation related to the matrix multiplication.

[0006] According to a second aspect of the present application, there is also provided a method for synchronizing a matrix multiplication instruction, the method comprising: a matrix multiplication instruction management unit of a tensor core sending, after confirming that information indicating that input data in a shared memory is ready is received, an address of a data block required for each loop calculation to the shared memory for a matrix multiplication operation unit to perform each loop calculation on matrix multiplication; and the matrix multiplication instruction management unit synchronously sending, when the address of the data block required for the last loop calculation is sent, information indicating that the shared memory space has been released to a data synchronization management unit.

[0007] According to a third aspect of the present application, there is also provided a computing device. The computing device comprises: at least one processor; and a memory communicatively connected with the at least one processor; wherein the memory stores instructions executable by the at least one processor, the instructions being executed by the at least one processor to enable the computing device to perform the method of the first aspect of the present application.

[0008] According to a fourth aspect of the present application, there is also provided a computer readable storage medium. The computer readable storage medium stores a computer program, the computer program being executed by a machine to perform the method of the first aspect of the present application.

[0009] According to a fifth aspect of the present application, there is also provided a computer program product comprising a computer program, the computer program being executed by a machine to perform the method of the first aspect of the present application.

[0010] In some embodiments, the matrix multiplication operation unit is further configured to: send instruction synchronization address information indicating that the matrix multiplication instruction operation result output is complete to an instruction synchronization management unit in the scalar core while outputting the matrix multiplication instruction operation result to the register of the scalar core; and the artificial intelligence chip further comprises the scalar core, the scalar core comprising: the instruction synchronization management unit, the register, and the instruction queue.

[0011] In some embodiments, the instruction synchronization management unit is configured to: after receiving the instruction synchronization address information, cause the instruction queue to send other scalar instructions to process the data output to the register of the scalar core.

[0012] In some embodiments, the matrix multiplication operation unit sends instruction synchronization address information indicating that the matrix multiplication instruction operation result output is complete to an instruction synchronization management unit in the scalar core while outputting the matrix multiplication instruction operation result to the register of the scalar core.

[0013] In some embodiments, the matrix multiplication instruction management unit of the tensor core obtains a matrix multiplication instruction, the matrix multiplication instruction indicating at least: first shared memory barrier address information for synchronizing input data readiness, second shared memory barrier address information for synchronizing memory space readiness, and instruction synchronization address information for synchronizing matrix multiplication instruction operation result output completion.

[0014] In some embodiments, after receiving the matrix multiplication instruction, the matrix multiplication instruction management unit sends the first shared memory barrier address information to the data synchronization management unit of the shared memory; after receiving the first shared memory barrier address information, the data synchronization management unit of the shared memory confirms whether the data in the shared memory corresponding to the first shared memory barrier address information is ready; in response to confirming that the data in the shared memory corresponding to the first shared memory barrier address information is ready, the data synchronization management unit of the shared memory sends information indicating that the input data in the shared memory is ready to the matrix multiplication instruction management unit; and in response to confirming that the data in the shared memory corresponding to the first shared memory barrier address information is not ready, the data synchronization management unit of the shared memory sends information indicating that the input data in the shared memory is not ready to the matrix multiplication instruction management unit.

[0015] In some embodiments, sending the address of the data block required for each loop calculation to the shared memory for the matrix multiplication operation unit to perform each loop calculation on the matrix multiplication includes: after confirming that the information indicating that the input data in the shared memory is ready is received from the data synchronization management unit, the matrix multiplication instruction management unit of the tensor core starts calculating the address of the data block required for each loop calculation to be sent to the shared memory; and sending the address of the data block required for each loop calculation to the shared memory for each loop, so that the shared memory provides the data block corresponding to the address to the matrix multiplication operation unit to perform each loop calculation.

[0016] In some embodiments, the matrix multiplication instruction management unit synchronously sends information indicating that the shared memory space has been released to the data synchronization management unit when the address of the data block required for the last loop calculation is sent includes: the matrix multiplication instruction management unit synchronously sends the second shared memory barrier address information to the data synchronization management unit of the shared memory when the address of the data block required for the last loop calculation is sent, so as to inform the shared memory space that has been released, so that other data loading instructions can store data to the shared memory space that has been released.

[0017] In some embodiments, after the instruction synchronization management unit in the scalar core receives the instruction synchronization address information, the instruction queue is caused to send other scalar instructions to process the data output to the register of the scalar core.

[0018] The application can release the corresponding storage space in the shared memory immediately after the input data in the shared memory is used up, without waiting for the matrix multiplication instruction to be executed, thereby separating the release of the shared memory space and the completion of the matrix multiplication instruction, and thus significantly improving the execution efficiency of the matrix multiplication instruction and the utilization rate of the shared memory space.

[0019] It should be understood that the content described in this part is not intended to identify key or important features of the embodiments of the application, nor to limit the scope of the application. Other features of the application will become apparent from the following description. BRIEF DESCRIPTION OF DRAWINGS

[0020] The above and other features, advantages, and aspects of embodiments of the application will become more apparent by describing in detail preferred embodiments thereof with reference to the attached drawings in which:

[0021] Figure 1 A schematic diagram of a computing device implementing a method for synchronizing a matrix multiplication instruction according to embodiments of the application is shown schematically.

[0022] Figure 2 A flowchart of a method for synchronizing a matrix multiplication instruction according to some embodiments of the application is shown.

[0023] Figure 3 A structural schematic diagram of an artificial intelligence chip according to some embodiments of the application is shown.

[0024] Figure 4 A schematic diagram of synchronization between a shared memory, a tensor core, and a register according to some embodiments of the application is shown.

[0025] Figure 5 A timing diagram comparing the execution of actions of a method for synchronizing a matrix multiplication instruction according to some embodiments of the application is shown schematically.

[0026] Figure 6 A flowchart of a method for sending information indicating that input data in a shared memory is ready to a matrix multiplication instruction management unit according to some embodiments of the application is shown.

[0027] In the various drawings, like or corresponding elements are denoted by like or corresponding reference numerals. DETAILED DESCRIPTION

[0028] Preferred embodiments of the present application will be described herein below with reference to the accompanying drawings. While preferred embodiments of the present application are shown in the drawings, it is understood that the present application can be embodied in various forms and should not be limited by the embodiments set forth herein. Rather, these embodiments are provided so that this disclosure will be thorough and complete, and will fully convey the scope of the present application to those skilled in the art.

[0029] The term "comprising" and variations thereof as used herein are intended to cover a non-exclusive inclusion, such that a process, method, article, or apparatus that comprises a list of elements is not necessarily limited to those elements, but can include other elements not expressly listed or inherent to such process, method, article, or apparatus. Unless otherwise defined, all terms (including technical and scientific terms) used herein have the same meaning as commonly understood by one of ordinary skill in the art to which this application belongs. It will be further understood that terms, such as those defined in commonly used dictionaries, should be interpreted as having a meaning that is consistent with their meaning in the context of the relevant art and the present disclosure, and will not be interpreted in an overly literal or overly formal sense unless expressly so defined herein. The terminology used herein is for the purpose of describing particular embodiments only and is not intended to be limiting of the present application. As used herein, the singular forms "a", "an" and "the" are intended to include the plural forms as well, unless the context clearly indicates otherwise. It will be further understood that the terms "comprises" and / or "comprising," when used in this specification, specify the presence of stated features, integers, steps, operations, elements, and / or components, but do not preclude the presence or addition of one or more other features, integers, steps, operations, elements, components, and / or groups thereof. As used herein, the term "and / or" includes any and all combinations of one or more of the associated listed items. As used herein, the term "based on" means "based, at least in part, on" unless expressly specified otherwise. As used herein, the term "one embodiment" is not used to limit the number of embodiments specifi ed in the claim to a single embodiment, but rather used to clarify that the embodiment so specifi ed is among many alternatives that can differ from that specifi ed. As used herein, the term "another embodiment" is not used to limit the number of embodiments specifi ed in the claim to a single embodiment, but rather used to clarify that the embodiment so specifi ed is among many alternatives that can differ from that specifi ed. As used herein, the terms "first," "second," and the like, do not imply

[0030] As described previously, the conventional method for synchronizing matrix multiplication instructions has the disadvantage of reducing the execution efficiency of the matrix multiplication instruction and wasting the space in the shared memory.

[0031] To at least partially solve one or more of the above problems and other potential problems, example embodiments of the present application at least propose an artificial intelligence chip and a method for synchronizing matrix multiplication instructions. In the method, by causing the matrix multiplication instruction management unit of the tensor core to send the address of the data block required for each loop calculation to the shared memory after confirming that the information indicating that the input data in the shared memory is ready is received, for the matrix multiplication operation unit to perform each loop calculation on the matrix multiplication; and the matrix multiplication instruction management unit synchronously sends the information indicating that the shared memory space has been released to the data synchronization management unit when the address of the data block required for the last loop calculation is sent, so that the input data in the shared memory of the present application can be released immediately after it is used up, without waiting for the matrix multiplication instruction to be completed, thereby separating the release of the shared memory space and the completion of the matrix multiplication instruction, so as to significantly improve the execution efficiency of the matrix multiplication instruction and the utilization rate of the shared memory space.

[0032] Figure 1 A schematic diagram of a computing device 100 implementing a method for synchronizing matrix multiplication instructions according to an embodiment of the present application is schematically shown. As Figure 1As shown, the computing device 100 can have one or more processing units, and a memory communicatively connected with the one or more processing units; wherein the memory stores instructions executable by the at least one processor. The processing units include special-purpose processing units such as a Graphics Processing Unit (GPU), a Field Programmable Gate Array (FPGA), and an Application Specific Integrated Circuit (ASIC), a General-purpose computing on graphics processing units (GPGPU), and general-purpose processing units such as a CPU. The computing device 100 further includes at least a data address sending module 110, a spatial information indicating module 120.

[0033] Regarding the data address sending module 110, it is configured to send, to the shared memory, an address of a data block required for each loop calculation, for the matrix multiplication operation unit to perform each loop calculation on matrix multiplication, when the matrix multiplication instruction management unit of the tensor core confirms that the information indicating that the input data in the shared memory is ready is received.

[0034] Regarding the spatial information indicating module 120, it is configured to send, to the data synchronization management unit, information indicating that the shared memory space has been released, when the matrix multiplication instruction management unit of the tensor core sends the address of the data block required for the last loop calculation.

[0035] Embodiments of the present application will be described below in conjunction with Figure 2 and Figure 3 a method 200 for synchronizing matrix multiplication instructions. Figure 2 A flowchart of the method 200 for synchronizing matrix multiplication instructions according to some embodiments of the present application is shown. Figure 3 A structural schematic diagram of an artificial intelligence chip 300 according to some embodiments of the present application is shown. It should be understood that the method 200 can be executed, for example, at the computing device 100 described above. Figure 1 The method 200 can further include additional actions not shown and / or can omit the actions shown, and the scope of the present application is not limited in this regard.

[0036] At step 202, the matrix multiplication instruction management unit of the tensor core of the computing device 100 sends, after confirming the reception of the information indicating that the input data in the shared memory is ready, the address of the data block required for each loop calculation to the shared memory for the matrix multiplication operation unit to perform each loop calculation on the matrix multiplication.

[0037] It should be understood that, Figure 3 The illustrated artificial intelligence chip is merely a specific implementation of the computing device 100, and the computing device 100 is not limited to Figure 3 The specific structure illustrated. The computing device 100 can also include Figure 3 Other components in addition to the artificial intelligence chip illustrated.

[0038] As Figure 3 illustrated, the matrix multiplication instruction management unit 322 of the tensor core 320 starts calculation and sends, after confirming the reception of the information from the data synchronization management unit 332 indicating that the input data in the shared memory 330 is ready, the address of the data block required for each loop calculation to the shared memory 330; and sends the address of the data block required for each loop calculation to the shared memory 330 at each loop to the shared memory 330 for the shared memory 330 to provide the data block corresponding to the address to the matrix multiplication operation unit 324 to perform each loop calculation.

[0039] Regarding the tensor core 320, it includes, for example, a matrix multiplication instruction management unit 322 and a matrix multiplication operation unit 324. Regarding the matrix multiplication operation unit 324, it is, for example, a general matrix multiplication (GEMM) calculation unit.

[0040] Regarding the shared memory 330, it is used to store at least the data block required for each loop calculation. In some embodiments, the shared memory 330 includes, for example, a data synchronization management unit 332 and a shared memory space 334.

[0041] Regarding the artificial intelligence chip 300, it also includes, for example, a scalar core 310. Regarding the scalar core 310, it includes, for example, an instruction queue 312, a register 314, and an instruction synchronization management unit 316.

[0042] In some embodiments, before step 202, the matrix multiplication instruction management unit 322 of the tensor core 320 obtains a matrix multiplication instruction (referred to as "MMA instruction" for short), which indicates at least: first shared memory barrier address information for synchronizing data readiness, second shared memory barrier address information for synchronizing memory space readiness, and instruction synchronization address information for synchronizing matrix multiplication instruction operation result output completion. As Figure 3As shown, the instruction queue 312 of the scalar core 310 sends an MMA instruction to the matrix multiplication instruction management unit 322 of the tensor core 320, which indicates a first shared memory barrier address (e.g., gsmb0), a second shared memory barrier address (e.g., gsmb1), and an instruction synchronization address (e.g., SC0). It should be understood that the scalar core 310 can also send other instructions to the matrix multiplication instruction management unit 322 through the instruction queue 312.

[0043] As to the method of sending information indicating that the input data in the shared memory 330 is ready to the matrix multiplication instruction management unit 322, it for example comprises: after receiving the matrix multiplication instruction, the matrix multiplication instruction management unit 322 sends the first shared memory barrier address information to the data synchronization management unit 332 of the shared memory 330; after receiving the first shared memory barrier address information, the data synchronization management unit 332 of the shared memory 330 confirms whether the data in the shared memory 330 corresponding to the first shared memory barrier address information is ready; in response to confirming that the data in the shared memory 330 corresponding to the first shared memory barrier address information is ready, the data synchronization management unit 332 of the shared memory 330 sends information indicating that the input data in the shared memory 330 is ready to the matrix multiplication instruction management unit 322; and in response to confirming that the data in the shared memory 330 corresponding to the first shared memory barrier address information is not ready, the data synchronization management unit 332 of the shared memory 330 sends information indicating that the input data in the shared memory 330 is not ready to the matrix multiplication instruction management unit 322. Details will be described below in conjunction with Figure 6 The method of sending information indicating that the input data in the shared memory is ready to the matrix multiplication instruction management unit will not be described again in detail here.

[0044] As to the method of sending the address of the data block required for each loop calculation to the shared memory for the matrix multiplication operation unit to perform each loop calculation of the matrix multiplication, it for example at least comprises: after confirming that the information indicating that the input data in the shared memory 330 is ready is received from the data synchronization management unit 332 of the shared memory 330, the matrix multiplication instruction management unit 322 of the tensor core 320 of the computing device 100 starts to calculate the address of the data block required for each loop calculation to be sent to the shared memory 330; and the address of the data block required for each loop calculation is sent to the shared memory in each loop for the shared memory 330 to provide the data block corresponding to the address to the matrix multiplication operation unit 324 of the tensor core 320 to perform each loop calculation.

[0045] At step 204, the matrix multiplication instruction management unit 322 of the computing device 100 synchronously sends information indicating that the shared memory space has been released to the data synchronization management unit 332 when the address of the data block required for the last loop calculation is sent.

[0046] For example, the matrix multiplication instruction management unit 322 of the tensor core 320 synchronously sends information indicating that the shared memory space 334 has been released to the data synchronization management unit 332 of the shared memory 330 when the address of the data block required for the last loop calculation is sent, for example, sends the second shared memory barrier address information (for example, shared memory barrier address gsmb1) to inform that the shared memory space 334 has been released, so that other data loading instructions can store data to the released shared memory space 334. Thus, the present application sends the address of the data block required for the last loop calculation, and the data synchronization management unit informs that the shared memory space has been released, so that other data loading instructions can load data to the released shared memory space without waiting for the MMA instruction to be executed, thereby greatly improving the utilization of the shared memory space.

[0047] In some embodiments, the matrix multiplication operation unit 324 of the tensor core 320 sends instruction synchronization address information (for example, instruction synchronization address SCO) for synchronizing the output of the matrix multiplication instruction operation result to the instruction synchronization management unit 316 in the scalar core 310 when the matrix multiplication instruction operation result is output to the register 314 of the scalar core 310. In some embodiments, after the instruction synchronization management unit 316 in the scalar core 310 receives the instruction synchronization address information (for example, instruction synchronization address SCO), the instruction queue 312 sends other scalar instructions to process the data output to the register 314 of the scalar core. Thus, the above-mentioned scheme sends the "instruction synchronization address information for synchronizing the output of the matrix multiplication instruction operation result" to the instruction synchronization management unit 316 in the scalar core 310 when the matrix multiplication operation unit outputs the matrix multiplication instruction operation result to the register of the scalar core, so that other instructions can be sent in advance, and thus the corresponding data loading can be performed in advance, realizing more efficient data processing.

[0048] In the above scheme, by causing the matrix multiplication instruction management unit of the tensor core to send the address of the data block required for each loop calculation to the shared memory for the matrix multiplication operation unit to perform each loop calculation on matrix multiplication after confirming that the information indicating that the input data in the shared memory is ready is received; and the matrix multiplication instruction management unit synchronously sends the information indicating that the shared memory space has been released to the data synchronization management unit when the address of the data block required for the last loop calculation is sent, so that the corresponding storage space can be released immediately after the input data in the shared memory of the application is used up, without waiting for the execution complete instruction to be executed, thereby separating the release of the shared memory space and the completion of the execution of the matrix multiplication instruction, so as to significantly improve the execution efficiency of the matrix multiplication instruction and the utilization rate of the shared memory space.

[0049] The structure of the artificial intelligence chip of the application and the process for synchronizing the matrix multiplication instruction of the application will be further described below in combination with Figure 3 Figure 3 The structure of the artificial intelligence chip of the application and the process for synchronizing the matrix multiplication instruction of the application will be further described below in combination with Figure 3 The data flow between the related components of the artificial intelligence chip 300 in the process of synchronizing the matrix multiplication instruction is also shown in

[0050] The structure of the artificial intelligence chip of the application and the process for synchronizing the matrix multiplication instruction of the application will be further described below in combination with Figure 3 The artificial intelligence chip 300 at least includes: a tensor core 320 and a shared memory 330.

[0051] Regarding the shared memory 330, it for example includes: a data synchronization management unit 332. The data synchronization management unit 332 is for example configured to send the information indicating that the input data in the shared memory 330 is ready to the matrix multiplication instruction management unit 322.

[0052] Regarding the shared memory 330, it for example also includes a shared memory space 334. Regarding the shared memory space 334, it is for example a static random access memory (SRAM), and the shared memory space 334 is configured to store the data required for matrix multiplication calculation.

[0053] Regarding the tensor core 320, it for example includes: a matrix multiplication instruction management unit 322 and a matrix multiplication operation unit 324.

[0054] Regarding the matrix multiplication operation unit 324, it is for example a general matrix multiplication (GEMM) calculation unit.

[0055] ​Regarding the matrix multiplication instruction management unit 322 of the tensor core 320, it is configured, for example, to send the address of the data block required for each loop calculation to the shared memory 330 after receiving the information indicating that the input data in the shared memory 330 is ready, and to send the information indicating that the shared memory space 334 has been released to the data synchronization management unit 332 of the shared memory 330 when the address of the data block required for the last loop calculation is sent, and to synchronize.

[0056] Regarding the matrix multiplication operation unit 324, it is configured to perform each loop calculation regarding matrix multiplication.

[0057] In some embodiments, the artificial intelligence chip 300 also includes, for example, a scalar core 310. Regarding the scalar core 310, it includes, for example, an instruction queue 312, a register 314, and an instruction synchronization management unit 316.

[0058] In some embodiments, the matrix multiplication operation unit 324 of the tensor core 320 is also configured to send instruction synchronization address information for synchronizing the output completion of the matrix multiplication instruction operation result to the instruction synchronization management unit 316 in the scalar core 310 while outputting the matrix multiplication instruction operation result to the register 314 of the scalar core 310. For example, the instruction synchronization management unit 316 of the scalar core 310 is configured to cause the instruction queue 312 to send other scalar instructions to process the data output to the register 314 of the scalar core 310 after receiving the instruction synchronization address information.

[0059] For example, as Figure 3As shown, in the process of executing the matrix multiplication instruction and its synchronization, the instruction queue 312 of the scalar core 310 sends the matrix multiplication instruction (referred to as "MMA instruction" for short) to the matrix multiplication instruction management unit 322 of the tensor core 320 (as indicated by the mark 410). After the matrix multiplication instruction management unit 322 of the tensor core 320 receives the MMA instruction, it needs to confirm whether the input data corresponding to the current MMA instruction in the shared memory 330 is ready. The matrix multiplication instruction management unit 322 of the tensor core 320 sends the first shared memory barrier address information (such as the first shared memory barrier address GSMB0) for synchronizing data readiness to the data synchronization management unit 332 of the shared memory 330 (as indicated by the mark 422); After the data synchronization management unit 332 of the shared memory 330 receives the first shared memory barrier address information, based on the received first shared memory barrier address information, it queries whether the data corresponding to the first shared memory barrier address information in the shared memory space 334 of the shared memory 330 is ready, if it is ready, the data synchronization management unit 332 of the shared memory 330 sends information indicating that the input data in the shared memory 330 is ready to the matrix multiplication instruction management unit 322 of the tensor core 320 (as indicated by the mark 424).

[0060] For example, if the data corresponding to the first shared memory barrier address GSMB0 is not ready, the data synchronization management unit 332 of the shared memory 330 sends information indicating that the input data in the shared memory 330 is not ready to the matrix multiplication instruction management unit 322 of the tensor core 320. It should be understood that for the case where the data is not ready, the data synchronization management unit 332 of the shared memory 330 can also represent that the data in the shared memory 330 is not ready by not sending information indicating that the input data in the shared memory 330 is ready.

[0061] For example, if the data corresponding to the first shared memory barrier address GSMB0 is not ready, after a predetermined time, the matrix multiplication instruction management unit 322 of the tensor core 320 re-sends the first shared memory barrier address information about the matrix multiplication instruction to the data synchronization management unit 332 of the shared memory 330 (for example, re-sending the first shared memory barrier address GSMB0).

[0062] By using the above means, the matrix multiplication instruction management unit of the tensor core confirms that the information indicating that the input data in the shared memory is ready is received, that is, the "data is ready" is confirmed, and then the data for executing the matrix multiplication instruction is obtained by the matrix multiplication operation unit, and if the data is not ready, the matrix multiplication instruction management unit processes other matrix multiplication instructions, so that the application can parallel multiple matrix multiplication instructions, and the utilization rate of the matrix multiplication operation unit is improved.

[0063] For example, after receiving the information indicating that the input data in the shared memory is ready, the matrix multiplication instruction management unit 322 of the tensor core 320 calculates the address of the data block required for each loop calculation, and sends the address of the data block required for each loop calculation to the shared memory space 334 of the shared memory 330 (as indicated by the mark 426), and when the address of the data block required for the last loop calculation is sent, the matrix multiplication instruction management unit 322 of the tensor core 320 synchronously sends the information indicating that the shared memory space 334 has been released to the data synchronization management unit 332 of the shared memory 330 (for example, indicating that the space corresponding to the first shared memory barrier address GSMB0 is released).

[0064] By using the above means, the input data in the shared memory of the application can be released immediately after being used up, without waiting for the matrix multiplication instruction to be executed, thereby separating the release of the shared memory space from the execution of the matrix multiplication instruction, so that the execution efficiency of the matrix multiplication instruction and the utilization rate of the shared memory space can be significantly improved.

[0065] For example, based on the data block required for loop calculation obtained from the shared memory space 334 (as indicated by the mark 430), the matrix multiplication operation unit 324 of the tensor core 320 performs loop calculation on matrix multiplication. After completing the calculation, the matrix multiplication operation unit 324 of the tensor core 320 sends the matrix multiplication instruction operation result to the register 314 of the scalar core 310 (as indicated by the mark 450); at the same time, the matrix multiplication operation unit 324 of the tensor core 320 sends the instruction synchronization address information indicating that the matrix multiplication instruction operation result output is completed to the instruction synchronization management unit 316 of the scalar core 310 (as indicated by the mark 460). The instruction synchronization address information is, for example, the instruction synchronization address SCO. For example, after receiving the message 460, the instruction synchronization management unit 316 of the scalar core 310 makes the instruction queue 312 send other scalar instructions to process the data output to the register 314 of the scalar core 310 (for example, the matrix multiplication instruction operation result).

[0066] Thus, the above scheme, while the matrix multiplication operation unit outputs the matrix multiplication instruction operation result to the register of the scalar core, also sends the "instruction synchronization address information for synchronizing the matrix multiplication instruction operation result output completion" to the instruction synchronization management unit in the scalar core, so that other instructions can be sent in advance and the corresponding data processing can be performed in advance.

[0067] The method for synchronizing matrix multiplication instructions of the present application is further described below in combination with Figure 3 , Figure 4 and Figure 5 Further, the method for synchronizing matrix multiplication instructions of the present application is further described below in combination with Figure 4 FIG. 1 shows a schematic diagram of synchronization between shared memory, tensor core and register according to some embodiments of the present application.

[0068] As shown in Figure 4 , during the execution of the matrix multiplication operation, four aspects of synchronization can be completed between the shared memory 330, the tensor core 320 and the register 314. Between the shared memory 330 and the tensor core 320, after the shared memory 330 loads the input data (e.g., input data A about the left matrix and input data B about the right matrix) for the matrix multiplication, the shared memory 330 synchronizes the information of "input data preparation completed" to the tensor core 320, as indicated by the mark 424 in Figure 4 . When the address of the data block required for the last loop calculation is sent, the tensor core 320 synchronizes the information of "shared memory space 334 has been released" to the shared memory 330, as indicated by the mark 428 in Figure 4 . The tensor core 320 performs calculation based on the data block required for the last loop calculation, so as to generate the matrix multiplication instruction operation result.

[0069] In addition, between the tensor core 320 and the register 314, before the tensor core 320 outputs the storage matrix multiplication instruction operation result, the register 314 synchronizes the information of "storage space of the register 314 is ready" to the tensor core 320, as indicated by the mark 452 in Figure 4 . When the matrix multiplication instruction operation result is output, the tensor core 320 synchronizes the information of "matrix multiplication instruction operation result output completion" to the register 314, as indicated by the mark 454 in Figure 4 . As shown in Figure 4As shown, the present application separates the synchronization about "shared memory space 334 has been released" indicated by mark 428 from the synchronization about "matrix multiplication instruction operation result output completion" indicated by mark 454. In contrast, the traditional method for synchronizing the matrix multiplication instruction makes the execution thread wait until the completion of the waiting matrix multiplication instruction by using the instruction "wgmma.wait_group". There is no distinction between the memory space release and the completion of the matrix multiplication instruction, resulting in a significant reduction in the execution efficiency of the matrix multiplication instruction and a waste of space in the shared memory. The present application introduces the new shared memory data synchronization mechanism and the instruction synchronization mechanism within the thread bundle described above, which separates the release of the shared memory space and the completion of the matrix multiplication instruction.

[0070] Figure 5 The action execution timing comparison diagram of the method of synchronizing the matrix multiplication instruction about some embodiments of the present application is schematically shown.

[0071] Figure 5 The upper half of the figure schematically shows the time sequence of some actions in the matrix multiplication instruction synchronization method of some embodiments of the present application. Figure 5 The lower half of the figure shows the time sequence of some actions in the traditional matrix multiplication instruction synchronization method.

[0072] As shown, the block indicated by mark 502 represents the time of "sending the matrix multiplication instruction". For example, at this time, the matrix multiplication instruction is sent to the matrix multiplication instruction management unit 322 of the tensor core 320. In the traditional matrix multiplication instruction synchronization method (as indicated by the lower half of the figure), the time of "sending the matrix multiplication instruction" needs to be after "loading the input data" and "confirming that the input data is ready". In comparison, as indicated by mark 502, the present application can "send the matrix multiplication instruction" at the same time as "loading the input data", so the matrix multiplication instruction can be sent earlier. Figure 5 Figure 5 The lower half of the figure shows the time sequence of some actions in the traditional matrix multiplication instruction synchronization method.

[0073] The block indicated by mark 504 represents the time of "loading the input data", for example, at this time, the shared memory 330 loads the input data related to the current matrix multiplication instruction (such as the input data A about the left matrix and the input data B about the right matrix) in the shared memory space 334.

[0074] The block indicated by mark 506 represents the time of "executing other scalar instructions", for example, at this time, the artificial intelligence chip 300 can execute other scalar instructions. It should be understood that the present application does not need to wait for "the matrix multiplication instruction operation result output completion" after "sending the matrix multiplication instruction", but can "execute other scalar instructions". As Figure 5 ​In the present embodiment, there is an idle time between the action of "sending the matrix multiplication calculation instruction" and the action of "matrix multiplication instruction operation result output completion", and other scalar instructions can be executed in the idle time. Thus, the present embodiment can realize more efficient data processing, which is beneficial to parallel execution of multiple matrix multiplication instructions.

[0075] The block indicated by the label 508 represents the time of "confirming that the input data is ready". In some embodiments, the block indicated by the label 508 represents the time of "confirming that the input data is ready" and "the storage space of the register 314 is ready". For example, at this time, the shared memory 330 confirms that the data in the shared memory 330 corresponding to the memory barrier address information is ready, and sends information indicating that the input data in the shared memory 330 is ready to the matrix multiplication instruction management unit 322. In addition, the register 314 synchronizes the information of "the storage space of the register is ready" to the tensor core 320.

[0076] The block indicated by the label 510 represents the time of "performing matrix multiplication calculation on the input data". For example, at this time, the matrix multiplication operation unit 324 of the tensor core 320 acquires the data block corresponding to the address of the data block required for each loop calculation to perform each loop calculation on the matrix multiplication. As Figure 5 As indicated in the upper half of the figure, since the present embodiment can send the matrix multiplication instruction earlier, the time of "performing matrix multiplication calculation on the input data" can also be advanced.

[0077] The block indicated by the label 512 represents the time of "releasing the shared memory space occupied by the input data". For example, at this time, the matrix multiplication instruction management unit 322 synchronizes the sending of information indicating that the shared memory space 334 has been released to the data synchronization management unit 332 when the address of the data block required for the last loop calculation is sent, for indicating the release of the shared memory space 334 occupied by the input data. It should be understood that, since the matrix multiplication instruction management unit 322 of the tensor core 320 sends information indicating that the shared memory space 334 has been released to the data synchronization management unit 332 of the shared memory 330 when the address of the data block required for the last loop calculation is sent, the shared memory 330 can not need to wait for the "output matrix multiplication instruction operation result" of the tensor core 320, i.e., can perform "release the shared memory space occupied by the input data" to further perform "load other input data". However, in the conventional matrix multiplication instruction synchronization method (such as the method shown in the lower half of the figure), the shared memory 330 needs to wait for the "output matrix multiplication instruction operation result" of the tensor core 320, and then performs "release the shared memory space occupied by the input data" to further perform "load other input data". Figure 5As indicated in the lower part, the time for "releasing the shared memory space occupied by the input data" needs to be after "outputting the result of the matrix multiplication instruction". In comparison, the present invention can release the shared memory space occupied by the input data of the current matrix multiplication instruction earlier during the execution of the current matrix multiplication instruction, and can load other input data (e.g., the input data of other matrix multiplication instructions) earlier, so as to realize the parallelism of multiple matrix multiplication instructions and improve the utilization of shared memory.

[0078] The box indicated by marker 514 represents the time when the matrix multiplication instruction operation result is output. For example, at this time, the matrix multiplication operation unit 324 of the tensor core 320 generates the matrix multiplication instruction operation result.

[0079] The box indicated by marker 516 represents the time when "other input data is loaded". For example, at this time, other input data (e.g., input data for other matrix multiplication instructions) is loaded for the already freed shared memory space 334. Thus, the present invention can achieve parallelism of multiple matrix multiplication instructions and improve the utilization of shared memory.

[0080] The box indicated by marker 518 represents the time when the matrix multiplication instruction operation result is completed. For example, at this time, the matrix multiplication operation unit 324 of the tensor core 320 outputs the matrix multiplication instruction operation result to the register 314 of the scalar core 310, and at the same time sends the instruction synchronization address information used to synchronize the completion of the matrix multiplication instruction operation result output to the instruction synchronization management unit 316 in the scalar core 310.

[0081] In summary, this solution can separate the release of shared memory space from the completion of matrix multiplication instructions, thereby significantly improving the execution efficiency of matrix multiplication instructions and the utilization rate of shared memory space.

[0082] The following will combine Figure 3 and Figure 6 A method 600 for sending information indicating that input data in shared memory is ready to be used to a matrix multiplication instruction management unit, according to an embodiment of the present invention. Figure 6 A flowchart illustrating a method 600 for sending information indicating that input data in shared memory is ready to be used to a matrix multiplication instruction management unit according to some embodiments of the present invention is shown. It should be understood that method 600 can, for example, be used in... Figure 1 The described computing device 100 performs the operation. Method 600 may also include additional actions not shown and / or the actions shown may be omitted; the scope of the invention is not limited in this respect.

[0083] At step 602, the matrix multiplication instruction management unit of the computing device 100 sends, to the data synchronization management unit of the shared memory, first shared memory barrier address information for synchronizing data ready after receiving the matrix multiplication instruction.

[0084] For example, the matrix multiplication instruction management unit 322 of the tensor core 320 sends, to the data synchronization management unit 332 of the shared memory 330, first shared memory barrier address information for synchronizing data ready after receiving the matrix multiplication instruction. The first shared memory barrier address information is, for example, first shared memory barrier address GSMB0.

[0085] At step 604, the data synchronization management unit of the shared memory of the computing device 100 confirms whether the data in the shared memory corresponding to the first shared memory barrier address information is ready after receiving the first shared memory barrier address information.

[0086] For example, as shown in FIG. 3, the data synchronization management unit 332 of the shared memory 330 confirms whether the data in the shared memory 330 corresponding to the first shared memory barrier address GSMB0 is ready after receiving the first shared memory barrier address GSMB0. Figure 3

[0087] At step 606, the computing device 100 sends, to the matrix multiplication instruction management unit, information for indicating that the input data in the shared memory is ready if the data in the shared memory corresponding to the first shared memory barrier address information is ready.

[0088] For example, the matrix multiplication instruction management unit 322 of the shared memory 330 sends, to the matrix multiplication instruction management unit 322, information for indicating that the input data in the shared memory 330 is ready if the data in the shared memory 330 corresponding to the first shared memory barrier address information is ready.

[0089] At step 608, the computing device 100 sends, to the matrix multiplication instruction management unit, information for indicating that the input data in the shared memory is not ready if the data in the shared memory corresponding to the first shared memory barrier address information is not ready.

[0090] For example, the matrix multiplication instruction management unit 322 of the shared memory 330 sends, to the matrix multiplication instruction management unit 322, information for indicating that the input data in the shared memory 330 is not ready if the data in the shared memory 330 corresponding to the first shared memory barrier address information is not ready. It should be understood that the data synchronization management unit 332 can also represent that the data in the shared memory 330 is not ready by not sending information for indicating that the input data in the shared memory 330 is ready.​

[0091] In some embodiments, the matrix multiplication instruction management unit 322 receives the information that the input data is not ready, and the matrix multiplication instruction management unit 322 can process other matrix multiplication instructions. After a predetermined time, the matrix multiplication instruction management unit 322 re-sends the first shared memory barrier address information for synchronizing the data readiness of the matrix multiplication instruction to the data synchronization management unit 332 of the shared memory 330.

[0092] By using the above-mentioned means, the matrix multiplication instruction management unit of the tensor core confirms that the information indicating that the input data in the shared memory is ready is received, that is, the "data is ready" is confirmed, and then the data for executing the matrix multiplication instruction is obtained by the matrix multiplication operation unit. If the data is not ready, the matrix multiplication instruction management unit processes other matrix multiplication instructions, so that the present application can parallel multiple matrix multiplication instructions, and the utilization rate of the matrix multiplication operation unit is improved.

[0093] In the above-mentioned scheme, the present application can quickly synchronize to the matrix multiplication instruction management unit after confirming that the shared memory data is ready, so that the matrix multiplication operation unit can efficiently obtain the data required for executing the MMA instruction.

[0094] The various processes and procedures described above, such as the method 200, 600, can be executed at a computing device. The computing device includes, for example, at least one processor (at least one graphics processor and at least one central processor), and a memory communicatively connected with 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. In some embodiments, the method 200, 600 can be implemented as a computer software program or program product, which is tangibly contained in a machine-readable medium. In some embodiments, part or all of the computer program can be loaded and / or installed on the computing device via a read-only memory (ROM) and / or a communication unit. When the computer program is loaded into a random access memory (RAM) and executed by a GPU and a CPU, one or more actions of the method 200, 600 described above can be performed.

[0095] The present application can be a method, an apparatus, a system, and / or a computer program product. The computer program product can include a computer readable storage medium (or media) having computer readable program instructions thereon for performing various aspects of the present application. The computer readable storage medium can be a tangible device that can retain and store instructions for execution by a processor. The computer readable storage medium can be, for example, but is not limited to, an electronic storage device, a magnetic storage device, an optical storage device, an electromagnetic storage device, a semiconductor storage device, or any suitable combination of the foregoing.

[0096] Computer readable program instructions described herein can be downloaded to respective computing / processing devices from a computer readable storage medium or to an external computer or external storage device via a network, for example, the Internet, a local area network, a wide area network and / or a wireless network. The computer readable program instructions can be embodied on one or more computer readable storage media, for example, a magnetic or optical disk, a magnetic tape, a semiconductor memory device, an application specific integrated circuit (ASIC) or a compact disc read only memory (CD-ROM). The computer readable program instructions described herein can be implemented in a variety of programming languages, including but not limited to LISP, PERL, C, C++, Java, Visual Basic, and JavaScript. The computer readable program instructions described herein can be implemented as a routine in a general purpose computer, a routine in a dedicated computer, or as a dedicated hardware component. The computer readable program instructions described herein can be implemented in a combination of one or more of the above.

[0097] These computer readable program instructions can be provided to a processor of a general purpose computer, special purpose computer, or other programmable data processing apparatus to produce a machine, such that the instructions, which execute via the processor of the computer or other programmable data processing apparatus, create means for implementing the functions / acts specified in the flowchart and / or block diagram block or blocks. These computer readable program instructions can also be stored in a computer readable storage medium that can be coupled to a computer or other programmable data processing apparatus, which can be used to implement the functions / acts specified in the flowchart and / or block diagram block or blocks.

[0098] The computer readable program instructions described herein can be implemented in a variety of programming languages, including but not limited to LISP, PERL, C, C++, Java, Visual Basic, and JavaScript. The computer readable program instructions described herein can be implemented as a routine in a general purpose computer, a routine in a dedicated computer, or as a dedicated hardware component. The computer readable program instructions described herein can be implemented in a combination of one or more of the above.

[0099] It should be understood that the various forms of flow shown above can be re-ordered, steps added or removed. For example, the steps loaded in the present application can be executed in parallel, can be executed sequentially, or can be executed in a different order, as long as the desired results of the technology disclosed in the present application can be achieved, which is not limited herein.

[0100] The above detailed description does not constitute a limitation on the scope of protection of the present application. Those skilled in the art should understand that various modifications, combinations, sub-combinations and substitutions can be made according to design requirements and other factors.

Claims

1. An artificial intelligence chip, characterized in that, The artificial intelligence chip includes: Shared memory, including a data synchronization management unit configured to send information indicating that input data in the shared memory is ready to a matrix multiplication instruction management unit; and The tensor kernel includes a matrix multiplication instruction management unit and a matrix multiplication operation unit. The matrix multiplication instruction management unit is configured to send the address of the data block required for each loop calculation to the shared memory after receiving information indicating that the input data in the shared memory is ready, and synchronously send information indicating that the shared memory space has been released to the data synchronization management unit when the address of the data block required for the last loop calculation is sent, so that the corresponding storage space in the shared memory is released immediately after the input data in the shared memory is used. The matrix multiplication operation unit is configured to perform each loop calculation of matrix multiplication.

2. The artificial intelligence chip according to claim 1, characterized in that, The matrix multiplication unit is also configured to: while outputting the result of the matrix multiplication instruction to the register of the scalar core, send the instruction synchronization address information indicating the completion of the matrix multiplication instruction output to the instruction synchronization management unit in the scalar core; The artificial intelligence chip also includes a scalar core, which includes an instruction synchronization management unit, registers, and an instruction queue.

3. The artificial intelligence chip according to claim 2, characterized in that, The instruction synchronization management unit is configured to, upon receiving instruction synchronization address information, cause the instruction queue to send other scalar instructions to process the data output to the registers of the scalar core.

4. A method for synchronizing matrix multiplication instructions, characterized in that, The method includes: After confirming receipt of information indicating that the input data in shared memory is ready, the matrix multiplication instruction management unit of the tensor kernel sends the addresses of the data blocks required for each loop calculation to the shared memory, so that the matrix multiplication unit can perform each loop calculation related to matrix multiplication; and When the address of the data block required for the last loop calculation is sent, the matrix multiplication instruction management unit synchronously sends information to the data synchronization management unit indicating that the shared memory space has been released, so that the corresponding storage space in the shared memory is released immediately after the input data in the shared memory is used up.

5. The method according to claim 4, characterized in that, Also includes: While outputting the result of the matrix multiplication instruction to the register of the scalar core, the matrix multiplication unit also sends the instruction synchronization address information indicating the completion of the matrix multiplication instruction output to the instruction synchronization management unit in the scalar core.

6. The method according to claim 4, characterized in that, Also includes: The matrix multiplication instruction management unit of the tensor kernel acquires the matrix multiplication instruction, which indicates at least: the first shared memory barrier address information for synchronizing input data readiness, the second shared memory barrier address information for synchronizing memory space readiness, and the instruction synchronization address information for synchronizing the completion of the matrix multiplication instruction operation result output.

7. The method according to claim 4, characterized in that, Also includes: After receiving the matrix multiplication instruction, the matrix multiplication instruction management unit sends the first shared memory barrier address information to the shared memory data synchronization management unit; After receiving the first shared memory barrier address information, the shared memory data synchronization management unit confirms whether the data in the shared memory corresponding to the first shared memory barrier address information is ready. In response to confirming that the data in the shared memory corresponding to the first shared memory barrier address information is ready, information indicating that the input data in the shared memory is ready is sent to the matrix multiplication instruction management unit. as well as In response to the confirmation that the data in the shared memory corresponding to the first shared memory barrier address information is not yet ready, information indicating that the input data in the shared memory is not yet ready is sent to the matrix multiplication instruction management unit.

8. The method according to claim 4, characterized in that, Sending the addresses of the data blocks required for each loop calculation to the shared memory for the matrix multiplication unit to perform each loop calculation related to matrix multiplication includes: After confirming that it has received information from the data synchronization management unit indicating that the input data in the shared memory is ready, the matrix multiplication instruction management unit of the tensor kernel begins to calculate the address of the data block required to send to the shared memory for each loop calculation. In each loop, the address of the data block required for each loop calculation is sent to the shared memory, so that the shared memory can provide the data block corresponding to the address to the matrix multiplication unit to perform the calculation for each loop.

9. The method according to claim 4, characterized in that, When the address of the data block required for the last loop calculation is sent, the matrix multiplication instruction management unit synchronously sends information to the data synchronization management unit indicating that the shared memory space has been released, thereby ensuring that the corresponding storage space in the shared memory is released immediately after the input data in the shared memory is used. When the address of the data block required for the last loop calculation is sent, the matrix multiplication instruction management unit synchronously sends the second shared memory barrier address information to the shared memory data synchronization management unit to inform that the shared memory space has been released, so that other data loading instructions can store data in the released shared memory space.

10. The method according to claim 5, characterized in that, Also includes: After receiving the instruction synchronization address information, the instruction synchronization management unit in the scalar core causes the instruction queue to send other scalar instructions to process the data output to the registers of the scalar core.

11. A computing device, characterized in that, include: At least one processor; as well as A memory that is communicatively connected to the at least one processor; in The memory stores instructions that can be executed by the at least one processor to enable the at least one processor to perform the method of any one of claims 4-10.

12. A computer-readable storage medium, characterized in that, The computer-readable storage medium stores a computer program that, when executed by a machine, performs the method according to any one of claims 4-10.

13. A computer program product, characterized in that, Includes a computer program, which, when executed by a machine, performs the method according to any one of claims 4-10.

Images