System and method for network multicasting using a set of alternate indications

Abstract

Embodiments of the present disclosure relate to network multicasting using a set of fallback indications. Apparatuses, systems, and techniques for multicasting a transaction to a target group. In at least one embodiment, a set is selected from a set of fallback indications associated with the target group, and the transaction is transmitted to the target group in accordance with the selected set.

Description

TECHNICAL FIELD

[0001] At least one embodiment is directed to multicasting communications from a source device to a plurality of target devices. For example, at least one embodiment is directed to a switch circuit that implements such multicasting. BACKGROUND

[0002] Multicasting is a process in which a communication is sent by a source device to a routing device (sometimes referred to as a "switch"), which in turn sends the communication over a network to a plurality of targets. Because the switch sends multiple copies of the information to the targets, the source device only needs to send the communication once to reach multiple targets. In addition, some routing devices perform traffic shaping and help distribute load so that network traffic will exhibit desired static properties. Thus, multicasting is a useful tool for network communications. BRIEF DESCRIPTION OF DRAWINGS

[0003] Figure 1 A block diagram of a system including a switch circuit positioned between a set of source devices and a group of targets is shown in accordance with at least one embodiment;

[0004] Figure 2 A block diagram of an internal switch of the switch circuit of Figure 1 is shown in accordance with at least one embodiment;

[0005] Figure 3 A block diagram of an example implementation of the switch circuit of Figure 1 including interconnections between each odd column and one of its adjacent even columns is shown in accordance with at least one embodiment;

[0006] Figure 4A A block diagram of an example implementation of a routing data structure is shown in accordance with at least one embodiment;

[0007] Figure 4B An example format that can be used to encode each tree in the indication data portion of the routing data structure of Figure 4A is shown in accordance with at least one embodiment;

[0008] Figure 5 A block diagram of the routing data structure of Figure 4A linked to a shared routing data structure is shown in accordance with at least one embodiment;

[0009] Figure 6 A flow diagram of a method that can be performed by a switch circuit when the switch circuit receives a transaction is shown in accordance with at least one embodiment;

[0010] Figure 7 A flow diagram of a method that can be performed by a switch circuit when the switch circuit is performing a reduction operation is shown in accordance with at least one embodiment;

[0011] Figure 8 An exemplary data center is shown in accordance with at least one embodiment;

[0012] Figure 9 A processing system is shown in accordance with at least one embodiment;

[0013] Figure 10 A computer system is shown in accordance with at least one embodiment;

[0014] Figure 11 A system is shown in accordance with at least one embodiment;

[0015] Figure 12 An exemplary integrated circuit is shown in accordance with at least one embodiment;

[0016] Figure 13 A computing system is shown in accordance with at least one embodiment;

[0017] Figure 14 An APU is shown in accordance with at least one embodiment;

[0018] Figure 15 A CPU is shown in accordance with at least one embodiment;

[0019] Figure 16 An exemplary accelerator integration slice is shown in accordance with at least one embodiment;

[0020] Figures 17A-17B An exemplary graphics processor is shown in accordance with at least one embodiment;

[0021] Figure 18A A graphics core is shown in accordance with at least one embodiment;

[0022] Figure 18B A GPGPU is shown in accordance with at least one embodiment;

[0023] Figure 19A A parallel processor is shown in accordance with at least one embodiment;

[0024] Figure 19B A processing cluster is shown in accordance with at least one embodiment;

[0025] Figure 19C A graphics multiprocessor is shown in accordance with at least one embodiment;

[0026] Figure 20 A graphics processor is shown in accordance with at least one embodiment;

[0027] Figure 21 A processor is shown in accordance with at least one embodiment;

[0028] Figure 22 A processor is shown in accordance with at least one embodiment;

[0029] Figure 23 A graphics processor core is shown in accordance with at least one embodiment;

[0030] Figure 24 A PPU is shown in accordance with at least one embodiment;

[0031] Figure 25 A GPC is shown in accordance with at least one embodiment;

[0032] Figure 26 A streaming multiprocessor is shown in accordance with at least one embodiment;

[0033] Figure 27 A software stack of a programming platform is shown in accordance with at least one embodiment;

[0034] Figure 28 A CUDA implementation of the software stack of Figure 27 is shown in accordance with at least one embodiment;

[0035] Figure 29 A ROCm implementation of the software stack of Figure 27 is shown in accordance with at least one embodiment;

[0036] Figure 30 An OpenCL implementation of the software stack of Figure 27 is shown in accordance with at least one embodiment;

[0037] Figure 31 Software supported by a programming platform is shown in accordance with at least one embodiment;

[0038] Figure 32 Compiled code executed on the programming platform of Figures 27-30 is shown in accordance with at least one embodiment;

[0039] Figure 33 More detailed compiled code executed on the programming platform of Figures 27-30 is shown in accordance with at least one embodiment;

[0040] Figure 34 Conversion of source code prior to compiling the source code is shown in accordance with at least one embodiment;

[0041] Figure 35A A system configured to compile and execute CUDA source code using different types of processing units is shown in accordance with at least one embodiment;

[0042] Figure 35BA system configured to compile and execute CUDA source code using a CPU and a CUDA-enabled GPU is shown in accordance with at least one embodiment; Figure 35A

[0043] Figure 35C A system configured to compile and execute CUDA source code using a CPU and a non-CUDA-enabled GPU is shown in accordance with at least one embodiment; Figure 35A

[0044] Figure 36 An exemplary kernel converted by a CUDA to HIP conversion tool in accordance with at least one embodiment is shown; Figure 35C

[0045] Figure 37 A non-CUDA-enabled GPU in accordance with at least one embodiment is shown in more detail; Figure 35C

[0046] Figure 38 How threads of an exemplary CUDA grid are mapped to different compute units in accordance with at least one embodiment is shown; and Figure 37

[0047] Figure 39 How to migrate existing CUDA code to data-parallel C++ code in accordance with at least one embodiment is shown.DETAILED DESCRIPTION

[0048] In the following description, numerous specific details are set forth to provide a more thorough understanding of the embodiments. However, it will be apparent to one of skill in the art that the inventive concept can be practiced without one or more of these specific details.

[0049] Figure 1 A block diagram of a system 100 including a switch fabric 110 (sometimes referred to as a "fabric") positioned between a set of source devices 112 and a set of targets 114 is shown in accordance with at least one embodiment. Each of the set of source devices 112 and the set of targets 114 can be implemented as a GPU, a CPU, a controller, a switch, a switch fabric, a memory device, or the like. As a non-limiting example, the set 112 has been shown to include a number of source devices equal to "M*X", which are shown as source devices SD1-1 through SDM-X. However, the set 112 can include any number of source devices, including a single source device. In Figure 1 ​​​​​In the diagram, source devices SD1-1 to SDM-X are shown arranged in a two-dimensional array, comprising M rows and X columns. For example, source devices SD1-1 to SDM-X can be arranged or described as comprising rows S1-SM, each row comprising X source devices. Figure 1 In the example shown, row S1 includes source devices SD1-1 to SD1-X, row S2 includes source devices SD2-1 to SD2-X, and row SM includes source devices SDM-1 to SDM-X.

[0050] For ease of illustration, group 114 has been shown as including a quantity of "Y" targets, shown as targets T1-TY, but group 114 may include any number of targets. Furthermore, one or more devices may function as both source devices and targets. Therefore, set 112 and group 114 are not mutually exclusive, but may include one or more of the same devices. Thus, in at least one embodiment, switching circuitry 110 may allow any of source devices SD1-1 to SDM-X to communicate with any of source devices SD1-1 to SDM-X and / or any of targets T1-TY. Furthermore, in at least one embodiment, switching circuitry 110 may allow any of targets T1-TY to communicate with targets T1-TY and / or any of source devices SD1-1 to SDM-X.

[0051] Switching circuit 110 can implement multicast within network 116, which includes set 112 and group 114. Therefore, switching circuit 110 can interconnect GPUs, CPUs, and / or other types of processing units. As a non-limiting example, system 100 can be implemented as a GPU-to-GPU link system (e.g., A GPU-to-GPU interconnect system, wherein the switching circuit 110 can be implemented as a GPU-to-GPU switch (e.g., NVSWITCH). TM (Switch). Switching circuit 110 can be used to interconnect two or more parallel processing units (such as GPUs). Such parallel processing units can implement other systems (e.g., be components of other systems), such as autonomous vehicles, medical imaging equipment, and the like. Switching circuit 110 can be integrated into graphics cards and / or similar graphics processing units.

[0052] The switch circuit 110 includes a number "N" of outbound switches OSW1-OSWN, a predetermined number "M*N" of internal switches, and a routing block 118 that controls operation of the internal switches. The routing block 118 has been shown to include a clock 120, at least one processor 122, and a memory 124. In at least one embodiment, each processor 122 can be implemented as one or more hardware state machines, one or more microprocessors, one or more microcontrollers, one or more controllers, or the like. As a non-limiting example, the switch circuit 110 has been shown to include internal switches IS1C1-ISMCN; however, the switch circuit 110 can include any number of internal switches, including a single internal switch. Further, the internal switches IS1C1-ISMCN have been shown to be arranged in a two-dimensional array that includes a number "M" of rows and a number "N" of columns. For example, the internal switches IS1C1-ISMCN can be arranged in rows R1-RM and columns C1-CN. However, this is not required and the internal switches IS1C1-ISMCN can be positioned in an alternative arrangement to implement any network topology, including a three-dimensional array. As a non-limiting example, the internal switches IS1C1-ISMCN can be arranged to define a Clos structure or the like. For example, each of the rows R1-RM can be implemented as a different Clos. One or more of the internal switches IS1C1-ISMCN can be coupled to one another. For example, from the perspective of one of the internal switches IS1C1-ISMCN, one or more of the other internal switches IS1C1-ISMCN can be one of the targets T1-TY. Further, the switch circuit 110 can be coupled to one or more switches (e.g., a switch circuit like the switch circuit 110) that are external to the switch circuit 110. For example, one or more of the source devices SD1-1 to SDM-X and / or the targets T1-TY can each be implemented as an external switch.

[0053] The routing block 118 can be connected to each of the internal switches IS1C1-ISMCN by one or more buses. For example, the clock 120, the processor 122, the memory 124, and the internal switches IS1C1-ISMCN can be connected to one another by one or more buses 126. The processor 122 can be implemented as one or more microprocessors, one or more microcontrollers, one or more controllers, or the like. The memory 124 stores instructions 130 and data 132 that are executable by the processor 122. The data 132 can include routing information stored in a routing data structure (e.g., a routing table).

[0054] The internal switch IS1C1-ISMCN includes, or is connected to, a first set of internal ports I-T1C1 to I-TMCN, and a second set of internal ports O-T1C1 to O-TMCN, through which communication can be received and / or transmitted. The first set of internal ports I-T1C1 to I-TMCN and the second set of internal ports O-T1C1 to O-TMCN may each include one or more ports. Although shown and described as ports, the first set of internal ports I-T1C1 to I-TMCN and / or the second set of internal ports O-T1C1 to O-TMCN may be implemented as a bus or other type of communication connection (e.g., carrying signals on a silicon chip, printed circuit board "PCB", or the like).

[0055] As described above, in the illustrated embodiment, both the source devices SD1-1 to SDM-X and the internal switches IS1C1-ISMCN are arranged, or can be described as comprising a number of rows "M". Rows S1-SM can be described as corresponding to rows R1-RM respectively. Within each row of S1-SM, the source device in the row is connected to each of the internal switches in the corresponding row of R1-RM. For example, as Figure 1 As shown, source devices SD1-1 to SD1-X in row S1 are respectively connected to internal switches IS1C1 to IS1CN in row R1. The first sets I-T1C1 to I-TMCN may each include different ports, which are connected to each of at least a portion of the source devices SD1-1 to SDM-X via one or more buses 134. For example, the first sets I-T1C1 to I-T1CN of internal switches IS1C1 to IS1CN in row R1 may each include different internal ports, which are connected to each of the source devices SD1-1 to SD1-X in row S1 via bus 134. Similarly, the first sets I-T2C1 to I-T2CN of internal switches IS2C1 to IS2CN in row R2 may each include different internal ports, which are connected to each of the source devices SD2-1 to SD2-X in row S2 via bus 134. Furthermore, the first set I-TMC1 to I-TMCN of the internal switches ISMC1-ISMCN in the row RM can each include different internal ports, which are respectively connected to each of the source devices SDM-1 to SDM-X in the row SM via bus 134.

[0056] The second sets O-T1C1 through O-TMCN can each include different ports that are connected to each of at least a portion of the targets T1-TY by a bus 140. In the example shown, the bus 140 includes buses 140-1 through 140-N. For example, in column C1, the second set O-T1C1 through O-TMC1 is connected to each of at least a portion of the targets T1-TY by the bus 140-1. Similarly, in columns C2-CN, the second sets O-T1C2 through O-TMCN are connected to each of at least a portion of the targets T1-TY by the buses 140-2 through 140-N, respectively.

[0057] As noted above, the internal switches IS1C1-ISMCN can be arranged in rows R1-RM and columns C1-CN. Along each row in each of the columns C1-CN, a bus 142 connects the internal switches in that column to one of the outbound switches OSW1-OSWN. For example, in Figure 1 In the example shown, the bus 142-1 connects the internal switches IS1C1-ISMC1 to the outbound switch OSW1, the bus 142-2 connects the internal switches IS1C2-ISMC2 to the outbound switch OSW2, and the bus 142-N connects the internal switches IS1CN-ISMCN to the outbound switch OSWN. The bus 140-1 connects the outbound switch OSW1 to a first portion of the targets T1-TY, the bus 140-2 connects the outbound switch OSW2 to a second portion of the targets T1-TY, and the bus 140-N connects the outbound switch OSWN to a third portion of the targets T1-TY.

[0058] Along each column in C1-CN, up to a number "M" of the internal switches can attempt to transmit a communication to the same target. The outbound switches OSW1-OSWN queue any transmissions directed to the same target so that only one of the communications is sent to the same target at a time. For example, if a source device SD1-1 sends a first communication to the internal switches IS1C1-IS1CN addressed to a target T3 and a source device SD2-X sends a second communication to the internal switches IS2C1-IS2CN addressed to the target T3, the first and second communications can arrive at the outbound switch OSW1 at about the same time. The outbound switch OSW1 will queue one of the first and second communications until the other of the first and second communications is sent. The outbound switch OSW1 will then send the other communication to the target T3. Thus, the outbound switches OSW1-OSWN help control outbound transmissions.

[0059] In Figure 1In the illustrated example, source devices SD1-2 are multicasting communications 150 to targets T1-T3. In this example, internal switches IS1C1, IS1C2, and IS1CN of first row R1 route communications 150 to targets T3, T2, and T1, respectively. Thus, one port in each of first sets I-T1C1, I-T1C2, and I-T1CN receives communications 150. Routing block 118 then determines through which port or ports of second sets O-T1C1, O-T1C2, and O-T1CN to send communications 150 to reach targets T3, T2, and T1, respectively. For example, instructions 130 can cause processor 122 to determine through which port or ports of second sets O-T1C1, O-T1C2, and O-T1CN to send communications 150 to reach targets T3, T2, and T1, respectively. In Figure 1 In the illustrated example, one of the ports of second set O-T1C1 sends communications 150 to outbound switch OSW1, which forwards communications 150 to target T3 via at least one of buses 140-1. If outbound switch OSW1 receives other communications addressed to target T1 prior to communications 150, outbound switch OSW1 will store communications 150 until those communications have been sent. One of the ports of second set O-T1C2 sends communications 150 to outbound switch OSW2, which forwards communications 150 to target T2 via at least one of buses 140-2. If outbound switch OSW2 receives other communications addressed to target T2 prior to communications 150, outbound switch OSW2 will store communications 150 until those communications have been sent. One of the ports of second set O-T1CN sends communications 150 to outbound switch OSWN, which forwards communications 150 to target T1 via at least one of buses 140-N. If outbound switch OSWN receives other communications addressed to target T1 prior to communications 150, outbound switch OSWN will store communications 150 until those communications have been sent. Thus, outbound switches OSW1-OSWN help manage congestion along columns C1-CN, respectively, and allow internal switches along columns C1-CN to share buses 140-1 through 140-N, respectively. Thus, in at least one embodiment, outbound switches OSW1-OSWN can operate as or be connected to outbound ports of switching circuit 110.

[0060] Figure 2 A block diagram of internal switch IS1C1 is shown, in accordance with at least one embodiment. Each of other internal switches IS2C1-ISMCN (see Figure 1 and Figure 3 ) can be substantially identical to internal switch IS1C1. Turning to Figure 2The internal switch IS1C1 includes internal switch circuitry 200 that is connected to a first set I-T1C1 of ports and a second set O-T1C1 of ports. In this example, the first set I-T1C1 includes a number "P" of ports 201-1 through 201-P that are connected to a portion of the source devices SD1-1 through SD1-X, respectively, through the bus 134. For example, the number "P" can equal the number "X", however, this is not required. In at least one embodiment, the number "P" can be greater than or less than the number "X". Referring to Figure 1 Similarly, each of the first sets I-T2C1 through I-TMCN includes a number (e.g., a number "P") of ports. Thus, each of the first sets I-T1C1 through I-TMCN can be connected to all or a subset of the source devices SD1-1 through SDM-X. In the embodiment shown in Figure 2 In the embodiment shown in

[0061] In the embodiment shown in Figure 1 and Figure 2 The source device SD1-2 is multicasting a communication 150 to the targets T1-T3. Referring to Figure 1 The source device SD1-2 transmits the communication 150 through the bus 134 to the switch circuitry 110, which routes the communication 150 to each of the internal switches IS1C1, IS1C2, and IS1CN (see Figure 1 and Figure 3 ) connected to the source device SD1-2 by the bus 134. Referring to Figure 2 The bus 134 delivers the communication 150 to the port 201-2 and to the ports in the first sets I-T1C2 and I-T1CN (see Figure 1 and Figure 3 ) that are also connected to the source device SD1-2.

[0062] The second set O-T1C1 can include different ports that are connected to each of at least a portion of the targets T1-TY through the bus 140-1. For example, the second set O-T1C1 can include different ports that are connected to each of the targets T1-TY. In the example shown in Figure 2 The second set O-T1C1 includes a number "O" of ports 211-1 through 211-O that are connected to the targets T1-TY, respectively. Thus, in this example, the number "O" equals the number "Y", however, this is not required. In at least one embodiment, the number "O" can be greater than or less than the number "Y". Referring to Figure 1Similarly, the second set O-T2C1 through O-TMCN each includes a number (e.g., quantity "O") of ports. Thus, each of the second set O-T1C1 through O-TMCN can be connected to all or a subset of the targets T1-TY.

[0063] The value of the quantity "N" need not equal the value of the quantity "M," the value of the quantity "X," the value of the quantity "Y," the value of the quantity "P," or the value of the quantity "O." As a non-limiting example, the value of the quantity "M" can be 6, the value of the quantity "N" can be 6, the value of the quantity "X" can be 8, the value of the quantity "Y" can be 8, the value of the quantity "P" can be 11, and the value of the quantity "O" can be 11.

[0064] As described above, in Figure 2 , the source device SD1-2 is multicasting a communication 150 to the targets T1-T3. In the illustrated example, the source device SD1-2 transmits the communication 150 to the internal switches IS1C1, IS1C2, and IS1CN (see Figure 1 and Figure 3 ). Upon receiving the communication 150, one or more of the internal switches IS1C1, IS1C2, and IS1CN can notify the processor 122 that the communication 150 has been received. Alternatively, the instructions 130 can cause the processor 122 to know that the current transaction has been received. For example, the instructions 130 can cause the processor 122 to poll the internal switches IS1C1-ISMCN for transactions.

[0065] After the processor 122 knows that the communication 150 has been received, the instructions 130 can cause the processor 122 to instruct the internal switch circuit 200 to route the communication 150 through which of the ports of the second set O-T1C1. In the illustrated example, the instructions 130 cause the processor 122 to select the port 211-3 and instruct the internal switch circuit 200 to route the communication 150 through the port 211-3 to the target T3. Similarly, with reference to Figure 1 , when the internal switches IS1C2 and IS1CN receive the communication 150, the instructions 130 cause the processor 122 to select a second port of the second set O-T1C2 that is connected to the target T2 and a third port of the second set O-T1CN that is connected to the target T1. The instructions 130 then cause the processor 122 to instruct the internal switch circuit of the internal switch IS1C2 (e.g., the internal switch circuit 200) to route the communication 150 through the second port of the second set O-T1C2 to the target T2 and to instruct the internal switch circuit of the internal switch IS1CN (e.g., the internal switch circuit 200) to route the communication 150 through the third port of the second set O-T1CN to the target T1. Figure 2The internal switch circuit (e.g., internal switch circuit 200) of the internal switch IS1CN with the second port sends the communication 150, and instructs the internal switch circuit (e.g., internal switch circuit 200) of the internal switch IS1CN with the third port to send the communication 150 through the third port. In other words, the instructions 130 can cause the processor 122 to select one of the O-T1C2 to route the communication 150 to the target T2, and select one of the O-T1CN to route the communication 150 to the target T1. In the illustrated example, if the communication 150 is received by the internal switches IS1C1- ISMCN that do not transmit the communication 150 to either of the targets T1-TY, these internal switches can discard the communication 150. For example, the instructions 130 can cause the processor 122 to instruct any of the internal switches IS1C1- ISMCN other than the internal switches IS1C1, IS1C2, and IS1CN to drop (or discard) the communication 150.

[0066] Referring to Figure 2 In at least one embodiment, the communication 150 can be divided into a plurality of transactions 152 (such as data packets) according to a communication protocol. The transactions 152 can be sent by the source device (e.g., source device SD1-2) to the switch circuit 110 one at a time, serially. While Figure 2 While the transactions 152 are shown to include transaction 152A and transaction 152B, the transactions 152 can include any number of transactions. In some networks, the order of the transactions 152 (e.g., data packets) must be maintained to ensure that the transactions 152 arrive at the target in the same order that they were received by the switch circuit 110. To help ensure that this occurs, the switch circuit 110 can send the transactions 152 along the same path through the network 116. Thus, to maintain the order of the transactions 152, the switch circuit 110 can send each of the transactions 152 via the same port (e.g., port 211-3) of the same internal switch (e.g., internal switch IS1C1).

[0067] Referring to Figure 3 As described above, in at least one embodiment, the internal switches IS1C1- ISMCN can be arranged in columns C1-CN and rows R1-RM. The columns C1-CN, when numbered, alternate between odd columns (e.g., column C1) and even columns (e.g., column C2). Figure 3 A block diagram of an example implementation of the switch circuit 110 including interconnections between each odd column (e.g., column C1) and one of its adjacent even columns (e.g., column C2) is shown in accordance with at least one embodiment. As shown, the switch circuit 110 includes a plurality of internal switches IS1C1- ISMCN arranged in columns C1-CN and rows R1-RM. The internal switches IS1C1- ISMCN are interconnected by a plurality of ports 211-3. In the illustrated example, the switch circuit 110 includes a plurality of internal switches IS1C1- ISMCN arranged in columns C1-CN and rows R1-RM. The internal switches IS1C1- ISMCN are interconnected by a plurality of ports 211-3. Figure 3As shown, each of columns C1-CN can be interconnected with one or more of its adjacent columns. For example, interconnect conductor 302 can connect internal switches IS1C1- ISMC1 to outbound switch OSW2, and interconnect conductor 304 can connect internal switches IS1C2- ISMC2 to outbound switch OSW1. To facilitate illustration, Figure 3 Buses 142-1 and 142-2 are shown using dashed lines in FIG. 1. Interconnect conductor 302 transmits signals from internal switches IS1C1- ISMC1 to outbound switch OSW2, and interconnect conductor 304 transmits signals from internal switches IS1C2- ISMC2 to outbound switch OSW1. Thus, transactions (e.g., transaction 152A and transaction 152B shown in FIG. 1) that are routed along internal switches IS1C1- ISMC1 (in odd columns C1) can be transmitted to those of targets T1-TY connected to outbound switch OSW1 (in odd columns C1) via bus 142-1, or to those of targets T1-TY connected to outbound switch OSW2 in even columns C2 via interconnect conductor 302. Further, transactions that are routed along internal switches IS1C2- ISMC2 (in even columns C2) can be transmitted to those of targets T1-TY connected to outbound switch OSW2 (in even columns C2) via bus 142-2, or to those of targets T1-TY connected to outbound switch OSW1 in odd columns C1 via interconnect conductor 304. Figure 2

[0068] Even if internal switches IS1C1- ISMC1 and internal switches IS1C2- ISMC2 are connected to the same targets T1-TY, the interconnect allows multiple transactions to be sent to the same or different targets at the same time. In this way, the interconnect can reduce the amount of time required to send transactions to targets T1-TY. Each of the odd columns can be similarly interconnected with one of its adjacent even columns. As a non-limiting example, each odd column can be cross-coupled or interconnected with an even column that is assigned a column number that is one greater than the number of columns assigned to the odd column.

[0069] In Figure 3 ​In the illustrated embodiment, each of the internal switches IS1C1-ISMCN can output transactions along either odd or even columns because pairs of adjacent columns are interconnected. Therefore, routing block 118 can route communication 150 along a primary path on one or more of the buses 142 or an alternate path on one or more interconnect conductors (e.g., interconnect conductors 302 or 304). When routing block 118 routes communication 150 through an even-numbered internal switch (e.g., outbound switch IS1C1), the primary path extends through the outbound switch of that even-numbered column (e.g., outbound switch OSW1), and the alternate path extends through an outbound switch of an even-numbered column interconnected with an odd-numbered column (e.g., outbound switch OSW2). Similarly, when routing block 118 routes communication 150 through an internal switch in an even-numbered column (e.g., outbound switch IS1C2), the primary path extends through the outbound switch in that even-numbered column (e.g., outbound switch OSW2), and the backup path extends through the outbound switches in odd-numbered columns interconnected with the even-numbered columns (e.g., outbound switch OSW1). For example, if routing block 118 routes communication 150 through at least one of the internal switches IS1C1-ISMC1 in the odd-numbered column C1, then communication 150 has a primary path through outbound switch OSW1 and a backup path through outbound switch OSW2. Similarly, if routing block 118 routes communication 150 through at least one of the internal switches IS1C2-ISMC2 in the even-numbered column C2, then communication 150 has a primary path through outbound switch OSW2 and a backup path through outbound switch OSW1.

[0070] For example, in Figure 3 In the illustrated embodiment, source device SD1-2 can send communication 150 to target T3 via two different paths. First, source device SD1-2 can route communication 150 to outbound switch OSW1 via one of the internal ports I-T1C1 of internal switch IS1C1 and one or more buses 142-1. Second, source device SD1-2 can route communication 150 to outbound switch OSW1 via one of the internal ports I-T1C2 of internal switch IS1C2 and one or more interconnecting conductors 304.

[0071] Reference Figure 3Such interconnections can reduce the amount of time required to send transactions to targets T1-TY by reducing the number of rounds required. As noted above, if only one of the internal switches IS1C1-ISMCN is used to send a transaction, the transaction will have to be transmitted by the internal switch serially once for each target. Thus, if the internal switch IS1C1 is used to send a transaction to targets T3-T5, the internal switch IS1C1 will have to send the transaction in three separate and serial rounds. In other words, the tree will need to include three round directives for the internal switch IS1C1 to transmit the transaction to targets T3-T5. Each of these serial transmissions will occur after an increment of the clock 120. Thus, if the internal switch IS1C1 transmits the transaction to targets T3-T5, it will require at least two increments of the clock 120 (e.g., a first increment between transmission to targets T3 and T4, and a second increment between transmission to targets T4 and T5). This time can be reduced by implementing interconnections that provide multiple paths to a particular target. For example, referring to Figure 3 If the transmission is sent to both internal switches IS1C1 and IS1C2 simultaneously, the internal switch IS1C1 can transmit the transaction to target T3 via the outbound switch OSW1 (and bus 142-1), and the internal switch IS1C2 can transmit the transaction to target T4 via the outbound switch OSW1 (and interconnection conductor 304). Then, after an increment of the clock 120, the internal switch IS1C1 can transmit the transaction to target T5 via the outbound switch OSW1 (and bus 142-1). In this example, the use of the interconnection conductor 304 reduces the number of rounds by one round and reduces the number of clock increments by one increment.

[0072] In other words, the number of rounds represents the serialization cost (e.g., the time required for multicasting), so reducing the number of rounds also reduces this serialization cost. Generally, interconnecting a pair of even and odd columns can reduce the number of rounds required by about half. However, more of the columns can be interconnected to further reduce the number of rounds required for the multicast communication 150.

[0073] The bus 126, the bus 140, the bus 142, the interconnection conductor 302, and the interconnection conductor 304 can each be implemented as a signal-conducting medium, such as one or more wires, one or more signal traces, and the like. As a non-limiting example, the bus 140, the bus 142, the interconnection conductor 302, and / or the interconnection conductor 304 can each be implemented as one or more GPU-to-GPU links (e.g., PCIe links), one or more buses, one or more traces, and the like. This can be implemented using one or more links in a GPU-to-GPU interconnect system. As another non-limiting example, bus 140, bus 142, interconnect conductor 302, and / or interconnect conductor 304 can each be implemented using one or more pairs of differential signal conductors. For example, refer to... Figure 2 Bus 142-1 may include a first pair and a second pair of differential signal conductors for each of ports 201-1 to 201-P. The first pair conducts signals from outbound switch OSW1 to a specific port among ports 201-1 to 201-P, and the second pair conducts signals from that specific port to outbound switch OSW1. Similarly, bus 140-1 may include a third pair and a fourth pair of differential signal conductors for each of at least a portion of targets T1-TY. The third pair conducts signals from outbound switch OSW1 to a specific target in target T1-TY, and the fourth pair conducts signals from that specific target to outbound switch OSW1. In at least one embodiment, each of buses 140 is substantially similar to bus 140-1, and each of buses 142 is substantially similar to bus 142-1.

[0074] Reference Figure 1 Internal switches IS1C1-ISMCN can provide multiple paths through which communication 150 can reach at least some of the targets T1-TY. For example, two or more of the internal switches IS1C1-ISMCN can be connected directly or via other circuits (e.g., another internal switch IS1C1-ISMCN) to the same target T1-TY. However, if source device SD1-2 uses internal switch IS1C1 to multicast transaction 152 of communication 150 to more than one target T1-TY, then internal switch IS1C1 will have to serially send one of transaction 152 for each of the multiple targets each time. For example, if source device SD1-2 instructs to send transactions to three targets T1-T3 (e.g., Figure 2 For transaction 152A shown, internal switch IS1C1 will have to send the transaction in three rounds (referred to as "rounds"). Similarly, if any resource of switching circuit 110 (e.g., internal switches IS1C1-ISMCN) needs to send transactions serially multiple times, these resources can send transactions in rounds. Therefore, if one of the internal switches IS1C1-ISMCN needs to send a transaction to the same destination, that internal switch can send the transaction in multiple rounds.

[0075] Reference Figure 2To send the communication 150 to multiple targets among the targets T1-TY, the communication 150 includes an identifier MCID (e.g., a multicast identifier) that the switch circuit 110 uses to identify those of the targets T1-TY (e.g., the targets T1-T3) that are to receive the communication 150. As shown in Figure 2 When the communication 150 is divided into multiple transactions 152, each of the transactions 152 can include the same identifier MCID. Referring to Figure 4A As noted above, the data 132 can include routing information stored in a routing data structure 410 (e.g., a routing table) that maps the identifier MCID (see Figure 2 ) to an instruction set that, if executed, will deliver the communication 150 (e.g., divided into transactions 152) to those of the targets T1-TY (see Figures 1-3 ) that are associated with the identifier MCID. However, as noted above, referring to Figure 1 The internal switches IS1C1-ISMCN can provide multiple paths through which the communication 150 can reach each of the targets T1-TY. Thus, referring to Figure 2 The routing data structure 410 (see Figure 4A ) can store multiple instruction sets for delivering the transactions 152 to those of the targets T1-TY that are associated with the identifier MCID. Each of the multiple instruction sets will be referred to as a "tree."

[0076] Figure 4A A block diagram illustrating an example implementation of a routing data structure 410 according to at least one embodiment is shown. The routing data structure 410 includes an index data portion 412, an instruction data portion 414, and a communication data portion 416. In the example shown, Figure 4A The routing data structure 410 includes data fields (e.g., columns) shown along a first dimension identified by double arrow AR1, and entries (e.g., rows) shown along a second dimension identified by double arrow AR2. The routing data structure 410 can include a different entry (e.g., row) for each unique value of the identifier MCID. To facilitate illustration, Figure 4A only three entries 421-423 are shown. However, the routing data structure 410 can include any number of entries. For example, the routing data structure 410 can include a predetermined number (e.g., 128) of entries.

[0077] Referring to Figure 2instructions 130 can cause the processor 122 to identify, along the second dimension (identified by double arrow AR2), a particular entry (e.g., row) of the values of the identifier MCID that are included in each transaction 152 of the communication 150. For example, with reference to Figure 4A The processor 122 can identify the second entry 422 as the particular entry. The particular entry includes a particular portion of each of the index data portion 412, the indication data portion 414, and the communication data portion 416. Thus, after having received one of the transactions 152, the instructions 130 cause the processor 122 to obtain the identifier MCID from the transaction and to look up the identifier MCID in the routing data structure 410 to identify the particular entry associated with the identifier MCID. The instructions 130 can cause the processor 122 to perform a function (e.g., a hash function) on information included in the communication 150. For example, the information can include a field, such as an address (e.g., a memory address), that is unique to the communication 150 and / or each transaction 152 in the communication 150.

[0078] The index data portion 412 stores links or pointers to the indication data portion 414. The index data portion 412 can include a predetermined number (e.g., 16) of pointer fields along the first dimension, each having a predetermined length (e.g., 5 bits). For ease of illustration, only two of these pointer fields are shown within the second entry 422 in Figure 4A , labeled PTR[0] and PTR[1].

[0079] The indication data portion 414 can store up to a predetermined number (e.g., 32) of indication fields along the first dimension, each having a predetermined length (e.g., 19 bits). Each indication field in the indication data portion 414 stores an indication. Within each entry, the indications define one or more trees (or indication sets), each identifying one or more port sets through which a transaction 152 can be sent (see Figure 2 ) to reach a target associated with the identifier MCID included in the transaction 152. As mentioned above, sometimes, each transaction 152 has to be sent in multiple rounds to one or more of the targets T1-TY. Thus, a tree can include multiple rounds in which the transactions 152 are each sent two or more times through one or more ports.

[0080] With reference to Figure 4A , in the illustrated example, the indication data portion 414 includes indication fields storing indications D1-D12 that define trees 430 and rounds 432. Each tree 430 includes at least one round 432. In Figure 4AIn the illustrated example, the second entry 422 includes two trees, labeled "TREE0" and "TREE1." The first and second trees "TREE0" and "TREE1" are alternate routes to reach the same target associated with the identifier MCID. The pointer field PTR[0] stores a first pointer to a first instruction Dl in the first tree "TREE0," and the pointer field PTR[l] stores a second pointer to a first instruction D8 in the second tree "TREE1." Thus, in the second entry 422, the values in the pointer fields PTR[0] and PTR[l] point the processor 122 to the trees "TREE0" and "TREE1," respectively. The instructions 130 cause the processor 122 to identify one of the pointer fields PTR[0] and PTR[l] of the index data portion 412 to thereby select one of the trees "TREE0" and "TREE1." The instructions 130 also cause the processor 122 to read (and / or parse) and implement the instructions of each round of the selected tree.

[0081] The trees 430 can each include any number of instructions, and at least two of the trees 430 can include different numbers of instructions. However, in some embodiments, the trees 430 can store up to a maximum number of instructions according to the size of a particular entry (e.g., the second entry 422). The instructions of the trees 430 can be organized into any number of rounds, and at least two of the trees 430 can include different numbers of rounds. Moreover, two or more rounds within the same tree can include different numbers of instructions. For example, the first tree "TREE0" includes seven instructions Dl-D7, which are organized into four rounds, labeled "TORndO," "TORndl," "TORnd2," and "TORnd3." The first round "TORndO" of the first tree "TREE0" includes the first instruction Dl, the second round "TORndl" includes instructions D2 and D3, the third round "TORnd2" includes instructions D4-D6, and the fourth round "TORnd3" includes a single instruction D7. As non-limiting examples, the instruction Dl can instruct the internal switch ISlCl to transmit a transaction on a first output port (e.g., port 211-1 in FIG. 2) of the second set O-TlCl, the instruction D2 can instruct the internal switch ISlCl to transmit a transaction on a different second output port (e.g., port 211-2 in FIG. 2) of the second set O-TlCl, the instruction D3 can instruct the internal switch ISlC2 to transmit a transaction on a first output port of the second set O-TlC2, the instruction D4 can instruct the internal switch ISlCl to transmit a transaction on a different third output port (e.g., port 211-3 in FIG. 2) of the second set O-TlCl, the instruction D5 can instruct the internal switch ISlC2 to transmit a transaction on a different fourth output port (e.g., port 211-4 in FIG. 2) of the second set O-TlC2, the instruction D6 can instruct the internal switch ISlC3 to transmit a transaction on a first output port of the second set O-TlC3, and the instruction D7 can instruct the internal switch ISlC3 to transmit a transaction on a different second output port (e.g., port 211-5 in FIG. 2) of the second set O-TlC3. Figure 2 Figure 2 Figure 2 ​​indication D5 can instruct the internal switch IS1C2 to transmit the transaction on a different second output port of the second set O-T1C2, indication D6 can instruct the internal switch IS1CN to transmit the transaction on the first output port of the second set O-T1CN, and indication D7 can instruct the internal switch IS1CN to transmit the transaction on a different second output port of the second set O-T1CN. Thus, within each of rounds "TO RndO" through "TO Rnd3", the transaction is sent only once through any of the internal switches IS1C1 - IS1CN. Further, in the example above, within the first tree "TREEO", the transaction is sent only once on each used port, but this is not required. In at least one embodiment, the transaction can be sent more than once through the same port within a particular tree.

[0082] As another non-limiting example, in Figure 4A , the second tree "TREE1" includes five indications D8 - D12, which are organized into two rounds, labeled "T1 RndO" and "T1 Rndl". The first round "T1 RndO" in the second tree "TREE1" includes indications D8 - D10 and the second round "T1 Rndl" includes indications D11 and D12. As a non-limiting example, indication D8 can instruct the internal switch IS1C1 to transmit the transaction on a first output port of the second set O-T1C1 (e.g., port 211-1 as shown in FIG. 2), indication D9 can instruct the internal switch IS1C2 to transmit the transaction on a first output port of the second set O-T1C2, indication D10 can instruct the internal switch IS1CN to transmit the transaction on a first output port of the second set O-T1CN, indication D11 can instruct the internal switch IS1C2 to transmit the transaction on a different second output port of the second set O-T1C2, and indication D12 can instruct the internal switch IS1CN to transmit the transaction on a different second output port of the second set O-T1CN. Figure 2 indication D9 can instruct the internal switch IS1C2 to transmit the transaction on a first output port of the second set O-T1C2, indication D10 can instruct the internal switch IS1CN to transmit the transaction on a first output port of the second set O-T1CN, indication D11 can instruct the internal switch IS1C2 to transmit the transaction on a different second output port of the second set O-T1C2, and indication D12 can instruct the internal switch IS1CN to transmit the transaction on a different second output port of the second set O-T1CN. Thus, within each of rounds "T1 RndO" and "T1 Rndl", the transaction is sent only once through any of the internal switches IS1C1 - IS1CN. Further, in the example above, within the second tree "TREE1", the transaction is sent only once through each used port, but this is not required. As described above, in at least one embodiment, the transaction can be sent more than once through the same port within a particular tree.

[0083] If the first tree "TREEO" is used to send the transaction 152 (see Figure 2 ), the internal switches IS1C1 - IS1CN (see Figure 1 and 3After receiving one of transactions 152, internal switch IS1C1 transmits the transaction according to instruction D1 in the first round "T0Rnd0". Next, after clock 120 increments, internal switch IS1C1 transmits the transaction according to instruction D2, and internal switch IS1C2 transmits the transaction according to instruction D3 in the second round "T0Rnd1". Then, after clock 120 increments again, internal switch IS1C1 transmits the transaction according to instruction D4, internal switch IS1C2 transmits the transaction according to instruction D5, and internal switch IS1CN transmits the transaction according to instruction D6 in the third round "T0Rnd2". Finally, after clock 120 increments again, internal switch IS1CN transmits the transaction according to instruction D7 in the fourth round "T0Rnd3". Similarly, if the second tree "TREE1" is used instead to transmit transactions, after internal switches IS1C1-IS1CN receive the transaction, internal switch IS1C1 transmits the transaction according to indication D8, internal switch IS1C2 transmits the transaction according to indication D9, and internal switch IS1CN transmits the transaction according to indication D10 in the first round "T1Rnd0". Then, after clock 120 increments, internal switch IS1C2 transmits the transaction according to indication D11, and internal switch IS1CN transmits the transaction according to indication D12 in the second round "T1Rnd1".

[0084] Figure 4B A portion 414, which can be used to encode instruction data according to at least one embodiment, is shown (see Figure 4A Each of the trees in the tree (e.g., Figure 4A Example format 460 shows the alternative trees “TREE0” and “TREE1”. As mentioned above, each tree represents the target T1-TY by the identifier MCID (see [link to example format]). Figure 2 Alternative routes to the identified target, and one or more in the tree may optionally include one or more rounds. As a non-limiting example, format 460 may include outermost left and right curly braces ("{" and "}") 462 and 464, which enclose a series of strings separated by commas. For ease of illustration, Figure 4B The text only shows format 466 for one of the strings in the series, but format 466 can be used to implement one of the strings in the series. In format 460, the outermost left and right curly braces 462 and 464 identify one of the alternative strings, and each string within the outermost left and right curly braces 462 and 464 (e.g., separated by commas) encodes one of the indicators (e.g., in...). Figure 4A (One of the indicators D1-D12 shown). Optionally, each string may be enclosed in a left curly brace and a right curly brace "{" and "}".

[0085] The format 466 can include a field or variable "last_rnd" that indicates whether the indication is a member of the last round in the tree. A string formatted according to the format 466 will include a value for the variable "last_rnd" that indicates whether the indication encoded in the string is a member of the last round in the tree. For example, the string can include a flag (e.g., implemented as a single bit) that is set (e.g., equals 1) when the indication encoded in the string is a member of the last round in the tree and is not set (e.g., equals 0) when the indication encoded in the string is not a member of the last round in the tree. Thus, the value for the variable "last_rnd" informs the processor 122 when the processor 122 reads and / or processes an indication of the last round in the tree.

[0086] The format 466 can include a field or variable "rnd_cont" that indicates whether the current round continues to the next indication or whether the current indication is the last indication in the current round. A string formatted according to the format 466 will include a value for the variable "rnd_cont" that indicates whether the current round continues. For example, the string can include a flag (e.g., implemented as a single bit) that is set (e.g., equals 1) when the current round continues and is not set (e.g., equals 0) when the current indication is the last indication in the current round.

[0087] The format 466 can include a field or variable "tcp[1:0]" that identifies which of the internal switches IS1C1- ISMCN are used to implement the indication. A string formatted according to the format 466 will include a value for the variable "tcp[1:0]" that identifies one of the internal switches IS1C1- ISMCN.

[0088] Each of the indications can store or encode a plurality of sub-indications that can be performed during the same increment of the clock 120 (see Figures 1 to 3 ) rather than during different increments. To facilitate illustration, the plurality of sub-indications will be described as including a pair of indications, referred to as an odd indication and an even indication. Reference will be made to Figure 4BIn the illustrated embodiment, odd instructions have been encoded by the variables "o_port[3:0]", "o_altpath", and "o_req_vchop[l:0]" and even instructions have been encoded by the variables "e_port[3:0]", "e_altpath", and "e_req_vchop[l:0]". The value of the variable "o_port[3:0]" indicates which port and / or bus (referred to as an odd port) through the internal switch (identified by the value of the variable "tcp[l:0]") to output the transaction 152. The value of the variable "e_port[3:0]" indicates which port and / or bus (referred to as an even port) through the internal switch (identified by the value of the variable "tcp[l:0]") to output the transaction 152.

[0089] The values of the variables "o_altpath" and "e_altpath" indicate which of the odd and / or even ports, respectively identified by the variables "o_port[3:0]" and "e_port[3:0]", through which the internal switch (identified by the value of the variable "tcp[l:0]") outputs the transaction 152. The string can include a flag (e.g., implemented as a single bit) for each of the variables "o_altpath" and "e_altpath" that indicates whether the transaction is sent through the odd port identified by the variable "o_port[3:0]" and / or the even port identified by the variable "e_port[3:0]". For example, when the internal switch is used to send the transaction through the odd port and not the even port, the flag for the variable "o_altpath" can be set (e.g., set equal to one) and the flag for the variable "e_altpath" can not be set (e.g., set equal to zero). As another non-limiting example, when the internal switch is used to send the transaction through the even port and not the odd port, the flag for the variable "o_altpath" can not be set (e.g., set equal to zero) and the flag for the variable "e_altpath" can be set (e.g., set equal to one). As another non-limiting example, when the internal switch sends the transaction through both the odd and even ports during the same clock increment, the flags for the variables "o_altpath" and "e_altpath" can both be set (e.g., set equal to one).

[0090] Depending on implementation details, the network 116 (see Figure 1 and Figure 3) can include one or more cycles that can cause a deadlock. A deadlock occurs when different transactions each wait for the other to release a resource (e.g., a port of one of the internal switches IS1C1 - ISMCN). The internal switches IS1C1 - ISMCN can each implement one or more virtual channels ("VCs") to help avoid such deadlocks. The variables "o_req_vchop[1 :0]" and "e_req_vchop[1 :0]" indicate whether a transaction will change or hop to a different VC from a current VC to help avoid a deadlock. For example, these variables can indicate that, if a transaction is routed along a backup path, the transaction hops to another VC from its current VC. As another non-limiting example, when a transaction crosses a date change line or transitions a period, the transaction can hop to another VC from its current VC.

[0091] Figure 5 A block diagram of a routing data structure 410 linked to a shared routing data structure 510 is shown in accordance with at least one embodiment. The communication data portion 416 and the shared routing data structure 510 (e.g., a routing table) can be stored in the data 132. At times, two or more different target groups can be reachable through the same tree (referred to as a shared backup tree). When this occurs, by storing the shared backup tree in a different shared routing data structure 510 (e.g., a routing table) rather than storing the shared backup tree in multiple entries of the routing data structure 410, space in the memory 124 (see Figures 1-3 ) can be conserved.

[0092] In the illustrated embodiment, the communication data portion 416 can include fields 520-530 defined along a first dimension (identified by the arrow AR1). For each entry (e.g., entries 421-423) in the routing data structure 410 (see Figure 4A and 5 ), the communication data portion 416 can store field values for the fields 520-530 used to implement multicasting within the system 100 (see Figure 1 and Figure 3). The first field 520 can store a value indicating whether the routing data structure 410 and the shared routing data structure 510 include entries of valid routing data. Thus, if the value of the first field 520 indicates that the routing data structure 410 and the shared routing data structure 510 do not include entries of valid routing data, the processor 122 can generate an error. The second field 522 can store a value indicating whether to use the interconnect conductors (e.g., the interconnect conductors 302 and 304) for that entry. In other words, the value of the second field 522 can be used to disable the interconnect conductors. The third field 524 can store a value indicating how many pointers are stored in the index data portion 412 and, thus, how many trees are stored in the indication data portion 414. The fourth field 526 can store a value indicating how many targets are included in the multicast. The value of the fourth field 526 can be used by the processor 122 to collect all of the responses from those targets of the targets T1-TY that are sent a transaction so that the processor 122 can send a single response to the source device (e.g., the source devices SD1-2). The fifth field 528 can store a link or pointer identifying the entry (e.g., row) in the shared routing data structure 510. The sixth field 530 can store a value indicating whether the value in the fifth field 528 is valid.

[0093] The shared routing data structure 510 can have a structure similar to the routing data structure 410. For example, the shared routing data structure 510 can be implemented using a linear table structure that can be described as representing two dimensions. The shared routing data structure 510 includes data fields (e.g., columns) shown along a first dimension identified by the double arrow AR4, and entries (e.g., rows) shown along a second dimension identified by the double arrow AR5. The shared routing data structure 510 can include a predetermined number (e.g., 16) of entries. For ease of illustration, Figure 5 Only three entries 531-533 are shown in the shared routing data structure 510. For example, the shared routing data structure 510 can include a predetermined number (e.g., 16) of entries. In the illustrated example, the fifth field 528 stores a link or pointer (shown as the arrow 536) identifying the third entry 533.

[0094] Like the routing data structure 410, the shared routing data structure 510 can include a shared index data portion 512 and a shared indication data portion 514 that are substantially similar to the index data portion 412 and the indication data portion 414, respectively. For example, the shared index data portion 512 can store links or pointers to the shared indication data portion 514. The shared index data portion 512 can include a predetermined number of pointer fields (e.g., 16) each having a predetermined length (e.g., 5 bits). For ease of illustration, Figure 5 Only two of these pointer fields are shown in the third entry 533, labeled S-prt[0] and S-prt[l], in the shared routing data structure 510.

[0095] The shared indication data portion 514 can store up to a predetermined number (e.g., 32) of indication fields for each entry, each indication field having a predetermined length (e.g., 19 bits). Each indication field in the shared indication data portion 514 stores an indication. For each entry, the indications define one or more shared trees 540, each shared tree having at least one round 542. In the illustrated example, the shared indication data portion 514 includes indication fields storing indications S-D1 through S-D7, which define shared trees 540 and rounds 532. Each shared tree 540 can be stored as a string or array using the format 460 (see Figure 4B ). In the illustrated example, the third entry 533 includes two trees, labeled "S-tree0" and "S-tree1". The first and second trees "S-tree0" and "S-tree1" are alternate routes to the same target. The pointer field S-prt[0] stores a pointer to the first indication S-D1 in the first tree "S-tree0" and the pointer field S-prt[1] stores a pointer to the first indication S-D5 in the second tree "S-tree1". Figure 5

[0096] The first tree "S-tree0" includes four indications S-D1 through S-D5, which are organized into two rounds, labeled "S0Rnd0" and "S0Rnd1". The first round "S0Rnd0" includes indications SD-1 and SD-2, and the second round "S0Rnd1" includes indications SD-3 and SD-4. As a non-limiting example, the indication S-D1 can indicate that the internal switch IS1C1 transmit a transaction on a first output port of the second set O-T1C1 (e.g., port 211-1 shown), the indication S-D2 can indicate that the internal switch IS1C2 transmit a transaction on a first output port of the second set O-T1C2, the indication S-D3 can indicate that the internal switch IS1C1 transmit a transaction on a different second output port of the second set O-T1C1 (e.g., port 211-2 shown), and the indication S-D4 can indicate that the internal switch IS1C2 transmit a transaction on a different second output port of the second set O-T1C2. Thus, in each of the rounds "S0Rnd0" through "S0Rnd3", a transaction is sent only once through any of the internal switches IS1C1-IS1CN. Further, in the example above, within the first tree "S-tree0", a transaction is sent only once on each used port, but this is not required. In at least one embodiment, a transaction can be sent more than once through the same port within a particular shared tree. Figure 2 Figure 2 ​​​

[0097] The second tree "S-treei" includes three directives S-D5 through S-D7, which are organized into two rounds, labeled "S1RndO" and "S1Rndi". The first round "S1RndO" includes the first directive S-D5 and the second round "S1Rndi" includes the directives S-D6 and S-D7. As a non-limiting example, the directive S-D5 can direct the internal switch IS1CN to transmit the transaction on a first output port of the second set O-T1CN, the directive S-D6 can direct the internal switch IS1C2 to transmit the transaction on an output port of the second set O-T1C2, and the directive S-D7 can direct the internal switch IS1CN to transmit the transaction on a different second output port of the second set O-T1CN. Thus, within each of the rounds "S1RndO" and "S1Rndi", the transaction is sent only once through any of the internal switches IS1C1 - IS1CN. Further, in the example above, within the second tree "S-treei", the transaction is sent only once through each port used, but this is not required. In at least one embodiment, the transaction can be sent more than once through the same port within a particular shared tree.

[0098] When the value of the first field 520 indicates that the shared routing data structure 510 includes entries of valid routing data, the fifth field 528 stores a shared link or shared pointer that identifies an entry (e.g., row) in the shared routing data structure, and the value of the sixth field 530 indicates that the value in the fifth field 528 is valid, the instructions 130 can cause the processor 122 to use the shared pointer (as indicated by arrow 536) to identify an entry (e.g., third entry 533) in the shared routing data structure 510. The instructions 130 can then cause the processor 122 to select one of the pointer fields (e.g., one of the pointer fields S-prt[0] and S-prt[l]) included in the particular portion of the index data portion 412. The values of the pointer fields S-prt[0] and S-prt[l] point the processor 122 to the trees "S-treeO" and "S-tree 1," respectively. For example, the values of the pointer fields S-prt[0] and S-prt[l] can point to the first indications S-Dl and S-D5 of the trees "S-treeO" and "S-tree 1," respectively. The instructions 130 can cause the processor 122 to execute a selection method (e.g., a hash function, a static method, a random number generator, and the like) to select one of the pointer fields S-prt[0] and S-prt[l] to thereby select one of the trees 540. As an example, the selection method can use information included in the transaction as input and can output an identifier of one of the pointers in the shared index data portion 512 that points to one of the backup trees 540. Such information can include an identifier of the source device SDl-2, a group identifier (e.g., the identifier MCID), and optionally, a memory address, when present. By using information in the transaction that is common to all transactions in the same communication, the processor 122 will select the same tree for all transactions in the same communication. Thus, the order of the transactions will be preserved. In other words, the transactions will be received by the destination Tl-TY associated with the group identifier in the same order that the transactions were received by the switch fabric 110, as all of the transactions will be sent by the switch fabric 110 through the same port and / or bus.

[0099] Figure 6 A flowchart of a method 600 that can be performed by the switch fabric 110 as the switch fabric 110 receives transactions is shown in accordance with at least one embodiment. For ease of illustration, the method 600 will be described as being performed, at least in part, by the internal switches ISIC1, ISIC2, and ISICN (see Figure 1 and Figure 3 ). However, the method 600 can be performed, at least in part, by any of the internal switches ISIC1- ISMCN (see Figure 1 and Figure 3 ). Prior to the method 600 beginning, the source devices SD1-1 through SD1-X (see Figure 1and Figure 3 ) to the switch circuit 110 (see Figures 1-3 ). For ease of illustration, the source devices SD1-2 will be described as transactions 152 (see Figure 2 ) sending communications 150 to the switch circuit 110. Each transaction 152 includes or is addressed to an identifier MCID, which in this example identifies a target T1-T3.

[0100] In a first block 602, the internal switches IS1C1, IS1C2, and IS1CN each receive a transaction 152A (see Figure 2 ) from the source devices SD1-2 as a current transaction. Alternatively, with reference to Figure 1 and Figure 3 , the current transaction can be routed to each of the internal switches IS1C1-IS1CN connected to the source devices SD1-2 by the bus 134. In block 602, after the internal switches IS1C1, IS1C2, and IS1CN receive the current transaction, they can each notify the processor 122 that they have received the current transaction. For example, with reference to Figure 2 , the internal switch circuit 200 of each of the internal switches IS1C1, IS1C2, and IS1CN can notify the processor 122. Alternatively, the instructions 130 can cause the processor 122 to know that the current transaction has been received. For example, the instructions 130 can cause the processor 122 to poll the internal switches IS1C1-ISMCN for transactions.

[0101] Then, in block 604 (see Figure 6 ), the instructions 130 cause the processor 122 to use information included in the current transaction (e.g., in a packet header) to identify a relevant portion of the data 132. In at least one embodiment, the instructions 130 can cause the processor 122 to identify an entry (e.g., Figure 4A and Figure 5 ) in the routing data structure 410 (see Figure 4A and Figure 5(One of the entries shown). For example, the current transaction may include group identifiers (e.g., identifier MCID) that identify two or more of the targets T1-TY. Optionally, the information may include the identifier of the source device SD1-2. If the current transaction is to be written to memory, then the current transaction may include a memory address that the internal switch IS1C1 may also use to identify the entry. In any case, this information (e.g., address) identifies communication 150 and prevents communication 150 from being confused with any other communication. Instruction 130 may cause processor 122 to execute a function (such as a hash function) on the information to obtain an index value, which processor 122 uses to identify the relevant portion of data 132 (e.g., to look up an entry in routing data structure 410).

[0102] In box 606 (see Figure 6 In block 604, instruction 130 causes processor 122 to use information included in the current transaction (e.g., in the packet header) to select a tree within the entry selected in block 604. For example, in block 606, instruction 130 causes processor 122 to use information included in the current transaction (e.g., in the packet header) to identify a tree 430 pointing to the routing data structure 410 (see [link to routing data structure 410]). Figure 4A One of the shared routing data structures 510 or the shared tree 540 (see) Figure 5 The instruction 130 causes the processor 122 to use a selection method (e.g., hash function, static method, random number generation, and similar) to select a pointer to either tree 430 or shared tree 540. As an example, the selection method can use information included in the current transaction as input and can output index data portion 412 (see [link to relevant documentation]). Figure 4A ) or shared index data section 512 (see Figure 5 The pointer is an identifier in one of the pointers, which points to one of the trees 430 or 540. Such information may include the identifier of the source devices SD1-2, the group identifier (e.g., identifier MCID), and optionally, the memory address, if present. By using information that identifies the communication in the current transaction and is common to all transactions in the same communication, the processor 122 will select the same tree for all transactions in the same communication. Therefore, the order of transactions will be maintained. In other words, transactions will be received by the targets T1-T3 in the same order as they were received by the switching circuit 110, because all transactions will be sent by the switching circuit 110 through the same port and / or bus.

[0103] The data in the routing data structure 410 and the optional shared routing data structure 510, when present, allow the switch circuit 110 to perform traffic shaping and help distribute load so that traffic will exhibit the desired static characteristics. For example, the selection method can use information other than that included in the current transaction as input, such as one or more previous trees selected, transmission latency times, and the like. For example, the switch circuit 110 can send different communications by using different trees in the different backup trees 430 to help load balance communications over the network 116. However, as noted above, transactions within a single communication can use the same tree to send to ensure that their order is maintained.

[0104] The routing data structure 410 and the optional shared routing data structure 510, when present, can be described as compressed representations of trees that are easily parsed by the hardware of the switch circuit 110. When present, the routing data structure 410 and the optional shared routing data structure 510 can each store the backup trees 430 and 540 as strings or arrays using a format 460 (see Figure 4B ) that can be read by the processor 122. Each indication informs the switch circuit 110 through which of the ports in the second set O-T1C1 to O-TMCN to send the current transaction.

[0105] In block 608, the instructions 130 cause the processor 122 to use the pointer selected in block 606 to locate and read the current round in the selected tree, which at this point is the first round in the tree selected in block 606. In other words, the processor 122 will read the first indication in the tree selected in block 606 and every indication that follows until the value of the variable "rnd_cont" indicates that the current round no longer continues. For example, the pointer field PTR[0] can have been selected in block 606, in which case the instructions 130 cause the processor 122 to read the first indication D1. Because the first indication D1 is the last indication in the first round "TO RndO", then only the first indication D1 will be read in block 608.

[0106] In optional block 610, the instructions 130 can cause the processor 122 to insert information, referred to as "breadcrumbs", into transactions sent according to any of the indications in the current round (e.g., in packet headers). As described below, the breadcrumbs can be used to perform a reduction operation. In embodiments that omit optional block 610, the instructions 130 can cause the processor 122 to proceed to block 612 after block 608.

[0107] Then, in block 612, the instructions 130 cause the processor 122 to send the current transaction according to any of the indications in the current round. For example, if the pointer field PTR[0] is selected in block 606, in block 612 the processor 122 can instruct the internal switch IS1C1 to send the current transaction through the output port of the second set O-T1C1 (e.g., port 211-1 as shown in FIG. 2B). Figure 2

[0108] In decision block 614, the instructions 130 cause the processor 122 to determine whether the current round is the last round in the selected tree. The determination in decision block 614 is "yes" when the current round is the last round in the selected tree. Otherwise, the determination in decision block 614 is "no". As a non-limiting example, the processor 122 can determine that the current indication is the last indication in the current round when the value of the variable "last_rnd" in the current indication indicates that the current indication is a member of the last round in the selected tree.

[0109] When the determination in decision block 614 is "no", in block 616, the instructions 130 cause the processor 122 to locate and read the next round in the selected tree, which becomes the current round. The processor 122 can then return to optional block 610, when present. In embodiments that omit optional block 610, the instructions 130 can cause the processor 122 to return to block 612.

[0110] On the other hand, when the determination in decision block 614 is "yes", the processor 122 proceeds to block 618. In block 618, the instructions 130 cause the processor 122 to wait for a new transaction (e.g., transaction 152B). Upon receiving the new transaction, the processor 122 returns to block 604.

[0111] Referring to Figure 1 ​, system 100 can use multicasting to implement one or more reduction operations. Reduction operations perform mathematical operations, such as an ADD operation, on multiple operands (e.g., data included in a transaction) to create a single result. For example, multiple operands can be received from different GPUs that are combined together to solve a particular problem, such as training a neural network. While reduction operations can be performed in the GPUs, reduction operations can also be performed, at least in part, by switch circuit 110. Using switch circuit 110 to perform reduction operations can free up the GPUs for other tasks. Moreover, the GPUs do not have to locally collect all of the operands, which can save network bandwidth. Non-limiting examples of reduction operations that can be performed by switch circuit 110 include a MIN operation that returns the minimum value of multiple values, a MAX operation that returns the maximum value of multiple values, an ADD operation that returns the sum of multiple values, an AND operation that returns the aggregate (e.g., concatenated string, array, and the like) of multiple values, an OR operation, an XOR operation, and custom operations.

[0112] Figure 7 A flowchart of a method 700 that can be performed by switch circuit 110 (see Figures 1-3 ) when switch circuit 110 performs a reduction operation is shown in accordance with at least one embodiment. Some reduction operations, such as floating point addition, are not associative operations; the order of operations can matter. Thus, when using multicasting to perform such reduction operations, switch circuit 110 can be used to combine responses to a transaction in the same order, regardless of the order in which the responses are received by switch circuit 110. Doing so makes these reduction operations repeatable, which is especially important to programmers.

[0113] In a first block 702, switch circuit 110 performs method 600 (see Figure 6 ) and sends the current transaction received from a source device (e.g., source devices SD1-2) to each of the portion of targets T1-TY. As described above, in optional block 610 (see Figure 6 ), instructions 130 can cause processor 122 to add a breadcrumb to each copy of the current transaction sent to the portion of targets T1-TY (e.g., in a packet header). The breadcrumb can indicate the order specified by the tree selected in block 606 (see Figure 6 ). The breadcrumb can also indicate which reduction operation to perform on the responses.

[0114] Next, in block 704, the switch circuit 110 receives a current response to one of the copies of the current transaction sent in block 702. The breadcrumb included in the copy of the current transaction can also be included in the current response by the target that receives the copy of the current transaction. Eventually, the switch circuit 110 can receive a response for each copy of the current transaction sent in block 702.

[0115] In block 706, the instructions 130 cause the processor 122 to read the breadcrumb in the current response received in block 704.

[0116] In block 708, the instructions 130 cause the processor 122 to store the current response in a location according to the information read from the breadcrumb in block 706. In other words, the breadcrumb can be used to specify the position that the current response occupies in the order specified by the tree selected in block 606 (see Figure 6 ). The instructions 130 cause the processor 122 to place the current response in a location corresponding to its position in the order. For example, the current response and other responses to the current transaction can be stored in particular memory locations that indicate the order of the responses specified by the tree selected in block 606.

[0117] In decision block 710, the instructions 130 cause the processor 122 to determine whether all responses to the current transaction have been received. When all responses to the current transaction have been received, the determination in decision block 710 is "yes." Otherwise, the determination in decision block 710 is "no." As a non-limiting example, the processor 122 can determine that all responses have been received when the number of responses received to the current transaction matches the value of the fourth field 526 (see Figure 5 ) of the entry used to send the current transaction in block 702.

[0118] When the determination in decision block 710 is "no," the switch circuit 110 returns to block 704 and receives another response to the current transaction.

[0119] On the other hand, when the determination in decision block 710 is "yes," in block 712, the instructions 130 cause the processor 122 to combine the responses and produce a single response. The breadcrumb can also indicate how the responses are to be combined. For example, the breadcrumb can identify a particular reduction operation to be performed on the responses. Thus, after the responses are collected, the processor 122 can perform a reduction operation on the responses in the order specified by the breadcrumb to produce a single response. For example, if the reduction operation is floating point addition, the processor 122 can add the values included in the responses (e.g., the packet payloads) according to the order specified by the breadcrumb.

[0120] Next, in block 714, the instructions 130 cause the processor 122 to send a single response to the source device (e.g., source device SD1-2). Then, the method 700 terminates.

[0121] In the method 700, the indication (e.g., in the string of the format 460 shown in Figure 4B Figure 6B) provides the order of the reduction sequence. When the processor 122 parses the string (e.g., in blocks 608 and 616 of Figure 6A) to multicast each transaction, the processor 122 can add a breadcrumb to the transaction (e.g., in optional block 610), and can assign reduction resources (e.g., memory locations) using the order specified by the tree selected in block 606 (see Figure 6A), which will ensure that the reduction operation performed using the multicast is performed the same way each time. Thus, the processor 122 can assign the reduction resources in a precise way that ensures the order. Figure 6 Figure 6 In the method 700, the indication (e.g., in the string of the format 460 shown in

[0122] Data Center

[0123] Figure 8 An example data center 800 is shown in accordance with at least one embodiment. In at least one embodiment, data center 800 includes, without limitation, a data center infrastructure layer 810, a framework layer 820, a software layer 830, and an application layer 840.

[0124] In at least one embodiment, as shown in Figure 8 Figure 6B, the data center infrastructure layer 810 can include a resource orchestrator 812, grouped computing resources 814, and node computing resources (“node C.R.”) 816(1)-816(N), where “N” represents any whole, positive integer. In at least one embodiment, node C.R.s 816(1)-816(N) can include, without limitation, any number of central processing units (“CPUs” or other processors (including accelerators, field programmable gate arrays (“FPGAs”), data processing units (DPUs) in network devices, graphics processors, etc.), memory devices (e.g., dynamic random access memory), storage devices (e.g., solid state drives or disk drives), network input / output (“NW I / O”) devices, network switches, virtual machines (“VMs”), power modules, and cooling modules, etc. In at least one embodiment, one or more of node C.R.s 816(1)-816(N) can be a server having one or more of the above computing resources.

[0125] ​In at least one embodiment, the grouped computing resources 814 may include individual groups (not shown) of node CRs housed in one or more racks, or a plurality of racks (also not shown) housed in data centers in various geographical locations. The individual groups of node CRs within the grouped computing resources 814 may include computing, networking, memory, or storage resources that can be configured or allocated to support groups of one or more workloads. In at least one embodiment, several node CRs, including CPUs or processors, may be grouped within one or more racks to provide computing resources to support one or more workloads. In at least one embodiment, the one or more racks may also include any number of power modules, cooling modules, and network switches, in any combination.

[0126] In at least one embodiment, resource coordinator 812 may configure or otherwise control one or more nodes CR816(1)-816(N) and / or grouped computing resources 814. In at least one embodiment, resource coordinator 812 may include a software design infrastructure (“SDI”) management entity for data center 800. In at least one embodiment, resource coordinator 812 may include hardware, software, or some combination thereof.

[0127] In at least one embodiment, such as Figure 8 As shown, the framework layer 820 includes, but is not limited to, a job scheduler 832, a configuration manager 834, a resource manager 836, and a distributed file system 838. In at least one embodiment, the framework layer 820 may include a framework of software 852 supporting the software layer 830 and / or one or more applications 842 of the application layer 840. In at least one embodiment, the software 852 or application 842 may respectively include web-based service software or applications, such as services or applications provided by Amazon Web Services, Google Cloud, and Microsoft Azure. In at least one embodiment, the framework layer 820 may be, but is not limited to, a free and open-source software web application framework, such as Apache Spark, which can utilize the distributed file system 838 for large-scale data processing (e.g., "big data"). TM(“Spark”). In at least one embodiment, job scheduler 832 can include a Spark driver to facilitate scheduling of workloads supported by various tiers of data center 800. In at least one embodiment, configuration manager 834 can be capable of configuring different tiers, such as software tier 830 and framework tier 820 including Spark and a distributed file system 838 for supporting large scale data processing. In at least one embodiment, resource manager 836 can be capable of managing clustered or grouped computing resources mapped to or allocated for supporting distributed file system 838 and job scheduler 832. In at least one embodiment, clustered or grouped computing resources can include grouped computing resources 814 on data center infrastructure layer 810. In at least one embodiment, resource manager 836 can coordinate with resource orchestrator 812 to manage these mapped or allocated computing resources.

[0128] In at least one embodiment, software 852 included in software tier 830 can include software used by at least a portion of node C.R.s 816(1)-816(N), grouped computing resources 814, and / or distributed file system 838 of framework tier 820. One or more types of software can include, but are not limited to, Internet web page search software, email virus scanning software, database software, and streaming video content software.

[0129] In at least one embodiment, one or more application programs 842 included in application tier 840 can include one or more types of application programs used by at least a portion of node C.R.s 816(1)-816(N), grouped computing resources 814, and / or distributed file system 838 of framework tier 820. One or more types of application programs can include, but are not limited to, CUDA application programs.

[0130] In at least one embodiment, any of configuration manager 834, resource manager 836, and resource orchestrator 812 can implement any number and type of self-modifying actions based on any number and type of data acquired in any technically feasible fashion. In at least one embodiment, self-modifying actions can mitigate data center operators of data center 800 making possibly poor configuration decisions and can avoid underutilized and / or poorly performing portions of a data center.

[0131] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement data center 800. For example, cluster 112 and / or group 114 can include one or more of grouped computing resources 814 and / or one or more of C.R.s 816(1)-816(N).

[0132] Computer-based system

[0133] The following figures present, without limitation, exemplary computer-based systems that can be used to implement at least one embodiment.

[0134] Figure 9 A processing system 900, in accordance with at least one embodiment, is shown. In at least one embodiment, system 900 includes one or more processor(s) 902 and one or more graphics processor(s) 908, and can be a single processor desktop system, a multiprocessor workstation system, or a server system having many processor(s) 902 or processor core(s) 907. In at least one embodiment, processing system 900 is a processing platform incorporated within a system- on-a-chip (SoC) integrated circuit for use in mobile, handheld, or embedded devices.

[0135] In at least one embodiment, processing system 900 can include, or be incorporated within a server-based gaming platform, a game console, a media console, a mobile gaming console, a handheld game console, or an online game console. In at least one embodiment, processing system 900 is a mobile phone, a smart phone, a tablet device, or a mobile internet device. In at least one embodiment, processing system 900 can also include, couple with, or be integrated within a wearable device, such as a smart watch wearable device, a smart eyewear device, an augmented reality device, or a virtual reality device. In at least one embodiment, processing system 900 is a television or set-top box device having one or more processors 902 and a graphical interface generated by one or more graphics processors 908.

[0136] In at least one embodiment, one or more processor(s) 902 each include one or more processor core(s) 907 to process instructions which, when executed, implement the operations for system and user software. In at least one embodiment, each of the one or more processor core(s) 907 are configured to process a specific instruction set 909. In at least one embodiment, instruction set 909 can facilitate Complex Instruction Set Computing (CISC), Reduced Instruction Set Computing (RISC), or computing via a Very Long Instruction Word (VLIW). In at least one embodiment, multiple processor core(s) 907 can each process a different instruction set 909, which can include instructions to facilitate emulation of other instruction sets. In at least one embodiment, processor core(s) 907 can also include other processing devices, such as a digital signal processor (DSP).

[0137] In at least one embodiment, processor 902 includes cache memory 904. In at least one embodiment, processor 902 can have a single internal cache or multiple levels of internal caches. In at least one embodiment, cache memory is shared among multiple components of processor 902. In at least one embodiment, processor 902 also uses an external cache (e.g., a level three (L3) cache or last level cache (LLC)) (not shown), which can be shared between processor cores 907 using known cache coherency techniques. In at least one embodiment, additionally included in processor 902 are register file 906, which processor 902 can include different types of registers to store different types of data (e.g., integer registers, floating point registers, status registers, and instruction pointer registers). In at least one embodiment, register file 906 can include general registers or other registers.

[0138] In at least one embodiment, one or more processors 902 are coupled with one or more interface buses 910 for communicating information to and from system 900 and other components in system 900. In at least one embodiment, interface bus 910 can be implemented using one or more of various types of communication buses, including an address bus for communicating addresses, a data bus for communicating data, and a control bus for communicating controls. In at least one embodiment, interface bus 910 is not limited to the DMI bus, and can include one or more of a Peripheral Component Interconnect bus (e.g., a PCI, a PCI Express), a memory bus, or another type of interface bus. In at least one embodiment, processor 902 includes an integrated memory controller 916 and platform controller hub 930. In at least one embodiment, memory controller 916 facilitates communication between memory 920 and other components of system 900, while platform controller hub 930 provides connections to input / output (I / O) devices to local I / O bus.

[0139] In at least one embodiment, storage device 920 can be a dynamic random access memory (DRAM) device, a static random access memory (SRAM) device, flash memory device, or a phase change memory device, among others. In at least one embodiment, storage device 920 can be used as a main memory for processing system 900. In at least one embodiment, storage device 920 can be used for storage of data 922 and instructions 921 for use when implementing applications or processes by one or more processors 902. In at least one embodiment, a memory controller 916 can also be coupled with the optional external graphics processor 912, which can communicate with one or more graphics processors 908 within one or more processors 902 to perform graphics and media operations. In at least one embodiment, a display device 911 can be coupled to one or more processors 902. In at least one embodiment, display device 911 can include one or more of an internal display device, as in a mobile electronic device or a laptop computer or an external display device attached via a display interface (e.g., DisplayPort, etc.). In at least one embodiment, display device 911 can include a head mounted display (HMD) such as a stereoscopic display device for use in virtual reality (VR) applications or augmented reality (AR) applications.

[0140] In at least one embodiment, platform controller hub 930 enables peripherals to connect to storage device 920 and processor 902 via a high-speed I / O bus. In at least one embodiment, I / O peripherals include, without limitation, audio controller 946, network controller 934, firmware interface 928, wireless transceiver 926, touch sensors 925, data storage device 924 (e.g., hard disk drive, flash memory, etc.). In at least one embodiment, data storage device 924 can connect to the storage interface via a storage interface bus, e.g., a SATA, or a peripheral bus, e.g., a peripheral component interconnect bus (such as PCI, PCI Express). In at least one embodiment, touch sensors 925 can include touch screen sensors, pressure sensors, or fingerprint sensors. In at least one embodiment, wireless transceiver 926 can be a Wi-Fi transceiver, a Bluetooth transceiver, or a mobile network transceiver such as a 3G, 4G, or Long Term Evolution (LTE) transceiver. In at least one embodiment, firmware interface 928 enables communication with system firmware, and can be, e.g., a unified extensible firmware interface (UEFI) according to at least one embodiment. In at least one embodiment, network controller 934 can enable network connectivity to a wired network. In at least one embodiment, a high-performance network controller (not shown) is coupled

[0141] In at least one embodiment, memory controller 916 and instances of platform controller hub 930 can be integrated into a discrete external graphics processor, such as external graphics processor 912. In at least one embodiment, platform controller hub 930 and / or memory controller 916 can be external to one or more processor(s) 902. For example, in at least one embodiment, processing system 900 can include an external memory controller 916 and platform controller hub 930, which can be configured as a memory controller hub and a peripheral controller hub in a system chipset that is in communication with processor(s) 902.

[0142] In at least one embodiment, system 100 (see Figure 1 and 3) can be used to implement processing system 900. In at least one embodiment, set 112 and / or group 114 can include one or more of processors 902, one or more of processor cores 907, and / or one or more of graphics processors 908. In at least one embodiment, interface bus 910 can include switch circuitry 110 (see Figures 1-3

[0143] Figure 10 A computer system 1000 according to at least one embodiment is shown. In at least one embodiment, computer system 1000 can be a system with interconnected devices and components, a SOC, or some combination. In at least one embodiment, computer system 1000 is formed from a processor 1002 that can include execution units to execute an instruction. In at least one embodiment, computer system 1000 can include, without limitation, components such as processor 1002 that employ execution units including logic to perform algorithms for processing data. In at least one embodiment, computer system 1000 can include processors such as Pentium®, Core®, Xenon®, Xeon®, Itanium®, XScale™, and / or StrongARM™, ARM® Cortex-M1™, which can be obtained from Core TM or Nervana TM microprocessors, though the scope of a embodiment is not so limited. In at least one embodiment, computer system 1000 can execute a version of the WINDOWS operating system available from Microsoft Corporation of Redmond, Wash. In at least one embodiment, other operating systems, including UNIX and Linux, as well as embedded software, can also be used. In at least one embodiment, computer system 1000 can be a server, desktop, laptop, or other computing system.

[0144] ​In at least one embodiment, the computer system 1000 can be used in other devices, such as handheld devices and embedded applications. Some examples of handheld devices include cellular phones, Internet Protocol (IP) devices, digital cameras, personal digital assistants (“PDAs”), and handheld PCs. In at least one embodiment, the embedded application may include a microcontroller, a digital signal processor (“DSP”), a system-on-a-chip (SoC), a network computer (“NetPC”), a set-top box, a network hub, a wide area network (“WAN”) switch, or any other system capable of executing one or more instructions according to at least one embodiment.

[0145] In at least one embodiment, the computer system 1000 may include, but is not limited to, a processor 1002, which may include, but is not limited to, one or more execution units 1008 configured to execute a Computational Unified Device Architecture (“CUDA”). (Developed by NVIDIA Corporation, Santa Clara, California) In at least one embodiment, the CUDA program is at least a part of a software application written in the CUDA programming language. In at least one embodiment, the computer system 1000 is a single-processor desktop or server system. In at least one embodiment, the computer system 1000 may be a multiprocessor system. In at least one embodiment, the processor 1002 may include, but is not limited to, a CISC microprocessor, a RISC microprocessor, a VLIW microprocessor, a processor implementing instruction set combinations, or any other processor device, such as a digital signal processor. In at least one embodiment, the processor 1002 may be coupled to a processor bus 1010, which can transmit data signals between the processor 1002 and other components in the computer system 1000.

[0146] In at least one embodiment, processor 1002 may include, but is not limited to, a Level 1 (“L1”) internal cache memory (“cache”) 1004. In at least one embodiment, processor 1002 may have a single internal cache or multiple levels of internal cache. In at least one embodiment, the cache memory may reside external to processor 1002. In at least one embodiment, processor 1002 may include a combination of internal and external caches. In at least one embodiment, register file 1006 may store different types of data in various registers, including but not limited to integer registers, floating-point registers, status registers, and instruction pointer registers.

[0147] In at least one embodiment, execution unit 1008 includes, without limitation, logic to perform integer and floating point operations, including the execution of intemiediate and floating point operations in accordance with an instruction set architecture. The processor 1002 can also include a microcode (“ucode”) read only memory (“ROM”) that stores microcode for certain macroinstructions. In at least one embodiment, execution unit 1008 can also include logic to handle a packed data instruction set 1009. In at least one embodiment, by including the packed data instruction set 1009 in an instruction set architecture for a general -purpose processor 1002, many multimedia applications can be accelerated via the use of this extended instruction set without the need for specific multimedia processors. In at least one embodiment, 1009 can be an instruction set that includes instructions for operating on multi-byte operands in a single instruction. This can be performed on the full width of the processor’s data bus in a single clock cycle. In at least one embodiment, this can eliminate the need to use multiple instructions to process multi-byte data in multiple clock cycles.

[0148] In at least one embodiment, execution unit 1008 can also be used in a microcontroller, embedded processor, graphics device, DSP, and other types of logic circuits. In at least one embodiment, computer system 1000 can include, without limitation, a memory 1020. In at least one embodiment, memory 1020 can be implemented as a DRAM device, SRAM device, flash memory device, or other storage device. Memory 1020 can store data signals represented by a data signal in a processor 1002 executable instructions 1019 and / or data 1021.

[0149] In at least one embodiment, a system logic chip can be coupled to processor bus 1010 and memory 1020. In at least one embodiment, system logic chip can include, without limitation, a memory controller hub (“MCH”) 1016, and processor 1002 can communicate with MCH 1016 via processor bus 1010. In at least one embodiment, MCH 1016 can provide a high bandwidth memory path 1018 to memory 1020 for instruction and data storage and for storage of graphics commands, data, and textures. In at least one embodiment, MCH 1016 can also include an integrated graphics processing unit (“GPU”) 1014, or MCH 1016 can be separate from the integrated GPU 1014. In at least one embodiment, MCH 1016 can enable processor 1002 to access I / O devices 1022 via memory 1020. In at least one embodiment, MCH 1016 can bridge between the processor bus 1010 and the memory 1020, and the system I / O 1022.

[0150] In at least one embodiment, system logic chip can provide a graphics port for coupling to a graphics controller. In at least one embodiment, MCH 1016 can be coupled to memory 1020 through a high bandwidth memory path 1018, and a graphics / video card 1012 can be coupled to MCH 1016 through an Accelerated Graphics Port (“AGP”) interconnect 1014.

[0151] In at least one embodiment, computer system 1000 can use system I / O 1022 as a proprietary hub interface bus to couple MCH 1016 to I / O controller hub (“ICH”) 1030. In at least one embodiment, ICH 1030 can provide direct connections to some I / O devices and high-speed

[0152] In at least one embodiment, Figure 10 A system including interconnected hardware devices or “chips” is shown. In at least one embodiment, Figure 10 An exemplary SoC can be shown. In at least one embodiment, Figure 10 Devices shown in FIG. 1 can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, one or more components of system 1000 are interconnected using a Compute Express Link (CXL) interconnect.

[0153] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement processing system 1000. In at least one embodiment, set 112 and / or group 114 can include processor 1002 and / or graphics / video card 1012. In at least one embodiment, processor bus 1010 can include switch circuitry 110 (see Figures 1-3 ), bus 134, and / or bus 140.

[0154] Figure 11A system 1100 is shown, in accordance with at least one embodiment. In at least one embodiment, system 1100 is an electronic device that utilizes a processor 1110. In at least one embodiment, system 1100 can be, for example and without limitation, a laptop, a tower server, a rack server, a blade server, an edge device communicatively coupled to one or more local or cloud service providers, a laptop, a desktop, a tablet, a mobile device, a phone, an embedded computer, or any other suitable electronic device.

[0155] In at least one embodiment, system 1100 can include, without limitation, a processor 1110 communicatively coupled to any suitable number or kind of components, peripherals, modules, or devices. In at least one embodiment, processor 1110 is coupled using a bus or interface such as an Industry Standard 2 C bus, a System Management Bus (“SMBus”), a Low Pin Count (LPC) bus, a Serial Peripheral Interface (“SPI”), a High Definition Audio (“HDA”) bus, a Serial Advanced Technology Attachment (“SATA”) bus, a USB (versions 1, 2, 3), or a Universal Asynchronous Receiver / Transmitter (“UART”) bus. In at least one embodiment, system 1100 can include, without limitation, a processor 1110 communicatively coupled to any suitable number or kind of components, peripherals, modules, or devices. In at least one embodiment, processor 1110 is coupled using a bus or interface such as an Industry Standard Figure 11 A system is shown that includes interconnected hardware devices or “chips.” In at least one embodiment, system 1100 is an SoC. Figure 11 An exemplary SoC can be shown. In at least one embodiment, system 1100 is an SoC. Figure 11 Devices shown in FIG. 11 can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, one or more components of system 1100 are interconnected using Compute Express Link (CXL) interconnects. Figure 11 Devices shown in FIG. 11 can be interconnected with proprietary interconnects, standardized interconnects (e.g., PCIe), or some combination thereof. In at least one embodiment, one or more components of system 1100 are interconnected using Compute Express Link (CXL) interconnects.

[0156] In at least one embodiment, system 1100 can include, without limitation, a processor 1110 communicatively coupled to any suitable number or kind of components, peripherals, modules, or devices. In at least one embodiment, processor 1110 is coupled using a bus or interface such as an Industry Standard Figure 11 may include a display 1124, a touchscreen 1125, a touchpad 1130, a near field communication unit (“NFC”) 1145, a sensor hub 1140, a thermal sensor 1146, an Express Chipset (“EC”) 1135, a Trusted Platform Module (“TPM”) 1138, a BIOS / Firmware / Flash (“BIOS, FW Flash”) 1122, a DSP 1160, a Solid State Disk (“SSD”) or Hard Disk Drive (“HDD”) 1120, a Wireless Local Area Network unit (“WLAN”) 1150, a Bluetooth unit 1152, a Wireless Wide Area Network unit (“WWAN”) 1156, a Global Positioning System (GPS) 1155, a camera (“USB 3.0 camera”) 1154 (e.g., a USB 3.0 camera), or a Low Power Double Data Rate (“LPDDR”) memory unit (“LPDDR3”) 1115 implemented in, for example, LPDDR3 standard. These components can each be implemented in any suitable manner.

[0157] In at least one embodiment, other components can be communicatively coupled to processor 1110 by components discussed above. In at least one embodiment, an accelerometer 1141, an ambient light sensor (“ALS”) 1142, a compass 1143, and a gyroscope 1144 can be communicatively coupled to a sensor hub 1140. In at least one embodiment, a thermal sensor 1139, a fan 1137, a keyboard 1136, and a touchpad 1130 can be communicatively coupled to EC 1135. In at least one embodiment, a speaker 1163, a headphone 1164, and a microphone (“mic”) 1165 can be communicatively coupled to an audio unit (“audio codec and class D amplifier”) 1162, which in turn can be communicatively coupled to a DSP 1160. In at least one embodiment, audio unit 1162 can include, for example and without limitation, an audio coder / decoder (“codec”) and a class D amplifier. In at least one embodiment, a SIM card (“SIM”) 1157 can be communicatively coupled to a WWAN unit 1156. In at least one embodiment, components such as WLAN unit 1150 and Bluetooth unit 1152, as well as WWAN unit 1156, can be implemented in a next generation form factor (NGFF).

[0158] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement system 1100.

[0159] Figure 12 An exemplary integrated circuit 1200, in accordance with at least one embodiment, is shown. In at least one embodiment, exemplary integrated circuit 1200 is a SoC, which can be fabricated using one or more IP cores. In at least one embodiment, integrated circuit 1200 includes one or more application processors 1205 (e.g., CPUs, DPUs), at least one graphics processor 1210, and can additionally include an image processor 1215 and / or a video processor 1220, any of which can be a modular IP core. In at least one embodiment, integrated circuit 1200 includes peripheral or bus logic including USB controller(s) 1225, UART controller(s) 1230, SPI / SDIO controller(s) 1235, and I2S / I2C bus 2 S / I 2C controller 1240. In at least one embodiment, integrated circuit 1200 can include a display device 1245 coupled to one or more of a high-definition multimedia interface (HDMI) controller 1250 and a mobile industry processor interface (MIPI) display interface 1255. In at least one embodiment, storage can be provided by a flash memory subsystem 1260 including flash memory and a flash memory controller. In at least one embodiment, a memory interface can be provided via a memory controller 1265 for accessing SDRAM or SRAM memory devices. In at least one embodiment, some integrated circuits also include an embedded security engine 1270.

[0160] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement integrated circuit 1200. In at least one embodiment, clusters 112 and / or groups 114 can include one or more of application processors 1205, one or more of graphics processors 1210, image processors 1215, and / or video processors 1220. In at least one embodiment, refer to Figure 1 and Figure 3 , in at least one embodiment, peripheral or bus logic of integrated circuit 1200 (see Figure 12 ) can include switching circuitry 110, bus 134, and / or bus 140.

[0161] Figure 13 A computing system 1300, in accordance with at least one embodiment, is shown. In at least one embodiment, computing system 1300 includes a processing subsystem 1301 with one or more processor(s) 1302 and system memory 1304 communicating via an interconnection path 1305 that can include a memory hub 1305. In at least one embodiment, memory hub 1305 can be a separate component coupled with one or more processors 1302 via individual communication links 1307A to 1307N. In at least one embodiment, memory hub 1305 can be integrated into one or more processors 1302 or can be a standalone component.

[0162] In at least one embodiment, processing subsystem 1301 includes one or more parallel processors 1312 coupled to memory hub 1305 via a bus or other communication link 1313. In at least one embodiment, communication link 1313 can be one of many based on well-known bus or communications technologies, such as PCI (Peripheral Component Interconnect) based links, or other communications interfaces and protocols. In at least one embodiment, one or more parallel processors 1312 form a programmable computing platform, which can include a large number of processing cores and / or processing clusters, such as a many integrated core (MIC) processor. In at least one embodiment, one or more parallel processors 1312 form a graphics processing subsystem that can output pixels to one of one or more display devices 1310A coupled via I / O Hub 1307. In at least one embodiment, one or more parallel processors 1312 can also include a display controller and display interface (not shown) to enable a direct connection to one or more display devices 1310B.

[0163] In at least one embodiment, system storage 1314 can be connected to I / O Hub 1307 to provide a storage mechanism for computing system 1300. In at least one embodiment, I / O

[0164] In at least one embodiment, computing system 1300 can include other components not explicitly shown, including USB or other port connections, optical storage drives, video capture devices, and like, which can also be connected to I / O Hub 1307. In at least one embodiment, communication paths interconnecting various components in computing system 1300 can implement one or more busses that can include any number of bus standards, including a PCI bus, a PCI-Express bus and the like. Figure 13 In at least one embodiment, communication paths interconnecting various components in computing system 1300 can use any suitable protocol, such as a protocol based on the PCI (Peripheral Component Interconnect) standard, for example, PCI Express, or the like, or other bus or point-to-point communication interfaces and / or protocols (e.g., NVLink high-speed interconnect, or interconnect protocols).

[0165] In at least one embodiment, parallel processor(s) 1312 include circuitry optimized for graphics and video processing, including for example video output circuitry, and are configured for use in a gaming console, a mobile phone, a personal computer, or other application. In at least one embodiment, parallel processor(s) 1312 include circuitry optimized for general use applications, including for example high-precision floating point, integer and Boolean operations on a large number of operands. In at least one embodiment, computing system 1300 can include multiple parallel processor(s) 1312. In at least one embodiment, one or more of parallel processor(s) 1312 can be configured for use in a server. In at least one embodiment, one or more of parallel processor(s) 1312 can be configured for use in a cloud computing platform.

[0166] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement computing system 1300. In at least one embodiment, cluster 112 and / or group 114 can include one or more of processor(s) 1302 and / or one or more of parallel processor(s) 1312. In at least one embodiment, communication link(s) 1313 can include switch circuitry 110, bus 134, and / or bus 140.

[0167] Processing system

[0168] The following figures illustrate, but are not limited to, exemplary processing systems that can be used to implement at least one embodiment.

[0169] Figure 14An accelerated processing unit (“APU”) 1400, in accordance with at least one embodiment, is shown. In at least one embodiment, APU 1400 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, APU 1400 can be configured to execute application programs such as CUDA programs. In at least one embodiment, APU 1400 includes, without limitation, core complex 1410, graphics complex 1440, fabric 1460, I / O interface 1470, memory controllers 1480, display controllers 1492, and multimedia engine 1494. In at least one embodiment, APU 1400 can include, without limitation, any combination of any number of core complexes 1410, any number of graphics complexes 1440, any number of display controllers 1492, and any number of multimedia engines 1494. For purposes of illustration, multiple instances of like objects are denoted with reference numerals including a leading asterisk and a number identifying the object, and a number in parentheses identifying the instance.

[0170] In at least one embodiment, core complex 1410 is a CPU, graphics complex 1440 is a GPU, and APU 1400 is a processing unit that integrates, without limitation, core complex 1410 and graphics complex 1440 onto a single chip. In at least one embodiment, some tasks can be assigned to core complex 1410 while other tasks can be assigned to graphics complex 1440. In at least one embodiment, core complex 1410 is configured to execute host control software associated with APU 1400, such as an operating system. In at least one embodiment, core complex 1410 is a master processor of APU 1400 that controls and coordinates the operation of other processors. In at least one embodiment, core complex 1410 issues commands that control the operation of graphics complex 1440. In at least one embodiment, core complex 1410 can be configured to execute host executable code derived from CUDA source code, and graphics complex 1440 can be configured to execute device executable code derived from CUDA source code.

[0171] In at least one embodiment, core complex 1410 includes, without limitation, cores 1420(1)-1420(4) and L3 cache 1430. In at least one embodiment, core complex 1410 can include, without limitation, any number of cores 1420 and any number and type of caches in any combination. In at least one embodiment, cores 1420 are configured to execute instructions of a particular instruction set architecture (“ISA”). In at least one embodiment, each core 1420 is a CPU core.

[0172] In at least one embodiment, each core 1420 includes, without limitation, a fetch / decode unit 1422, an integer execution engine 1424, a floating point execution engine 1426, and an L2 cache 1428. In at least one embodiment, fetch / decode unit 1422 fetches instructions, decodes such instructions, generates micro-operations, and dispatches individual micro-instructions to integer execution engine 1424 and floating point execution engine 1426. In at least one embodiment, fetch / decode unit 1422 can simultaneously dispatch one micro-instruction to integer execution engine 1424 and another micro-instruction to floating point execution engine 1426. In at least one embodiment, integer execution engine 1424 executes not limited to integer and memory operations. In at least one embodiment, floating point engine 1426 executes not limited to floating point and vector operations. In at least one embodiment, fetch-decode unit 1422 dispatches micro-instructions to a single execution engine in place of both integer execution engine 1424 and floating point execution engine 1426.

[0173] In at least one embodiment, each core 1420(i) has access to an L2 cache 1428(i) included in core 1420(i), where i is an integer representing a particular instance of core 1420. In at least one embodiment, each core 1420 included in core complex 1410(j) is connected to other cores 1420 included in core complex 1410(j) via an L3 cache 1430(j) included in core complex 1410(j), where j is an integer representing a particular instance of core complex 1410. In at least one embodiment, cores 1420 included in core complex 1410(j) have access to all L3 caches 1430(j) included in core complex 1410(j), where j is an integer representing a particular instance of core complex 1410. In at least one embodiment, L3 cache 1430 can include, without limitation, any number of slices.

[0174] In at least one embodiment, graphics complex 1440 can be configured to perform compute operations in a highly parallel manner. In at least one embodiment, graphics complex 1440 is configured to perform graphics pipeline operations such as draw commands, pixel operations, geometric computations, and other operations associated with rendering images to a display. In at least one embodiment, graphics complex 1440 is configured to perform operations that are not graphics related. In at least one embodiment, graphics complex 1440 is configured to perform both graphics related operations and operations that are not graphics related.

[0175] In at least one embodiment, the graphics complex 1440 includes, but is not limited to, any number of computing units 1450 and an L2 cache 1442. In at least one embodiment, the computing units 1450 share the L2 cache 1442. In at least one embodiment, the L2 cache 1442 is partitioned. In at least one embodiment, the graphics complex 1440 includes, but is not limited to, any number of computing units 1450 and any number (including zero) and type of cache. In at least one embodiment, the graphics complex 1440 includes, but is not limited to, any number of dedicated graphics hardware.

[0176] In at least one embodiment, each computing unit 1450 includes, but is not limited to, any number of SIMD units 1452 and shared memory 1454. In at least one embodiment, each SIMD unit 1452 implements a SIMD architecture and is configured to execute operations in parallel. In at least one embodiment, each computing unit 1450 may execute any number of thread blocks, but each thread block executes on a single computing unit 1450. In at least one embodiment, a thread block includes, but is not limited to, any number of execution threads. In at least one embodiment, a workgroup is a thread block. In at least one embodiment, each SIMD unit 1452 executes a different warp. In at least one embodiment, a warp is a group of threads (e.g., 16 threads), where each thread in the warp belongs to a single thread block and is configured to process different datasets based on a single instruction set. In at least one embodiment, prediction can be used to disable one or more threads in a warp. In at least one embodiment, a channel is a thread. In at least one embodiment, a work item is a thread. In at least one embodiment, a wavefront is a warp. In at least one embodiment, different wavefronts in a thread block can be synchronized together and communicate via shared memory 1454.

[0177] In at least one embodiment, fabric 1460 is a system interconnect that facilitates data and control transmissions across core complex 1410, graphics complex 1440, I / O interface 1470, memory controllers 1480, display controller 1492, and multimedia engine 1494. In at least one embodiment, APU 1400 can include, without limitation, any number and type of system interconnects in addition to or instead of fabric 1460 that facilitate data and control transmissions across any number and type of directly or indirectly linked components that can be internal or external to APU 1400. In at least one embodiment, I / O interface 1470 represents any number and type of I / O interface (e.g., PCI, PCI-Extended (“PCI-X”), PCIe, Gigabit Ethernet (“GBE”), USB, etc.). In at least one embodiment, various types of peripheral devices are coupled to I / O interface 1470. In at least one embodiment, peripheral devices coupled to I / O interface 1470 can include, without limitation, a keyboard, a mouse, a printer, a scanner, a joystick or other type of game controller, a media recording device, an external storage device, a network interface card, etc.

[0178] In at least one embodiment, display controller 1492 displays images on one or more display devices, such as liquid crystal display (“LCD”) devices. In at least one embodiment, multimedia engine 1494 includes, without limitation, any number and type of multimedia-related circuitry, such as a video decoder, a video encoder, an image signal processor, etc. In at least one embodiment, memory controllers 1480 facilitate data transfers between APU 1400 and unified system memory 1490. In at least one embodiment, core complex 1410 and graphics complex 1440 share unified system memory 1490.

[0179] In at least one embodiment, APU 1400 implements a memory subsystem that includes, without limitation, any number and type of memory controllers 1480 and memory devices (e.g., shared memory 1454) that can be dedicated to one component or shared among multiple components. In at least one embodiment, APU 1400 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 1528, L3 cache 1430, and L2 cache 1442), each of which can be private to a component or shared among any number of components (e.g., core 1420, core complex 1410, SIMD unit 1452, compute unit 1450, and graphics complex 1440).

[0180] In at least one embodiment, system 100 (see Figure 1 andFigure 3 ) can be used to implement the APU 1400. In at least one embodiment, the set 112 and / or the set 114 can include one or more of one or more components of the core complex 1410 and / or one or more components of the graphics complex 1440. In at least one embodiment, the fabric 1460 can include the switch circuit 110.

[0181] Figure 15 A CPU 1500, in accordance with at least one embodiment, is shown. In at least one embodiment, the CPU 1500 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, the CPU 1500 can be configured to execute application programs. In at least one embodiment, the CPU 1500 is configured to execute host control software, such as an operating system. In at least one embodiment, the CPU 1500 issues commands that control the operation of an external GPU (not shown). In at least one embodiment, the CPU 1500 can be configured to execute host executable code derived from CUDA source code, and the external GPU can be configured to execute device executable code derived from such CUDA source code. In at least one embodiment, the CPU 1500 includes, without limitation, any number of core complexes 1510, a fabric 1560, I / O interfaces 1570, and memory controllers 1580.

[0182] In at least one embodiment, the core complex 1510 includes, without limitation, cores 1520(1)-1520(4) and an L3 cache 1530. In at least one embodiment, the core complex 1510 can include, without limitation, any number of cores 1520 and any combination and type of caches. In at least one embodiment, the cores 1520 are configured to execute instructions of a particular ISA. In at least one embodiment, each core 1520 is a CPU core.

[0183] In at least one embodiment, each core 1520 includes, without limitation, a fetch / decode unit 1522, an integer execution engine 1524, a floating point execution engine 1526, and an L2 cache 1528. In at least one embodiment, fetch / decode unit 1522 fetches instructions, decodes such instructions, generates micro-operations, and dispatches individual micro-instructions to integer execution engine 1524 and floating point execution engine 1526. In at least one embodiment, fetch / decode unit 1522 can simultaneously dispatch one micro-instruction to integer execution engine 1524 and another micro-instruction to floating point execution engine 1526. In at least one embodiment, integer execution engine 1524 executes integer and memory operations, without limitation. In at least one embodiment, floating point engine 1526 executes floating point and vector operations, without limitation. In at least one embodiment, fetch-decode unit 1522 dispatches micro-instructions to a single execution engine in place of both integer execution engine 1524 and floating point execution engine 1526.

[0184] In at least one embodiment, each core 1520(i) has access to an L2 cache 1528(i) included in core 1520(i), where i is an integer representing a particular instance of core 1520. In at least one embodiment, each core 1520 included in core complex 1510(j) is connected to the other cores 1520 in core complex 1510(j) via an L3 cache 1530(j) included in core complex 1510(j), where j is an integer representing a particular instance of core complex 1510. In at least one embodiment, cores 1520 included in core complex 1510(j) have access to all L3 caches 1530(j) included in core complex 1510(j), where j is an integer representing a particular instance of core complex 1510. In at least one embodiment, L3 cache 1530 can include, without limitation, any number of slices.

[0185] In at least one embodiment, fabric 1560 is a system interconnect that facilitates data and control transfers across core complexes 1510(1)-1510(N) (where N is an integer greater than zero), I / O interface 1570, and memory controllers 1580. In at least one embodiment, CPU 1500 can include, without limitation, any number and type of system interconnects in addition to or instead of fabric 1560 that facilitate data and control transfers across any number and type of directly or indirectly linked components that can be internal or external to CPU 1500. In at least one embodiment, I / O interface 1570 represents any number and type of I / O interface (e.g., PCI, PCI-X, PCIe, GBE, USB, etc.). In at least one embodiment, various types of peripheral devices are coupled to I / O interface 1570. In at least one embodiment, peripheral devices coupled to I / O interface 1570 can include, without limitation, a display, a keyboard, a mouse, a printer, a scanner, a joystick or other type of game controller, a media recording device, an external storage device, a network interface card, etc.

[0186] In at least one embodiment, memory controllers 1580 facilitate data transfers between CPU 1500 and system memory 1590. In at least one embodiment, core complexes 1510 and graphics complex 1540 share system memory 1590. In at least one embodiment, CPU 1500 implements a memory subsystem that includes, without limitation, any number and type of memory controllers 1580 and memory devices that can be dedicated to one component or shared among multiple components. In at least one embodiment, CPU 1500 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 1528 and L3 cache 1530), each of which can be private to a component or shared among any number of components (e.g., cores 1520 and core complexes 1510).

[0187] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement CPU 1500. In at least one embodiment, sets 112 and / or groups 114 can include one or more of core complexes 1510. In at least one embodiment, fabric 1560 can include switch circuitry 110.

[0188] Figure 16An exemplary accelerator integration slice 1690 is shown in accordance with at least one embodiment. As used herein, a “slice” includes a specified portion of processing resources of an accelerator integration circuit. In at least one embodiment, an accelerator integration circuit provides cache management, memory access, environment management, and interrupt management services on behalf of a number of graphics processing engines that are part of a graphics acceleration module. Graphics processing engines can each comprise a separate GPU. Alternatively, graphics processing engines can include different types of graphics processing engines within a GPU, such as graphics execution units, media processing engines (e.g., video encoders / decoders), samplers, and blit engines. In at least one embodiment, a graphics acceleration module can be a GPU with a plurality of graphics processing engines. In at least one embodiment, a graphics processing engine can be a separate GPU integrated on a common package, line card, or chip as the application processor 1607.

[0189] An application effective address space 1682 within system memory 1614 stores process elements 1683. In one embodiment, process elements 1683 are stored in response to GPU invocations 1681 from an application 1680 executing on processor 1607. Process elements 1683 contain processing state for a corresponding application 1680. A work descriptor (WD) 1684 contained in process element 1683 can be a single job requested by an application or can contain pointers to a queue of jobs. In at least one embodiment, WD 1684 is a pointer to a job request queue in application effective address space 1682.

[0190] Graphics acceleration module 1646 and / or individual graphics processing engines can be shared by all or a subset of processes in a system. In at least one embodiment, a process can include infrastructure for setting up processing state and sending WDs 1684 to graphics acceleration module 1646 to start jobs in a virtualized environment.

[0191] In at least one embodiment, a dedicated process programming model is implemented for owned. In this model, a single process owns a graphics acceleration module 1646 or individual graphics processing engines. As graphics acceleration module 1646 is owned by a single process, a hypervisor initializes the accelerator integration circuit for the owning partition and an operating system initializes the accelerator integration circuit for the owning partition when assigning graphics acceleration module 1646.

[0192] In operation, a WD fetch unit 1691 in accelerator integration slice 1690 fetches the next WD 1684, which includes an indication of work to be completed by one or more graphics processing engines of graphics acceleration module 1646. Data from WD 1684 can be stored in registers 1645 used by memory management unit (MMU) 1639, interrupt management circuit 1647, and / or environment management circuit 1648, as shown. For example, one embodiment of MMU 1639 includes segment / page walk circuitry to access segment / page tables 1686 within an OS virtual address space 1685. Interrupt management circuit 1647 can handle interrupt events (INT) 1692 received from graphics acceleration module 1646. Effective addresses 1693 produced by the graphics processing engines, when executing graphics operations, are translated to real addresses by MMU 1639.

[0193] In one embodiment, the same set of registers 1645 is replicated for each graphics processing engine and / or graphics acceleration module 1646 and can be initialized by a hypervisor or operating system. Each of these replicated registers can be included in accelerator integration slice 1690. Exemplary registers that can be initialized by a hypervisor are shown in Table 1.

[0194] Table 1 - Hypervisor Initialized Registers

[0195]

[0196]

[0197] Exemplary registers that can be initialized by an operating system are shown in Table 2.

[0198] Table 2 - Operating System Initialized Registers

[0199] 1 Process and thread identification 2 Effective address (EA) environment save / restoration pointer 3 Virtual address (VA) accelerator utilization record pointer 4 Virtual address (VA) storage segment table pointer 5 Authority mask 6 Work descriptor

[0200] In one embodiment, each WD 1684 is specific to a particular graphics acceleration module 1646 and / or a particular graphics processing engine. It contains all information needed for the graphics processing engines to do the work or it can be a pointer to a memory location where an application has set up a command queue of work to be completed.

[0201] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement accelerator integration slice 1690. In at least one embodiment, set 112 and / or group 114 can include processor 1607, graphics acceleration module 1646, and / or individual graphics processing engines.

[0202] Figures 17A-17BExemplary graphics processors in accordance with at least one embodiment are shown. In at least one embodiment, any of the exemplary graphics processors can be fabricated using one or more IP cores. In addition to the illustrations, other logic and circuitry can be included in at least one embodiment, including additional graphics processors / cores, peripheral interface controllers or general purpose processor cores. In at least one embodiment, exemplary graphics processors are used within SoCs.

[0203] Figure 17A An exemplary graphics processor 1710 of an SoC integrated circuit in accordance with at least one embodiment is shown, which can be fabricated using one or more IP cores. Figure 17B An additional exemplary graphics processor 1740 of an SoC integrated circuit in accordance with at least one embodiment is shown, which can be fabricated using one or more IP cores. In at least one embodiment, Figure 17A Graphics processor 1710 is a low power graphics processor core. In at least one embodiment, Figure 17B Graphics processor 1740 is a higher performance graphics processor core. In at least one embodiment, each graphics processor 1710, 1740 can be Figure 12 Variations of graphics processor 1210.

[0204] In at least one embodiment, graphics processor 1710 includes a vertex processor 1705 and one or more fragment processor(s) 1715A-1715N (e.g., 1715A, 1715B, 1715C, 1715D, through 1715N-1, and 1715N). In at least one embodiment, graphics processor 1710 can execute different shader programs via separate logical

[0205] In at least one embodiment, the graphics processor 1710 additionally includes one or more MMUs 1720A-1720B, caches 1725A-1725B, and circuit interconnects 1730A-1730B. In at least one embodiment, one or more MMUs 1720A-1720B provide a virtual-to-physical address mapping for the graphics processor 1710, including for vertex processors 1705 and / or fragment processors 1715A-1715N, which can reference vertex or image / texture data stored in memory, in addition to vertex or image / texture data stored in one or more caches 1725A-1725B. In at least one embodiment, one or more MMUs 1720A-1720B can be synchronized with other MMUs within the system, including with... Figure 12 One or more application processors 1205, image processors 1215, and / or video processors 1220 are associated with one or more MMUs, enabling each processor 1205-1220 to participate in a shared or unified virtual memory system. In at least one embodiment, one or more circuit interconnects 1730A-1730B enable the graphics processor 1710 to connect to other IP cores within the SoC via the SoC's internal bus or via a direct connection.

[0206] In at least one embodiment, the graphics processor 1740 includes Figure 17A The graphics processor 1710 includes one or more MMUs 1720A-1720B, caches 1725A-1725B, and circuit interconnects 1730A-1730B. In at least one embodiment, the graphics processor 1740 includes one or more shader cores 1755A-1755N (e.g., 1755A, 1755B, 1755C, 1755D, 1755E, 1755F, to 1755N-1 and 1755N) that provide a unified shader core architecture, wherein a single core or type of core can execute all types of programmable shader code, including shader program code for implementing vertex shaders, fragment shaders, and / or compute shaders. In at least one embodiment, the number of shader cores may vary. In at least one embodiment, the graphics processor 1740 includes an inter-core task manager 1745 that acts as a thread dispatcher to assign execution threads to one or more shader cores 1755A-1755N and a tile unit 1758 to accelerate tile-based rendering operations, wherein rendering operations of a scene are subdivided in image space, for example, to take advantage of local spatial consistency within the scene or to optimize the use of internal caches.

[0207] In at least one embodiment, set 112 (see Figure 1 and Figure 3) and / or group 114 (see Figure 1 and Figure 3 The set 112 and / or group 114 may include one or more of the graphics core 1700 and / or one or more of the graphics processor 1740. In at least one embodiment, the set 112 and / or group 114 may include one or more of the vertex processor 1705 and / or the fragment processors 1715A-1715N. In at least one embodiment, the set 112 and / or group 114 may include one or more of the inter-kernel task manager 1745 and / or the shader cores 1755A-1755N.

[0208] Figure 18A A graphics core 1800 according to at least one embodiment is shown. In at least one embodiment, the graphics core 1800 may include... Figure 12 The graphics processor 1210 is located within it. In at least one embodiment, the graphics core 1800 may be... Figure 17B The graphics core 1800 uses a unified shader core 1755A-1755N. In at least one embodiment, the graphics core 1800 includes a shared instruction cache 1802, texture units 1818, and cache / shared memory 1820, which are common to execution resources within the graphics core 1800. In at least one embodiment, the graphics core 1800 may include multiple slices 1801A-1801N or partitions of each core, and the graphics processor may include multiple instances of the graphics core 1800. Slices 1801A-1801N may include supporting logic, including local instruction caches 1804A-1804N, thread schedulers 1806A-1806N, thread dispatchers 1808A-1808N, and a set of registers 1810A-1810N. In at least one embodiment, slices 1801A-1801N may include a set of additional function units (AFU) 1812A-1812N, floating-point units (FPU) 1814A-1814N, integer arithmetic logic units (ALU) 1816A-1816N, address calculation units (ACU) 1813A-1813N, double-precision floating-point units (DPFPU) 1815A-1815N, and matrix processing units (MPU) 1817A-1817N.

[0209] In one embodiment, FPUs 1814A-1814N can perform single-precision (32-bit) and half-precision (16-bit) floating point operations, while DPFPUs 1815A-1815N can perform double-precision (64-bit) floating point operations. In at least one embodiment, ALUs 1816A-1816N can perform variable precision integer operations at 8-bit, 16-bit, and 32-bit precision, and can be configured for mixed precision operations. In at least one embodiment, MPUs 1817A-1817N can also be configured for mixed precision matrix operations, including half-precision floating point operations and 8-bit integer operations. In at least one embodiment, MPUs 1817A-1817N can perform a variety of matrix operations to accelerate CUDA programs, including enabling support for accelerated General Matrix to Matrix multiplication (GEMM). In at least one embodiment, AFUs 1812A-1812N can perform additional logical operations not supported by floating point or integer units, including trigonometric operations (e.g., Sine, Cosine, etc.).

[0210] Figure 18B A general purpose graphics processing unit (GPGPU) 1830 is shown in at least one embodiment. In at least one embodiment, GPGPU 1830 is highly parallel and is suitable for deployment on a multi-chip module. In at least one embodiment, GPGPU 1830 can be configured to enable highly parallel compute operations to be performed by a GPU array. In at least one embodiment, GPGPU 1830 can be directly linked to other instances of GPGPU 1830 to create a multi-GPU cluster to improve execution time for CUDA programs. In at least one embodiment, GPGPU 1830 includes a host interface 1832 to enable connections with host processors. In at least one embodiment, host interface 1832 is a PCIe interface. In at least one embodiment, host interface 1832 can be a vendor-specific communications interface or communications fabric. In at least one embodiment, GPGPU 1830 receives commands from a host processor to perform operations associated with those commands using a global scheduler 1834 to dispatch execution threads associated with those commands to a group of compute clusters 1836A-1836H. In at least one embodiment, compute clusters 1836A-1836H share a cache memory 1838. In at least one embodiment, cache memory 1838 can be used as an upper level cache for cache memory within compute clusters 1836A-1836H.

[0211] In at least one embodiment, GPGPU 1830 includes memory 1844A-1844B coupled to compute clusters 1836A-1836H via a set of memory controllers 1842A-1842B. In at least one embodiment, memory 1844A-1844B can include various types of memory devices including dynamic random access memory (DRAM) or graphics random access memory, such as synchronous graphics random access memory (SGRAM), including graphics double data rate (GDDR) memory.

[0212] In at least one embodiment, compute clusters 1836A-1836H each include a group of graphics cores, such as graphics core 1800, which can include multiple types of integer and floating point logic units that can perform computational operations at various precisions, including suitable for computations related to CUDA programs. For example, in at least one embodiment, at least a subset of floating point units in each compute cluster 1836A-1836H can be configured to perform 16- or 32-bit floating point operations, while a different subset of floating point units can be configured to perform 64-bit floating point operations. Figure 18A

[0213] In at least one embodiment, multiple instances of GPGPU 1830 can be configured to operate as compute clusters. Compute clusters 1836A-1836H can implement any technically feasible communication technology for synchronization and data exchange. In at least one embodiment, multiple instances of GPGPU 1830 communicate over host interface 1832. In at least one embodiment, GPGPU 1830 includes I / O hub 1839 that couples GPGPU 1830 with GPU link 1840, enabling a direct connection to other instances of GPGPU 1830. In at least one embodiment, GPU link 1840 is coupled to a specialized GPU-to-GPU bridge that enables communication and synchronization between multiple instances of GPGPU 1830. In at least one embodiment, GPU link 1840 is coupled with a high-speed interconnect to transmit and receive data to other GPGPUs or parallel processors. In at least one embodiment, multiple instances of GPGPU 1830 are located in separate data processing systems and communicate over a network device accessible via host interface 1832. In at least one embodiment, GPU link 1840 can be configured to connect to a host processor, in addition to or in place of host interface 1832. In at least one embodiment, GPGPU 1830 can be configured to execute CUDA programs.

[0214] Figure 19A ​A parallel processor 1900 is shown according to at least one embodiment. In at least one embodiment, various components of parallel processor 1900 can be implemented using one or more integrated circuits, for example, programmable processors, application specific integrated circuits (ASICs), or FPGAs.

[0215] In at least one embodiment, parallel processor 1900 includes a parallel processing unit 1902. In at least one embodiment, parallel processing unit 1902 includes an I / O unit 1904 that enables communication with other devices, including other instances of parallel processing unit 1902. In at least one embodiment, I / O unit 1904 can be directly connected to the other devices. In at least one embodiment, I / O unit 1904 connects with other devices using a hub or switch interface, for example, memory hub 1905. In at least one embodiment, connections between memory hub 1905 and I / O unit 1904 form a communication link. In at least one embodiment, I / O unit 1904 connects with a host interface 1906 and a memory crossbar switch 1916, where host interface 1906 receives commands directed to the processing operations and memory crossbar switch 1916 receives commands directed to the memory operations.

[0216] In at least one embodiment, when host interface 1906 receives a command buffer via I / O unit 1904, host interface 1906 can direct work operations to execute those commands to front end 1908. In at least one embodiment, front end 1908 couples with a scheduler 1910, which is configured to assign commands or other work items to processing array 1912. In at least one embodiment, scheduler 1910 ensures that processing array 1912 is properly configured and in an active state before assigning tasks to processing array 1912 of processing array 1912. In at least one embodiment, scheduler 1910 is implemented by firmware logic executing on a microcontroller. In at least one embodiment, microcontroller- implemented scheduler 1910 is configurable to perform complex scheduling and work distribution operations at a coarse and fine grain level to enable fast preemption and context switching for threads executing on processing array 1912. In at least one embodiment, host software can prove a workload for scheduling on processing array 1912 through one of a number of graphics processing doorbell signals. In at least one embodiment, workload can then be automatically distributed on processing array 1912 by scheduler 1910 logic within microcontroller that includes scheduler 1910.

[0217] In at least one embodiment, processing array 1912 can include up to “N” processing clusters (e.g., cluster 1914A, cluster 1914B, through cluster 1914N). In at least one embodiment, each cluster 1914A-1914N of processing array 1912 can execute a large number of concurrent threads. In at least one embodiment, scheduler 1910 can allocate work to clusters 1914A-1914N of processing array 1912 using various scheduling and / or work distribution algorithms, which can be determined at least in part by workload generated by each program or computational type. In at least one embodiment, scheduling can be handled dynamically by scheduler 1910, or can be aided in part by compiler logic during compilation of program logic configured for execution by processing array 1912. In at least one embodiment, different clusters 1914A-1914N of processing array 1912 can be allocated for processing different types of programs or for performing different types of computations.

[0218] In at least one embodiment, processing array 1912 can be configured to perform various types of parallel processing operations. In at least one embodiment, processing array 1912 is configured to perform general purpose parallel compute operations. For example, in at least one embodiment, processing array 1912 can include logic to perform processing tasks including filtering of video and / or audio data, performing modeling operations including physical operations, and performing data transformations.

[0219] In at least one embodiment, processing array 1912 is configured to perform parallel graphics processing operations. In at least one embodiment, processing array 1912 can include additional logic to support performance of such graphics processing operations, including but not limited to texture sampling logic to perform texture operations, tessellation logic, and other vertex processing logic. In at least one embodiment, processing array 1912 can be configured to execute graphics processing related shader programs, such as, but not limited to, vertex shaders, tessellation shaders, geometry shaders, and pixel shaders. In at least one embodiment, parallel processing unit 1902 can transfer data from system memory for processing via I / O unit 1904. In at least one embodiment, during processing, results can be stored to on-chip memory (e.g., parallel processor memory 1922) for write back to system memory.

[0220] In at least one embodiment, when parallel processing unit 1902 is used to perform graphics processing, scheduler 1910 can be configured to divide incoming workloads into tasks of approximately equal size to better enable distribution of graphics processing operations across multiple clusters 1914A-1914N of processing array 1912. In at least one embodiment, portions of processing array 1912 can be configured to perform different types of processing. For example, in at least one embodiment, a first portion can be configured to

[0221] In at least one embodiment, processing array 1912 can receive processing tasks to be executed from scheduler 1910, which receives commands defining the processing tasks from front end 1908. In at least one embodiment, a processing task can include an index into a data store that is to be processed, which can include surface (patch) data, raw data, vertex data, and / or pixel data, for example, as well as state parameters and commands defining how the data is to be processed (e.g., what program is to be executed). In at least one embodiment, scheduler 1910 can be configured to fetch the index corresponding to a task, or can receive the index from front end 1908. In at least one embodiment, front end 1908 can be configured to ensure that processing array 1912 is configured in an effective state before launching a workload specified by an incoming command buffer (e.g., a batch-buffer, a push buffer, etc.).

[0222] In at least one embodiment, each of one or more instances of parallel processing unit 1902 can be coupled to a parallel processor memory 1922. In at least one embodiment, parallel processor memory 1922 can be accessed by the processing array 1912, as well as the I / O unit 1904, via a memory crossbar 1916. In at least one embodiment, memory crossbar 1916 can be used to transfer data between memory elements and the processing array 1912, I / O unit 1904, and / or other components of the parallel processing unit 1902. In at least one embodiment, a memory controller 1918 can be included in the parallel processing unit 1902 to perform memory access operations. In at least one embodiment, memory controller 1918 can be a dedicated controller, or can be integrated within another component (e.g., an I / O unit 1904).

[0223] In at least one embodiment, memory units 1924A-1924N can include various types of memory devices including dynamic random access memory (DRAM) or graphics random access memory, such as synchronous graphics random access memory (SGRAM), including graphics double data rate (GDDR) memory. In at least one embodiment, memory units 1924A-1924N can also include 3D stacked memory, including but not limited to high bandwidth memory (HBM). In at least one embodiment, rendering targets such as frame buffers or texture maps can be stored across memory units 1924A-1924N, allowing partition units 1920A-1920N to write portions of each rendering target in parallel to effectively use available bandwidth of parallel processor memory 1922. In at least one embodiment, local instances of parallel processor memory 1922 can be excluded from a unified memory design that utilizes system memory in combination with local cache memory.

[0224] In at least one embodiment, any of clusters 1914A-1914N of processing array 1912 can process data that is to be written into any of memory units 1924A-1924N within parallel processor memory 1922. In at least one embodiment, memory crossbar 1916 can be configured to transmit outputs of each cluster 1914A-1914N to any partition unit 1920A-1920N or another cluster 1914A-1914N, which can perform other processing operations on the outputs. In at least one embodiment, each cluster 1914A-1914N can communicate with memory interface 1918 through memory crossbar 1916 to read from or write to various external memory devices. In at least one embodiment, memory crossbar 1916 has a connection to memory interface 1918 to communicate with I / O unit 1904 and to local instances of parallel processor memory 1922 to enable processing elements within different processing clusters 1914A-1914N to communicate with system memory or other memories not local to parallel processing unit 1902. In at least one embodiment, memory crossbar 1916 can use virtual channels to separate traffic streams between clusters 1914A-1914N and partition units 1920A-1920N.

[0225] In at least one embodiment, multiple instances of parallel processing unit 1902 can be provided on a single add-in card, or multiple add-in cards can be interconnected. In at least one embodiment, different instances of parallel processing unit 1902 can be configured to operate in coordination with each other to enable single program multi-processing (SPMP). In at least one embodiment, SPMP enables a single program to utilize the resources of multiple parallel processing units 1902 by running all of the program’s threads on the different parallel processing units 1902. For example, in at least one embodiment, some instances of parallel processing unit 1902 can include a greater number of processing cores relative to other instances, a greater amount of local parallel processor memory, and / or other configuration differences.

[0226] In at least one embodiment, clusters 112 (see Figure 1 and Figure 3 ) and / or groups 114 (see Figure 1 and Figure 3 ) can include one or more of parallel processor 1900.

[0227] Figure 19BA processing cluster 1994 is shown in accordance with at least one embodiment. In at least one embodiment, processing cluster 1994 is included in a parallel processing unit. In at least one embodiment, processing cluster 1994 is an example of one of processing clusters 1914A-1914N. In at least one embodiment, processing cluster 1994 can be configured to execute many threads in parallel, where the term “thread” refers to a particular instance of a particular program executing on a particular set of input data. In at least one embodiment, Single Instruction, Multiple Data (SIMD) instruction issue techniques are used to support parallel execution of a large number of threads, without requiring multiple independent instruction units. In at least one embodiment, Single Instruction, Multiple Thread (SIMT) techniques are used to support parallel execution of a large number of generally synchronous threads, using a common instruction unit configured to issue instructions to a set of processing engines within each processing cluster 1994. Figure 19A

[0228] In at least one embodiment, operation of processing cluster 1994 can be controlled via a pipeline manager 1932 that is assigned to a SIMT parallel processor. In at least one embodiment, pipeline manager 1932 receives instructions from scheduler 1910 of a graphics processing unit, and manages execution of those instructions via a graphics multiprocessor 1934 and / or a texture unit 1936. In at least one embodiment, graphics multiprocessor 1934 is an exemplary instance of a SIMT parallel processor. However, in at least one embodiment, various types of SIMT parallel processors of differing architectures can be included within processing cluster 1994. In at least one embodiment, one or more instances of graphics multiprocessor 1934 can be included within processing cluster 1994. In at least one embodiment, graphics multiprocessor 1934 can process data, and a data crossbar 1940 can be used to distribute processed data to one of a number of possible destinations, including other shader units. In at least one embodiment, pipeline manager 1932 can facilitate distribution by specifying destinations for processed data as a function of the processed data. Figure 19A

[0229] In at least one embodiment, each graphics multiprocessor 1934 within processing cluster 1994 can include an identical set of functional execution logic (e.g., arithmetic logic units, load store units (LSUs), etc.). In at least one embodiment, functional execution logic can be configured in a pipelined manner, where new instructions can be issued before previous instructions are complete. In at least one embodiment, functional execution logic supports a variety of operations including integer and floating point arithmetic, comparison operations, Boolean operations, shift operations, and the like. In at least one embodiment, same functional-unit hardware can be leveraged to perform ​​

[0230] In at least one embodiment, instructions delivered to processing cluster 1994 form a thread. In at least one embodiment, a set of threads executing across a set of parallel processing engines is a thread group. In at least one embodiment, a thread group executes a program on different input data. In at least one embodiment, each thread within a thread group can be assigned to a different processing engine within graphics multiprocessor 1934. In at least one embodiment, a thread group can include fewer threads than are present in a plurality of processing engines within graphics multiprocessor 1934. In at least one embodiment, when a thread group includes fewer threads than the number of processing engines present in graphics multiprocessor 1934, one or more of the processing engines can be idle during the execution of the threads in the thread group. In at least one embodiment, a thread group can also include more threads than are present in a plurality of processing engines within graphics multiprocessor 1934. In at least one embodiment, when a thread group includes more threads than the number of processing engines present in graphics multiprocessor 1934, multiple threads can be executed concurrently on the processing engines present in graphics multiprocessor 1934. In at least one embodiment, a plurality of thread groups can be executed concurrently by graphics multiprocessor 1934.

[0231] In at least one embodiment, graphics multiprocessor 1934 includes internal cache memory, to perform load and store operations. In at least one embodiment, graphics multiprocessor 1934 can bypass internal cache and use register file memory, e.g., LI cache 1948, within processing cluster 1994. In at least one embodiment, each graphics multiprocessor 1934 can also have access to L2 cache within a partition unit (e.g., partition units 1920A-1920N) that is shared among multiple processing clusters 1994 and can be used to transfer data between threads. Figure 19A In at least one embodiment, graphics multiprocessor 1934 can also have access to off-chip global memory, which can include one or more of local parallel processor memory and / or system memory. In at least one embodiment, any memory outside of parallel processor 1902 can be considered global memory. In at least one embodiment, processing cluster 1994 includes multiple instances of graphics multiprocessor 1934, which can share common instructions and data stored in LI cache 1948.

[0232] In at least one embodiment, each processing cluster 1994 can include an MMU 1945 that is configured to translate virtual addresses into physical addresses, as is known to those skilled in the art. In at least one embodiment, one or more instances of MMU 1945 can reside in Figure 19Awithin memory interface 1918. In at least one embodiment, MMU 1945 includes a set of page table entries (PTEs) for mapping virtual addresses into physical addresses at which tiles (more on tiles below) are stored and optionally into cache line indices. In at least one embodiment, MMU 1945 can include an address translation lookaside buffer (TLB), or cache, which can reside within graphics multiprocessor 1934 or within LI cache 1948 or processing cluster 1994. In at least one embodiment, processing physical addresses to allocate surface data access locality for efficient request interweaving among partition units. In at least one embodiment, cache line indices can be used to determine whether a request for a cache line is a hit or miss.

[0233] In at least one embodiment, processing cluster 1994 can be configured such that each graphics multiprocessor 1934 is coupled to a texture unit 1936 for performing texture mapping operations, e.g., determining texture sample positions, reading texture data, and filtering texture data. In at least one embodiment, texture data is read from an internal texture LI cache (not shown) or from an LI cache within graphics multiprocessor 1934 as needed, and texture data is fetched from an L2 cache, local parallel processor memory, or system memory, as needed. In at least one embodiment, each graphics multiprocessor 1934 outputs processed tasks to data crossbar 1940 to provide processed tasks to another processing cluster 1994 for further processing or to store processed task data in an L2 cache, local parallel processor memory, or system memory via memory crossbar 1916. In at least one embodiment, a raster operations pre-unit (Pre-ROP) 1942 is configured to receive data from graphics multiprocessor 1934, direct data to ROP units located within sub-regions of processing cluster 1994, which ROP units can be the same as or separate from ROP units 1920A-1920N as discussed elsewhere herein. In at least one embodiment, Pre-ROP 1942 can perform optimizations to minimize or eliminate certain exceptions that can invalidate a processor pipeline. In at least one embodiment, Pre-ROP 1942 instructs a memory crossbar to send primitive information to an appropriate ROP unit located within processing cluster 1994. Figure 19A

[0234] Figure 19C A graphics multiprocessor 1996, according to at least one embodiment, is shown. In at least one embodiment, graphics multiprocessor 1996 is a GPC as described herein, which can be one member of a group of GPCs including a number of graphics multiprocessors 1996. Figure 19B ​graphics processor 1934. In at least one embodiment, graphics processor 1996 is coupled with pipeline manager 1932 of processing cluster 1994. In at least one embodiment, graphics processor 1996 has a thread execution pipeline that includes, without limitation, an instruction cache 1952, an instruction unit 1954, an address mapping unit 1956, a register file 1958, one or more GPGPU cores 1962, and one or more LSU’s 1966. GPGPU cores 1962 and LSUs 1966 are coupled with cache memory 1972 and shared memory 1970 via memory and cache interconnect 1968.

[0235] In at least one embodiment, instruction cache 1952 receives a stream of instructions 1950 to execute from pipeline manager 1932. In at least one embodiment, instructions are cached in instruction cache 1952 and dispatched for execution by instruction unit 1954. In one embodiment, instruction unit 1954 can dispatch instructions to the threads of a thread group, with each thread in the thread group assigned to a different execution lane within GPGPU core 1962. In at least one embodiment, instructions can access any of the local, shared, or global address spaces by specifying an address in a unified address space. In at least one embodiment, address mapping unit 1956 can be used to convert an address in the unified address space into an address that can be accessed by LSU 1966.

[0236] In at least one embodiment, register file 1958 provides a set of registers for functional units of graphics processor 1996. In at least one embodiment, register file 1958 provides temporary storage for operands of the data

[0237] In at least one embodiment, GPGPU cores 1962 can each include FPUs and / or ALUs for executing instructions for graphics processing. GPGPU cores 1962 can be similar to each other in architecture or can differ from each other in architecture. In at least one embodiment, a first portion of GPGPU cores 1962 includes single precision FPUs and integer ALUs, while a second portion of GPGPU cores 1962 includes double precision FPUs. In at least one embodiment, FPUs can implement IEEE 754-2008 standard for floating point arithmetic or enable variable precision floating point arithmetic. In at least one embodiment, graphics processor 1900 can additionally include one or more fixed function or special purpose logic units to perform specific functions such as copy rectangle or pixel blending operations. In at least one embodiment one or more of GPGPU cores 1962 can also include fixed or special function logic.

[0238] In at least one embodiment, GPGPU cores 1962 include SIMD logic capable of

[0239] In at least one embodiment, memory and cache interconnect 1968 is an interconnect network that connects each functional unit of graphics multiprocessor 1996 to register file 1958 and shared memory 1970. In at least one embodiment, memory and cache interconnect 1968 is a crossbar interconnect that allows LSUs 1966 to enable load and store operations between shared memory 1970 and register file 1958. In at least one embodiment, register file 1958 can operate at same frequency as GPGPU cores 1962, making the transfer of data between GPGPU cores 1962 and register file 1958 very low latency. In at least one embodiment, shared memory 1970 can be used to enable communication between threads executing on functional units within graphics multiprocessor 1996. In at least one embodiment, cache memory 1972 can be used to cache data stored in shared memory 1970, such as texture data, for example, that is communicated between graphics processing engines 1936. In at least one embodiment, shared memory 1970 can also be used as program managed cache. In at least one embodiment, in addition to automatically cached data that is stored in cache memory 1972, threads executing on GPGPU cores 1962 can store data in shared memory through programmatic code.

[0240] In at least one embodiment, parallel processor or GPGPU as described herein is communicatively coupled to host / processor cores to accelerate graphics operations, machine learning operations, pattern analysis operations, and various general purpose GPU (GPGPU) functions. In at least one embodiment, GPU can be communicatively coupled to host processor / cores over a bus or other interconnect (e.g., a high-speed

[0241] Figure 20A graphics processor 2000 according to at least one embodiment is shown. In at least one embodiment, graphics processor 2000 includes ring interconnect 2002, front-end 2004, media engine 2037, and graphics cores 2080A-2080N. In at least one embodiment, ring interconnect 2002 couples graphics processor 2000 to other processing units including other graphics processors or one or more general-purpose processor cores. In at least one embodiment, graphics processor 2000 is one of many processors integrated within a multi-core processing system.

[0242] In at least one embodiment, graphics processor 2000 receives batches of commands via ring interconnect 2002. In at least one embodiment, incoming commands are interpreted by command streamer 2003 in front-end 2004. In at least one embodiment, graphics processor 2000 includes scalable execution logic to perform 3D geometry processing and media processing via graphics cores 2080A-2080N. In at least one embodiment, for 3D geometry processing commands, command streamer 2003 supplies commands to geometry pipeline 2036. In at least one embodiment, for at least some media processing commands, command streamer 2003 supplies commands to video front end 2034, which couples with media engine 2037. In at least one embodiment, media engine 2037 includes a video quality engine (VQE) 2030 for video and image post-processing, and a multi-format encode / decode (MFX) 2033 engine to

[0243] In at least one embodiment, graphics processor 2000 includes a scalable thread execution resource featuring a modular graphics core 2080A-2080N (sometimes referred to as a core slice), each having a plurality of sub-cores 2050A-2050N, 2060A-2060N (sometimes referred to as a core sub-slice). In at least one embodiment, graphics processor 2000 can have any number of graphics cores 2080A-2080N. In at least one embodiment, graphics processor 2000 includes graphics core 2080A having at least first and second sub-cores 2050A, 2060A. In at least one embodiment, graphics processor 2000 is a low power processor with a single sub-core (e.g., 2050A). In at least one embodiment, graphics processor 2000 includes a plurality of graphics cores 2080A-2080N each including a set of first sub-cores 2050A-2050N and a set of second sub-cores 2060A-2060N. In at least one embodiment, each of the first sub-cores 2050A-2050N includes at least a first set of execution units (EUs) 2052A-2052N and a media / texture sampling unit 2054A-2054N. In at least one embodiment, each of the second sub-cores 2060A-2060N includes at least a second set of execution units 2062A-2062N and a sampler 2064A-2064N. In at least one embodiment, each of the sub-cores 2050A-2050N, 2060A-2060N share a set of shared resources 2070A-2070N. In at least one embodiment, shared resources include shared cache memory and pixel operation logic.

[0244] In at least one embodiment, set 112 (see Figure 1 and Figure 3 ) and / or group 114 (see Figure 1 and Figure 3 ) can include one or more of graphics processor 2000.

[0245] Figure 21A processor 2100 according to at least one embodiment is shown. In at least one embodiment, processor 2100 can include, without limitation, a logic circuit that executes instructions. In at least one embodiment, processor 2100 can execute instructions including x86 instructions, ARM instructions, special purpose instructions for ASICs, etc. In at least one embodiment, processor 2110 can include registers to store packed data, for example, 64-bit wide MMX™ registers in microprocessors enabled with MMX technology by Intel Corporation of Santa Clara, California. In at least one embodiment, MMX registers available in integer and floating point form can operate with packed data elements that accompany SIMD and Streaming SIMD Extensions (“SSE”) instructions. In at least one embodiment, 128-bit wide XMM registers related to SSE2, SSE3, SSE4, AVX or higher (generically referred to as “SSEx”) technology can hold such packed data operands. In at least one embodiment, processor 2110 can execute instructions to accelerate CUDA programs.

[0246] In at least one embodiment, processor 2100 includes an in-order front-end (“front-end”) 2101 to fetch instructions to be executed and to prepare instructions for execution by the processor pipeline. In at least one embodiment, front-end 2101 can include several units. In at least one embodiment, an instruction prefetcher 2126 fetches instructions from memory and provides instructions to an instruction decoder 2128 that, in turn, decodes or interprets instructions. In at least one embodiment, instruction decoder 2128 decodes a received instruction into one or more operations called “micro-instructions” or “micro-operations” (also called “uops” or “microinstructions”) for execution. In at least one embodiment, instruction decoder 2128 parses instructions into operation codes and corresponding data and control fields which can be used by micro-architecture to perform operations using registers and flags. In at least one embodiment, a trace cache 2130 can assemble decoded uops into program ordered sequences or traces in a uop queue 2134 for execution. In at least one embodiment, when trace cache 2130 encounters a complex instruction, a microcode ROM 2132 provides uops needed to complete the operation.

[0247] In at least one embodiment, some instructions can be converted into a single micro- operation, while others can require several micro-operations to complete. In at least one embodiment, if more than four micro-instructions are needed to complete a single instruction, then instruction decoder 2128 can access microcode ROM 2132 to execute the instruction. In at least one embodiment, instructions can be decoded into a small number of micro-instructions to be processed at instruction decoder 2128. In at least one embodiment, if multiple micro-instructions are needed to complete an operation, then the instructions can be stored in microcode ROM 2132. In at least one embodiment, a trace cache 2130 references an entry point programmable logic array (“PLA”) to determine a correct micro-instruction pointer for reading a microcode sequence from microcode ROM 2132 to complete one or more instructions, in accordance with at least one embodiment. In at least one embodiment, after microcode ROM 2132 completes sequencing of micro-operations for an instruction, a front end 2101 of a machine can resume fetching micro-operations from trace cache 2130.

[0248] In at least one embodiment, out-of-order execution engine (“out-of-order engine”) 2103 can prepare instructions for execution. In at least one embodiment, out-of-order execution logic has multiple buffers to smooth and reorder instruction flow to optimize performance as instructions are pipelined down and dispatched for execution. Out-of-order execution engine 2103 includes, without limitation, an allocator / register renamer 2140, a memory micro instruction queue 2142, an integer / floating point micro instruction queue 2144, a memory scheduler 2146, a fast scheduler 2102, a slow / general floating point scheduler (“slow / general FP scheduler”) 2104, and a simple floating point scheduler (“simple FP scheduler”) 2106. In at least one embodiment, fast scheduler 2102, slow / general floating point scheduler 2104, and simple floating point scheduler 2106 are also collectively referred to as “micro instruction schedulers 2102, 2104, 2106.” Allocator / register renamer 2140 allocates machine buffers and resources needed for each micro instruction to execute in order. In at least one embodiment, allocator / register renamer 2140 renames logical registers to entries in a register file. In at least one embodiment, allocator / register renamer 2140 also allocates entries for each micro instruction in one of two micro instruction queues, memory micro instruction queue 2142 for memory operations and integer / floating point micro instruction queue 2144 for non-memory operations, in front of memory scheduler 2146 and micro instruction schedulers 2102, 2104, 2106. In at least one embodiment, micro instruction schedulers 2102, 2104, 2106 determine when micro instructions are ready to execute based on readiness of their dependent input register operand sources and availability of execution resource micro instructions needed to complete. In at least one embodiment, fast scheduler 2102 of at least one embodiment can schedule on every half of a main clock cycle, while slow / general floating point scheduler 2104 and simple floating point scheduler 2106 can schedule once per main processor clock cycle. In at least one embodiment, micro instruction schedulers 2102, 2104, 2106 arbitrate for a dispatch port to dispatch micro instructions for execution.

[0249] In at least one embodiment, execution block 2111 includes, without limitation, an integer register file / bypass network 2108, a floating point register file / bypass network (“FP register file / bypass network”) 2110, address generation units (“AGUs”) 2112 and 2114, fast ALUs 2116 and 2118, a slow ALU 2120, a floating point ALU (“FP”) 2122, and a floating point move unit (“FP move”) 2124. In at least one embodiment, integer register file / bypass network 2108 and floating point register file / bypass network 2110 are also referred to herein as “register files 2108, 2110.” In at least one embodiment, AGUs 2112 and 2114, fast ALUs 2116 and 2118, slow ALU 2120, floating point ALU 2122, and floating point move unit 2124 are also referred to herein as “execution units 2112, 2114, 2116, 2118, 2120, 2122, and 2124.” In at least one embodiment, execution block can include, without limitation, any number (including zero) and type of register files, bypass networks, address generation units, and execution units (in any combination).

[0250] In at least one embodiment, register files 2108, 2110 can be arranged between microinstruction schedulers 2102, 2104, 2106 and execution units 2112, 2114, 2116, 2118, 2120, 2122, and 2124. In at least one embodiment, integer register file / bypass network 2108 performs integer operations. In at least one embodiment, floating point register file / bypass network 2110 performs floating point operations. In at least one embodiment, each of register files 2108, 2110 can include, without limitation, a bypass network that can bypass or forward a just-completed result that has not yet been written into a register file to a new dependee. In at least one embodiment, register files 2108, 2110 can communicate data with each other. In at least one embodiment, integer register file / bypass network 2108 can include, without limitation, two separate register files, one for lower 32 bits of data and a second for upper 32 bits of data. In at least one embodiment, floating point register file / bypass network 2110 can include, without limitation, 128 bit wide entries because floating point instructions typically have operands that are 64 to 128 bits wide.

[0251] In at least one embodiment, execution units 2112, 2114, 2116, 2118, 2120, 2122, 2124 can execute instructions. In at least one embodiment, register files 2108, 2110 store integer and floating point data operand values upon which microinstructions require execution. In at least one embodiment, processor 2100 can include, without limitation, any number of execution units 2112, 2114, 2116, 2118, 2120, 2122, 2124 and combinations thereof. In at least one embodiment, floating point ALU 2122 and floating point move unit 2124 can execute floating point, MMX, SIMD, AVX and SSE, or other operations, including specialized machine learning instructions. In at least one embodiment, floating point ALU 2122 can include, without limitation, a 64 bit by 64 bit floating point divider to execute divide, square root, and remainder micro-ops. In at least one embodiment, instructions for floating point

[0252] In at least one embodiment, micro-instruction schedulers 2102, 2104, 2106 schedule dependent operations prior to completion of parent load execution. In at least one embodiment, because micro-instructions can be speculatively scheduled and executed in processor 2100, processor 2100 can also include logic to handle memory misses. In at least one embodiment, if a data cache data load miss occurs, there can be dependent operations running in the pipeline that cause the scheduler to temporarily have incorrect data. In at least one embodiment, a replay mechanism tracks and re-executes instructions that use incorrect data. In at least one embodiment, dependent operations can need to be replayed and independent operations can be allowed to complete. In at least one embodiment, a scheduler and replay mechanism of at least one embodiment of a processor can also be designed to capture instruction sequences for text string compare operations.

[0253] In at least one embodiment, the term “register” can refer to an on-board processor storage location that can be used as part of an instruction that identifies an operand. In at least one embodiment, a register can be one that can be used from outside of a processor (from a programmer’s perspective). In at least one embodiment, a register can not be limited to a particular type of circuit. Rather, in at least one embodiment, a register can store data, provide data, and perform functions described herein. In at least one embodiment, registers described herein can be implemented by circuitry within a processor using a variety of different techniques, such as dedicated physical registers, physical registers allocated dynamically with register renaming, a combination of dedicated and dynamically allocated physical registers, etc. In at least one embodiment, an integer register stores 32-bit integer data. A register file of at least one embodiment also contains eight multimedia SIMD registers for packing data.

[0254] In at least one embodiment, set 112 (see Figure 1 and Figure 3 ) and / or group 114 (see Figure 1 and Figure 3 ) can include one or more of processor 2100.

[0255] Figure 22A processor 2200 according to at least one embodiment is shown. In at least one embodiment, processor 2200 includes, without limitation, one or more processor cores (cores) 2202A-2202N, an integrated memory controller 2214, and an integrated graphics processor 2208. In at least one embodiment, processor 2200 can include additional cores up to and including the additional processor cores 2202N represented by the dashed line. In at least one embodiment, each processor core 2202A-2202N includes one or more internal cache units 2204A-2204N. In at least one embodiment, each processor core can also access one or more shared cache units 2206.

[0256] In at least one embodiment, internal cache units 2204A-2204N and shared cache units 2206 represent a cache memory hierarchy within processor 2200. In at least one embodiment, cache memory units 2204A-2204N can include one or more levels of cache, such as L2, L3, 4-way (L4), or other levels of cache, within each processor core and a shared mid-level cache, for example, with the highest level of cache being classified as an LLC before external memory. In at least one embodiment, cache coherence logic maintains coherence between various cache units 2206 and 2204A-2204N.

[0257] In at least one embodiment, processor 2200 can also include a set of one or more bus controller units 2216 and a system agent core 2210. In at least one embodiment, one or more bus controller units 2216 manage a set of peripheral buses, such as one or more PCI or PCI Express buses. In at least one embodiment, system agent core 2210 provides management functionality for various processor components. In at least one embodiment, system agent core 2210 includes one or more integrated memory controllers 2214 to manage access to various external memory devices (not shown), such as one or more dynamic random access memory DRAM or static RAM (SRAM) devices.

[0258] In at least one embodiment, one or more processor cores 2202A-2202N include support for simultaneous multithreading. In at least one embodiment, system agent core 2210 includes components for coordination and operation of processor cores 2202A-2202N during multithreading. In at least one embodiment, system agent core 2210 can additionally include a power control unit (PCU), including logic and components to regulate one or more power states of processor cores 2202A-2202N and graphics processor 2208.

[0259] In at least one embodiment, processor 2200 additionally includes a graphics processor 2208 to perform graphics processing operations. In at least one embodiment, graphics processor 2208 couples with shared cache unit 2206 and system agent core 2210 including one or more integrated memory controllers 2214. In at least one embodiment, system agent core 2210 also includes a display controller 2211 to drive one or more coupled displays to display graphics processor output. In at least one embodiment, display controller 2211 can also be a separate module coupled with graphics processor 2208 via at least one interconnect, or can be integrated within graphics processor 2208.

[0260] In at least one embodiment, ring based interconnect unit 2212 is used to couple internal components of processor 2200. In at least one embodiment, alternative interconnect units can be used, such as a point-to-point interconnect, a switched interconnect, or other technology. In at least one embodiment, graphics processor 2208 couples with ring interconnect 2212 via I / O link 2213.

[0261] In at least one embodiment, I / O link 2213 represents at least one of a variety of I / O interconnects, including a package I / O interconnect facilitating communication between various processor components and a high performance embedded memory module 2218, such as an eDRAM module. In at least one embodiment, each of processor cores 2202A-2202N and graphics processor 2208 uses embedded memory module 2218 as a shared LLC.

[0262] In at least one embodiment, processor cores 2202A-2202N are homogenous cores executing a common instruction set architecture. In at least one embodiment, processor cores 2202A-2202N are heterogeneous in terms of ISA, with one or more processor cores 2202A-2202N executing a common instruction set, while one or more other processor cores 2202A-2202N execute a subset of the common instruction set or a different instruction set. In at least one embodiment, processor cores 2202A-2202N are heterogeneous in terms of microarchitecture, with one or more cores having a relatively high power consumption coupled with one or more power cores having a lower power consumption. In at least one embodiment, processor 2200 can be implemented on one or more chips or as a SoC integrated circuit.

[0263] In at least one embodiment, set 112 (see Figure 1 and Figure 3 ) and / or set 114 (see Figure 1 and Figure 3) can include one or more of the processor(s) 2200. In at least one embodiment, cluster 112 (see Figure 1 and Figure 3 ) and / or group 114 (see Figure 1 and Figure 3 ) can include one or more of integrated graphics processor 2208 and / or processor core(s) 2202A-2202N. In at least one embodiment, ring interconnect 2212 can include switch circuitry 110.

[0264] Figure 23 A graphics processor core 2300 according to at least one embodiment described is shown. In at least one embodiment, graphics processor core 2300 is included within a graphics core array. In at least one embodiment, graphics processor core 2300 (sometimes called a core slice) can be one or more graphics cores within a modular graphics processor. In at least one embodiment, graphics processor core 2300 is an example of a graphics core slice, and a graphics processor described herein can include multiple graphics core slices based on target power and performance envelopes. In at least one embodiment, each graphics core 2300 can include fixed function and programmable processing logic in a number of sub-cores 2301A-2301F coupled with a fixed function block 2330, also known as a sub-slice.

[0265] In at least one embodiment, fixed function block 2330 includes a geometry / fixed function pipeline 2336, for example, that can be shared by all of the sub-cores in graphics processor 2300 in lower performance and / or lower power graphics processor implementations. In at least one embodiment, geometry / fixed function pipeline 2336 includes a 3D fixed function pipeline, a video front-end unit, a thread generator and thread dispatcher, and a unified return buffer manager that manages a unified return buffer.

[0266] In at least one embodiment, fixed function block 2330 also includes a graphics SoC interface 2337, a graphics microcontroller 2338, and a media pipeline 2339. Graphics SoC interface 2337 provides an interface between graphics core 2300 and other processor cores within a SoC integrated circuit. In at least one embodiment, graphics microcontroller 2338 is a programmable sub-processor that is configurable to manage graphics processing engine 2300. In at least one embodiment, graphics microcontroller 2338 includes programmable logic to facilitate distribution of graphics processing operations to graphics processing engine 2300 and / or to facilitate integration of graphics processor 2300 with other processors or cores. In at least one embodiment, media pipeline 2339 includes logic to assist with decoding, encoding, pre-processing, and / or post-processing of multimedia data like images and video. In at least one embodiment, media pipeline 2339 implements media operations via requests to compute or sample logic within sub-cores 2301-2301F.

[0267] In at least one embodiment, SoC interface 2337 enables graphics core 2300 to communicate with general application processor cores (e.g., CPUs) and / or other components within the SoC, including memory hierarchy elements such as shared Ll cache, system RAM, and / or embedded on-chip or packaged DRAM. In at least one embodiment, SoC interface 2337 can also enable communication with fixed function devices within the SoC, such as a camera image signal processor (ISP) and / or a video codec, and to use and / or implement global memory atoms that can be shared between graphics core 2300 and a CPU within the SoC. In at least one embodiment, SoC interface 2337 can also implement power management controls for graphics core 2300 and enable an interface between a clock domain of graphics core 2300 and other clock domains within the SoC. In at least one embodiment, SoC interface 2337 enables receiving command buffers from a command streamer and global thread dispatcher that are configured to provide commands and instructions to each of one or more graphics cores within a graphics processor. In at least one embodiment, commands and instructions can be dispatched to a media pipeline 2339 when media operations are to be performed, or to a geometry and fixed function pipeline (e.g., geometry and fixed function pipeline 2336, geometry and fixed function pipeline 2314) when graphics processing operations are to be performed.

[0268] In at least one embodiment, graphics microcontroller 2338 can be configured to perform various scheduling and management tasks for graphics core 2300. In at least one embodiment, graphics microcontroller 2338 can perform graphics and / or compute workload scheduling on various graphics processing engines within execution unit (EU) arrays 2302A-2302F, 2304A-2304F in sub-cores 2301A-2301F. In at least one embodiment, host software executing on a CPU core of a SoC including graphics core 2300 can submit workloads for one of graphics processing units that invoke scheduling operations on appropriate graphics engines. In at least one embodiment, scheduling operations include determining which workload to run next, submitting a workload to a command streamer, pre-empting existing workloads running on an engine, monitoring progress of a workload, and notifying host software when a workload completes. In at least one embodiment, graphics microcontroller 2338 can also facilitate low power or idle states for graphics core 2300, providing the ability to save and restore registers across low power states independently of operating systems and / or graphics driver software on the system.

[0269] In at least one embodiment, graphics core 2300 can have more or less than the illustrated sub-cores 2301 A-2301F, up to N modular sub-cores. For each set of N sub-cores, graphics core 2300 can also include, in at least one embodiment, shared function logic 2310, shared and / or cache memory 2312, geometry / fixed function pipeline 2314, and additional fixed function logic 2316 to accelerate various graphics and compute processing operations. In at least one embodiment, shared function logic 2310 can include logic units (e.g., samplers, math, and / or inter-thread communication logic) that are shareable across each N sub-core within graphics core 2300. Shared and / or cache memory 2312 can be an LLC for N sub-cores 2301A-2301F within graphics core 2300, and can also serve as shared memory accessible by multiple sub-cores. In at least one embodiment, geometry / fixed function pipeline 2314 can be included in place of geometry / fixed function pipelines 2336 within fixed function block 2330, and can include the same or similar logic units.

[0270] In at least one embodiment, graphics core 2300 includes additional fixed function logic 2316, which can include various fixed function acceleration logic used by graphics core 2300. In at least one embodiment, additional fixed function logic 2316 includes an additional geometry pipeline used in position-only shading. In position-only shading, there are at least two geometry pipelines, while in a full geometry pipeline and cull pipeline within geometry / fixed function pipeline 2316, 2336, which is an additional geometry pipeline that can be included in additional fixed function logic 2316. In at least one embodiment, the cull pipeline is a trimmed down version of the full geometry pipeline. In at least one embodiment, the full pipeline and the cull pipeline can execute different instances of an application, each with a separate environment. In at least one embodiment, position-only shading can hide long cull runs of triangles that are discarded, which can complete shading earlier in some cases. For example, in at least one embodiment, cull pipeline logic in additional fixed function logic 2316 can execute position shaders in parallel with a main application, and often generate critical results faster than the full pipeline because the cull pipeline takes and shades position attributes of vertices without performing rasterization and rendering pixels to a frame buffer. In at least one embodiment, the cull pipeline can use generated critical results to compute visibility information for all triangles, regardless of whether those triangles are culled or not. In at least one embodiment, the full pipeline, which can be referred to as a replay pipeline in this case, can consume the visibility information to skip culled triangles to only shade visible triangles that are ultimately passed to a rasterization stage.

[0271] In at least one embodiment, additional fixed function logic 2316 can also include general purpose processing acceleration logic, such as fixed function matrix multiplication logic, for implementing a reduced CUAD program.

[0272] In at least one embodiment, within each graphics sub-core 2301A-2301F includes a set of execution resources which can be used to perform graphics, media, and compute operations in response to requests by graphics pipeline, media pipeline, or shader programs. In at least one embodiment, graphics sub-cores 2301A-2301F include multiple arrays of execution units 2302A-2302F, 2304A-2304F, thread dispatch and inter-thread communication (TD / IC) logic 2303A-2303F, 3D (e.g., texture) samplers 2305A-2305F, media samplers 2306A-2306F, shader processors 2307A-2307F, and shared local memory (SLM) 2308A-2308F. Arrays of execution units 2302A-2302F, 2304A-2304F each include multiple execution units (EUs) which are GPGPUs capable of performing floating point and integer / fixed point logic operations in service of a graphics, media, or compute application running on a thread or a thread group. In at least one embodiment, TD / IC logic 2303A-2303F performs local thread dispatch and thread control operations for execution units within a sub-core and facilitate communication between threads executing on execution units of a sub-core. In at least one embodiment, 3D samplers 2305A-2305F can read and return data from 3D textures or other 3D graphics related data structures. In at least one embodiment, 3D samplers can read different types of data from a texture based on a configured sampling state and a texture format associated with a given texture. In at least one embodiment, media samplers 2306A-2306F can perform similar read operations on media data associated with media units or other media data structures. In at least one embodiment, each graphics sub-core 2301A-2301F can alternatively include unified 3D and media samplers. In at least one embodiment, threads executing on execution units within each sub-core 2301A-2301F can utilize shared local memory 2308A-2308F within each sub-core for storage of thread private data and / or shared data between threads executing on execution units within the same sub-core.

[0273] In at least one embodiment, set 112 (see Figure 1 and Figure 3 ) and / or set 114 (see Figure 1 and Figure 3 ) can include one or more of graphics processor cores 2300.

[0274] Figure 24A parallel processing unit (“PPU”) 2400, in accordance with at least one embodiment, is shown. In at least one embodiment, PPU 2400 is configured with machine-readable code that, if executed by PPU 2400, causes PPU 2400 to perform some or all of the processes and techniques described herein. In at least one embodiment, PPU 2400 is a multi-threaded processor implemented on one or more integrated circuit devices and utilizes multi-threading as a latency-hiding technique designed to process computer-readable instructions (also referred to as machine-readable instructions or simply instructions) that are executed in parallel on multiple threads. In at least one embodiment, a thread refers to an execution thread and is an instance of a set of instructions configured to be executed by PPU 2400. In at least one embodiment, PPU 2400 is a graphics processing unit (“GPU”) configured to implement a graphics rendering pipeline for processing three-dimensional (“3D”) graphics data in order to generate two-dimensional (“2D”) image data for display on a display device, such as an LCD device. In at least one embodiment, PPU 2400 is used to perform computations, such as linear algebra operations and machine learning operations. Figure 24 The example parallel processor is shown for illustrative purposes only and should be construed as a non-limiting example of a processor architecture implemented in at least one embodiment.

[0275] In at least one embodiment, PPU(s) 2400 are configured to accelerate high- performance computing (“HPC”), datacenter, and machine learning applications. In at least one embodiment, PPU(s) 2400 are configured to accelerate CUDA programs. In at least one embodiment, PPU 2400 includes, without limitation, I / O unit 2406, front-end unit 2410, scheduler unit 2412, work distribution unit 2414, hub 2416, crossbar (“Xbar”) 2420, one or more general processing clusters (“GPCs”) 2418, and one or more partition units (“memory partition units”) 2422. In at least one embodiment, PPU(s) 2400 connect to a host processor or other PPU(s) 2400 by one or more high-speed GPU interconnects 2408. In at least one embodiment, PPU(s) 2400 connect to host processors or other peripherals by a system bus or interconnect 2402. In at least one embodiment, PPU(s) 2400 connect to a local memory comprising one or more memory devices (“memory”) 2404. In at least one embodiment, memory devices 2404 include, without limitation, one or more dynamic random access memory (“DRAM”) devices. In at least one embodiment, one or more DRAM devices are configured and / or configurable as high-bandwidth memory (“HBM”) subsystems with multiple DRAM dies stacked

[0276] In at least one embodiment, high-speed GPU interconnect 2408 can refer to a link-based parallel computer bus that systems use to scale and includes one or more PPUs 2400 in conjunction with one or more CPUs (“CPU(s)”), supports cache coherency between PPUs 2400 and CPUs, and CPU mastering. In at least one embodiment, high-speed GPU interconnect 2408 transports data and / or commands through hub 2416 to other units of PPU(s) 2400, such as one or more copy engines, video encoders, video decoders, power management units, and / or other components not explicitly shown in Figure 24

[0277] In at least one embodiment, I / O unit 2406 is configured to facilitate communication between PPU(s) 2400 and a host processor (not shown), other PPU(s) 2400, and / or one or more Figure 24 ​The I / O units 2406 send and receive communications (e.g., commands, data) to and from the system bus 2402. In at least one embodiment, the I / O units 2406 communicate directly with the host processor(s) via the system bus 2402 or through one or more intermediate devices such as a memory hub. In at least one embodiment, the I / O units 2406 can communicate with one or more other processors, such as one or more PPUs 2400, via the system bus 2402. In at least one embodiment, the I / O units 2406 implement a PCIe interface for communications over a PCIe bus. In at least one embodiment, the I / O units 2406 implement interfaces for communicating with external devices.

[0278] In at least one embodiment, the I / O units 2406 decode packets received via the system bus 2402. In at least one embodiment, at least some packets represent commands configured to cause the PPU 2400 to perform various operations. In at least one embodiment, the I / O units 2406 send decoded commands to various other units of the PPU 2400 as designated by the commands. In at least one embodiment, commands are sent to the front-end unit 2410 and / or to the hub 2416 or other units of the PPU 2400 such as one or more copy engines, a video encoder, a video decoder, a power management unit, etc. Figure 24 In at least one embodiment, the I / O units 2406 are not explicitly shown in FIG. 2. In at least one embodiment, the I / O units 2406 are configured to route communications between various logical units of the PPU 2400.

[0279] In at least one embodiment, a program executed by the host processor encodes a command stream in a buffer that provides a workload to the PPU 2400 for processing. In at least one embodiment, the workload includes instructions and data to be processed by those instructions. In at least one embodiment, the buffer is a region in memory that is accessible (e.g., read / write) by both the host processor and the PPU 2400 - the host interface unit can be configured to access memory requests transmitted by the I / O units 2406 over the system bus 2402 to the buffer in system memory connected to the system bus 2402. In at least one embodiment, the host processor writes the command stream to the buffer and then sends a pointer to the start of the command stream to the PPU 2400 so that the front-end unit 2410 receives the one or more command stream pointers and manages the one or more command streams, reading commands from the command stream and forwarding the commands to various units of the PPU 2400.

[0280] In at least one embodiment, front-end unit 2410 is coupled to a scheduler unit 2412 which configures various GPCs 2418 to process tasks defined by one or more command streams. In at least one embodiment, scheduler unit 2412 is configured to track state information related to various tasks managed by scheduler unit 2412, where state information can indicate which task is assigned to which GPC 2418, whether task is active or inactive, priority of task associated with it, and so forth. In at least one embodiment, scheduler unit 2412 manages multiple tasks that are executed on one or more GPCs 2418.

[0281] In at least one embodiment, scheduler unit 2412 is coupled to a work distribution unit 2414, which is configured to apportion tasks to be performed by GPCs 2418. In at least one embodiment, work distribution unit 2414 tracks multiple scheduled tasks received from scheduler unit 2412 and manages a pending task pool and an active task pool for each GPC 2418. In at least one embodiment, pending task pool includes multiple slots (e.g., 32 slots) that contain tasks assigned to be processed by a particular GPC 2418; active task pool can include multiple slots (e.g., 4 slots) for tasks that are actively being processed by GPC 2418, such that as one of GPCs 2418 completes processing a task, that task is evicted from active task pool for GPC 2418 and one of other tasks from pending task pool is selected and scheduled for execution on GPC 2418. In at least one embodiment, if an active task is idle, for example, while waiting for a data dependency to resolve, active task is evicted from GPC 2418 and returned to pending task pool, while another task from pending task pool is selected and scheduled for execution on GPC 2418.

[0282] In at least one embodiment, work distribution unit 2414 communicates with one or more GPCs 2418 via XBar 2420. In at least one embodiment, XBar 2420 is an interconnect network that couples many units of PPU 2400 to other units of PPU 2400 and can be configured to couple work distribution unit 2414 to a particular GPC 2418. In at least one embodiment, other units of one or more PPU(s) 2400 can also be connected to XBar 2420 via hub 2416.

[0283] In at least one embodiment, tasks are managed by a scheduler unit 2412 and dispatched to one of GPCs 2418 by a work distribution unit 2414. GPCs 2418 are configured to process tasks and generate results. In at least one embodiment, results can be consumed by other tasks within GPC 2418, routed to different GPCs 2418 via XBar 2420 or stored in memory 2404. In at least one embodiment, results can be written to memory 2404 via a partition unit 2422, which implements a memory interface for reading and writing data to memory 2404. In at least one embodiment, results can be transmitted to another PPU 2400 or CPU via a high-speed GPU interconnect 2408. In at least one embodiment, PPU 2400 includes, without limitation, U partition units 2422 equal to a number of separate and distinct memory devices 2404 coupled to PPU 2400.

[0284] In at least one embodiment, a host processor executes a driver core that implements an application programming interface (API) that enables one or more applications executing on a host processor to schedule operations to be performed on PPU 2400. In one embodiment, multiple compute applications are executed simultaneously by PPU 2400 and PPU 2400 provides isolation, quality of service (“QoS”), and independent address spaces for multiple compute applications. In at least one embodiment, an application generates instructions (e.g., in the form of API calls) that cause a driver core to generate one or more tasks for execution by PPU 2400 and driver core outputs tasks to one or more streams processed by PPU 2400. In at least one embodiment, each task includes one or more related thread groups, which can be referred to as warps. In at least one embodiment, a warp includes multiple related threads (e.g., 32 threads) that can be executed in parallel. In at least one embodiment, cooperating threads can refer to multiple threads including instructions for performing a task and exchanging data via shared memory.

[0285] In at least one embodiment, set 112 (see Figure 1 and Figure 3 ) and / or group 114 (see Figure 1 and Figure 3 ) can include one or more of PPU 2400.

[0286] Figure 25 A GPC 2500 according to at least one embodiment is shown. In at least one embodiment, GPC 2500 is a Figure 24GPC 2418. In at least one embodiment, each GPC 2500 includes, without limitation, a number of hardware units for handling compute tasks and each GPC 2500 includes, without limitation, a pipeline manager 2502, a pre-raster operations unit (“PROP”) 2504, a raster engine 2508, a work distribution crossbar (“WDX”) 2516, a memory management unit (“MMU”) 2518, one or more data processing clusters (“DPCs”) 2506, and any suitable combination of such components.

[0287] In at least one embodiment, operation of GPC 2500 is controlled by pipeline manager 2502. In at least one embodiment, pipeline manager 2502 manages configuration of one or more DPCs 2506 to process tasks assigned to GPC 2500. In at least one embodiment, pipeline manager 2502 configures at least one of one or more DPCs 2506 to implement at least a portion of a graphics rendering pipeline. In at least one embodiment, DPC 2506 is configured to execute vertex shader programs on a programmable streaming multi-processor (“SM”) 2514. In at least one embodiment, pipeline manager 2502 is configured to route packets received from a work distribution unit to appropriate logical units within GPC 2500, and in at least one embodiment, some packets can be routed to fixed function hardware units in PROP 2504 and / or raster engine 2508 while other packets can be routed to DPCs 2506 for processing by a primitive engine 2512 or SM 2514. In at least one embodiment, pipeline manager 2502 configures at least one of DPCs 2506 to implement a neural network model and / or compute pipeline. In at least one embodiment, pipeline manager 2502 configures at least one of DPCs 2506 to execute at least a portion of a CUDA program.

[0288] In at least one embodiment, PROP unit 2504 is configured to route data generated by raster engine 2508 and DPCs 2506 to a pixel stage in a partition unit, such as a render output unit (“ROP”) of GPC 2418 as discussed above in relation to FIG. 24A. In at least one embodiment, ROP 2510 is configured to perform graphics pixel operations related to two-dimensional rendering of graphics primitives. Figure 24Memory partition unit 2422, among other things, is described in greater detail below. In at least one embodiment, PROP unit 2504 is configured to perform optimizations for color blending, organize pixel data, perform address translations, and the like. In at least one embodiment, raster engine 2508 includes, without limitation, a number of fixed function hardware units configured to perform various raster operations, and in at least one embodiment, raster engine 2508 includes, without limitation, a setup engine, a coarse raster engine, a cull engine, a clip engine, a fine raster engine, a tile aggregation engine, and any suitable combinations thereof. In at least one embodiment, setup engine receives transformed vertices and generates plane equations associated with geometric primitives defined by the vertices; plane equations are passed to coarse raster engine to generate coverage information (e.g., x, y coverage masks for tiles) for the primitive; output from coarse raster engine is passed to cull engine, where fragments associated with primitives that fail z-test are culled, and to clip engine, where fragments that are outside of a viewing frustum are clipped. In at least one embodiment, clipped and culled fragments are passed to fine raster engine to generate attributes of pixel fragments based on plane equations generated by setup engine. In at least one embodiment, output from raster engine 2508 includes fragments to be processed by any suitable entity, such as by a fragment shader implemented within DPC 2506.

[0289] In at least one embodiment, each DPC 2506 included in GPC 2500 includes, without limitation, an M-Pipe Controller (“MPC”) 2510; a primitive engine 2512; one or more SMs 2514; and any suitable combinations thereof. In at least one embodiment, MPC 2510 controls operation of DPC 2506, routing received packets from pipeline manager 2502 to appropriate units in DPC 2506. In at least one embodiment, packets associated with vertices are routed to primitive engine 2512, which is configured to fetch vertex attributes associated with the vertices from memory; in contrast, packets associated with a shader program can be transmitted to SM 2514.

[0290] In at least one embodiment, SM 2514 comprises, without limitation, a programmable streaming processor configured to process tasks represented by a plurality of threads. In at least one embodiment, SM 2514 is multi-threaded and configured to execute a plurality of threads (e.g., 32 threads) from a particular group of threads concurrently and implements a single-instruction, multiple-data (“SIMD”) architecture wherein each thread in group of threads is configured to process a different data set based on same set of instructions. In at least one embodiment, all threads in a group of threads execute same instructions. In at least one embodiment, SM 2514 implements a single-instruction, multiple-thread (“SIMT”) architecture wherein each thread in a group of threads is configured to process a different data set based on same set of instructions, but where threads in group of threads are allowed to diverge during execution. In at least one embodiment, program counter, call stack, and execution state are maintained for each threadlet, enabling concurrency between threadlets and serial execution within threadlet as threads diverge. In another embodiment, program counter, call stack, and execution state are maintained for each individual thread, enabling equal concurrency between all threads within and across threadlets. In at least one embodiment, execution state is maintained for each individual thread and threads executing same instructions can be converged and executed in parallel for better efficiency. Details regarding SM 2514 are provided below in conjunction with FIG. 27. Figure 26 At least one embodiment of SM 2514 is described in greater detail.

[0291] In at least one embodiment, MMU 2518 interfaces with GPC 2500 and memory partition unit (e.g., partition unit 2422 of FIG. 24) and MMU 2518 provides translations of virtual addresses to physical addresses, memory protection, and demand paging. Figure 24 In at least one embodiment, MMU 2518 provides one or more translation lookaside buffers (TLBs) for handling translation of virtual addresses to physical addresses in memory.

[0292] In at least one embodiment, clusters 112 (see Figure 1 and Figure 3 ) and / or groups 114 (see Figure 1 and Figure 3 ) can comprise one or more of GPCs 2500. In at least one embodiment, Figure 25 Xbar shown in FIG. 25A can be implemented as switch circuit 110.

[0293] Figure 26 A streaming multiprocessor (“SM”) 2600, according to at least one embodiment, is shown. In at least one embodiment, SM 2600 is a Figure 25SM 2514. In at least one embodiment, SM 2600 includes, without limitation, an instruction cache 2602; one or more scheduler units 2604; a register file 2608; one or more processing cores (“cores”) 2610; one or more special-function units (“SFUs”) 2612; one or more load / store units (“LSUs”) 2614; an interconnect network 2616; shared memory / level-one (“LI”) cache 2618; and any suitable combination thereof. In at least one embodiment, a work distribution unit dispatches tasks for execution on general processing clusters (“GPCs”) of parallel processing units (“PPUs”) and each task is assigned a specific data processing cluster (“DPC”) within a GPC and, if task is associated with a shader program, to one of SMs 2600. In at least one embodiment, scheduler units 2604 receive tasks from work distribution unit and manage scheduling of instruction dispatch to one or more thread blocks assigned to SM 2600. In at least one embodiment, scheduler units 2604 schedule thread blocks to be executed to be executed as warps of parallel threads, with each thread block allocated at least one warp. In at least one embodiment, each warp executes a thread. In at least one embodiment, scheduler units 2604 manage a plurality of different thread blocks, allocating warps of threads to different thread blocks, and then dispatching instructions from different cooperating groups of warps to various functional units (e.g., processing cores 2610, SFUs 2612, and LSUs 2614) during each clock cycle.

[0294] In at least one embodiment, a “cooperative group” can refer to a programming model for organizing groups of communication threads that allows developers to express the granularity at which threads are communicating, enabling richer, more efficient parallel decomposition. In at least one embodiment, a cooperative launch API supports synchronization between thread blocks to execute parallel algorithms. In at least one embodiment, an API of a conventional programming model provides a single, simple construct for synchronizing cooperative threads: a barrier across all threads of a thread block (e.g., a syncthreads() function). However, in at least one embodiment, a programmer can define thread groups at less than a thread block granularity and synchronize within defined groups to achieve higher performance, design flexibility, and software reuse in the form of collective group-wide function interfaces. In at least one embodiment, cooperative groups enable programmers to explicitly define thread groups at sub-block and multi-block granularity and perform collective operations, such as synchronizing threads in a cooperative group. In at least one embodiment, sub-block granularity is as small as a single thread. In at least one embodiment, a programming model supports clean composition across software boundaries, so that library and utility functions can safely synchronize in their local environment without having to make assumptions about convergence. In at least one embodiment, cooperative group primitives enable new patterns of cooperative parallelism, including but not limited to producer-consumer parallelism, opportunistic parallelism, and global synchronization across a grid of thread blocks.

[0295] In at least one embodiment, dispatch units 2606 are configured to send instructions to one or more of the functional units, and a scheduler unit 2604 includes, without limitation, two dispatch units 2606 that enable two different instructions from the same thread to be dispatched in each clock cycle. In at least one embodiment, each scheduler unit 2604 includes a single dispatch unit 2606 or an additional dispatch unit 2606.

[0296] In at least one embodiment, each SM 2600 includes, without limitation, a register file 2608 that provides a set of registers for functional units of the SM 2600. In at least one embodiment, register file 2608 is split between functional units as is needed to perform the computational and / or logical operations of the SM 2600. In at least one embodiment, register file 2608 is partitioned between different thread blocks being executed by the SM 2600 and is further partitioned between different warps being executed by the SM 2600. In at least one embodiment, register file 2608 provides temporary storage for operands of the SM 2600’s data paths. In at least one embodiment, each SM 2600 includes, without limitation, a plurality L of processing cores 2610. In at least one embodiment, SM 2600 includes, without limitation, a large number (e.g., 128 or more) of different processing cores 2610. In at least one embodiment, each processing core 2610 includes, without limitation, a full-pipe, single precision, double precision, and / or mixed precision processing unit including, without limitation, a floating point arithmetic logic unit and an integer arithmetic logic unit. In at least one embodiment, floating point arithmetic logic units implement IEEE 754-2008 standard for floating point arithmetic. In at least one embodiment, processing cores 2610 include, without limitation, 64 single precision (32-bit) floating point cores, 64 integer cores, 32 double precision (64-bit) floating point cores, and 8 tensor cores.

[0297] In at least one embodiment, tensor cores are configured to perform matrix operations. In at least one embodiment, one or more tensor cores are included in processing cores 2610. In at least one embodiment, tensor cores are configured to perform deep learning matrix arithmetic, such as convolution operations for neural network training and inferencing. In at least one embodiment, each tensor core operates on 4x4 matrices and performs matrix multiplication and accumulation operations D = A x B + C, where A, B, C, and D are 4x4 matrices.

[0298] In at least one embodiment, matrix multiplication inputs A and B are 16-bit floating point matrices, and accumulation matrices C and D are 16-bit floating point or 32-bit floating point matrices. In at least one embodiment, a tensor core performs 32-bit floating point accumulation operations on 16-bit floating point input data. In at least one embodiment, 16-bit floating point multiplication uses 64 operations and results in a full precision product, which is then accumulated with other intermediate products using 32-bit floating point addition for 4x4x4 matrix multiplication. In at least one embodiment, tensor cores are used to perform larger two-dimensional or higher dimensional matrix operations composed of these smaller elements. In at least one embodiment, an API such as a CUDA-C++ API exposes specialized matrix load, matrix multiply and accumulate, and matrix store operations to efficiently use tensor cores from a CUDA-C++ program. In at least one embodiment, at a CUDA level, a warp level interface assumes a 16x16 size matrix across all 32 warp threads.

[0299] In at least one embodiment, each SM 2600 includes, without limitation, M SFUs 2612 to perform special functions (e.g., certain math functions, atomics, bit endcomplement, etc.). In at least one embodiment, SFUs 2612 include, without limitation, tree traversal units configured to traverse a hierarchical tree data structure. In at least one embodiment, SFUs 2612 include, without limitation, texture units configured to perform texture mapping operations. In at least one embodiment, texture units are configured to load a texture map (e.g., a 2D array of texture pixels) from memory and sample the texture map to produce sampled texture values for use by a shader program executed by SM 2600. In at least one embodiment, texture maps are stored in shared memory / L1 cache 2618. In at least one embodiment, texture units use mip-maps (e.g., different levels of detail for a texture map) to perform texture operations such as filtering operations. In at least one embodiment, each SM 2600 includes, without limitation, two texture units.

[0300] In at least one embodiment, each SM 2600 includes, without limitation, N LSUs 2614 that implement load and store operations between shared memory / L1 cache 2618 and register file 2608. In at least one embodiment, each SM 2600 includes, without limitation, interconnect network 2616 that connects each of the functional units to register file 2608 and LSUs 2614 to register file 2608 and shared memory / L1 cache 2618. In at least one embodiment, interconnect network 2616 is a cross-bar switch that can be configured to connect any of the functional units to any of the registers in register file 2608 and connect LSUs 2614 to registers in register file 2608 and memory locations in shared memory / L1 cache 2618.

[0301] In at least one embodiment, shared memory / L1 cache 2618 is an array of on-chip memory that, in at least one embodiment, allows data storage and communication between SMs 2600 and graphics primitives engines and threads within SMs 2600. In at least one embodiment, shared memory / L1 cache 2618 includes, without limitation, 128 KB of storage and is located on a path from SM 2600 to partition units. In at least one embodiment, shared memory / L1 cache 2618 is used for caching reads and writes, in at least one embodiment. In at least one embodiment, one or more of shared memory / L1 cache 2618, L2 cache, and memory are backing stores.

[0302] In at least one embodiment, combining data cache and shared memory functionality into a single memory block provides improved performance for both types of memory access. In at least one embodiment, capacity is used by programs that do not use shared memory or used as a cache, e.g., if shared memory is configured to use half of capacity, then textures and load / store operations can use remaining capacity. According to at least one embodiment, integration within shared memory / L1 cache 2618 enables shared memory / L1 cache 2618 to be used as a high-throughput pipe for streaming data while providing high bandwidth and low latency access to frequently reused data. In at least one embodiment, when configured for general purpose parallel computation, a simpler configuration can be used compared to graphics processing. In at least one embodiment, fixed function GPU is bypassed, creating a more straightforward programming model. In at least one embodiment, in a general purpose parallel computation configuration, work distribution unit allocates and distributes blocks of threads directly to DPCs. In at least one embodiment, threads in a block execute the same program, use a unique thread ID in a computation to ensure that each thread generates a unique result, use SM 2600 to execute the program and perform the computation, use shared memory / L1 cache 2618 to communicate between threads, and use LSU 2614 to read and write global memory through shared memory / L1 cache 2618 and memory partition unit. In at least one embodiment, when configured for general purpose parallel computation, SM 2600 writes commands to scheduler unit 2604 that can be used to launch new work on DPCs.

[0303] In at least one embodiment, PPU is included in a desktop computer, laptop computer, tablet computer, server computer, supercomputer, smart- phone (e.g., a wireless, hand-held device), PDA, digital camera, vehicle, head mounted display, hand-held electronic device, etc. or is coupled to such devices. In at least one embodiment, PPU is implemented on a single semiconductor

[0304] In at least one embodiment, PPU can be included on a graphics card that includes one or more memory devices. In at least one embodiment, graphics card can be configured to connect with a PCIe slot on a motherboard of a desktop computer. In at least one embodiment, PPU can be an integrated GPU (“iGPU”) included in a chipset of a motherboard.

[0305] In at least one embodiment, set 112 (see Figure 1 and Figure 3 ) and / or group 114 (see Figure 1and Figure 3 ) can include one or more of the SMs 2600.

[0306] Software constructs for general computing

[0307] The following figures set forth, without limitation, example software constructs for implementing at least one embodiment.

[0308] Figure 27 A software stack of a programming platform, in accordance with at least one embodiment, is shown. In at least one embodiment, a programming platform is a platform for utilizing hardware on a computing system to accelerate computational tasks. In at least one embodiment, a software developer can access a programming platform through libraries, compiler directives, and / or extensions to a programming language. In at least one embodiment, a programming platform can be, but is not limited to, CUDA, Radeon Open Compute Platform (“ROCm”), OpenCL (OpenCL TM ), SYCL, or Intel One API.

[0309] In at least one embodiment, a software stack 2700 of a programming platform provides an execution environment for an application 2701. In at least one embodiment, an application 2701 can include any computer software capable of launching on a software stack 2700. In at least one embodiment, an application 2701 can include, without limitation, an artificial intelligence (“AI”) / machine learning (“ML”) application, a high performance computing (“HPC”) application, a virtual desktop infrastructure (“VDI”), or a data center workload.

[0310] In at least one embodiment, an application 2701 and a software stack 2700 run on hardware 2707. In at least one embodiment, hardware 2707 can include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of computing devices that support a programming platform. In at least one embodiment, for example with CUDA, a software stack 2700 can be vendor-specific and only compatible with devices from a particular vendor. In at least one embodiment, for example in OpenCL, a software stack 2700 can be used with devices from different vendors. In at least one embodiment, hardware 2707 includes a host connected to one or more devices that can be accessed via an application programming interface (API) call to perform a computational task. In at least one embodiment, in contrast to a host within hardware 2707, which can include, without limitation, a CPU (but can also include a computing device) and its memory, a device within hardware 2707 can include, without limitation, a GPU, FPGA, AI engine, or other computing device (but can also include a CPU) and its memory.

[0311] In at least one embodiment, software stack 2700 of a programming platform includes, without limitation, a plurality of libraries 2703, a runtime 2705, and a device kernel driver 2706. In at least one embodiment, each of libraries 2703 can include data and programming code that can be used by computer programs and leveraged during software development. In at least one embodiment, libraries 2703 can include, without limitation, pre-written code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. In at least one embodiment, libraries 2703 include functions that are optimized for execution on one or more types of devices. In at least one embodiment, libraries 2703 can include, without limitation, functions for performing mathematical, deep learning, and / or other types of operations on a device. In at least one embodiment, libraries 2703 are associated with corresponding APIs 2702, which can include one or more APIs that expose functions implemented in libraries 2703.

[0312] In at least one embodiment, application 2701 is written as source code that is compiled into executable code, as discussed in more detail below in conjunction with FIG. 27B. In at least one embodiment, executable code of application 2701 can run, at least partially, on an execution environment provided by software stack 2700. In at least one embodiment, during execution of application 2701, code can be derived that needs to run on a device (as opposed to a host). In this case, in at least one embodiment, runtime 2705 can be invoked to load and launch the necessary code on a device. In at least one embodiment, runtime 2705 can include any technically feasible runtime system capable of supporting execution of application 2701. Figures 32-34 In at least one embodiment, runtime 2705 is implemented as one or more runtime libraries associated with corresponding APIs (which are shown as APIs 2704). In at least one embodiment, one or more such runtime libraries can include, without limitation, functions for memory management, execution control, device management, error handling, and / or synchronization, among others. In at least one embodiment, memory management functions can include, without limitation, functions for allocating, deallocating, and copying device memory, as well as transferring data between host memory and device memory. In at least one embodiment, execution control functions can include, without limitation, functions for launching functions on a device (sometimes referred to as “kernels” when functions are global functions that are callable from a host), and functions for setting attribute values in buffers maintained by a runtime library for a given function to be executed on a device.

[0313] In at least one embodiment, runtime 2705 is implemented as one or more runtime libraries associated with corresponding APIs (which are shown as APIs 2704). In at least one embodiment, one or more such runtime libraries can include, without limitation, functions for memory management, execution control, device management, error handling, and / or synchronization, among others. In at least one embodiment, memory management functions can include, without limitation, functions for allocating, deallocating, and copying device memory, as well as transferring data between host memory and device memory. In at least one embodiment, execution control functions can include, without limitation, functions for launching functions on a device (sometimes referred to as “kernels” when functions are global functions that are callable from a host), and functions for setting attribute values in buffers maintained by a runtime library for a given function to be executed on a device.

[0314] In at least one embodiment, runtime libraries and corresponding APIs 2704 can be implemented in any technically feasible manner. In at least one embodiment, one (or any number) of APIs can expose a low-level set of functions for fine-grained control of a device, while another (or any number) of APIs can expose a higher-level set of functions. In at least one embodiment, high-level runtime APIs can be built on top of low-level APIs. In at least one embodiment, one or more runtime APIs can be language-specific APIs layered on top of language-independent runtime APIs.

[0315] In at least one embodiment, a device kernel driver 2706 is configured to facilitate communication with underlying devices. In at least one embodiment, device kernel driver 2706 can provide low-level functions relied upon by APIs such as APIs 2704 and / or other software. In at least one embodiment, device kernel driver 2706 can be configured to compile intermediate representation (“IR”) code into binary code at runtime. In at least one embodiment, for CUDA, device kernel driver 2706 can compile non-hardware-specific parallel thread execution (“PTX”) IR code into binary code for a particular target device at runtime (caching compiled binary code), which is sometimes also referred to as “final” code. In at least one embodiment, doing so can allow final code to run on a target device that can not have existed when source code was initially compiled into PTX code. Alternatively, in at least one embodiment, device source code can be compiled into binary code offline without requiring device kernel driver 2706 to compile IR code at runtime.

[0316] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement at least a portion of software stack 2700.

[0317] Figure 28 A CUDA implementation of software stack 2700 is shown in accordance with at least one embodiment of Figure 27 In at least one embodiment, CUDA software stack 2800 on which application 2801 can be launched includes CUDA libraries 2803, CUDA runtime 2805, CUDA driver 2807, and device kernel driver 2808. In at least one embodiment, CUDA software stack 2800 executes on hardware 2809, which can include a CUDA-enabled GPU developed by NVIDIA Corporation of Santa Clara, California.

[0318] In at least one embodiment, application 2801, CUDA runtime 2805, and device kernel driver 2808 can perform similar functions as application 2701, runtime 2705, and device kernel driver 2706, respectively, described above in connection with FIG. 27. Figure 27 which is described above. In at least one embodiment, CUDA driver 2807 includes a library (libcuda.so) that implements CUDA driver API 2806. In at least one embodiment, similar to CUDA runtime API 2804 implemented by CUDA runtime library (cudart), CUDA driver API 2806 can expose, without limitation, functions for memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability, etc. In at least one embodiment, CUDA driver API 2806 differs from CUDA runtime API 2804 in that CUDA runtime API 2804 simplifies device code management by providing implicit initialization, context (similar to a process) management, and module (similar to a dynamically loaded library) management. In contrast to high-level CUDA runtime API 2804, in at least one embodiment, CUDA driver API 2806 is a low-level API that provides more fine-grained control over a device, particularly with respect to context and module loading. In at least one embodiment, CUDA driver API 2806 can expose functions for context management that are not exposed by CUDA runtime API 2804. In at least one embodiment, CUDA driver API 2806 is also language agnostic and supports, for example, OpenCL in addition to CUDA runtime API 2804. Further, in at least one embodiment, development libraries including CUDA runtime 2805 are considered separate from driver components, including user-mode CUDA driver 2807 and kernel-mode device driver 2808 (sometimes also referred to as a “display” driver).

[0319] In at least one embodiment, CUDA libraries 2803 can include, without limitation, mathematical libraries, deep learning libraries, parallel algorithm libraries, and / or signal / image / video processing libraries that can be utilized by parallel computing applications such as application 2801. In at least one embodiment, CUDA libraries 2803 can include mathematical libraries such as a cuBLAS library, which is an implementation of basic linear algebra subprograms (“BLAS”) for performing linear algebra operations; a cuFFT library for computing fast Fourier transforms (“FFTs”), and a cuRAND library for generating random numbers, etc. In at least one embodiment, CUDA libraries 2803 can include deep learning libraries such as a cuDNN library for primitives of deep neural networks and a TensorRT platform for high-performance deep learning inference, etc.

[0320] In at least one embodiment, system 100 (see Figure 1 and Figure 3 ) can be used to implement at least a portion of CUDA software stack 2800.

[0321] Figure 29 ROCm implementation of software stack 2700 is shown, in accordance with at least one embodiment. In at least one embodiment, ROCm software stack 2900 on which application 2901 can be launched includes language runtime 2903, system runtime 2905, thunk 2907, and ROCm kernel driver 2908. In at least one embodiment, ROCm software stack 2900 executes on hardware 2909, which can include a GPU that supports ROCm, which was developed by AMD Corporation of Santa Clara, CA. Figure 27

[0322] In at least one embodiment, application 2901 can perform similar functionality as application 2701 discussed above in connection with Figure 27 In addition, in at least one embodiment, language runtime 2903 and system runtime 2905 can perform similar functionality as runtime 2705 discussed above in connection with Figure 27 In at least one embodiment, language runtime 2903 and system runtime 2905 differ in that system runtime 2905 is a language-agnostic runtime that implements ROCr system runtime API 2904 and utilizes a Heterogeneous System Architecture (“HSA”) runtime API. In at least one embodiment, the HSA runtime API is a thin user-mode API that exposes interfaces for accessing and interacting with AMD GPUs, including functions for memory management, execution control dispatching of kernels through the architecture, error handling, system and agent information, and runtime initialization and shutdown, among others. In at least one embodiment, language runtime 2903 is an implementation of a language-specific runtime API 2902 layered on top of ROCr system runtime API 2904, as compared to system runtime 2905. In at least one embodiment, a language runtime API can include, without limitation, a Portable Compute Interface (“HIP”) language runtime API, a Heterogeneous Compute Compiler (“HCC”) language runtime API, or an OpenCL API, among others. In particular, the HIP language is an extension of the C++ programming language with functionally similar versions of CUDA mechanisms, and in at least one embodiment, a HIP language runtime API includes functions similar to CUDA runtime API 2804 discussed above in connection with Figure 28

[0323] ​​In at least one embodiment, thunk (ROCt) 2907 is an interface 2906 that can be used to interface with underlying ROCm drivers 2908. In at least one embodiment, ROCm drivers 2908 are ROCk drivers, which are a combination of AMDGPU drivers and HSA kernel drivers (amdkfd). In at least one embodiment, AMDGPU drivers are device kernel drivers for GPUs developed by AMD that perform similar functions to those discussed above in connection with Figure 27 HSA kernel drivers. In at least one embodiment, HSA kernel drivers are drivers that allow different types of processors to more efficiently share system resources via hardware features.

[0324] In at least one embodiment, various libraries (not shown) can be included in ROCm software stack 2900 above language runtime 2903 and provide similar functionality to CUDA libraries 2803 discussed above in connection with Figure 28 In at least one embodiment, various libraries can include, but are not limited to, math, deep learning, and / or other libraries such as a hipBLAS library that implements similar functions to CUDA cuBLAS, a rocFFT library similar to CUDA cuFFT for computing FFTs, etc.

[0325] Figure 30 FIG. 29 illustrates an OpenCL implementation of software stack 2900 in accordance with at least one embodiment. Figure 27 In at least one embodiment, OpenCL software stack 3000 on which application 3001 can be launched includes an OpenCL framework 3010, an OpenCL runtime 3006, and a driver 3007. In at least one embodiment, OpenCL software stack 3000 executes on hardware 2809 that is not vendor-specific. In at least one embodiment, because OpenCL is supported by devices developed by different vendors, specific OpenCL drivers can be required to interoperate with hardware from such vendors.

[0326] In at least one embodiment, application 3001, OpenCL runtime 3006, device kernel driver 3007, and hardware 3008 can perform similar functions to those discussed above in connection with Figure 27 application 2701, runtime 2705, device kernel driver 2706, and hardware 2707 discussed above. In at least one embodiment, application 3001 also includes an OpenCL kernel 3002 with code that is to be executed on a device.

[0327] In at least one embodiment, OpenCL defines a “platform” that allows a host to control devices connected to that host. In at least one embodiment, OpenCL framework provides a platform layer API and a runtime API, shown as platform API 3003 and runtime API 3005. In at least one embodiment, runtime API 3005 uses a context to manage execution of kernels on a device. In at least one embodiment, each identified device can be associated with a respective context, which runtime API 3005 can use to manage a command queue, program and kernel objects, shared memory objects, etc. for that device. In at least one embodiment, platform API 3003 exposes functions that allow device contexts to be used for selecting and initializing devices, submitting work to devices via command queues, and enabling data transfers to and from devices, among other things. Additionally, in at least one embodiment, OpenCL framework provides various built-in functions (not shown), including mathematical functions, relational functions, and image processing functions, among others.

[0328] In at least one embodiment, compiler 3004 is also included in OpenCL framework 3010. In at least one embodiment, source code can be compiled offline before executing an application or online during execution of an application. In contrast to CUDA and ROCm, OpenCL applications in at least one embodiment can be compiled online by compiler 3004, which is included to represent any number of compilers that can be used to compile source and / or IR code (e.g., Standard Portable Intermediate Representation (“SPIR-V”) code) into binary code. Alternatively, in at least one embodiment, OpenCL applications can be compiled offline before executing such applications.

[0329] Figure 31 Software supported by a programming platform according to at least one embodiment is shown. In at least one embodiment, programming platform 3104 is configured to support various programming models 3103, middleware and / or libraries 3102, and frameworks 3101 that applications 3100 can rely on. In at least one embodiment, applications 3100 can be AI / ML applications implemented using, for example, deep learning frameworks (e.g., MXNet, PyTorch, or TensorFlow) that can rely on libraries such as cuDNN, NVIDIA Collective Communications Library (“NCCL”), and / or NVIDIA Developer Data Loading Library (“DALI”) CUDA libraries to provide accelerated computing on underlying hardware.

[0330] In at least one embodiment, programming platform 3104 can be any of the programming platforms described above in connection with FIGS. 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97, 98, 99, and 100 respectively.Figure 28 , Figure 29 and Figure 30 One of the CUDA, ROCm, or OpenCL platforms described above. In at least one embodiment, programming platform 3104 supports multiple programming models 3103, which are abstractions of the underlying computing system that allow for expression of algorithms and data structures. In at least one embodiment, programming models 3103 can expose features of the underlying hardware in order to improve performance. In at least one embodiment, programming models 3103 can include, but are not limited to, CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism (“C++ AMP”), Open Multi-Processing (“OpenMP”), Open Accelerators (“OpenACC”), and / or Vulcan Compute.

[0331] In at least one embodiment, libraries and / or middleware 3102 provide implementations of the abstractions of programming models 3104. In at least one embodiment, such libraries include data and programming code that can be used by computer programs and utilized during software development. In at least one embodiment, such middleware includes software that provides services to applications in addition to those that can be available from programming platform 3104. In at least one embodiment, libraries and / or middleware 3102 can include, but are not limited to, cuBLAS, cuFFT, cuRAND, and other CUDA libraries, or rocBLAS, rocFFT, rocRAND, and other ROCm libraries. Additionally, in at least one embodiment, libraries and / or middleware 3102 can include NCCL and ROCm Communication Collection Library (“RCCL”) libraries, which provide communication routines for GPUs, MIOpen libraries for deep learning acceleration, and / or Eigen libraries for linear algebra, matrix and vector operations, geometric transformations, numerical solvers, and related algorithms.

[0332] In at least one embodiment, application frameworks 3101 rely on libraries and / or middleware 3102. In at least one embodiment, each application framework 3101 is a software framework used to implement a standard structure for application software. Returning to the AI / ML example discussed above, in at least one embodiment, an AI / ML application can be implemented using a framework such as a Caffe, Caffe2, TensorFlow, Keras, PyTorch, or MxNet deep learning framework.

[0333] Figure 32 shows compiled code to run on Figures 27-30on one of the programming platforms. In at least one embodiment, compiler 3201 receives source code 3200, which includes both host code and device code. In at least one embodiment, compiler 3201 is configured to convert source code 3200 into host executable code 3202 for execution on a host and device executable code 3203 for execution on a device. In at least one embodiment, source code 3200 can be compiled offline, prior to execution of an application, or online during execution of an application.

[0334] In at least one embodiment, source code 3200 can include code in any programming language supported by compiler 3201, such as C++, C, Fortran, etc. In at least one embodiment, source code 3200 can include a single-source file with a mix of host code and device code, with locations of device code indicated therein. In at least one embodiment, the single-source file can be a.cu file including CUDA code or a.hip.cpp file including HIP code. Alternatively, in at least one embodiment, source code 3200 can include multiple source code files, rather than a single-source file, with host code and device code separated.

[0335] In at least one embodiment, compiler 3201 is configured to compile source code 3200 into host executable code 3202 for execution on a host and device executable code 3203 for execution on a device. In at least one embodiment, compiler 3201 performs operations including parsing source code 3200 into an abstract syntax tree (AST), performing optimizations, and generating executable code. In at least one embodiment in which source code 3200 includes a single-source file, compiler 3201 can separate device code from host code in such single-source file, compile device code and host code into device executable code 3203 and host executable code 3202, respectively, and link device executable code 3203 and host executable code 3202 together in a single file, as discussed below with respect to FIG. 4. Figure 33 discussed in more detail.

[0336] In at least one embodiment, host executable code 3202 and device executable code 3203 can be in any suitable format, such as binary code and / or IR code. In the case of CUDA, in at least one embodiment, host executable code 3202 can include native object code, while device executable code 3203 can include PTX intermediate representation code. In the case of ROCm, in at least one embodiment, both host executable code 3202 and device executable code 3203 can include object binary code.

[0337] Figure 33 is a more detailed illustration of compiling code to execute on one of the programming platforms of Figures 27-30 In at least one embodiment, compiler 3301 is configured to receive source code 3300, compile source code 3300, and output executable 3310. In at least one embodiment, source code 3300 is a single source file, such as a.cu file, a.hip.cpp file, or other format of file, that includes both host code and device code. In at least one embodiment, compiler 3301 can be, without limitation, an NVIDIA CUDA compiler (“NVCC”) for compiling CUDA code in a.cu file, or an HCC compiler for compiling HIP code in a.hip.cpp file.

[0338] In at least one embodiment, compiler 3301 includes a compiler front end 3302, a host compiler 3305, a device compiler 3306, and a linker 3309. In at least one embodiment, compiler front end 3302 is configured to separate device code 3304 from host code 3303 in source code 3300. In at least one embodiment, device code 3304 is compiled by device compiler 3306 into device executable code 3308, which can include binary code or IR code, as described. In at least one embodiment, host code 3303 is separately compiled by host compiler 3305 into host executable code 3307. In at least one embodiment, for NVCC, host compiler 3305 can be, without limitation, a general C / C++ compiler that outputs native object code, while device compiler 3306 can be, without limitation, a low-level virtual machine (“LLVM”) based compiler that forks the LLVM compiler infrastructure and outputs PTX code or binary code. In at least one embodiment, for HCC, both host compiler 3305 and device compiler 3306 can be, without limitation, LLVM based compilers that output target binary code.

[0339] In at least one embodiment, after source code 3300 is compiled into host executable code 3307 and device executable code 3308, linker 3309 links the host and device executable code 3307 and 3308 together in executable 3310. In at least one embodiment, native object code for the host and PTX or binary code for the device can be linked together in an executable and linkable format (“ELF”) file, which is a container format for storing object code.

[0340] Figure 34Conversion of source code prior to compilation is shown in accordance with at least one embodiment. In at least one embodiment, source code 3400 is passed through a conversion tool 3401 that converts source code 3400 to converted source code 3402. In at least one embodiment, a compiler 3403 is used to compile converted source code 3402 to host executable 3404 and device executable 3305, a process similar to that of compiler 3201 compiling source code 3200 to host executable 3202 and device executable 3203, as discussed above in conjunction with FIG. 3. In at least one embodiment, compiler 3403 can be, without limitation, a CUDA compiler, a DPC++ compiler, and / or a HIP compiler. Figure 32

[0341] In at least one embodiment, the conversion performed by conversion tool 3401 is used to port source code 3400 to execute in a different environment than originally intended. In at least one embodiment, conversion tool 3401 can include, without limitation, a HIP Transpiler, which is used to “hipify” CUDA code intended for a CUDA platform to HIP code that can be compiled and executed on a ROCm platform. In at least one embodiment, conversion of source code 3400 can include parsing source code 3400 and converting calls to APIs provided by one programming model (e.g., CUDA) to corresponding calls to APIs provided by another programming model (e.g., HIP), as discussed in more detail below in conjunction with FIG. 4. Returning to the example of porting CUDA code, in at least one embodiment, calls to CUDA runtime APIs, CUDA driver APIs, and / or CUDA libraries can be converted to corresponding HIP API calls. In at least one embodiment, automatic conversion performed by conversion tool 3401 can sometimes be incomplete, requiring additional human effort to fully port source code 3400. Figure 35A Figure 36 In at least one embodiment, conversion tool 3401 can be used to convert source code 3400 to a different programming model, such as from CUDA to DPC++ or from CUDA to HIP. In at least one embodiment, conversion tool 3401 can be used to convert source code 3400 to a different version of a programming model, such as from CUDA 11 to CUDA 12. In at least one embodiment, conversion tool 3401 can be used to convert source code 3400 to a different version of a programming model, such as from DPC++ 2020 to DPC++ 2022.

[0342] Configuring a GPU for general-purpose computing

[0343] The following figures set forth, without limitation, exemplary architectures for compiling and executing compute source code in accordance with at least one embodiment.

[0344] Figure 35A ​​A system 3500 configured to compile and execute CUDA source code 3510 using different types of processing units is shown, in accordance with at least one embodiment. In at least one embodiment, system 3500 includes, without limitation, CUDA source code 3510, a CUDA compiler 3550, host executable code 3570(1), host executable code 3570(2), CUDA device executable code 3584, a CPU 3590, a CUDA-enabled GPU 3594, a GPU 3592, a CUDA to HIP translation tool 3520, a HIP source code 3530, a HIP compiler driver 3540, an HCC 3560, and HCC device executable code 3582.

[0345] In at least one embodiment, CUDA source code 3510 is a collection of human-readable code in the CUDA programming language. In at least one embodiment, CUDA code is human-readable code in the CUDA programming language. In at least one embodiment, the CUDA programming language is an extension of the C++ programming language that includes, without limitation, mechanisms to define device code and to distinguish between device code and host code. In at least one embodiment, device code is source code that, after compilation, can be executed in parallel on a device. In at least one embodiment, a device can be a processor optimized for parallel instruction processing, such as CUDA-enabled GPU 3590, GPU 3592, or another GPGPU, etc. In at least one embodiment, host code is source code that, after compilation, can be executed on a host. In at least one embodiment, a host is a processor optimized for sequential instruction processing, such as CPU 3590.

[0346] In at least one embodiment, CUDA source code 3510 includes, without limitation, any number (including zero) of global functions 3512, any number (including zero) of device functions 3514, any number (including zero) of host functions 3516, and any number (including zero) of host / device functions 3518. In at least one embodiment, global functions 3512, device functions 3514, host functions 3516, and host / device functions 3518 can be intermixed in CUDA source code 3510. In at least one embodiment, each global function 3512 is executable on a device and is callable from a host. Thus, in at least one embodiment, one or more of global functions 3512 can serve as an entry point for a device. In at least one embodiment, each global function 3512 is a kernel. In at least one embodiment, and in a technique known as dynamic parallelism, one or more global functions 3512 define a kernel that is executable on a device and is callable from such a device. In at least one embodiment, a kernel is executed N times in parallel by N different threads on a device during execution (where N is any positive integer).

[0347] In at least one embodiment, each device function 3514 executes on a device and is callable only from such a device. In at least one embodiment, each host function 3516 executes on a host and is callable only from such a host. In at least one embodiment, each host / device function 3516 defines both a host version of a function executable on a host and callable only from such a host, as well as a device version of a function executable on a device and callable only from such a device.

[0348] In at least one embodiment, CUDA source code 3510 can also include, without limitation, any number of calls to any number of functions defined by CUDA runtime API 3502. In at least one embodiment, CUDA runtime API 3502 can include, without limitation, any number of functions that execute on a host for allocating and deallocating device memory, transferring data between host memory and device memory, managing systems with multiple devices, etc. In at least one embodiment, CUDA source code 3510 can also include any number of calls to any number of functions specified in any number of other CUDA APIs. In at least one embodiment, a CUDA API can be any API designed to be used by CUDA code. In at least one embodiment, CUDA APIs include, without limitation, CUDA runtime API 3502, CUDA driver APIs, APIs for any number of CUDA libraries, etc. In at least one embodiment and relative to CUDA runtime API 3502, CUDA driver APIs are lower-level APIs, but can provide more fine-grained control over a device. In at least one embodiment, examples of CUDA libraries include, without limitation, cuBLAS, cuFFT, cuRAND, cuDNN, etc.

[0349] In at least one embodiment, CUDA compiler 3550 compiles input CUDA code (e.g., CUDA source code 3510) to generate host executable code 3570(1) and CUDA device executable code 3584. In at least one embodiment, CUDA compiler 3550 is NVCC. In at least one embodiment, host executable code 3570(1) is a compiled version of host code included in input source code that is executable on CPU 3590. In at least one embodiment, CPU 3590 can be any processor optimized for sequential instruction processing.

[0350] In at least one embodiment, CUDA device executable code 3584 is a compiled version of device code included in input source code executable on a CUDA-enabled GPU 3594. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, IR code, such as PTX code, which is further compiled into binary code for a particular target device (e.g., CUDA-enabled GPU 3594) at runtime by a device driver. In at least one embodiment, CUDA-enabled GPU 3594 can be any processor optimized for parallel instruction processing and supporting CUDA. In at least one embodiment, CUDA-enabled GPU 3594 is developed by NVIDIA Corporation of Santa Clara, CA.

[0351] In at least one embodiment, CUDA to HIP translation tool 3520 is configured to translate CUDA source code 3510 into functionally similar HIP source code 3530. In at least one embodiment, HIP source code 3530 is a collection of human-readable code in the HIP programming language. In at least one embodiment, HIP code is human-readable code in the HIP programming language. In at least one embodiment, the HIP programming language is an extension of the C++ programming language that includes, without limitation, functionally similar versions of CUDA mechanisms for defining device code and distinguishing device code from host code. In at least one embodiment, the HIP programming language can include a subset of functionality of the CUDA programming language. In at least one embodiment, for example, the HIP programming language includes, without limitation, mechanisms to define global functions 3512, but such a HIP programming language can lack support for dynamic parallelism, and thus global functions 3512 defined in HIP code are only callable from a host.

[0352] In at least one embodiment, HIP source code 3530 includes, without limitation, any number (including zero) of global functions 3512, any number (including zero) of device functions 3514, any number (including zero) of host functions 3516, and any number (including zero) of host / device functions 3518. In at least one embodiment, HIP source code 3530 can also include any number of calls to any number of functions specified in a HIP runtime API 3532. In one embodiment, HIP runtime API 3532 includes, without limitation, functionally similar versions of a subset of functions included in CUDA runtime API 3502. In at least one embodiment, HIP source code 3530 can also include any number of calls to any number of functions specified in any number of other HIP APIs. In at least one embodiment, a HIP API can be any API designed for use with HIP code and / or ROCm. In at least one embodiment, a HIP API includes, without limitation, HIP runtime API 3532, a HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, and the like.

[0353] In at least one embodiment, CUDA to HIP translation tool 3520 translates each kernel call in CUDA code from CUDA syntax to HIP syntax, and translates any number of other CUDA calls in CUDA code to any number of other functionally similar HIP calls. In at least one embodiment, a CUDA call is a call to a function specified in a CUDA API, and a HIP call is a call to a function specified in a HIP API. In at least one embodiment, CUDA to HIP translation tool 3520 translates any number of calls to functions specified in CUDA runtime API 3502 to any number of calls to functions specified in HIP runtime API 3532.

[0354] In at least one embodiment, CUDA to HIP translation tool 3520 is a tool known as hipify-perl, which performs a text-based translation process. In at least one embodiment, CUDA to HIP translation tool 3520 is a tool known as hipify-clang, which performs a more complex and robust translation process relative to hipify-perl, involving parsing of CUDA code using clang (a compiler frontend), followed by translation of resulting symbols. In at least one embodiment, in addition to modifications performed by CUDA to HIP translation tool 3520, proper translation of CUDA code to HIP code can also require modifications (e.g., manual edits).

[0355] In at least one embodiment, HIP compiler driver 3540 is a front end that determines target device 3546 and then configures a compiler compatible with target device 3546 to compile HIP source code 3530. In at least one embodiment, target device 3546 is a processor optimized for parallel instruction processing. In at least one embodiment, HIP compiler driver 3540 can determine target device 3546 in any technically feasible manner.

[0356] In at least one embodiment, if target device 3546 is compatible with CUDA (e.g., a CUDA-enabled GPU 3594), then HIP compiler driver 3540 generates HIP / NVCC compilation commands 3542. In at least one embodiment and in conjunction with Figure 35B In more detail, HIP / NVCC compilation commands 3542 configure CUDA compiler 3550 to compile HIP source code 3530 using, without limitation, a HIP-to-CUDA translation header and a CUDA runtime library. In at least one embodiment and in response to HIP / NVCC compilation commands 3542, CUDA compiler 3550 generates host executable code 3570(1) and CUDA device executable code 3584.

[0357] In at least one embodiment, if target device 3546 is not compatible with CUDA, then HIP compiler driver 3540 generates HIP / HCC compilation commands 3544. In at least one embodiment and in conjunction with Figure 35C In more detail, HIP / HCC compilation commands 3544 configure HCC 3560 to compile HIP source code 3530 using a HCC header and a HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation commands 3544, HCC 3560 generates host executable code 3570(2) and HCC device executable code 3582. In at least one embodiment, HCC device executable code 3582 is a compiled version of device code contained in HIP source code 3530 that is executable on GPU 3592. In at least one embodiment, GPU 3592 can be any processor optimized for parallel instruction processing that is not compatible with CUDA and is compatible with HCC. In at least one embodiment, GPU 3592 is developed by AMD Corporation of Santa Clara, California. In at least one embodiment, GPU 3592 is a GPU 3592 that is not CUDA-enabled.

[0358] For illustrative purposes only, in Figure 35ACUDA source code 3510 to execute on CPU 3590 and a different device in at least one embodiment. In at least one embodiment, a direct CUDA flow compiles CUDA source code 3510 to execute on CPU 3590 and a CUDA-enabled GPU 3594 without converting CUDA source code 3510 to HIP source code 3530. In at least one embodiment, an indirect CUDA flow converts CUDA source code 3510 to HIP source code 3530 and then compiles HIP source code 3530 to execute on CPU 3590 and a CUDA-enabled GPU 3594. In at least one embodiment, a CUDA / HCC flow converts CUDA source code 3510 to HIP source code 3530 and then compiles HIP source code 3530 to execute on CPU 3590 and a GPU 3592.

[0359] A direct CUDA flow that can be implemented in at least one embodiment can be depicted by the dashed line and series of bubble annotations Al-A3. In at least one embodiment, and as shown by bubble annotation Al, a CUDA compiler 3550 receives CUDA source code 3510 and a CUDA compile command 3548 that configures CUDA compiler 3550 to compile CUDA source code 3510. In at least one embodiment, CUDA source code 3510 used in a direct CUDA flow is written in a CUDA programming language that is based on a programming language other than C++ (e.g., C, Fortran, Python, Java, etc.). In at least one embodiment, and in response to CUDA compile command 3548, CUDA compiler 3550 generates host executable code 3570(1) and CUDA device executable code 3584 (denoted by bubble annotation A2). In at least one embodiment and as shown by bubble annotation A3, host executable code 3570(1) and CUDA device executable code 3584 can be executed on CPU 3590 and CUDA-enabled GPU 3594, respectively. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0360] The indirect CUDA flow that can be implemented in at least one embodiment can be described by the dashed line and series of bubble annotations Bl- B6. In at least one embodiment and as shown by bubble annotation Bl, CUDA to HIP translation tool 3520 receives CUDA source code 3510. In at least one embodiment and as shown by bubble annotation B2, CUDA to HIP translation tool 3520 translates CUDA source code 3510 to HIP source code 3530. In at least one embodiment and as shown by bubble annotation B3, HIP compiler driver 3540 receives HIP source code 3530 and determines whether the target device 3546 has CUDA enabled.

[0361] In at least one embodiment and as shown by bubble annotation B4, HIP compiler driver 3540 generates HIP / NVCC compilation commands 3542 and sends both HIP / NVCC compilation commands 3542 and HIP source code 3530 to CUDA compiler 3550. In at least one embodiment and as described in greater detail below in connection with FIG. 36, HIP / NVCC compilation commands 3542 configure CUDA compiler 3550 to compile HIP source code 3530 using, without limitation, a HIP to CUDA translation header and a CUDA runtime library. Figure 35B In at least one embodiment and in response to HIP / NVCC compilation commands 3542, CUDA compiler 3550 generates host executable code 3570(1) and CUDA device executable code 3584 (represented by bubble annotation B5). In at least one embodiment and as shown by bubble annotation B6, host executable code 3570(1) and CUDA device executable code 3584 can be executed on CPU 3590 and CUDA-enabled GPU 3594, respectively. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0362] The CUDA / HCC flow that can be implemented in at least one embodiment can be described by the solid line and series of bubble annotations Cl- C6. In at least one embodiment and as shown by bubble annotation Cl, CUDA to HIP translation tool 3520 receives CUDA source code 3510. In at least one embodiment and as shown by bubble annotation C2, CUDA to HIP translation tool 3520 translates CUDA source code 3510 to HIP source code 3530. In at least one embodiment and as shown by bubble annotation C3, HIP compiler driver 3540 receives HIP source code 3530 and determines that the target device 3546 does not have CUDA enabled. In at least one embodiment and as shown by bubble annotation C4, HIP compiler driver 3540 generates HCC compilation commands 3544 and sends both HCC compilation commands 3544 and HIP source code 3530 to HCC compiler 3560. In at least one embodiment and as described in greater detail below in connection with FIG. 36, HCC compilation commands 3544 configure HCC compiler 3560 to compile HIP source code 3530 using, without limitation, a HIP to HCC translation header and a HCC runtime library.

[0363] In at least one embodiment, the HIP compiler driver 3540 generates HIP / HCC compilation commands 3544 and sends both the HIP / HCC compilation commands 3564 and the HIP source code 3530 to the HCC 3560 (indicated by bubble comment C4). In at least one embodiment and as in combination Figure 35C In more detail, HIP / HCC compilation command 3564 configures HCC 3560 to compile HIP source code 3530 using, but not limited to, the HCC header and HIP / HCC runtime library. In at least one embodiment and in response to HIP / HCC compilation command 3544, HCC 3560 generates host executable code 3570(2) and HCC device executable code 3582 (indicated by bubble comment C5). In at least one embodiment and as shown by bubble comment C6, host executable code 3570(2) and HCC device executable code 3582 can be executed on CPU 3590 and GPU 3592, respectively.

[0364] In at least one embodiment, after converting CUDA source code 3510 to HIP source code 3530, the HIP compiler driver 3540 can then be used to generate executable code for a CUDA-enabled GPU 3594 or GPU 3592 without re-executing CUDA to the HIP conversion tool 3520. In at least one embodiment, the CUDA to HIP conversion tool 3520 converts CUDA source code 3510 to HIP source code 3530 and then stores it in memory. In at least one embodiment, the HIP compiler driver 3540 then configures HCC 3560 to generate host executable code 3570(2) and HCC device executable code 3582 based on the HIP source code 3530. In at least one embodiment, the HIP compiler driver 3540 then configures CUDA compiler 3550 to generate host executable code 3570(1) and CUDA device executable code 3584 based on the stored HIP source code 3530.

[0365] Figure 35B The diagram illustrates a configuration, according to at least one embodiment, to compile and execute using a CPU 3590 and a CUDA-enabled GPU 3594. Figure 35A The system 3504 includes, but is not limited to, CUDA source code 3510, CUDA to HIP conversion tool 3520, HIP source code 3530, HIP compiler driver 3540, CUDA compiler 3550, host executable code 3570(1), CUDA device executable code 3584, CPU 3590 and CUDA-enabled GPU 3594.

[0366] In at least one embodiment and as previously described herein in connection with Figure 35A CUDA source code 3510 includes, without limitation, any number (including zero) of global functions 3512, any number (including zero) of device functions 3514, any number (including zero) of host functions 3516, and any number (including zero) of host / device functions 3518. In at least one embodiment, CUDA source code 3510 also includes, without limitation, any number of calls to any number of functions specified in any number of CUDA APIs.

[0367] In at least one embodiment, CUDA to HIP translation tool 3520 translates CUDA source code 3510 into HIP source code 3530. In at least one embodiment, CUDA to HIP translation tool 3520 translates each kernel call in CUDA source code 3510 from CUDA syntax to HIP syntax, and translates any number of other CUDA calls in CUDA source code 3510 to any number of other functionally similar HIP calls.

[0368] In at least one embodiment, HIP compiler driver 3540 determines that target device 3546 is CUDA-enabled and generates a HIP / NVCC compilation command 3542. In at least one embodiment, HIP compiler driver 3540 then configures CUDA compiler 3550 via HIP / NVCC compilation command 3542 to compile HIP source code 3530. In at least one embodiment, as part of configuring CUDA compiler 3550, HIP compiler driver 3540 provides access to a HIP to CUDA translation front end 3552. In at least one embodiment, HIP to CUDA translation front end 3552 converts any number of mechanisms (e.g., functions) specified in any number of HIP APIs to any number of mechanisms specified in any number of CUDA APIs. In at least one embodiment, CUDA compiler 3550 uses HIP to CUDA translation front end 3552 in conjunction with a CUDA runtime library 3554 corresponding to CUDA runtime API 3502 to generate host executable code 3570(1) and CUDA device executable code 3584. In at least one embodiment, host executable code 3570(1) and CUDA device executable code 3584 can then be executed on CPU 3590 and CUDA-enabled GPU 3594, respectively. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, binary code. In at least one embodiment, CUDA device executable code 3584 includes, without limitation, PTX code and is further compiled into binary code for a particular target device at runtime.

[0369] Figure 35C FIG. 35 shows a system 3506 configured to compile and execute CUDA source code 3510 using CPU 3590 and a GPU 3592 that is not CUDA-enabled, according to at least one embodiment. Figure 35A In at least one embodiment, system 3506 includes, without limitation, CUDA source code 3510, a CUDA to HIP translation tool 3520, HIP source code 3530, a HIP compiler driver 3540, HCC 3560, host executable code 3570(2), HCC device executable code 3582, CPU 3590, and GPU 3592.

[0370] In at least one embodiment, and as previously discussed herein in conjunction with Figure 35AAs described, the CUDA source code 3510 includes, without limitation, any number (including zero) of global functions 3512, any number (including zero) of device functions 3514, any number (including zero) of host functions 3516, and any number (including zero) of host / device functions 3518. In at least one embodiment, the CUDA source code 3510 also includes, without limitation, any number of calls to any number of functions specified in any number of CUDA APIs.

[0371] In at least one embodiment, the CUDA to HIP translation tool 3520 translates the CUDA source code 3510 into HIP source code 3530. In at least one embodiment, the CUDA to HIP translation tool 3520 translates each kernel call in the CUDA source code 3510 from CUDA syntax to HIP syntax, and translates any number of other CUDA calls in the source code 3510 to any number of other functionally similar HIP calls.

[0372] In at least one embodiment, the HIP compiler driver 3540 then determines that the target device 3546 is not CUDA-enabled, and generates a HIP / HCC compilation command 3544. In at least one embodiment, the HIP compiler driver 3540 then configures the HCC 3560 to perform the HIP / HCC compilation command 3544, thereby compiling the HIP source code 3530. In at least one embodiment, the HIP / HCC compilation command 3544 configures the HCC 3560 to use, without limitation, the HIP / HCC runtime library 3558 and the HCC header 3556 to generate host executable code 3570(2) and HCC device executable code 3582. In at least one embodiment, the HIP / HCC runtime library 3558 corresponds to the HIP runtime API 3532. In at least one embodiment, the HCC header 3556 includes, without limitation, any number and type of interoperability mechanisms for HIP and HCC. In at least one embodiment, the host executable code 3570(2) and the HCC device executable code 3582 can be executed on the CPU 3590 and the GPU 3592, respectively.

[0373] In at least one embodiment, the set 112 (see Figure 1 and Figure 3 ) and / or the group 114 (see Figure 1 and Figure 3 ) can include one or more of the CPU 3590, one or more of the GPU 3592, and / or one or more of the CUDA-enabled GPU 3594.

[0374] Figure 36 FIG. 13 illustrates a system including a CPU 3590, a GPU 3592, and a CUDA-enabled GPU 3594, in accordance with at least one embodiment. Figure 35Can exemplary kernel converted by the CUDA to HIP translation tool 3520. In at least one embodiment, CUDA source code 3510 divides the overall problem that a given kernel is designed to solve into relatively coarse sub-problems that can be solved independently using thread blocks. In at least one embodiment, each thread block includes, without limitation, any number of threads. In at least one embodiment, each sub-problem is divided into relatively fine pieces that can be solved in parallel by threads in a thread block in collaboration. In at least one embodiment, threads within a thread block can collaborate by sharing data through shared memory and by synchronizing execution to coordinate memory access.

[0375] In at least one embodiment, CUDA source code 3510 organizes thread blocks associated with a given kernel into a one-, two-, or three-dimensional grid of thread blocks. In at least one embodiment, each thread block includes, without limitation, any number of threads, and a grid includes, without limitation, any number of thread blocks.

[0376] In at least one embodiment, a kernel is a function in device code that is defined using a “__global__” declaration specifier. In at least one embodiment, a CUDA kernel launch syntax 3610 is used to specify the size of a grid of kernels to execute for a given kernel invocation, as well as an associated stream. In at least one embodiment, CUDA kernel launch syntax 3610 is specified as “KernelName<<<GridSize, BlockSize, SharedMemorySize, Stream>>>(KernelArguments);”. In at least one embodiment, an execution configuration syntax is the “<<<...>>>” construct, which is inserted between a kernel name (“KernelName”) and a parenthetical list of kernel arguments (“KernelArguments”). In at least one embodiment, CUDA kernel launch syntax 3610 includes, without limitation, a CUDA launch function syntax instead of an execution configuration syntax.

[0377] In at least one embodiment, “GridSize” is of type dim3 and specifies the size and dimensions of a grid. In at least one embodiment, type dim3 is a CUDA-defined structure that includes, without limitation, unsigned integers x, y, and z. In at least one embodiment, if z is not specified, it defaults to 1. In at least one embodiment, if y is not specified, it defaults to 1. In at least one embodiment, a number of thread blocks in a grid is equal to a product of GridSize.x, GridSize.y, and GridSize.z. In at least one embodiment, “BlockSize” is of type dim3 and specifies the size and dimensions of each thread block. In at least one embodiment, a number of threads per thread block is equal to a product of BlockSize.x, BlockSize.y, and BlockSize.z. In at least one embodiment, each thread executing a kernel is given a unique thread ID that is accessible within a kernel via a built-in variable, such as “threadIdx”.

[0378] In at least one embodiment, with respect to CUDA kernel launch syntax 3610, “SharedMemorySize” is an optional parameter that specifies a number of bytes of dynamic allocation of shared memory per thread block for a given kernel call, in addition to statically allocated memory. In at least one embodiment and with respect to CUDA kernel launch syntax 3610, SharedMemorySize defaults to zero. In at least one embodiment and with respect to CUDA kernel launch syntax 3610, “stream” is an optional parameter that specifies an associated stream and defaults to zero to specify the default stream. In at least one embodiment, a stream is a sequence of commands (which can be issued by different host threads) that are executed in-order. In at least one embodiment, different streams can execute commands out-of-order or simultaneously with respect to each other.

[0379] In at least one embodiment, CUDA source code 3510 includes, without limitation, a kernel definition and a host function for an exemplary kernel “MatAdd”. In at least one embodiment, a host function is host code that executes on a host and includes, without limitation, a kernel call that causes kernel MatAdd to execute on a device. In at least one embodiment, as shown, kernel MatAdd adds two matrices A and B of size NxN, where N is a positive integer, and stores a result in matrix C. In at least one embodiment, a host function defines a threadsPerBlock variable to be 16x 16 and a numBlocks variable to be N / 16 x N / 16. In at least one embodiment, host function then specifies a kernel call “MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);”. In at least one embodiment, and in accordance with CUDA kernel launch syntax 3610, kernel MatAdd is executed using a thread block grid of size N / 16 x N / 16, where each thread block is of size 16 x 16. In at least one embodiment, each thread block includes 256 threads, a grid with enough blocks is created so that each matrix element has one thread, and each thread in that grid executes kernel MatAdd to perform one pairwise addition.

[0380] In at least one embodiment, while converting CUDA source code 3510 to HIP source code 3530, CUDA to HIP translation tool 3520 converts each kernel call in CUDA source code 3510 from CUDA kernel launch syntax 3610 to HIP kernel launch syntax 3620 and converts any number of other CUDA calls in source code 3510 to any number of other functionally similar HIP calls. In at least one embodiment, HIP kernel launch syntax 3620 is specified as “hipLaunchKernelGGL(KernelName, GridSize, BlockSize, SharedMemorySize, Stream, KernelArguments);”. In at least one embodiment, each of KernelName, GridSize, BlockSize, ShareMemorySize, Stream, and KernelArguments has the same meaning in HIP kernel launch syntax 3620 as in CUDA kernel launch syntax 3610 (described previously herein). In at least one embodiment, parameters SharedMemorySize and Stream are required in HIP kernel launch syntax 3620, but are optional in CUDA kernel launch syntax 3610.

[0381] In at least one embodiment, in addition to the kernel call to cause kernel MatAdd to execute on a device, Figure 36 The portion of HIP source code 3530 depicted in FIG. 33A is the same as the portion of CUDA source code 3510 depicted in FIG. 31 A. In at least one embodiment, kernel MatAdd is defined in HIP source code 3530 with the same “__global__” declaration specifier as kernel MatAdd is defined in CUDA source code 3510. In at least one embodiment, the kernel call in HIP source code 3530 is “hipLaunchKernelGGL(MatAdd, numBlocks, threadsPerBlock, 0, 0, A, B, C);”, while the corresponding kernel call in CUDA source code 3510 is “MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);”. Figure 36

[0382] Figure 37 FIG. 34 illustrates a portion of CUDA source code 3510 according to at least one embodiment. Figure 35C FIG. 34 illustrates a portion of CUDA source code 3510 according to at least one embodiment.In at least one embodiment, GPU 3592 is developed by AMD Corporation of Santa Clara, CA. In at least one embodiment, GPU 3592 can be configured to perform computational operations in a highly parallel manner. In at least one embodiment, GPU 3592 is configured to perform graphics pipeline operations, such as draw commands, pixel operations, geometry calculations, and other operations associated with rendering images to a display. In at least one embodiment, GPU 3592 is configured to perform operations that are not graphics-related. In at least one embodiment, GPU 3592 is configured to perform both graphics-related operations and operations that are not graphics-related. In at least one embodiment, GPU 3592 can be configured to execute device code included in HIP source code 3530.

[0383] In at least one embodiment, GPU 3592 includes, without limitation, any number of programmable processing units 3720, command processor 3710, L2 cache 3722, memory controllers 3770, DMA engine 3780(1), system memory controller 3782, DMA engine 3780(2), and GPU controller 3784. In at least one embodiment, each programmable processing unit 3720 includes, without limitation, workload manager 3730 and any number of compute units 3740. In at least one embodiment, command processor 3710 reads commands from one or more command queues (not shown) and distributes commands to workload managers 3730. In at least one embodiment, for each programmable processing unit 3720, associated workload manager 3730 distributes work to compute units 3740 included in programmable processing unit 3720. In at least one embodiment, each compute unit 3740 can execute any number of thread blocks, but each thread block executes on a single compute unit 3740. In at least one embodiment, a workgroup is a thread block.

[0384] In at least one embodiment, each compute unit 3740 includes, without limitation, any number of SIMD units 3750 and shared memory 3760. In at least one embodiment, each SIMD unit 3750 implements a SIMD architecture and is configured to execute operations in parallel. In at least one embodiment, each SIMD unit 3750 includes, without limitation, a vector ALU 3752 and a vector register file 3754. In at least one embodiment, each SIMD unit 3750 executes a different thread bundle. In at least one embodiment, a thread bundle is a group of threads (e.g., 16 threads), where each thread in a thread bundle belongs to a single thread block and is configured to process a different set of data based on a single instruction set. In at least one embodiment, one or more threads in a thread bundle can be disabled using predication. In at least one embodiment, a lane is a thread. In at least one embodiment, a work item is a thread. In at least one embodiment, a wavefront is a thread bundle. In at least one embodiment, different wavefronts in a thread block can be synchronized together and communicate via shared memory 3760.

[0385] In at least one embodiment, programmable processing units 3720 are referred to as “shader processing units” or “shaders.” In at least one embodiment, each programmable processing unit 3720 includes, without limitation, any number of specialized graphics hardware in addition to compute units 3740. In at least one embodiment, each programmable processing unit 3720 includes, without limitation, any number (including zero) of geometry processors, any number (including zero) of rasterizers, any number (including zero) of render back-ends, workload manager 3730, and any number of compute units 3740.

[0386] In at least one embodiment, compute units 3740 share L2 cache 3722. In at least one embodiment, L2 cache 3722 is partitioned. In at least one embodiment, all compute units 3740 in GPU 3592 have access to GPU memory 3790. In at least one embodiment, memory controllers 3770 and system memory controllers 3782 facilitate data transfers between GPU 3592 and a host, and DMA engines 3780(1) enable asynchronous memory transfers between GPU 3592 and the host. In at least one embodiment, memory controllers 3770 and GPU controllers 3784 facilitate data transfers between GPU 3592 and other GPUs 3592, and DMA engines 3780(2) enable asynchronous memory transfers between GPU 3592 and other GPUs 3592.

[0387] In at least one embodiment, GPU 3592 includes, without limitation, any number and type of system interconnects that facilitate data and control transmissions between any number and type of directly or indirectly linked components within or external to GPU 3592. In at least one embodiment, GPU 3592 includes, without limitation, any number and type of I / O interfaces (e.g., PCIe) that are coupled to any number and type of peripheral devices. In at least one embodiment, GPU 3592 can include, without limitation, any number (including zero) of display engines and any number (including zero) of multimedia engines. In at least one embodiment, GPU 3592 implements a memory subsystem that includes, without limitation, any number and type of memory controllers (e.g., memory controllers 3770 and system memory controllers 3782) and memory devices that are dedicated to one component or shared between multiple components (e.g., shared memory 3760). In at least one embodiment, GPU 3592 implements a cache subsystem that includes, without limitation, one or more cache memories (e.g., L2 cache 3722), each of which can be private or shared between any number of components (e.g., SIMD units 3750, compute units 3740, and programmable processing units 3720).

[0388] Figure 38 FIG. 38 illustrates how threads of an exemplary CUDA grid 3820 are mapped to Figure 37different compute units 3740. In at least one embodiment, and for purposes of illustration only, grid 3820 has a GridSize of BX by BY by 1 and a BlockSize of TX by TY by 1. Thus, in at least one embodiment, grid 3820 includes, without limitation, (BX*BY) thread blocks 3830, each thread block 3830 including, without limitation, (TX*TY) threads 3840. Threads 3840 are depicted in Figure 38

[0389] In at least one embodiment, grid 3820 is mapped to programmable processing unit 3720(1), which includes, without limitation, compute units 3740(1)-3740(C). In at least one embodiment and as shown, (BJ*BY) thread blocks 3830 are mapped to compute unit 3740(1), and the remaining thread blocks 3830 are mapped to compute unit 3740(2). In at least one embodiment, each thread block 3830 can include, without limitation, any number of thread warps, and each thread warp is mapped to Figure 37 different SIMD units 3750.

[0390] In at least one embodiment, thread warps in a given thread block 3830 can synchronize together and communicate through shared memory 3760 included in associated compute unit 3740. For example and in at least one embodiment, thread warps in thread block 3830(BJ,1) can synchronize together and communicate through shared memory 3760(1). For example and in at least one embodiment, thread warps in thread block 3830(BJ+1,1) can synchronize together and communicate through shared memory 3760(2).

[0391] Figure 39 ​How existing CUDA code can be migrated to data parallel C++ code is shown in accordance with at least one embodiment. Data Parallel C++ (DPC++) can refer to an open, standards-based alternative to single-architecture proprietary languages that allows developers to reuse code across hardware targets (CPUs and accelerators, such as GPUs and FPGAs), and also perform custom tuning for specific accelerators. DPC++ uses similar and / or the same C and C++ constructs that developers can be familiar with in accordance with ISOC++. DPC++ incorporates the SYCL standard by The Khronos Group to support data parallelism and heterogeneous programming. SYCL refers to an abstraction layer that is cross-platform, which builds on the underlying concepts, portability, and efficiency of OpenCL, which enables code for heterogeneous processors to be written in “single-source” style using standard C++. SYCL can enable single-source development, where C++ template functions can contain both host and device code to build complex algorithms that use OpenCL acceleration, and then reuse them throughout the source code for different types of data.

[0392] In at least one embodiment, a DPC++ compiler is used to compile DPC++ source code that can be deployed across a variety of hardware targets. In at least one embodiment, a DPC++ compiler is used to generate DPC++ applications that can be deployed across a variety of hardware targets, and DPC++ compatibility tools can be used to migrate CUDA applications to multi-platform programs in DPC++. In at least one embodiment, a DPC++ foundation toolset includes: a DPC++ compiler to deploy applications across a variety of hardware targets; DPC++ libraries to improve productivity and performance for CPUs, GPUs, and FPGAs; DPC++ compatibility tools to migrate CUDA applications to multi-platform applications; and any suitable combination thereof.

[0393] In at least one embodiment, a DPC++ programming model is used to simplify one or more aspects related to programming CPUs and accelerators by expressing parallelism with a programming language called Data Parallel C++ using modern C++ features. DPC++ programming language can be used for code reuse for both hosts (e.g., CPUs) and accelerators (e.g., GPUs or FPGAs) using a single-source language, and clearly communicates execution and memory dependencies. Mapping within DPC++ code can be used to convert an application to run on hardware or a set of hardware devices that best accelerates a workload. Hosts can be used to simplify development and debugging of device code even on platforms where no accelerators are available.

[0394] In at least one embodiment, CUDA source code 3900 is provided as input to a DPC++ compatibility tool 3902 to generate human readable DPC++ 3904. In at least one embodiment, human readable DPC++ 3904 includes inline comments generated by DPC++ compatibility tool 3902 that guide a developer how and / or where to modify DPC++ code to complete coding and tuning to desired performance 3906, resulting in DPC++ source code 3908.

[0395] In at least one embodiment, CUDA source code 3900 is or includes a collection of human readable source code in a CUDA programming language. In at least one embodiment, CUDA source code 3900 is human readable source code in CUDA programming language. In at least one embodiment, CUDA programming language is an extension of C++ programming language that includes, without limitation, mechanisms to define device code and to distinguish between device code and host code. In at least one embodiment, device code is source code that, when compiled, is executable on a device (e.g., GPU or FPGA) and can include one or more parallel workstreams that are executable on one or more processor cores of a device. In at least one embodiment, a device can be a processor that is optimized for parallel instruction processing, such as a CUDA-enabled GPU, GPU, or another GPGPU, etc. In at least one embodiment, host code is source code that, when compiled, is executable on a host. In at least one embodiment, some or all of host code and device code can be executed in parallel across a CPU and GPU / FPGA. In at least one embodiment, a host is a processor that is optimized for sequential instruction processing, such as a CPU. In conjunction with Figure 39 Described CUDA source code 3900 can be consistent with that discussed elsewhere in this document.

[0396] In at least one embodiment, DPC++ compatibility tool 3902 refers to an executable tool, program, application, or any other suitable type of tool for facilitating migration of CUDA source code 3900 to DPC++ source code 3908. In at least one embodiment, DPC++ compatibility tool 3902 is a command-line based code migration tool that is available as part of a DPC++ toolchain for porting existing CUDA sources to DPC++. In at least one embodiment, DPC++ compatibility tool 3902 converts some or all of source code of a CUDA application from CUDA to DPC++ and generates a result file written at least partially in DPC++, referred to as human-readable DPC++ 3904. In at least one embodiment, human-readable DPC++ 3904 includes annotations generated by DPC++ compatibility tool 3902 to indicate places where user intervention can be necessary. In at least one embodiment, user intervention is necessary when CUDA source code 3900 calls a CUDA API that does not have a similar DPC++ API; other examples of where user intervention is needed are discussed in more detail later.

[0397] In at least one embodiment, a workflow for migrating CUDA source code 3900 (e.g., an application or portions thereof) includes creating one or more compilation database files; using DPC++ compatibility tool 3902 to migrate CUDA to DPC++; completing the migration and verifying correctness, resulting in DPC++ source code 3908; and compiling DPC++ source code 3908 using a DPC++ compiler to generate a DPC++ application. In at least one embodiment, the compatibility tool provides a utility that intercepts commands used when Makefiles are executed and stores them in a compilation database file. In at least one embodiment, the file is stored in JSON format. In at least one embodiment, intercepting build commands converts Makefile commands to DPC compatibility commands.

[0398] In at least one embodiment, intercept-build is a utility script that intercepts the build process to capture compilation options, macro definitions, and include paths, and writes this data to a compilation database file. In at least one embodiment, the compilation database file is a JSON file. In at least one embodiment, DPC++ compatibility tool 3902 parses the compilation database and applies the options when migrating input sources. In at least one embodiment, use of intercept-build is optional, but is highly recommended for Make or CMake based environments. In at least one embodiment, a migration database includes commands, directories, and files: commands can include necessary compilation flags; directories can include paths to header files; and files can include paths to CUDA files.

[0399] In at least one embodiment, DPC++ compatibility tool 3902 migrates CUDA code (e.g., applications) written in CUDA to DPC++ by generating DPC++ wherever possible. In at least one embodiment, DPC++ compatibility tool 3902 is available as part of a toolchain. In at least one embodiment, DPC++ toolchain includes an intercept-build tool. In at least one embodiment, intercept-build tool creates a compilation database that captures compilation commands to migrate CUDA files. In at least one embodiment, DPC++ compatibility tool 3902 uses the compilation database generated by intercept-build tool to migrate CUDA code to DPC++. In at least one embodiment, non-CUDA C++ code and files are migrated as-is. In at least one embodiment, DPC++ compatibility tool 3902 generates human-readable DPC++ 3904, which can be DPC++ code as generated by DPC++ compatibility tool 3902, cannot be compiled by a DPC++ compiler and requires additional pipeline to validate portions of code that were not migrated correctly, and can involve manual intervention, e.g., by a developer. In at least one embodiment, DPC++ compatibility tool 3902 provides hints or tools embedded in code to help a developer manually migrate additional code that could not be automatically migrated. In at least one embodiment, migration is a one-time activity for a source file, project, or application.

[0400] In at least one embodiment, DPC++ compatibility tool 39002 is able to successfully migrate all portions of CUDA code to DPC++, and can simply exist as an optional step for manual verification and adjustment of performance of generated DPC++ source code. In at least one embodiment, DPC++ compatibility tool 3902 directly generates DPC++ source code 3908 that is compiled by a DPC++ compiler without requiring or utilizing human intervention to modify DPC++ code generated by DPC++ compatibility tool 3902. In at least one embodiment, DPC++ compatibility tool generates DPC++ code t...

Claims

1. An apparatus comprising: One or more circuits, which include circuitry for the following operations: Retrieve the transaction addressed to the target group; An entry associated with the target group is selected from at least one routing data structure, the entry identifying an indication set, each set transmitting the transaction to the target group, wherein a first indication in the first set instructs a first switch among a plurality of switches to transmit the transaction on the output of the first switch, and a second indication in the first set instructs a second switch among the plurality of switches to transmit the transaction on the output of the second switch; Select the selected set from the indicated set; and The transaction is transmitted to the target group according to the selected set.

2. The device of claim 1, wherein before the circuit selects the selected set, the circuit is configured to identify the specific entry using information included in the transaction.

3. The device according to claim 2, wherein the information includes an identifier of the target group.

4. The device according to claim 2, wherein the transaction is a first transaction. The information mentioned is the first information, and The circuit is used to acquire a second transaction including second information, and to use the second information to identify the specific entry and select the selected set.

5. The device of claim 1, wherein the circuitry is configured to select the selected set using a hash function, a static method, or a random number generator.

6. The device of claim 1, wherein each of the indicator sets is identified by an indicator string or array.

7. The device of claim 1, wherein the one or more circuits comprise: At least one processor; as well as A memory for storing instructions executable by the at least one processor, which, when executed by the at least one processor, causes the at least one processor to select the selected set and causes the transaction to be transferred to the target group according to the selected set.

8. The device of claim 7, wherein the one or more circuits comprise a plurality of ports, and The indication of the selected set causes the circuit to transmit the transaction to the target group through at least a portion of the plurality of ports.

9. The device of claim 8, wherein the memory includes the routing data structure that associates the target group with the indication set, and The routing data structure includes a string or array that identifies the part.

10. The device of claim 7, wherein the instructions, when executed by the at least one processor, cause the at least one processor to select the selected set based on information included in the transaction, the information including an identifier of the source device and an identifier of the group.

11. The device of claim 10, wherein the information includes an identifier of a memory location.

12. The device of claim 1, wherein at least two of the indicators in the set of indicators comprise different numbers of indicators.

13. The device of claim 1, wherein at least two specific indications in the selected set each comprise an indication set, and The instruction set of the first of the two specific instructions includes a different number of instructions than the instruction set of the second of the two specific instructions.

14. The device of claim 1, wherein at least a portion of the indications in the selected set specifies different routing characteristics.

15. The device of claim 1, wherein at least a portion of the indications in the selected set specifies that the transaction is to be transferred from the first virtual channel to a different second virtual channel.

16. The device of claim 1, wherein the indication set is associated with the target group via a first entry of a first data structure of the at least one routing data structure and a second entry of a second data structure of the at least one routing data structure, and The second entry is shared by multiple target groups.

17. The device of claim 16, wherein the indicator set comprises a first set stored in the first entry of the first data structure and a second set stored in the second entry of the second data structure, and One set in the first set and one set in the second set include different numbers of indicators.

18. The device of claim 16, wherein the indicator set comprises a first set stored in the first entry of the first data structure and a second set stored in the second entry of the second data structure, and Each of the first sets has an indication of the maximum number of first maximum values. Each of the second sets has an indication of the maximum number of second maximum values, and The indications of the first maximum value number and the indications of the second maximum value number are different from each other.

19. The device of claim 18, wherein the second entry of the second data structure includes the selected set.

20. The device of claim 1, wherein the transaction is obtained from the source device, and The circuitry is configured to receive multiple responses to the transaction from the target group, combine the multiple responses to obtain a combined response, and transmit the combined response to the source device.

21. The device of claim 20, wherein the circuitry is configured to determine the order based on the received order information included in the plurality of responses, and The circuit is used to combine the plurality of responses to obtain the combined response by performing a reduction operation on the plurality of responses in the order stated therein.

22. The apparatus of claim 21, wherein the circuitry is configured to insert transmission order information into the transaction before transmitting the transaction to the target group, and the target group copies the transmission order information as the received order information into the response before the target group sends the response.

23. The device of claim 22, wherein the selected set comprises two or more indications having an indication order. The transmission sequence information identifies the indicated sequence, and The order determined by the circuit based on the received sequence information is the indicated order.

24. The apparatus of claim 21, wherein the reduction operation comprises adding data included in the plurality of responses in the order stated therein.

25. The apparatus of claim 20, wherein the circuitry is configured to combine the plurality of responses by performing a reduction operation on the plurality of responses to obtain the combined response, and The reduction operations include addition, minimum value, maximum value, AND, OR, or XOR.

26. The device of claim 1, wherein the circuit comprises: A first internal switch, which is connected to a first outbound switch via a first main path; as well as A second internal switch is connected to a second outbound switch via a second primary path. A first internal switch is connected to the second outbound switch via a first backup path. The second internal switch is connected to the first outbound switch via a second backup path. The selected set includes a third indication and a fourth indication. The third indication indicates that the transaction will be transmitted as a first copy by the first internal switch to the first outbound switch via the first primary path. The fourth indication indicates that the transaction will be transmitted as a second copy by the second internal switch to the first outbound switch via the second backup path. The first outbound switch transmits the first copy and the second copy to different first and second targets in the target group, respectively.

27. A method comprising: Retrieve the transaction addressed to the target group; An entry associated with the target group is selected from at least one routing data structure, the entry identifying an indication set, each set transmitting the transaction to the target group, wherein a first indication in a first set instructs a first switch among a plurality of switches to transmit the transaction on the output of the first switch, and a second indication in the first set instructs a second switch among the plurality of switches to transmit the transaction on the output of the second switch; Select the selected set from the indicated set; and The transaction is transmitted to the target group according to the selected set.

28. The method of claim 27, wherein the transaction is a first transaction including first information, and the method further comprises: The first information is obtained, and the entries and the selected set are selected based on the first information; Acquire a second transaction that includes second information; Based on the second information, the entry and the selected set are selected; and The second transaction is transmitted to the target group according to the selected set.

29. The method of claim 27, wherein the transaction is transmitted to the target group via a plurality of transmission resources, and The selected set includes at least two indications that define multiple rounds during which the transaction is to be sent multiple times via at least one of the multiple transport resources.

30. The method of claim 27, wherein the selected set is selected by performing an operation on information included in the transaction, the operation including a hash function, a static method, or generating a value using a random number generator.

31. The method of claim 30, wherein the information includes an identifier of the source device and an identifier of the group.

32. The method of claim 31, wherein the information includes an identifier of a memory location.

33. The method of claim 27, wherein the at least one routing data structure comprises a first routing data structure and a second routing data structure. The first routing data structure includes the entry and a first part of the indicator set. The second routing data structure includes a second portion of the indication set, and The entry includes a pointer to the second part of the indicator set.

34. The method of claim 33, wherein the selected set is a set in the second portion of the indication set.

35. The method of claim 33, wherein the second portion of the indication set is shared by a plurality of target groups.

36. The method of claim 27, wherein the transaction is obtained from the source device, and the method further comprises: Receive multiple responses to the transaction from the target group; The order is determined based on the order information received in the plurality of responses; The multiple responses are combined in the specified order to obtain a combined response; as well as The combined response is transmitted to the source device.

37. The method of claim 36, further comprising: Before transmitting the transaction to the target group, the transmission order information is inserted into the transaction. Before the target group sends the response, the target group copies the transmission order information as the received order information into the response.

38. The method of claim 37, wherein the selected set comprises two or more indications having an indicative order. The transmission sequence information identifies the indicated sequence, and The order determined based on the received order information is the indicated order.

39. The method of claim 36, wherein combining the plurality of responses in the order comprises performing a reduction operation on data included in the plurality of responses in the order to obtain the combined response.

40. The method of claim 39, wherein the reduction operation comprises adding the data included in the plurality of responses in the order to obtain the combined response.

41. The method of claim 27, wherein the transaction is obtained from the source device, and the method further comprises: Receive multiple responses to the transaction from the target group; A combined response is obtained by performing a reduction operation on the data included in the plurality of responses; as well as The combined response is transmitted to the source device.

42. The method according to claim 41, wherein the reduction operation includes an addition operation, a minimum value operation, a maximum value operation, an AND operation, an OR operation, or an XOR operation.

43. The method of claim 27, used in conjunction with a first internal switch connected to a first outbound switch via a first primary path and a second internal switch connected to a second outbound switch via a second primary path, the first internal switch connected to the second outbound switch via a first alternate path and the second internal switch connected to the first outbound switch via a second alternate path, wherein transmitting the transaction to the target group according to the selected set comprises: The first internal switch transmits the transaction to the first outbound switch via the first main path according to the third instruction in the selected set. as well as The second internal switch transmits the transaction to the first outbound switch via the second backup path according to the fourth instruction in the selected set. After receiving the transaction from the first internal switch and the second internal switch, the first outbound switch transmits the transaction to the first target and the second target in the target group.

44. One or more non-transitory processor-readable media storing at least one routing data structure and instructions, said instructions being executable by at least one processor and, when executed by said at least one processor, causing said at least one processor to: An entry is selected from the at least one routing data structure based on a group identifier transmitted in a transaction, the entry identifying an alternate set of indications for transmitting the transaction to a target group, wherein a first indication in the first set instructs a first switch among a plurality of switches to transmit the transaction on the output of the first switch, and a second indication in the first set instructs a second switch among the plurality of switches to transmit the transaction on the output of the second switch. Select the selected set from the alternative instruction set of the transaction; and The transaction is transmitted to the target group according to the selected set.

Images