Systems, methods, computer devices, readable storage media, and program products for synchronization of tensor memory accelerator instructions

Abstract

The application relates to a synchronization system, method, computer device, readable storage medium and program product of a tensor memory accelerator instruction. The system comprises an execution unit and a tensor core, the tensor core is integrated with a synchronization control module and a tensor memory accelerator, the execution unit is used for sending a tensor memory accelerator instruction to the tensor core, the tensor memory accelerator instruction carries synchronization information, the synchronization information comprises a to-be-monitored state of a first target signal and a to-be-updated state of a second target signal after the tensor memory accelerator instruction is completed; and the synchronization control module is used for controlling state synchronization of the first target signal and state updating of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information. The system can effectively improve the efficiency of tensor operation.

Description

TECHNICAL FIELD

[0001] The present application relates to the technical field of artificial intelligence chips, in particular to a tensor memory accelerator instruction synchronization system and method, computer equipment, readable storage medium and program product. BACKGROUND

[0002] In the field of deep learning, large language models, etc., tensor operation becomes a core computing task, and the access synchronization mechanism of the tensor memory accelerator (TMA) directly affects the system performance as the key hardware for optimizing tensor data memory access.

[0003] The access synchronization of the existing TMA relies on software explicit control, and the execution unit (EU) needs to intervene in the synchronization logic throughout the process, and needs to insert instructions in the instruction stream of the execution unit to guarantee the operation sequence, such as inserting a waiting instruction to check the memory idle before the load / store instruction, inserting a state update instruction to mark the transmission completion after the operation is completed, and inserting an instruction to confirm that the data is ready before subsequent calculation.

[0004] However, the existing access synchronization method requires the execution unit to intervene in the synchronization logic throughout the process, which causes the execution unit to block when executing the synchronization instruction, and cannot process other tasks in parallel, resulting in idle computing resources and ultimately low efficiency of tensor operation. SUMMARY

[0005] Therefore, it is necessary to provide a tensor memory accelerator instruction synchronization system and method, computer equipment, readable storage medium and program product capable of improving the efficiency of tensor operation.

[0006] In a first aspect, the present application provides a tensor memory accelerator instruction synchronization system, comprising an execution unit and a tensor core, and the tensor core is integrated with a synchronization control module and a tensor memory accelerator, wherein:

[0007] The execution unit is configured to send a tensor memory accelerator instruction to the tensor core, and the tensor memory accelerator instruction carries synchronization information, and the synchronization information includes a to-be-monitored state of a first target signal and a to-be-updated state of a second target signal after the tensor memory accelerator instruction is completed.

[0008] The synchronization control module is configured to receive the tensor memory accelerator instruction, parse the synchronization information from the tensor memory accelerator instruction, and control state synchronization of the first target signal and state update of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information.

[0009] In one of the embodiments, the controlling the state synchronization of the first target signal and the state update of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information comprises:

[0010] monitoring the state of the first target signal, and controlling the tensor memory accelerator to execute the accelerator instruction when the state of the first target signal is monitored to be the to-be-monitored state;

[0011] updating the state of the second target signal to the to-be-updated state after the tensor memory accelerator completes the accelerator instruction.

[0012] In one of the embodiments, the system further comprises a synchronization state storage module configured to store the state of the first target signal and the state of the second target signal.

[0013] In one of the embodiments, the monitoring the state of the first target signal and controlling the tensor memory accelerator to execute the accelerator instruction when the state of the first target signal is monitored to be the to-be-monitored state comprises:

[0014] The synchronization control module is configured to send a monitoring signal to the synchronization state storage module, the monitoring signal carrying the to-be-monitored state of the first target signal, and the monitoring signal being used to instruct the synchronization state storage module to monitor the state of the first target signal.

[0015] The synchronization state storage module is configured to feed back a response signal to the synchronization control module in response to the monitoring signal when the state of the first target signal is the to-be-monitored state.

[0016] The synchronization control module is further configured to send an execution notification to the tensor memory accelerator in response to the response signal, and the execution notification is used to instruct the tensor memory accelerator to execute the tensor memory accelerator instruction.

[0017] In one of the embodiments, the updating the state of the second target signal to the to-be-updated state after the tensor memory accelerator completes the accelerator instruction comprises:

[0018] The synchronization control module is further configured to send an update signal to the synchronization state storage module in response to the instruction completion notification fed back by the tensor memory accelerator, the update signal carrying the to-be-updated state of the second target signal and used to instruct the synchronization state storage module to update the state of the second target signal.

[0019] The synchronization state storage module is further configured to update the state of the second target signal to the to-be-updated state in response to the update signal.

[0020] In one of the embodiments, in the case where the tensor memory accelerator instruction is a load instruction, the to-be-monitored state of the first target signal is an idle state of a storage space, and the to-be-updated state of the second target signal is a ready state of operation data.

[0021] In the case where the tensor memory accelerator instruction is a store instruction, the to-be-monitored state of the first target signal is a ready state of operation data, and the to-be-updated state of the second target signal is an idle state of a storage space.

[0022] In a second aspect, the application further provides a synchronization method of a tensor memory accelerator instruction, applied to a synchronization control module integrated in a tensor core, the tensor core further deploying a tensor memory accelerator, and the method comprising:

[0023] receiving a tensor memory accelerator instruction sent by an execution unit, the tensor memory accelerator instruction carrying synchronization information, the synchronization information including a to-be-monitored state of a first target signal and a to-be-updated state of a second target signal after the tensor memory accelerator instruction is completed;

[0024] parsing the synchronization information from the tensor memory accelerator instruction, and based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information, controlling state synchronization of the first target signal and state updating of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction.

[0025] In a third aspect, the application further provides a synchronization device of a tensor memory accelerator instruction, applied to a synchronization control module integrated in a tensor core, the tensor core further deploying a tensor memory accelerator, and the device comprising:

[0026] The receiving module is configured to receive a tensor memory accelerator instruction sent by an execution unit, the tensor memory accelerator instruction carrying synchronization information, the synchronization information including a to-be-monitored state of a first target signal and a to-be-updated state of a second target signal after the tensor memory accelerator instruction is completed.

[0027] The parsing module is configured to parse the synchronization information from the tensor memory accelerator instruction, and control state synchronization of the first target signal and state updating of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information.

[0028] In a fourth aspect, the present application further provides a computer device, including a memory and a processor, the memory stores a computer program, and the processor implements the steps of the above method when executing the computer program.

[0029] In a fifth aspect, the present application further provides a computer readable storage medium, which stores a computer program, and the computer program implements the steps of the above method when executed by a processor.

[0030] In a sixth aspect, the present application further provides a computer program product, which includes a computer program, and the computer program implements the steps of the above method when executed by a processor.

[0031] The aforementioned synchronization system, method, computer device, readable storage medium, and program product for tensor memory accelerator instructions include a synchronization system comprising an execution unit and a tensor core. The tensor core integrates a synchronization control module and a tensor memory accelerator. The execution unit can send tensor memory accelerator instructions to the tensor core. The tensor memory accelerator instructions carry synchronization information, including the state to be monitored for a first target signal and the state to be updated for a second target signal after the tensor memory accelerator instructions are executed. The synchronization control module can receive the tensor memory accelerator instructions, parse the synchronization information from the instructions, and, based on the state to be monitored for the first target signal and the state to be updated for the second target signal in the synchronization information, control the synchronization of the state of the first target signal and the updating of the state of the second target signal before and after the tensor memory accelerator executes the instructions. The tensor memory accelerator instruction synchronization system, method, computer device, readable storage medium, and program product provided in this application automatically manage data synchronization and updates through a synchronization control module deployed in the tensor core. This decouples the synchronization logic from the execution unit, thereby eliminating the resource occupation of the execution unit by the execution process of the tensor memory accelerator instructions. This can eliminate the blocking of the execution unit, reduce instruction overhead, and significantly improve the efficiency of tensor operations. At the same time, it eliminates the need to insert instructions for state synchronization and updates in the execution unit, thereby reducing programming complexity. Attached Figure Description

[0032] To more clearly illustrate the technical solutions in the embodiments of this application or related technologies, the drawings used in the description of the embodiments of this application or related technologies will be briefly introduced below. Obviously, the drawings described below are only some embodiments of this application. For those skilled in the art, other related drawings can be obtained based on these drawings without creative effort.

[0033] Figure 1 This is a block diagram of a synchronization system for tensor memory accelerator instructions in one embodiment;

[0034] Figure 2 This is a flowchart illustrating a method for synchronizing instructions in a tensor memory accelerator in one embodiment.

[0035] Figure 3 This is a schematic diagram illustrating the execution of instructions in one embodiment;

[0036] Figure 4 This is a schematic diagram illustrating instruction execution in another embodiment;

[0037] Figure 5 This is a block diagram of a synchronization device for tensor memory accelerator instructions in one embodiment;

[0038] Figure 6 Figure 1 is a diagram of an internal structure of a computer device in one embodiment. DETAILED DESCRIPTION

[0039] For the purposes of the present application, the technical solutions and advantages thereof are more clearly apparent, the following will be further described in detail in conjunction with the drawings and embodiments. It should be understood that the specific embodiments described herein are only used to explain the present application, and are not intended to limit the present application.

[0040] It should be noted that the terms "first", "second" and the like used herein can be used to describe various elements, but these elements are not limited by these terms. These terms are only used to distinguish the first element from the second element. The terms "include" and "have" and any variations thereof used herein are intended to cover non-exclusive inclusion. The term "multiple" used herein refers to two or more. The term "and / or" used herein refers to one of the options or any combination of multiple options.

[0041] In one exemplary embodiment, as shown in Figure 1 A synchronization system 100 of a tensor memory accelerator instruction is provided, the system includes an execution unit 102 and a tensor core 104, the tensor core 104 is integrated with a synchronization control module 106 and a tensor memory accelerator 108, wherein:

[0042] The execution unit 102 is configured to send a tensor memory accelerator instruction to the tensor core, the tensor memory accelerator instruction carries synchronization information, and the synchronization information includes a to-be-monitored state of a first target signal and a to-be-updated state of a second target signal after the tensor memory accelerator instruction is completed.

[0043] The synchronization control module 106 is configured to receive the tensor memory accelerator instruction, parse the synchronization information from the tensor memory accelerator instruction, and based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information, control the state synchronization of the first target signal and the state update of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction.

[0044] In the embodiment of the present application, the execution unit 102 is the instruction initiator of the synchronization system, responsible for sending tensor memory accelerator instructions such as load instructions, store instructions, etc. to the tensor core (also represented as Tensor Core) 104. Among them, the tensor memory accelerator instruction not only contains data transmission information, but also carries complete synchronization information, including the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal. Among them, the to-be-monitored state of the first target signal is the precondition that needs to be met before the tensor memory accelerator executes, such as the idle state of the target shared memory block. The to-be-updated state of the second target signal is the post-state that needs to be set after the tensor memory accelerator executes, such as the ready state of the data loading completion.

[0045] For example, when the execution unit 102 needs to trigger an operation of loading tensor data from the global memory to the shared memory once, the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal will be carried in the sent tensor memory accelerator instruction. Among them, the to-be-monitored state of the first target signal is to monitor the state of the shared memory address 0x1000 to be idle. The to-be-updated state of the second target signal is to update the state of the address 0x2000 to be ready after loading is completed. After sending the instruction, the execution unit 102 does not need to wait for the synchronization result, and can immediately execute other data-independent computing tasks, such as executing another group of tensor operation instructions.

[0046] The synchronization control module 106 in the tensor core 104 will first extract the synchronization information in the tensor memory accelerator instruction received by the execution unit 102, and parse out the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal.

[0047] The synchronization control module 106 will perform the synchronization of the pre-state and the update of the post-state according to the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal before and after the execution of the tensor memory accelerator instruction, that is, the state synchronization of the first target signal and the state update of the second target signal. In an exemplary embodiment, the above process includes:

[0048] The state of the first target signal is monitored, and when the state of the first target signal is monitored to be in the to-be-monitored state, the tensor memory accelerator is controlled to execute the accelerator instruction;

[0049] After the tensor memory accelerator completes the accelerator instruction, the state of the second target signal is updated to the to-be-updated state.

[0050] Before the tensor memory accelerator 108 executes the data transmission, the synchronization control module 106 will monitor the current state of the first target signal in real time, and when the state of the first target signal is monitored to be in the to-be-monitored state, the tensor memory accelerator is controlled to execute the tensor memory accelerator instruction.

[0051] The synchronization control module 106 is a function module implemented by a hardware circuit. Assuming that the tensor memory accelerator instruction is a load instruction, the synchronization control module 106 parses the to-be-monitored state of the first target signal from the tensor memory accelerator instruction as the state of the shared memory block 0x001 being an idle state, i.e., the state being binary 1, indicating that the memory block is not occupied. Subsequently, the hardware monitoring circuit in the synchronization control module 106 periodically reads the current state of the first target signal and compares it with the to-be-monitored state in real time. When the current state of the first target signal is consistent with the to-be-monitored state, such as the state of the shared memory block being 1 (i.e., the idle state), the synchronization control module 106 immediately sends a signal to the tensor memory accelerator 108 to trigger it to execute a specific accelerator instruction.

[0052] After the tensor memory accelerator 108 completes execution of the corresponding instruction, the synchronization control module 106 updates the state of the second target signal to the to-be-updated state specified in the instruction, such as changing the value of the data ready signal register from 0 to 1 after the load operation is completed, to provide a clear available flag for the subsequent calculation unit. For example, when the tensor memory accelerator completes the load operation of a 1024x1024 tensor, the synchronization control module 106 automatically changes the data ready flag from 0 to 1 (i.e., ready), and at this time the tensor core can directly read the flag without waiting for the notification of the execution unit to start the calculation.

[0053] In the synchronization system of the tensor memory accelerator instruction described above, the synchronization system includes an execution unit and a tensor core, and the tensor core integrates a synchronization control module and a tensor memory accelerator. The execution unit can send a tensor memory accelerator instruction to the tensor core, and the tensor memory accelerator instruction carries synchronization information, which includes a to-be-monitored state of a first target signal and a to-be-updated state of a second target signal after the tensor memory accelerator instruction is completed. The synchronization control module can receive the tensor memory accelerator instruction, parse the synchronization information from the tensor memory accelerator instruction, and control the state synchronization of the first target signal and the state update of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information. The synchronization system of the tensor memory accelerator instruction provided in the embodiments of the present application automatically manages data synchronization and update by deploying the synchronization control module in the tensor core, decouples the synchronization logic from the execution unit, thereby releasing the resource occupation of the execution unit in the execution process of the tensor memory accelerator instruction, so as to eliminate the blocking of the execution unit, reduce the instruction overhead, significantly improve the efficiency of tensor operation, and at the same time, without inserting instructions for state synchronization and update in the execution unit, the programming complexity is also reduced.

[0054] In an example embodiment, referring to Figure 1 The system further includes a synchronization state storage module 110, which is configured to store the state of the first target signal and the state of the second target signal.

[0055] In an example embodiment, the synchronization state storage module 110 is a hardware storage unit for storing the state of the synchronization signal in the system, and stores the mark of the free and ready state of the data space. The module can be a separate module, or can be integrated into the shared memory, the execution unit, or other modules, and provides data support for the automatic synchronization logic of the synchronization control module 106. For example, the synchronization state storage module 110 can be used to store the real-time state of the first target signal and the second target signal.

[0056] For example, the synchronization state storage module 110 directly interacts with the synchronization control module 106 through a hardware connection. When the synchronization control module 106 listens to the to-be-monitored state of the first target signal, it will access the storage address corresponding to the first target signal in the synchronization state storage module 110, read the current state, and compare it with the to-be-monitored state. After the tensor memory accelerator 108 completes the instruction execution, the synchronization control module 106 will send a state update signal to the synchronization state storage module 110, and rewrite the value of the storage address corresponding to the second target signal to the to-be-updated state. Subsequent tensor cores can directly read the state value in the synchronization state storage module 110 without the need to forward through the execution unit 102, thereby achieving hardware-level sharing of state information.

[0057] For example, in a load operation, after the synchronization control module 106 reads the state of the shared memory block A from the address 0x0001 of the synchronization state storage module 110 as the free state, it triggers the tensor memory accelerator to execute. After the tensor memory accelerator completes the data loading, the synchronization control module 106 immediately updates the state of the address 0x0002 in the synchronization state storage module 110 from 0 (in transmission) to 1 (ready). Subsequent tensor cores can directly read the state of this address and start the calculation after confirming that the data is available.

[0058] In an example embodiment, the state of the first target signal is monitored, and when the state of the first target signal is monitored, the tensor memory accelerator is controlled to execute the accelerator instruction, including:

[0059] The synchronization control module 106 is configured to send a listening signal to the synchronization state storage module, and the listening signal carries the to-be-monitored state of the first target signal, which is used to instruct the synchronization state storage module to listen to the state of the first target signal.

[0060] The synchronization state storage module 110 is configured to, in response to the monitoring signal, feed back a response signal to the synchronization control module when the state of the first target signal is the to-be-monitored state.

[0061] The synchronization control module 106 is further configured to, in response to the response signal, send an execution notification to the tensor memory accelerator, where the execution notification is used to instruct the tensor memory accelerator to execute the tensor memory accelerator instruction.

[0062] In the embodiment, the synchronization control module 106 and the synchronization state storage module 110 interact through a hardware signal, to realize the monitoring of the state of the first target signal and the triggered execution of the tensor memory accelerator instruction. The synchronization control module 106, as the initiator of the monitoring logic, first sends a monitoring signal to the synchronization state storage module 110. The signal is not a software instruction, but an electrical signal transmitted through a hardware connection, which carries the to-be-monitored state of the first target signal.

[0063] After receiving the monitoring signal, the synchronization state storage module 110 starts real-time monitoring. When the state of the first target signal is the to-be-monitored state, the synchronization state storage module 110 sends a response signal to the synchronization control module 106, to inform that the target state has been met. After receiving the response signal fed back by the synchronization state storage module 110, the synchronization control module 106 directly sends an execution notification to the tensor memory accelerator 108. The notification carries the core information of the tensor memory accelerator instruction, such as the source address, the target address, the data length, etc., and explicitly instructs the tensor memory accelerator 108 to start a specific data transmission operation, such as a load operation or a store operation.

[0064] For example, if the tensor memory accelerator instruction is to load a tensor from the global memory to the shared memory block B, the first target signal is the occupancy state of the shared memory block B, and the to-be-monitored state is 0 (unoccupied). At this time, the monitoring signal sent by the synchronization control module 106 carries the address 0x0002 (the state address of the shared memory block B) and the to-be-monitored state 0, to instruct the synchronization state storage module 110 to perform real-time monitoring on the state change of the address. After receiving the monitoring signal, the synchronization state storage module 110 starts real-time monitoring through the built-in state monitoring circuit, such as reading the current state of the storage address 0x0002 corresponding to the first target signal at a period of 1 ns. When the state of the shared memory block B changes from 1 (occupied) to 0 (unoccupied), the synchronization state storage module 110 immediately sends a response signal, such as a high-level pulse, to the synchronization control module 106 through the hardware connection, to confirm that the shared memory block B is unoccupied and the load operation can be performed. After receiving the response signal, the synchronization control module 106 sends an execution notification to the tensor memory accelerator 108, and the tensor memory accelerator 108 will immediately perform the load operation according to the instruction parameters.

[0065] In an exemplary embodiment, after the tensor memory accelerator completes the accelerator instruction, the state of the second target signal is updated to a to-be-updated state, including:

[0066] The synchronization control module 106 is further configured to send an update signal to the synchronization state storage module in response to the instruction completion notification fed back by the tensor memory accelerator, the update signal carrying the to-be-updated state of the second target signal, for instructing the synchronization state storage module to update the state of the second target signal.

[0067] The synchronization state storage module 110 is further configured to update the state of the second target signal to the to-be-updated state in response to the update signal.

[0068] In the embodiment, when the tensor memory accelerator 108 completes all data transmission according to the instruction (for example, the tensor data is successfully moved to the shared memory by the load operation or the data is written into the global memory by the store operation), the tensor memory accelerator 108 sends an instruction completion notification to the synchronization control module 106 through internal hardware connection. After receiving the completion notification, the synchronization control module 106 immediately sends an update signal to the synchronization state storage module 110. The update signal carries the to-be-updated state of the second target signal. After receiving the update signal, the synchronization state storage module 110 directly executes state update through the built-in hardware write circuit. The state update is a hardware-level direct write, and the time consumption is in nanoseconds. During the write process, the address is locked to avoid state conflicts caused by simultaneous operations of other modules.

[0069] For example, the tensor memory accelerator executes a store instruction for writing the 128x128 tensor data after calculation into the global memory, and the second target signal is the occupancy state of the global memory region, and the to-be-updated state is 0 (idle). At this time, the update signal sent by the synchronization control module 106 carries the address 0x0003 (the state address of the memory region) and the to-be-updated state 0 (idle), for instructing the synchronization state storage module 110 to execute state update. After receiving the update signal, the synchronization state storage module 110 changes the state of the address 0x0003 from 1 (occupied) to 0 (idle). Subsequently, the tensor core depending on the data can directly read the state of the address, and start calculation immediately after confirming that the data is available, without waiting for the forwarding notification of the execution unit.

[0070] In an exemplary embodiment, in the case where the tensor memory accelerator instruction is a load instruction, the to-be-monitored state of the first target signal is the idle state of the storage space, and the to-be-updated state of the second target signal is the ready state of the operation data.

[0071] In the case where the tensor memory accelerator instruction is a store instruction, the to-be-monitored state of the first target signal is the ready state of the operation data, and the to-be-updated state of the second target signal is the idle state of the storage space.

[0072] In the embodiments of the present application, the system differentiates the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal according to the specific type of the tensor memory accelerator instruction, load or store, to ensure the orderly cooperation of data transmission and storage space.

[0073] For example, when the execution unit 102 sends a load instruction (such as transferring tensor data in the global memory to the shared memory for use by the tensor core and other computing units), the to-be-monitored state of the first target signal is the idle state of the storage space, such as the unoccupied state of the target shared memory block, which is coded as binary 1. This means that the synchronization control module 106 will first listen to the state of the target shared memory block, and only when it is confirmed that the space is not occupied by other tasks, will it trigger the tensor memory accelerator to perform the load operation, avoiding data overwrite or transmission conflict. When the tensor memory accelerator completes the data loading, the synchronization control module 106 sends an update signal to the synchronization state storage module 110, changing the ready state of the corresponding tensor data, i.e., the second target signal, from 0 (in transmission) to 1 (ready). The subsequent tensor core can directly read this state and start the calculation as soon as the data is available, without additional notification from the execution unit 102.

[0074] For example, when the execution unit 102 sends a store instruction (such as writing the calculated tensor data in the shared memory to the global memory for permanent storage), the to-be-monitored state of the first target signal is the ready state of the operation data, such as the calculated completed state of the tensor to be stored in the shared memory, which is coded as binary 1. This means that the synchronization control module 106 will first listen to the state of the data to be stored, and only when it is confirmed that the data has been calculated, will it trigger the tensor memory accelerator to perform the storage operation, avoiding writing dirty data that is not completed. When the tensor memory accelerator completes the data calculation and writes it to the global memory, the synchronization control module 106 sends an update signal to the synchronization state storage module 110, changing the occupied state of the corresponding shared memory block, i.e., the second target signal, from 1 (occupied) to 0 (idle), releasing the space for use by other tasks.

[0075] Referring to Figure 2 As shown in the figure, a synchronization method of a tensor memory accelerator instruction is provided, which is applied to a synchronization control module integrated in a tensor core, and a tensor memory accelerator is also deployed in the tensor core. In the embodiments, the method includes the following steps 202 to 204, wherein:

[0076] Step 202, receiving the tensor memory accelerator instruction sent by the execution unit, the tensor memory accelerator instruction carrying synchronization information, the synchronization information including the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal after the completion of the tensor memory accelerator instruction;

[0077] In step 204, the synchronization information is parsed from the tensor memory accelerator instruction, and based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information, the state synchronization of the first target signal and the state update of the second target signal are controlled before and after the tensor memory accelerator executes the tensor memory accelerator instruction.

[0078] In the embodiments of the present application, the execution process of the synchronization method of the tensor memory accelerator instruction can refer to the related description of the foregoing embodiments of the synchronization system of the tensor memory accelerator instruction, which will not be repeated here.

[0079] The synchronization method of the tensor memory accelerator instruction provided in the embodiments of the present application is that the synchronization control module integrated in the tensor core receives the tensor memory accelerator instruction sent by the execution unit, the tensor memory accelerator instruction carries synchronization information, the synchronization information includes the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal after the tensor memory accelerator instruction is completed, the synchronization information is parsed from the tensor memory accelerator instruction, and based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information, the state synchronization of the first target signal and the state update of the second target signal are controlled before and after the tensor memory accelerator executes the tensor memory accelerator instruction. By using the synchronization method of the tensor memory accelerator instruction provided in the embodiments of the present application, the synchronization control module deployed in the tensor core automatically manages data synchronization and update, the synchronization logic is decoupled from the execution unit, and then the resource occupation of the execution unit in the execution process of the tensor memory accelerator instruction is released, so that the blocking of the execution unit can be eliminated, the instruction overhead is reduced, the efficiency of the tensor operation is significantly improved, and at the same time, the programming complexity can be reduced without inserting instructions for state synchronization and update in the execution unit.

[0080] In order for those skilled in the art to better understand the embodiments of the present application, the embodiments of the present application are described below through specific examples.

[0081] In one example, the application automatically manages data synchronization through a hardware state machine (i.e., a synchronization control module), decouples synchronization logic from the execution unit, polls a synchronization signal state (idle / ready) through a built-in module of a tensor memory accelerator instruction, implements hardware signal monitoring, automatically completes waiting and state updating before and after the execution of a load / store instruction, and enables the execution unit to continue processing other tasks during the execution of the tensor memory accelerator instruction, thereby achieving efficient parallelism. The synchronization signal used can employ synchronization strategies such as mbarrier (Memory Barrier), tgbar (Thread Group Barrier), and syncchannel (Synchronization Channel). The key hardware components mainly include a synchronization state storage module and a synchronization control module. Taking mbarrier as an example, the synchronization among the execution unit, Tensor Core (which integrates the synchronization control module and the tensor memory accelerator), and the synchronization state storage module is shown in FIG. 7. The execution unit only needs to send a tensor memory accelerator instruction to the Tensor Core and does not need to participate in synchronization. The Tensor Core (or the synchronization control module therein) initiates a query state and waits operation to the mbarrier (i.e., the synchronization state storage module) actively. When the synchronization condition of the mbarrier is met, the Tensor Core automatically continues to execute subsequent operations, and the execution unit can execute other independent tasks at the same time, thereby avoiding blocking. Figure 3

[0082] For example, as shown in FIG. 7, taking mbarrier as an example, under the automatic synchronization mechanism of hardware, the cooperation process of the execution unit and the tensor core regarding mbarrier synchronization is as follows: Figure 4

[0083] 1. The execution unit sends a tensor memory accelerator instruction to the tensor core. The instruction includes mbarrier (memory barrier) information (i.e., the state to be monitored of the first target signal) that needs to be waited for and mbarrier information (i.e., the state to be updated of the second target signal) that needs to be updated. Then, the execution unit continues to execute other subsequent instructions without data dependency.

[0084] 2. The synchronization control module in the tensor core receives the instruction, checks the mbarrier information that needs to be waited for, and polls the mbarrier ready (i.e., waits for the state of the first target signal to be updated to the state to be monitored).

[0085] 3. After the mbarrier is ready, i.e., the state is met, the response signal is fed back to the synchronization control module in the tensor core.

[0086] ​​4. The tensor memory accelerator starts to execute the load or store operation.

[0087] 5. The data processing is completed.

[0088] 6. An mbarrier signal is sent (i.e., an update signal including the to-be-updated state of the second target signal is sent), and mbarrier state information that needs to be updated is set.

[0089] The synchronization method of the tensor memory accelerator instruction provided in the embodiments of the present application automatically manages data synchronization through a hardware state machine, the synchronization waiting is handled by hardware, the execution unit can execute other instructions in parallel, thereby eliminating the blocking of the execution unit, the synchronization state update is automatically completed by hardware, the execution unit cycle is saved, the instruction overhead is reduced, the developer does not need to write tensor memory accelerator synchronization code, only needs to call the tensor memory accelerator instruction, the programming complexity is reduced, and the efficiency of tensor operation is significantly improved.

[0090] It should be understood that, although each step in the flowchart involved in each embodiment as described above is displayed in sequence according to the arrow, these steps are not necessarily executed in sequence according to the arrow. Unless otherwise specified herein, the execution of these steps is not strictly limited in sequence, and these steps can be executed in other sequences. Moreover, at least part of the steps in the flowchart involved in each embodiment as described above can include multiple steps or multiple stages, which are not necessarily executed at the same time, but can be executed at different times, and the execution sequence of these steps or stages is not necessarily sequential, but can be executed in rotation or alternation with at least part of other steps or steps or stages in other steps. It can be understood that the steps in different embodiments can be freely combined as needed, and various non-contradictory schemes formed by the combination all belong to the scope of protection of the present application.

[0091] Based on the same inventive concept, the embodiments of the present application also provide a synchronization device of a tensor memory accelerator instruction for implementing the synchronization method of the tensor memory accelerator instruction involved above. The implementation scheme for solving the problem provided by the device is similar to the implementation scheme described in the above method, and therefore the specific limitations in one or more synchronization device embodiments of the tensor memory accelerator instruction provided below can refer to the limitations of the synchronization method of the tensor memory accelerator instruction described above, which will not be repeated here.

[0092] In one exemplary embodiment, as shown in Figure 5 A synchronization device 500 of a tensor memory accelerator instruction is provided, which is applied to a synchronization control module integrated in a tensor core, and the tensor core also has a tensor memory accelerator deployed therein, and includes a receiving module 502 and an analysis module 504, wherein:

[0093] The receiving module 502 is configured to receive a tensor memory accelerator instruction sent by an execution unit, and the tensor memory accelerator instruction carries synchronization information, and the synchronization information includes a to-be-monitored state of a first target signal and a to-be-updated state of a second target signal after the tensor memory accelerator instruction is completed.

[0094] The parsing module 504 is configured to parse the synchronization information from the tensor memory accelerator instruction, and based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information, control state synchronization of the first target signal and state updating of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction.

[0095] The synchronization device for the tensor memory accelerator instruction provided by the embodiment of the present application is configured to receive, by a synchronization control module integrated in a tensor core, a tensor memory accelerator instruction sent by an execution unit, and the tensor memory accelerator instruction carries synchronization information, and the synchronization information includes a to-be-monitored state of a first target signal and a to-be-updated state of a second target signal after the tensor memory accelerator instruction is completed, parse the synchronization information from the tensor memory accelerator instruction, and based on the to-be-monitored state of the first target signal and the to-be-updated state of the second target signal in the synchronization information, control state synchronization of the first target signal and state updating of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction. By using the synchronization device for the tensor memory accelerator instruction provided by the embodiment of the present application, the synchronization control module deployed in the tensor core automatically manages data synchronization and updating, the synchronization logic is decoupled from the execution unit, and then the resource occupation of the execution unit in the execution process of the tensor memory accelerator instruction is released, so that the blocking of the execution unit can be eliminated, the instruction overhead is reduced, the efficiency of the tensor operation is significantly improved, and meanwhile, the programming complexity can be reduced without inserting instructions for state synchronization and updating in the execution unit.

[0096] The modules in the synchronization device for the tensor memory accelerator instruction described above can be all or partially implemented by software, hardware, or a combination thereof. The modules described above can be embedded in or independent of a processor in a computer device in a hardware form, or can be stored in a memory in the computer device in a software form, so as to be called and executed by a processor to perform the operations corresponding to the modules.

[0097] In one exemplary embodiment, a computer device, which can be a terminal, is provided, and an internal structure diagram of the computer device can be as shown in Figure 6The computer device shown in the figure includes a processor, a memory, an input / output interface, a communication interface, a display unit and an input device. Among them, the processor, the memory and the input / output interface are connected through a system bus, and the communication interface, the display unit and the input device are connected to the system bus through the input / output interface. Among them, the processor of the computer device is used to provide computing and control capability. The memory of the computer device includes a non-volatile storage medium and an internal memory. The non-volatile storage medium stores an operating system and a computer program. The internal memory provides an environment for the operating system and the computer program in the non-volatile storage medium to run. The input / output interface of the computer device is used to exchange information between the processor and external devices. The communication interface of the computer device is used to communicate with external terminals in a wired or wireless manner. The wireless manner can be realized through WIFI, mobile cellular network, Near Field Communication (NFC) or other technologies. The computer program is executed by the processor to realize a synchronization method of a tensor memory accelerator instruction. The display unit of the computer device is used to form a visually visible picture, which can be a display screen, a projection device or a virtual reality imaging device. The display screen can be a liquid crystal display screen or an electronic ink display screen. The input device of the computer device can be a touch layer overlaid on the display screen, or a key, trackball or touchpad arranged on the shell of the computer device, or an external keyboard, touchpad or mouse, etc.

[0098] Those skilled in the art can understand that, Figure 6 The structure shown in the figure is only a block diagram of part of the structure related to the scheme of the present application, and does not constitute a limitation on the computer device to which the scheme of the present application is applied. The specific computer device can include more or fewer components than those shown in the figure, or combine certain components, or have a different component arrangement.

[0099] In one embodiment, a computer device is also provided, including a memory and a processor, the memory storing a computer program, and the processor executing the computer program to realize the steps in each of the above method embodiments.

[0100] In one embodiment, a computer readable storage medium is provided, which stores a computer program, and the computer program is executed by a processor to realize the steps in each of the above method embodiments.

[0101] In one embodiment, a computer program product is provided, including a computer program, and the computer program is executed by a processor to realize the steps in each of the above method embodiments.

[0102] It should be noted that the user information (including but not limited to user equipment information, user personal information, etc.) and data (including but not limited to data for analysis, stored data, displayed data, etc.) involved in the present application are all information and data authorized by the user or authorized by all parties, and the collection, use and processing of related data need to comply with relevant regulations.

[0103] It can be understood by those skilled in the art that all or part of the processes in the above-mentioned embodiment methods can be completed by instructing related hardware through a computer program. The computer program can be stored in a non-volatile computer readable storage medium. When the computer program is executed, it can include the processes of the above-mentioned embodiments of each method. In the embodiments provided in the present application, any reference to memory, database or other medium can include at least one of non-volatile memory and volatile memory. Non-volatile memory can include read-only memory (ROM), magnetic tape, floppy disk, flash memory, optical storage, high-density embedded non-volatile memory, resistive random access memory (ReRAM), magnetoresistive random access memory (MRAM), ferroelectric random access memory (FRAM), phase change memory (PCM), graphene memory, etc. Volatile memory can include random access memory (RAM) or external cache memory, etc. As an illustration but not limitation, RAM can be in various forms, such as static random access memory (SRAM) or dynamic random access memory (DRAM), etc. The database involved in the embodiments provided in the present application can include at least one of a relational database and a non-relational database. The non-relational database can include a distributed database based on a block chain, etc., without being limited thereto. The processor involved in the embodiments provided in the present application can be a general-purpose processor, a central processing unit, a graphics processing unit, a digital signal processor, a programmable logic device, a data processing logic device based on quantum computing, an artificial intelligence (AI) processor, etc., without being limited thereto.

[0104] Any technical features in the above embodiments can be combined, and for the sake of brevity, not all possible combinations are described above, however, any combination of these technical features is deemed to be within the scope of the present application.

[0105] The above embodiments only express several implementation manners of the present application, and the description is relatively specific and detailed, but it should not be understood as a limitation on the patent scope of the present application. It should be pointed out that, for ordinary skilled persons in the art, several modifications and improvements can be made without departing from the concept of the present application, and these all belong to the protection scope of the present application. Therefore, the protection scope of the present application should be subject to the appended claims.

Claims

1. A synchronization system for tensor memory accelerator instructions, characterized in that, The system includes: an execution unit and a tensor core, wherein the tensor core integrates a synchronization control module and a tensor memory accelerator, wherein: The execution unit is used to send a tensor memory accelerator instruction to the tensor core. The tensor memory accelerator instruction carries synchronization information, which includes the monitoring status of the first target signal and the update status of the second target signal after the tensor memory accelerator instruction is completed. The synchronization control module is used to receive the tensor memory accelerator instruction, parse the synchronization information from the tensor memory accelerator instruction, and based on the monitoring state of the first target signal and the update state of the second target signal in the synchronization information, control the state synchronization of the first target signal and the state update of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction.

2. The system according to claim 1, characterized in that, Based on the monitored state of the first target signal and the updated state of the second target signal in the synchronization information, the control of state synchronization of the first target signal and state update of the second target signal before and after the tensor memory accelerator executes the tensor memory accelerator instruction includes: The state of the first target signal is monitored, and when the state of the first target signal is detected to be the state to be monitored, the tensor memory accelerator is controlled to execute the accelerator instruction. After the tensor memory accelerator completes the accelerator instruction, the state of the second target signal is updated to the state to be updated.

3. The system according to claim 2, characterized in that, The system further includes a synchronization state storage module, which is used to store the state of the first target signal and the state of the second target signal.

4. The system according to claim 3, characterized in that, The step of monitoring the state of the first target signal, and controlling the tensor memory accelerator to execute the accelerator instructions when the state of the first target signal is detected to be the state to be monitored, includes: The synchronization control module is used to send a monitoring signal to the synchronization state storage module. The monitoring signal carries the monitoring state of the first target signal and is used to instruct the synchronization state storage module to monitor the state of the first target signal. The synchronization state storage module is used to respond to the monitoring signal and, when the state of the first target signal is the state to be monitored, to feed back a response signal to the synchronization control module. The synchronization control module is further configured to send an execution notification to the tensor memory accelerator in response to the response signal, the execution notification being used to instruct the tensor memory accelerator to execute the tensor memory accelerator instruction.

5. The system according to claim 3 or 4, characterized in that, The step of updating the state of the second target signal to the state to be updated after the tensor memory accelerator completes the accelerator instruction includes: The synchronization control module is further configured to send an update signal to the synchronization state storage module in response to the instruction completion notification fed back by the tensor memory accelerator. The update signal carries the pending update state of the second target signal and is used to instruct the synchronization state storage module to update the state of the second target signal. The synchronization state storage module is further configured to update the state of the second target signal to the state to be updated in response to the update signal.

6. The system according to claim 1, characterized in that, When the tensor memory accelerator instruction is a load instruction, the monitoring state of the first target signal is the idle state of the storage space, and the updating state of the second target signal is the ready state of the operation data. When the tensor memory accelerator instruction is a storage instruction, the monitoring state of the first target signal is the ready state of the operation data, and the updating state of the second target signal is the idle state of the storage space.

7. A method for synchronizing instructions in a tensor memory accelerator, characterized in that, A synchronization control module integrated in a tensor core, wherein a tensor memory accelerator is also deployed in the tensor core, the method comprising: The system receives a tensor memory accelerator instruction sent by the execution unit. The tensor memory accelerator instruction carries synchronization information, which includes the monitoring status of the first target signal and the update status of the second target signal after the tensor memory accelerator instruction is completed. The synchronization information is parsed from the tensor memory accelerator instructions, and based on the monitoring state of the first target signal and the update state of the second target signal in the synchronization information, the state synchronization of the first target signal and the state update of the second target signal are controlled before and after the tensor memory accelerator executes the tensor memory accelerator instructions.

8. A computer device comprising a memory and a processor, wherein the memory stores a computer program, characterized in that, When the processor executes the computer program, it implements the steps of the method of claim 7.

9. A computer-readable storage medium having a computer program stored thereon, characterized in that, When the computer program is executed by the processor, it implements the steps of the method of claim 7.

10. A computer program product, comprising a computer program, characterized in that, When the computer program is executed by the processor, it implements the steps of the method of claim 7.

Images