Instructions for copying stack operands

Abstract

This disclosure relates to instructions for copying stack operands, and more specifically to apparatus, systems, and techniques for translating instructions using operand stack structures into instructions using registers and / or memory to store operands. In at least one embodiment, if one or more branch instructions use one or more instruction operands stored in one or more stack structures, then one or more instruction operand stack structures are copied.

Description

Technical Field

[0001] At least one embodiment relates to executing stack-based program instructions to translate graphics processing unit (GPU) program instructions. For example, at least one embodiment relates to a processor or circuit for translating program code instructions that access one or more operands on the stack into program instructions that access one or more operands in a register. Background Technology

[0002] Programs written in virtual machine (VM) languages ​​(such as Python) are translated into VM instructions by a compiler (such as a Just-In-Time (JIT) compiler), which are then executed by the VM. This processing model differs from that used by processors such as central processing units (CPUs) and graphics processing units (GPUs), where operands are provided to the CPU or GPU instructions along with the corresponding instructions, rather than being stored on a stack. Furthermore, operands for CPU or GPU instructions can include one or more registers, memory addresses, and / or constants. Running stack-based VM programs on a CPU or GPU (methods involving translating VM instructions into CPU instructions) is computationally time-consuming and resource-intensive and can be improved. Attached Figure Description

[0003] Figure 1 A system for translating a program written in a VM-based language into instructions that use operands in registers and / or memory, according to at least one embodiment, is shown.

[0004] Figure 2 An exemplary VM-to-GPU instruction translator according to at least one embodiment is shown;

[0005] Figure 3 Exemplary VM-oriented instructions and corresponding GPU-oriented instructions according to at least one embodiment are shown, wherein VM-oriented instructions access operands on the stack, while GPU-oriented instructions access operands in registers and / or memory.

[0006] Figures 4 to 11 An exemplary translation of VM-oriented instructions into GPU-oriented instructions is shown according to at least one embodiment;

[0007] Figure 12 A method for translating a VM program into a GPU program according to at least one embodiment is illustrated;

[0008] Figure 13 An example of a system including a processor and modules according to at least one embodiment is shown;

[0009] Figure 14 An example data center system according to at least one embodiment is shown;

[0010] Figure 15 A system-on-a-chip (SOC) according to at least one embodiment is shown;

[0011] Figure 16A A parallel processor according to at least one embodiment is shown;

[0012] Figure 16B A processing cluster according to at least one embodiment is shown;

[0013] Figure 16C A graphics multiprocessor according to at least one embodiment is shown;

[0014] Figure 17 An accelerator processor according to at least one embodiment is shown;

[0015] Figure 18A A central processing unit according to at least one embodiment is shown;

[0016] Figure 18B The illustration shows an embodiment according to at least one of the embodiments. Figure 18A The core of the central processing unit;

[0017] Figure 19 Another accelerator processor according to at least one embodiment is shown;

[0018] Figure 20 A neuromorphic processor according to at least one embodiment is shown;

[0019] Figure 21 A supercomputer according to at least one embodiment is shown;

[0020] Figure 22 Another accelerator processor according to at least one embodiment is shown;

[0021] Figure 23 Another processor according to at least one embodiment is shown;

[0022] Figure 24 Another accelerator processor according to at least one embodiment is shown;

[0023] Figure 25 A tensor processing unit according to at least one embodiment is shown;

[0024] Figure 26 A RISC-V compatible processor according to at least one embodiment is shown;

[0025] Figure 27A and Figure 27BA language processing unit according to at least one embodiment is shown;

[0026] Figure 28 A software stack of a programming platform according to at least one embodiment is shown;

[0027] Figure 29 Software supported by a programming platform according to at least one embodiment is shown;

[0028] Figure 30 A method for using at least one embodiment is shown. Figure 23 Compiled code executed on the programming platform;

[0029] Figure 31 An example of an autonomous vehicle and its system architecture according to at least one embodiment is shown;

[0030] Figure 32A The inference and / or training logic according to at least one embodiment is illustrated;

[0031] Figure 32B The inference and / or training logic according to at least one embodiment is shown; and

[0032] Figure 32C Training and deployment of a neural network according to at least one embodiment are illustrated. Detailed Implementation

[0033] As will be described below, numerous specific details will be set forth in order to provide a more comprehensive understanding of at least one embodiment. However, it will be apparent to those skilled in the art that the concept of the invention can be practiced without one or more of these specific details.

[0034] Figure 1 A system 100 is illustrated, according to at least one embodiment, for translating a program written in a VM-based language into instructions that use operands in registers and / or memory. In at least one embodiment, the system 100 is capable of translating or otherwise converting program source code 102 written in a VM-based language into GPU instructions 142 that can run on a GPU 140. In at least one embodiment, the VM-based language is Python, C#, Java, or another language that uses a stack-based virtual machine to execute program instructions. In at least one embodiment, the system 100 can be used to adapt an existing native parallel processing API 136 that is compatible with a language (e.g., C++) for use by VM-based languages.

[0035] In at least one embodiment, the virtual machine (VM) uses a stack-based processing model, where instruction operands (e.g., input parameters and / or output parameters) are stored in a stack data structure. The stack data structure is a last-in-first-out (LIFO) list of items, from which a retrieval operation retrieves the most recently added item. Therefore, for an instruction with two or more operands, retrieving the operands from the stack yields a list of operands in reverse order compared to the order in which they were added to the stack. Furthermore, the operands for an instruction are stored in the stack before the instruction is executed. VM instructions may store operands (such as variables (e.g., memory addresses) or constants) in the stack, and other program instructions may subsequently retrieve their input operands from the stack and store their output operands on the stack.

[0036] In at least one embodiment, executing a program using a stack with GPU 140 may be less efficient than executing a program using GPU 140 that performs similar operations but uses registers and / or memory instead of a stack to store instruction operands. In at least one embodiment, accessing registers is more efficient than accessing the stack with GPU 140 because accessing the stack may involve synchronization between multiple threads.

[0037] In at least one embodiment, system 100 includes a language compiler 104 that translates program source code 102 of a VM-based, stack-based language into program bytecode 106. In at least one embodiment, the program bytecode 106 includes one or more VM instructions (e.g., add, load from memory, store to memory, etc.) that retrieve input operands from a stack, remove the input operands from the stack, perform one or more operations, and store output operands onto the stack. In at least one embodiment, system 100 includes a language-specific managed layer 108 that provides a set of language-specific parallel processing APIs 110 (e.g., CUDA APIs) for use by the program bytecode 106 generated by the language compiler 104 for a specific VM-based language (e.g., Python or C#). In at least one embodiment, the language-specific managed layer uses a VM-based language runtime 122. In at least one embodiment, the language-specific managed layer 108 is implemented for each VM-based language, and system 100 translates the program source code 102 of that VM-based language into GPU instructions 142.

[0038] In at least one embodiment, a programmer may implement the language-specific managed layer 108 by implementing a language-specific parallel processing API (e.g., in a VM-based language such as Python), which converts data from a language-specific format to a language-independent format used as input and output to a language-managed API 114, which is language-independent (e.g., using bytecode from a language-independent common runtime 134). In at least one embodiment, the language-specific managed layer 108 includes API translation functions (e.g., wrapper functions) implemented by a language-specific program to adapt the language-specific parallel processing API 110 to the language-managed API 114. In at least one embodiment, the language-specific managed layer may provide an interface between garbage-collected data structures used by program bytecode 106 and memory management operations for data structures used by the native parallel processing API 136.

[0039] In at least one embodiment, the language-managed API 114 includes a language-independent parallel processing API 118. In at least one embodiment, the language-independent parallel processing API 118 is a language-independent CUDA API that a programmer can use (e.g., call) to implement a language-specific managed layer 108 as described herein (e.g., write source code for the language-specific managed layer 108). In at least one embodiment, the language-specific managed layer 108 is bytecode generated by a language compiler 104 from the source code implemented by the programmer. In at least one embodiment, the language-specific managed layer 108 is relatively simple to implement for a programmer because the language-independent parallel processing API 118 performs more complex tasks involved in adapting different languages ​​to existing parallel processing API implementations (such as the native parallel processing API 136).

[0040] In at least one embodiment, the language-independent parallel processing API 118 uses a common core 112 (e.g., a bytecode library) that provides relatively complex program code to adapt the native parallel processing API to a VM-based language. In at least one embodiment, the language-independent parallel processing API 118 and the common core 112 are implemented in a language such as C#, which provides the API in a language-independent API format (e.g., using a common runtime 134). In at least one embodiment, the language-independent parallel processing API 118 can be called from any language compatible with the language-independent API format. In at least one embodiment, a language-specific managed layer 108 performs data type conversions between VM-based languages ​​(e.g., Python, C#, or Java) to adapt the VM-based language to use the language-managed API 114. In at least one embodiment, the common core 112 performs conversions between the language-managed API 114 and the language implemented by the native parallel processing API 136 (e.g., C or C++).

[0041] In at least one embodiment, the language-managed API 114 includes a kernel startup API 116, which can be invoked by a language-specific managed layer to enable the execution of a parallel processing kernel implemented in a VM-based language (e.g., as stack-based bytecode instructions) on the GPU 140. In at least one embodiment, the common core 112 includes a VM bytecode-to-GPU instruction translator 130 that translates the GPU kernel bytecode 120 into GPU instructions 132. In at least one embodiment, because the common runtime 134 performs translation between the bytecode instructions of the common core 112 and specific processor architectures and / or instruction sets, the common core 112 uses the common runtime (executed on the CPU), making the common core 112 independent of specific processor architectures and / or instruction sets. In at least one embodiment, the GPU-specific common runtime 144 on the GPU 140 provides GPU-specific operations invoked by GPU instructions 142 and / or the native parallel processing API 136.

[0042] Figure 2 An exemplary VM-to-GPU instruction translator 130 according to at least one embodiment is shown. In at least one embodiment, the VM-to-GPU instruction translator 130 corresponds to Figure 1A VM-to-GPU instruction translator 130 is provided. In at least one embodiment, the VM-to-GPU instruction translator 130 translates bytecode operations 220 into GPU instructions 132. In at least one embodiment, each bytecode operation 220 includes an offset (e.g., the number of bytes from the start of the bytecode program), a tag 224 (if present, the tag is an identifier representing the offset, e.g., an alphanumeric string), an action statement 226, or an operand instruction 228. In at least one embodiment, the action statement 226 may be a bytecode instruction that retrieves an operand from the VM stack, performs the operation, and stores the result on the VM stack. In at least one embodiment, as an alternative to the action statement 226, the bytecode operation 220 may include an operand instruction 228 to push (e.g., store) an operand onto the VM stack or pop (e.g., retrieve and remove) an operand from the VM stack.

[0043] In at least one embodiment, the VM-to-GPU instruction translator 130 performs the instructions described herein (see reference 1). Figure 12 The operations described in the process of translating a VM program into a GPU program are illustrated. In at least one embodiment, the VM-to-GPU instruction translator 130 retrieves bytecode operation 220. In at least one embodiment, the bytecode operation 220 retrieved by the VM-to-GPU instruction translator 130 from the bytecode operation 220 is shown as the current bytecode operation 240. In at least one embodiment, the current bytecode operation 240 includes the offset 222 and label 224 (if present) as described above. In at least one embodiment, the current bytecode operation 240 also includes a branch action statement 226A, a non-branch action statement 226B, or an operand instruction 228. In at least one embodiment, the branch action statement 226A is a branch instruction that causes the VM program to execute the bytecode operation 220 indicated by the specified label and continue execution of subsequent bytecode instructions after the specified label (or at another specified label if the specified label indicates a branch instruction). In at least one embodiment, the branch action statement 226A specified by bytecode operation 240 causes the program counter specifying the next instruction to be executed to be set to an offset 222 associated with the label 224 specified by the bytecode operation 240. In at least one embodiment, the label specified by the branch instruction is referred to herein as a branch target label or a branch target instruction. In at least one embodiment, the branch action statement 226A causes the execution of the program containing the branch action statement 226A to continue at the target instruction indicated by the label 224 associated with the branch action statement 226A.

[0044] In at least one embodiment, branch action statement 226A may be an unconditional branch action statement 226A, which causes the VM program to continue execution at the target instruction indicated by label 224. In at least one embodiment, branch action statement 226A may be a conditional branch action statement 226A, which causes the VM program to continue execution at the target instruction indicated by the true branch target label 242 when the condition is determined to be true, or to continue execution at the target instruction indicated by the false branch target label 244 when the condition is determined to be false. In at least one embodiment, the condition is a Boolean condition, such as a comparison between two values. In at least one embodiment, the two values ​​are retrieved from the VM stack by the branch action statement 226A. In at least one embodiment, the current bytecode operation 240 includes a non-branch action statement 226B, which may be any action statement other than a branch action statement. In at least one embodiment, the non-branch action statement 226B is an addition statement or other statement that retrieves one or more operands from the VM stack, performs an action, and stores one or more results on the VM stack. In at least one embodiment, the non-branch action statement 226B does not retrieve operands from the VM stack and / or does not store results on the VM stack. In at least one embodiment, the action performed by the non-branch action statement 226B is any action other than branching, storing a value on the VM stack, or retrieving a value from the VM stack. In at least one embodiment, the current bytecode operation 240 may be an operand instruction 228. In at least one embodiment, the operand instruction stores a value on the VM stack (e.g., a push instruction). In at least one embodiment, the operand instruction retrieves a value from the VM stack and removes the value from the VM stack (e.g., a pop instruction).

[0045] In at least one embodiment, the VM-to-GPU instruction translator 130 generates a reparse point stack 250 containing at least one reparse point 252. In at least one embodiment, the reparse point 252 is a data structure that includes a branch target instruction label 254 and an associated saved operand instruction stack 256. In at least one embodiment, the VM-to-GPU instruction translator 130 translates bytecode operations 220 in offset order (e.g., sequentially), except for operations 220 including branch action statement 226A. In at least one embodiment, upon encountering a current bytecode operation 240 specifying a branch action statement 226A, the VM-to-GPU instruction translator 130 creates a reparse point 252, which includes a label 254 associated with the branch action statement 226A and a saved operand instruction stack 256, which is a copy of the operand instruction stack 270. In at least one embodiment, when retrieving reparse point 252 from the reparse point stack 250, the VM-to-GPU instruction translator 130 translates the branch action statement 226A into GPU instruction 132, and then retrieves reparse point 252 from the reparse point stack 250. In at least one embodiment, the VM-to-GPU instruction translator 130 translates each action statement 226 in the bytecode operation 220 into a corresponding GPU instruction (e.g., an addition action is translated into an addition instruction to be executed by the GPU). In at least one embodiment, the VM-to-GPU instruction translator 130 generates a GPU instruction 132 for each bytecode operation 220 containing an action statement 226, and the GPU instruction contains an operation 260 (e.g., an addition operation) corresponding to the action statement and one or more operands 262. The VM-to-GPU instruction translator 130 retrieves the one or more operands from the operand instruction stack 270 based on the input used by the action statement 226 (e.g., for an addition action statement 226 with two input operands, two operand instructions 272 are retrieved).

[0046] In at least one embodiment, upon retrieving a reparse point 252 from the reparse point stack 250, the VM-to-GPU instruction translator 130 loads a saved operand instruction stack 256 from the reparse point 252 and uses the saved operand instruction stack 256 as an operand instruction stack 270 to translate bytecode operations 220. In at least one embodiment, furthermore, upon retrieving a reparse point 252 from the reparse point stack 250, the VM-to-GPU instruction translator 130 begins translating instructions at the instruction specified by the tag 254 indicated by the reparse point. In at least one embodiment, upon retrieving a reparse point 252 from the reparse point stack 250, the VM-to-GPU instruction translator 130 removes the reparse point 252 from the reparse point stack 250.

[0047] In at least one embodiment, the VM-to-GPU instruction translator 130 continues to translate instructions sequentially from the tag 254 until it encounters another operation 220 containing a branch action statement 226A (in which case, another re-parse point 252 is added to the re-parse point stack 250, then the re-parse point is retrieved from the re-parse point stack 250, and translation continues at the tag 254 specified by the re-parse point). In at least one embodiment, the VM-to-GPU instruction translator 130 stops translating when the re-parse point stack 250 is empty and no more bytecode operations 220 are available.

[0048] In at least one embodiment, upon encountering a bytecode operation 240 with a specified operand instruction 228, the VM-to-GPU instruction translator 130 stores the operand instruction on an operand instruction stack 270. In at least one embodiment, the operand instruction stack is a simulated VM stack 250, which the VM-to-GPU instruction translator 130 uses to translate the bytecode operation 220 into GPU instructions 132. In at least one embodiment, upon encountering an action statement 226 with one or more input operands (e.g., input to an addition action statement), the VM-to-GPU instruction translator 130 removes an operand instruction from the operand instruction stack 270 and uses the removed operand instruction as an input operand for the action statement 226. In at least one embodiment, the VM-to-GPU instruction translator 130 does not simulate the action specified by the action statement 226 (e.g., not calculating the sum of the inputs of an addition instruction), but instead translates the action statement 226 into GPU instructions 132 to execute the action when executed by the GPU 140.

[0049] Figure 3Exemplary VM-oriented instructions and corresponding GPU-oriented instructions according to at least one embodiment are illustrated, wherein VM-oriented instructions access operands on the stack, while GPU-oriented instructions access operands in registers and / or memory. In at least one embodiment, the VM-oriented instructions (e.g., VM operations) are illustrated as VM program 310, which pushes a constant value 1 onto the VM stack, pushes a condition onto the VM stack, and performs a conditional branch operation. In at least one embodiment, if the condition is true, the conditional branch operation transfers program control to an operation labeled "lbl1"; if the condition is false, the conditional branch operation transfers control to an operation labeled "lbl2". In at least one embodiment, the operation at label "lbl1" pushes a constant value 2 onto the VM stack, and the subsequent operation branches to label "lbl3". In at least one embodiment, the operation at "lbl3" executes an addition action statement that retrieves and removes two values ​​from the VM stack and adds the two values. In at least one embodiment, the operation at label "lbl2" pushes the constant value 3 onto the VM stack, followed by the addition action statement. In at least one embodiment, a preprocessing operation is performed to add a branch instruction before each operation (i.e., an operation that would result in statement penetration if there were no branching statements). In at least one embodiment, as shown in the preprocessed VM program 312, the operation "branch lbl3" is added at offset 6.5.

[0050] In at least one embodiment, the VM-to-GPU instruction translator 130 translates the preprocessed VM program 312 into a GPU program 124. In at least one embodiment, the VM-to-GPU instruction translator 130 translates the preprocessed VM program 312 into an intermediate program 312 that does not use a stack. In at least one embodiment, as shown, the intermediate program 312 includes addition instructions 304 and 305, whose operands are provided as parameters instead of on the VM stack. In at least one embodiment, the operands of the addition instructions 304 and 305 may be stored in registers and / or memory locations instead of on the VM stack. In at least one embodiment, the intermediate program 314 is as described above. Figures 4 to 11 Example generation. In at least one embodiment, the VM-to-GPU instruction translator 130 uses appropriate compiler optimization techniques to translate the intermediate program 314 into an optimized intermediate program 316, for example, removing unnecessary operations, such as goto operations 302, 303 after branch operation 301.

[0051] Figures 4 to 11An exemplary translation of VM-oriented instructions into GPU-oriented instructions is illustrated according to at least one embodiment. In at least one embodiment, the VM-oriented instructions are an exemplary VM program 310, and the GPU-oriented instructions are an optimized intermediate program 316. In at least one embodiment, the exemplary translation is based on... Figure 12 The process 1200 is described herein. In at least one embodiment, process 1200 has an outer loop that is executed in one or more outer loop iterations. In at least one embodiment, each outer loop iteration begins at condition box 1206 and ends at box 1226 (if the operation is a branching action) or 1214 (if the operation is not a branching action and there are no more operations to decode), as referenced herein. Figure 12 Furthermore, in at least one embodiment, process 1200 has an inner loop, which is executed in one or more inner loop iterations for each outer loop iteration. In at least one embodiment, each inner loop iteration begins at block 1214 and ends at block 1230 or 1232 (depending on whether the operation is an operand instruction).

[0052] Figure 4 An exemplary translation of VM-oriented instructions for program 310, starting from the first outer loop iteration (outer loop iteration #1) 401 of translator 130, is shown according to at least one embodiment. In at least one embodiment, the first outer loop iteration 401 is used to process an initial re-resolution point 252, which occurs before the first outer loop iteration 401 (e.g., in...). Figure 12 The initial re-parse point 252 (at box 1202) is added to the re-parse point stack 250. In at least one embodiment, the initial re-parse point 252 corresponds to the initial bytecode operation 220 of the program 310.

[0053] In at least one embodiment, in each outer loop iteration, the translator 130 retrieves a re-parse point 252 from the re-parse point stack 250 and removes the retrieved re-parse point 252 from the re-parse point stack 250 (e.g., in...). Figure 12 (at box 1208). In at least one embodiment, the translator 130 sets the current tag to the tag 254 of the retrieved re-parsing point 252, and sets the contents of the operand instruction stack 270 (e.g., operand instruction 272) according to the saved operand instruction stack 256 of the retrieved re-parsing point 252 (e.g., in the box 1208). Figure 12 (at box 1210). In at least one embodiment, the translator 130 sets the bytecode index to the bytecode offset corresponding to the current tag (in input program 310) (e.g., in...). Figure 12(at box 1212). In at least one embodiment, the translator 130 outputs the current label, for example, by adding the current label to the list of output GPU instructions 132.

[0054] In at least one embodiment, each outer loop iteration (e.g., 401) performs one or more inner loop iterations (e.g., 401A, 401B). In at least one embodiment, each inner loop iteration processes the bytecode operation 220 of the input program 310 (e.g., in...). Figure 12 (in boxes 1208-1232). In at least one embodiment, each inner loop iteration decodes bytecode operation 220 at the bytecode index and updates the bytecode index to point to the next bytecode instruction (e.g., in...). Figure 12 (at box 1214). In at least one embodiment, bytecode operation 220 includes action statement 225 or operand instruction 228. In at least one embodiment, the inner loop iteration is performed to transform action statement 225 and operand instruction 228 until a bytecode operation 220 containing branch action statement 226A is encountered.

[0055] In at least one embodiment, upon encountering a branch action statement 226A, the translator 130 creates a new re-parsing point 252 for each branch case (e.g., true and false) of the branch action statement 226A, and stores each new re-parsing point 252 (including the contents of the operand instruction stack 270) in the re-parsing point stack 250 (e.g., in...). Figure 12 (See boxes 1218-1226). In at least one embodiment, the translator 130 then generates an output GPU instruction 132 having an operation 260 (e.g., a "branch" operation) determined according to the branch action statement 226B and operands 262 (e.g., condition and target labels if the branch action is a conditional branch) determined according to operand instructions 272 of the operand instruction stack 270. In at least one embodiment, after translating the branch action statement 226A, the translator 130 begins the next outer loop iteration.

[0056] In at least one embodiment, when an operation containing operand instructions is encountered, the translator 130 adds the operand instruction 228 to the operand instruction stack 270 (e.g., in...). Figure 12(at box 1230). In at least one embodiment, upon encountering an operation containing a non-branching action statement 226B, the translator 130 generates an output GPU instruction 132 having an operation 260 determined according to the non-branching action statement 226B and an operand 262 determined according to operand instructions 272 of the operand instruction stack 270. In at least one embodiment, after translating the non-branching action statement 226B, the translator 130 begins the next outer loop iteration. In at least one embodiment, the next outer loop iteration retrieves and removes a re-parse point 252 from the re-parse point stack 250 (e.g., from the top of the re-parse point stack 250) and performs one or more inner loop iterations.

[0057] In at least one embodiment, reference Figure 4 The first outer loop iteration 401 processes the initial re-parse point 252, which corresponds to the initial bytecode operation 220 of the program 310 (for example, the initial re-parse point 252 may be...). Figure 12 (Created at box 1202). In at least one embodiment, the initial re-parse point 252 includes a label 254 (e.g., "_start") corresponding to the initial bytecode operation (e.g., "push constant 1") and has an empty saved operand instruction stack 256. In at least one embodiment, to process the re-parse point, the translator 130 retrieves and removes the re-parse point 252 from the re-parse point stack 250. In at least one embodiment, the re-parse point stack 250 contains the initial re-parse point 252, which has a label 254 corresponding to the initial bytecode operation and also has an empty saved operand instruction stack 256. In at least one embodiment, the translator 130 sets the bytecode index to a bytecode offset of "1" corresponding to the label 254 (e.g., "_start") of the re-parse point retrieved from the re-parse point stack.

[0058] In at least one embodiment, translator 130 begins a first inner loop iteration ((inner loop iteration #1)) 401A to process the bytecode operation 220 by decoding the bytecode operation 220 at bytecode index "1". In at least one embodiment, translator 130 then advances the bytecode index to the next bytecode index "2" (e.g., by adding the size of the decoded bytecode operation 220 to the bytecode index). In at least one embodiment, the advanced bytecode index "2" is used in subsequent inner loop iterations to decode the next instruction. In at least one embodiment, the current bytecode operation 411 has an offset 421 (e.g., the same as bytecode index "1") and an operand instruction 431 ("push constant 1"). In at least one embodiment, translator 130 determines whether the current bytecode operation 411 contains a branch action statement 226A, a non-branch action statement 226B, or an operand instruction 228. Since the current bytecode operation 411 contains operand instruction 431 (“pushconstant 1”), the translator 130 adds operand instruction 431 to the operand instruction stack 441 (e.g., to the bottom of the operand instruction stack 441). In at least one embodiment, the first inner loop iteration 401A ends, and the translator 130 attempts to decode another operation to process in another inner loop iteration (e.g., in...). Figure 12 (at box 1214).

[0059] In at least one embodiment, translator 130 initiates a second inner loop iteration 401B to process another bytecode operation after successfully decoding bytecode operation 220 at bytecode index "2". In at least one embodiment, translator 130 then advances the bytecode index to the next bytecode index "3". In at least one embodiment, during the second inner loop iteration 401B, translator 130 decodes the current bytecode operation 412 at bytecode index "2". In at least one embodiment, the current bytecode operation 412 has an offset 422 (e.g., bytecode index "2") and an operand instruction 432 ("push condition"). In at least one embodiment, translator 130 determines whether the current bytecode operation 412 contains a branch action statement 226A, a non-branch action statement 226B, or an operand instruction 228. Since the current bytecode operation 412 contains operand instruction 432 (“push condition”), the translator 130 adds operand instruction 432 to the operand instruction stack 441 (e.g., to the bottom of the operand instruction stack 441). That is, in at least one embodiment, the second inner loop iteration 401B ends, and the translator 130 attempts to decode another operation for processing in another inner loop iteration (e.g., in...). Figure 12 (at box 1214).

[0060] Figure 5 An exemplary translation of a VM instruction-oriented program 310 according to at least one embodiment is shown, which continues the first outer loop iteration 501 of the translator 130 to process another input program operation after the initial re-parsing point. In at least one embodiment, as Figure 5 The third inner loop iteration 501C shown transforms the bytecode operation at line 3 of program 310. In at least one embodiment, the bytecode operation at line 3 is a conditional branch operation (“conditional branch lbl1: lbl2”), which causes program 310 to branch to the instruction preceded by label “lbl1” when the condition is true, or to the instruction preceded by label “lbl2” when the condition is false.

[0061] In at least one embodiment, translator 130 initiates a third inner loop iteration 501C to process another bytecode operation upon successfully decoding the current bytecode operation 513 at bytecode index "3". In at least one embodiment, translator 130 then advances the bytecode index to the next bytecode index "4". In at least one embodiment, the current bytecode operation 513 has an offset 523 (e.g., bytecode index "3") and a branch action statement 533 ("conditional branch lbl1: lbl2"). In at least one embodiment, translator 130 determines whether the current bytecode operation 513 contains a branch action statement, a non-branch action statement, or an operand instruction. Since the current bytecode operation 513 contains a conditional branch action statement 533 ("conditional branch lbl1: lbl2"), translator 130 retrieves the branch condition from the operand instruction stack 543. In at least one embodiment, to retrieve the branch condition, translator 130 retrieves and removes an operand instruction ("push condition") from the operand instruction stack 543. In at least one embodiment, retrieving the operand instruction leaves a single operand instruction 543A, namely "push constant 1", on the operand instruction stack 543.

[0062] In at least one embodiment, the translator 130 identifies the branch target labels of the true and false branches of the conditional branch action statement (e.g., the true branch is "lbl1" and the false branch is "lbl2"). In at least one embodiment, the translator 130 creates a true branch re-parse point and a false branch re-parse point (arrow 534) for the true branch label "lbl1" and the false branch label "lbl2" of the branch action. In at least one embodiment, the translator 130 creates a true branch re-parse point 563A named "Point1", which has a target label 573A corresponding to the branch target label ("lbl1") of the true branch and a saved operand instruction stack 583A ("push constant 1") copied from the operand instruction stack 543. In at least one embodiment, translator 130 creates a pseudo branch re-parsing point 563B named "Point2," which has a target label 573B corresponding to the branch target label ("lbl2") of the pseudo branch, and a saved operand instruction stack 583B ("pushconstant 1") copied from operand instruction stack 543. In at least one embodiment, translator 130 adds the true branch re-parsing point 563A ("Point1") and the pseudo branch re-parsing point 563B ("Point2") to the re-parsing point stack 553 (e.g., such that the last added re-parsing point 563B is at the bottom of the re-parsing point stack 553).

[0063] In at least one embodiment, translator 130 then generates an output GPU instruction 593 having an operation 594 (e.g., a "conditional branch lbl1: lbl2" operation) determined according to the branch action statement 533 and an operand 596 (e.g., a condition) determined according to an operand instruction 272 retrieved from the operand instruction stack 270. In at least one embodiment, translator 130 replaces the branch label with the re-parse point name 563 corresponding to the branch label. In at least one embodiment, translator 130 converts the branch label "lbl1" to the re-parse point name "Point1" and the branch label "lbl2" to the re-parse point name "Point2". In at least one embodiment, translator 130 generates an output instruction 598 where the branch label has been replaced by the re-parse point name ("branch (condition) Point1: Point2"). In at least one embodiment, after converting the branch action statement 533, translator 130 begins a new outer loop iteration to process another re-parse point, so no further inner loop iterations are performed in the first outer loop iteration 501.

[0064] Figure 6An exemplary translation of VM-oriented instructions for an input program 310 according to at least one embodiment is shown, the translation beginning with the second outer loop iteration 602 of the translator 130 to process re-parsing point 563A (“Point1”). In at least one embodiment, in the second outer loop iteration 602, the translator 130 retrieves and removes the popped re-parsing point 563A (“Point1”) from the re-parsing point stack 553 (e.g., in...). Figure 12 (at box 1208). In at least one embodiment, the re-parse point 563A corresponds to the branch target instruction at label "lbl1", and Figure 5 In the third inner loop iteration 501C, it is added to the re-parsing point stack. In at least one embodiment, the translator 130 sets the current label to the label 573A (“lbl1”) of the popped re-parsing point 563A, and sets the contents of the operand instruction stack 651 (e.g., operand instructions) according to the saved operand instruction stack 583A of the popped re-parsing point 563A. Figure 12 (at box 1210). In at least one embodiment, the translator 130 sets the index of the recovered parsed bytecode to the bytecode offset corresponding to tag 573A (“lbl1”) of the pop re-parsing point 563A, which is “4” in the input program 310 (e.g., in...). Figure 12 (at box 1212). In at least one embodiment, translator 130 outputs a label corresponding to the pop re-parsing point 563A. In at least one embodiment, the label is the name of the pop re-parsing point 563A (e.g., "Point1"). In at least one embodiment, to output the label, translator 130 generates an output instruction 792 having the name of the pop re-parsing point 563A (e.g., "Point1"), such as... Figure 7 As shown.

[0065] In at least one embodiment, the second outer loop iteration 602 executes one or more inner loop iterations, including the first inner loop iteration 602A, to process the operation of the input program 310, starting from the bytecode index (corresponding to label “lbl1”) and using the operand instruction stack 651 recovered from re-parse point 563A (“Point1”). In at least one embodiment, the first inner loop iteration 602A processes the operation of the input program 310 at label “lbl1” by adding “push constant 2” to the operand instruction stack 651, which is “push constant 2”.

[0066] In at least one embodiment, translator 130 initiates the first inner loop iteration 602A by decoding bytecode operation 611 at bytecode index "4". In at least one embodiment, translator 130 then advances the bytecode index to the next bytecode index "5". In at least one embodiment, the current bytecode operation 611 has an offset 621 (e.g., the same as bytecode index "4"), a label 631 ("lbl1"), and an operand instruction 641 ("push constant 2"). In at least one embodiment, translator 130 determines whether the current bytecode operation 611 contains a branching action statement, a non-branching action statement, or an operand instruction. Since the current bytecode operation 611 contains operand instruction 641 ("push constant 2"), translator 130 adds operand instruction 641 to the operand instruction stack 651 (e.g., to the bottom of the operand instruction stack 651). In at least one embodiment, the first inner loop iteration 602A ends, and the re-parse point stack 661 does not contain the popped re-parse point 563A (“Point1”) that has been removed from the stack 661. In at least one embodiment, the translator 130 attempts to decode another operation for processing in another inner loop iteration (e.g., in...). Figure 12 (at box 1214).

[0067] Figure 7 An exemplary transformation is shown for the second outer loop iteration 602 of a VM-oriented instruction continuation translator 130 of program 310 according to at least one embodiment, to process another input program operation following re-parse point 563A (“Point 1”). In at least one embodiment, the second inner loop iteration 702B transforms the bytecode operation at line 5 of program 310. In at least one embodiment, the bytecode operation at line 5 is an unconditional branch operation (“branch lbl3”) used to branch program 310 to an instruction preceded by the label “lbl3”.

[0068] In at least one embodiment, translator 130 initiates a second inner loop iteration 702B to process another bytecode operation after successfully decoding bytecode operation 712 at bytecode index "5". In at least one embodiment, translator 130 then advances the bytecode index to the next bytecode index "6". In at least one embodiment, the current bytecode operation 712 has an offset 722 (the same as bytecode index "5") and a branch action statement 742 ("branch lbl3"). In at least one embodiment, translator 130 determines whether the current bytecode operation 712 contains a branch action statement, a non-branch action statement, or an operand instruction. Since the current bytecode operation 712 contains an unconditional branch action statement 742 ("branch lbl3"), translator 130 identifies a branch target label (e.g., "lbl3"). In at least one embodiment, translator 130 creates a re-parse point (arrow 744) for the label "lbl3" of the branch action. In at least one embodiment, translator 130 creates a reparse point 762A named “Point3”, which has a target label 772A (“lbl3”) and a saved operand instruction stack 782A copied from operand instruction stack 752 (“push constant 1” and “push constant 2”). In at least one embodiment, translator 130 adds the reparse point 762A (“Point3”) to the reparse point stack 762 (e.g., at the bottom of the reparse point stack 762).

[0069] In at least one embodiment, the translator 130 then generates an output instruction 792, which has a re-parsing point 563A (e.g., "Point 1", as referenced) Figure 6 The corresponding label and the operation 796 determined according to the branch action statement 742 (e.g., "Goto Point 3"). In at least one embodiment, the output instruction 792 is displayed as a label and instruction combination 798 ("Point 1: Goto Point 3"). In at least one embodiment, after translating the branch action statement 742, the translator 130 begins a new outer loop iteration to process another re-parse point, so no further inner loop iterations are performed in the second outer loop iteration 702.

[0070] Figure 8 An exemplary translation of VM-oriented instructions for an input procedure 310 according to at least one embodiment is shown, the translation beginning at the third outer loop iteration 803 of the translator 130 to process re-parsing point 563B (“Point 2”). In at least one embodiment, Figure 8 The third outer loop iteration 803 is similar to Figure 6In the second outer loop iteration 602, except for the re-parse point 563B "Point2" (for the branch false case) with label "lbl2",... Figure 8 The process is carried out. In at least one embodiment, Figure 8 The first inner loop iteration of 803A is similar to Figure 6 In the first inner loop iteration 602A, except for the bytecode operation "push constant 3" at the label "lbl2",... Figure 8 It was processed in the middle.

[0071] In at least one embodiment, in the third outer loop iteration 803, the translator 130 retrieves and removes the popped re-parsing point 563B (“Point2”) from the re-parsing point stack 762 (e.g., in...). Figure 12 (at box 1208). In at least one embodiment, the re-parse point 563B corresponds to the branch target instruction at label "lbl2", and Figure 5 In the third inner loop iteration 501C, it is added to the re-parsing point stack. In at least one embodiment, the translator 130 sets the current label to the label 573B (“lbl2”) of the popped re-parsing point 563B, and sets the contents of the operand instruction stack 851 (e.g., operand instructions) according to the saved operand instruction stack 583B of the popped re-parsing point 563B. Figure 12 (at box 1210). In at least one embodiment, the translator 130 sets the index of the recovered bytecode to the bytecode offset corresponding to the tag 573B (“lbl2”) of the pop re-parse point 563B, which is “6” in the input program 310 (e.g., in...). Figure 12 (at box 1212). In at least one embodiment, translator 130 outputs a label corresponding to the pop re-parsing point 563B. In at least one embodiment, the label is the name of the pop re-parsing point 563B (e.g., "Point2"). In at least one embodiment, to output the label, translator 130 generates an output instruction 992 having the name of the pop re-parsing point 563B (e.g., "Point2"), as shown in box 1212. Figure 9 As shown.

[0072] In at least one embodiment, the third outer loop iteration 803 executes one or more inner loop iterations, including the first inner loop iteration 803A, to process the operation of the input program 310, starting from the bytecode index (corresponding to label “lbl2”) and using the operand instruction stack 851 recovered from reparse point 563B (“Point2”). In at least one embodiment, the first inner loop iteration 803A processes the operation of the input program 310 at label “lbl2” by adding “push constant 3” to the operand instruction stack 851.

[0073] In at least one embodiment, translator 130 initiates the first inner loop iteration 803A by decoding bytecode operation 811 at bytecode index "6". In at least one embodiment, translator 130 then advances the bytecode index to the next bytecode index "6.5". In at least one embodiment, the current bytecode operation 811 has an offset 821 (e.g., the same as bytecode index "6"), a label 831 ("lbl2"), and an operand instruction 841 ("push constant 3"). In at least one embodiment, translator 130 determines whether the current bytecode operation 811 contains a branching action statement, a non-branching action statement, or an operand instruction. Since the current bytecode operation 811 contains operand instruction 841 ("push constant 3"), translator 130 adds operand instruction 841 to the operand instruction stack 851 (e.g., to the bottom of the operand instruction stack 851). In at least one embodiment, the first inner loop iteration 803A ends, and the re-parse point stack 861 does not contain the popped re-parse point 563B (“Point2”) that has been removed from the stack 861. In at least one embodiment, the translator 130 attempts to decode another operation for processing in another inner loop iteration (e.g., in...). Figure 12 (at box 1214).

[0074] Figure 9 An exemplary transformation is shown in the third outer loop iteration 903 of the VM-oriented instruction continuation translator 130 of program 310 according to at least one embodiment, to process another input program operation following re-parsing point 563B (“Point 2”). In at least one embodiment, Figure 9 The 3rd outer loop iteration 903 corresponds to Figure 8 The third outer loop iteration 803. In at least one embodiment, Figure 9 The second inner loop iteration 903B executes and Figure 7 The second inner loop iteration of 702B involves operations similar to those, except... Figure 9The code being processed is the bytecode operation "branch lbl3" at offset "6.5", not the bytecode operation "branch lbl3" at offset "5". Furthermore, Figure 9 With Figure 7 Different operand instruction stack contents and different reparsing point stack contents result in different... Figure 7 The output instructions are different.

[0075] In at least one embodiment, the second inner loop iteration 903B translates the bytecode operation at line 6.5 of program 310. In at least one embodiment, the bytecode operation at line 6.5 is an unconditional branch operation (“branchlbl3”) used to branch program 310 to an instruction preceded by the label “lbl3”.

[0076] In at least one embodiment, translator 130 initiates a second inner loop iteration 903B to process another bytecode operation after successfully decoding bytecode operation 912 at bytecode index "6.5". In at least one embodiment, translator 130 then advances the bytecode index to the next bytecode index "7". In at least one embodiment, the current bytecode operation 912 has an offset 922 (the same as bytecode index "6.5") and a branch action statement 942 ("branch lbl3"). In at least one embodiment, translator 130 determines whether the current bytecode operation 912 contains a branch action statement, a non-branch action statement, or an operand instruction. Since the current bytecode operation 912 contains an unconditional branch action statement 942 ("branch lbl3"), translator 130 identifies a branch target label (e.g., "lbl3"). In at least one embodiment, translator 130 creates a re-parse point (arrow 944) for the label "lbl3" of the branch action. In at least one embodiment, translator 130 creates a reparse point 962A named “Point4”, which has a target label 972A (“lbl3”) and a saved operand instruction stack 982A copied from operand instruction stack 952 (“push constant 1” and “push constant 3”). In at least one embodiment, translator 130 adds the reparse point 962A (“Point4”) to the reparse point stack 962 (e.g., at the bottom of the reparse point stack 962).

[0077] In at least one embodiment, the translator 130 then generates an output GPU instruction 992, which has a stack re-parsing point 563B (e.g., "Point 2", as referenced). Figure 8The corresponding label and the operation 996 determined according to the branch action statement 942 (e.g., "Goto Point 3"). In at least one embodiment, the output instruction 992 is displayed as a label and instruction combination 998 ("Point 2: Goto Point 4"). In at least one embodiment, after translating the branch action statement 942, the translator 130 begins a new outer loop iteration to process another re-parse point, so no further inner loop iterations are performed in the third outer loop iteration 903.

[0078] Figure 10 An exemplary translation of VM-oriented instructions for an input program 310 according to at least one embodiment is shown, the translation beginning in the fourth outer loop iteration 1004 of the translator 130 to process re-parsing point 762A (“Point 3”). In at least one embodiment, in the fourth outer loop iteration 1004, the translator 130 retrieves and removes the re-parsing point 762A (“Point 3”) from the re-parsing point stack 962 (e.g., in...). Figure 12 (at box 1208). In at least one embodiment, the re-parse point 762A corresponds to the branch target instruction at label "lbl3", and Figure 7 In the second inner loop iteration 702B, it is added to the re-parsing point stack. In at least one embodiment, the translator 130 sets the current label to the label 772A (“lbl3”) of the popped re-parsing point 762A, and sets the contents of the operand instruction stack 1051 (e.g., operand instructions) according to the saved operand instruction stack 782A of the popped re-parsing point 762A. Figure 12 (at box 1210). In at least one embodiment, the translator 130 sets the index of the recovered bytecode to the bytecode offset corresponding to the tag 772A (“lbl3”) of the pop re-parse point 762A, which is “7” in the input program 310 (e.g., in...). Figure 12 (See box 1212). In at least one embodiment, translator 130 outputs a label 1094 corresponding to the pop re-parsing point 762A. In at least one embodiment, the label 1094 is the name of the pop re-parsing point 762A. In at least one embodiment, to output the label 1094, translator 130 generates an output instruction 1092 having the label 1094 (e.g., "Point3") specifying the name of the pop re-parsing point 762A.

[0079] In at least one embodiment, the fourth outer loop iteration 1004 executes the first inner loop iteration 1004A to process the operation of the input program 310, starting from the bytecode index (corresponding to label "lbl3") and using the operand instruction stack 1051 recovered from the re-parse point 762A ("Point3"). In at least one embodiment, the first inner loop iteration 1004A processes the operation of the input program 310 at label "lbl3", which is an "add" operation.

[0080] In at least one embodiment, translator 130 initiates the first inner loop iteration 1004A by decoding the current bytecode operation 1011 at bytecode index "7". In at least one embodiment, the current bytecode operation 1011 has an offset 1021 (e.g., the same as bytecode index "7"), a label 1031 ("lbl3"), and a non-branching action statement 1041 ("add"). In at least one embodiment, translator 130 determines whether the current bytecode operation 1011 contains a branching action statement, a non-branching action statement, or an operand instruction. Since the current bytecode operation 1011 contains the non-branching action statement 1041 ("add"), translator 130 includes the non-branching action statement 1041 in the output instruction 1092 (e.g., in...). Figure 12(at box 1232). In at least one embodiment, the output instruction 1092 includes a label 1094 corresponding to the name of the pop re-parsing point 762A as described herein. In at least one embodiment, in order to include the non-branch action statement 1041 in the output instruction 1092, the translator 130 includes in the output instruction 1092 an operation name 1096 (“add”) corresponding to the non-branch action statement 1041 (“add”). In at least one embodiment, the translator 130 includes one or more operands 1098 in the output instruction 1092. In at least one embodiment, in order to include operands 1098 in the output instruction 1092, the translator retrieves and removes each operand instruction 1051A, 1051B from the operand instruction stack 1051 (e.g., in FIFO order) and adds each operand instruction 1051A, 1051B as operands 1098A, 1098B to the output instruction 1092. In at least one embodiment, operation 1096 is "add", operand 1098A is "push constant 1", and operand 1098B is "push constant 2". In at least one embodiment, operands "push constant 1" and "push constant 2" can be further simplified to "1" and "2", respectively. In at least one embodiment, the output instruction 1092 is displayed as a label and instruction combination 1099 ("Point3: Add(push constant 1)(push constant 2)").

[0081] In at least one embodiment, the first inner loop iteration 1004A ends, and the re-parse point stack 1061 does not contain the popped re-parse point 762A (“Point 3”) that has been removed from the stack 1061. In at least one embodiment, the translator 130 attempts to decode another operation to process in another inner loop iteration (e.g., in...). Figure 12 (at box 1214), but there are no further operations in input procedure 310, so no further inner loop iterations are performed in the fourth outer loop iteration 1004. In at least one embodiment, the translator begins the fifth outer loop iteration, which refers to Figure 11 As stated above.

[0082] Figure 11An exemplary translation of VM-oriented instructions for an input program 310 according to at least one embodiment is shown, the translation beginning at the fifth outer loop iteration 1104 of the translator 130 to process re-parsing point 962A (“Point 4”). In at least one embodiment, in the fourth outer loop iteration 1104, the translator 130 retrieves and removes the popped re-parsing point 962A (“Point 4”) from the re-parsing point stack 1061 (e.g., in...). Figure 12 (at box 1208). In at least one embodiment, the re-parse point 962A corresponds to the branch target instruction at label "lbl3", and Figure 9 In the second inner loop iteration 902B, it is added to the re-parsing point stack. In at least one embodiment, the translator 130 sets the current label to label 972A (“lbl3”) of the popped re-parsing point 962A, and sets the contents of the operand instruction stack 1151 (e.g., operand instructions) according to the saved operand instruction stack 982A of the popped re-parsing point 962A. Figure 12 (at box 1210). In at least one embodiment, the translator 130 sets the index of the recovered parsed bytecode to the bytecode offset corresponding to the tag 972A (“lbl3”) of the pop re-parsing point 962A, which is “7” in the input program 310 (e.g., in...). Figure 12 (See box 1212). In at least one embodiment, translator 130 outputs a label 1194 corresponding to a pop re-parsing point 962A. In at least one embodiment, the label 1194 is the name of the pop re-parsing point 962A. In at least one embodiment, to output the label 1194, translator 130 generates an output instruction 1192 having the label 1194 (e.g., "Point4") specifying the name of the pop re-parsing point 962A.

[0083] In at least one embodiment, the fifth outer loop iteration 1104 executes the first inner loop iteration 1104A to process the operation of the input program 310, starting from the bytecode index (corresponding to label "lbl3") and using the operand instruction stack 1151 recovered from re-parse point 962A ("Point4"). In at least one embodiment, the first inner loop iteration 1104A processes the operation of the input program 310 at label "lbl3", which is an "add" operation.

[0084] In at least one embodiment, translator 130 initiates the first inner loop iteration 1104A by decoding the current bytecode operation 1111 at bytecode index "7". In at least one embodiment, the current bytecode operation 1111 has an offset 1121 (e.g., the same as bytecode index "7"), a label 1131 ("lbl3"), and a non-branching action statement 1141 ("add"). In at least one embodiment, translator 130 determines whether the current bytecode operation 1111 contains a branching action statement, a non-branching action statement, or an operand instruction. Since the current bytecode operation 1111 contains the non-branching action statement 1141 ("add"), translator 130 includes the non-branching action statement 1141 in the output instruction 1192 (e.g., in...). Figure 12 (at box 1232). In at least one embodiment, the output instruction 1192 includes a label 1194 corresponding to the name of the pop re-parsing point 962A as described herein. In at least one embodiment, in order to include a non-branching action statement 1141 in the output instruction 1192, the translator 130 includes the operation name 1196 ("add") corresponding to the non-branching action statement 1141 ("add") in the output instruction 1192. In at least one embodiment, the translator 130 includes one or more operands 1198 in the output instruction 1192. In at least one embodiment, in order to include operands 1198 in the output instruction 1192, the translator retrieves and removes individual operand instructions 1151A, 1151B from the operand instruction stack 1151 (e.g., in FIFO order) and adds each operand instruction 1151A, 1151B as operands 1198A, 1198B to the output instruction 1192. In at least one embodiment, operation 1196 is "add", operand 1198A is "push constant 1", and operand 1198B is "push constant 3". In at least one embodiment, operands "push constant 1" and "push constant 3" can be further simplified to "1" and "3", respectively. In at least one embodiment, the output instruction 1192 is displayed as a label and instruction combination 1199 ("Point4: Add(push constant 1)(push constant 3)").

[0085] In at least one embodiment, the first inner loop iteration 1104A ends, and the re-parse point stack (not shown) does not contain the popped re-parse point 962A (“Point4”), which has been removed from the re-parse point stack. In at least one embodiment, the translator 130 attempts to decode another operation to process in another inner loop iteration (e.g., in...). Figure 12 (at box 1214), but there are no further operations in input program 310, so no further inner loop iterations are performed in the fourth outer loop iteration 1104. In at least one embodiment, since the re-parsing point stack is empty, no further transformations are performed in this example. In at least one embodiment, Figures 4 to 11 The output instructions generated by the first to fourth outer loop iterations are in Figure 3 The intermediate program 314 is shown as such. In at least one embodiment, the intermediate program 314 can be translated into an optimized intermediate program 316 using suitable optimization techniques. In at least one embodiment, the optimized intermediate program 316 can be translated into a GPU program (suitable for running on a GPU) using suitable translation and / or compilation techniques.

[0086] Figure 12 A process 1200 for translating a VM program into a GPU program according to at least one embodiment is illustrated. In at least one embodiment, by executing process 1200, a processor including one or more circuits or a system including one or more processors performs the operations described herein, such as causing one or more instruction operand stack structures to be copied if one or more branch instructions use one or more instruction operands stored in the one or more stack structures. In at least one embodiment, by executing process 1200, a processor including one or more circuits or a system including one or more processors performs the operations described herein, such as translating VM-oriented (e.g., stack-based) instructions into GPU-oriented instructions. In at least one embodiment, reference is made to... Figures 1 to 11 The system and components described herein may perform part or all of process 1200 or be integrated into process 1200. In at least one embodiment, reference is made to... Figures 13 to 32C The system and processor, as described in various ways, execute part or all of process 1200.

[0087] In at least one embodiment, part or all of process 1200 (or any other process described herein, or variations and / or combinations thereof) is executed under the control of one or more computer systems including computer-executable instructions and is implemented by hardware, software, or a combination thereof as code (e.g., computer-executable instructions, one or more computer programs, or one or more applications) that executes jointly on one or more processors. In at least one embodiment, the code is stored in the form of a computer program on a computer-readable storage medium comprising a plurality of computer-readable instructions executable by one or more processors. In at least one embodiment, the computer-readable storage medium is a non-transitory computer-readable medium. In at least one embodiment, at least some of the computer-readable instructions available for executing process 1200 are not stored using only transient signals (e.g., propagating transient electrical or electromagnetic transmissions). In at least one embodiment, the non-transitory computer-readable medium does not necessarily include non-transitory data storage circuitry (e.g., buffers, caches, and queues) within a transient signal transceiver. In at least one embodiment, process 1200 is executed at least partially on a computer system such as those described elsewhere in this disclosure. In at least one embodiment, process 1200 is executed by logic (e.g., hardware, software, or a combination of hardware and software). In at least one embodiment, one or more methods of process 1200 are executed in any suitable order, including sequential, parallel, and / or variations thereof, and using any suitable processing unit, such as a CPU, GPGPU, GPU, PPU, and / or variations thereof. In at least one embodiment, part or all of process 1200 is executed by one or more neural networks (e.g., simultaneously). In at least one embodiment, one or more neural networks executing part or all of the steps of process 1200 employ the above-described combination. Figure 1 The architecture described above is implemented.

[0088] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for at least creating an initial re-parse point object having a target label corresponding to a first instruction in the input program, and also having an empty instruction stack 1202.

[0089] In at least one embodiment, the system performing at least a portion of process 1200 includes executable code for at least adding the initial re-parse point object to the re-parse point stack 1204. In at least one embodiment, the re-parse point stack is... Figure 2 The re-parsing point stack 250. In at least one embodiment, the addition of the initial re-parsing point object is performed using software (e.g., code, algorithm).

[0090] In at least one embodiment, the system performing at least a portion of process 1200 includes executable code for at least inferring whether the re-parsing point stack at 1206 is empty (e.g., does not contain re-parsing points). In at least one embodiment, if it is inferred that the re-parsing point stack at 1206 is not empty, then the current re-parsing point is set as a re-parsing point object 1208 from the re-parsing point stack. In at least one embodiment, the re-parsing point object is retrieved from the re-parsing point stack in a first-in-first-out (FIFO) manner. In at least one embodiment, the re-parsing point object is retrieved from the re-parsing point stack in a last-in-first-out (LIFO) manner (e.g., by popping the re-parsing point object from the re-parsing point stack). In at least one embodiment, the re-parsing point object is removed from the re-parsing point stack. In at least one embodiment, the re-parsing point stack is... Figure 2 The re-parsing point stack 250, and the re-parsing point object from the re-parsing point stack is Figure 2 The re-parsing point 252. In at least one embodiment, if it is inferred that the re-parsing point stack described in 1206 is empty, then process 1200 ends.

[0091] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for loading a reparse point tag and instruction stack 1210 from at least the current reparse point. In at least one embodiment, the tag is tag 224, and is from, as... Figure 2 The label 254 of the current re-parsing point shown is loaded (for example, assuming it is as follows). Figure 2 The label 254 of the current re-parsing point is shown. In at least one embodiment, the instruction stack is... Figure 2 The operand instruction stack 270 is loaded from the saved operand instruction stack 256 of the current reparsing point (for example, the saved operand instruction stack 256 of the current reparsing point).

[0092] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for setting at least a bytecode index to the re-parse point label loaded from the current re-parse point and outputting the re-parse point label (e.g., output to an output statement list) 1212. In at least one embodiment, the set bytecode index indicates the position (e.g., offset) of the next bytecode instruction to be decoded (e.g., translated) in the input program.

[0093] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for at least decoding the operation at the bytecode index and advancing the bytecode index (e.g., advancing to the index of the next bytecode instruction to be decoded) 1214. In at least one embodiment, the bytecode operation is bytecode operation 240, which originates from the input program and is located at, for example... Figure 2 The bytecode operation 220 is retrieved at the offset indicated by the bytecode index shown. In at least one embodiment, to decode the bytecode operation 240, the system identifies the offset 222 of the bytecode operation 240 (e.g., the bytecode index), the tag 224 (if present), and the branch action statement 226A, the non-branch action statement 226B, or the operand instruction 228. In at least one embodiment, if there is no operation to be decoded at the bytecode index (e.g., because there are no more operations in the input program), process 1200 ends.

[0094] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for at least inferring whether the operation 1216 specifies a branch action. In at least one embodiment, if it is inferred that the operation 1216 specifies a branch action, then one or more branch target labels 1218 are identified for true and false branches (if present) of the branch action. In at least one embodiment, if it is inferred that the operation 1216 does not specify a branch action, then it is determined whether the operation specifies an operand instruction 1228. In at least one embodiment, the operation is a bytecode operation 220, and if the bytecode operation 220 contains a branch action statement 226A, then the bytecode operation 220 specifies a branch action.

[0095] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for at least creating a true branch re-parse point object having a target label corresponding to the branch target label of the true branch and having a copy 1220 of the current instruction stack. In at least one embodiment, the true branch re-parse point object is a branch action statement 226A having a true branch target label 242 (e.g., ...). Figure 2 The re-parsing point 252 (as shown). In at least one embodiment, the current instruction stack is Figure 2 The operand instruction stack 270. In at least one embodiment, if the branch action statement 226A is a conditional branch or an unconditional branch (e.g., an unconditional branch has a single target label, and the true branch re-parsing point object is created for the single target label), then the true branch re-parsing point object is created.

[0096] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for at least creating a false branch re-parse point object, which has a target label corresponding to the branch target label of the false branch and has a copy 1222 of the current instruction stack. In at least one embodiment, the false branch re-parse point object is a branch action statement 226A with a false branch target label 244 (e.g., ...). Figure 2 The re-parse point 252 (as shown). In at least one embodiment, if the branch action statement 226A is a conditional branch, the pseudo-branch re-parse point object is created. In at least one embodiment, if the branch action statement 226A is an unconditional branch, the pseudo-branch re-parse point object is created.

[0097] In at least one embodiment, if one or more branch instructions use one or more instruction operands stored in one or more stack structures, creating the dummy branch reparsing point object with a copy of the current instruction stack 1222 causes the operand stack structure to be copied. In at least one embodiment, branch instructions (such as...) Figure 3 The conditional branches (“cond branch lbl1: lbl2”) of the VM program 310 use the instruction operand “pushconstant 1” (or in other words, “constant 1”) stored on the VM stack. In at least one embodiment, process 1200 makes copies of the instruction operand for both the true branch and the false branch. In at least one embodiment, a copy of the instruction operand for the false branch is made and stored in a copy of the current instruction stack, which is stored in the reparse point object corresponding to the false branch.

[0098] In at least one embodiment, if a copy of the current instruction stack containing the instruction operand (e.g., "push constant1") is not made, then when translating VM program 310, process 1200 removes the instruction operand from operand instruction stack 270 when translating the true branch of the conditional branch instruction (the 7th line addition instruction in the true branch, which is reachable via lbl1 followed by lbl3, will have operands "push constant 1" and "push constant 2" removed from operand instruction stack 270). However, in at least one embodiment, if a copy of the current instruction stack is not made, then when process 1200 continues translating the VM program 310, at the moment of converting the false branch of the conditional branch instruction (the 7th line addition instruction in the false branch, reachable via lbl2 after lbl2, attempts to remove operands "pushconstant 1" and "push constant 3" from the operand instruction stack, but operand "push constant 1" no longer exists in the operand instruction stack 270 because it was removed during the translation of the true branch), in at least one embodiment, if a copy of the current instruction stack containing the instruction operand (e.g., "push constant 1") is made before translating the false branch, for example by creating the false branch re-parsing point object containing the copy of the current instruction stack, 1222 is implemented.

[0099] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for adding at least the true branch re-parse point object and the false branch re-parse point object to the re-parse point stack 1224. In at least one embodiment, the false branch re-parse point object is a branch action statement 226A with a false branch target label 244 (e.g., ...). Figure 2 The re-parsing point 252 (as shown). In at least one embodiment, if the branch action statement 226A is an unconditional branch, the pseudo-branch re-parsing point object is not added to the re-parsing point stack. In at least one embodiment, the re-parsing point stack is... Figure 2 The re-parsing point stack 250, and the true branch re-parsing point object and the false branch re-parsing point object respectively correspond to Figure 2 The reanalyzed points are shown as 252A and 252N.

[0100] In at least one embodiment, the re-parsing point object is retrieved from the re-parsing point stack in a first-in, first-out (FIFO) manner. In at least one embodiment, the re-parsing point object is retrieved from the re-parsing point stack in a last-in, first-out (LIFO) manner (e.g., by popping the re-parsing point object from the re-parsing point stack).

[0101] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for at least adding the branch operation to output statement list 1226. In at least one embodiment, the output statement list is... Figure 2 GPU instruction 132.

[0102] In at least one embodiment, the system executing at least a portion of process 1200 includes executable code for at least inferring whether the operation 1228 specifies an operand instruction (e.g., pushing a value onto or popping a value from the VM stack). In at least one embodiment, if 1228 infers that the operation specifies an operand instruction, then the operand instruction is added to the operand instruction stack 1230. In at least one embodiment, the operand instruction stack is... Figure 2 Operand instruction stack 270. In at least one embodiment, after the operand instruction is added to the operand instruction stack 1230, process 1200 continues by decoding another operation at the bytecode index 1214, advancing the bytecode index, and further continues as described above.

[0103] In at least one embodiment, if the operation described in inference 1228 does not specify an operand instruction, then the operation specifies an action statement, and the operation is added to the output statement list 1232. In at least one embodiment, the operation is added to the output statement list 1218 using an operand instruction retrieved and removed from the operand instruction stack (e.g., as an operand of the operation). In at least one embodiment, after the operand instruction is added to the output statement list 1232, process 1200 continues by decoding another operation at the bytecode index of 1214, advancing the bytecode index, and further continues as described above.

[0104] In at least one embodiment, condition box 1206 corresponds to the start of an outer loop iteration, and each outer loop iteration processes a re-parse point 252 from the re-parse point stack 250. In at least one embodiment, an initial re-parse point corresponding to the first bytecode operation of the input program is created 1202 and added to the re-parse point stack 1204. In at least one embodiment, a subsequent re-parse point is created for each branch operation in the input program and added to the re-parse point stack 250. In at least one embodiment, each outer loop iteration ends when: a branch instruction is added to the output statement list 1226, or if no more bytecode operations are available for decoding at the current bytecode index 1214 (e.g., if the end of the input program is reached).

[0105] In at least one embodiment, condition 1216 corresponds to the start of an inner loop iteration, and each inner loop iteration processes bytecode operation 220 starting at re-parse point 252. In at least one embodiment, as described herein, there is an initial re-parse point corresponding to a first instruction in the input program, and subsequent re-parse points corresponding to branch instructions in the input program. In at least one embodiment, each inner loop iteration ends when: an operand instruction 1230 is added to the operand instruction stack, an operation 1232 is added to the output statement list, or if no more bytecode operations are available for decoding at the current bytecode index 1214 (e.g., if the end of the input program has been reached).

[0106] Figure 13 An example of a system 1300 according to at least one embodiment is shown. This system may include software and hardware to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or to perform any of the operations described herein. System 1300 may include a storage device 1302 and a processor 1308. Storage device 1302 may include, for example, memory, cache, or other storage devices further described herein. Storage device 1302 may be separate from processor 1308, or storage device 1302 may be included in processor 1308 (e.g., in storage device 1312). In at least one embodiment, software program 1304 and / or software library (or instructions) may be stored in memory, cache, or other storage devices and provided to processor 1308 to cause one or more circuits of processor 1308 to perform the operations described herein. In at least one embodiment, software program 1304 and / or software library (or instructions) 1306 may be integrated into one or more circuits of processor 1308. Software program 1304 that can be used to perform any of the operations described herein may be stored on storage device 1302.

[0107] In at least one embodiment, software program 1304 may include one or more software modules. In at least one embodiment, software program 1304 includes an instruction translation module. In at least one embodiment, the instruction translation module is a module that, when executed, causes one or more processors to translate one or more processor instructions that use a stack to process instruction operands into one or more processor instructions that use registers and / or memory addresses to process instruction operands and / or perform the operations described herein. In at least one embodiment, the instruction translation module causes one or more processors 1308 to execute one or more processes described herein, which is implemented at least by including or otherwise encoding related instructions that cause one or more processors 1308 to execute the one or more processes, or are otherwise available for execution by one or more processors 1308.

[0108] In at least one embodiment, as used in any implementation described herein, unless the context explicitly states otherwise, a module refers to any combination of software logic, firmware logic, hardware logic, and / or circuitry configured to provide the functionality described herein. In at least one embodiment, software is embodied as a software package, code, and / or instruction set or instructions, and "hardware" as used in any implementation described herein includes, for example, hardwired circuitry, programmable circuitry, state machine circuitry, fixed-function circuitry, execution unit circuitry, and / or firmware storing instructions executed by programmable circuitry, either individually or in any combination. In at least one embodiment, modules are collectively or individually embodied as circuitry forming part of a larger system, such as an integrated circuit (IC), a system-on-a-chip (SoC), etc. In at least one embodiment, a module performs one or more methods in conjunction with any suitable processing unit and / or combination of processing units (such as one or more CPUs, GPUs, GPGPUs, PPUs, and / or variants thereof, including those further described herein).

[0109] In at least one embodiment, the software program 1304 may include a series of software codes, commands, instructions, or other text sequences to instruct the computing device to perform one or more computing operations and / or invoke one or more other instruction sets, such as APIs or API functions or instruction set architecture (ISA) level instructions, to be run or executed.

[0110] In at least one embodiment, system 1300 may include one or more software libraries 1306, which may, for example, provide one or more API and / or ISA instructions. In at least one embodiment, if one or more branch instructions use one or more instruction operands stored in the one or more stack structures, the one or more API and / or ISA instructions may be used to cause the one or more instruction operand stack structures to be copied. In at least one embodiment, one or more software libraries 1306 may be included in a driver and / or runtime. In at least one embodiment, software library 1306 (e.g., including one or more API and / or ISA instructions) may include multiple sets of software instructions that, if executed or otherwise implemented, cause processor 1308 to perform one or more computational operations, such as any operations described herein. In at least one embodiment, one or more API and / or ISA instructions may be distributed or otherwise provided as part of one or more software libraries 1306, runtimes, drivers, and / or any other grouping of software and / or executable code further described herein. In at least one embodiment, one or more API and / or ISA instructions may perform one or more computational operations in response to a call from software program 1304.

[0111] Processor 1308 may include any number of processors and any suitable processing units and / or combinations of processing units, such as, but not limited to, a central processing unit (“CPU”), a graphics processing unit (“GPU”), or other processors (including accelerators, field-programmable gate arrays (FPGAs), GPUs, parallel processors, GPGPUs, DPUs, and / or variations thereof, including those further described herein), including any processors described herein, such as, but not limited to, those described herein. Figures 15 to 27BThe processor 1308 is a processor in the memory. In at least one embodiment, the processor 1308 may, for example, use instruction fetch 1316 (e.g., for an instruction fetch phase) to retrieve or fetch instructions (e.g., one or more API and / or ISA instructions) from the memory device 1302. Instructions may include instructions that cause one or more instruction operand stack structures to be copied if one or more branch instructions use one or more instruction operands stored in the one or more stack structures. In at least one embodiment, the processor 1308 may include a memory device 1312 and an instruction queue 1310 to store and queue instructions fetched from the memory device 1302. In at least one embodiment, the fetched instructions may be decoded by decode 1318 to determine what operation the processor 1308 should perform (e.g., in an instruction decoding phase). In at least one embodiment, the processor 1308 may fetch additional operands (data) that can be used for instructions, and the operands may be stored, for example, in a register or memory device 1312. In at least one embodiment, micro-operation 1320 may perform operations on the data stored in one or more registers or memory device 1312. For example, each step of the instruction fetched by processor 1308 can be broken down during execution, allowing processor 1308 to execute the instruction step by step through a series of micro-operations 1320. In at least one embodiment, program counter (PC) 1314 can store the address of the next instruction and can be updated to point to the next instruction to be executed by processor 1308.

[0112] In at least one embodiment, processor 1308 may execute instructions (e.g., during execution phase). For example, processor 1308 may perform operations specified by the instructions, such as arithmetic operations, logical operations, or data transfers. In at least one embodiment, computing unit 1322 may execute instructions to perform any of the operations described herein. In at least one embodiment, computing unit may include ALU 1324 (Arithmetic Logic Unit), which may be used to perform arithmetic and logical operations. In at least one embodiment, computing unit may include FPU (Floating Point Unit) 1326, which may be used to perform floating-point calculations. In at least one embodiment, other circuitry 1328 may be used to perform other operations, such as vector and / or scalar operations. In at least one embodiment, accelerator 1330 may include one or more matrix multiplication accelerators, one or more parallel processing units (PPUs) (such as GPUs), or any other accelerators or processors further described herein. In at least one embodiment, software program 1304 may utilize one or more API and / or ISA instructions to perform various computational operations, such as matrix multiplication, arithmetic operations, or any other computational operations further described herein, via accelerator 1330. In at least one embodiment, one or more computational operations using accelerator 1330 may include at least one or more sets of computational operations that will be accelerated by being executed at least partially by accelerator 1330, including: if one or more branch instructions use one or more instruction operands stored in the one or more stack structures, then causing one or more instruction operand stack structures to be copied.

[0113] In at least one embodiment, system 1300 can be used to execute one or more instructions comprising functions or operations, such as combining Figures 1 to 12 The instructions described herein. In at least one embodiment, a system 1300 including one or more processors, if one or more branch instructions use one or more instruction operands stored in the one or more stack structures, causes one or more circuits to copy the one or more instruction operand stack structures and / or perform other operations described herein. In at least one embodiment, system 1300 is included in Figures 1 to 12 The system shown contains and / or includes Figures 1 to 12 The system shown, if one or more branch instructions use one or more instruction operands stored in the one or more stack structures, then system 1300 causes one or more circuits to copy the one or more instruction operand stack structures and / or perform other operations described herein. In at least one embodiment, system 1300 includes Figures 14 to 32CThe one or more hardware components shown are used to copy one or more instruction operand stack structures and / or perform other operations described herein if one or more branch instructions use one or more instruction operands stored in one or more stack structures.

[0114] Data Center

[0115] Figure 14 An example data center 1400 according to at least one embodiment is illustrated. The data center 1400 may include one or more rooms having racks 1402 and auxiliary equipment for housing one or more racks 1402 and one or more substrates 1404. Racks 1402 may include one or more substrates 1404. Racks 1402 may include enclosures for housing and supporting individual substrates 1404. Operational aspects of racks 1402 may be adjustable at the rack level (corresponding to a group of substrates 1404) or at the substrate level (corresponding to an individual substrate 1404), among other options. Racks 1402 or substrates 1404 may have specific selected maximum operating parameters, such as, but not limited to, power consumption, operating frequency, etc. The data center 1400 may be supported by various cooling systems, such as, but not limited to, cooling towers, cooling loops, pumps, and other support systems. The cooling system may include sensors and controllers for monitoring and managing the cooling characteristics of racks 1402. The substrates 1404 within racks 1402 may draw operating power from one or more power distribution units (PDUs; not shown). PDUs can be arranged within racks 1402, for example, between racks 1402 that include substrates 1404, or within racks 1402 that also house substrates 1404.

[0116] Rack 1402 and substrate 1404 may include subsystems, modules, add-in cards, and other semiconductor components. Substrate 1404 may include one or more computing units 1406, each computing unit 1406 may include one or more processors 1408, one or more memories 1410, and an interface controller 1412. Computing unit 1406 may include any number of processors, such as, but not limited to, a central processing unit (“CPU”), a graphics processing unit (“GPU”), or other processors (including accelerators, field-programmable gate arrays (FPGAs), graphics processors, etc.), including any processors described herein, such as, but not limited to, those described herein. Figure 15-27BThe processor in the computing unit 1406 may include one or more memory storage devices 1410 (e.g., dynamic read-only memory, solid-state storage devices, or disk drives), as well as network input / output (“NW I / O”) devices, network switches, virtual machines (“VMs”), power supply modules, and cooling modules, etc. One or more computing units 1406 may be a server having one or more of the aforementioned computing resources.

[0117] Computing unit 1406 may include individual computing unit groups housed in one or more racks (not shown), or in numerous racks within data centers in different geographical locations (also not shown). Individual computing unit groups may include grouped computing, networking, memory, or storage resources that can be configured or allocated to support one or more workloads. Several computing units (e.g., including CPUs and / or other processors) may be grouped within one or more racks to provide computing resources to support one or more workloads. Resource coordinator 1414 may configure or otherwise control one or more computing units 1406 or computing unit groups. Resource coordinator 1414 may include a Software Design Infrastructure (“SDI”) management entity for data center 1400. Resource coordinator 1414 may include hardware, software, or some combination thereof.

[0118] Data center 1400 may include any one or any combination of framework layer 1420, software layer 1430, and application layer 1440. For example... Figure 14 As shown, framework layer 1420 includes a job scheduler 1422, a configuration manager 1424, a resource manager 1426, and a distributed file system 1428. Framework layer 1420 may include a framework for supporting software 1432 of software layer 1430 and / or one or more applications 1442 of application layer 1440. Software 1432 or application 1442 may respectively include web-based service software or applications, such as, but not limited to, software or applications provided by Amazon Web Services, Google Cloud, and Microsoft Azure. Framework layer 1420 may be a type of free and open-source software web application framework, such as, but not limited to, Apache Spark. TM(Hereinafter referred to as "Spark"), which can utilize the distributed file system 1428 for large-scale data processing (e.g., "big data"). The job scheduler 1422 may include Spark drivers, which facilitate the scheduling of workloads supported by various layers of the data center 1400. The configuration manager 1424 may be able to configure different layers, such as, but not limited to, the software layer 1430 and the framework layer 1420 (which includes Spark and the distributed file system 1428 for supporting large-scale data processing). The resource manager 1426 may be able to manage clustered or grouped compute units 1406 mapped to or allocated to support the distributed file system 1428 and the job scheduler 1422. The resource manager 1426 may coordinate with the resource coordinator 1414 to manage these mapped or allocated compute resources.

[0119] Software 1432 may be included in software layer 1430, and may include software used by at least a portion of computing units 1406, one or more computing units 1406, groups of computing units 1406, and / or the distributed file system 1428 of framework layer 1420. One or more types of software may include, but are not limited to, internet web search software, email virus scanning software, database software, and streaming video content software.

[0120] Application 1442 may be included in application layer 1440 and may include one or more types of applications used by at least portions of computing unit 1406, one or more computing units 1406, groups of computing units 1406, and / or the distributed file system 1428 of framework layer 1420. One or more types of applications may include, but are not limited to, any number of genomics applications, cognitive computing applications, and machine learning applications, including training or inference software, machine learning framework software (e.g., PyTorch, TensorFlow, Caffe, etc.), or other machine learning applications used in conjunction with one or more embodiments.

[0121] Any of the Configuration Manager 1424, Resource Manager 1426, and Resource Coordinator 1414 can implement any number and type of self-modification actions based on any amount and type of data obtained in any technically feasible manner. Self-modification actions can alleviate the burden on data center operators of data center 1400 to make potentially erroneous configuration decisions and may prevent underutilized and / or poorly performing portions of the data center.

[0122] Data center 1400 may include tools, services, software, or other resources for training one or more machine learning models according to one or more embodiments described herein, or for using one or more machine learning models to predict or infer information. For example, a machine learning model can be trained by calculating weight parameters based on a neural network architecture using the software and computing resources described above regarding data center 1400. The trained machine learning model corresponding to one or more neural networks can be used with the resources described above regarding data center 1400 to infer or predict information using weight parameters calculated through one or more training techniques described herein.

[0123] Data Center 1400 can use CPUs, application-specific integrated circuits (ASICs), GPUs, FPGAs, or other hardware (e.g., Figure 15-27B The embodiments described herein can be used to perform some or all of the processes and techniques described elsewhere, such as, but not limited to, training and / or inference using the resources described above. Furthermore, one or more of the software and / or hardware resources described above can be configured as a service to allow a user to train or perform information inference, such as, but not limited to, image recognition, speech recognition, or other artificial intelligence services.

[0124] In at least one embodiment, processor 1408 may include one or more processors and / or circuitry for causing one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise performing any operation described above or elsewhere herein. In at least one embodiment, processor 1408 is configured by software 1432 to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any operation described above or elsewhere herein. Data center 1400 may use logic, CPU, application-specific integrated circuit (ASIC), GPU, FPGA, or other hardware (e.g., Figure 15-27B (The embodiments in the document) can be used to perform any of the operations described above or elsewhere in this document.

[0125] processor

[0126] The figure below illustrates, without limitation, example processors and processing systems that can be used to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform some or all of the processes, operations, and / or techniques described elsewhere herein. The example processors and processing systems can be software-configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. Processors and processing systems can include logic, central processing units (CPUs), application-specific integrated circuits (ASICs), graphics processing units (GPUs), field-programmable gate arrays (FPGAs), XPUs (i.e., any computing architecture best suited to the needs of the application), or other hardware (e.g., Figure 15-27B The processor and / or processing system described herein may include one or more circuitry configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. As used herein, one or more circuitry may be software-configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. Figure 32A and Figure 32B Logic 3215, as described elsewhere herein, is illustrated and can be used in one or more devices to perform, for example (but not limited to), the operations discussed herein according to at least one embodiment. Logic can refer to, for example, any combination of software logic, hardware logic, and / or firmware logic for providing the functionality and / or operations described herein, wherein the logic can be collectively or individually embodied as circuitry forming part of a larger system, such as an integrated circuit (IC), an application-specific integrated circuit (ASIC), a field-programmable gate array (FPGA), a system-on-a-chip (SoC), or one or more processors (e.g., a CPU, a GPU).

[0127] Figure 15A processor according to at least one embodiment is illustrated, which is a system-on-a-chip (SOC) 900 (which may be referred to as a system-on-a-chip, superchip, or other names). SOC 1500 may include processor complex 1510 and processor complex 1540. SOC 1500 may include any number of processor complexes 1510 and / or processor complexes 1540, which may include any number of processors described herein in any combination, such as, but not limited to, processors described herein in any combination. Figure 15-27B The processor in the system. For example, processor 1510 may include a central processing unit (CPU), and processor 1540 may include a graphics processor. Alternatively, processor 1510 may include a graphics processor, and processor 1540 may include a graphics processor. The SOC 1500 may include any number of display controllers 1592, any number of multimedia engines 1594, any number of I / O interfaces 1570, any number of memory controllers 1580, and any number of fabrics 1560 in any combination. For ease of explanation, this document uses reference numbers identifying objects and bracket numbers identifying instances (if needed) to denote multiple instances of similar objects. The SOC 1500 may include a processor from Broadcom Inc., Palo Alto, California.

[0128] Processor complex 1510 may include a CPU, processor complex 1540 may include a GPU, and SOC 1500 may be a processing unit integrating processor complex 1510 and processor complex 1540 onto a single chip. Certain tasks may be assigned to processor complex 1510, while other tasks may be assigned to processor complex 1540. Processor complex 1510 may be configured to execute main control software associated with SOC 1500, such as, but not limited to, an operating system. Processor complex 1510 may be the main processor of SOC 1500, controlling and coordinating the operation of other processors. Processor complex 1510 may issue commands that control the operation of processor complex 1540 to perform some or all of the operations described herein. Processor complex 1510 may be configured to execute host-executable code derived from CUDA or other source code (e.g., HIP source code), while processor complex 1540 may be configured to execute device-executable code derived from CUDA or other source code to perform any of the operations described herein.

[0129] Processor complex 1510 may include cores 1520(1)-1520(4) and cache (e.g., L3 cache) 1530 for storing information for performing the operations described herein. Processor complex 1510 may include any number of cores 1520 in any combination and any number and type of cache. Cores 1520 may be configured to execute instructions of a specific instruction set architecture (“ISA”) to perform some or all of the operations described herein. Each core 1520 may include a CPU core. Cores 1520(1)-920(4) may be referred to as compute units or arithmetic units. SOC 1500 may include any number of processor complexes 1510, architecture 1560, I / O interface 1570, and memory controller 1580.

[0130] Each core 1520 may include a fetch / decode unit 1522, an integer execution engine 1524, a floating-point execution engine 1526, and an L2 cache 1528. The fetch / decode unit 1522 may fetch instructions to perform some or all of the operations described herein (e.g., but not limited to APIs compiled into instructions) and decode those instructions, generate micro-operations, and dispatch individual micro-instructions to the integer execution engine 1524 and / or the floating-point execution engine 1526. The fetch / decode unit 1522 may concurrently dispatch one micro-instruction to the integer execution engine 1524 and another micro-instruction to the floating-point execution engine 1526. The integer execution engine 1524 may perform integer and memory operations. The floating-point engine 1526 may perform floating-point and vector operations. The fetch / decode unit 1522 may dispatch micro-instructions to one or more execution engines, which may replace both the integer execution engine 1524 and the floating-point execution engine 1526.

[0131] Each core 1520(i) (where i is an integer representing a specific instance of core 1520) can access the L2 cache 1528(i) included in core 1520(i). Each core 1520 included in core complex 1510(j) (where j is an integer representing a specific instance of core complex 1510) can be connected to other cores 1520 included in core complex 1510(j) via the L3 cache 1530(j) included in core complex 1510(j). The cores 1520 included in core complex 1510(j) (where j is an integer representing a specific instance of core complex 1510) can access all L3 caches 1530(j) included in core complex 1510(j). The L3 cache 1530 can include any number of slices.

[0132] Processor complex 1540 may be a graphics complex that can be configured to perform computational operations (e.g., the computational operations described herein) in a highly parallel manner. Processor complex 1540 may be configured to perform graphics pipeline operations, such as, but not limited to, drawing commands, pixel operations, geometric calculations, and other operations associated with rendering an image to a display. Processor complex 1540 may be configured to perform graphics-independent operations, such as, but not limited to, neural network training and / or simulation. Processor complex 1540 may be configured to perform both graphics-related and graphics-independent operations.

[0133] Processor complex 1540 may include any number of compute units 1550(1)-1550(N) (where N is any integer greater than 1) and L2 cache 1542. The compute units 1550 may share L2 cache 1542, which may store information that will be used to perform some or all of the operations described herein. L2 cache 1542 may be partitioned. Processor complex 1540 may include any number of compute units 1550 and any number (including zero) and type of cache. Processor complex 1540 may include any number of dedicated graphics hardware.

[0134] Each compute unit 1550 may include any number of SIMD units 1552(1)-1552(N) (where N is any integer greater than 1) and shared memory 1554. Each SIMD unit 1552 may implement a SIMD architecture and may be configured to perform some or all of the operations described herein in parallel. Each compute unit 1550 may execute any number of thread blocks, but each thread block may execute on a single compute unit 1550, although in some embodiments, the thread block may execute on multiple compute units. A thread block may include any number of execution threads. A workgroup may be a thread block. Each SIMD unit 1552 may execute a set of threads. A set of threads (e.g., 16 threads), also referred to as a warp, subgroup, or wavefront (e.g., used by AMD and Intel), may belong to a single thread block and be configured to process different datasets based on a single instruction set. Prediction may be used to disable one or more threads in a warp, subgroup, or wavefront. A lane may be a thread. A work item can be a thread, such as (but not limited to) an OpenCL thread. Different thread bundles, subgroups, or wavefronts within a thread block can be synchronized together and communicate via shared memory 1554. Each compute unit 1550 can include one or more thread block clusters, where thread block clusters can implement programmable control over locality at a larger granularity than a single thread block of a single streaming multiprocessor (SM). Thread block clusters (also referred to as “clusters”) can support multiple thread blocks running concurrently across streaming multiprocessors, thereby synchronously and cooperatively acquiring, exchanging, or otherwise using data. In at least one embodiment, a streaming multiprocessor (“SM”) can refer to a streaming microprocessor, a streaming processor (“SP”), a streaming processing unit (“SPU”), a compute unit (“CU”), an execution unit (“EU”), and / or a slice, where a slice in this context can refer to a portion of the processing resources within a processing unit (e.g., 16 cores, a ray tracing unit, a thread bootstrap, or a scheduler).

[0135] Structure 1560 may be a system interconnect that facilitates data and control transfers across processor complex 1510, processor complex 1540, I / O interface 1570, memory controller 1580, display controller 1592, and multimedia engine 1594, for example, to perform some or all of the operations described herein. SOC 1500 may include any number and type of system interconnects other than or replacing structure 1560, facilitating data and control transfers across any number and type of directly or indirectly linked components, either internally or externally to SOC 1500. I / O interface 1570 may represent any number and type of I / O interfaces (e.g., PCI, PCI expansion (“PCI-X”), PCIe, Gigabit Ethernet (“GBE”), USB, etc.). Various types of peripheral devices may be coupled to I / O interface 1570. Peripheral devices that may be coupled to I / O interface 1570 may include keyboards, mice, printers, scanners, joysticks or other types of game controllers, media recording devices, external storage devices, network interface cards, etc.

[0136] Display controller 1592 can display images on one or more display devices, such as, but not limited to, liquid crystal displays (“LCD”) devices. Multimedia engine 1594 can include any number and type of circuitry related to multimedia, such as, but not limited to, video decoders, video encoders, image signal processors, etc. Memory controller 1580 can facilitate data transfer between SOC 1500 and unified system memory 1590. Processor complex 1510 and processor complex 1540 can share unified system memory 1590. Unified system memory 1590 can include various types of memory devices, including dynamic random access memory (DRAM) or graphics random access memory, such as, but not limited to, synchronous graphics random access memory (SGRAM), including graphics double data rate (GDDR) memory. Unified system memory 1590 can include 3D stacked memory, including but not limited to high bandwidth memory (HBM), HBM2e, or HDM3.

[0137] The SOC 1500 can implement a memory subsystem including any number and type of memory controllers 1580 and memory devices (e.g., shared memory 1554), which may be dedicated to a single component or shared among multiple components to perform any of the operations described herein. The SOC 1500 can implement a cache subsystem including one or more cache memories (e.g., L2 cache 1528, L3 cache 1530, and L2 cache 1542), each cache memory may be dedicated to any number of components (e.g., core 1520, core complex 1510, SIMD unit 1552, compute unit 1550, and processor complex 1540), or may be shared among any number of components (e.g., core 1520, core complex 1510, SIMD unit 1552, compute unit 1550, and processor complex 1540).

[0138] In at least one embodiment, the SOC 1500 may include one or more circuitry for causing one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise performing any of the operations described above or elsewhere herein. One or more circuitry may be software-configurable for causing one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise performing any of the operations described above or elsewhere herein.

[0139] Figure 16A A parallel processor 1600 according to at least one embodiment is illustrated. The parallel processor 1600 may be implemented using one or more circuits and may be referred to as a programmable processor (e.g., CPU and / or GPU), logic, application-specific integrated circuit (ASIC), field-programmable gate array (FPGA), or other hardware (e.g., Figure 15-27B (The embodiments in this document) are used to perform any of the operations described above or elsewhere herein.

[0140] Parallel processor 1600 may include parallel processing unit 1602 for performing any of the operations described above or elsewhere herein. Parallel processing unit 1602 may include I / O unit 1604 that enables communication with other devices, including other instances of parallel processing unit 1602. I / O unit 1604 may be directly connected to other devices. I / O unit 1604 may be connected to other devices via the use of a hub or switch interface, such as, but not limited to, memory hub 1605. The connection between memory hub 1605 and I / O unit 1604 may form a communication link 1613. I / O unit 1604 may be connected to host interface 1606 and memory crossbar switch 1616, wherein host interface 1606 receives commands directed to perform processing operations, and memory crossbar switch 1616 receives commands directed to perform memory operations.

[0141] When host interface 1606 receives a command buffer via I / O unit 1604, host interface 1606 can route the work operations that execute these commands to front-end 1608. Front-end 1608 can be coupled to scheduler 1610 (which may be referred to as sequencer), which is configured to distribute commands or other work items to processing cluster array 1612. Scheduler 1610 can ensure that processing cluster array 1612 is correctly configured and in an active state before tasks are distributed to the cluster of processing cluster array 1612. Scheduler 1610 can be implemented via firmware logic executed on a microcontroller. The microcontroller-implemented scheduler 1610 can be configured to perform complex scheduling and work distribution operations at both coarse and fine granular levels, thereby enabling fast preemption and context switching of threads executing on processing array 1612. Host software can validate workloads scheduled on processing cluster array 1612 via one of multiple graphics processing paths. The workload can then be automatically distributed to the processing array cluster 1612 by the scheduler 1610 logic within the microcontroller, which includes the scheduler 1610.

[0142] Processing cluster array 1612 can perform any of the operations described above or elsewhere herein, and may include up to “N” processing clusters (e.g., clusters 1614A, 1614B through 1614N), where “N” represents a positive integer (which may be a different integer “N” than used in other diagrams). Each cluster 1614A-1614N in processing cluster array 1612 can execute a large number of concurrent threads. Scheduler 1610 may use various scheduling and / or work distribution algorithms to distribute work to clusters 1614A-1614N in processing cluster array 1612, which may vary depending on the workload generated by each type of program or computation. Scheduling may be dynamically handled by scheduler 1610 or may be assisted by compiler logic during the compilation of program logic configured to be executed by processing cluster array 1612. Different clusters 1614A-1614N of processing cluster array 1612 may be assigned to process different types of programs or perform different types of computations.

[0143] The processing cluster array 1612 can be configured to perform various types of parallel processing operations, such as, but not limited to, any of the operations described above or elsewhere herein. The processing cluster array 1612 can be configured to perform general-purpose parallel computing operations. For example, the processing cluster array 1612 may include logic for performing processing tasks, including filtering video and / or audio data, performing modeling operations (including physical operations), and performing data transformations.

[0144] Processing cluster array 1612 can be configured to perform parallel graphics processing operations. Processing cluster array 1612 may include additional logic for supporting the execution of such graphics processing operations, including but not limited to texture sampling logic for performing texture operations, as well as tessellation logic and other vertex processing logic. Processing cluster array 1612 can be configured to execute shader programs related to graphics processing, such as, but not limited to, vertex shaders, tessellation shaders, geometry shaders, and pixel shaders. Parallel processing unit 1602 can transfer data from system memory via I / O unit 1604 for processing. During processing, the transferred data may be stored in on-chip memory (e.g., parallel processor memory 1622) during processing and then written back to system memory.

[0145] When the parallel processing unit 1602 is used to perform graphics processing, the scheduler 1610 can be configured to divide the processing workload into tasks of approximately equal size to better distribute graphics processing operations to multiple clusters 1614A-1614N of the processing cluster array 1612. Each part of the processing cluster array 1612 can be configured to perform different types of processing. For example, a first part can be configured to perform vertex shading and topology generation, a second part can be configured to perform tessellation and geometry shading, and a third part can be configured to perform pixel shading or other screen-space operations to produce a rendered image for display. Intermediate data generated by one or more clusters 1614A-1614N can be stored in a buffer to allow intermediate data to be transferred between clusters 1614A-1614N for further processing.

[0146] Processing cluster array 1612 can receive processing tasks to be executed via scheduler 1610, which receives commands defining the processing tasks from front end 1608. Processing tasks may include indexes of data to be processed, such as surface (patch) data, primitive data, vertex data, and / or pixel data, as well as state parameters and commands defining how to process the data (e.g., which program to execute). Scheduler 1610 can be configured to retrieve the index corresponding to the task, or can receive the index from front end 1608. Front end 1608 can be configured to ensure that processing cluster array 1612 is configured to be active before the workload specified by the incoming command buffer (e.g., batch buffer, push buffer, etc.) is initiated.

[0147] Each instance of one or more instances of parallel processing unit 1602 may be coupled to parallel processor memory 1622 to perform any of the operations described above or elsewhere herein. Parallel processor memory 1622 may be accessed via memory crossbar switch 1616, which may receive memory requests from processing cluster array 1612 and I / O unit 1604. Memory crossbar switch 1616 may access parallel processor memory 1622 via memory interface 1618. Memory interface 1618 may include multiple partition units (e.g., partition units 1620A, 1620B through 1620N), each partition unit may be coupled to a portion (e.g., a memory cell) of parallel processor memory 1622. The number of partition units 1620A-1620N can be configured to be equal to the number of memory units, such that the first partition unit 1620A has a corresponding first memory unit 1624A, the second partition unit 1620B has a corresponding memory unit 1624B, and the Nth partition unit 1620N has a corresponding Nth memory unit 1624N. The number of partition units 1620A-1620N may not be equal to the number of memory units.

[0148] Memory cells 1624A-1624N may include various types of memory devices, including dynamic random access memory (DRAM) or graphics random access memory, such as, but not limited to, synchronous graphics random access memory (SGRAM), which includes graphics double data rate (GDDR) memory. Memory cells 1624A-1624N may also include 3D stacked memory, including but not limited to high-bandwidth memory (HBM), HBM2e, or HDM3. Render targets (e.g., but not limited to framebuffers or texture maps) may be stored in memory cells 1624A-1624N, allowing partitioning cells 1620A-1620N to write portions of each render target in parallel to efficiently utilize the available bandwidth of the parallel processor memory 1622. A local instance of the parallel processor memory 1622 may not be included to support a unified memory design that combines system memory with local cache memory.

[0149] Any cluster 1614A-1614N in the processing cluster array 1612 can process data to be written to any memory cell 1624A-1624N within the parallel processor memory 1622. The memory crossbar switch 1616 can be configured to transfer the output of each cluster 1614A-1614N to any partition cell 1620A-1620N, or to another cluster 1614A-1614N on which additional processing operations can be performed. Each cluster 1614A-1614N can communicate with the memory interface 1618 via the memory crossbar switch 1616 to read from or write to various external memory devices. The memory crossbar switch 1616 can be connected to the memory interface 1618 to communicate with the I / O unit 1604, or to a local instance of the parallel processor memory 1622, enabling processing units within different processing clusters 1614A-1614N to communicate with system memory or other memory local to the non-parallel processing unit 1602. The memory crossbar switch 1616 can use virtual channels to separate traffic flows between clusters 1614A-1614N and partition units 1620A-1620N.

[0150] Multiple instances of the parallel processing unit 1602 can be mounted on a single add-in card, or multiple add-in cards can be interconnected. Even if different instances of the parallel processing unit 1602 have different numbers of processing cores, different amounts of local parallel processor memory, and / or other configuration differences, these different instances can be configured to interoperate. For example, some instances of the parallel processing unit 1602 may include higher-precision floating-point units relative to other instances. Systems including one or more instances of the parallel processing unit 1602 or the parallel processor 1600 can be implemented in a variety of configurations and form factors, including but not limited to desktop, laptop, or handheld personal computers, servers, workstations, game consoles, and / or embedded systems.

[0151] Figure 16A It also includes a block diagram of a partitioning unit 1620 according to at least one embodiment. The partitioning unit 1620 is... Figure 16AAn example of one of the partition units 1620A-1620N in the parallel processor memory. Partition unit 1620 may include an L2 cache 1621, a frame buffer interface 1625, and a ROP 1626 (raster operation unit). The L2 cache 1621 may be a read / write cache configured to perform load and store operations received from the memory crossbar switch 1616 and the ROP 1626. Read misses and urgent write-back requests may be output from the L2 cache 1621 to the frame buffer interface 1625 for processing. Updates may also be sent to the frame buffer via the frame buffer interface 1625 for processing. The frame buffer interface 1625 may interface with one of the memory cells in the parallel processor memory, such as, but not limited to, Figure 16A The memory cells 1624A-1624N (shown as 1624) are located (e.g., within the parallel processor memory 1622).

[0152] ROP 1626 can be a processing unit that performs raster operations, such as, but not limited to, stenciling, z-testing, blending, etc. ROP 1626 can then output processed graphics data stored in graphics memory. ROP 1626 may include compression logic for compressing depth or color data written to memory and decompressing depth or color data read from memory. The compression logic can be lossless compression logic that utilizes one or more compression algorithms. The type of compression performed by ROP 1626 can vary based on the statistical characteristics of the data to be compressed. For example, incremental color compression is performed on depth and color data on a per-tile basis.

[0153] ROP 1626 can be included in each processing cluster (e.g., Figure 16A The data is stored within clusters 1614A-1614N, not within partition units 1620. Read and write requests for pixel data (not pixel fragment data) can be transferred via memory crossbar switch 1616. Processed graphics data can be displayed on a monitor and routed for further processing by the processor, or routed to... Figure 16A One of the processing entities within the parallel processor 1600 in the system is further processed.

[0154] In at least one embodiment, the parallel processor 1600 may include one or more circuitry configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuitry may be software-configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0155] Figure 16B A block diagram including a processing cluster 1614 within a parallel processing unit according to at least one embodiment. The processing cluster may be... Figure 16A An instance of one of the processing clusters 1614A-1614N is provided, which can be used to perform any of the operations described above or elsewhere herein. The processing cluster 1614 can be configured to execute many threads in parallel, where a “thread” refers to an instance of a specific program executed on a specific input dataset. Single Instruction Multiple Data (SIMD) instruction issuing techniques can be used to support the parallel execution of a large number of threads without providing multiple independent instruction units. Single Instruction Multiple Thread (SIMT) techniques can be used to support the parallel execution of a large number of typically synchronous threads using a common instruction unit configured to issue instructions to a set of processing engines within each processing cluster.

[0156] The operation of cluster 1614 can be controlled via pipeline manager 1632, which distributes processing tasks to SIMT parallel processors. Pipeline manager 1632 can... Figure 16A The scheduler 1610 receives instructions and manages the execution of these instructions via the graphics multiprocessor 1634 and / or texture unit 1636. The graphics multiprocessor 1634 may be an example instance of a SIMT parallel processor. However, the processing cluster 1614 may include various types of SIMT parallel processors with different architectures. The processing cluster 1614 may include one or more instances of the graphics multiprocessor 1634. The graphics multiprocessor 1634 can process data and can use the data cross switch 1640 to distribute the processed data to one of several possible destinations, including other shader units. The pipeline manager 1632 can facilitate the distribution of processed data by specifying the destination of the processed data to be distributed via the data cross switch 1640.

[0157] Each graphics multiprocessor 1634 within the processing cluster 1614 may include a set of identical functional execution logic (e.g., arithmetic logic units, load-memory units, etc.) for performing computations for any of the operations described above or elsewhere herein. The functional execution logic can be configured in a pipelined manner, where new instructions can be issued before previous instructions complete. The functional execution logic can support a wide range of operations, including integer and floating-point arithmetic, comparison operations, Boolean operations, bit shifting, and computation of various algebraic functions. Different operations can be performed using the same functional unit hardware, and arbitrary combinations of functional units are possible.

[0158] Instructions transmitted to the processing cluster 1614 can form threads, which may also be called thread bundles, subgroups, waves, or wavefronts. A group of threads executing across a set of parallel processing engines can be called a thread group. Thread groups can execute a common program on different input data. Each thread within a thread group can be assigned to a different processing engine within the graphics multiprocessor 1634. The number of threads in a thread group can be less than the number of processing engines within the graphics multiprocessor 1634. When the number of threads in a thread group is less than the number of processing engines, one or more processing engines may be idle during the processing cycle of that thread group. The number of threads in a thread group can also be more than the number of processing engines within the graphics multiprocessor 1634. When the number of threads in a thread group is more than the number of processing engines within the graphics multiprocessor 1634, processing can be performed in consecutive clock cycles. Multiple thread groups can execute concurrently on the graphics multiprocessor 1634.

[0159] The graphics multiprocessor 1634 includes an internal cache memory for performing load and store operations, such as, but not limited to, any of the operations described above or elsewhere herein. The graphics multiprocessor 1634 may forgo the internal cache and instead use a cache memory within the processing cluster 1614 (e.g., L1 cache 1648). Each graphics multiprocessor 1634 may also access partition units that can be shared across all processing clusters 1614 (e.g., ...). Figure 16A The graphics multiprocessor 1634 has an L2 cache within partition units 1620A-1620N and can be used to transfer data between threads. The graphics multiprocessor 1634 can also access off-chip global memory, which may include one or more of the local parallel processor memory and / or system memory. Any memory outside of the parallel processing unit 1602 can be used as global memory. The processing cluster 1614 may include multiple instances of the graphics multiprocessor 1634 and can share common instructions and data, which can be stored in the L1 cache 1648.

[0160] Each processing cluster 1614 may include an MMU 1645 (Memory Management Unit), which can be configured to map virtual addresses to physical addresses. One or more instances of the MMU 1645 may reside in... Figure 16A The MMU 1645 is located within the memory interface 1618. It may include a set of page table entries (PTEs) for mapping virtual addresses to physical addresses on tiles, and optional cache line indexes. The MMU 1645 may include address translation lookup buffers (TLBs) or caches that may reside within the graphics multiprocessor 1634 or L1 1648 cache, or within the processing cluster 1614. Physical addresses can be processed to distribute surface data access locally, allowing for efficient request interleaving between partition units. The cache line indexes can be used to determine whether a request for a cache line is a hit or a miss.

[0161] Processing cluster 1614 can be configured such that each graphics multiprocessor 1634 is coupled to a texture unit 1636 to perform texture mapping operations, such as determining texture sample locations, reading texture data, and filtering texture data. Texture data can be read from an internal texture L1 cache (not shown) or an L1 cache within the graphics multiprocessor 1634, and can be retrieved as needed from an L2 cache, local parallel processor memory, or system memory. Each graphics multiprocessor 1634 can output a processed task to a data crossbar switch 1640 to provide the processed task to another processing cluster 1614 for further processing, or store the processed task in an L2 cache, local parallel processor memory, or system memory via a memory crossbar switch 1616. A preROP (pre-raster operation unit) 1642 can be configured to receive data from the graphics multiprocessor 1634 and direct the data to ROP units, which can be associated with partitioning units (e.g., ...) described herein. Figure 16A The PreROP 1642 unit is located together with partition units 1620A-1620N. The PreROP 1642 unit can perform color blending optimization, organize pixel color data, and perform address translation.

[0162] In at least one embodiment, the processing cluster 1614 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described herein above or elsewhere. The one or more circuits may be software-configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described herein above or elsewhere.

[0163] Figure 16C A graphics multiprocessor 1634 according to at least one embodiment is illustrated, for example, to perform any of the operations described above or elsewhere herein. The graphics multiprocessor 1634 may be coupled to a pipeline manager 1632 of a processing cluster 1614. The graphics multiprocessor 1634 may include an execution pipeline including, but not limited to, an instruction cache 1652 (e.g., which may store instructions, such as, but not limited to, compiled API instructions), an instruction unit 1654, an address mapping unit 1656, a register file 1658, one or more general-purpose graphics processing unit (GPGPU) cores 1662, and one or more load / store units 1666, one or more of which may perform load / store operations to load / store instructions corresponding to the execution operations. The GPGPU cores 1662 and load / store units 1666 may be coupled to a cache memory 1672 and a shared memory 1670 via a memory and cache interconnect 1068. The GPGPU cores 1662 may be part of a SoC, such as, but not limited to, [other components]. Figure 15 It is part of the integrated circuit 1500.

[0164] Instruction cache 1652 can receive a stream of instructions to be executed (e.g., perform any of the operations described above or elsewhere herein) from pipeline manager 1632. Instructions can be cached in instruction cache 1652 and dispatched for execution by instruction unit 1654. Instruction unit 1654 can dispatch instructions into thread groups (e.g., thread bundles, subgroups, wavefronts, or waves), with each thread in the thread group assigned to a different execution unit within GPGPU core 1662. Instructions can access any of the local, shared, or global address spaces by specifying an address within a unified address space. Address mapping unit 1656 can be used to translate addresses in the unified address space into different memory addresses accessible by load / store unit 1666.

[0165] Register file 1658 provides a set of registers for the functional units of graphics multiprocessor 1634. Register file 1658 provides temporary storage for operands on data paths connected to functional units of graphics multiprocessor 1634 (e.g., GPGPU core 1662, load / store unit 1666). Register file 1658 can be partitioned among functional units such that each functional unit is allocated a dedicated portion of register file 1658. Register file 1658 can be partitioned among different thread bundles (which may be referred to as wavefronts, subgroups, and / or waves or threads) executed by graphics multiprocessor 1634.

[0166] Each GPGPU core 1662 may include a floating-point unit (FPU) and / or an integer arithmetic logic unit (ALU) for executing instructions of the graphics multiprocessor 1634. The architectures of the GPGPU cores 1662 may be similar or different. A first part of the GPGPU core 1662 may include a single-precision FPU and an integer ALU, while a second part of the GPGPU core may include a double-precision FPU. The FPU may implement IEEE 754-2008 standard floating-point arithmetic or enable variable-precision floating-point arithmetic. The graphics multiprocessor 1634 may also include one or more fixed-function or special-function units for performing specific functions, such as, but not limited to, copying rectangles or pixel blending operations. One or more of the GPGPU cores 1662 may also include fixed-function or special-function logic.

[0167] The GPGPU core 1662 may include SIMD logic capable of executing a single instruction on multiple sets of data. The GPGPU core 1662 can physically execute SIMD4, SIMD8, and SIMD16 instructions, and logically execute SIMD1, SIMD2, and SIMD32 instructions. The SIMD instructions for the GPGPU core can be generated by the shader compiler at compile time, or automatically generated when executing programs written and compiled for Single Program Multiple Data (SPMD) or SIMT architectures. Multiple threads of a program can be configured for a SIMT execution model that can be executed via a single SIMD instruction. For example, eight SIMT threads performing the same or similar operations can be executed in parallel via a single SIMD8 logic unit.

[0168] The memory and cache interconnect 1668 may include an interconnect network that connects each functional unit of the graphics multiprocessor 1634 to the register file 1658 and shared memory 1670. The memory and cache interconnect 1668 may be a cross-switch interconnect that allows the load / store unit 1666 to perform load and store operations between the shared memory 1670 and the register file 1658. The register file 1658 may operate at the same frequency as the GPGPU core 1662, thus data transfer between the GPGPU core 1662 and the register file 1658 can have very low latency. The shared memory 1670 can be used to implement communication between threads executing on functional units within the graphics multiprocessor 1634. The cache memory 1672 can be used as a data cache, for example, for caching texture data transferred between functional units and texture units 1636. The shared memory 1670 can also be used as a program-managed cache. In addition to automatically caching the data stored in cache memory 1672, threads executing on GPGPU core 1662 can also programmatically store data in shared memory.

[0169] The parallel processor or GPGPU described herein can be communicatively coupled to a host / processor core to accelerate graphics operations, machine learning operations, pattern analysis operations, and various general-purpose GPU (GPGPU) functions. The GPU can be communicatively coupled to the host processor / core via a bus or other interconnect (e.g., high-speed interconnects, such as, but not limited to, PCIe or NVLink). A System-on-a-Chip (SoC) may include the parallel processor or GPGPU described herein, which executes on the SoC. The GPU may be integrated as a core on a package or chip and communicatively coupled to the core via an internal processor bus / interconnect within the package or chip. Regardless of the GPU's connection method, the processor core can assign work to the GPU in the form of a sequence of commands / instructions contained in a job descriptor. The GPU can then use dedicated circuitry / logic to efficiently process these commands / instructions to perform any of the operations described above or elsewhere herein.

[0170] In at least one embodiment, the graphics multiprocessor 1634 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0171] Figure 17A processor 1700 according to at least one embodiment is illustrated. The processor 1700 may include a hybrid architecture processor (e.g., Lunar Lake or Meteor Lake) from Intel Corporation, Santa Clara, California, or other processors sharing at least some of the components described herein. The processor 1700 may include one or more central processing units (CPU 1702), one or more graphics processing units (GPU 1706), and / or one or more neural processing units (NPU 1708), which may be, for example, dedicated AI accelerators for offloading artificial intelligence (AI) workloads from the CPU 1702 and GPU 1706. The processor 1700 may use instructions that, if executed, cause the processor 1700 and / or any of its components to perform some or all of the processes and techniques described elsewhere herein. The processor 1700 may include any number of memory and cache units 1710 for facilitating processing between different components of the processor 1700. The memory and cache 1710 on processor 1700 may include one or more levels of cache (e.g., L1, L2, L3, and / or last-level cache) and high-bandwidth memory (e.g., HBM2e or HBM3) in any combination. Regarding processor 1700 and any components described above or elsewhere herein, one or more APIs described herein may, for example, be compiled into instructions that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) internal or external to processor 1700. The results of APIs may be stored in storage devices internal or external to processor 1700, including registers, DRAM, flash memory, SRAM, cache, or other memory. One or more APIs described herein may include calls.

[0172] The processor 1700 may include the computing engine as the CPU 1702, and may include any number of cores, such as, but not limited to, up to 16 cores / 22 threads. The cores in the CPU 1702 may include P-cores (performance), E-cores (high efficiency), and LP-E cores (low-power, high-efficiency). Performance cores can be used for low-latency, single-threaded, compute-intensive workloads, while high-efficiency cores can be used for multi-threaded, less compute-intensive workloads. Low-power, high-efficiency cores can be used for scalable multi-threaded execution and offloading background tasks. P-cores can be used for single-threaded and limited-threaded execution, while E-cores and LP-E cores are used for multi-threaded throughput and power efficiency.

[0173] The GPU 1706 can include any number of graphics engines, such as, but not limited to, those with 8 Xe cores (up to 128 execution units or EUs). Arc TM Graphics engine (Xe LPG). For example... Figure 17 As shown, GPU 1706 may include a vector engine 1710 and a matrix engine 1712, which, for example, can run FP, INT, and matrix operation tasks simultaneously, individually, or in batches. GPU 1706 may include a load / store unit 1714, as well as other memories, such as, but not limited to, an instruction cache (I$) 1716 and an L1 cache / subsystem local memory (SLM) 1718, which may, for example, store instructions for performing any of the operations described above or elsewhere herein.

[0174] The NPU 1704 may include one or more AI Boost integrates a Neural Processing Unit (NPU). The NPU1704 can be enumerated as an integrated PCIe device to the host processor. The NPU 1704 can include one or more (e.g., two) Neural Computation Engine (NCE) tiles 1730. Each tile can be configured with any combination of, but not limited to, the following: (e.g., 2000) Multiply-Accumulate (MAC) engines 1734, a post-processing engine (not shown), an AEP processor (not shown), and memory per tile (2MB dedicated SRAM), such as... Figure 17 As shown. For general computing needs, the neural computing engine 1730 may include an inference pipeline 1732 for high-performance parallel computing, an activation function (AF) 1736, a data transformation 1738, a load / store 1740, and a streaming hybrid architecture vector engine (SHAVE) 1728, which may include a DMA (Direct Memory Access) engine 1724 for transporting data between system memory DRAM (Dynamic Random Access Memory) 1726 and a software-managed cache. The built-in device MMU (Memory Management Unit) 1722, plus the IOMMU (Input-Output Memory Management Unit) (not shown), can support multiple concurrent hardware contexts and provide secure isolation between execution contexts according to the MCDM (Microsoft Computing Driver Model) architecture. The processor 1700 may also include a media unit (not shown), which may be included on or separate from the XCD or other components of the processor 1700 to enable video playback and video processing of compressed or uncompressed data, such as using HEVC, AV1, VP9, ​​and AVC hardware-accelerated decoding support and HEVC, VP9, ​​and AVC hardware-accelerated encoding support.

[0175] Thread bootstrap ( The Thread Director (which includes firmware built into the processor 1700) can prioritize and manage the distribution of workloads, sending tasks to optimized cores. For example, the Thread Director can tie P cores, E cores, and / or LP-E cores (as described above) together with task scheduling capabilities and the ability to send less demanding tasks to E cores or LP-E cores. Deep learning acceleration ( DLBoost (not shown) can provide built-in AI acceleration for training and inference workloads and may include support for VNNI (for CPU) and DP4a (for GPU) instruction sets. This instruction set can be used with OpenVINO. TM The toolkit and oneAPI are optimized to accelerate INT8 inference. For example, the software stack described elsewhere in this document can be used to leverage OpenVINO. TM The toolkit enables AI inference. The processor 1700 can be configured to execute applications, such as, but not limited to, CUDA programs.

[0176] In at least one embodiment, processor 1700 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any operations described herein above or elsewhere. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any operations described herein above or elsewhere.

[0177] Processor 1700 may alternatively include a processor based on Qualcomm's AIEngine Direct architecture from Santa Clara, California, or other processors sharing at least some of the components described herein. It may include any number of NPUs, GPUs, CPUs, and other associated components, such as, but not limited to, an NPU 1704 as a Hexagon NPU, a GPU 1706 as an Adreno GPU, a CPU 1702 as a Kryo or Qualcomm Oryon CPU, and a Qualcomm Sensing Hub (not shown) and a memory subsystem 1710 in any combination. Hexagon NPU 1704 may include power rails, micro-tile inference units, hardware acceleration units, tensor units, scalar units, and vector units (all not shown), which may have dedicated or shared memory (e.g., cache or memory, such as HBM3) for storing, for example, instructions for performing any of the operations described above or elsewhere herein. The Adreno GPU 1706 can provide graphics and parallel processing for AI, in formats including but not limited to 32-bit floating-point (FP32), 16-bit floating-point (FP16), and 8-bit integer (INT8). The Kryo or Qualcomm Oryon CPU 1702 can execute AI workloads and handle the contextualization of ubiquitous generative AI applications. The CPU 1702 may also include an instruction fetch unit, a renaming and dismiss unit, a memory management unit, a vector execution unit, an integer execution unit, and a load and store unit for processing and instruction management. Regarding the processor 1700 and any of its components described above or elsewhere herein, one or more APIs described herein can, for example, be compiled into instructions that can be fetched by the instruction fetch unit, decoded by the processor decoder or equivalent, scheduled (e.g., sequentially or out of order) for execution by the scheduler or equivalent, executed by execution logic or equivalent, reordered, and then dismissed by the renaming and dismiss unit. The API (and / or compiled instructions including the API) can be stored in any storage device (e.g., cache and / or memory) inside or outside the processor 1700. An arbitrary number of CPU cores 1702 can be included in an arbitrary number of CPU clusters, which can be coupled to memory and / or cache, such as, but not limited to, a shared L2 cache. Memory can be separate or shared; for example, the CPU clusters of CPU cores 1702 can be coupled to a memory subsystem 1710, which can include structures capable of reading and writing to memory (e.g., DRAM), system-level caches, and an arbitrary number of memory management units.The Qualcomm sensing hub (not shown) includes a miniature NPU, power rails, and conventional sensors (such as gyroscopes, accelerometers, or even barometers) that support voice and data streaming. The memory subsystem 1710 may include memory and cache on the processor 1700, which may include L1 or more levels of cache (e.g., L1, L2, L3, and / or last-level cache) and high-bandwidth memory (e.g., HBM2e or HBM3) in any combination, for example, for storing information and / or instructions for performing any of the operations described above or elsewhere herein. All or part of the memory and / or cache in the memory subsystem 1710 may be shared or used individually by any component or combination of components on the processor 1700 (e.g., GPU 1706, NPU 1704, and CPU 1702).

[0178] The Qualcomm AI Engine 1700 can be programmed and controlled using a software stack to perform some or all of the operations described herein, including, for example... A neural processing SDK is provided for inference on Android, Linux, and Windows. Developer libraries and services support programming languages, virtual platforms, and compilers. At lower levels of the software stack, system software includes a basic real-time operating system (RTOS), system interfaces, and drivers. The software stack supports various operating systems, including Android, Windows, Linux, and QNX, as well as deployment and monitoring infrastructures such as Prometheus, Kubernetes, and Docker. OpenCL and DirectML are supported for direct cross-platform access to the GPU 1706. For the CPU 1702, LLVM compiler infrastructure optimizations enable accelerated and efficient AI inference. Regarding the Qualcomm AI Engine 1700 and any of its components described above or elsewhere herein, one or more APIs described herein can, for example, be compiled into instructions that can be fetched by instruction fetching logic or equivalents, decoded by processor decoders or equivalents, scheduled (e.g., sequentially or out of order) for execution by a scheduler or equivalent, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. The API (and / or compiled instructions including the API) can be stored in any storage device (e.g., cache and / or memory) inside or outside the Qualcomm AI Engine 1700. The results of the API can be stored in storage devices inside or outside the Qualcomm AI Engine 1700, including registers, DRAM, flash memory, SRAM, cache, or other memory.

[0179] In at least one embodiment, the processor 1700 or the Qualcomm AI engine 1700 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0180] Figure 18A A processor 1800 according to at least one embodiment is illustrated. The processor 1800 may include a Scalable family of processors from Intel Corporation, Santa Clara, California, or other processors that share at least some of the components described herein. The processor 1800 may include one or more cores 1812(1)-1812(N) capable of performing the operations described elsewhere herein, where N is any integer greater than 1. Cores 1812(1)-1812(N) may be interconnected using ring and / or mesh interconnects. Utilizing a mesh interconnect architecture, arrays of vertical and horizontal communication paths may allow traversal from one core to another 1812(1)-1812(N) via the shortest path (jumping to the correct row along the vertical path and to the correct column along the horizontal path). For the mesh interconnect, a die may accommodate cores 1812(1)-1812(N) and may include a Converged Grid Stop Point (CMS) grid that may be associated (e.g., 1:1) with cores 1812(1)-1812(N). Each core can be associated with a low-level cache (LLC) slice 1814(1)-1814(N), or cores 1812(1)-1812(N) can share a cache, such as a low-level cache. LLC 1814(1)-1814(N) can be inclusive or non-inclusive (having blocks not present in the high-level cache) by merging blocks from a higher-level cache (e.g., L2 cache). Each core and LLC slice can include a caching and home agent (CHA) (not shown), which can maintain cache consistency by providing resource scalability via the mesh interconnect. Super Path Interconnect ( UPI 1816 provides cache coherency capabilities. UPI 1816 enables coherent interconnects for scalable systems and allows multiple processors to share a single shared address space via links, such as, but not limited to, two or three UPI links per processor.

[0181] The processor 1800 may also include a system agent 1810, which may house and / or perform various functions, such as, but not limited to, memory management, display functions, and / or input / output (I / O) functions. For example, the processor 1800 may include one or more integrated memory controllers (IMCs) 1808. The IMC 1808 may control and manage memory, such as, but not limited to, different memory types, such as DDR RAM, such as DDR4, or other memory described elsewhere herein. The system agent 1810 may include a display controller (not shown) for supporting one or more displays. The system agent 1810 may also integrate a PCIe 1804 (e.g., up to 20 PCIe lanes), which may, for example, be connected to an external dedicated graphics connector via a DMI bus (e.g., Intel's DMI 3.0 bus) 1206. The system agent 1810 may include an image processing unit (IPU) (not shown) that integrates an on-die image signal processor (ISP). Structure 1802 provides scalability for connecting to other nodes (e.g., processors, such as processor 1800) and can, for example, be connected to Cornelis Networks (…). It can be used together with elements of a scalable system framework that provides performance for high-performance computing (HPC) workloads and the ability to scale to tens of thousands of nodes.

[0182] Figure 18BComponents within a core 1812 according to at least one embodiment are illustrated. The core 1812 may include a front-end 1818, a back-end or execution engine 1832, and a memory subsystem 1842. The front-end 1818 can provide operations (e.g., operations described elsewhere herein) to the execution engine 1832 by decoding instructions stored in memory. For example, the front-end 1818 may include micro-operation (μOps) cache paths and / or conventional paths, and a branch prediction unit 1820 that can determine path instructions. A conventional path of instructions might involve fetching variable-length (e.g., x86) instructions from an L1 instruction cache using instruction fetching and pre-decoding 1822, queuing these instructions into an instruction queue 1824, and decoding the instructions into μOps that can be provided to an allocation queue 1828 using a decoder 1826. Alternatively, the μOps cache path may include a cache that includes decoded μOps (μOps 1830) that can be sent to the allocation queue 1828. The allocation queue 1828 can act as an interface between the front end 1818 and the execution engine 1832, and can provide instructions to the execution engine 1832. For example, one or more APIs described herein can be compiled into instructions that can be stored, processed, and executed by the front end 1818 and the execution engine 1832, and stored in the memory subsystem 1842.

[0183] Execution engine 1832 can receive micro-operations into reordering buffer 1834, which can register, rename, and terminate μOPs. μOPs can be sent from the reordering buffer to scheduler 1836, which can be connected to one or more different execution units 1838, which can be connected to address generation units (AGUs) 1840. Execution units 1838 can perform operations such as basic arithmetic logic unit (ALU) operations, multiplication, division, and / or more complex operations, such as, but not limited to, various vector operations. Scheduler 1836 can manage the queuing of μOPs for one or more execution units 1838 based on, for example, the operations that need to be performed.

[0184] The memory subsystem 1842 can handle load and store requests as well as sorting operations. For example, μOPs may be associated with memory accesses (e.g., load and store), and these μOPs can be sent through dedicated scheduler ports that can perform these memory operations. For example, store and load operations can be sent to load and store buffers 1844. The memory subsystem 1842 may also include shared or separate L1 data and instruction caches 1846, and an L2 cache 1848 that can be used and shared by the L1 data and instruction caches 1846. (As described above regarding...) Figure 18AEach core 1812 can be connected to a slice of a third-level cache (e.g., LLC 1814), which can be shared by all cores 1812.

[0185] In at least one embodiment, processor 1800 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0186] Figure 19 An AI accelerator 1900 according to at least one embodiment is illustrated. The processor 1900 may include a processor with an AI accelerator architecture manufactured by Intel Corporation, Santa Clara, California, or other processors sharing at least some of the components described herein. The AI ​​accelerator 1900 may use instructions that, if executed by the AI ​​accelerator 1900, cause the AI ​​accelerator 1900 to perform some or all of the processes and techniques described elsewhere herein. For example, with respect to the AI ​​accelerator 1900 and any components described above or elsewhere herein, one or more APIs described herein may, for example, be compiled into instructions that may be fetched by instruction fetching logic or equivalents, decoded by processor decoders or equivalents, scheduled (e.g., sequentially or out of order) by a scheduler or equivalent for execution, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. The APIs (and / or compiled instructions including the APIs) may be stored in any storage device internal or external to the AI ​​accelerator 1900 (e.g., in cache and / or memory). The results of the API can be stored in internal or external storage devices of the AI ​​accelerator 1900, including registers, DRAM, flash memory, SRAM, cache, or other memory. The AI ​​accelerator 1900 may include one or more compute dies, which may include homogeneous or heterogeneous processors. The compute dies may include one or more central processing units (CPUs), one or more graphics processing units (GPUs), or a combination of both.

[0187] In at least one embodiment, the computational die may include a computational engine for performing AI computations. In at least one embodiment, the computational die of the AI ​​accelerator 1900 may be split into any number (e.g., four) clusters, which may be referred to as DCORE (Deep Learning Core) 1906, and include any number of matrix multiplication engines (MME) 1908, tensor processor cores (TPC) 1910, memory management units 1912, and L2 caches 1914 in any combination. The MME 1908 may perform operations using matrix multiplication, such as fully connected layers, convolutions, and batch general matrix multiplication (GEMM). The MME 1908 may be equipped with a multiplication-accumulation unit (MAC) (not shown), which may perform general matrix multiplication (GEMM) operations, such as, but not limited to, AxB multiplication, which involves generating a tensor C [NxM] from two input tensors A [NxK] and B [KxN]. The MME 1908 may be programmed with array dimensions, positions, data types, and various operands. The MME 1908 can retrieve tensors A and B from memory and pull them into its streaming buffer for parallel matrix multiplication by the MAC. After completion, the MME 1908 can push tensor C back to memory. The TPC 1910 may include any number of scalar units for performing scalar operations, any number of vector units for performing vector operations, any number of register files or local memory units (e.g., vector local memory), and load and store components for instructions, which may be coupled to memory or caches (e.g., HBM, L3 cache, and / or L2 cache) (all not shown). The TPC can support different types of parallel processing, such as Very Long Instruction Word (VLIW) Single Instruction Multiple Data (SIMD) data types such as, but not limited to, FP32, BF16, FP16, and FP8 (both E4M3 and E5M2), UINT32, INT32, UINT16, INT16, UINT8, and INT8 data types. Any number of computational dies can be interconnected. Interconnects that can connect computing dies can be via intermediate bridges, which are transparent to software, for example.

[0188] The memory on the AI ​​accelerator 1900 may include one or more levels of cache (e.g., L1, L2, L3, and / or last-level cache) and high-bandwidth memory (e.g., HBM2e or HBM3) in any combination. The memory and / or cache system may be unified or separate. The compute die of the AI ​​accelerator 1900 may include on-chip memory comprising one or more levels (e.g., two levels) of cache. On-chip SRAM or other memory described elsewhere herein may be used as a unified last-level cache (L3) or split into multiple slices of L2 cache accessible to the MME 1908 and TPC 1910 groups. Using on-chip memory as an L2 or L3 cache is entirely software-configurable, with the software dynamically determining its optimal cache allocation based on I / O tensors. The AI ​​accelerator 1900 may include one or more memory management units (MMUs) 1922 for managing memory, such as allowing the AI ​​accelerator 1900 memory subsystem to run in virtual space when accessing VRAM.

[0189] AI accelerator 1900 may include a communication port (e.g., a PCIe Gen5 x16 port) 1902 for communicating with the host and scheduling and synchronization unit 1904. AI accelerator 1900 may include a media unit 1916, which may include any number or combination of media decoder engines (DECs) 1920 and rotation engines (ROTs) 1918. AI accelerator 1900 may include a network unit 1924, which may include any number or combination of network ports 1926 and an accompanying RDMA engine 1928, L2 cache, and memory (e.g., HBM2e or HBM3) stack. AI accelerator 1900 may include a programmable control path entity (not shown) for managing the parallel and efficient execution of the various engines. The control path may include a submission queue (SQ) that can be issued by the runtime system, a completion queue (CQ) that can be used for job completion reporting, a programmable scheduling mechanism that can be used for task scheduling, a programmable hardware synchronization mechanism or "synchronization manager (SM)" that can be used for hardware synchronization, and a programmable interrupt service mechanism or "interrupt manager (INTR)" that can pass asynchronous events to drivers.

[0190] The AI ​​Accelerator 1900 may include media decoding units supporting video formats such as, but not limited to, HEVC, Progressive H.264, SVC base layer, MVC, VP9, ​​JPEG, and Progressive JPEG. The AI ​​Accelerator 1900 may support post-processing of the decoded media stream, such as, but not limited to, image downsizing (image resizing), vertical and horizontal scaling at different scaling ratios, image enlargement, image cropping, bilinear scaling, and Lancos scaling. The AI ​​Accelerator 1900 may implement two post-processing channels per decoder unit, one for scalar (up and down) and the other solely for outputting the original image. The AI ​​Accelerator 1900 may include a hardware rotation engine that performs the following transformations on the input image: 2D rotation, 3D rotation, projection, image warping and de-warping, resampling of the input data at user-defined coordinates, and rescaling.

[0191] The RDMA 1928 based on converged Ethernet on the AI ​​accelerator 1900 enables scaling from a single node (i.e., from a single AI accelerator 1900 to hundreds or thousands of nodes or AI accelerators 1900). The network subsystem 1924 may include... The accelerator 1900 includes an In-Gigabit Ethernet Communication Library (IGCL), a master controller coordinating data movement, and a programmable scheduling mechanism that enables smooth engine activation while maintaining task dependencies. The accelerator network subsystem may include a Gigabit Ethernet NIC port 1926, a Layer 2 MAC (not shown), and an RDMA engine 1928. The AI ​​accelerator 1900 may include an aggregation engine for performing summation activities. All engines in the processor 1900 can operate in parallel; for example, the MME 1908, TPC 1910, and NIC 1926 can all operate simultaneously. Dependencies may exist between operations running on different engines; for example, the output of one engine may be used as the input of another engine, and / or the MME, TPC, and NIC may be scheduled to run in parallel. When one engine completes its execution, another engine can be scheduled to begin working on the next operation (executed immediately after its input is ready).

[0192] The AI ​​accelerator 1900 can be operated and controlled using software layer 1928, which may include low-level components such as, but not limited to, a graph compiler, an automatic kernel fusionist and pre-compiled kernel libraries, and integrations with the AI ​​ecosystem such as, but not limited to, PyTorch, DeepSpeed, Hugging Face, vLLM, Ray, etc., or as described elsewhere in this document regarding software and programming platforms. Software layer 1928 may include implementations of algorithms such as, but not limited to, paged attention, flash attention, etc. Software layer 1928 can generate optimized binary code that implements a given model topology, such as, but not limited to, performing operator fusion, data layout management, parallelization, pipeline and memory management, and graph-level optimization.

[0193] In at least one embodiment, the AI ​​accelerator 1900 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0194] This paper describes a neuromorphic computing system employing a multi-core architecture, where each core houses computing elements including neurons, synapses with on-chip learning capabilities, and local memory for storing synaptic weights and routing tables. Figure 20This is a simplified block diagram 2000 illustrating at least a portion of an example of such a neuromorphic computing device 2005 according to at least one embodiment. The neuromorphic computing device 2005 may include a neuromorphic processor from Intel Corporation, Santa Clara, California, or other processors that include at least a portion of the components described herein. As shown in this example, the device 2005 may be equipped with a network 2010 consisting of multiple neural network cores interconnected by a network on the device, thereby potentially defining multiple distinct connections between the cores. For example, the device 2005 may provide a network 2010 of spiking neural network cores, each core communicating via short packet spiking messages sent from one core to another through network channels. Each core (e.g., 2015) may have processing and memory resources, as well as logic, for implementing a number of primitive nonlinear time computation elements, such as, but not limited to, multiple (e.g., more than 1000) distinct artificial neurons (referred to herein as “neurons”). For example, each core may be able to implement multiple neurons concurrently, such that the neuromorphic core can implement many, many neurons using the device 2005. With respect to the neuromorphic computing device 2005 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) internal or external to the neuromorphic computing device 2005. The results of the APIs may be stored in storage devices internal or external to the neuromorphic computing device 2005, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0195] continue Figure 20For example, the neuromorphic computing device 2005 may also include a processor 1420 and system memory 2025 for implementing one or more components to manage and provide the functionality of the neuromorphic computing device 2005. For instance, a system manager 2030 may be provided to manage the global attributes and operations of the neuromorphic computing device 2005 (e.g., attributes affecting the core network 2010, multiple cores in the network 2010, interconnections of the neuromorphic computing device 2005 with other devices, managing access to the global system memory 2025, and other potential examples). In one example, the system manager 2030 may manage the definition and configuration of specific routing tables for individual routers in the network 2010, the orchestration of network definitions and attributes to be applied to the network (e.g., weights, attenuation rates, etc.), core synchronization and time multiplexing management, routing input to the appropriate core, and other potential functions.

[0196] As another example, the neuromorphic computing device 2005 may also include a programming interface 2035 through which a user or system can specify the neural network definition to be applied (e.g., via routing tables and individual neuron attributes), implemented by the neuromorphic core grid 2010. A software-based programming tool may be provided with or separately from the neuromorphic computing device 2005, through which a user can provide a definition for a specific neural network to be implemented using the neuromorphic core network 2010. The programming interface 2035 can receive input from the programmer, then generate the corresponding routing table and populate the specified parameters into the local memory of each neuromorphic core (e.g., 2015) to implement the corresponding custom artificial neural network implemented by the neuromorphic core 2015.

[0197] In certain circumstances, the neuromorphic computing device 2005 can advantageously engage and interoperate with other devices, including general-purpose computing devices, to enable specific applications and use cases. Therefore, in some cases, external interface logic 2040 may be provided to communicate with one or more other devices (e.g., via one or more defined communication protocols). External interface 2040 may be used to accept input data from another device or an external memory controller used as an input data source. External interface 2040 may additionally or alternatively be used to allow the results or outputs of computations performed using the neural network implemented using the neuromorphic computing device 2005 to be provided to another device (e.g., another general-purpose processor implementing machine learning algorithms) to enable additional applications and enhancements, among other examples.

[0198] like Figure 20The diagram illustrates a network 2010 interconnected by a network on a device, showing a portion of a network structure interconnecting multiple neuromorphic cores (e.g., 2015a-2015d). For example, several neuromorphic cores (e.g., 2015a-2015d) can be provided in a mesh, each core interconnected via a network including multiple routers (e.g., 2050). In one implementation, each neuromorphic core (e.g., 2015a-2015d) can be connected to a single router (e.g., 2050) in the routers, and the router can be connected to at least one other router (e.g., [other router]). Figure 20 (As shown at 2010 in [reference]). As an example, in one particular implementation, four neuromorphic cores (e.g., 2015a-2015d) can be connected to a single router (e.g., 2050), and each router 2050 can be connected to two or more other routers to form a multi-core mesh, thereby allowing each neuromorphic core to interconnect with every other neuromorphic core in the neuromorphic computing device 2005. Furthermore, since each neuromorphic core can be configured to implement multiple different neurons, the router network of the neuromorphic computing device 2005 can similarly implement connections or artificial synapses (or simply "synapses") defined between any two of the potentially many (e.g., 30,000+) neurons defined using the neuromorphic core 2010 network provided in the neuromorphic computing device 2005.

[0199] Figure 20A block diagram of the internal components of an example implementation of a neuromorphic core 2015 is shown. In one example, a single neuromorphic core may implement a number of neurons (e.g., 1024), which share the architectural resources of the neuromorphic core 2015 in a time-division multiplexing manner. In one example, each neuromorphic core 2015 may include a processor block 2055 capable of executing arithmetic functions and routing related to the implementation of the digitally implemented artificial neurons, such as, but not limited to, those explained herein. Each neuromorphic core 2015 may also provide local memory in which routing tables of the neural network can be stored and accessed, accumulated potentials of each cell body of each neuron implemented using core 2015 can be tracked, parameters of each neuron implemented by core 2015 can be recorded, and other data and usage can be recorded. Components or architectural resources of the neuromorphic core 2015 may also include: an input interface 2065 for receiving input spike messages generated by other neurons on other neuromorphic cores; and an output interface 2070 for sending spike messages to other neuromorphic cores via a mesh network 2010. In some instances, the routing logic of the neuromorphic core 2015 can be implemented at least partially using the output interface 2070. Furthermore, in some cases, the core (e.g., 2015) can implement multiple neurons within an example SNN, and some of these neurons can be interconnected. In this case, spiking messages sent between neurons hosted on the core 2015 can forgo communication via the routing structure of the neuromorphic computing device 2005 and can be managed locally within the specific neuromorphic core 2015.

[0200] Each neuromorphic core may also include logic for implementing artificial dendrites 2080 and artificial cell bodies 2085 (hereinafter referred to as “dendrites” and “cell bodies”, respectively) for each neuron 2075. Dendrites 2080 may be hardware-implemented processes for receiving impulses from the network. Cell bodies 2085 may be hardware-implemented processes for receiving the current time-accumulated neurotransmitter mass of each dendrite and evolving the potential states of each dendrite and cell body to generate outgoing impulse messages at appropriate times. Dendrites 2080 may be defined for each connection receiving input from another source (e.g., another neuron). In one implementation, the dendritic process 2080 may receive and process the impulse message as it arrives serially from network 2010 in a time-division multiplexed manner. With the reception of impulses, neuronal activation (tracked using cell bodies 2085 (and local memory 2060)) may increase. When the activation of a neuron 2075 exceeds a threshold set for that neuron, neuron 2075 generates a spike message, which is propagated via output interface 2070 to a fixed set of fan-out neurons. The network distributes the spike messages to all destination neurons, which in turn can update their activation in a transient, time-dependent manner in response. This can lead to some of the destination neurons also exceeding their corresponding thresholds and triggering further spike messages, just as in real biological neural networks.

[0201] As described above, the neuromorphic computing device 2005 can reliably implement spiking-based neural computing models. Such models are also known as spiking neural networks (SNNs). In addition to neuronal and synaptic states, SNNs incorporate temporal concepts. For example, in SNNs, communication occurs via event-driven action potentials or spiking, which convey no explicit information other than the spiking time and the implicit source and destination neuron pairs corresponding to the spiking transmission. The computation of the result of a dynamic nonlinear integral as a weighted spiking input occurs in each neuron. In some implementations, recurrent and dynamic feedback can be incorporated into the SNN computation model. Furthermore, various network connectivity models can be employed to model a wide range of real-world networks or relationships, including fully connected (all-to-all) networks, feedforward trees, completely random projections, “small-world” networks, and other examples. Isomorphic two-dimensional networks at the neuromorphic core (e.g., but not limited to...) Figure 20 The network shown in the example can advantageously support all these network models. Since all cores of the neuromorphic computing device 2005 are connectable, some or all neurons defined in a core can also be fully connected via a certain number of router hops. The neuromorphic computing device 2005 can also include fully configurable routing tables for defining various neural networks by allowing neurons in each core to distribute their spiking to any number of cores in the grid 2010 to achieve a completely arbitrary connection graph.

[0202] In improved implementations of systems capable of supporting SNNs, for example, but not limited to... Figure 20 The example illustrates a very large-scale integrated circuit (VLSI) hardware device that can provide high-speed, reliable circuitry to implement SNNs (Spiritual Neural Networks) to model the information processing algorithms used by the brain, but in a more programmable way. For instance, while a biological brain can only perform a specific set of defined behaviors (a consequence of years of development), a neuromorphic processor device can provide the ability to rapidly reprogram all neural parameters. Therefore, a single neuromorphic processor can be used to implement a wider range of behaviors than a single slice of biological brain tissue. This distinction can be achieved by employing neuromorphic processors with neuromorphic designs that are radically different from those found in natural neural circuitry.

[0203] As an example, a neuromorphic processor can implement a spontaneous neural network (SNN) using time-multiplexed computation in both a spiking communication network and the neuronal mechanism of a neuromorphic computing device 2005. Therefore, the physical circuitry of the neuromorphic computing device 2005 can be shared by many neurons to achieve a higher neuron density. Through time multiplexing, the network can connect N cores with a total wiring length of O(N), while the length of discrete point-to-point wiring will be extended to O(N). 2 This significantly reduces wiring resources to accommodate planar and non-plastic VLSI routing techniques, among other examples. In the neuromorphic core, time multiplexing can be implemented through dense memory allocation, for example, using static random access memory (SRAM) with a shared bus, address decoding logic, and other multiplexed logic elements. The state of each neuron can be stored in the processor's memory, where data describing the state of each neuron includes the state of the collective synapse of each neuron, all currents and voltages on its membrane, and other example information (e.g., but not limited to configuration and other information).

[0204] Neuromorphic processors can be implemented in a “digital” manner, unlike other processors that employ more “analog” or “isomorphic” neuromorphic approaches. For example, a digital implementation can use digital adder and multiplier circuitry to integrate synaptic currents, in contrast to an analog isomorphic neuromorphic approach that accumulates charge on capacitors in a manner similar to how neurons accumulate synaptic charge on their lipid membranes. For instance, the accumulated synaptic charge for each neuron can be stored in the local memory of the corresponding core. Furthermore, at the architectural level of an example digital neuromorphic processor, reliable and deterministic operation can be achieved through time synchronization across the core network, ensuring that any two executions of the design, given the same initial conditions and configuration, will produce the same results. Asynchronicity can be reserved at the circuit level to allow individual cores to operate as quickly and freely as possible while maintaining determinism at the system level. Therefore, in neural computing, the concept of time as a time variable can be abstracted away from the “wall clock” time used by the hardware to perform computations. Thus, in some implementations, a time synchronization mechanism can be provided that globally synchronizes the neuromorphic cores at discrete time intervals. The synchronization mechanism allows neural computation to be completed at the fastest speed allowed by the circuit, and there is a difference between the runtime and the biological time for modeling neuromorphic systems.

[0205] In operation, the neuromorphic computing device 2005 can start in an idle state when all neuromorphic cores are inactive. As each core asynchronously loops through its neurons, it generates impulse messages, which are routed by the mesh interconnect to the appropriate destination core containing all destination neurons. The implementation of multiple neurons on a single neuromorphic core can be time-multiplexed, and time steps can be defined, where all impulses involving multiple neurons can be processed and considered using the shared resources of the respective cores. When each core completes its service to its neurons within the corresponding time step, in some implementations, the core can communicate with neighboring cores using synchronization messages (e.g., using a handshake) to refresh the mesh of all transmitted impulse messages, allowing the core to safely determine that all impulses have been serviced within a given time step. At this point, all cores can be considered synchronized, allowing them to advance their time steps and return to their initial state to begin the next time step.

[0206] Given this context, as described above, an apparatus (e.g., 2005) can be provided to implement an interconnected neuromorphic core grid 2010, wherein core 2015 can implement multiple artificial neurons capable of interconnecting to realize an SNN. Each neuromorphic core (e.g., 2015) can provide two loosely coupled asynchronous processes: an input dendrite process (e.g., 2080) that receives impulses from network 2010 and applies them to the appropriate destination dendritic chamber at an appropriate future time; and an output cell body process (e.g., 2085) that receives the current-time accumulated neurotransmitter mass of each dendritic chamber and evolves the membrane potential state of each dendrite and cell body to generate an outgoing impulse message at an appropriate time (e.g., when the threshold potential of the cell body is reached). It should be noted that, from a biological perspective, the names of dendrites and cell bodies used herein are only approximate to the function of these features and should not be interpreted too literally.

[0207] In at least one embodiment, the neuromorphic computing device 2005 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described herein above or elsewhere. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described herein above or elsewhere.

[0208] Figure 21 This is a block diagram of an embodiment of a multi-node network capable of enabling remote memory computing according to any embodiment. System 2100 may represent the node network described herein, for example, the node network can be used to perform some or all of the operations described herein. System 2100 may represent a data center. System 2100 may represent a server farm. System 2100 may represent a data cloud or processing cloud. System 2100 may represent a supercomputer. System 2100 may include tens, hundreds, or thousands of nodes. The nodes of System 2100 may include processors, such as, but not limited to, a central processing unit (CPU), a graphics processing unit (GPU), or any combination of processors described herein, such as, but not limited to, processors described herein. Figure 15-27BOther processors in the system. For any processor in system 2100 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) inside or outside the processor or node. The results of APIs may be stored in storage devices inside or outside the processor or node, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents. System 2100 may include more than nine thousand nodes, each node comprising two Intel Xeon Max processors, six Intel Max series GPUs, and a unified memory architecture, such as, but not limited to, the architecture used in Intel's Aurora supercomputer in Santa Clara, California, or other supercomputers that share at least some of the components described herein.

[0209] One or more clients 2102 send requests to system 2100 via network 2104. Network 2104 represents one or more local area networks, wide area networks, or a combination of both. Clients 2102 can be human or machine clients that generate requests for operations to be performed by system 2100. System 2100 executes the application or data computation task requested by client 2102.

[0210] System 2100 may include one or more racks, which represent structural and interconnect resources for housing and interconnecting multiple computing nodes. Rack 2110 may include multiple nodes 2130. Rack 2110 may carry multiple blade assemblies 2120(0) to 2120(N-1), where N is an integer greater than or equal to 2. Carrying may refer to providing power, structural or mechanical support, and interconnection. Blades 2120(0) to 2120(N-1) may refer to computing resources on a printed circuit board (PCB), where the PCB houses the hardware components of one or more nodes 2130. Blades 2120(0) to 2120(N-1) may or may not include a chassis, enclosure, or other “box” besides those provided by rack 2110. Blades 2120(0) to 2120(N-1) may include an enclosure with exposed connectors for connection to rack 2110. System 2100 may or may not include rack 2110, and each blade (e.g., 2120(0)) may include a chassis or housing that may be stacked or otherwise closely proximate with other blades and allow nodes 2130 to interconnect. System 2100 may include 10,624 compute blades, comprising 63,744 Intel Max series GPUs and 21,248 Intel Xeon Max CPUs on 166 racks.

[0211] System 2100 may include architecture 2170, which represents one or more interconnectors of nodes 2130. Architecture 2170 may include multiple switches 2172 or routers or other hardware for routing signals between nodes 2130. Furthermore, architecture 2170 may couple system 2100 to network 2104 for access by client 2102. In addition to routing devices, architecture 2170 may also be considered to include cables or ports or other hardware devices for coupling nodes 2130 together. Architecture 2170 may have one or more associated protocols for managing signal routing through system 2100. One or more protocols are at least partially dependent on the hardware devices used in system 2100.

[0212] As shown in the figure, rack 2110 may include N blades (e.g., 2120(0) to 2120(N-1)). In addition to rack 2110, system 2100 may also include rack 1550. As shown in the figure, rack 1550 may include M blades (e.g., 2160(0) to 2160(M-1)). M is not necessarily the same as N; therefore, it is understood that various different hardware device components may be used and coupled together into system 2100 via structure 2170. Blades 2160(0) to 2160(M-1) may be the same as or similar to blades 2120(0) to 2120(N-1). Node 2130 may be any type of node described herein and is not necessarily of the same type. System 2100 is not limited to homogeneous or non-homogeneous.

[0213] The nodes in blade 2120(0) are shown in detail. However, other nodes in system 2100 may be the same or similar. At least some nodes 2130 may be compute nodes, having processor 2132 and memory 2140. A compute node is a node having processing resources (e.g., one or more processors) that executes an operating system and can receive and process one or more tasks. At least some nodes 2130 may include storage server nodes, which have servers as processing resources 2132 and memory 2140. A storage server is a node having more storage resources than a compute node, and instead of having processors for performing tasks, a storage server includes processing resources for managing access to storage nodes within the storage server.

[0214] Node 2130 may include interface controller 2134, which may represent logic for controlling node 2130's access to structure 2170. The logic may include hardware resources for interconnecting to physical interconnect hardware. The logic may include software or firmware logic for managing the interconnect. Interface controller 1534 may include a host structure interface, which may include a structure interface according to any embodiment described herein.

[0215] Node 2130 may include a memory subsystem 2140. Memory 2140 may include a memory computation resource (comp) 2142, which represents the ability of memory 2140 to perform one or more memory computations. System 2100 supports remote memory operations, such as, but not limited to, those described elsewhere herein. Thus, node 2130 may request a remote node to perform a memory computation, wherein the data used for the computation remains local to the executing node and is not sent via structure 2170 or from memory to the structure interface. In response to the execution of the memory computation, the executing node may provide the result to the requesting node.

[0216] Processor 2132 may include one or more individual processors. Each individual processor may include a single processing unit, a multi-core processing unit, or a combination thereof. A processing unit may include a main processor, such as, but not limited to, a CPU (Central Processing Unit), a peripheral processor (such as, but not limited to, a GPU (Graphics Processing Unit)), or a combination thereof. Memory 2140 may be or include a memory device and a memory controller.

[0217] The term "memory device" can refer to different types of memory. Memory devices generally refer to volatile memory technology. Volatile memory is memory whose state (and the data stored within it) is uncertain if power is interrupted. Non-volatile memory is memory whose state is deterministic even if power is interrupted. Dynamically volatile memory can refresh the data stored in the device to maintain its state. An example of dynamically volatile memory includes DRAM (Dynamic Random Access Memory) or variations thereof, such as, but not limited to, Synchronous DRAM (SDRAM). The memory subsystem described in this article is compatible with a variety of memory technologies, such as, but not limited to, DDR3 (Double Data Rate version 3, originally released by JEDEC (Joint Electron Device Engineering Committee) on June 27, 2007, currently version 21), DDR4 (DDR version 4, initial specification released by JEDEC in September 2012), DDR4E (DDR version 4, extended version, currently under discussion by JEDEC), LPDDR3 (Low Power DDR version 3, JESD209-3B, released by JEDEC in August 2013), and LPDDR4 (Low Power Double Data Rate (LPDDR) version 4). JESD209-4 (originally released by JEDEC in August 2014), WIO2 (Wide I / O2), JESD229-2 (originally released by JEDEC in August 2014), HBM (High Bandwidth DRAM), JESD235 (originally released by JEDEC in October 2013), DDR5 (DDR version 5, currently under discussion by JEDEC), LPDDR5 (currently under discussion by JEDEC), HBM2 (HBM version 2, currently under discussion by JEDEC) or combinations of other memory technologies, as well as technologies derived from or extended based on such specifications.

[0218] In addition to or as an alternative to volatile memory, in one embodiment, a reference to a memory device may refer to a non-volatile memory device whose state is deterministic even when power is interrupted. In one embodiment, a non-volatile memory device is a block-addressable memory device, such as, but not limited to, NAND or NOR technology. Therefore, the memory device may also include future-generation non-volatile devices, such as, but not limited to, three-dimensional cross-point (3DXP) memory devices, other byte-addressable non-volatile memory devices, or memory devices using chalcogenide phase change materials (e.g., chalcogenide glasses). In one embodiment, the memory device may be or include multi-threshold NAND flash memory, NOR flash memory, single-level or multi-level phase change memory (PCM) or switched phase change memory (PCMS), resistive memory, nanowire memory, ferroelectric transistor random access memory (FeTRAM), magnetoresistive random access memory (MRAM) incorporating memristor technology, or spin-transfer torque (STT)-MRAM, or any combination of the foregoing, or other memories.

[0219] In at least one embodiment, system 2100 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied, or otherwise perform any operation described herein above or elsewhere, when one or more branch instructions use one or more instruction operands stored in one or more stack structures. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied, or otherwise perform any operation described herein above or elsewhere, when one or more branch instructions use one or more instruction operands stored in one or more stack structures.

[0220] Figure 22An accelerated processing unit 2200 according to at least one embodiment is illustrated. The accelerated processing unit 2200 may include a processor based on the CDNA architecture of AMD Inc., Santa Clara, California, or other processors sharing at least some of the components described herein. The accelerated processing unit 2200 may include one or more accelerator complex dies (XCDs) 2204 for performing the operations described elsewhere in this document, such as, but not limited to, graphics processing and / or parallel processing and instruction-level parallel computing, including support for multiple precisions (INT8, FP8, BF16, FP16, TF32, FP32, and FP64) and sparse matrix data (i.e., sparsity). In some cases, the XCD may be referred to as a graphics computing die (GCD). The accelerated processing unit 2200 may include one or more complex computing dies (CCDs) 2206 for performing the operations described elsewhere in this document, such as, but not limited to, operations performed by a host processor. In some cases, the CCD may be referred to as a core complex or CCX, such as, but not limited to, the CCX used in AMD Ryzen processors. XCDs and CCDs can share any type of cache or memory (e.g., one or more memory cells 2202), or a cache or memory can be allocated to each XCD or CCD or group of XCDs or CCDs. For example, AMD Infinity Fabric within the package connects XCDs and CCDs to a shared AMD Infinity Cache 2208, and in some embodiments, to high-bandwidth memory (e.g., HMB3). Accelerated processing unit 2200 may include an AMD MI300a processor comprising three CPU die (or CCD) and six accelerator die (XCD) on top of four input-output dies (IODs) that may be layered on a single silicon die (e.g., via AMD Infinity Fabric) and linked together to eight high-bandwidth DRAM stacks in a ring to form a superchip. For systems using only accelerators, the AMD MI300x processor replaces the CCD with two or more XCDs.

[0221] Accelerated processing unit 2200 may include one or more input / output (I / O) interfaces. For example, XCD 2204 and CCD 2206 may reside together on one or more input-output dies (IODs) 2210, which may include one or more I / O interfaces. IOD 2210 may include any number and type of I / O interfaces (e.g., PCI, PCI expansion (“PCI-X”), PCIe, Gigabit Ethernet (“GBE”), USB, etc.). Various types of peripheral devices may be coupled to I / O interface 2270. The I / O interface of IOD 2210 may also be used to connect one or more accelerated processing units 2200, for example, in a server architecture.

[0222] Accelerated processing unit 2200 may include one or more memory units 2202 for storing instructions and other information for performing the operations described in other parts of this document. Memory units 2202 may include any volatile memory, such as, but not limited to, the memory types described in other parts of this document, and may include, for example, high-bandwidth memory (e.g., HMB3) or high-bandwidth DRAM. The memory associated with accelerated processing unit 2200 (e.g., memory unit 2202) may include system memory, which can be used for, for example, commands, instructions, and constants, as well as input and output. Memory unit 2202 may also include device memory, which can be used for storage and, for example, for commands, instructions, and constants, as well as input and output, as a return buffer, and for private data. Memory unit 2202 may be linked to one or more IODs 2210. In at least one embodiment, L1 cache 2220 initiates a memory hierarchy including a shared L2 cache 2228 (e.g., within an XCD). AMD Infinity Cache TM It is the last-level cache (LLC) located on the active I / O die (IOD). The CCD 2206 and XCD2204 can have independent or shared memory. AMD Infinity architecture and AMD Infinity Fabric TM The technology enables consistent, high-throughput unification of GPU and CPU chip technologies (such as XCD, CCD, and / or CCX) with memory (such as stacked HBM3 memory) in a single device and across multiple device platforms.

[0223] like Figure 22As shown, XCD 2204 may include a set of shared global resources 2230, which may include a hardware scheduler 1632 and an asynchronous compute engine (ACE) 2224. The ACE 2224 sends tasks (e.g., compute shader workgroups) to compute units (CUs or cores) 2234. ACEs 2224 (e.g., four) may each be associated with CUs 2234 (e.g., 40 CUs), and some CUs 2234 may be disabled for yield management. CUs 2234 may have dedicated caches or shared caches (e.g., L2 caches) 2228 for consolidating all memory traffic on a single die. CU2234 may include threads and parallel processor cores, including instruction fetching and scheduling using a scheduler (S) 2212, a matrix core unit (MCU) 2216, and a shader core (SC) 2218 (e.g., execution units for scalar, vector, and matrix data types), and a load / store pipeline with an L1 cache 2220 and a local data share (LDS) 2214. The local data share may include, for example, sticky-note RAM with built-in arithmetic capabilities, allowing data sharing between threads in a workgroup. Instruction cache 2240 (e.g., for storing and providing instructions for performing operations described elsewhere in this document) and constant cache 2238 may be connected to one or more CUs and may be shared between two CUs. Matrix core 2216 can handle various data types, such as, but not limited to, INT8, FP8, FP16, BF16, and TF32 data types. Accelerated processing unit 2200 may include computation units 2234, which may be arranged in an array format, such as as a data parallel processor (DPP) array. The hyper-threaded dispatch processor 2242 can communicate with the compute unit 2234, and the command processor 2244 can read commands written by the host to memory-mapped registers in the system memory address space (not shown). When a command is completed, the command processor 2244 can send a hardware-generated interrupt to the host processor (e.g., a CCD). The memory controller 2236 can also directly access all device memories and system memory regions specified by the host. To satisfy read / write requests, the memory controller 2236 can perform the functions of a direct memory access (DMA) controller, including calculating the memory address offset based on the format of the requested data in memory. For example, one or more APIs described herein can be compiled into instructions that can be stored in the instruction cache 2240, then fetched by the instruction fetch logic in the processor 2240, decoded by the processor decoder or equivalent, scheduled (e.g., sequentially or out of order) by the scheduler or equivalent for execution, executed by the execution logic or equivalent, reordered, and then retired by the retirement logic or equivalent.The API (and / or compiled instructions including the API) can be stored in any storage device, either inside or outside the processor 2200 (e.g., in a cache and / or memory). The results of the API can be stored in storage devices, either inside or outside the processor 2200, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0224] Applications may include programs running on the main processor (e.g., a CCD) and programs running on one or more XCDs (referred to as kernels). Programs can be controlled by host commands that set internal base addresses and other configuration registers, specify data fields on which the accelerator processing unit 2200 can run, invalidate and flush caches on the accelerator processing unit 2200, and cause the accelerator processing unit 2200 to begin executing a program. A kernel can be referred to as a program executed by the accelerator processing unit 2200. Kernels can execute independently on each work item or as a group of work items, referred to as a wavefront, which can execute kernels on all (e.g., 64) work items in a single pass. The computation unit 2234 may include: a scalar arithmetic logic unit (ALU) that can operate on a single value for each wavefront (shared by all work items); a vector ALU that can operate on a unique value for each work item; a local data share 2214 that allows work items within a workgroup to communicate and share data; a scalar memory (not shown) that can transfer data between the scalar general-purpose registers (SGPRs) and memory via cache; and a vector memory that can transfer data between the vector general-purpose registers (VGPRs) and memory, including sampling of texture maps. Kernel control flow can be manipulated using scalar ALU instructions, which may include if / else statements, branches, and loops. Scalar ALU (SALU) and memory instructions can operate on the entire wavefront and operate on one or more SGPRs. Vector memory and ALU instructions can operate on all work items in the wavefront simultaneously.

[0225] In at least one embodiment, the acceleration processing unit 2200 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any operation described herein above or elsewhere. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any operation described herein above or elsewhere.

[0226] Figure 23 Processor 2300 is shown, including, but not limited to, a Zen architecture-based processor (e.g., Zen1, 2, 3, 4, 5, or other architectures) from AMD Inc. of Santa Clara, California, or other processors that share at least some of the components described herein. Processor 2300 includes one or more CPU dies 2302(1)-2302(N), where N is any integer greater than 1. CPU die 2302 may include any number of processor cores 2316 (e.g., for performing any operations described elsewhere herein) and any number of cache memories (e.g., for storing instructions and other information to perform any operations described elsewhere herein). For example, L2 cache unit 2318 may be coupled to processor core 2316, and processor core 2316 may share and / or be individually coupled to L2 cache unit 2318. Processor core 2316 can be coupled to L3 cache 2322 and / or a shared L3 cache, which can be the lowest level cache (LLC) 2322 used to access data and other information used by processor core 2316. One or more processor cores 2316 and one or more L2 cache units 2318 can be included in a core complex (CCX) 2320, which can include (e.g., 32MB) a shared cache (e.g., L3 cache 2322). Core complex 2320 can be manufactured onto a die (CCD or CPU die) 2302. For example, up to 12 core complexes 2320 can be configured in a processor along with 8 CPU dies 2302, thereby providing the processor with up to 96 processor cores 2316. For example, the "Zen 4c" core complex 2320 can include up to 8 cores 2316 and a shared 16MB L3 cache 2322. Two core complexes of these core complexes 2320 can be combined onto a single CPU die 2302, resulting in 16 cores per die and a total of 32MB of L3 cache per die 2322. Up to eight CPU dies 2302 can be combined with I / O units 2304 to provide up to 128 processor cores 2316 for the CPU. Up to four "Zen 4c" dies mentioned above can be combined to provide up to 64 processor cores 2316 for the CPU.

[0227] Processor 2300 may include various configurations for input / output operations, which will be further described herein. I / O unit 2304 may include one or more memory controllers 2306 capable of managing the memory usage of processor 2300 (e.g., DDR5 memory). I / O unit 2304 may include one or more SATA disk controllers for managing storage device 2312, and one or more Compute Express Links (CXL) for providing CPU-to-device and CPU-to-memory connectivity, which can be flexibly assigned to specific functions during server design. TM 1.1+ Memory controller 2314. I / O unit 2304 may include PCIe controller 2308 for connecting peripherals and other components connected to processor 2300. I / O unit 2304 may also include USB port 2310 for connecting to other components separate from processor 2300. CPU die 2302 may support any number of connections to I / O unit 2304, for example, one or two connections. As shown, I / O unit 2304 may include components further described herein, and I / O unit 2304 may be an I / O die accommodating several different components. Memory controller 2306, PCIe controller 2308, USB port 2310, SATA controller 2312 and / or CXL controller 2314 may be individually integrated into any location within processor 2300, or integrated in any group or combination.

[0228] Processor 2300 may include an Infinity Fabric 2324 interconnect (which may be similar to or based on a PCIe architecture) that provides connectivity between the CPU (e.g., CPU die 2302(1)-2302(N)), graphics processor 2326, inference engine 2332, and other components in a multi-chip architecture (e.g., security processor 2328 and I / O unit 2304). One or more AMD Infinity Fabrics TM Interconnect 2310 can be connected to CPU dies 2302(1)-2302(N) and used as a connection between CPUs. One or more Infinity Fabric connections 2310 can connect each CPU die 2302 to the I / O unit 2310.

[0229] In at least one embodiment, processor 2300 may include a central processing unit (CPU) and other related hardware and software described above and further herein. Processor 2300 may also include a graphics processor 2326. Graphics processor 2326 may be used for image generation and processing, as well as other computations and operations described further herein. Graphics processor 2326 may be based on AMD's RDNA 3 or 3.5 architecture, located in Santa Clara, California. Graphics processor 2326 may include a graphics computing die (GCD) and a memory cache die (MCD). The GCD may include any number of computing units (CUs) for graphics or other processing, such as operations performed by an arithmetic logic unit (ALU) described further herein. Graphics processor 2326 may include an L2 cache available for use by the computing units. The MCD (not shown) may include any number of memory cells and may include a cache (e.g., an L3 cache) and a memory interface for coupling to memory (e.g., memory 2342(1)-(N), where N is an integer). Components within the graphics processor 2326 can be connected using various methods, such as using Infinity Fabric 2324 interconnects, either internally or externally to the graphics processor 2326.

[0230] Inference engine 2332 can provide neural processing capabilities to processor 2300 for computational processes used in neural networks, deep learning, and other artificial intelligence-related operations, which will be further described herein. Processor 2300 may include: a security processor 2328 for managing the security of processor 2300; a display controller 2330 for controlling the display; a system management unit 2334 for managing and operating some or all components on processor 2300; a multimedia engine 2336 for audio and video operations; a fusion controller hub 2338 for managing USB, SATA, and PCIe connections to the processor; and a sensor fusion hub 2340 for managing sensors (e.g., accelerometers). Processor 2300 may also include memory 2342(1)-2342(N), where N is any integer. Memory may include different memory types, such as LPDDR5 and / or DDR5, or other memory described elsewhere herein.

[0231] To perform the operations further described herein, processor 2300 may include an execution pipeline including a front-end that may include a cache for storing instructions (e.g., an L1 cache) (not shown). A branch predictor may modify the instruction stream. Instructions may be decoded by a decoder, dispatched to a back-end for execution, and renamed. For example, the instruction fetch and decode pipeline may be dispatched to integer or floating-point execution operations, which may be scheduled by a scheduler and passed to vectors and / or general-purpose registers. Floating-point multipliers and / or addition operations may be processed, and an arithmetic logic unit (ALU) may also be used to perform computations, such as arithmetic and logical operations. The output of the computation unit may be coupled to a load / store queue that may be connected to a cache, such as an L1 cache and / or an L2 cache.

[0232] With respect to processor 2300 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents (e.g., AVX-512 instructions based on a SIMD model). These instructions or equivalents may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) internal or external to processor 2300. The results of the APIs may then be stored in storage devices internal or external to processor 2300, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0233] In at least one embodiment, processor 2300 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0234] Figure 24 An example of a processing core 2400 is shown, which may implement an Arm architecture (e.g., v9.0-A) or other processors that share at least some of the components described herein. Neoverse TMThe V2 core 2400 can be implemented within a DynamIQ shared unit (DSU) cluster via a DSU-110 interconnect 2454 for connecting one or more cores, for example, for parallel processing. Neoverse TM The V2 core can be implemented as a single core in a DSU cluster configured for direct interconnection, with or without L3 cache, listener filter, or listener control unit (SCU) logic (not shown). Neoverse TMThe V2 core may include a CPU bridge 2452 for connecting core 2400 to the DSU-110 interconnect. This bridge may also connect core 2400 to an external memory system and the remainder of the on-chip system. The L1 instruction memory system 2402 may fetch instructions from instruction cache 2404 and deliver instructions (e.g., one or more compileable APIs described herein) to instruction decoding unit 2410, for example, to perform some or all of the operations described above or elsewhere herein. The L1 instruction memory system 2402 may include L1 instruction cache 2404 (e.g., with 64-byte cache lines), L1 instruction translation back buffer (TLB) 2406 (e.g., natively supporting 4KB, 16KB, 64KB, and 2MB page sizes), and macro operation cache (MOP) 2408 (e.g., a 1536-entry, 4-way skewed associated L0 MOP cache), which may include decoded and optimized instructions for higher performance. Instruction decoding unit 2410 may decode AArch64 instructions into its internal format. Register renaming unit 2412 can perform register renaming to facilitate out-of-order execution and dispatch decoded instructions to various issue queues. Instruction issue unit 2414 can control when decoded instructions are dispatched to the execution pipeline, and it can include issue queues for storing instructions to be dispatched to the execution pipeline. Integer execution pipeline 2416 can be included in the execution pipeline and includes integer execution unit 2418 capable of performing arithmetic and logical data processing operations. Vector execution unit 2420 can be included in the execution pipeline and can execute advanced SIMD and floating-point arithmetic (FPU) 2422, execute Scalable Vector Extension (SVE) and Scalable Vector Extension 2 (SVE2) instructions 2424, and can also selectively execute cryptographic instructions 2426. Advanced SIMD may include a media and signal processing architecture that primarily adds instructions for audio, video, 3D graphics, image, and speech processing. The floating-point architecture provides support for single-precision and double-precision floating-point operations. The L1 data memory system 2430 executes load and store instructions, as well as service memory coherence requests. The L1 data memory system 2430 may include an L1 data cache 2432 and a fully associative L1 data TLB 2434, natively supporting 4KB, 16KB, and 64KB page sizes and 2MB and 512MB block sizes. A memory management unit (MMU) 2428 provides fine-grained memory system control through a set of virtual-to-physical address mappings and memory attributes, which can be stored in a translation table and saved to the TLB 2434 after address translation. The L2 memory system 2436 may include an L2 cache 2438 and can be connected to a DSU-110 2454 via an asynchronous CPU bridge 2452. TMThe V2 core 2400 supports a range of debugging, testing, and tracing options, including the tracing unit 2442, the tracing buffer 2440, and the embedded logic analyzer (ELA) 2448. Neoverse TM The V2 core 2400 implements the Statistical Analysis Extension (SPE) 2444, which provides a statistical view of the performance characteristics of the executed instructions. Software writers can leverage these views to optimize code for better performance. The Performance Monitoring Unit (PMU) 2446 provides a performance monitor that can be configured to collect statistics on the operation of each core and memory system. This information can be used for debugging and code analysis. The General Purpose Interrupt Controller (GIC) CPU interface 2450, when integrated with external allocator components, serves as a resource for supporting and managing interrupts in a cluster system. In a cluster, each Neoverse... TM A CPU bridge 2452 can exist between the V2 core 2400 and the DSU-110 2454. The CPU bridge 2452 controls buffering and synchronization between the core 2400 and the DSU-110 2454. The CPU bridge 2452 can be asynchronous to allow each core 2400 to use a different frequency, power, and area implementation point. The CPU bridge 2452 can operate synchronously without affecting other interfaces, such as, but not limited to, asynchronous debug and tracing interfaces.

[0235] In at least one embodiment, core 1800 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied, or otherwise perform any operation described herein above or elsewhere, when one or more branch instructions use one or more instruction operands stored in one or more stack structures. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied, or otherwise perform any operation described herein above or elsewhere, when one or more branch instructions use one or more instruction operands stored in one or more stack structures.

[0236] Figure 25 One or more chips including one or more tensor processing units (TPUs) 2500 are shown according to at least one embodiment. Figure 25The TPU 2500 may include an application-specific integrated circuit (ASIC), for example, for performing some or all of the operations described above or elsewhere herein, such as, but not limited to, machine learning workloads that accelerate matrix operations. The TPU 2500 may be an ASIC from Alphabet Corporation, Mountain View, California. The cloud TPU includes a cloud service that enables the TPU to be used as a scalable resource for processing tasks, such as, but not limited to, machine learning workloads that can run on frameworks such as, but not limited to, TensorFlow, PyTorch, and JAX.

[0237] Chip 2500 may include any number of TPUs, which may include tensor cores 2506. Tensor core 2506 may include one or more core sequencers 2508, vector processing units (VPUs) 2510, matrix multiplication units (MXUs) 2512(A)-2514(N) (where N is any integer greater than 1), and transpose permutation units 2516. Core sequencer 2508 may fetch instructions (e.g., VLIW (Very Long Instruction Word)) from the instruction memory (Imem) of core 2506, perform scalar operations using scalar data memory (Smem) and scalar registers (Sregs) (not shown), and forward vector instructions to vector processing units (VPUs) 2510. For example, an instruction may initiate eight operations: two scalar operations, two vector ALU operations, vector loading and storing, and queuing data into matrix multiplication and transpose units and a pair of slots for queuing data from them. The VPU2510 can perform vector operations using a large on-chip vector memory (Vmem) and vector registers (Vregs). The VPU 2510 can stream data to or from the MXU via a decoupled FIFO. The VPU 2510 can collect and distribute data to the Vmem using both data-level parallelism (2D matrix and vector function units) and instruction-level parallelism (8 operations per instruction). Large two-dimensional matrix multiplication units (MXUs) 2512(A)–2512(N) can, for example, use a systolic array to reduce area and power consumption, and use a large, software-controlled on-chip memory instead of a cache. The transpose-reduction-permute unit 2516 can perform matrix transpose, reduction, and permute operations (e.g., 128x128) on the VPU 2510 channels. A high-bandwidth memory 2504 can be used for on-chip applications and can be coupled to the host queue 2502, for example, via PCIe. One or more chips 2500 can be connected together for computing. For example, one or more chips 2500 can be connected to form a torus, such as a 2D torus. Chips 2500 can also include any number (e.g., four) of inter-core interconnect (ICI) links 2518, which can enable direct connections between chips to form a supercomputer.

[0238] For any processor in chip 2500 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device (e.g., cache and / or memory) external to or internal to any processor in chip 2500. The results of the APIs may then be stored in any storage device internal to or external to any processor in chip 2500, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0239] In at least one embodiment, chip 2500 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any operation described herein above or elsewhere. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any operation described herein above or elsewhere.

[0240] Figure 26A vector processor according to at least one embodiment is illustrated. The vector processor 2600 may support the RISC-V standard. The vector processor 2600 may include one or more cores 2610 (e.g., scalar units) and one or more vector processing units (VPUs) 2642 (e.g., vector units), which may, for example, perform some or all of the operations described above or elsewhere herein. The core 2610 may include an Andes Custom Extension (ACE) 2616, which can be used to deliver custom instructions to the processor 2600, for example, via ACP 2638. The core 2610 may include a 1-cycle multiplier and a 1-cycle instruction / data local memory (ILM / DLM) for improving parallelism by allowing simultaneous instruction fetching and data access. A memory management unit (MMU) 2624 manages system memory and cache, and provides branch execution, instruction pair issuance, L1 instruction / data cache, and local memory storage. The core 2610 may include a physical memory protection and programmable physical memory attribute unit (PMP / PPMA) 2622. Core 2610 may include a digital signal processor (DSP) 2628 and a floating-point unit (FPU) 2626, as well as a load-memory unit (LSU) 2632 for interaction with memory hierarchies (D$2634 and I$2630). Core 2610 may include a branch prediction unit 2618 and a multiplier unit 2620.

[0241] The vector processing unit (VPU) 2642 may include one or more vector function units (FU) 2646(A)-2646(N) (these vector function units may be linked together for parallel processing), a separate memory path for loading / storing RISC-V vectors (RVVs) via the ACE-RVV 2648 and the AndesStreaming port (ASP) 2644, and a vector load / store unit (VLSU) 2650.

[0242] The vector processor 2600 may include bus interfaces such as, but not limited to, a cache-accessible L2 cache port 2656, a non-cache-accessible MMIO port 2654, an input-output coherence port (IOCP) 2658 for a cacheless bus master, a local memory access port for accessing the ILM / DLM 2612 (which may be coupled to SRAM 2606) and the high-bandwidth vector memory (HVM) 2636, and a shared peripheral port (SPP) 2652 for external peripherals. Other memory ports include the LM slave port AXI 2602 and the HVM subordinate ports AXI 2604, MEM (AXI) 2662, and AXI 2660. The trace I / F 2614 can capture, encode, and transfer records of, for example, executed processor instructions via the Inst.Trace I / F 2608, and software tools can use the records of executed processor instructions to reconstruct the exact execution sequence of a program.

[0243] For any processor in processor 2600 and any components described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. APIs (and / or compiled instructions including APIs) may be stored in any storage device external to or internal to processor 2600 (e.g., in cache and / or memory). The results of the APIs may then be stored in storage devices internal to or external to processor 2600, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0244] In at least one embodiment, the vector processor 2600 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0245] Figure 27AA schematic diagram of an example multi-core tiled processor microarchitecture is shown. Figure 27A Multi-core tiled processors in a system can include language processing processors. For example... Figure 27A As shown, each “tile” in the processor architecture is a processing element bundled together using an on-chip network (NoC), which can be used to perform some or all of the operations described above or elsewhere in this document. For example, each tile may have an instruction dispatch 2704 and integer (INT) units 2706 and floating-point (FP) units 2708, a load-memory unit (LSU) 2712 for engaging with the memory hierarchy (data cache (D$) 2710 and instruction cache (I$) 2714), and a network (NET) interface 2716 for communicating with other tiles. Some tiles in processor 2700 may include a memory controller 2702 for managing and controlling memory, as further described herein. Processor 2700 may have a functionally sliced ​​architecture. Processor 2700 may reside on an application-specific integrated circuit (ASIC). Figure 27A The layout of an ASIC can be represented. Processor 2700 may include a coprocessor designed to execute instructions for a predictive model. A predictive model refers to any model configured to make predictions based on input data. The predictive model can use a classifier for classification predictions. The predictive model can be a machine learning model, such as, but not limited to, a tensor flow model, and processor 2700 is a tensor flow processor.

[0246] The processor 2700 can use different microarchitectures, and these microarchitectures will Figure 27B Each tile in the process represents a separate functional unit. Conversely, the functional tiles of the processor 2700 can be aggregated into multiple functional processing units (hereinafter referred to as "slices") 2704, each corresponding to a specific functional type (e.g., FP / INT 2718, NET 2720, MEM 2722). For example, as... Figure 27B As shown, each slice may correspond to a row of functional tiles extending in a north-south direction. Furthermore, the processor 2700 may also include communication channels for carrying data between tiles in different slices, each communication channel extending horizontally in an east-west direction. Each communication channel may be connected to each slice 2704 of the processor 2700.

[0247] The slices 2704 of processor 2700 may each correspond to different functions and may include arithmetic logic slices (e.g., FP / INT 2718), channel switching slices (e.g., NET 2720), and memory slices (e.g., MEM 2722). Arithmetic logic units can perform one or more arithmetic and / or logical operations on data received via communication channels to generate output data. Examples of arithmetic logic units may be matrix multiplication units and vector multiplication units. Memory slices include memory units that store data. Memory slices can provide data to other slices via communication channels. Memory slices can also receive data from other slices via communication channels. Channel switching slices can configurably route data from one communication channel to any other communication channel. For example, data from a first channel can be provided to a second channel via a channel switching slice. In some embodiments, a channel switching slice can be implemented as a crossbar switch. Each slice 2704 also includes its own instruction queue (not shown) for storing instructions and an instruction control unit (ICU) for controlling instruction execution. Instructions in a given instruction queue can only be executed by a tile in its associated functional slice and not by other slices of processor 2700.

[0248] By arranging the tiles of processor 2700 into different functional slices 2704, the on-chip instruction and control flow of processor 2700 can be separated from the data flow. For example, according to some embodiments, Figure 27B One of the arrows illustrates the instruction flow within the processor architecture. According to at least one embodiment, Figure 27B Another arrow in the diagram illustrates the data flow within the processor architecture. As shown, instruction and control flow can flow across the tiles of processor 2700 in a first direction (e.g., north-south along the length of a functional slice, as indicated by the first arrow), while data flow can flow across the tiles of processor 2700 in a second direction (e.g., east-west across a functional slice, as indicated by the second arrow), which is perpendicular to the first direction.

[0249] Different functional slices of processor 2700 can correspond to MEM 2722 (memory), VXM (vector execution module), MXM (matrix execution module), NIM (numerical interpretation module), and SXM (swapping and permutation module). Each slice can include N tiles, all of which can be controlled by the same instruction control unit (ICU) (not shown). Each slice can operate completely independently and can only be coordinated using barrier-like synchronization primitives or by the compiler using tractable determinism. Each tile of processor 2700 can correspond to an execution unit organized as ×MSIMD tiles. For example, each tile of the on-chip memory of processor 2700 can be organized to atomically store L-element vectors. Therefore, MEM slices with N tiles can work together to store or process large vectors (e.g., with a total of N×M elements).

[0250] Tiles within a slice can execute instructions in an "interleaved" manner, where instructions can be issued tile-by-tile within the slice over N cycle periods. Functional slices can be physically arranged on the chip to allow for efficient data flow for pipelined execution over hundreds of cycles for common patterns. The data flow can perform a single "u-turn" (direction change) corresponding to a single matrix operation before being written back to memory; in some embodiments, a particular data flow can change direction multiple times before writing the resulting data back to memory (due to multiple matrix and vector operations).

[0251] When using a processor 2700 with a function slice architecture (e.g., a TSP), the TSP compiler (not shown) generates an explicit plan of how the processor 2700 can execute programs (e.g., microprograms). The compiler can specify when each operation will be executed, which function slices will perform the work, and which STREAM registers will hold operands. The compiler can maintain a high-fidelity (cycle-accurate) model of the processor 2700 (e.g., the TSP) hardware state so that microprograms can coordinate data flow.

[0252] Processor 2700 (e.g., TSP) can use a web-hosted compiler that takes a model (e.g., an ML model, such as but not limited to a TensorFlow model) as input and issues a proprietary stream of instructions for processor 2700 (e.g., TSP). The compiler is responsible for coordinating the program's control and data flow and specifying any instruction-level parallelism by explicitly bundling instructions that can and should be executed concurrently so that they can be dispatched together. The main hardware architecture includes the architecture-visible streaming register file (STREAM), described in more detail below, which acts as a conduit for operands to flow from MEM slices (e.g., SRAM) to functional slices (and vice versa).

[0253] The MEM 2722 of the processor 2700 can be used as: (1) a storage for model parameters, microprograms, and data on which they operate; and (2) an on-chip network (NoC) for transferring data operands from the MEM to functional slices and returning computation results to the MEM. In some embodiments, the on-chip memory may consume approximately 75% of the chip area of ​​the processor 2700. In some embodiments, the on-chip memory of the MEM tile may include SRAM instead of DRAM due to the bandwidth requirements of the processor 2700. The on-chip memory capacity of the processor 2700 can be determined by: (i) the number of ML models that can reside on the chip simultaneously, (ii) the size of any given model, and (iii) partitioning of large models for adaptation to a multi-chip system. In some embodiments, the MEM system of the processor 2700 may provide multiple memory slices organized into two distinct hemispheres (referred to as “MEM WEST” and “MEM EAST”, respectively).

[0254] The memory slices in each hemisphere can be mirrored such that the slices are physically numbered {0,...L} in the eastern hemisphere and {L,...0} in the western hemisphere, such that memory slice 0 in each hemisphere corresponds to the slice of the VXM slice closest to the hemisphere, where each hemisphere comprises L slices. Data transfer towards the chip center can be referred to as inward, while data transfer towards the outer edge of the chip (easternmost or westernmost) can be referred to as outward. Although the memory hemispheres of processor 2700 can be referred to as east and west, it is understood that other names may be used to refer to different memory hemispheres in other embodiments.

[0255] In some embodiments, streaming register files (referred to as STREAMS) transfer operands and results between the SRAM of a MEM slice of the processor 2700 and the functional slice. In some embodiments, multiple MEM slices (e.g., 2 to 10 adjacent MEM slices) can be physically organized into sets. Each slice set can be located between a pair of STREAMS register files, allowing each slice to read from or write to the STREAMS registers in either direction. By placing STREAMS register files between MEM slice sets, the number of cycles required to transfer data operands across hemispheres can be reduced (e.g., reduced by a factor corresponding to the number of slices per set). The number of slices per set can be configured based on the distance of data transfer within a single clock cycle.

[0256] for Figure 27A Any processor and any component described above or elsewhere herein, one or more APIs or equivalents described herein may, for example, be compiled into instructions or equivalents that may be fetched by instruction fetching logic or equivalents, decoded by processor decoder or equivalents, scheduled (e.g., sequentially or out of order) for execution by scheduler or equivalents, executed by execution logic or equivalents, reordered, and then retired by retirement logic or equivalents. The API (and / or compiled instructions including the API) may be stored in any storage device internal or external to processor 2700 (e.g., in cache and / or memory). The results of the API may then be stored in storage devices internal or external to processor 2700, including registers, DRAM, flash memory, SRAM, cache, or other memory equivalents.

[0257] In at least one embodiment, processor 2100 may include one or more circuits configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein. The one or more circuits may be software-configurable to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0258] Software Structure

[0259] The following figures illustrate, in a non-limiting manner, examples of software structures for implementing at least one embodiment.

[0260] Figure 28A software stack of a programming platform according to at least one embodiment is illustrated. The programming platform may include a platform for accelerating computational tasks by utilizing hardware on a computing system. In at least one embodiment, software developers can access the programming platform through libraries, compiler instructions, and / or extensions to programming languages. The programming platform may be CUDA, Radeon Open Computing Platform (“ROCm”), OpenCL (OpenCL... TM Developed by the Khronos Group), SYCL, or Intel OneAPI.

[0261] The software stack 2800 of the programming platform can provide an execution environment for application 2801. Application 2801 may include any computer software capable of being launched on the software stack 2800. Application 2801 may include artificial intelligence (“AI”) / machine learning (“ML”) applications, high-performance computing (“HPC”) applications, virtual desktop infrastructure (“VDI”), or data center workloads.

[0262] Application 2801 and software stack 2800 run on hardware 2808. Hardware 2808 may include one or more GPUs, CPUs, FPGAs, AI engines, and / or other types of computing devices supporting programming platforms. Software stack 2800 may be vendor-specific and compatible only with vendor-specific devices, such as CUDA, ROCm, OneAPI, OpenCL, or other implementations. Hardware 2808 may include a host connected to one or more devices that can be accessed via application programming interface (“API”) calls to perform computational tasks. In at least one embodiment, the devices within hardware 2808 may include GPUs, FPGAs, AI engines, or other computing devices (but may also include CPUs) and their memory, while the host within hardware 2808 may include CPUs (but may also include computing devices) and their memory. For any hardware 2808 described above or elsewhere herein, the one or more APIs described herein may, for example, be compiled into instructions that may be fetched by instruction fetching logic, decoded by a processor decoder, scheduled (e.g., sequentially or out of order) for execution by a scheduler, executed by execution logic, reordered, and then retired by retirement logic. The API (and / or compiled instructions including the API) may be stored in any storage device (e.g., cache and / or memory) internal or external to the hardware 2808. The results of the API may be stored in storage devices internal or external to the hardware 2808, including registers, DRAM, flash memory, SRAM, cache, or other memory. One or more APIs described herein may receive calls. One or more APIs described herein may communicate with a library or a portion of a library to perform the function described by the call. One or more APIs described herein may receive calls and communicate with a library or a portion of a library to perform the function described by the call.

[0263] The software stack 2800 of the programming platform may include multiple libraries 2803, a runtime 2805, optional drivers / interfaces 2807, and device kernel drivers 2808. Each library 2803 may include data and programming code that can be used by computer programs and utilized during software development. Library 2803 may include pre-written code and subroutines, classes, values, type specifications, configuration data, documentation, help data, and / or message templates. Library 2803 may include functionality optimized for execution on one or more types of devices. Library 2803 may include functionality for performing mathematical, deep learning, and / or other types of operations on the device. Library 2803 may be associated with corresponding APIs 2802, which may include one or more APIs for exposing the functionality implemented in library 2803. A processor (e.g., CPU, GPU) may execute, call, or otherwise use one or more APIs to determine kernel priority. For example, a first kernel (e.g., a parent kernel) may launch a second kernel (e.g., a child kernel), and the processor may use the second kernel to launch an additional kernel (e.g., a grandchild kernel) independent of the first kernel. The processor can execute APIs or call APIs from memory to support dynamic stream priorities (e.g., updating priorities when performing operations using streams). For example, when the processor executes the API, it allows the programmer to copy stream priorities from one stream to one or more other streams.

[0264] Software stack 2800 may include APIs that support dynamic stream prioritization (e.g., updating priority while performing operations on the stream), allowing programmers to set the stream's priority at any time after the stream is created. Software stack 2800 may include APIs that support dynamic stream prioritization (e.g., updating priority while performing operations on the stream), allowing programmers to obtain the stream's current priority, where the priority is one of several attributes of the stream. Software stack 2800 may include APIs that support dynamic stream prioritization (e.g., updating priority while performing operations on the stream), allowing programmers to obtain the stream's current priority as a single attribute. Software stack 2800 may include APIs that support dynamic stream prioritization (e.g., updating priority while performing operations on the stream), allowing programmers to start the kernel to perform operations on the stream at a set priority, which may be different from the stream priority. The software stack 2800 may include an API for indicating whether an object (e.g., a thread synchronization object, such as, but not limited to, a barrier) tracks whether all data movement operations of a set of threads running on the GPU have a specified state after a specified time period, wherein the specified state may be a state indicating that data has been moved and is ready for use, and is specified using expected parity values ​​as input to the API.

[0265] Software stack 2800 may include one or more APIs for updating the kernel. The processor may execute APIs or call APIs from memory to update existing APIs, thereby supporting a context-independent kernel. This allows programmers to add kernel nodes to a graph without a graph context, so that the graph context can be dynamically associated with the kernel at runtime. Software stack 2800 may include one or more APIs that allow programmers to obtain kernel identifiers and graph contexts as separate parameters from kernel nodes, thereby enabling parameter retrieval from both the kernel and the context-independent kernel. Software stack 2800 may include one or more APIs that use parallel processors (e.g., but not limited to one or more graphics processing units) to initiate task graphs (e.g., task graphs) and execute one or more task graphs (e.g., including one or more programs).

[0266] Software stack 2800 may include one or more APIs for associating one or more instructions with one or more memory sorting operations (e.g., but not limited to fence or memory barrier operations). Instructions may be associated with one or more domains, causing memory sorting operations to execute in relation to one or more specific domains without interfering with instructions in other domains. APIs may indicate that a thread has reached (e.g., reached a thread synchronization barrier) or completed a certain phase of work associated with an asynchronous data movement operation on the GPU. Software stack 2800 may include one or more APIs that allow programmers to manually indicate an expected transaction count when a thread completes a certain phase of work; this transaction count can be used to update an object used to track whether all data movement operations for a set of threads have been completed.

[0267] Application 2801 can be written as source code and then compiled into executable code, as described below. Figure 29 and Figure 30 As discussed in more detail, the executable code of application 2801 can run, at least partially, within the execution environment provided by software stack 2800. During the execution of application 2801, code that needs to run on the device (rather than the host) may be encountered. In this case, runtime 2805 can be invoked to load and launch the required code on the device. Runtime 2805 can include any technically feasible runtime system capable of supporting the execution of application 2801.

[0268] Runtime 2805 can be implemented as one or more runtime libraries associated with the corresponding API (which is shown as API 2804). One or more such runtime libraries may include functions for memory management, execution control, device management, error handling, and / or synchronization, etc. Memory management functions may include functions for allocating, dealing with, and copying device memory, as well as functions for transferring data between host memory and device memory. Execution control functions may include functions for starting functions on the device (sometimes referred to as "kernels" when the function is a global function that can be called from the host) and setting attribute values ​​in buffers maintained by the runtime libraries so that a given function can be executed on the device.

[0269] Runtime libraries and the corresponding API 2804 can be implemented in any technically feasible manner. One (or any number) APIs can expose a set of low-level functions for fine-grained control of the device, while another (or any number) APIs can expose a set of high-level functions for such functions. High-level runtime APIs can be built on top of low-level APIs. One or more runtime APIs can be language-specific APIs, which can be layered on top of language-independent runtime APIs.

[0270] Optional drivers or interfaces 2807 can be implemented, for example, for CUDA and ROCm implementations, which will be described further below. Optional drivers / interfaces 2807 can be associated with optional driver or interface APIs, such as, but not limited to, the CUDA and / or ROCm APIs.

[0271] One or more processors disclosed in the “processing system” may execute, access, or otherwise use the software stack 2800. For example, the system-on-a-chip 1500, parallel processor 1600, graphics multiprocessor 1634, processor 1700, processor 1800, accelerator 1900, neuromorphic processor 2005, supercomputer 2100, acceleration processing unit 2200, processor 2300, processor 2400, tensor processing unit 2500, processor 2600, and language processing unit 2700 may execute, use, call, or otherwise implement (e.g., by accessing memory) one or more APIs included in the software stack 2800.

[0272] Device kernel driver 2808 can be configured to facilitate communication with underlying devices. Device kernel driver 2808 can provide low-level functionality for APIs (such as, but not limited to, API 2804) and / or other software. Device kernel driver 2808 can be configured to compile intermediate representation (“IR”) code into binary code at runtime. For CUDA or other implementations (such as, but not limited to, ROCm, OneAPI, or OpenCL), device kernel driver 2808 can compile non-hardware-specific parallel thread execution (“PTX”) IR code into binary code for a specific target device at runtime (and cache the compiled binary code), which is sometimes referred to as “finalized” code. Doing so allows the finalized code to run on a target device that may not have existed when the source code was initially compiled into PTX code. Alternatively, device source code can be compiled into binary code offline without device kernel driver 2808 compiling the IR code at runtime.

[0273] Processors described elsewhere in this document (e.g., but not limited to) Figure 15-27B The processor (in the document) may include one or more circuits for causing one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be configured by software (e.g., software stack 2800) to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0274] According to at least one embodiment, Figure 28 The software stack 2800 can execute in a CUDA implementation. The CUDA software stack 2800, on which application 2801 can be launched, may include CUDA libraries 2803, CUDA runtime 2805, CUDA driver 2807, and device kernel driver 2808. The CUDA software stack 2800 can execute on hardware (e.g., a graphics multiprocessor 1634, which may include a CUDA-enabled GPU developed by NVIDIA Corporation in Santa Clara, California).

[0275] Application 2801, CUDA runtime 2805, and device kernel driver 2808 can perform the functions described above and elsewhere in this document. CUDA driver 2807 may include a library (libcuda.so) that implements CUDA driver API 2806. Similar to CUDA runtime API 2804 implemented by the CUDA runtime library (cudart), CUDA driver API 2806 exposes functions for memory management, execution control, device management, error handling, synchronization, and / or graphics interoperability. 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 processes), and module (similar to dynamically loaded libraries) management. Compared to the high-level CUDA runtime API 2804, CUDA driver API 2806 can serve as a low-level API, providing finer-grained device control, especially in terms of context and module loading. CUDA Driver API 2806 can expose context management functions not exposed in CUDA Runtime API 2804. CUDA Driver API 2806 can also be language-agnostic, supporting technologies such as OpenCL in addition to CUDA Runtime API 2804. Furthermore, development libraries, including CUDA Runtime 2805, can be considered separate from driver components, including user-mode CUDA Driver 2807 and kernel-mode device driver 2808 (sometimes referred to as the "display" driver).

[0276] CUDA library 2803 may include mathematical libraries, deep learning libraries, parallel algorithm libraries, and / or signal / image / video processing libraries, which parallel computing applications (such as, but not limited to, application 2801) may utilize. CUDA library 2803 may include mathematical libraries, such as, but not limited to, the cuBLAS library (an implementation of the basic linear algebra subroutine (“BLAS”) for performing linear algebra operations), the cuFFT library (for computing the Fast Fourier Transform (“FFT”)), and the cuRAND library (for generating random numbers), etc. CUDA library 2803 may include deep learning libraries, such as, but not limited to, the cuDNN primitive library for deep neural networks and the TensorRT platform for high-performance deep learning inference, etc.

[0277] In at least one embodiment, the processor described elsewhere herein (e.g., but not limited to Figure) Figure 15-27BThe processor (in the document) may include one or more circuits for causing one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be configured by software (e.g., software stack 2800) to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0278] According to at least one embodiment, Figure 28 The software stack 2800 can execute in the ROCm implementation. Application 2801 can be launched on the ROCm software stack 2800, which includes a language runtime 2803, a system runtime 2805, a thunk 2807, and a ROCm kernel driver 2808. The ROCm software stack 2800 executes on hardware 2809, which may include a GPU that supports ROCm, developed by AMD Inc. of Santa Clara, California.

[0279] Application 2801 is executable and can be combined with the above. Figure 28 Similar functionality to that discussed. Furthermore, the language runtime 2803 and system runtime 2805 can execute functions related to the above. Figure 28 The discussed runtime 2805 has similar functionality. The difference between the language runtime 2803 and the system runtime 2805 is that the system runtime 2805 is a language-independent runtime that implements the ROCr System Runtime API 2804 and uses the Heterogeneous System Architecture (“HSA”) runtime API. The HSA runtime API may include a streamlined user-mode API that exposes interfaces for accessing and interacting with the AMD GPU, including functions for memory management, execution control via kernel architecture dispatch, error handling, system and agent information, and runtime initialization and shutdown. Unlike the system runtime 2805, the language runtime 2803 may be an implementation of the language-specific runtime API 2802, which sits in a layer above the ROCr System Runtime API 2804. The language runtime API may include the Heterogeneous Computing Portable Interface (“HIP”) language runtime API, the Heterogeneous Computing Compiler (“HCC”) language runtime API, or the OpenCL API, etc. HIP is an extension of the C++ programming language, offering a functionally similar version to the CUDA mechanism. Furthermore, the HIP runtime API can include features related to the above. Figure 28The discussion covers functions similar to those in the CUDA runtime API, such as, but not limited to, memory management, execution control, device management, error handling, and synchronization.

[0280] The Thunk(ROCt)2807 can be interface 2806, which is used to interact with the underlying ROCm driver 2808. The ROCm driver 2808 can be the ROCk driver, a combination of the AMD GPU driver and the HSA core driver (amdkfd). The AMD GPU driver can be a device core driver developed by AMD for GPUs, performing functions similar to those described above. Figure 28 The device kernel driver discussed is similar to the 2809. An HSA kernel driver can be a driver that allows different types of processors to share system resources more efficiently through hardware features.

[0281] Various libraries (not shown) may be included in the ROCm software stack 2800 on top of the language runtime 2803, and provide integration with the above. Figure 28 The discussion focuses on functionality similar to CUDA library 2803. Various libraries can include mathematical libraries, deep learning libraries, and / or other libraries, such as, but not limited to, the hipBLAS library which implements functionality similar to CUDA cuBLAS, the rocFFT library for computing FFTs similar to CUDA cuFFT, etc.

[0282] Processors described elsewhere in this document (e.g., but not limited to) Figure 15-27B The processor (in the document) may include one or more circuits for causing one or more instruction operand stack structures to be copied or otherwise performing any operation described above or elsewhere herein when one or more branch instructions use one or more instruction operands stored in one or more stack structures. One or more circuits may be configured by software (e.g., software stack 2800) to cause one or more instruction operand stack structures to be copied or otherwise perform any operation described above or elsewhere herein when one or more branch instructions use one or more instruction operands stored in one or more stack structures.

[0283] According to at least one embodiment, Figure 28The software stack 2800 can execute in an OpenCL implementation. The OpenCL software stack 2800, on which the application 2801 can be launched, may include the OpenCL framework 2803, the OpenCL runtime 2805, and a driver 2808. The OpenCL software stack 2800 can execute on non-vendor-specific hardware 2809. Because devices developed by different vendors support OpenCL, specific OpenCL drivers may be required for interoperability with the hardware of those vendors.

[0284] Application 2801, OpenCL runtime 2805, device kernel driver 2808, and hardware 2809 can execute in combination with the above. Figure 28 Other implementations of the application 2801, runtime 2805, device kernel driver 2808, and hardware 2809 discussed have similar functionality. The application 2801 may also include an OpenCL kernel (not shown), whose code will execute on the device.

[0285] OpenCL can define a "platform" that allows a host to control devices connected to it. The OpenCL framework provides platform-level APIs and runtime APIs, shown as Platform API 2802 and Runtime API 2804, respectively. Runtime API 2804 uses contexts to manage kernel execution on devices. Each identified device can be associated with a corresponding context, which Runtime API 2804 uses to manage the device's command queue, program objects, kernel objects, shared memory objects, and so on. Platform API 2802 exposes functions that allow the use of device contexts to select and initialize devices, submit work to devices via command queues, and enable data transfer with devices. In addition, the OpenCL framework provides various built-in functions (not shown), including mathematical functions, relational functions, and image processing functions.

[0286] The OpenCL framework 2803 may also include a compiler (not shown). Source code can be compiled offline before application execution or online during application execution. Unlike CUDA and ROCm, OpenCL applications can be compiled online by a compiler, which represents any number of compilers that can be used to compile source code and / or IR code (e.g., but not limited to Standard Portable Intermediate Representation (“SPIR-V”) code) into binary code. Alternatively, OpenCL applications can be compiled offline before execution.

[0287] In at least one embodiment, the processor described elsewhere herein (e.g., but not limited to) Figure 15-27BThe processor (in the document) may include one or more circuits for causing one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise performing any of the operations described above or elsewhere herein. One or more circuits may be configured by software (e.g., software stack 2800) to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0288] According to at least one embodiment, the software may be supported by a programming platform configured to support various programming models, middleware and / or libraries, and frameworks that the application may rely on. The application may be an AI / ML application implemented using, for example, a deep learning framework (e.g., but not limited to, MXNet, PyTorch, or TensorFlow), which may rely on libraries such as, but not limited to, cuDNN, the NVIDIA Collective Communication Library (“NCCL”), and / or the NVIDIA Developer Data Loading Library (“DALI”) CUDA library to provide accelerated computation on the underlying hardware.

[0289] The programming platform can be a combination of the above. Figure 28 The platform described is one of CUDA, ROCm, or OpenCL. The programming platform can support various programming models, which can be abstractions of the underlying computing system that allow the expression of algorithms and data structures. Programming models can expose features of the underlying hardware to improve performance. Programming models may include CUDA, HIP, OpenCL, C++ Accelerated Massive Parallelism (“C++AMP”), Open Multiprocessing (“OpenMP”), Open Accelerators (“OpenACC”), and / or Vulkan Compute.

[0290] Libraries and / or middleware can provide abstract implementations of programming models. Such libraries may include data and programming code that computer programs can use and leverage during software development. Such middleware may include software that provides services to applications beyond those offered by the programming platform. Libraries and / or middleware may include cuBLAS, cuFFT, cuRAND, and other CUDA libraries, or rocBLAS, rocFFT, rocRAND, and other ROCm libraries. Furthermore, libraries and / or middleware may include the NCCL and ROCm communication collection library (“RCCL”) libraries that provide communication routines for GPUs, the MIOpen library for accelerating deep learning, and / or the Eigen library for linear algebra, matrix and vector operations, geometric transformations, numerical solvers, and related algorithms.

[0291] Application frameworks may depend on libraries and / or middleware. Each application framework can be a software framework that provides a standard structure for implementing application software. Returning to the AI / ML example discussed above, AI / ML applications can be implemented using frameworks such as, but not limited to, deep learning frameworks like Caffe, Caffe2, TensorFlow, Keras, PyTorch, or MxNet.

[0292] In at least one embodiment, the processor described elsewhere herein (e.g., but not limited to) Figure 15-27B The processor (in the document) may include one or more circuitry for causing one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise performing any of the operations described above or elsewhere herein. One or more circuitry may be configured by software (e.g., the programming platform described herein) to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures, or otherwise perform any of the operations described above or elsewhere herein.

[0293] Figure 29 The method for using at least one embodiment in the above is shown. Figure 28Compiled code executed on one of the programming platforms shown. Compiler 2901 is configured to receive source code 2900, compile source code 2900, and output executable file 2910. Compiler 2901 can be configured to convert source code 2900 into host executable code 2907 for execution on a host and device executable code 2908 for execution on a device. Source code 2900 can be compiled offline before executing the application or online during application execution. Source code 2900 can include code in any programming language supported by compiler 2901, such as, but not limited to, C++, C, Fortran, etc. Source code 2900 can be included in a single source file containing both host code and device code, indicating the location of the device code. The single source file can be a .cu file including CUDA code, or a .hip.cpp file including HIP code, or a file in other formats including both host code and device code. Alternatively, source code 2900 can include multiple source code files instead of a single source file, with the host code and device code separated into these files. Compiler 2901 includes or can access one or more libraries to identify API call sequences to execute a single fused API, wherein the single fused API is a combination of two or more APIs. In at least one embodiment, compiler 2901 may be an NVIDIA CUDA compiler (“NVCC”) for compiling CUDA code in .cu files, or an HCC compiler for compiling HIP code in .hip.cpp files, or other compilers.

[0294] Compiler 2901 can be configured to compile source code 2900 into host executable code 2907 for execution on a host and device executable code 2908 for execution on a device. The operations performed by compiler 2901 include parsing source code 2900 into an abstract system tree (AST), performing optimizations, and generating executable code. When source code 2900 comprises a single source file, compiler 2901 can separate the device code from the host code in that single source file, compile the device code and host code into device executable code 2908 and host executable code 2907 respectively, and link the device executable code 2908 and host executable code 2907 together to form a single file.

[0295] Compiler 2901 may include compiler front-end 2902, host compiler 2905, device compiler 2906, and linker 2909. Compiler front-end 2902 may be configured to separate device code 2904 from host code 2903 in source code 2900. In at least one embodiment, device code 2904 may be compiled by device compiler 2906 into device executable code 2908, which may include binary code or IR code as described above. Separately, host code 2903 may be compiled by host compiler 2905 into host executable code 2907. For other compilers such as NVCC (e.g., but not limited to oneAPI, ROCm, and OpenCL compilers), host compiler 2905 may be a general-purpose C / C++ compiler that outputs native object code, while device compiler 2906 may be a low-level virtual machine (“LLVM”) based compiler that forks the LLVM compiler infrastructure and outputs PTX code or binary code. For HCC, both host compiler 2905 and device compiler 2906 may be LLVM based compilers that output object binary code.

[0296] After compiling source code 2900 into host executable code 2907 and device executable code 2908, linker 2909 can link host executable code 2907 and device executable code 2908 together to form executable file 2910. The host's native object code and the device's PTX or binary code can be linked together in an executable and linkable format (“ELF”) file, a container format for storing object code. Host executable code 2907 and device executable code 2908 can take any suitable format, such as, but not limited to, binary code and / or IR code. In at least one embodiment, for CUDA, host executable code 2907 may include native object code, while device executable code 2908 may include code in a PTX intermediate representation. In at least one embodiment, for ROCm, both host executable code 2907 and device executable code 2908 can include object binary code. Other implementations (e.g., but not limited to oneAPI, OpenCL) are considered and can be performed similarly to the CUDA and ROCm implementations described above.

[0297] Source code 2900 can be transformed before compilation. The source code is passed through a transformation tool (not shown) that converts source code 2900 into transformed source code. Compiler 2901 can be used to compile the transformed source code into host-executable code 2907 and device-executable code 2908, a process similar to compiler 2901 compiling source code 2900 into host-executable code 2907 and device-executable code 2908, as described above. Figure 29The discussion.

[0298] The transformations performed by the transformation tools can be used to port Source Code 2900 to environments different from the environment it was originally intended to run in. The transformation tools may include a HIP converter, which is used to "hipify" CUDA code intended for the CUDA platform into HIP code that can be compiled and executed on the ROCm platform. The transformation of Source Code 2900 may include parsing Source Code 2900 and translating calls to APIs provided by one programming model (e.g., CUDA) into corresponding calls to APIs provided by another programming model (e.g., HIP), as described below. Figure 30 Let's discuss this in more detail. Returning to the example of HIP-enhanced CUDA code, calls to the CUDA runtime API, CUDA driver API, and / or CUDA libraries can be translated into corresponding HIP API calls. The automatic conversion performed by conversion tool 2901 may sometimes be incomplete, requiring additional manual intervention to fully port the source code 2900.

[0299] One or more techniques described herein can leverage other methods of converting one type of code into another to achieve interoperability between different device architectures. In at least one embodiment, an application for one platform (e.g., a CUDA application) can be compiled into code for implementation on another platform (e.g., an AMD processor, an Intel processor, or another processor). For example, source code 2900 may include source code for one platform (e.g., CUDA). Compiler 2901 can compile source code 2900 into an executable file 2910 that can be used by another platform (e.g., AMD or Intel). Programming toolkits can allow applications for one platform (e.g., CUDA) to be compiled (e.g., natively compiled) against another platform (e.g., AMD or Intel). For example, a GPGPU programming toolkit can allow CUDA applications to be natively compiled against AMD GPUs. Programs (e.g., CUDA programs) or their build systems do not require modification or conversion into other languages ​​before being compiled into code for another platform. The compiler can accept the same command-line options and programming cognates (e.g., CUDA cognates) as another compiler (e.g., nvcc for CUDA), serving as a direct replacement for the installation of emulation toolkits (e.g., the NVIDIA CUDA Toolkit), so existing build tools and scripts (e.g., cmake) will work without further modification. In at least one embodiment, an nvcc-cognate CUDA can be compiled for AMD GPUs (including PTX asm) using an nvcc-compatible compiler. Implementations of the CUDA runtime and driver APIs for AMD GPUs can be used. Libraries (e.g., open-source wrapper libraries) can provide APIs by delegating to the corresponding ROCm libraries, such as the "CUDA-X" API. Example implementations include SCALE from Spectral Compute in London, UK. SCALE does not provide a new method for writing GPGPU software, but rather allows programs written in the popular CUDA language to be directly compiled for AMD GPUs. Additional implementations could include the Clang compiler, which provides language front-ends and tooling infrastructure for languages ​​in the C family of languages ​​(C, C++, Objective C / C++, OpenCL, CUDA, and RenderScript).In at least one embodiment, the compiler described herein (e.g., but not limited to compiler 2901, compiler 2905, and / or compiler 2906) may include one or more circuits for compiling code (e.g., CUDA, HIP, OpenCL, OneAPI, or others) to cause one or more instruction operand stack structures to be copied in cases where one or more branch instructions use one or more instruction operands stored in one or more stack structures, or to otherwise perform any of the operations described above or elsewhere herein.

[0300] Figure 30 A system 3000 is shown, configured to compile and execute CUDA source code 3010 using different types of processing units according to at least one embodiment. The system 3000 includes CUDA source code 3010, CUDA compiler 3050, host executable code 3070(1), host executable code 3070(2), CUDA device executable code 3084, CPU 3090, CUDA-enabled GPU 3094, GPU 3092, CUDA to HIP conversion tool 3020, HIP source code 3030, HIP compiler driver 3040, HCC 3060, and HCC device executable code 3082.

[0301] CUDA source code 3010 can be a collection of human-readable code in the CUDA programming language. The CUDA programming language can be an extension of the C++ programming language, including mechanisms for defining device code and distinguishing between device code and host code. Device code can include source code that, after compilation, can be executed in parallel on a device. A device can be a processor optimized for parallel instruction processing, such as, but not limited to, a CUDA-enabled GPU 3090, GPU 3092, or other GPGPUs. Host code is source code that, after compilation, can be executed on a host machine. A host machine is a processor optimized for sequential instruction processing, such as, but not limited to, a CPU 3090.

[0302] CUDA source code 3010 may include any number (including zero) of global functions 3012, any number (including zero) of device functions 3014, any number (including zero) of host functions 3016, and any number (including zero) of host / device functions 3018. Global functions 3012, device functions 3014, host functions 3016, and host / device functions 3018 can be mixed within CUDA source code 3010. Each global function 3012 can be executed on the device and called from the host. Therefore, one or more global functions 3012 can act as entry points for the device. Each global function 3012 can be a kernel. In a technique called dynamic parallelism, one or more global functions 3012 can define a kernel that can be executed on the device and called from the device. During execution, the kernel can be executed in parallel N times by N different threads on the device (where N is any positive integer).

[0303] Each device function 3014 can be executed on a device and can only be called from that device. Each host function 3016 can be executed on a host and can only be called from that host. Each host / device function 3016 can define a host version of a function that can be executed on a host and called only from that host, and a device version of a function that can be executed on a device and called only from that device.

[0304] CUDA source code 3010 can also include any number of calls to any number of functions, which can be defined through CUDA runtime API 3002. CUDA runtime API 3002 can include any number of functions that execute on the host to allocate and deallocate device memory, transfer data between host memory and device memory, manage systems with multiple devices, etc. CUDA source code 3010 can also include any number of calls to any number of functions, which can be specified in any number of other CUDA APIs. CUDA APIs can be any API designed for use by CUDA code. CUDA APIs can include CUDA runtime API 3002, CUDA driver APIs, APIs for any number of CUDA libraries, etc., including any APIs described elsewhere in this document. Compared to CUDA runtime API 3002, CUDA driver APIs can be lower-level APIs but can provide finer-grained control over devices. Examples of CUDA libraries include cuBLAS, cuFFT, cURAND, cuDNN, etc.

[0305] The CUDA compiler 3050 can compile input CUDA code (e.g., CUDA source code 3010) to generate host executable code 3070(1) and CUDA device executable code 3084. The CUDA compiler 3050 can be, but is not limited to, an NVCC. The host executable code 3070(1) can be a compiled version of the host code included in the input source code, which can be executed on a CPU 3090. The CPU 3090 can be any processor optimized for sequential instruction processing.

[0306] CUDA device executable code 3084 may be a compiled version of the device code included in the input source code, which can be executed on a CUDA-enabled GPU 3094. CUDA device executable code 3084 may include binary code. CUDA device executable code 3084 may include IR code (such as, but not limited to, PTX code), which is further compiled at runtime by the device driver into binary code for a specific target device (e.g., a CUDA-enabled GPU 3094). CUDA-enabled GPU 3094 may include any processor optimized for parallel instruction processing and supporting CUDA. CUDA-enabled GPU 3094 may be developed by NVIDIA Corporation, located in Santa Clara, California.

[0307] The CUDA to HIP conversion tool 3020 can be configured to convert CUDA source code 3010 into functionally similar HIP source code 3030. The HIP source code 3030 may include a collection of human-readable code written in the HIP programming language. The HIP code may include human-readable code written in the HIP programming language. The HIP programming language may include extensions to the C++ programming language that include functionally similar versions of the CUDA mechanisms for defining device code and distinguishing it from host code. The HIP programming language may include a subset of the functionality of the CUDA programming language. For example, the HIP programming language may include a mechanism for defining a global function 3012, but such a HIP programming language may lack support for dynamic parallelism, so the global function 3012 defined in the HIP code may only be called from the host.

[0308] HIP source code 3030 may include any number (including zero) of global functions 3012, any number (including zero) of device functions 3014, any number (including zero) of host functions 3016, and any number (including zero) of host / device functions 3018. HIP source code 3030 may also include any number of calls to any number of functions specified in the HIP runtime API 3032. The HIP runtime API 3032 may include functionally similar versions of a subset of functions contained in the CUDA runtime API 3002. HIP source code 3030 may also include any number of calls to any number of functions specified in any number of other HIP APIs. The HIP API can be any API designed for use by HIP code and / or ROCm. The HIP API may include the HIP runtime API 3032, the HIP driver API, APIs for any number of HIP libraries, APIs for any number of ROCm libraries, etc.

[0309] The CUDA to HIP conversion tool 3020 can convert every kernel call in CUDA code from CUDA syntax to HIP syntax, and can convert any number of other CUDA calls in CUDA code into any number of other functionally similar HIP calls. CUDA calls can include calls to functions specified in the CUDA API, and HIP calls can include calls to functions specified in the HIP API. The CUDA to HIP conversion tool 3020 can convert any number of calls to functions specified in the CUDA runtime API 3002 into any number of calls to functions specified in the HIP runtime API 3032.

[0310] The CUDA to HIP conversion tool 3020 may include a tool called hipify-perl, which performs a text-based conversion process. The CUDA to HIP conversion tool 3020 may also include a tool called hipify-clang, which, compared to hipify-perl, performs a more complex and robust conversion process, including parsing the CUDA code using clang (a compiler front-end) and then converting the resulting symbols. Converting CUDA code to HIP code may include modifications beyond those performed by the CUDA to HIP conversion tool 3020 (e.g., manual editing).

[0311] HIP compiler driver 3040 may include a front-end that identifies target device 3046 and then configures a compiler compatible with target device 3046 to compile HIP source code 3030. Target device 3046 may include a processor optimized for parallel instruction processing. HIP compiler driver 3040 may identify target device 3046 in any technically feasible manner.

[0312] If the target device 3046 is CUDA compatible (e.g., a CUDA-enabled GPU 3094), the HIP compiler driver 3040 can generate HIP / NVCC compilation command 3042. HIP / NVCC compilation command 3042 can configure the CUDA compiler 3050 to compile HIP source code 3030 using a HIP-to-CUDA translation header and the CUDA runtime library. In response to HIP / NVCC compilation command 3042, the CUDA compiler 3050 can generate host executable code 3070(1) and CUDA device executable code 3084.

[0313] If the target device 3046 is incompatible with CUDA, the HIP compiler driver 3040 can generate HIP / HCC compilation command 3044. HIP / HCC compilation command 3044 can configure HCC 3060 to compile HIP source code 3030 using the HCC header and HIP / HCC runtime library. In response to HIP / HCC compilation command 3044, HCC 3060 can generate host executable code 3070(2) and HCC device executable code 3082. HCC device executable code 3082 can be a compiled version of the device code included in HIP source code 3030, which can be executed on GPU 3092. GPU 3092 can be any processor optimized for parallel instruction processing, incompatible with CUDA, and compatible with HCC. GPU 3092 can be developed by AMD Inc., located in Santa Clara, California. GPU 3092 can include GPU 3092 that does not support CUDA.

[0314] For illustrative purposes only, Figure 30Three different flows, which can be implemented in at least one embodiment, are described for compiling CUDA source code 3010 for execution on CPU 3090 and various devices. A direct CUDA stream can compile CUDA source code 3010 for execution on CPU 3090 and CUDA-enabled GPU 3094 without converting CUDA source code 3010 to HIP source code 3030. An indirect CUDA stream can convert CUDA source code 3010 to HIP source code 3030 and then compile HIP source code 3030 for execution on CPU 3090 and CUDA-enabled GPU 3094. A CUDA / HCC stream can convert CUDA source code 3010 to HIP source code 3030 and then compile HIP source code 3030 for execution on CPU 3090 and GPU 3092.

[0315] The achievable direct CUDA stream is represented by dashed lines and a series of bubbles labeled A1-A3. As shown in bubble A1, the CUDA compiler 3050 can receive CUDA source code 3010 and CUDA compilation command 3048, which configures the CUDA compiler 3050 to compile the CUDA source code 3010. The CUDA source code 3010 available for the direct CUDA stream can be written in a CUDA programming language based on a programming language other than C++ (e.g., C, Fortran, Python, Java, etc.). In response to CUDA compilation command 3048, the CUDA compiler 3050 can generate host executable code 3070(1) and CUDA device executable code 3084 (shown in bubble A2). As shown in bubble A3, the host executable code 3070(1) and CUDA device executable code 3084 can be executed on CPU 3090 and CUDA-enabled GPU 3094, respectively. The CUDA device executable code 3084 may include binary code. CUDA device executable code 3084 may include PTX code and may be further compiled at runtime into binary code for a specific target device.

[0316] The achievable indirect CUDA streams are represented by dashed lines and a series of bubbles labeled B1-B6. As shown in bubble B1, the CUDA-to-HIP conversion tool 3020 can receive CUDA source code 3010. As shown in bubble B2, the CUDA-to-HIP conversion tool 3020 can convert CUDA source code 3010 into HIP source code 3030. As shown in bubble B3, the HIP compiler driver 3040 can receive HIP source code 3030 and determine that the target device 3046 supports CUDA.

[0317] As indicated by bubble B4, the HIP compiler driver 3040 can generate HIP / NVCC compilation command 3042 and transfer HIP / NVCC compilation command 3042 and HIP source code 3030 to CUDA compiler 3050. HIP / NVCC compilation command 3042 can configure CUDA compiler 3050 to compile HIP source code 3030 using HIP to CUDA translation header and CUDA runtime library. HIP to CUDA translation header can convert 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. CUDA compiler 3050 can combine HIP to CUDA translation header and CUDA runtime library corresponding to CUDA runtime API 3002 to generate host executable code 3070(1) and CUDA device executable code 3084. In response to HIP / NVCC compilation command 3042, CUDA compiler 3050 can generate host executable code 3070(1) and CUDA device executable code 3084 (represented by bubble B5). As shown in bubble B6, host executable code 3070(1) and CUDA device executable code 3084 can be executed on CPU 3090 and CUDA-enabled GPU 3094, respectively. CUDA device executable code 3084 may include binary code. CUDA device executable code 3084 may include PTX code and can be further compiled at runtime into binary code for a specific target device.

[0318] The implementable CUDA / HCC streams are represented by solid lines and a series of bubbles labeled C1-C6. As shown in bubble C1, the CUDA-to-HIP conversion tool 3020 can receive CUDA source code 3010. As shown in bubble C2, the CUDA-to-HIP conversion tool 3020 can convert CUDA source code 3010 into HIP source code 3030. As shown in bubble C3, the HIP compiler driver 3040 can receive HIP source code 3030 and can determine that the target device 3046 does not support CUDA.

[0319] HIP compiler driver 3040 can generate HIP / HCC compilation command 3044 and transfer HIP / HCC compilation command 3044 and HIP source code 3030 to HCC 3060 (shown as bubble C4). HIP / HCC compilation command 3044 can configure HCC 3060 to compile HIP source code 3030 using HCC header files and HIP / HCC runtime library. HIP / HCC runtime library can correspond to HIP runtime API 3032. HCC header can include any number and type of HIP and HCC interoperability mechanisms. In response to HIP / HCC compilation command 3044, HCC 3060 can generate host executable code 3070(2) and HCC device executable code 3082 (shown as bubble C5). As shown in the bubble labeled C6, host executable code 3070(2) and HCC device executable code 3082 can be executed on CPU 3090 and GPU 3092, respectively.

[0320] After converting CUDA source code 3010 to HIP source code 3030, executable code for a CUDA-enabled GPU 3094 or GPU 3092 can be generated using HIP compiler driver 3040 without re-executing the CUDA to HIP conversion tool 3020. CUDA to HIP conversion tool 3020 can convert CUDA source code 3010 to HIP source code 3030 and then store the HIP source code 3030 in memory. HIP compiler driver 3040 can then configure HCC 3060 to generate host executable code 3070(2) and HCC device executable code 3082 based on the HIP source code 3030. In at least one embodiment, HIP compiler driver 3040 subsequently configures CUDA compiler 3050 to generate host executable code 3070(1) and CUDA device executable code 3084 based on the stored HIP source code 3030.

[0321] According to at least one embodiment, the example kernel may be derived from... Figure 30 The CUDA to HIP conversion tool 3020 is used for conversion. CUDA source code 3010 divides the overall problem that a given kernel is designed to solve into relatively coarse subproblems, which can be solved independently using thread blocks. Each thread block contains any number of threads. Each subproblem can be divided into relatively fine pieces, which can be solved collaboratively in parallel by threads within the thread block. Threads within a thread block can cooperate by sharing data via shared memory and synchronizing execution to coordinate memory access.

[0322] CUDA source code 3010 can organize thread blocks associated with a given kernel into one-dimensional, two-dimensional, or three-dimensional thread block grids. Each thread block contains any number of threads, and the grid contains any number of thread blocks.

[0323] A kernel can be a function defined in device code using the "__global__" declaration specifier. The dimensions of the raster of kernels and their associated streams that execute a given kernel call can be specified using the CUDA kernel startup syntax. The CUDA kernel startup syntax is specified as "KernelName <<<GridSize,BlockSize,SharedMemorySize,Stream> >>

[0324] (KernelArguments);". The execution configuration syntax can include a "<<<...>>>" structure between the kernel name ("KernelName") and the parenthesized list of kernel parameters ("KernelArguments"). The CUDA kernel boot syntax can include the CUDA boot function syntax, instead of the execution configuration syntax.

[0325] "GridSize" can be of type dim3 and specifies the dimensions and size of the grid. The dim3 type can be a CUDA-defined structure containing unsigned integers x, y, and z. If z is not specified, it defaults to 1. If y is not specified, it defaults to 1. The number of thread blocks in the grid can be equal to the product of GridSize.x, GridSize.y, and GridSize.z. "BlockSize" can be of type dim3 and specifies the dimensions and size of each thread block. The number of threads per thread block can be equal to the product of BlockSize.x, BlockSize.y, and BlockSize.z. Each thread executing the kernel can be assigned a unique thread ID, which can be accessed in the kernel via built-in variables such as "threadIdx".

[0326] Regarding the CUDA kernel startup syntax, "SharedMemorySize" is an optional argument that, in addition to statically allocated memory, can specify the number of bytes in shared memory dynamically allocated per thread block for a given kernel call. The default value for SharedMemorySize is zero. Regarding the CUDA kernel startup syntax, "Stream" is also an optional argument that specifies the associated stream; the default value is zero to specify the default stream. A stream can be a sequence of commands executed sequentially (which may be emitted by different host threads). Different streams can execute commands out of order or concurrently.

[0327] CUDA source code 3010 may include the kernel definition and main function of the example kernel "MatAdd". The main function may be host code that executes on the host machine and includes a kernel call that causes the MatAdd kernel to execute on the device. The MatAdd kernel can add two NxN matrices A and B, where N is a positive integer, and store the result in matrix C. The main function can define the threadsPerBlock variable as 16x16 and the numBlocks variable as N / 16 x N / 16. Then, the main function can specify the kernel call "MatAdd<<<numBlocks,threadsPerBlock> >>(A,B,C);”. According to the CUDA kernel startup syntax, the kernel MatAdd can be executed using a thread block grid of dimensions N / 16 x N / 16, where each thread block is 16x16. Each thread block can contain 256 threads, and a grid with enough thread blocks can be created so that there is one thread per matrix element, and each thread in such a grid can execute the kernel MatAdd to perform a pairwise addition.

[0328] When converting CUDA source code 3010 to HIP source code 3030, the CUDA to HIP conversion tool 3020 can convert each kernel call in CUDA source code 3010 from CUDA kernel startup syntax to HIP kernel startup syntax, and can convert any number of other CUDA calls in source code 3010 into any number of other functionally similar HIP calls. The HIP kernel startup syntax can be specified as "hipLaunchKernelGGL(KernelName,GridSize,BlockSize,SharedMemorySize,Stream,KernelArguments);". Each of KernelName, GridSize, BlockSize, ShareMemorySize, Stream, and KernelArguments has the same meaning in the HIP kernel startup syntax as it does in the CUDA kernel startup syntax (described earlier in this document). The arguments SharedMemorySize and Stream can be required in the HIP kernel startup syntax, but optional in the CUDA kernel startup syntax.

[0329] A portion of the HIP source code 3030 can be identical to a portion of the CUDA source code 3010 shown, except for the kernel call that causes the kernel MatAdd to execute on the device. The kernel MatAdd can be defined in the HIP source code 3030 using the same "__global__" declaration specifier as used in the CUDA source code 3010. The kernel call in the HIP source code 3030 can be "hipLaunchKernelGGL(MatAdd,numBlocks,threadsPerBlock,0,0,A,B,C);", while the corresponding kernel call in the CUDA source code 3010 is "MatAdd<<<numBlocks,threadsPerBlock> >>(A,B,C);

[0330] Other implementations are conceivable, and these can be implemented similarly to the CUDA and HIP implementations described above, such as oneAPI, OpenCL, and other programming platforms. Code can be converted in any direction. For example, CUDA can be converted to HIP, and CUDA can be converted to OpenCL. SnuCL-Tr and CUCL can be used to convert OpenCL to CUDA or CUDA to OpenCL, respectively. Compiled code or intermediate representations (such as CUDA PTX code) can also be converted to run on other processor platforms (such as AMD or Intel). For example, conversion tools (such as ZLUDA) can be used to convert PTX code to run on Intel or AMD processors. [033...

Claims

1. A processor, comprising: One or more circuits, said one or more circuits being configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures.

2. The processor according to claim 1, wherein, The one or more circuits are used to translate stack-based bytecode into instructions to be executed by the graphics processing unit (GPU) using the one or more instruction operand stack structure.

3. The processor according to claim 1, wherein, The one or more circuits are used to generate one or more instructions for using one or more registers by using a copied one or more instruction operand stack structure.

4. The processor according to claim 1, wherein, The one or more circuits are used to convert one or more first instructions using one or more operands from the one or more instruction operand stack structure into one or more second instructions using one or more operands stored in one or more registers.

5. The processor according to claim 1, wherein, The one or more circuits are used to generate a statement list based at least in part on the stack-based intermediate language bytecode using the one or more instruction operand stack structure.

6. The processor according to claim 1, wherein, The one or more instruction operand stack structures to be copied are part of a larger instruction operand stack structure.

7. The processor according to claim 1, wherein, The one or more circuits are used to generate one or more instructions corresponding to conditional branches of stack-based bytecode using a copied stack structure of one or more instruction operands.

8. A system comprising: One or more processors, wherein the one or more processors are configured to cause one or more instruction operand stack structures to be copied when one or more branch instructions use one or more instruction operands stored in one or more stack structures.

9. The system according to claim 8, wherein, The one or more processors are used to convert stack-based bytecode into one or more instructions to be executed by the accelerator using an instruction operand stack structure.

10. The system according to claim 8, wherein, The one or more processors are used to generate at least one instruction for using one or more registers by using a copied instruction operand stack structure.

11. The system according to claim 8, wherein, The one or more processors are configured to modify one or more first instructions that use one or more operands from the one or more instruction operand stack structure to use one or more second instructions that use one or more operands stored in one or more registers.

12. The system according to claim 8, wherein, The one or more processors are used to generate statement sequences, at least in part, based on the bytecode of a stack-based language, using the one or more instruction operand stack structure.

13. The system according to claim 8, wherein, The one or more instruction operand stack structures to be copied are less than the entirety of the larger instruction operand stack structure.

14. The system according to claim 8, wherein, The one or more processors are used to generate one or more instructions corresponding to branches of stack-based bytecode using a copied one or more instruction operand stack structure.

15. A method comprising: When one or more branch instructions use one or more instruction operands stored in one or more stack structures, the one or more instruction operand stack structures are copied.

16. The method of claim 15, further comprising: The copied one or more instruction operand stack structure is used to generate one or more instructions that use one or more registers.

17. The method of claim 15, further comprising: One or more first instructions using one or more operands from the one or more instruction operand stack structure are translated into one or more second instructions using one or more operands stored in one or more registers.

18. The method of claim 15, further comprising: The one or more instruction operand stack structure is used to generate a statement list at least in part based on the bytecode of a stack-based intermediate language.

19. The method according to claim 15, wherein, The one or more instruction operand stack structures to be copied are part of a larger instruction operand stack structure.

20. The method of claim 15, further comprising: The copied one or more instruction operand stack structure is used to generate one or more instructions corresponding to conditional branches of stack-based bytecode.

Images