Synchronous processing method and device, computer equipment, readable storage medium and program product

Abstract

The invention relates to a synchronous processing method and device, computer equipment, a computer readable storage medium and a computer program product. The method comprises the steps of obtaining a to-be-processed task, splitting the to-be-processed task into at least one to-be-processed instruction, and distributing the to-be-processed instructions to different thread processing units; obtaining an access sequence barrier corresponding to the to-be-processed task, and sending the access sequence barrier to each thread processing unit; when the execution of the at least one to-be-processed instruction is completed, generating access sequence barrier execution completion feedback information, and storing the access sequence barrier execution completion feedback information to a target barrier cache region corresponding to the thread processing unit; if the corresponding access sequence barrier execution completion feedback information is stored in each target barrier cache region, it is determined that the to-be-processed task is completed, and the target barrier cache regions correspond to the thread processing units. By adopting the method, the processing efficiency can be improved.

Description

Technical Field

[0001] The present application relates to the field of computer technology, and in particular to a synchronization processing method, apparatus, computer equipment, computer-readable storage medium, and computer program product. Background Art

[0002] A barrier, also called a fence in some architectures, is an operation that explicitly enforces a certain type of memory access order. In multithreaded operations, barriers are often used to ensure memory access order and prevent out-of-order memory access caused by unsafe optimizations by the compiler and hardware.

[0003] In related art, when using barrier technology to synchronize multiple threads, after receiving a barrier instruction, instructions following it are executed only after all preceding instructions of the specified type have completed execution. However, this can cause instructions in threads following the barrier to be blocked. If there are many instructions preceding the barrier and they execute slowly, this can significantly impact the performance of threads following the barrier. Consequently, thread execution efficiency in related art is low. Summary of the Invention

[0004] Based on this, it is necessary to provide a synchronization processing method, apparatus, computer equipment, computer-readable storage medium and computer program product that can improve thread processing efficiency in order to address the above technical problems.

[0005] In a first aspect, the present application provides a synchronization processing method, comprising:

[0006] Obtaining a task to be processed, splitting the task to be processed into at least one instruction to be processed, and allocating the instruction to be processed to different thread processing units;

[0007] Obtaining an access order barrier corresponding to a task to be processed, and sending the access order barrier to each thread processing unit, wherein each thread processing unit is configured to process the corresponding at least one instruction to be processed and the access order barrier;

[0008] When at least one pending instruction is executed, generating access order barrier execution completion feedback information, and storing the access order barrier execution completion feedback information in a target barrier buffer corresponding to the thread processing unit;

[0009] If each target barrier buffer stores corresponding access sequence barrier execution completion feedback information, it is determined that the to-be-processed task has been completed, and the target barrier buffer corresponds to each thread processing unit.

[0010] In one embodiment, after each pending instruction is sent to the thread processing unit, it also includes: determining the first path control unit based on the task tag of the pending instruction through the thread processing unit, and updating the count value of the current context number in the first path control unit, and the count value of the current context number is used to indicate whether the pending instruction has been executed.

[0011] In one embodiment, after sending the access order barrier to the thread processing unit, it also includes: determining a second path control unit based on the task tag of the access order barrier by the thread processing unit, and updating the current context number in the second path control unit to facilitate receiving the next instruction to be processed.

[0012] In an optional embodiment, it also includes: assigning a packet label to the instruction to be processed through the thread processing unit, and associating and storing the packet label, the path identifier of the first path control unit corresponding to the instruction to be processed, and the context number corresponding to the instruction to be processed; sending the instruction to be processed to the memory unit for processing; upon receiving a processing completion signal returned by the memory unit, obtaining the path identifier and context number corresponding to the packet label; based on the path identifier and context number, updating the count value of the corresponding context number; when the count value of the context number is the target value, determining that the execution of the instruction to be processed is completed.

[0013] In one embodiment, the access sequence barrier execution completion feedback information is stored in the target barrier cache area corresponding to the thread processing unit to be processed, including: determining the path identifier and the thread processing unit identifier of the path control unit corresponding to the access sequence barrier completion feedback information; determining the target storage pointer according to the thread processing unit identifier and the path identifier of the path control unit, and determining the current storage position corresponding to the target storage pointer as the target barrier cache area; wherein the position pointed to by the target storage pointer is updated after each access sequence barrier completion feedback information is stored; and storing the access sequence barrier completion feedback information in the target barrier cache area.

[0014] In one embodiment, before storing the access sequence barrier completion feedback information in the target barrier cache, the method further includes: obtaining a credit value corresponding to the thread processing unit identifier; the credit value is used to indicate whether the thread processing unit is authorized to store the access sequence barrier completion feedback information; when the credit value table indicates that the thread processing unit is authorized, executing the step of storing the access sequence barrier completion feedback information in the target barrier cache, and updating the credit value if the access sequence barrier completion feedback information is successfully stored.

[0015] In one embodiment, before storing the access order barrier execution completion feedback information in the target barrier cache area corresponding to the thread processing unit, the method further includes: obtaining a credit value corresponding to the thread processing unit; when the current credit value indicates that the thread processing unit is not authorized, detecting whether the credit value indicates that the thread processing unit is authorized, and executing the step of storing the access order barrier completion feedback information in the target barrier cache area when the credit value indicates that the thread processing unit is authorized.

[0016] In one embodiment, the access sequence barrier completion feedback information also includes a task tag; after storing the access sequence barrier execution completion feedback information in the target barrier cache corresponding to the thread processing unit, it also includes: obtaining all thread processing unit identifiers corresponding to the task tag, and determining the target barrier cache corresponding to each thread processing unit identifier; when each target barrier cache stores the corresponding access sequence barrier execution completion feedback information, clearing the target barrier cache and updating the credit value corresponding to the thread processing unit.

[0017] In an optional embodiment, before the target barrier cache is cleared, it also includes: when the number of thread processing units corresponding to the task tag is less than the target number, determining the cache bar where the target barrier cache is located, and updating the position pointed to by the target pointer corresponding to the barrier cache that does not store the access order barrier execution completion feedback information in the cache bar.

[0018] In a second aspect, the present application further provides a synchronization processing device, comprising:

[0019] An instruction splitting and distributing module is used to obtain a task to be processed, split the task to be processed into at least one instruction to be processed, and distribute the instruction to be processed to different thread processing units;

[0020] a barrier module configured to obtain an access sequence barrier corresponding to a task to be processed and send the access sequence barrier to each thread processing unit, wherein each thread processing unit is configured to process the corresponding at least one task to be processed and the access sequence barrier;

[0021] a storage module, configured to generate access order barrier execution completion feedback information when at least one pending instruction is executed, and store the access order barrier execution completion feedback information in a target barrier buffer corresponding to the thread processing unit;

[0022] The task completion determination module is used to determine that the to-be-processed task has been completed if each target barrier buffer area stores corresponding access sequence barrier execution completion feedback information, and the target barrier buffer area corresponds to each thread processing unit.

[0023] In a third aspect, the present application further provides a computer device comprising a memory and a processor, wherein the memory stores a computer program, and the processor implements the steps of the above-mentioned method embodiments when executing the computer program.

[0024] In a fourth aspect, the present application further provides a computer-readable storage medium on which a computer program is stored, and when the computer program is executed by a processor, the steps of the above-mentioned method embodiments are implemented.

[0025] In a fifth aspect, the present application also provides a computer program product, including a computer program, which implements the steps of the above-mentioned method embodiments when executed by a processor.

[0026] The above-mentioned synchronous processing method, device, computer equipment, computer-readable storage medium and computer program product generate access sequence barrier execution completion feedback information when determining that at least one pending instruction is executed. The access sequence barrier execution completion feedback information can be stored in the target cache area, and when each target barrier cache area stores the corresponding access sequence barrier execution completion feedback information, it is determined that the pending task has been completed. When it is determined that each target barrier cache area stores the corresponding access sequence barrier execution completion feedback information, it is determined that the pending task has been completed. This can achieve parallel processing of each thread without interfering with each other, improve processing efficiency, and also achieve multi-task parallel processing, further improving processing efficiency. BRIEF DESCRIPTION OF THE DRAWINGS

[0027] In order to more clearly illustrate the technical solutions in the embodiments of the present application or related technologies, the following briefly introduces the drawings required for use in the embodiments of the present application or related technical descriptions. Obviously, the drawings described below are only some embodiments of the present application. For ordinary technicians in this field, other related drawings can be obtained based on these drawings without paying any creative work.

[0028] Figure 1 1 is a flowchart of a synchronization processing method in a related art according to an embodiment;

[0029] Figure 2 1 is a flow chart of a synchronization processing method according to an embodiment;

[0030] Figure 3 is a structural block diagram of a thread processing unit in one embodiment;

[0031] Figure 4 is a structural block diagram of a thread processing unit in another embodiment;

[0032] Figure 5 A structural block diagram of a thread processing unit in yet another embodiment;

[0033] Figure 6 A schematic diagram of barrier storage in one embodiment;

[0034] Figure 7 A schematic diagram of barrier storage in another embodiment;

[0035] Figure 8 A schematic diagram of barrier storage in yet another embodiment;

[0036] Figure 9 A schematic diagram of barrier storage in yet another embodiment;

[0037] Figure 10 A schematic diagram of barrier storage in yet another embodiment;

[0038] Figure 11 A schematic diagram of barrier storage in yet another embodiment;

[0039] Figure 12 A schematic diagram of barrier storage in another embodiment;

[0040] Figure 13 A schematic diagram of barrier storage in yet another embodiment;

[0041] Figure 14 is a flowchart of a synchronization processing method in another embodiment;

[0042] Figure 15 A schematic flow chart of a synchronization processing method in yet another embodiment;

[0043] Figure 16 is a structural block diagram of a synchronization processing module in one embodiment;

[0044] Figure 17 FIG. 1 is a diagram showing the internal structure of a computer device in one embodiment. DETAILED DESCRIPTION

[0045] In order to make the purpose, technical solutions and advantages of this application more clear, the following further describes this application in detail with reference to the accompanying drawings and embodiments. It should be understood that the specific embodiments described herein are only used to explain this application and are not intended to limit this application.

[0046] In the SIMT (Single Instruction, Multiple Threads) computing model, each thread executes independently. However, in real applications and algorithms, with the exception of fully parallelizable computations like vector addition, parallel threads often require some form of coordination and communication. For example:

[0047] ① A task depends on the results of another task, such as a producer-consumer relationship;

[0048] ② The intermediate results of several tasks need to be aggregated before processing, such as reduction operations.

[0049] This requires the introduction of some form of synchronization operation, taking the thread synchronization in Thread Block as an example (such as Figure 1 shown):

[0050] In the CUDA (Compute Unified Device Architecture) programming model, __syncthreads() can be used to synchronize threads within the same task. Its corresponding PTX instruction is the Fence instruction. This instruction creates a synchronization fence at its program counter (PC) location and requires all threads within the task to reach this fence before continuing execution. This is achieved by monitoring the thread PC.

[0051] In the existing Fence synchronization mechanism, counters are often used to achieve thread synchronization. Taking a 16-thread fence as an example, in each Thread Processing Unit (TPU), multiple groups of counters are set according to the instruction type of the Thread instruction. Taking 4 types as an example, there are 64 Fence counters in all TPUs. Each counter is responsible for counting instructions of a fixed type in the corresponding thread. It increases when an instruction is received and decreases when an instruction completion signal is received. When the counter returns to zero, it is considered that all instructions before the fence have been completed. At this time, the fence is considered complete, and the fence is sent to the central module for synchronization with other threads to complete the synchronization operation.

[0052] Due to thread synchronization requirements, even if some threads run faster and complete the Fence instruction first, they are still paused until all threads in the entire task complete the Fence instruction before the entire task can continue execution. Therefore, if the execution speeds of different threads vary due to thread imbalance, the overall execution speed of the task will be limited by the slowest thread, thereby reducing the overall parallelization capability of the threads.

[0053] To solve the above technical problems, in an exemplary embodiment, Figure 2 As shown, a synchronization processing method is provided. This embodiment uses the method applied to a server as an example for illustration. It is understandable that the method can also be applied to a terminal, and can also be applied to a system including a terminal and a server, and implemented through the interaction between the terminal and the server. In this embodiment, the method includes the following steps:

[0054] Step 202 : Obtain a task to be processed, split the task to be processed into at least one instruction to be processed, and distribute the instruction to be processed to different thread processing units.

[0055] In one embodiment, the pending tasks are distributed by a driver, received by a GPU (graphics processing unit), and split into at least one pending instruction. Exemplarily, the pending tasks can be split into input / output (I / O) instructions.

[0056] In one embodiment, different instructions may correspond to different instruction types. Optionally, the instruction types may include, but are not limited to, data transfer instructions, operation instructions, program control instructions, and processor control instructions. In one embodiment, data transfer instructions may include, but are not limited to, transfer instructions, exchange instructions, stack operation instructions, input / output instructions, etc. In one embodiment, operation instructions may include, but are not limited to, arithmetic operation instructions, logical operation instructions, and shift instructions. In one embodiment, program control instructions may include, but are not limited to, transfer instructions, loop control instructions, subroutine call and return instructions, etc. In one embodiment, processor control instructions may include, but are not limited to, halt instructions, no-operation instructions, and interrupt control instructions.

[0057] In one embodiment, a Command Stream Processor / Parser (CSP) may be provided to receive pending tasks distributed by the driver and split the pending tasks into at least one pending instruction.

[0058] The thread processing unit is used to process pending instructions. Exemplarily, the thread processing unit obtains pending instructions and sends the obtained pending instructions to the memory unit. The memory unit processes the pending instructions and generates a processing completion feedback signal after processing is completed. The processing completion feedback signal is then returned. The pending instructions returned by the memory unit correspond to the processing completion feedback signal.

[0059] Optionally, there is at least one thread processing unit, for example, there may be 16 thread processing units. By using multiple thread processing units to implement parallel processing, task processing efficiency can be improved.

[0060] In one embodiment, each thread processing unit has a corresponding thread processing unit identifier. For example, for 16 thread processing units, the thread processing unit identifiers may be TPU0 to TPU15.

[0061] In one embodiment, the thread processing unit transmits the instructions to be processed to the memory unit via an on-chip bus.

[0062] In an optional embodiment, the tasks to be processed may include, but are not limited to, image rendering tasks, computing tasks, etc. For example, graphics rendering tasks may include, but are not limited to, processing 3D graphics, video rendering, and game screens. For example, computing tasks may include, but are not limited to, executing parallel computing, using a GPU for non-graphics-related computing, etc.

[0063] In one embodiment, after each instruction to be processed is sent to the thread processing unit, the thread processing unit may further determine the first path control unit based on the task tag of the instruction to be processed and update the count value of the current context number in the first path control unit.

[0064] The task tag is used to mark different pending tasks. Optionally, the task tag can be a task number, or other forms, as long as it is used to mark the pending tasks, and there is no limitation on this.

[0065] Wherein, the route is used to mark the pending instructions corresponding to different pending tasks. Optionally, the route has a corresponding route identifier, such as Route0.

[0066] The Route Control Unit (RCU) is used to execute and record received instructions to be processed.

[0067] Optionally, the thread processing unit includes a path control unit, and the path control unit is used to record the received instructions to be processed.

[0068] In one embodiment, there is at least one path control unit. Optionally, each thread processing unit includes at least one path control unit.

[0069] In one embodiment, each path control unit includes a corresponding number to distinguish different path control units, for example, RCU0 to RCUn-1.

[0070] In an exemplary embodiment, a route configuration register (RouteConfig Register) may also be provided in the thread processing unit, for determining a route control unit corresponding to the instruction to be processed according to a task tag of the instruction to be processed.

[0071] In one embodiment, after receiving a pending instruction, the thread processing unit determines the task tag of the pending instruction and uses the path configuration register to determine the path control unit corresponding to the pending instruction. For example, for a pending instruction with a task tag of Type 0, the task tag of Type 0 is input into the path configuration register, and the path control unit corresponding to the pending instruction obtained from the path configuration register is RCU0.

[0072] The context number is used to distinguish between pending instructions corresponding to different pending tasks. For example, a first pending task may have its corresponding pending instruction recorded under a first context number, while a second pending task may have its corresponding pending instruction recorded under a second context number, where the first context number is different from the second context number.

[0073] In one embodiment, there may be at least one context number, for example, there may be two context numbers.

[0074] In one embodiment, the context number may be represented by a combination of letters and numbers, such as Context 0. Alternatively, the context number may be represented in other ways, as long as it can be distinguished, and there is no limitation on this.

[0075] The current context number is used to represent the context number currently used by the path control unit. In one embodiment, after determining the path control unit, the current context number corresponding to the path identifier of the path control unit can be obtained according to the path identifier of the path control unit, and the instruction to be processed can be recorded under the current context number.

[0076] The count value is used to represent the number of pending instructions recorded under the context number.

[0077] It is understandable that the count value of the current context number can also be used to indicate whether the pending instruction has been executed. For example, when the count value of the current context number is the target count value, it is considered that the pending instruction under the current context number has been executed.

[0078] In one embodiment, the target count value may be preset, for example, the target count value is 0.

[0079] In one embodiment, after the current context number is determined, a count value of the current context number may be determined, and when the instruction to be processed is recorded in the current context number, the count value of the current context number is updated.

[0080] In an exemplary embodiment, updating the count value may include adding 1 to the acquired count value.

[0081] For ease of understanding, the following examples are given:

[0082] like Figure 3As shown, a pending instruction is obtained, the task tag of the pending instruction is determined to be Type0, and Type0 is input into the path configuration register. The path configuration register determines the corresponding path as Route0 based on Type0, and the path control unit RCU0 corresponding to this path is further determined to be Context0. The current context number of RCU0 is further determined to be Context0, and the pending instruction can be recorded under Context0. Optionally, when the count value of the current context number Context0 is determined to be counter, when the pending instruction is recorded under Context0, counter is updated, and the updated count value is counter+1.

[0083] In an optional embodiment, a packet label is assigned to the instruction to be processed by the thread processing unit, and the packet label, the path identifier of the first path control unit corresponding to the instruction to be processed, and the context number corresponding to the instruction to be processed are associated and stored; the instruction to be processed is sent to the memory unit for processing; when a processing completion signal is received from the memory unit, the path identifier and context number corresponding to the packet label are obtained; based on the path identifier and context number, the count value corresponding to the context number is updated; when the count value of the context number is the target value, it is determined that the execution of the instruction to be processed is completed.

[0084] The transaction ID (TID) is used to identify pending instructions so as to distinguish different pending instructions. For example, for n pending instructions, when the transaction ID is used to distinguish the pending instructions, the transaction IDs may be TID0 to TIDn-1.

[0085] In one embodiment, after receiving the pending instructions, the thread processing unit transmits them to the memory unit so that the memory unit can identify the pending instructions before processing them, for example, using a packet tag to identify the pending instructions to distinguish different pending instructions.

[0086] In one embodiment, the packet tag of the instruction to be processed, the path identifier of the path control unit and the context number can be associated and stored. Optionally, it can be stored in the response data buffer (ResponseBuffer), such as Figure 3 shown.

[0087] In an optional embodiment, after the packet tag of the pending instruction, the path identifier of the path control unit, and the context number are associated and stored, the pending instruction is sent to the memory unit, which processes the pending instruction and returns processing feedback information based on the processing result of the pending instruction. The processing feedback information is used to indicate that the memory unit has completed processing the pending instruction or that the memory unit has failed to process the pending instruction. Exemplarily, the memory unit returns a processing completion signal for the pending instruction.

[0088] In one embodiment, when receiving the processing completion signal, the thread processing unit extracts the packet tag corresponding to the instruction to be processed from the processing completion signal, and can obtain the corresponding path identifier and context number according to the packet tag.

[0089] Optionally, after determining the path identifier and the context number, the count value of the context number may be updated, exemplarily by reducing the count value by one.

[0090] In one embodiment, the updated count value is obtained and it is determined whether the count value is the target value. If so, it indicates that all pending instructions recorded under the current context number have been processed, and access sequence barrier execution completion feedback information can be generated.

[0091] Optionally, the target value may be a preset fixed value, such as zero.

[0092] For example, Figure 4 As shown, the thread processing unit sends the pending instruction, along with its TID, to the memory unit via the on-chip bus. The memory unit processes the pending instruction and returns a processing completion signal via the on-chip bus. After receiving the processing completion signal, the thread processing unit retrieves the TID of the completed instruction from the signal and uses the retrieved TID to search the response buffer to obtain the corresponding path tag and context number. After obtaining the context number, the thread processing unit updates the corresponding count value in the context number.

[0093] For example, the instruction obtained is TID0. TID0 is used to search the Response Buffer, and the corresponding path is Route0 and the corresponding context number is Context0. Further, according to the path identifier Route0, the corresponding path control unit is determined to be RCU0. Then, using the context number Context0, the count value counter corresponding to Context0 is obtained, and the obtained counter is updated to obtain an updated counter, such as counter-1.

[0094] Step 204 : Obtain the access order barrier corresponding to the task to be processed, and send the access order barrier to each thread processing unit.

[0095] Each thread processing unit is used to process at least one corresponding instruction to be processed and an access order barrier.

[0096] Among them, the access order barrier (Fence) is used to limit the instruction execution order.

[0097] In one embodiment, the number of thread processing units having instructions to be processed is obtained, the same number of access order barriers is copied according to the obtained number of thread processing units, and the access order barriers are sent to the corresponding thread processing units.

[0098] Optionally, an access order barrier may be sent by the CSP to the thread processing unit.

[0099] In one embodiment, a barrier first-in-first-out queue (FenceFIFO) may also be provided in the thread processing unit to record access sequence barriers received by the thread processing unit.

[0100] In one embodiment, the access order barrier may be sent to the thread processing unit after the CSP sends the pending instruction to the corresponding thread processing unit.

[0101] In one embodiment, after the access order barrier is sent to the thread processing unit, the thread processing unit determines the second path control unit based on the task tag corresponding to the access order barrier, and updates the current context number in the second path control unit to facilitate receiving the next instruction to be processed.

[0102] Optionally, after receiving the access order barrier, the thread processing unit may use the path configuration register to determine the path corresponding to the access order barrier, thereby determining the corresponding path control unit.

[0103] In one embodiment, the task tag corresponding to the access order barrier may be determined, and the path corresponding to the task tag may be determined using a path configuration register to determine the corresponding path control unit.

[0104] In one embodiment, after the path control unit corresponding to the access order barrier is acquired, the access order barrier is stored in a barrier first-in-first-out queue corresponding to the path control unit.

[0105] Optionally, after the access order barrier is stored in the barrier first-in-first-out queue corresponding to the path control unit, the context number of the path control unit may be switched so as to use the switched context number to record the next instruction to be processed.

[0106] For example, Figure 5As shown, assuming that the task tag corresponding to the Fence received by the thread processing unit is Type0, the task tag Type0 corresponding to the Fence is input into the path configuration register, the path configuration register determines that the path corresponding to the Fence with the task tag Type0 is Route0, and the path control unit corresponding to Route0 is RCU0, then the Fence FiFO corresponding to RCU0 is recorded as Fence0, for example.

[0107] In one embodiment, after completing the recording of the Fence, the current context number of RCU0 is switched to Context 1. If another pending instruction is received, the pending instruction received at this time is recorded under Context 1. Optionally, the recording operation when a pending instruction is received again is described in step 202 and will not be repeated here.

[0108] For example, Figure 4 As shown, the thread processing unit receives an instruction to be processed, the TID corresponding to the instruction to be processed is TIDn+1, the path identifier determined by the task tag of the instruction to be processed is Route0, the current context number corresponding to Route0 is Context1, TIDn+1, Route0 and Context1 are set accordingly and stored in ResponseBuffer.

[0109] Step 206 : When at least one pending instruction is executed, generate access order barrier execution completion feedback information, and store the access order barrier execution completion feedback information in the target barrier buffer corresponding to the thread processing unit.

[0110] In one embodiment, when the count value reaches the target count value, it can be considered that at least one pending instruction has been executed, and access order barrier execution completion feedback information can be generated. For example, when the count value counter of context number Context0 reaches the target count value (e.g., 0), it can be considered that at least one pending instruction in Context0 has been executed, and access order barrier execution completion feedback information can be generated.

[0111] Optionally, it is possible to detect in real time whether the count value corresponding to each context number in the thread processing unit is the target count value. If so, it indicates that all pending instructions under the context number have been executed, and access sequence barrier execution completion feedback information can be generated.

[0112] In an exemplary embodiment, a trigger module may also be provided, and the trigger module is used to trigger the step of generating access sequence barrier execution completion feedback information when detecting that the count value is the target count value.

[0113] In one embodiment, after generating access order barrier execution completion feedback information, the thread processing unit can transmit the access order barrier execution completion feedback information to the central module, and the central module stores and processes the received access order barrier execution completion feedback information.

[0114] In an optional embodiment, the central module includes a barrier buffer (Fence Sync Buffer), wherein the barrier buffer is used to store access sequence barrier execution completion feedback information.

[0115] In one embodiment, the barrier buffer may include at least one sub-barrier buffer to store at least one access order barrier execution completion feedback information.

[0116] In one embodiment, a sub-barrier buffer may be associated with a route and a thread processing unit (TPU), so that a target barrier buffer may be determined through the route and the TPU.

[0117] In an optional embodiment, the barrier cache area corresponding to the same route may have multiple layers, for example, two layers. Figure 6 For Route0, there can be two layers of barrier buffers, namely Route0 FenceSyne Buffer0 and Route0 Fence Syne Buffer1.

[0118] In one embodiment, the access sequence barrier execution completion feedback information is stored in the target barrier cache area corresponding to the thread processing unit to be processed, including: determining the path identifier and the thread processing unit identifier of the path control unit corresponding to the access sequence barrier completion feedback information; determining the target storage pointer according to the thread processing unit identifier and the path identifier of the path control unit, and determining the current storage position corresponding to the target storage pointer as the target barrier cache area; wherein the position pointed to by the target storage pointer is updated after each access sequence barrier completion feedback information is stored; and storing the access sequence barrier completion feedback information in the target barrier cache area.

[0119] Optionally, a storage pointer may be used to indicate the current storable location.

[0120] In one embodiment, the starting position of the storage pointer is layer 0. For example, Figure 6 As shown, before storage occurs, the storage pointers are all at level 0.

[0121] In one embodiment, after receiving the access sequence barrier execution completion feedback information, the central module can obtain the corresponding Route and TPU from the access sequence barrier execution completion feedback information, determine the target barrier cache area based on the path identifier and TPU identifier of the Route, and store the access sequence barrier execution completion feedback information in the target barrier cache area.

[0122] For example, Figure 6 , assuming that TPU0 sends the Fence corresponding to Route0 to the central module, the central module extracts the path identifier of Route as Route0 and the thread processing unit identifier TPU0 from the received information, and determines the target storage pointer (such as Figure 6 The location pointed to by the target storage pointer at this time can be considered as the target barrier buffer area, and the central module can store the received Fence in the target barrier buffer area.

[0123] Optionally, after the received Fence is stored in the target barrier buffer, the target storage pointer may be updated. For example, Figure 6 As shown, the target storage pointer points to the buffer area corresponding to TPU0 in Route0 Fence Sync Buffer1.

[0124] In one embodiment, if new access order barrier completion feedback information is further received, the same steps may be used to implement storage.

[0125] For example, Figure 7 As shown, it is assumed that new access order barrier completion feedback information is received, and the path identifier corresponding to the new access order barrier completion feedback information is Route0, and the thread processing unit is TPU8. At this time, the new access order barrier completion feedback information is stored in the target cache area corresponding to TPU8, and the storage pointer is updated.

[0126] In one embodiment, a detection module can be set up to detect whether all barrier buffer areas corresponding to the same pending task have been occupied, that is, to detect whether all barrier buffer areas corresponding to the same pending task have stored corresponding access sequence barrier execution completion feedback information. If so, it indicates that the pending task has been processed; if not, it is considered that the pending task has not been completed.

[0127] For example, see Figure 8 , the target barrier buffers corresponding to TPU0~TPU14 are all occupied, while the target barrier buffer corresponding to TPU15 is not occupied. At this time, it is considered that the pending tasks have not been completed and the task completion signal cannot be returned to the CSP.

[0128] Understandably, Figure 8 As shown, the storage pointers corresponding to TPU0~TPU14 all point to the barrier buffer area of ​​Route0Fence Sync Buffer1, and the pointer of TPU15 still points to the barrier buffer area of ​​Route0Fence Sync Buffer0. At this time, the instructions of other tasks can continue to be executed, and the corresponding access sequence barrier execution completion feedback information can be stored in the barrier buffer area of ​​Route0Fence Sync Buffer1 without being affected by TPU15.

[0129] It can be understood that the number of fences that can be stored in the barrier buffer is limited. In order to avoid the transmission of the central module being blocked due to the full Fene Sync Buffer of a certain Route, a credit value can be set. The credit value is used to indicate whether the thread processing unit is authorized to store the access sequence barrier completion feedback information.

[0130] Exemplarily, a credit value corresponding to the thread processing unit identifier is obtained; when the credit value indicates that the thread processing unit is authorized, a step of storing access sequence barrier completion feedback information into a target barrier cache is performed, and if the access sequence barrier completion feedback information is successfully stored, the credit value is updated.

[0131] Optionally, when the current credit value indicates that the thread processing unit is not authorized, whether the credit value indicates that the thread processing unit is authorized is detected in real time, and when the credit value indicates that the thread processing unit is authorized, the step of storing the access order barrier completion feedback information in the target barrier buffer is performed.

[0132] In one embodiment, when the thread processing unit sends access order barrier completion feedback information to the central module, it may also send its corresponding credit value to the central module, so that the central module determines whether the thread processing unit is authorized based on the received credit value.

[0133] In one embodiment, after sending the access order barrier completion feedback information, the thread processing unit updates the credit value. For example, the credit value may be reduced by one to obtain an updated credit value.

[0134] In an optional embodiment, the initial credit value may correspond to the number of storable levels of the barrier buffer. For example, if the barrier buffer has two storable levels, the initial credit value may be set to 2. When the thread processing unit sends access order barrier completion feedback information, the updated credit value may be 2-1=1.

[0135] In one embodiment, when using a credit value to determine whether a thread processing unit is authorized, the credit value can be compared with a target credit value to determine whether the thread processing unit is authorized. For example, if the credit value is the same as the target credit value, the thread processing unit is unauthorized; otherwise, the thread processing unit is authorized. For example, if the target credit value is 0 and the credit value is 2, the thread processing unit is considered authorized. However, if the target credit value is 0 and the credit value is also 0, the thread processing unit is considered unauthorized.

[0136] In one embodiment, if the thread processing unit is not authorized, the thread processing unit cannot upload the access order barrier completion feedback information, and the thread processing unit temporarily does not perform the step of sending the access order barrier completion feedback information.

[0137] Optionally, if the thread processing unit is not authorized, the central module may temporarily not execute the step of storing the access order barrier completion feedback information.

[0138] In one embodiment, when the step of sending access order barrier completion feedback information or the step of storing access order barrier completion feedback information is not performed, the credit value of the thread processing unit can be monitored to determine whether the credit value of the thread processing unit has been updated. If the updated credit value indicates that the thread processing unit is authorized, the step of sending access order barrier completion feedback information or the step of storing access order barrier completion feedback information is continued. Optionally, after the step of sending access order barrier completion feedback information or the step of storing access order barrier completion feedback information is completed, the credit value is updated again.

[0139] For example, Figure 9 Taking TPU0 as an example, when the target storage pointer points to the cache area in Route0 Fence SyncBuffer1, the credit value of TPU0 is 1. When the access sequence barrier completion feedback information corresponding to TPU0 of Task 1 is stored, the target storage pointer points to the cache area of ​​Route0 Fence Sync Buffer0 again, and the credit value of TPU0 is 0. It can be considered that the target pointer points to the cache area at this time. However, the credit value of TPU0 is the target credit value, so TPU0 is not authorized and cannot upload the access sequence barrier execution completion feedback information, or the central module does not perform the storage operation for the access sequence barrier execution completion feedback information uploaded by TPU0 until the credit value of TPU0 is updated, and the updated credit value indicates that TPU0 is authorized.

[0140] It is understandable that the same processing method is used for other TPUs. For example, Figure 10As shown, TPU7 is specifically explained as follows. When the target storage pointer points to the cache area in Route0 Fence Sync Buffer1, the credit value of TPU7 is 1. When the access sequence barrier corresponding to TPU7 of Task 1 is completed and the feedback information is stored, the target storage pointer points to the cache area of ​​Route0 Fence Sync Buffer0 again, and the credit value of TPU7 is 0. It can be considered that the target pointer points to the cache area at this time, but the credit value of TPU0 is the target credit value, then TPU7 is not authorized and cannot upload the access sequence barrier execution completion feedback information, or the central module does not perform the storage operation for the access sequence barrier execution completion feedback information uploaded by TPU7 until the credit value of TPU7 is updated, and the updated credit value indicates that TPU7 is authorized.

[0141] In one embodiment, if the access sequence barrier execution completion feedback information uploaded by TPU15 corresponding to task 0 is received at this time, the target barrier buffer is determined according to the direction of the target storage pointer of TPU15, and the access sequence barrier execution completion feedback information is stored in the target barrier buffer, such as Figure 11 shown.

[0142] In step 208 , if each target barrier buffer stores the corresponding access sequence barrier execution completion feedback information, it is determined that the pending task has been completed.

[0143] The target barrier buffer area corresponds to each thread processing unit.

[0144] In one embodiment, a determination step may be performed to determine whether each target barrier buffer stores the corresponding access sequence barrier execution completion feedback information. If so, the pending task is determined to be completed. Optionally, after determining that the pending task is completed, a pending task completion signal may be sent to the CSP so that the next pending task can be executed.

[0145] In one embodiment, the access order barrier completion feedback information also includes a task tag. The task tag can be used to determine whether all thread processing units of the same task have generated access order barrier execution completion feedback information, and the access order barrier execution completion feedback information is sent to the central module. Optionally, all thread processing unit identifiers corresponding to the task tag can be obtained, and the target barrier buffer corresponding to each thread processing unit identifier can be determined; when each target barrier buffer stores the corresponding access order barrier execution completion feedback information, the target barrier buffer is cleared and the credit value corresponding to the thread processing unit is updated.

[0146] In one embodiment, a task tag of a pending task can be obtained, and all corresponding thread processing unit identifiers can be obtained based on the task tag. The thread processing unit identifiers can be used to determine whether access sequence barrier execution completion feedback information has been stored in the corresponding target barrier buffer. If so, it indicates that the pending task has been processed and a task completion signal can be generated and sent to the CSP so that the execution of the new pending task can continue.

[0147] Optionally, after determining that the task is completed and notifying the CSP, each target barrier buffer corresponding to the task tag may be cleared so that each barrier buffer can continue to be used for storing access sequence barrier execution completion feedback information.

[0148] For example, Figure 12 As shown, the target buffers corresponding to each TPU0-TPU15 in Route0 Fence Sync Buffer0 are all occupied. At this point, the task can be considered to have been completed, a task completion signal is generated, and the task completion signal is sent to the CSP. Optionally, the credit values ​​of each TPU0-TPU15 are updated, for example, by adding 1 to the current credit values ​​of TPU0-TPU15. Optionally, the target buffers corresponding to each TPU0-TPU15 in Route0 Fence Sync Buffer0 can also be cleared to receive new access order barrier execution completion feedback information.

[0149] In one of the embodiments, when the number of thread processing units corresponding to the task tag is less than the target number, the cache bar where the target barrier cache area is located is determined, and the position pointed to by the target pointer corresponding to the barrier cache area in the cache bar that does not store access order barrier execution completion feedback information is updated.

[0150] For example, Figure 13 As shown in the figure, assuming that the task only contains instructions for thread 0 and thread 1, only thread 0 and thread 1 need to be synchronized. After the central module collects the fences of thread 0 and thread 1, it completes synchronization based on the thread information carried in the fences. It also shifts the pointers corresponding to threads 2 to 15 to ensure that the fences in the next task can be written to the correct location.

[0151] In the above-mentioned synchronous processing method, after determining that the pending task has been completed, access sequence barrier execution completion feedback information is generated. When the corresponding access sequence barrier execution completion feedback information is stored in each target barrier cache, it is determined that the pending task has been completed. By counting whether each target barrier cache stores the corresponding access sequence barrier execution completion feedback information, the barrier can be used to define the instruction execution order. Secondly, whether the target barrier cache stores the corresponding access sequence barrier execution completion feedback information to determine whether the pending task is completed, the parallelism of each thread processing can be achieved without interfering with each other, thereby improving processing efficiency, and multi-task parallel processing can also be achieved, further improving processing efficiency.

[0152] In an exemplary embodiment, Figure 14 As shown, a synchronization processing method is provided. This embodiment uses the method applied to a server as an example for illustration. It is understandable that the method can also be applied to a terminal, and can also be applied to a system including a terminal and a server, and implemented through the interaction between the terminal and the server. In this embodiment, the method includes the following steps:

[0153] Step 302: Obtain a task to be processed, split the task to be processed into at least one instruction to be processed, and distribute the instruction to be processed to different thread processing units.

[0154] Step 304 : Determine a first path control unit based on the task tag of the instruction to be processed by the thread processing unit, and update the count value of the current context number in the first path control unit.

[0155] The count value of the current context number is used to indicate whether the pending instruction has been executed.

[0156] Step 306: Allocate a packet tag to the instruction to be processed by the thread processing unit, and associate and store the packet tag, the path identifier of the first path control unit corresponding to the instruction to be processed, and the context number corresponding to the instruction to be processed.

[0157] Step 308: Send the processing instruction to the memory unit for processing.

[0158] Step 310: Upon receiving the processing completion signal returned by the memory unit, obtain the path identifier and context number corresponding to the packet label.

[0159] Step 312: Based on the path identifier and the context number, update the count value corresponding to the context number.

[0160] Step 314 : When the count value of the context number reaches the target value, it is determined that the execution of the pending instruction is completed.

[0161] Step 316 : Obtain the access order barrier corresponding to the task to be processed, and send the access order barrier to each thread processing unit.

[0162] Each thread processing unit is used to process at least one corresponding instruction to be processed and an access order barrier.

[0163] In step 318 , the thread processing unit determines the second path control unit based on the task tag of the access order barrier, and updates the current context number in the second path control unit to facilitate receiving the next instruction to be processed.

[0164] Step 320 , determining whether at least one pending instruction has been executed. If so, executing step 322 ; otherwise, executing step 324 .

[0165] Step 322 , generating access order barrier execution completion feedback information, and storing the access order barrier execution completion feedback information in the target barrier buffer corresponding to the thread processing unit, and then proceeding to step 326 .

[0166] Step 324 , detecting whether at least one pending instruction has been executed, and continuing to step 322 until at least one pending instruction has been executed.

[0167] Step 326 , determining whether each target barrier buffer stores corresponding access sequence barrier execution completion feedback information, if so, executing step 328 , if not, executing step 330 .

[0168] The target barrier buffer area corresponds to each thread processing unit.

[0169] Step 328 : Determine that the pending task has been completed, and continue to execute step 302 .

[0170] Step 330 , detecting whether each target barrier buffer stores the corresponding access order barrier execution completion feedback information. Execution continues at step 328 until each target barrier buffer stores the corresponding access order barrier execution completion feedback information.

[0171] In this embodiment, in the above-mentioned synchronous processing method, after determining that the pending task has been completed, access sequence barrier execution completion feedback information is generated. When the corresponding access sequence barrier execution completion feedback information is stored in each target barrier cache, it is determined that the pending task has been completed. By counting whether each target barrier cache stores the corresponding access sequence barrier execution completion feedback information, the barrier can be used to define the instruction execution order. Secondly, whether the target barrier cache stores the corresponding access sequence barrier execution completion feedback information to determine whether the pending task is completed, the parallelism of each thread processing can be achieved without interfering with each other, thereby improving processing efficiency, and multi-task parallel processing can also be achieved, further improving processing efficiency.

[0172] In an exemplary embodiment, Figure 15 As shown, a synchronization processing method is provided. This embodiment uses the method applied to a server as an example for illustration. It is understandable that the method can also be applied to a terminal, and can also be applied to a system including a terminal and a server, and implemented through the interaction between the terminal and the server. In this embodiment, the method includes the following steps:

[0173] Step 402 : Obtain a task to be processed, split the task to be processed into at least one instruction to be processed, and distribute the instruction to be processed to different thread processing units.

[0174] Step 404 : Obtain the access order barrier corresponding to the task to be processed, and send the access order barrier to each thread processing unit.

[0175] Each thread processing unit is used to process at least one corresponding instruction to be processed and an access order barrier.

[0176] Step 406 , determining whether at least one pending instruction has been executed. If not, executing step 408 ; if so, executing step 410 .

[0177] Step 408 , detecting whether at least one pending instruction has been executed, and continuing to step 410 until at least one pending instruction has been executed.

[0178] Step 410: Generate access order barrier execution completion feedback information.

[0179] Step 412: Obtain the credit value corresponding to the thread processing unit identifier.

[0180] The credit value is used to indicate whether the thread processing unit is authorized to store access order barrier completion feedback information.

[0181] Step 414 , determining whether the credit value represents the thread processing unit is authorized, if so, executing step 416 , if not, executing step 422 .

[0182] Step 416: Determine the path identifier and thread processing unit identifier of the path control unit corresponding to the access order barrier completion feedback information.

[0183] Step 418 : Determine the target storage pointer according to the thread processing unit identifier and the path identifier of the path control unit, and determine that the current storage location corresponding to the target storage pointer is the target barrier buffer.

[0184] The location pointed to by the target storage pointer is updated after each access sequence barrier completes the storage of feedback information.

[0185] Step 420 : Store the access sequence barrier completion feedback information into the target barrier buffer.

[0186] Step 422 , detecting whether the credit value indicates that the thread processing unit is authorized. If the credit value indicates that the thread processing unit is authorized, step 416 is continued.

[0187] Step 424 , determining whether each target barrier buffer stores corresponding access sequence barrier execution completion feedback information, if so, executing step 426 , if not, executing step 428 .

[0188] The target barrier buffer area corresponds to each thread processing unit.

[0189] Step 426 , determining that the pending task has been completed, and continuing to step 402 .

[0190] Step 428 , detecting whether each target barrier buffer stores the corresponding access order barrier execution completion feedback information. Continuing to step 426 until each target barrier buffer stores the corresponding access order barrier execution completion feedback information.

[0191] In the above-mentioned synchronous processing method, after determining that the pending task has been completed, access order barrier execution completion feedback information is generated. When each target barrier buffer stores the corresponding access order barrier execution completion feedback information, the pending task is determined to be completed. By counting whether each target barrier buffer stores the corresponding access order barrier execution completion feedback information, the barrier can be used to define the instruction execution order. Secondly, whether the target barrier buffer stores the corresponding access order barrier execution completion feedback information can be used to determine whether the pending task has been completed. This can achieve parallel processing of each thread without interference, improving processing efficiency. It can also achieve multi-task parallel processing, further improving processing efficiency. In addition, multiple layers of barrier buffers can be provided, including thread processing units, to achieve multi-task and multi-thread processing simultaneously, effectively improving data processing efficiency. Furthermore, a credit value detection is provided, which uses the credit value to determine whether the thread processing unit can transmit the access order barrier execution completion feedback information to the barrier buffer. This ensures the orderliness of the thread processing unit uploading the access order barrier execution completion feedback information, avoids data transmission congestion caused by excessive data received by the central module, and improves data processing efficiency.

[0192] It should be understood that, although the steps in the flowcharts of the above embodiments are shown in sequence as indicated by the arrows, these steps are not necessarily performed in the order indicated by the arrows. Unless otherwise specified herein, there is no strict order restriction on the execution of these steps, and these steps can be performed in other orders. Moreover, at least a portion of the steps in the flowcharts of the above embodiments may include multiple steps or multiple stages, and these steps or stages are not necessarily performed at the same time, but can be performed at different times. The execution order of these steps or stages is not necessarily to be performed in sequence, but can be performed in turn or alternately with other steps or at least a portion of steps or stages in other steps.

[0193] Based on the same inventive concept, embodiments of the present application also provide a synchronization processing device for implementing the synchronization processing method involved above. The implementation solution provided by this device is similar to the implementation solution described in the above method. Therefore, the specific limitations of one or more synchronization processing device embodiments provided below can be found in the limitations of the synchronization processing method above and will not be repeated here.

[0194] In an exemplary embodiment, Figure 16 As shown, a synchronization processing device 1600 is provided, comprising: an instruction splitting and allocating module 1602, a barrier module 1604, a storage module 1606 and a task completion determination module 1608, wherein:

[0195] The instruction splitting and distributing module is used to obtain a task to be processed, split the task to be processed into at least one instruction to be processed, and distribute the instruction to be processed to different thread processing units.

[0196] The barrier module is used to obtain the access order barrier corresponding to the task to be processed and send the access order barrier to each thread processing unit, wherein each thread processing unit is used to process the corresponding at least one instruction to be processed and the access order barrier.

[0197] The storage module is used to generate access order barrier execution completion feedback information when at least one pending instruction is executed, and store the access order barrier execution completion feedback information in a target barrier buffer corresponding to the thread processing unit.

[0198] The task completion determination module is used to determine that the to-be-processed task has been completed if each target barrier buffer area stores corresponding access sequence barrier execution completion feedback information, and the target barrier buffer area corresponds to each thread processing unit.

[0199] In one embodiment, the synchronization processing device also includes an update module, which is used to determine the first path control unit based on the task tag of the instruction to be processed through the thread processing unit after sending each instruction to be processed to the thread processing unit, and update the count value of the current context number in the first path control unit. The count value of the current context number is used to indicate whether the execution of the instruction to be processed is completed.

[0200] In one embodiment, the update module is further used to send the access order barrier to the thread processing unit, determine the second path control unit based on the task tag of the access order barrier through the thread processing unit, and update the current context number in the second path control unit to facilitate receiving the next instruction to be processed.

[0201] In an optional embodiment, the synchronization processing device further includes:

[0202] The association module is used to assign a packet tag to the instruction to be processed through the thread processing unit, and associate and store the packet tag, the path identifier of the first path control unit corresponding to the instruction to be processed, and the context number corresponding to the instruction to be processed.

[0203] The instruction sending module is used to send the instructions to be processed to the memory unit for processing.

[0204] The acquisition module is used to acquire the path identifier and context number corresponding to the packet label when receiving the processing completion signal returned by the memory unit.

[0205] The updating module is further configured to update the count value corresponding to the context number based on the path identifier and the context number.

[0206] The instruction completion determination module is used to determine that the execution of the pending instruction is completed when the count value of the context number is the target value.

[0207] In one embodiment, the storage module is further used to determine the path identifier and thread processing unit identifier of the path control unit corresponding to the access sequence barrier completion feedback information; determine the target storage pointer according to the thread processing unit identifier and the path identifier of the path control unit, and determine that the current storage position corresponding to the target storage pointer is the target barrier cache area; wherein the position pointed to by the target storage pointer is updated after each access sequence barrier completion feedback information is stored; and store the access sequence barrier completion feedback information in the target barrier cache area.

[0208] In one embodiment, the synchronization device also includes a credit value acquisition module for acquiring a credit value corresponding to the thread processing unit identifier before storing the access sequence barrier completion feedback information in the target barrier cache area; the credit value is used to indicate whether the thread processing unit is authorized to store the access sequence barrier completion feedback information.

[0209] The updating module is further configured to, when the credit value indicates that the thread processing unit is authorized, execute the step of storing the access sequence barrier completion feedback information into the target barrier buffer area, and update the credit value if the access sequence barrier completion feedback information is stored successfully.

[0210] In an exemplary embodiment, the credit value acquisition module is further configured to acquire the credit value corresponding to the thread processing unit before storing the access order barrier execution completion feedback information in the target barrier buffer corresponding to the thread processing unit.

[0211] In one embodiment, the synchronization processing device further includes an execution module for detecting whether the credit value represents the thread processing unit when the current credit value represents that the thread processing unit is not authorized, and executing the step of storing the access sequence barrier completion feedback information into the target barrier buffer area when the credit value represents that the thread processing unit is authorized.

[0212] In one embodiment, the access sequence barrier completion feedback information also includes a task tag; the synchronization processing device also includes a target cache area determination module, which is used to obtain all thread processing unit identifiers corresponding to the task tag after storing the access sequence barrier execution completion feedback information in the target barrier cache area corresponding to the thread processing unit, and determine the target barrier cache area corresponding to each thread processing unit identifier.

[0213] Optionally, the updating module is further configured to clear the target barrier buffer and update the credit value corresponding to the thread processing unit when each target barrier buffer stores corresponding access sequence barrier execution completion feedback information.

[0214] In one embodiment, the update module is further used to determine the cache bar where the target barrier cache is located before the target barrier cache is cleared, when the number of thread processing units corresponding to the task tag is less than the target number, and update the position pointed to by the target pointer corresponding to the barrier cache that does not store the access order barrier execution completion feedback information in the cache bar.

[0215] Each module in the synchronization processing device can be implemented in whole or in part through software, hardware, or a combination thereof. Each module can be embedded in or independent of a processor in a computer device in the form of hardware, or can be stored in a memory in the computer device in the form of software, so that the processor can call and execute the corresponding operations of each module.

[0216] In an exemplary embodiment, a computer device is provided. The computer device may be a server, and its internal structure diagram may be as shown in FIG. Figure 17 As shown. The computer device includes a processor, a memory, an input / output interface (Input / Output, abbreviated as I / O) and a communication interface. The processor, memory and input / output interface are connected through a system bus, and the communication interface is connected to the system bus through the input / output interface. The processor of the computer device is used to provide computing and control capabilities. 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, a computer program and a database. The internal memory provides an environment for the operation of the operating system and computer program in the non-volatile storage medium. The database of the computer device is used to store data. The input / output interface of the computer device is used to exchange information between the processor and an external device. The communication interface of the computer device is used to communicate with an external terminal through a network connection. When the computer program is executed by the processor, a synchronous processing method is implemented.

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

[0218] In an exemplary embodiment, a computer device is provided, including a memory and a processor. The memory stores a computer program, and the processor implements the steps of the above method embodiments when executing the computer program.

[0219] In one embodiment, a computer-readable storage medium is provided, on which a computer program is stored. When the computer program is executed by a processor, the steps of the above-mentioned method embodiments are implemented.

[0220] In one embodiment, a computer program product is provided, including a computer program, which implements the steps of the above method embodiments when executed by a processor.

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

[0222] Those skilled in the art will appreciate that all or part of the processes in the above-mentioned embodiments can be implemented by instructing the relevant hardware through a computer program, and 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 embodiments of the above-mentioned methods. Among them, any reference to memory, database or other media used in the embodiments provided in this application may include at least one of non-volatile memory and volatile memory. Non-volatile memory may include read-only memory (ROM), magnetic tape, floppy disk, flash memory, optical memory, high-density embedded non-volatile memory, resistive random access memory (ReRAM), magnetic random access memory (MRAM), ferroelectric random access memory (FRAM), phase change memory (PCM), graphene memory, etc. Volatile memory may include random access memory (RAM) or external cache memory, etc. By way of illustration and not limitation, RAM can be in various forms, such as static random access memory (SRAM) or dynamic random access memory (DRAM). The database involved in the various embodiments provided herein may include at least one of a relational database and a non-relational database. Non-relational databases may include, but are not limited to, distributed databases based on blockchains. The processor involved in the various embodiments provided herein may be, but are not limited to, a general-purpose processor, a central processing unit (CPU), a graphics processing unit (GPU), a digital signal processor (DSP), a programmable logic unit (PLC), a data processing logic unit based on quantum computing, an artificial intelligence (AI) processor, and the like.

[0223] The technical features of the above embodiments can be combined arbitrarily. In order to make the description concise, not all possible combinations of the technical features in the above embodiments are described. However, as long as there is no contradiction in the combination of these technical features, they should be considered to be within the scope of this application.

[0224] The above embodiments merely illustrate several implementation methods of the present application. While the descriptions are relatively specific and detailed, they should not be construed as limiting the scope of the present invention. It should be noted that a person skilled in the art may make various modifications and improvements without departing from the spirit of the present invention, all of which fall within the scope of protection of the present application. Therefore, the scope of protection of the present application shall be determined by the appended claims.

Claims

1. A synchronization processing method, characterized in that: The method comprises: Acquire a task to be processed, split the task to be processed into at least one instruction to be processed, and distribute the instruction to be processed to different thread processing units; Obtaining an access order barrier corresponding to the task to be processed, and sending the access order barrier to each of the thread processing units, wherein each of the thread processing units is configured to process the corresponding at least one instruction to be processed and the access order barrier; When the at least one pending instruction is executed, generating access order barrier execution completion feedback information, and storing the access order barrier execution completion feedback information in a target barrier buffer corresponding to the thread processing unit; If each of the target barrier buffers stores corresponding access sequence barrier execution completion feedback information, it is determined that the to-be-processed task has been completed, and the target barrier buffer corresponds to each of the thread processing units.

2. The method according to claim 1, characterized in that After sending each instruction to be processed to the thread processing unit, the method further includes: The thread processing unit determines a first path control unit based on the task tag of the instruction to be processed, and updates the count value of the current context number in the first path control unit. The count value of the current context number is used to indicate whether the instruction to be processed is executed.

3. The method according to claim 2, characterized in that After sending the access order barrier to the thread processing unit, the method further includes: The thread processing unit determines a second path control unit based on the task tag of the access order barrier, and updates the current context number in the second path control unit to facilitate receiving the next instruction to be processed.

4. The method according to claim 2, characterized in that The method further comprises: assigning a packet tag to the instruction to be processed by the thread processing unit, and storing the packet tag, a path identifier of the first path control unit corresponding to the instruction to be processed, and a context number corresponding to the instruction to be processed in association with each other; sending the processing instruction to the memory unit for processing; Upon receiving a processing completion signal returned by the memory unit, obtaining a path identifier and a context number corresponding to the packet label; Based on the path identifier and the context number, updating the count value corresponding to the context number; When the count value of the context number is the target value, it is determined that the execution of the pending instruction is completed.

5. The method according to claim 1, wherein The step of storing the access sequence barrier execution completion feedback information into a target barrier buffer corresponding to the pending thread processing unit includes: Determine the path identifier and thread processing unit identifier of the path control unit corresponding to the access sequence barrier completion feedback information; determining a target storage pointer according to the thread processing unit identifier and the path identifier of the path control unit, and determining a current storage location corresponding to the target storage pointer as the target barrier buffer; wherein the location pointed to by the target storage pointer is updated after each access sequence barrier completes feedback information storage; The access sequence barrier completion feedback information is stored in the target barrier buffer area.

6. The method according to claim 5, characterized in that Before storing the access sequence barrier completion feedback information in the target barrier buffer area, the method further includes: Obtaining a credit value corresponding to the thread processing unit identifier; the credit value is used to indicate whether the thread processing unit is authorized to store the access sequence barrier completion feedback information; When the credit value indicates that the thread processing unit is authorized, the step of storing the access order barrier completion feedback information in the target barrier buffer is performed, and if the access order barrier completion feedback information is successfully stored, the credit value is updated.

7. The method according to claim 5, characterized in that Before storing the access order barrier execution completion feedback information in the target barrier buffer corresponding to the thread processing unit, the method further includes: Obtaining a credit value corresponding to the thread processing unit; When the current credit value indicates that the thread processing unit is not authorized, detecting whether the credit value indicates that the thread processing unit is authorized, and until the credit value indicates that the thread processing unit is authorized, performing the step of storing the access sequence barrier completion feedback information in the target barrier buffer area.

8. The method according to any one of claims 5 to 7, characterized in that: The access order barrier completion feedback information further includes a task tag; after storing the access order barrier execution completion feedback information in the target barrier buffer corresponding to the thread processing unit, the method further includes: Obtaining all thread processing unit identifiers corresponding to the task tag, and determining a target barrier buffer corresponding to each thread processing unit identifier; When each of the target barrier buffers stores the corresponding access sequence barrier execution completion feedback information, the target barrier buffer is cleared and the credit value corresponding to the thread processing unit is updated.

9. The method according to claim 8, characterized in that Before the target barrier buffer is cleared, the method further includes: When the number of thread processing units corresponding to the task tag is less than the target number, the cache bar where the target barrier cache area is located is determined, and the position pointed to by the target pointer corresponding to the barrier cache area in the cache bar that does not store the access sequence barrier execution completion feedback information is updated.

10. A synchronization processing device, characterized in that: The device comprises: An instruction splitting and distributing module is used to obtain a task to be processed, split the task to be processed into at least one instruction to be processed, and distribute the instruction to be processed to different thread processing units; a barrier module, configured to obtain an access sequence barrier corresponding to the task to be processed, and send the access sequence barrier to each of the thread processing units, wherein each of the thread processing units is configured to process the corresponding at least one instruction to be processed and the access sequence barrier; a storage module, configured to generate access order barrier execution completion feedback information when the at least one pending instruction is executed, and store the access order barrier execution completion feedback information in a target barrier buffer corresponding to the thread processing unit; The task completion determination module is configured to determine that the to-be-processed task has been completed if each of the target barrier buffers stores corresponding access sequence barrier execution completion feedback information, and the target barrier buffers correspond to each of the thread processing units.

11. A computer device comprising a memory and a processor, wherein the memory stores a computer program, wherein: When the processor executes the computer program, the steps of the method according to any one of claims 1 to 9 are implemented.

12. A computer-readable storage medium having a computer program stored thereon, characterized in that: When the computer program is executed by a processor, the steps of the method according to any one of claims 1 to 9 are implemented.

13. A computer program product comprising a computer program, characterized in that When the computer program is executed by a processor, the steps of the method according to any one of claims 1 to 9 are implemented.

Images