Data production subject unit, data consumption subject unit, related device and method

Abstract

The present disclosure provides a data production subject unit, a data consumption subject unit, related devices and methods. The data production subject unit comprises: a first thread group execution unit configured to execute a cross-subject unit write instruction to a specific data consumption subject unit for a thread group, and exchange first information between the thread group and other thread groups; and a first communication control unit configured to execute, for a thread group: if all write requests corresponding to all thread groups for the cross-subject unit write instruction have been sent from the data production subject unit, send a synchronization set instruction on each of a plurality of routes between the data production subject unit and the data consumption subject unit. The present disclosure can ensure that the use of data written by the data production subject unit by the data consumption subject unit is error-free in the case of instruction-level writing and multiple routes between the data production subject and the data consumption subject.

Description

Technical Field

[0001] This disclosure relates to the field of chips, and more specifically, to a data production unit, a data consumption unit, related apparatus and method. Background Technology

[0002] In the era of big data, the increasing scale of processing tasks typically requires processing units (including CPUs for routine computations and GPUs, NPUs, tensor processors, etc., for complex computations in graphics processing and deep neural networks) to collaboratively handle large-scale tasks. Each processing unit processes a portion of the task and stores task-related data. Since the parts processed by these units may be interconnected, a processing unit often encounters situations where it needs to access data generated and stored locally by other processing units while performing its own task. Traditionally, this requires the use of inter-host communication networks (ICNs) between hosts (computing devices) to retrieve the necessary data stored by other processing units. However, these ICNs are slow, and data transfer within the host requires traversing internal buses such as PCIe, resulting in low data retrieval speeds between processing units. To address this issue, Inter-Chip Interconnect Networks (ICNs) were proposed. ICNs directly interconnect the processing units of each host (computing device). ICNs enable rapid data retrieval between processing units.

[0003] Writing data from one processing unit to another via an inter-chip interconnect network can be done either by writing commands (called command-level writing) or by breaking down commands into instructions and writing them one instruction at a time (called instruction-level writing). An instruction is the smallest unit of command execution. This article primarily discusses instruction-level writing. In instruction-level writing, writing is typically performed in warp groups. Typically, acceleration units perform repetitive operations on batches of data. For example, in matrix multiplication, the instructions for multiplying and adding matrix elements are fixed; the only difference is that the values ​​of different elements in the matrix are substituted into the instructions for repeated calculations. Therefore, operations are performed in warp groups. A warp group consists of multiple threads. Each thread performs the same code operations on its assigned data. A large computational task can be broken down into many such warp groups. Once all warp groups have successfully completed their operations, the entire computational task is complete.

[0004] When performing instruction-level writes across processing units by thread group, a synchronization issue arises. A thread group successfully writing data to the target processing unit does not mean that the data is immediately usable. This is because the data written by that thread group is related to the data written by other thread groups in the same task. The target processing unit can only use this data after all related data in all thread groups within the task has been written; otherwise, it may read incorrect data. In the case of command-level writes, this synchronization can be achieved by setting a synchronization primitive on the source processing unit side and a synchronization wait primitive on the target processing unit side. The synchronization primitive on the source processing unit side is set after the batch write command to the target processing unit. The corresponding synchronization wait primitive on the target processing unit side is set before the command to use the data. Upon receiving the written data, the target processing unit receives the synchronization primitive and sets a flag. Then, it executes the synchronization wait primitive, and only allows subsequent commands to use the data to execute after detecting that the flag is set. In this way, in the case of command-level writes, the situation where the target processing unit starts reading data before the batch data is fully written, causing read errors, will not occur. However, as mentioned above, in the case of instruction-level writes, the entire write task is executed in groups of threads. This makes it difficult for the source processing unit to determine whether all data has been written. Furthermore, there are often multiple routes between the source and target processing units. Within a processing unit, different caches and paths may also be used. Therefore, write requests from different thread groups may arrive at the target processing unit via different routes, caches, and paths. Write requests from the same thread group may also arrive at the target processing unit via different routes, caches, and paths. The routes used to send the synchronization primitives may also differ from the routes, caches, and paths used to send each write request. This means that the synchronization primitives may arrive at the target processing unit first, while write requests from some thread groups may not arrive at the target processing unit yet due to congestion in their respective routes, caches, or paths. In this case, the target processing unit detects that the flag is set and considers synchronization complete, allowing execution of subsequent data access statements. However, this is incorrect because some thread group write requests have not yet arrived at the target processing unit, leading to data read errors. Summary of the Invention

[0005] In view of this, one objective of this disclosure is to ensure that, even with instruction-level writing and multiple routing between the data producer and the data consumer, the data consumer can still use the data written by the data producer without errors.

[0006] According to one aspect of this disclosure, a data production entity is provided, comprising: The first thread group execution unit is used to execute cross-subject unit write instructions to a specific data consuming subject unit for the thread group, and to make the thread group exchange first information with other thread groups. The first information indicates whether the write requests corresponding to the thread group have all been issued from the data producing subject unit. The first communication control unit is configured to execute the following for a thread group: if all write requests corresponding to the thread group for the cross-subject unit write instruction have been issued from the data production subject unit, a synchronization set instruction is issued on each of the multiple routes between the data production subject unit and the data consumption subject unit, so that after the data consumption subject unit receives the synchronization set instruction on each route, it ensures that all previously received write requests on that route have been successfully written, and allows the use of the written data after ensuring successful writing on each route.

[0007] Optionally, the data production unit may also include: The first multi-level cache, located between the first thread group execution unit and the first communication control unit, executes a flush instruction and a wait-for-flush-success instruction before issuing synchronization set instructions on the multiple routes. The flush instruction is used to flush out the data stored in the cache and send it from the first communication control unit to the data consumption unit through a higher-level cache. The wait-for-flush-success instruction is used to wait for all the data flushed out of the cache to be successfully sent out from the data production unit.

[0008] Optionally, after the first communication control unit issues a synchronization set instruction on each of the multiple routes, it executes a synchronization set waiting instruction, which is used to indicate that no subsequent instruction can be executed before the synchronization set instructions are issued on the multiple routes.

[0009] Optionally, after the first communication control unit issues synchronization set instructions on the multiple routes respectively, the first thread group execution unit also causes the thread group to exchange third information with other thread groups, the third information indicating that synchronization set instructions have been issued on the multiple routes respectively.

[0010] Optionally, the first communication control unit includes: Multiple direct storage access modules, corresponding to multiple sending ports, each store the write requests to be sent from the corresponding sending port; The instruction allocator, based on the routing table between the data production unit and the data consumption unit, allocates the write requests corresponding to the thread group to the direct storage access module of the sending port of the corresponding route. After determining that all write requests corresponding to the thread group have been sent from the direct storage access module of the sending port of the corresponding route, it reads the synchronization set instruction and copies the synchronization set instruction to the direct storage access module of the sending port of the corresponding route for issuance.

[0011] Optionally, the data production entity is a first processing unit, and the data consumption entity is a second processing unit. The first processing unit and the second processing unit communicate through an inter-chip interconnection network. The cross-entity unit write instruction writes the data in the first on-chip memory of the first processing unit to the second on-chip memory of the second processing unit through the inter-chip interconnection network.

[0012] Optionally, the first processing unit is a first acceleration unit, used to execute tasks assigned by the first scheduling unit; the second processing unit is a second acceleration unit, used to execute tasks assigned by the second scheduling unit.

[0013] According to one aspect of this disclosure, a data consumption subject unit is provided, comprising: The second communication control unit is configured to perform the following for a thread group: after receiving a synchronization set instruction on a single route between the data production main unit and the data consumption main unit, it determines whether the previously received write request on that route has been successfully written. After determining that the write request has been successfully written on all routes between the data production main unit and the data consumption main unit, it sets the specific flag bit to a specific value and executes a set-wait instruction and a data usage instruction. The set-wait instruction allows the data usage instruction to be executed after the specific flag bit becomes the specific value. The synchronization set instruction is issued on multiple routes between the data production main unit and the data consumption main unit after all write requests corresponding to all thread groups from the data production main unit to the data consumption main unit have been issued from the data production main unit. The second thread group execution unit is used to enable the thread group to exchange second information with other thread groups, wherein the second information indicates that the specific flag bit becomes a specific value.

[0014] Optionally, the data consumption entity may also include: The second multi-level cache is located between the second thread group execution unit and the second communication control unit. After the specific flag bit becomes a specific value, the second multi-level cache executes a cache invalidation instruction and a wait-for-invalidation-success instruction. The cache invalidation instruction is used to prevent the data stored in the second multi-level cache from being read by subsequent read instructions. The wait-for-invalidation-success instruction is used to wait for the invalidation instruction to be executed, and only after the invalidation instruction has been executed will the data be allowed to be used for instruction execution.

[0015] Optionally, the initial value of the specific flag bit is 0. After receiving the synchronization set instruction on a single route, the second communication control unit determines whether the previously received write request on that route has been successfully written. If the write is successfully completed, a fixed value is added to the specific flag bit. The specific value is equal to the product of the fixed value and the total number of routes.

[0016] Optionally, determining whether a previously received write request on the route has been successfully written includes: if the direct storage access module corresponding to the route has been cleared, then determining that a previously received write request on the route has been successfully written.

[0017] Optionally, the second communication control unit includes: Multiple direct storage access modules, corresponding to multiple receiving ports, are used to cache write requests received from each receiving port. For a thread group, after receiving a synchronization set instruction on the corresponding receiving port, it determines whether the write request previously received on that receiving port from the data production main unit has been successfully written. The receive controller is configured to, for the thread group, after each direct storage access module has determined that all write requests previously received from the data production host unit on the corresponding receive port have been successfully written, read the set-wait instruction and the data use instruction, and execute the set-wait instruction to allow the data use instruction to be executed after the specific flag bit becomes a specific value.

[0018] Optionally, the data production entity is a first processing unit, and the data consumption entity is a second processing unit. The first processing unit and the second processing unit communicate through an inter-chip interconnection network. The cross-entity unit write instruction writes the data in the first on-chip memory of the first processing unit to the second on-chip memory of the second processing unit through the inter-chip interconnection network.

[0019] Optionally, the first processing unit is a first acceleration unit, used to execute tasks assigned by the first scheduling unit; the second processing unit is a second acceleration unit, used to execute tasks assigned by the second scheduling unit.

[0020] According to one aspect of this disclosure, a system-on-a-chip is provided, including a data production main unit as described above and / or a data consumption main unit as described above.

[0021] According to one aspect of this disclosure, a computing apparatus is provided, including a data production unit as described above and / or a data consumption unit as described above.

[0022] According to one aspect of this disclosure, a cross-subject unit write method is provided at the data production subject unit end, including: Execute cross-subject write instructions for thread groups to specific data consuming subject units; This enables the thread group to exchange first information with other thread groups, the first information indicating whether all write requests corresponding to the thread group have been issued from the data production main unit. For a single thread group: If all write requests corresponding to the thread groups used for the cross-subject unit write instruction have been issued from the data production subject unit, a synchronization set instruction is issued on each of the multiple routes between the data production subject unit and the data consumption subject unit. This ensures that after the data consumption subject unit receives the synchronization set instruction on each route, all previously received write requests on that route are successfully written. Only after ensuring successful writing on each route can the written data be used.

[0023] According to one aspect of this disclosure, a cross-subject unit writing method is provided at the data consuming subject unit end, including: For a single thread group, after receiving a synchronization set instruction on a single route between the data production main unit and the data consumption main unit, it is determined whether the previously received write request on that route has been successfully written. After determining that the write request has been successfully written on all routes between the data production main unit and the data consumption main unit, the specific flag bit is set to a specific value, and a set-wait instruction and a data usage instruction are executed. The set-wait instruction is used to allow the data usage instruction to be executed after the specific flag bit becomes a specific value. The synchronization set instruction is issued on multiple routes between the data production main unit and the data consumption main unit after all write requests corresponding to all thread groups from the data production main unit to the data consumption main unit have been issued from the data production main unit. The thread group exchanges second information with other thread groups, the second information indicating that the specific flag bit becomes a specific value.

[0024] This embodiment of the disclosure considers instruction-level write operations, executing cross-subject unit write instructions for individual thread groups. First information is exchanged between thread groups, indicating whether all write requests for the corresponding thread group have been issued from the data production subject unit. Thus, when the first information of each thread group indicates that its write requests have all been issued from the data production subject unit, a synchronization set instruction is issued on each of the multiple routes between the data production subject unit and the data consumption subject unit. When the consumption subject unit receives the synchronization set instructions on these multiple routes, it has actually received all the write requests previously sent by all thread groups on all these routes. Once all these write requests have been completed, it indicates that the write operations of all thread groups have been completed, and the data can be used without errors. This embodiment ensures that all write requests from thread groups have been sent by exchanging first information among thread groups. By issuing a synchronization set instruction on each route after all write requests from thread groups have been sent, the data consumer unit can clearly determine whether the write requests from each thread group on each route have been successfully received and control them to complete the write operation. This ensures that even with instruction-level writing and multiple routes between the data producer and data consumer, the data consumer's use of the data written by the data producer is error-free. Furthermore, when write requests traverse multiple caches and paths within the data producer, since the first information indicates whether all write requests corresponding to the thread group have been sent from the data producer unit, the internal caches and paths do not affect the determination of whether all write requests from that thread group have been received. Therefore, even when the data reaches the data consumer through different routes, caches, and paths, the data consumer's use of the data written by the data producer is error-free. Attached Figure Description

[0025] The above and other objects, features, and advantages of this disclosure will become clearer from the following description of embodiments with reference to the accompanying drawings, in which: Figure 1 This is a structural diagram of a data center used in one embodiment of this disclosure; Figure 2 This is an internal structural diagram of a server (computing device) according to an embodiment of the present disclosure; Figure 3 It shows in Figure 2 The acceleration unit serves as the internal structure and external communication diagram of the data production main unit or data consumption main unit in this embodiment of the disclosure.

[0026] Figure 4 An embodiment according to this disclosure is shown. Figure 3The hierarchical processing and caching structure diagram within the data production unit or data consumption unit.

[0027] Figure 5 A detailed structural diagram of a first or second communication control unit according to an embodiment of the present disclosure is shown.

[0028] Figure 6 An example of compiling source code on a data production subject unit and a data consumption subject unit according to an embodiment of the present disclosure into assembly language is shown.

[0029] Figure 7 A flowchart of a cross-subject unit write method at the data production subject unit end according to an embodiment of the present disclosure is shown.

[0030] Figure 8 A flowchart of a cross-subject unit write method at the data consumption subject unit end according to an embodiment of the present disclosure is shown. Detailed Implementation

[0031] The present disclosure is described below based on embodiments, but it is not limited to these embodiments. In the detailed description of the present disclosure below, certain specific details are described in detail. Those skilled in the art will fully understand the present disclosure even without these details. To avoid obscuring the substance of the present disclosure, well-known methods, processes, and procedures are not described in detail. Furthermore, the accompanying drawings are not necessarily drawn to scale.

[0032] The following terms are used in this document.

[0033] Data production unit: The generator / sender of computational data. Specifically, it writes its stored data to other units (data consumption units) for their use. This unit can be an actual physical unit or an abstract software module, i.e., a software unit. When the unit is a physical unit, it can be a processing unit on a chip, such as a regular central processing unit (CPU), or a special-purpose computing acceleration unit (such as an NPU, GPU, or tensor processor). However, the unit can also be smaller or larger. Smaller units include microcontroller units (MCUs), which are essentially simplified versions of CPUs with some peripheral circuitry integrated onto a single chip. Larger units include systems-on-a-chip (SoCs) integrating multiple processing units and peripheral circuitry onto a single chip. Even larger units can be entire computing devices or even communication terminals. In the case of a communication terminal, the data production unit is the sending terminal, and the data consumption unit is the receiving terminal. When the unit is a software module, it can be a software entity running on hardware or a functional part of a software entity. For the sake of simplicity, the following example primarily uses an acceleration unit as the main body unit. Those skilled in the art, benefiting from the teachings of this disclosure, can extrapolate to other main body units.

[0034] Data consuming entity: The user / receiver of computational data. Specifically, it needs to use data from other entities (data producing entities), and therefore requires other entities to write data to it. Data consuming entities can also take various forms similar to those of data producing entities. The following example primarily uses an acceleration unit as the entity. Those skilled in the art can extrapolate the same logic to other entities.

[0035] Computing device: A device with computing or processing capabilities. It can take the form of a terminal, such as an IoT device, a mobile terminal, a desktop computer, or a laptop computer, or it can take the form of a server or a cluster of servers. In a data center environment, a computing device is a server within the data center.

[0036] System-on-a-Chip (SoC): A chip containing processing units and peripheral circuitry that can be inserted into or replaced from a computing device.

[0037] Scheduling Unit: In addition to performing traditional processing (excluding complex computations such as image processing and various deep learning models) on the computing device or on-chip system, this unit also undertakes the function of scheduling acceleration units. It assigns tasks to the acceleration units. Scheduling units can take various forms, such as Central Processing Units (CPUs), Application-Specific Integrated Circuits (ASICs), and Field-Programmable Gate Arrays (FPGAs).

[0038] Acceleration Unit: A processing unit designed to improve data processing speed in specialized fields (e.g., image processing, various operations for deep learning models, etc.) where traditional processing units are inefficient. Acceleration units include Central Processing Units (CPUs), Graphics Processing Units (GPUs), General-Purpose Graphics Processing Units (GPGPUs), Field-Programmable Gate Arrays (FPGAs), Application-Specific Integrated Circuits (ASICs), and dedicated intelligent acceleration hardware (e.g., Neural Processing Units, NPUs).

[0039] On-chip memory: Storage units located inside the processing unit. Compared to off-chip memory, it is used to directly access frequently used data within the processing unit, improving the efficiency of accessing frequently used data. Off-chip memory, on the other hand, is located outside the processing unit and is shared by all processing units, but accessing it requires going through an external bus, resulting in lower efficiency. It is generally used to store infrequently used data.

[0040] Processing unit: A unit with processing capabilities, including units that perform traditional processing (excluding complex operations such as image processing and various deep learning models), scheduling units that also have scheduling functions, and acceleration units.

[0041] On-chip interconnect network: A communication network that interconnects the various units within a processing unit, such as the internal bus network that interconnects the cores and on-chip caches within an acceleration unit.

[0042] Inter-chip interconnect network: Unlike on-chip interconnect network, it is a network that interconnects different processing units. A processing unit connected to it can directly transmit data to different processing units within the same computing device, or to processing units within different computing devices.

[0043] Command: A set of programming instructions issued by a computer program to perform a specific task. To complete a task, a series of instructions are required, therefore a command contains multiple instructions.

[0044] Instruction: An instruction is the smallest unit of command execution, generally referring to a single statement in computer programming.

[0045] A process is a single execution of a program on a set of data in a computer. It is the basic unit for system resource allocation and scheduling and the foundation of the operating system structure. A program is a description of instructions, data, and their organization; a process is the actual implementation of a program.

[0046] A thread is the smallest unit of computation that an operating system can schedule. It is contained within a process and is the actual unit of operation within that process. A thread refers to a single, sequential flow of control within a process. Multiple threads can run concurrently within a process, each executing different tasks in parallel.

[0047] A warp is a collection of threads derived from a process. These threads execute the same program code, but on different data. This is particularly common in large-scale data processing units (such as NPUs and GPUs) for neural network computing or graphics computing. For example, in matrix multiplication, the code for multiplying and adding the elements of the matrices is fixed; the only difference is that the values ​​of different parts of the matrix need to be substituted into the instruction code for repeated calculations. Therefore, different parts of the elements are assigned to different threads, but the instruction code executed by each thread in the warp is consistent. This allows the entire process to complete quickly.

[0048] A thread block is composed of multiple thread groups that execute in parallel within the thread block. The process is first divided into several thread blocks, and each thread block is further divided into several thread groups.

[0049] Acceleration Unit Layered Processing Architecture: Since large-scale data processing typically involves dividing processes into thread blocks, and then further subdividing these thread blocks into thread groups for processing, the corresponding hardware instruction set executing these processes also often needs to be configured with a layered processing architecture, such as... Figure 4 As shown, the instruction execution units in its architecture are divided into processing execution units 342 (PU, hereinafter referred to as thread group execution units), computing execution units 341 (CU, hereinafter referred to as thread block execution units), and computing execution elements 340 (CE, i.e. acceleration unit cores) in order of hierarchy from low to high.

[0050] Thread group execution unit 342: such as Figure 4 As shown, the lowest level instruction execution unit in the accelerated unit hierarchical processing architecture is the processing execution unit (PU). It performs processing at the level of the aforementioned thread groups (warps). It can run one thread group, or multiple thread groups sequentially, but the same thread group must be executed within a single PU and cannot run across PUs. Because of this characteristic, and to avoid confusion with the aforementioned processing unit concept, the PU in this embodiment is named the thread group execution unit. A thread group typically includes 16 threads.

[0051] Thread block execution unit 341: such as Figure 4As shown, in the hierarchical processing architecture of the acceleration unit, the layer above the thread group execution unit 342 is the thread block execution unit 341, or CU. It includes multiple thread group execution units 342 and a first-level cache 343 shared by the thread group execution units 342. A thread block execution unit 341 can run one thread block, or it can run multiple thread blocks, but the same thread block must run in one CU and cannot run across CUs. Figure 4 In this context, a thread block execution unit 341 includes four thread group execution units 342, but may also include other numbers of thread group execution units 342. The data cached in the first-level cache 343 is shared by all thread group execution units 342 within the thread block execution unit 341. Data written from one thread group execution unit 342 to the outside can be cached in the first-level cache 343 first. Then, if other thread group execution units 342 within the same thread block execution unit 341 need the data, they can retrieve it directly from this cache without going through layers of caching to the original target body to which the data was originally written.

[0052] Acceleration Unit Core 340: such as Figure 4 As shown, in the hierarchical processing architecture of the acceleration unit, the layer above the thread block execution unit 341 is the acceleration unit core 340, i.e., CE. An acceleration unit 230 typically contains 16 acceleration unit cores 340. An acceleration unit core 340 includes multiple (typically four, but other numbers are possible) thread block execution units 341 and a shared second-level cache 344. Data written from multiple thread block execution units 341 to the outside can be cached in the second-level cache 344. Then, if other thread block execution units 341 within the same acceleration unit core 340 need this data, it can be retrieved directly from this cache without going through multiple layers of caching to the original target body to which the data was originally written.

[0053] Synchronization: Since the process of writing data to the target processing unit is executed in thread groups, and the data written by one thread group is related to the data written by other thread groups in the same process, the target processing unit can only use this data after all the associated data of all thread groups in the process has been written; otherwise, it may read incorrect data. Synchronization is a mechanism used to achieve the above purpose, enabling the data consuming unit to know that the data associated with all thread groups in the writing process has been written to it, thereby correctly initiating data reading.

[0054] Routing: Data may arrive between two nodes (e.g., a data production unit and a data consumption unit) via different intermediate nodes. On each pair of adjacent nodes, since each node has multiple ports, different sending and receiving ports may also be selected for transmission. Therefore, data transmission between two nodes forms different paths, and each path is a route.

[0055] Path: The route through which data travels between different components within a node or between different nodes. Its meaning is broader than routing.

[0056] Synchronization set instruction: An instruction set on the data production main unit side to set a flag bit in order to achieve synchronization.

[0057] Multi-level caching: Accelerating the processing of different levels of cache in a hierarchical architecture, such as... Figure 4 The first-level cache 343, the second-level cache 344, and the last-level cache 345 are shown.

[0058] Emptying: The process of removing cached data from the cache and transferring it to the target body where the data actually needs it, so that the cache is empty in the end.

[0059] Set-and-wait instruction: This instruction is used to check if a bit has been set successfully (a specific flag becomes a specific value) and allows subsequent instructions to be executed after the bit has been set successfully.

[0060] Synchronous set wait instruction: An instruction that waits for the synchronous set instruction to complete, indicating that no subsequent instruction can be executed before the previous synchronous set instruction.

[0061] Application Environment of This Disclosure This disclosure proposes a communication control scheme for a data production unit and a data consumption unit. The entire communication control scheme is relatively universal and can be used in various scenarios. For example, it can be used for data communication control between processing units in servers within a data center, data communication control between MCUs in a microcontroller unit (MCU) architecture, data communication control between processing units in an IoT terminal, data communication control between nodes in the Internet, data communication control between different software modules within the same hardware, and data communication control between different functional parts within the same software module, etc. Due to the numerous application scenarios, the following description uses only a data center as an example. However, those skilled in the art should understand that this scheme is also applicable to other application scenarios, and the implementation details in other scenarios can be easily conceived from the description of this disclosure's embodiments in a data center scenario.

[0062] Data Center Data centers are globally collaborative networks of specialized equipment used to transmit, accelerate, display, compute, and store data information on the internet infrastructure. In the future, data centers will become a key competitive asset for businesses. With the widespread application of data centers, data security is receiving increasing attention.

[0063] In traditional large data centers, the network architecture is typically as follows: Figure 1 As shown, this is the hierarchical inter-networking model. This model includes the following components: Server 140: Each server 140 is a processing and storage entity in the data center. The processing and storage of a large amount of data in the data center are all done by these servers 140.

[0064] Access Switch 130: Access Switch 130 is used to connect Server 140 to the data center. One Access Switch 130 connects multiple Servers 140. Access Switches 130 are typically located at the top of the rack, so they are also called Top of Rack switches, and they physically connect the servers.

[0065] Aggregation Switch 120: Each aggregation switch 120 connects to multiple access switches 130 and provides other services such as firewall, intrusion detection, network analysis, etc.

[0066] Core Switch 110: Core Switch 110 provides high-speed forwarding for packets entering and leaving the data center and provides connectivity for Aggregation Switch 120. The entire data center network is divided into an L3 routing network and an L2 routing network. Core Switch 110 typically provides a flexible L3 routing network for the entire data center network.

[0067] Typically, aggregation switch 120 serves as the boundary between L2 and L3 layer routing networks. Below aggregation switch 120 is the L2 network, and above it is the L3 network. Each aggregation switch group manages one Point of Delivery (POD), and each POD contains an independent VLAN network. Server migration within a POD does not require modification of IP addresses and default gateways, as one POD corresponds to one L2 broadcast domain.

[0068] The aggregation switch 120 and access switch 130 typically use the Spanning Tree Protocol (STP). STP ensures that only one aggregation switch 120 is available for a given VLAN network; other aggregation switches 120 are only used in case of failure (dashed lines in the diagram above). In other words, horizontal scaling is not possible at the aggregation switch 120 level, because even if multiple aggregation switches 120 are added, only one will be operational.

[0069] Computing device In the era of artificial intelligence, data center servers (computing devices) 140 need to handle a large number of neural network operations. Neural network operations require large-scale parallel computing. Traditional central processing unit (CPU) architectures result in control units and storage units occupying a significant portion of the space, while the computing units themselves occupy insufficient space. Therefore, while efficient for logic control, they are inefficient for large-scale parallel computing. In this context, servers (computing devices) 140 with numerous acceleration units 230 are employed. Figure 2 As shown, this computing device 140 includes a memory 210, a scheduling unit cluster 270, and an acceleration unit cluster 280 connected by a bus. The scheduling unit cluster 270 includes multiple scheduling units 220. The acceleration unit cluster 280 includes multiple acceleration units 230.

[0070] As mentioned above, traditional central processing units (CPUs) are highly efficient in logic control but inefficient in large-scale parallel computing. Acceleration units 230 are designed to improve computational speed for different functions and domains, such as a dedicated processing unit (NPU) for accelerating neural network model computations. It employs a data-driven parallel computing architecture to handle the numerous operations (e.g., convolution, pooling) of each neural network node. Since the data and intermediate results of these operations are closely linked and frequently used throughout the computation process, existing CPU architectures suffer from inefficiencies due to the small memory capacity within the CPU core, necessitating frequent access to external memory. By employing this dedicated acceleration unit for neural network model computations, on-chip memory with sufficient storage capacity for neural network computations is provided, avoiding frequent access to external memory and significantly improving processing efficiency and computational performance. In this embodiment of the disclosure, the acceleration unit 230 may be embodied as a processing unit (NPU), graphics processing unit (GPU), application-specific integrated circuit (ASIC), or field-programmable gate array (FPGA) specifically designed for neural network operations.

[0071] The scheduling unit 220 is a processing unit that schedules the acceleration units 230 and allocates the sequences of instructions to be executed to each acceleration unit 230. It can take various forms, such as a central processing unit (CPU), an application-specific integrated circuit (ASIC), or a field-programmable gate array (FPGA). The scheduling unit 220 can also act as a central processing unit performing traditional logic control operations. The scheduling unit 220 reads instructions to be executed from the memory 210. If the instruction is a logic control instruction, the scheduling unit 220 executes it itself to complete the corresponding logic control operation. If the instruction requires processing by the acceleration unit 230, it is allocated to the acceleration unit 230 for execution. The acceleration unit 230 is subject to scheduling by the scheduling unit 220.

[0072] Take neural network operations as an example. The memory 210 stores various neural network models, including the nodes and weight data of these models. These neural network models are used when needed. Figure 2 A scheduling unit 220 is deployed to an acceleration unit 230. Specifically, the scheduling unit 220 can send the addresses of the model's parameters (such as the weights of each node) in memory 210 to the acceleration unit 230 via instructions. When the acceleration unit 230 actually uses the neural network model for computation, it directly addresses these parameters (e.g., weights) in memory 210 based on their addresses and temporarily stores them in its on-chip memory. When the acceleration unit 230 actually uses the neural network model for computation, the scheduling unit 220 also sends the model's input parameters to the acceleration unit 230 via instructions, which are then temporarily stored in the acceleration unit 230's on-chip memory. In this way, the acceleration unit 230 can perform inference computation based on these input parameters and the model's parameters (e.g., weights).

[0073] Acceleration unit Figure 3 The diagram illustrates the internal structure and external communication of an acceleration unit 230 according to an embodiment of the present disclosure.

[0074] like Figure 3 As shown, the acceleration unit 230 includes a bus channel 391, a direct memory access module 390, on-chip memory 370, a switching module 320, various ports 310, a core 340, a command processor 360, a last-level on-chip cache 345, a communication control unit 330, and an on-chip interconnect network 380. The core 340, command processor 360, on-chip cache 345, and communication control unit 330 are interconnected and communicate with each other via the on-chip interconnect network 380.

[0075] Bus channel 391 is the channel through which instructions enter and exit the acceleration unit 230. Instructions assigned to the acceleration unit 230 for execution by the scheduling unit 220 enter through bus channel 391.

[0076] The Direct Memory Access (DMA) module 390 is a module provided by some computer bus architectures that enables data to be written directly from an external device to the computer motherboard's memory. This method significantly improves data access efficiency compared to the previous method where all data transfers between devices had to go through a scheduling unit. Because of this mechanism, the core 340 of the acceleration unit 230 can directly access memory 210 to read parameters (such as the weights of each node) from the neural network model, greatly improving data access efficiency.

[0077] The command processor 360 is a device that divides the commands sent from the scheduling unit 220 to the acceleration unit 230 into instructions and allocates them to the cores 340 for execution. After the sequence of instructions to be executed allocated by the scheduling unit 220 enters from the bus channel 391, it is cached in the command processor 360. The command processor 360 then selects the cores 340 and allocates the instructions to the corresponding cores 340 for execution. The command processor 360 can, for example, use a load balancing algorithm to consider the processing load of each core 340 and allocate instructions to the cores 340 with lighter processing loads.

[0078] Core 340 is the core component of the acceleration unit 230, processing instructions internally. It receives instructions from the command processor 360 and executes them according to thread groups. Its internal structure is as follows... Figure 4 The hierarchical structure is shown. After instructions are divided into thread groups for execution, each thread group executes upwards along the hierarchical structure from the end (i.e., the first thread group execution unit 342) (for example, passing write requests upwards level by level). Regarding... Figure 4 The hierarchical structure and instruction execution will be discussed in detail later. Figure 4 Detailed explanation.

[0079] On-chip memory 370 is a shared memory among multiple cores 340 within the acceleration unit 230. The command processor 360 distributes the sequence of instructions to be executed to different cores 340 for execution. When different cores 340 execute different instructions, because the instructions are interrelated, a core 340 executing a certain instruction may need to use the results of other cores 340 executing other instructions. Therefore, some shared storage space is needed within the acceleration unit 230 to store data shared by multiple cores 340. On-chip memory 370 is such shared storage space. Figure 3 In this example, one acceleration unit 230 has four on-chip memory modules 370, which is merely an example. Those skilled in the art will understand that other numbers of on-chip memory modules 370 may also be used.

[0080] In the acceleration unit, the cache is as follows: Figure 4 The diagram shows a hierarchical structure, which will be discussed later. Figure 4 This will be explained in detail. The last-level cache 345 is the highest-level cache in the acceleration unit 230. It can store some data that the cores 340 frequently use. This data should originally be stored in the on-chip memory 370, but because it is frequently used, it can be stored in the last-level cache 345, which is closer to the cores 340 and has a faster read speed, so as to enable the cores 340 to access it quickly.

[0081] The on-chip interconnect network 380 is an interconnect network that enables high-speed communication between the command processor 360, cores 340, last-level cache 345, and communication control unit 330. Through it, the command processor 360 can allocate instructions to each core 340 more quickly, and each core 340 can also quickly establish communication with the communication control unit 330 when it needs to read or write data to other acceleration units 230, so that the communication control unit 330 can write data out or read data from other acceleration units 230.

[0082] Inter-chip interconnect network 150 is a network that interconnects acceleration units 230. Through it, one acceleration unit 230 can directly transmit data to other acceleration units 230 within the same computing device 140, or to acceleration units 230 within different computing devices 230. Since neural network operations involve large-scale parallel processing, each operation may be executed by different acceleration units 230 within the same computing device 140, or even by acceleration units 230 within different computing devices 140. Because instructions are often correlated, one acceleration unit 230 often needs the result of another acceleration unit 230 executing another instruction when executing one instruction; therefore, data exchange is necessary between acceleration units 230. In conventional technology, this data exchange is performed through a conventional inter-device interconnect network 160 between computing devices 140. This network can only transmit data at the computing device 140 level and cannot directly transmit data at the level of individual acceleration units 230. When an acceleration unit 230 within a computing device 140 needs to transfer data to another acceleration unit 230 within another computing device 140, it first sends the data to be transferred to a network interface card (NIC) 290 via the bus within the computing device 140. The NIC 290 is the entry and exit point for data entering and exiting the computing device 140. The NIC 290 then sends the data to be transferred to a conventional inter-device interconnection network (IoC network) 160, which then sends or writes the data to another acceleration unit 230 within that computing device 140 via the NIC. Thus, this process involves network transmission via the IoC network 160, bus transmission within the computing device 140 sending the data, and bus transmission within the computing device 140 receiving the data. This process is complex, and the network transmission via the IoC network 160 is slow, significantly impacting the efficiency of data transfer between acceleration units 230 of different computing devices 140. To improve the transmission efficiency, an inter-chip interconnection network 150 is used to directly interconnect the acceleration units 230 of each computing device 140. Since the acceleration units 230 of each computing device 140 are directly interconnected through this network, and the network has a high transmission speed, the efficiency of data transmission between the acceleration units 230 of different computing devices 140 is greatly improved.

[0083] Port 310 is the entry and exit point for the acceleration unit 230 to exchange data with the inter-chip interconnection network 150. An acceleration unit 230 may have multiple ports 310. Each port 310 may simultaneously serve as a transmitting port for sending data to the inter-chip interconnection network 150 and a receiving port for receiving data from the inter-chip interconnection network 150, or a portion of it may be dedicated to sending data to the inter-chip interconnection network 150 and another portion to receiving data from the inter-chip interconnection network 150.

[0084] The communication control unit 330 controls the transmission of data and instructions to other acceleration units 230 via the inter-chip interconnection network 150, and controls the reception of data and instructions from other acceleration units 230 via the inter-chip interconnection network 150. When controlling the transmission of data and instructions to other acceleration units 230 via the inter-chip interconnection network 150, it determines the route for data and instruction transmission and controls the transmission of data and instructions according to the determined route. When controlling the reception of data and instructions from other acceleration units 230 via the inter-chip interconnection network 150, it ensures that the received data is correctly written into the memory space of the current acceleration unit 230, while ensuring that instructions with sequential requirements entering from different ports 310 are executed in the required order.

[0085] The switching module 320 is a module that switches data to be sent to the inter-chip interconnection network 150 through each port 310 to the corresponding port 310 for transmission through that port 310, and aggregates data received from each port 310. The switching module 320 contains a routing table. Based on the route determined by the communication control unit 330 for the data to be sent, the switching module 320 looks up the routing table, determines the corresponding sending port 310, and then sends the corresponding data through the sending port 310.

[0086] Hierarchical structure of instruction execution unit and cache in acceleration unit 230 As mentioned above, large-scale data computation typically involves dividing the process into thread blocks, which are then further divided into warps for processing. In large-scale data computation, the computation rules and instructions are relatively fixed; the only difference lies in the repetitive processing of large amounts of different data. Therefore, different data are assigned to multiple threads that execute identical instructions. Since these threads execute the same instructions but deal with different data, they form a thread warp. Multiple thread warps then form a thread block, and many thread blocks complete the entire process. Correspondingly, the hardware instruction execution units that execute these units are also hierarchical, such as... Figure 4 As shown, the lowest level of execution unit is the thread group execution unit (PU) 342, the next level of execution unit is the thread block execution unit (CU) 341, and the next level of execution unit is the acceleration unit core (CE) 340.

[0087] For example Figure 4As shown, PU 342 is the lowest-level, most basic execution unit. Threads within the same group must run within the same PU 342 and cannot run across PU 342s, but a single PU 342 can execute multiple thread groups sequentially. The layer above PU 342 is CU 341. It includes multiple PU 342s and a shared first-level cache 343. A CU 341 can run one thread block, or multiple thread blocks, but the same thread block must run within the same CU 341. The layer above CU 341 is CE 340. A CE 340 includes multiple CU 341s and a shared second-level cache 344. Finally, multiple CE 340s share a last-level cache 345.

[0088] If an acceleration unit 230 receives a large-scale data write task to write data to another acceleration unit 230 (the target acceleration unit), the write task is decomposed into a series of thread groups within the acceleration unit 230 as described above. Each thread group includes multiple parallel threads, each executing the same write instruction for different data, i.e., issuing write requests for different data to the same target acceleration unit 230. Correspondingly... Figure 4 In the hierarchical structure, such a thread group is executed by a PU 342. PU 342 issues a write request for the corresponding data. If the write request is cacheable, it is cached in the first-level cache 343 and not passed up the hierarchy. At this point, although it is already stored in the first-level cache 343 and has not yet reached its intended destination, it is not passed to the destination temporarily because traditionally it is assumed that since the data can be found in the cache even though it has not reached its destination, it can still be retrieved from the cache later. However, in some applications, although the write request carries a destination address, it does not actually care where it is written; as long as the data can be retrieved, the requirement is met. In this case, if the data is cached in the first-level cache 343, other PUs 342 of the same CU 341 can quickly read the data directly from the first-level cache 343 without going through layers of caching to the destination, greatly reducing subsequent data read time. Furthermore, for write requests with strict requirements on the destination, a forced write flag can be set in the write request. For such write requests, storage at intermediate levels of cache is prohibited; the data must be written to the actual destination. In this way, both the actual writing of data with strict requirements on the write destination and the fast storage and retrieval of general data can be achieved.

[0089] If a write request is not cached in the first-level cache 343, it enters the second-level cache 344 shared by multiple CUs 341 in the CE 340. If the write request is also not cached in the second-level cache 344, it is sent to its actual destination acceleration unit 230 via the communication control unit 330.

[0090] Internal structure of the communication control unit Figure 5 The internal structure of the first communication control unit 330 in the data production unit or the second communication control unit 330 in the data consumption unit is shown. It includes an instruction distributor 530, a receiver controller 540, multiple direct memory access modules 520 corresponding one-to-one with multiple ports 310, and a switching module 510. Figure 5 The diagram also shows external devices such as the core 340, on-chip interconnect network 380, and on-chip memory 370, in order to clearly illustrate the communication relationship between the communication control unit 330 and the outside world.

[0091] Although Figure 5 Not shown in the text Figure 5 The communication control unit 330 may also include a command distributor for controlling command-level writes and reads. Since this disclosure primarily concerns instruction-level writes, the command distributor is not shown to avoid obscuring the focus of this disclosure. Unlike the command distributor, which reads commands sequentially from the command storage ring (not shown), in instruction-level writes, the commands have already been decomposed and executed in the core 340. What the core 340 delivers to the communication control unit 330 via the on-chip interconnect network 380 is the instruction, not the command itself. Therefore, the instruction distributor 530 directly assigns the instruction lookup routing table to different direct memory access modules 520, issuing from the corresponding routes.

[0092] Data transmission between two main units may occur via different intermediate main units. Furthermore, between any two adjacent main units, since each main unit has multiple ports 310, different sending and receiving ports 310 may be selected for transmission. Therefore, data transmission between the data producing main unit and the data consuming main unit will have multiple paths, each a route. The communication control unit 330 stores a pre-configured routing table to record routes from the data producing main unit to various data consuming main units. Given the required data consuming main unit, multiple routes to the required data consuming main unit can be determined by looking up the routing table.

[0093] The switching module 320 is a module that switches data to be sent to the inter-chip interconnection network 150 through each port 310 to the corresponding port 310 for transmission through that port 310, and also aggregates the data received from each port 310. The switching module 320 also has a routing table. Based on the route that the data to be sent by the communication control unit 330 will take, the switching module 320 looks up the routing table, determines the corresponding sending port 310, and then sends the corresponding data through the sending port 310.

[0094] Each direct storage access module 520 corresponds to a port 310 and maintains two sequences: one sequence stores write requests to be sent from the corresponding port 310, and the other sequence stores write requests received from the corresponding port 310. Write requests in the sequence to be sent are sent sequentially through the corresponding port 310 in the order they appear in the sequence. Write requests received from port 310 are queued in the order they are received and then processed by the receive controller 540 in the order they are queued.

[0095] The receiver controller 540 processes the instructions entering from the direct memory access module 520 sequentially. For example, for a write request, the data targeted by the write request is written to the on-chip memory 370 of the data consumer unit.

[0096] Cross-subject unit write process of this disclosure embodiment When performing instruction-level writes across different data processing units, a synchronization problem arises. As mentioned above, a large-scale data write task to the data consuming unit is executed by multiple thread groups. Since the data written by each thread group is related, the data successfully written to the data consuming unit by one thread group cannot be used immediately. The target processing unit can only use this data after all related data in all thread groups within the task has been written; otherwise, it may use incorrect data. This is the synchronization problem.

[0097] In command-level writes, synchronization is often addressed by setting synchronization primitives on the data production unit side and synchronization wait primitives on the target processing unit side. After all data to be written from the data production unit to the data consumption unit has been sent from the data production unit, if there is only one route between the two units, a synchronization primitive is sent to the data consumption unit along that route, setting a flag. The data consumption unit sets a synchronization wait primitive; upon detecting the flag being set, subsequent data is allowed to use the data written by the data production unit. Synchronization is achieved in this way. However, since command-level writes are performed by thread group, it is difficult to determine holistically whether all data to be written from the data production unit to the data consumption unit has been sent from the data production unit. Furthermore, there are often multiple routes between the data production unit and the data consumption unit. Moreover, data within the data production unit may also reach its sending port 310 via different caches and paths. Therefore, write requests from different thread groups may arrive at the data consumption unit through different routes, caches, and paths. Write requests from the same thread group may arrive at the data consumer unit through different routes, caches, and paths. The routes, caches, and paths used by the synchronization primitives may also differ from those used by individual write requests. This can lead to the synchronization primitives arriving at the data consumer unit first, while write requests from some thread groups may not reach the target processing unit even after the synchronization primitives have arrived due to congestion on their respective routes, caches, or paths. In this situation, the data consumer unit detects the flag being set and considers synchronization complete, allowing execution of subsequent data usage statements. However, because some thread group write requests may not have yet reached the data consumer unit, data read errors may occur during the execution of these data usage statements.

[0098] This disclosed embodiment is designed for situations where there is instruction-level writing and multiple routes exist between the data producer and the data consumer (even within the data producer, write requests need to pass through different caches and paths). It can ensure that the data consumer unit does not make mistakes in using the data written by the data producer unit when there is instruction-level writing and multiple routes exist between the data producer and the data consumer (as well as multiple caches and multiple paths within the subject).

[0099] To address the aforementioned issues, this embodiment considers instruction-level write operations and executes cross-subject unit write instructions on a thread-by-thread basis. First information is exchanged between thread groups, indicating whether all write requests for each thread group have been issued from the data production subject unit. Thus, when the first information for each thread group indicates that its write requests have been issued from the data production subject unit, it means that all write requests involved in the entire cross-subject unit write instruction have been successfully issued from the data production subject unit. The remaining issue is ensuring that the data is successfully received and written by the data consumption subject unit, thereby guaranteeing that subsequent data usage instructions can only use this data after it has been fully written. To this end, a synchronization set instruction is issued on each of the multiple routes between the data production subject unit and the data consumption subject unit. When the data consuming unit receives the synchronization set instructions on each of these multiple routes, it has actually received all the write requests previously sent by all thread groups on all these multiple routes. Once all these write requests have been completed, it indicates that the write operations of all thread groups have been executed, and the data can be used without errors. This successfully ensures that even with instruction-level writes and multiple routes between the data producing and consuming units, the data consuming unit can still use the data written by the data producing unit without errors. Specifically, exchanging first information between thread groups ensures that all thread group write requests have been sent; by issuing a synchronization set instruction on each route after all thread group write requests have been sent, the data consuming unit can clearly determine whether the write requests of each thread group on each route have been successfully received and control their completion. Furthermore, when write requests pass through multiple caches and paths within the data producing unit, since the first information indicates whether all write requests corresponding to the thread group have been sent from the data producing unit, the internal caches and paths do not affect the determination of whether all write requests of that thread group have been received.

[0100] The above process will be described in detail below.

[0101] In one embodiment, the first thread group execution unit 342 in the data production main unit executes cross-main unit write instructions to a specific data consumption main unit for the thread group, and causes the thread group to exchange first information with other thread groups. The first information indicates whether all write requests corresponding to the thread group have been issued from the data production main unit. Whether all write requests corresponding to the thread group have been issued from the data production main unit can be determined by whether an internal confirmation response is received from the first communication control unit 330. The first thread group execution unit 342 issues all write requests for a thread group. These write requests first pass through the first-level cache 343. If these write requests are mandatory writes, they will not stay in the first-level cache 343, but will enter the second-level cache 344 of the core 340. If these write requests are mandatory writes, they will not stay in the second-level cache 344. If these write requests are to other acceleration units 230 that are data consumption main units, they enter the first communication control unit 330 and are sent to the other acceleration units 230. Simultaneously, the first communication control unit 330 sends back an internal acknowledgment response, which is then relayed through the second-level cache 344 and the first-level cache 343 back to the first thread group execution unit 342. The first thread group execution unit 342, knowing that all write requests corresponding to the thread group have been sent from the data production main unit, begins exchanging first information with other thread groups. For example... Figure 6 As shown, `syncthreads()` is the code statement for exchanging initial information between thread groups, which compiles into assembly language as the assembly statement `S_BLKSYN` on the left. The assembly statement `S_WAIT.ST|LD` on the left indicates that the write requests corresponding to the thread group have all been issued from the data production unit.

[0102] The first communication control unit 330 executes the following for a single thread group: If all write requests corresponding to the thread groups used for the cross-subject unit write instruction have been sent from the data production subject unit, a synchronization set instruction is issued on each of the multiple routes between the data production subject unit and the data consumption subject unit. This instruction instructs the data consumption subject unit to consume data only after all write requests received on the multiple routes have completed data writing. The significance of issuing a synchronization set instruction on each route is as described above. Since write requests from different thread groups may have reached the data consumption subject unit via different routes, and different write requests from the same thread group may also have reached the data consumption subject unit via different routes, a synchronization set instruction must be issued on each route after all write requests from each thread group have been successfully sent. Each route is tracked separately. The data consumption subject unit checks whether the synchronization set instruction has been received for each route, and then determines whether the write requests sent before the synchronization set instruction on that route have completed writing. This ensures an accurate determination of whether all write requests for the entire cross-subject unit write instruction have been completed without omission.

[0103] like Figure 6 As shown, the assembly statement that issues the synchronization set instruction on each route is S_SIGNAL.SYSmarker.val, which is generated by... Figure 6 This is obtained by compiling the statement __signal_system(marker,val). Figure 6 In addition to compiling the __signal_system(marker,val) assembly statement, it is also compiled into two assembly statements: / / CACHE.CLN and / / S_WAIT.CACHE, which are detailed below.

[0104] exist Figure 6The statement `if(tid=0)` indicates that subsequent synchronization setting instructions are executed for thread group 0 (or even just thread 0 of thread group 0). This verifies that the first communication control unit 330 only performs subsequent operations such as sending synchronization setting instructions on each route for a single thread group. Sending synchronization setting instructions on each route, and the subsequent instructions to flush the first multi-level cache, incur significant execution overhead; requiring each thread group to perform these operations would be unacceptable. To reduce execution overhead, this embodiment of the disclosure allows a single thread group (or even just one thread within a thread group) to execute subsequent operations such as sending synchronization setting instructions and flushing the cache. The subsequent exchange of third information (indicating that synchronization setting instructions have been issued on the multiple routes and / or the first multi-level cache has been flushed) between this thread group and other thread groups informs each thread group that the synchronization setting instructions have been issued and / or the first multi-level cache has been flushed, thereby greatly reducing execution overhead. Figure 6 The ___syncwarp() function shown is the statement used to exchange third-party information between different thread groups.

[0105] As described above, the first communication control unit 330 includes components such as an instruction distributor 530, multiple direct memory access modules 520, and a switching module 510. The process of sending synchronization set instructions on each route can be executed by the aforementioned components. The instruction distributor 530 searches the routing table based on the data consuming entity to be written, determining multiple possible routes between the data producing entity and the data consuming entity. Then, the instruction distributor 530 assigns the write requests corresponding to the thread group to the corresponding routes. When assigning routes, principles such as load balancing can be used to improve the execution efficiency of the entire write task. For example, 95 write requests can be assigned to 10 routes, with 10 write requests assigned to 5 routes and 9 write requests assigned to the other 5 routes. Since each route corresponds to a unique sending port 310, and this port 310 is equipped with a unique direct memory access module 520, all write requests corresponding to the thread group can be assigned one by one to the corresponding direct memory access module 520. Each direct storage access module 520 maintains a sequence of write requests to be sent from its corresponding transmit port 310. When this sequence is empty, the direct storage access module 520 can report this to the instruction allocator 430. After receiving the above reports for all direct storage access modules 520 that have allocated write requests for all thread groups, the instruction allocator 430 can determine that the write requests for all thread groups have been sent from the direct storage access modules 520 of the corresponding route's transmit port. At this time, the first communication control unit 330 can read the synchronization set instruction and copy the synchronization set instruction to the direct storage access modules 520 of the transmit ports 310 of each reported route for issuance.

[0106] like Figure 4 As shown, the first multi-level cache includes a first-level cache 343, a second-level cache 344, and a last-level cache 345. As described above, if a write request issued by the first thread group execution unit 342 for a thread group is not a mandatory write, it will be cached when passing through the first-level cache 343 and will not reach the first communication control unit 330. While this greatly improves the reading speed by allowing other thread group execution units 342 within the same thread block execution unit 341 to read the data directly from the first-level cache 343 without needing to go to the data consumer unit, it is extremely detrimental to ensuring the actual writing of all data by the data consumer unit and preventing errors in subsequent data usage instructions. Therefore, in one embodiment, the first multi-level cache (specifically, the first-level cache 343 and the second-level cache 344) also executes a flush instruction and a wait-for-flush-success instruction, i.e., the above-mentioned __signal_system(marker,val) compiled into... Figure 6The assembly statements / / CACHE.CLN and / / S_WAIT.CACHE are used in the code. The flush command is used to flush the data stored in the cache and send it from the first communication control unit 330 to the data consumption unit via a higher-level cache. Specifically, for write requests cached in the first-level cache 343, the data is flushed from the first-level cache 343 and continuously flushed into the first communication control unit 330 before being sent out; for write requests cached in the second-level cache 344, the data is flushed from the second-level cache 344 and flushed into the first communication control unit 330 before being sent out. The wait-for-flush-success command is used to wait for all the data flushed from the cache to be successfully sent out from the data production unit. It can be determined by receiving an internal confirmation response from the first communication control unit 330 as described above.

[0107] The above embodiments further consider the impact of cache write requests in multi-level caching on non-real writes to data consumer units, and further ensure that the data consumer does not make mistakes in using the data written by the data producer when there are instruction-level writes and multiple routes exist between the data producer and the data consumer.

[0108] In one embodiment, after the first communication control unit 330 issues synchronization set commands on the multiple routes respectively, it executes a synchronization set wait command, such as... Figure 6 The / / S_WAIT.ST following S_SIGNAL.SYS marker.val. This synchronization set wait instruction is used to indicate that no subsequent instruction can be executed before the synchronization set instruction S_SIGNAL.SYS marker.val is issued on the multiple routes.

[0109] On the data consumption main unit side, the second communication control unit 330 performs the following for a thread group: after receiving a synchronization set instruction on a single route between the data production main unit and the data consumption main unit, it determines whether the previously received write request on that route has been successfully written, and after determining that the write request has been successfully written on all routes between the data production main unit and the data consumption main unit, it sets the specific flag bit to a specific value and executes a set wait instruction and a data use instruction. The set wait instruction is used to allow the data use instruction to be executed after the specific flag bit becomes a specific value.

[0110] When the communication control unit 330 shown in Figure 5 serves as the second communication control unit 330 of the data consumption subject unit, each port 310 is a receiving port 310. Correspondingly, the write requests received from the corresponding receiving port 310 are arranged in the received write request sequence in the direct storage access module 520. If the sequence is not empty, it indicates that the write request received on the corresponding route of that port 310 has not been successfully written to the data consumption subject unit. If the sequence is empty, it is determined that the write request received on the corresponding route of that port 310 has been successfully written to the data consumption subject unit. In this way, it can be determined whether the write request previously received on that route has been successfully written.

[0111] For a given thread group, after receiving a synchronization set instruction on the corresponding receive port 310, the direct storage access module 520 can determine whether previously received write requests from data production entities on that receive port have been successfully completed by checking if the sequence of received write requests is empty. An empty sequence indicates that all previously received write requests on that route have been successfully completed, including not only those from the current data production entity but also those from other data production entities. To specifically distinguish write requests from a particular data production entity, each write request can carry a flag indicating its sender. In this case, the success of previously received write requests from data production entities on the receive port can be determined by checking if any write requests with the corresponding flag are absent from the received write request sequence.

[0112] In one embodiment, one way to change the specific flag bit to a specific value is to: let the specific flag bit (in) Figure 6 The initial value of the marker is 0. Whenever a direct storage access module 520 determines that a write request previously received on the corresponding receive port 310 from the data production host unit has been successfully completed, it increments the specific marker by a fixed value ( ). Figure 6 The value in the middle is val). A direct storage access module 520 corresponding to a route determines a successful write and adds a fixed value to a specific flag. There are multiple routes between the data production unit and the data consumption unit. Figure 6 (This refers to num_MinLink_Src2Dst routes). Thus, after the direct storage access modules 520 corresponding to these multiple routes have confirmed successful writes, a specific flag should be changed to a specific value, which is equal to the product of the fixed value and the total number of the multiple routes (in...). Figure 6 The value in the middle is val·num_MinLink_Src2Dst).

[0113] exist Figure 6In the code, the set-wait instruction is while(...acquire_sys(marker)!=val*num_MinLink_Src2Dst), which compiles to assembly language as S_WAIT.LD. Its meaning is that after the marker becomes val*num_MinLink_Src2Dst, the data is executed using the instruction. Figure 6 In this context, the data usage instructions are LD…, &Data.

[0114] For this thread group, the receive controller 540 reads the set-wait instruction and the data use instruction, and executes the set-wait instruction to allow the data use instruction to execute after the specific flag bit becomes a specific value. The specific flag bit becoming a specific value means that each direct memory access module 520 has determined that the write requests previously received on the corresponding receive port from the data production host unit have been successfully written.

[0115] Furthermore, the data consuming unit, like the data producing unit, may also have a second multi-level cache similar to the first multi-level cache. This second multi-level cache includes a first-level cache 343, a second-level cache 344, and a last-level cache 345. Data to be written from the data consuming unit to its on-chip memory 370 may be cached by one of these caches during the propagation through the second multi-level cache, and thus not sent down to the on-chip memory 370. Therefore, when a data usage instruction is executed in the data consuming unit, the instruction first checks the first-level cache 343, the second-level cache 344, and the last-level cache 345 for cached data. If not found, it reads from the on-chip memory 370. If found, it reads directly from the cache. However, this process employs a complex mechanism to ensure that the data to be written from the data producing unit to the data consuming unit is ultimately written to the on-chip memory 370 of the data consuming unit. If subsequent data is read from the second multi-level cache using instructions, the result may be incorrect, rendering the aforementioned mechanism that ensures the data to be written by the data producing unit to the data consuming unit is actually written to the on-chip memory 370 of the data consuming unit meaningless. Therefore, in one embodiment, on the data consuming unit side, after the specific flag bit becomes a specific value, the second multi-level cache also executes a cache invalidation instruction and a wait-for-invalidation-success instruction. The cache invalidation instruction is used to prevent the data stored in the second multi-level cache from being read by subsequent read instructions, such as... Figure 6 The statement `CACHE.INV` in the code waits for an invalid success instruction to complete before allowing the data to be executed using instructions, as shown below. Figure 6The statement S_WAIT.CACHE in the code executes the cache invalidation instruction and the wait-for-invalidation-success instruction. Subsequent data usage instructions will no longer read the data cached in the second multi-level cache, but will instead use the actual data successfully written by the data producer unit to the on-chip memory 370 of the data consumer unit, thus avoiding data read errors.

[0116] After the above process is completed, the second thread group execution unit 342 causes the thread group to exchange second information with other thread groups. The second information indicates that the specific flag bit has become a specific value, meaning that subsequent data usage instructions can use the data successfully written by the data production main unit. The purpose of exchanging the second information is that the process of checking the synchronization set instruction on a single route, determining that all write requests before the synchronization set instruction on the single route have been successfully written, and executing the set wait instruction is only performed on one thread group, or even just one thread in a thread group. Exchanging the second information can save the huge overhead caused by each thread group executing the above process separately. One thread group has already changed the specific flag bit to a specific value, and subsequent data usage instructions can use the data successfully written by the data production main unit. However, other thread groups are unaware of this situation, which may cause them to hesitate to execute subsequent data usage instructions. This embodiment of the disclosure transmits the second information between thread groups. Through the second information, each thread group knows this situation, thereby allowing subsequent data usage instructions to confidently use the data successfully written by the data production main unit. The statement for exchanging the second information between thread groups is as follows: Figure 6 The ___syncthreads() function in the code compiles into assembly language as FENCE.BLK.

[0117] Communication control method for data production main unit like Figure 7 As shown, according to one embodiment of this disclosure, a cross-subject unit write method is provided at the data production subject unit end, which is executed by the data production subject unit. It includes: Step 710: Execute cross-subject write instructions to specific data consuming subject units for the thread group; Step 720: Enable the thread group to exchange first information with other thread groups, the first information indicating whether the write requests corresponding to the thread group have all been sent from the data production main unit; Step 730: For a thread group: If all write requests corresponding to the thread groups used for the cross-subject unit write instruction have been issued from the data production subject unit, a synchronization set instruction is issued on each of the multiple routes between the data production subject unit and the data consumption subject unit, so that after the data consumption subject unit receives the synchronization set instruction on each route, it ensures that all previously received write requests on that route have been successfully written, and allows the use of the written data after ensuring successful writing on each route.

[0118] Since the implementation details of the above process have been described in detail in the aforementioned data production device embodiment, they will not be repeated here to save space.

[0119] Communication control method for data consumption subject unit like Figure 8 As shown, according to one embodiment of this disclosure, a cross-subject unit write method is provided on the data consuming subject unit side, which is executed by the data producing subject unit, and includes: Step 810: Execute for a thread group: After receiving a synchronization set instruction on a single route between the data production main unit and the data consumption main unit, determine whether the previously received write request on that route has been successfully written. After determining that the write request has been successfully written on all routes between the data production main unit and the data consumption main unit, set the specific flag bit to a specific value and execute the set wait instruction and the data use instruction. The set wait instruction is used to allow the data use command to be executed after the specific flag bit becomes a specific value. The synchronization set instruction is issued on multiple routes between the data production main unit and the data consumption main unit after all write requests corresponding to all thread groups from the data production main unit to the data consumption main unit have been issued from the data production main unit. Step 820: The thread group exchanges second information with other thread groups, the second information indicating that the specific flag bit becomes a specific value.

[0120] Since the implementation details of the above process have been described in detail in the aforementioned data production main unit device embodiment, they will not be repeated here for the sake of saving space.

[0121] The disclosed commercial value This disclosure ensures that even with instruction-level writing, multiple routes between data production and consumption units, and data flowing through different caches and paths within the data production unit, the data consumption unit can still use the data written by the data production unit without errors, thus ensuring orderly data writing and usage. It can be applied to multi-route instruction-level writing between various entities such as acceleration units, processing units, on-chip systems, computing devices, and communication terminals. Experiments have shown that when applied to multi-route instruction-level writing between acceleration units, it can significantly improve the execution accuracy of instructions for using written data, demonstrating broad market prospects.

[0122] Those skilled in the art will understand that this disclosure can be implemented as a system, method, and computer program product. Therefore, this disclosure can be implemented as entirely hardware, entirely software (including firmware, resident software, and microcode), or a combination of software and hardware. Furthermore, in some embodiments, this disclosure can also be implemented as a computer program product contained in one or more computer-readable media, the computer-readable media containing computer-readable program code.

[0123] Any combination of one or more computer-readable media may be used. A computer-readable medium can be a computer-readable signal medium or a computer-readable storage medium. A computer-readable storage medium is, for example, but not limited to, an electrical, magnetic, optical, electromagnetic, infrared, or semiconductor system, apparatus, or device, or any combination thereof. More specific examples of a computer-readable storage medium include: an electrical connection of one or more wires, a portable computer disk, hard disk, random access memory (RAM), read-only memory (ROM), erasable programmable read-only memory (EPROM or flash memory), optical fiber, portable compact disk read-only memory (CD-ROM), optical storage, magnetic storage, or any suitable combination thereof. In this document, a computer-readable storage medium can be any tangible medium containing or storing a program that can be used by or in conjunction with a processing unit, apparatus, or device.

[0124] Computer-readable signal media may include data signals propagated in baseband or as part of a chopped signal, carrying computer-readable program code. Such propagated data signals may take various forms, including but not limited to electromagnetic signals, optical signals, or any other suitable combination. Computer-readable signal media may also be any computer-readable medium other than computer-readable storage media, capable of transmitting, propagating, or transmitting programs for use by or in connection with an instruction system, apparatus, or device.

[0125] The program code contained on a computer-readable medium may be transmitted using any suitable medium, including but not limited to wireless, wire, optical fiber, RF, and any suitable combination thereof.

[0126] Computer program code for executing embodiments of this disclosure can be written in one or more programming languages ​​or combinations thereof. The programming languages ​​include object-oriented programming languages ​​such as JAVA and C++, and may also include conventional procedural programming languages ​​such as C. The program code can be executed entirely on the user's computer, partially on the user's computer, as a standalone software package, partially on the user's computer and partially on a remote computer, or entirely on a remote computer or server. In cases involving remote computers, the remote computer can be connected to the user's computer via any type of network, including a local area network (LAn) or a wide area network (WAn), or it can be connected to an external computer (e.g., via the Internet using an Internet service provider).

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

Claims

1. A data production entity, comprising: The first thread group execution unit is used to execute cross-subject unit write instructions to a specific data consuming subject unit for a target thread group, and to make the target thread group exchange first information with other thread groups. The first information indicates whether the write requests corresponding to the target thread group have all been issued from the data producing subject unit. The first communication control unit is configured to execute the following for a thread group: if all write requests corresponding to the thread group for the cross-subject unit write instruction have been issued from the data production subject unit, a synchronization set instruction is issued on each of the multiple routes between the data production subject unit and the data consumption subject unit, so that after the data consumption subject unit receives the synchronization set instruction on each route, it ensures that all previously received write requests on that route have been successfully written, and allows the use of the written data after ensuring successful writing on each route.

2. The data production main unit according to claim 1 further includes: The first multi-level cache, located between the first thread group execution unit and the first communication control unit, executes a flush instruction and a wait-for-flush-success instruction before issuing synchronization set instructions on the multiple routes. The flush instruction is used to flush out the data stored in the cache and send it from the first communication control unit to the data consumption main unit through the higher-level cache. The wait-for-flush-success instruction is used to wait for all the data flushed out of the cache to be successfully sent out from the data production main unit.

3. The data production main unit according to claim 1 or 2, wherein, After the first communication control unit issues a synchronization set command on each of the multiple routes, it executes a synchronization set wait command. The synchronization set wait command is used to indicate that no subsequent command can be executed before the synchronization set command is issued on each of the multiple routes.

4. The data production main unit according to claim 1 or 2, wherein, After the first communication control unit issues synchronization set instructions on the multiple routes respectively, the first thread group execution unit also causes the target thread group to exchange third information with other thread groups, the third information indicating that synchronization set instructions have been issued on the multiple routes respectively.

5. The data production main unit according to claim 1, wherein, The first communication control unit includes: Multiple direct storage access modules, corresponding to multiple sending ports, each store the write requests to be sent from the corresponding sending port; The instruction allocator determines the route between the data production unit and the data consumption unit based on the routing table, allocates the write request corresponding to a thread group to the direct storage access module of the sending port of the corresponding route, and, after determining that all write requests corresponding to all thread groups have been sent from the direct storage access module of the sending port of the corresponding route, reads the synchronization set instruction and copies the synchronization set instruction to the direct storage access module of the sending port of the corresponding route for issuance.

6. The data production main unit according to claim 1 or 2, wherein, The data production entity is a first processing unit, and the data consumption entity is a second processing unit. The first processing unit and the second processing unit communicate through an inter-chip interconnection network. The cross-entity unit write instruction writes the data in the first on-chip memory of the first processing unit to the second on-chip memory of the second processing unit through the inter-chip interconnection network.

7. The data production main unit according to claim 6, wherein, The first processing unit is a first acceleration unit, used to execute tasks assigned by the first scheduling unit; the second processing unit is a second acceleration unit, used to execute tasks assigned by the second scheduling unit.

8. A data consumption subject unit, comprising: The second communication control unit is used to perform the following for the target thread group: after receiving a synchronization set instruction on a single route between the data production main unit and the data consumption main unit, it determines whether the previously received write request on that route has been successfully written, and after determining that the write request has been successfully written on all routes between the data production main unit and the data consumption main unit, it sets a specific flag bit to a specific value and executes a set-wait instruction and a data usage instruction. The set-wait instruction is used to allow the data usage instruction to be executed after the specific flag bit becomes a specific value. The synchronization set instruction is issued on multiple routes between the data production main unit and the data consumption main unit after all write requests corresponding to all thread groups from the data production main unit to the data consumption main unit have been issued from the data production main unit. The second thread group execution unit is used to enable the target thread group to exchange second information with other thread groups, wherein the second information indicates that the specific flag bit becomes a specific value.

9. The data consumption subject unit according to claim 8, further comprising: The second multi-level cache is located between the second thread group execution unit and the second communication control unit. After the specific flag bit becomes a specific value, the second multi-level cache executes a cache invalidation instruction and a wait-for-invalidation-success instruction. The cache invalidation instruction is used to prevent the data stored in the second multi-level cache from being read by subsequent read instructions. The wait-for-invalidation-success instruction is used to wait for the invalidation instruction to be executed, and only after the invalidation instruction has been executed will the data be allowed to be used for instruction execution.

10. The data consumption subject unit according to claim 8 or 9, wherein, The second communication control unit includes: Multiple direct storage access modules, corresponding to multiple receiving ports, are used to cache write requests received from each receiving port. For the target thread group, after receiving a synchronization set instruction on the corresponding receiving port, the module determines whether the write request previously received from the data production main unit on that receiving port has been successfully written. The receiving controller is configured to, for the target thread group, after each direct storage access module has determined that all write requests previously received from the data production host unit on the corresponding receiving port have been successfully written, read the set-wait instruction and the data use instruction, and execute the set-wait instruction to allow the data use instruction to be executed after the specific flag bit becomes a specific value.

11. A system-on-a-chip, comprising a data production unit according to any one of claims 1-7 and / or a data consumption unit according to any one of claims 8-10.

12. A computing device comprising a data production unit according to any one of claims 1-7 and / or a data consumption unit according to any one of claims 8-10.

13. A cross-subject unit write method at the data production subject unit end, comprising: Execute cross-subject write instructions to a specific data consuming subject unit for the target thread group; The target thread group exchanges first information with other thread groups, the first information indicating whether all write requests corresponding to the target thread group have been sent from the data production main unit. For a single thread group: If all write requests corresponding to the thread groups used for the cross-subject unit write instruction have been issued from the data production subject unit, a synchronization set instruction is issued on each of the multiple routes between the data production subject unit and the data consumption subject unit. This ensures that after the data consumption subject unit receives the synchronization set instruction on each route, all previously received write requests on that route are successfully written. Only after ensuring successful writing on each route can the written data be used.

14. A cross-subject unit write method at the data consuming subject unit end, comprising: For the target thread group: After receiving a synchronization set instruction on a single route between the data production main unit and the data consumption main unit, it is determined whether the previously received write request on that route has been successfully written. After determining that the write request has been successfully written on all routes between the data production main unit and the data consumption main unit, a specific flag bit is set to a specific value, and a set-wait instruction and a data usage instruction are executed. The set-wait instruction is used to allow the data usage instruction to be executed after the specific flag bit becomes a specific value. The synchronization set instruction is issued on multiple routes between the data production main unit and the data consumption main unit after all write requests corresponding to all thread groups from the data production main unit to the data consumption main unit have been issued from the data production main unit. The target thread group exchanges second information with other thread groups, the second information indicating that the specific flag bit becomes a specific value.

Images