Artificial intelligence chip, method for processing memory access related instructions, computing device, medium and program product

Abstract

The present application relates to an artificial intelligence chip, a method for processing memory access related instructions, a computing device, a computer readable storage medium and a computer program product. The artificial intelligence chip comprises a computing core, the computing core comprising: a computing unit, the computing unit comprising at least a load store cache module, the load store cache module comprising: a queuing module configured to process memory access related instructions for a predetermined domain segment without carrying a predetermined modification, the predetermined modification indicating whether the memory access related instruction is provided to a bypass module while bypassing the queuing module; and a bypass module configured to process memory access related instructions for a predetermined domain segment carrying a predetermined modification. The present application can significantly reduce the queuing and arbitration delay of memory access requests, and reduce the response time of memory access requests.

Description

TECHNICAL FIELD

[0001] Embodiments of the present application generally relate to the field of artificial intelligence, and more particularly, to an artificial intelligence chip, a method for processing memory access related instructions, a computing device, a computer readable storage medium, and a computer program product. BACKGROUND

[0002] In a conventional method for processing memory access related instructions, for example, tensor memory access instructions in an artificial intelligence chip (such as, but not limited to, a general-purpose graphics processor), such as "ldmatrix / stmatrix" instructions, are usually processed in order in a load store cache (LSC) module, and these instructions are uniformly entered into the same sorting, merging, or arbitration path. Taking a memory access request as an example, when it is queued in the LSC module, "head blocking" is prone to occur, thereby causing delay jitter and tail latency to rise. In addition, programmers cannot give higher priority to part of the "urgently needed return" memory access requests, so it is difficult to return critical data in time for the memory access requests that occur "head blocking". Taking the ldmatrix / stmatrix and other tensor memory access instructions as an example, they often have batch characteristics and high bandwidth occupation, so they are more likely to block subsequent small-granularity, low-latency-sensitive memory access requests. Therefore, in the conventional method for processing memory access related instructions, the strategy of processing in order and consistency in the LSC module lacks a low-overhead fast lane provided for scenarios that do not require strict order guarantee. Therefore, it leads to the queuing and arbitration delay of critical memory access requests, and increases the response time and tail latency of delay-sensitive memory access requests.

[0003] In summary, the conventional method for processing memory access related instructions has the disadvantage of easily leading to the queuing and arbitration delay of memory access requests, and increasing the response time of memory access requests. SUMMARY

[0004] The present application provides an artificial intelligence chip, a method for processing memory access related instructions, a computing device, a computer readable storage medium, and a computer program product, which can significantly reduce the queuing and arbitration delay of memory access requests, and reduce the response time of memory access requests.

[0005] According to a first aspect of the present application, there is provided an artificial intelligence chip. The artificial intelligence chip comprises: an artificial intelligence chip, characterized in that the artificial intelligence chip comprises a computing core, and the computing core comprises: a computing unit, the computing unit comprising: a load-store cache module, the load-store cache module comprising: a queuing module configured to process a memory access related instruction of a predetermined domain segment without carrying a predetermined modification, the predetermined modification indicating whether the memory access related instruction is provided to a bypass module without bypassing the queuing module; and a bypass module configured to process a memory access related instruction of a predetermined domain segment carrying a predetermined modification.

[0006] According to a second aspect of the present application, there is also provided a method for processing a memory access related instruction. The method comprises: configuring a predetermined domain segment in the memory access related instruction; determining whether the predetermined domain segment of a current memory access related instruction carries a predetermined modification, the predetermined modification indicating whether the memory access related instruction is provided to a bypass module without bypassing a queuing module; in response to determining that the predetermined domain segment of the current memory access related instruction carries the predetermined modification, causing the current memory access related instruction to be processed by the bypass module of the load-store cache module; and in response to determining that the predetermined domain segment of the current memory access related instruction does not carry the predetermined modification, causing the current memory access related instruction to be processed by the queuing module of the load-store cache module.

[0007] According to a third aspect of the present application, there is also provided a computing device. The computing device comprises: at least one processor; and a memory communicatively connected with the at least one processor; wherein the memory stores instructions executable by the at least one processor, and the instructions are executed by the at least one processor to enable the computing device to perform the method of the second aspect of the present application.

[0008] According to a fourth aspect of the present application, there is also provided a computer readable storage medium. The computer readable storage medium stores a computer program, and the computer program is executed by a machine to perform the method of the second aspect of the present application.

[0009] According to a fifth aspect of the present application, there is also provided a computer program product comprising a computer program, and the computer program is executed by a machine to perform the method of the second aspect of the present application.

[0010] In some embodiments, the load-store cache module further comprises: a judgment module located at an entrance of the load-store cache module and upstream of the queuing module and the bypass module, the judgment module being configured to determine whether the predetermined domain segment of the received current memory access related instruction carries the predetermined modification, and in response to determining that the predetermined domain segment of the current memory access related instruction carries the predetermined modification, providing the current memory access related instruction to the bypass module; and in response to determining that the predetermined domain segment of the current memory access related instruction does not carry the predetermined modification, providing the current memory access related instruction to the queuing module.

[0011] In some embodiments, the load-store buffer module comprises a front-end of the load-store buffer module, the queuing module, the bypass module and the determining module are comprised in the front-end of the load-store buffer module, the bypass module is configured to merge a plurality of thread memory access requests associated with the same address segment, and send the generated merged memory access request to a cache controller comprised in the load-store buffer module; the computing unit comprised in the computing core is a plurality of computing units, each of the plurality of computing units comprises the load-store buffer module.

[0012] In some embodiments, the bypass module is further configured to merge a plurality of thread memory access requests associated with the same address segment, so as to output the merged memory access request to the cache controller.

[0013] In some embodiments, the load-store buffer module further comprises: a cache controller configured to determine whether the merged memory access request output by the bypass module is provided to a static random access memory read-write module or an external system of the current computing unit; and a static random access memory read-write module configured to obtain returned mixed data based on the merged memory access request provided to the static random access memory read-write module, the mixed data comprising returned data corresponding to the plurality of thread memory access requests respectively.

[0014] In some embodiments, the load-store buffer module further comprises: a back-end of the load-store buffer module configured to remap the mixed data output by the static random access memory read-write module to threads corresponding to the plurality of thread memory access requests respectively.

[0015] In some embodiments, the memory access related instructions provided to the queuing module and the memory access related instructions provided to the bypass module are processed in parallel.

[0016] In some embodiments, causing the current memory access related instruction to be provided to the bypass module of the load-store buffer module for processing comprises: merging, via the bypass module comprised in the front-end of the load-store buffer module, a plurality of thread memory access requests associated with the same address segment, so as to output the merged memory access request to the cache controller; and the cache controller determining whether the merged memory access request is provided to the static random access memory read-write module or the external system of the current computing unit.

[0017] In some embodiments, the bypassing the current memory access related instruction to the load store buffer module for processing further comprises: in response to the cache controller determining that the merged memory access request is provided to the static random access memory read-write module, the static random access memory obtaining returned mixed data based on the merged memory access request, the mixed data comprising returned data corresponding to the plurality of thread memory access requests respectively; and remapping the mixed data to threads corresponding to the plurality of thread memory access requests respectively via a back end of the load store buffer module.

[0018] In some embodiments, the configuring the predetermined domain segment in the memory access related instruction comprises: configuring a flag bit in an extension modifier bitmap of the memory access related instruction for setting a predetermined modifier, the predetermined modifier being an additional instruction modifier.

[0019] In some embodiments, the determining whether the predetermined domain segment of the memory access related instruction carries the predetermined modifier comprises: decoding the memory access related instruction via a judging module of a front end of the load store buffer module to determine whether the predetermined domain segment of the memory access related instruction carries the predetermined modifier.

[0020] The present application can make the memory access related instruction carrying the predetermined modifier bypass the queuing constraint of the queuing module and enter the independent bypass module for processing, so as to significantly reduce the queuing and arbitration delay of the memory access request and reduce the response time of the memory access request.

[0021] It should be understood that the content described in this part is not intended to identify key or important features of the embodiments of the present application, nor to limit the scope of the present application. Other features of the present application will become apparent from the following description. BRIEF DESCRIPTION OF DRAWINGS

[0022] The above and other features, advantages and aspects of embodiments of the present application will become more apparent by describing in detail preferred embodiments thereof with reference to the attached drawings in which:

[0023] Figure 1 A schematic diagram of a computing device implementing a method for processing memory access related instructions according to an embodiment of the present application is schematically shown.

[0024] Figure 2A flowchart of a method for processing memory reference related instructions according to an embodiment of the present application is shown.

[0025] Figure 3 A schematic diagram of an artificial intelligence chip according to some embodiments of the present application is shown.

[0026] Figure 4 A schematic diagram of a load store buffer module and an external system according to some embodiments of the present application is shown.

[0027] Figure 5 A flowchart of a method of a bypass module processing memory reference related instructions carrying predetermined modifications for predetermined domain segments according to an embodiment of the present application is shown.

[0028] In the various drawings, like or corresponding elements are denoted by like or corresponding reference numerals. DETAILED DESCRIPTION

[0029] Preferred embodiments of the present application will be described herein below with reference to the accompanying drawings. While preferred embodiments of the present application are shown in the drawings, it is understood that the present application can be embodied in various forms and should not be limited by the embodiments set forth herein. Rather, these embodiments are provided so that this disclosure will be thorough and complete, and will fully convey the scope of the present application to those skilled in the art.

[0030] The term "comprising" and variations thereof as used herein are intended to cover a non-exclusive inclusion, such that a process, method, article, or apparatus that comprises a list of elements is not necessarily limited to those elements, but can include other elements not expressly listed or inherent to such process, method, article, or apparatus. Unless otherwise expressly stated, the term "or" refers to an inclusive "or" and not to an exclusive "or". The term "based on" means "based, at least in part, on". The terms "one example embodiment" and "an example embodiment" mean "at least one example embodiment". The term "another embodiment" means "at least one additional embodiment". The terms "a first", "a second", etc. do not require a strict ordering but rather indicate different or additional ones of the indicated objects.

[0031] As described previously, the conventional method of processing memory reference related instructions suffers from the drawbacks of causing queuing and arbitration latency of memory requests, and increasing the response time of memory requests.

[0032] To at least partially address one or more of the aforementioned problems and other potential issues, exemplary embodiments of the present invention provide an artificial intelligence chip. In this AI chip, by including a loading memory cache module within a computing unit of a computing core of the AI ​​chip, comprising a queuing module and a bypass module, and configuring the queuing module to process memory access instructions in a predetermined field without a predetermined modifier (the predetermined modifier indicating whether memory access instructions bypass the queuing module and are provided to the bypass module); and configuring the bypass module to process memory access instructions in the predetermined field with the predetermined modifier, the present invention enables memory access instructions in the predetermined field with the predetermined modifier to bypass the ordering constraints of the queuing module and enter an independent bypass module for processing. Therefore, it can significantly reduce the queuing and arbitration latency of memory access requests, and reduce the response time of memory access requests.

[0033] Figure 1 A schematic diagram illustrates a computing device 100 implementing a method for processing memory access-related instructions according to an embodiment of the present invention. Figure 1 As shown, the computing device 100 may have one or more processing units and a memory communicatively connected to one or more processors; wherein the memory stores instructions executable by at least one processor. The processing units include, for example, dedicated processing units such as Graphics Processing Units (GPUs), Field Programmable Gate Arrays (FPGAs), Application Specific Integrated Circuits (ASICs), General-purpose computing on graphics processing units (GPGPUs), and general-purpose processing units such as CPUs. The computing device 100 also includes at least: a predetermined domain segment configuration unit 102, a predetermined modification determination unit 104, a bypass module processing unit 106, and a queuing module processing unit 108.

[0034] Regarding the predefined domain segment configuration unit 102, it is used to configure a predefined domain segment in memory access related instructions.

[0035] Regarding the determination unit 104 for whether a predetermined modification is carried, it is used to determine whether a predetermined field of the current memory access related instruction carries a predetermined modification. The predetermined modification indicates whether the memory access related instruction bypasses the queuing module and is provided to the bypass module.

[0036] With regard to the bypass module processing unit 106, it is configured to cause the current memory access related instruction to be provided to the bypass module of the load store cache module for processing in response to determining that the predetermined domain segment of the current memory access related instruction carries the predetermined modification.

[0037] With regard to the queuing module processing unit 108, it is configured to cause the current memory access related instruction to be provided to the queuing module of the load store cache module for processing in response to determining that the predetermined domain segment of the current memory access related instruction does not carry the predetermined modification.

[0038] Embodiments of the present application will be described below with reference to Figure 2 , Figure 3 and Figure 4 a method 200 for processing a memory access related instruction. Figure 2 A flowchart of the method 200 for processing a memory access related instruction according to embodiments of the present application is shown. Figure 3 A schematic diagram of an artificial intelligence chip 300 according to some embodiments of the present application is shown. Figure 4 A schematic diagram of a load store cache module 332 and an external system 342 according to some embodiments of the present application is shown. It should be appreciated that the method 200 can be performed, for example, at the computing device 100 described. The method 200 can also include additional actions not shown and / or can omit actions shown, the scope of the present application not being limited in this respect. Figure 1

[0039] At step 202, the computing device 100 configures a predetermined domain segment in a memory access related instruction.

[0040] With regard to the method of configuring a predetermined domain segment in a memory access related instruction, it includes, for example, the computing device 100 configuring a flag bit in an extension modifier bitmap of the memory access related instruction for setting a predetermined modification, the predetermined modification indicating whether the memory access related instruction is provided to the bypass module without going through the queuing module. In some embodiments, the flag bit of the predetermined domain segment of the memory access related instruction for the predetermined modification is, for example, “1” or “true”, indicating that the predetermined modification is carried, and in turn indicating that the memory access related instruction is provided to the bypass module without going through the queuing module; the flag bit of the predetermined domain segment of the memory access related instruction for the predetermined modification is, for example, “0” or “false”, indicating that the predetermined modification is not carried, and in turn indicating that the memory access related instruction is provided to the queuing module.

[0041] With regard to the predetermined modification, it is, for example, an additional instruction modifier. The predetermined modification coexists with existing address space, alignment, shape, and other instruction modifiers, without changing the semantics of the other instruction modifiers. For example, a configured tensor access instruction is “ldmatrix.relax. ​The string "ldmatrix" indicates the tensor loading instruction, and ".relax" indicates predefined flags. "Indicates other instruction modifiers. In some embodiments, "relax" is a 1-bit flag.

[0042] At step 204, the computing device 100 determines whether a predetermined segment of the current memory access instruction carries a predetermined modifier, the predetermined modifier indicating whether the memory access instruction bypasses the queuing module and is provided to the bypass module.

[0043] For example, in some embodiments, the computing device 100 decodes memory access-related instructions via a determination module at the front end of the loaded memory cache module to determine whether a predetermined field of the memory access-related instruction carries a predetermined modification. For example, the determination module is configured at the entry point of the loaded memory cache module.

[0044] like Figure 4 As shown, the front end 410 of the load storage cache module 332 receives a memory access request. The judgment module 412 included in the front end 410 decodes the received access request to determine whether a predetermined field of the memory access-related instruction corresponding to the memory access request carries a predetermined modifier. For example, if the predetermined field of the memory access-related instruction carries a predetermined modifier (e.g., the relax flag indicates "1" or "true"), the memory access-related instruction is instructed to bypass the queuing module and be provided to the bypass module; if it does not carry a predetermined modifier (e.g., the relax flag indicates "0" or "false"), the memory access-related instruction is instructed to be provided to the queuing module.

[0045] In step 206, if the computing device 100 determines that a predetermined field of the current memory access instruction carries a predetermined modification, the current memory access instruction is provided to the bypass module of the loading memory cache module for processing.

[0046] In some embodiments, if the judgment module included in the front end of the loading storage cache module determines that a predetermined field of the current memory access related instruction of the judgment module carries a predetermined modification, the current memory access related instruction is provided to the bypass module, and the bypass module processes memory access related instructions whose predetermined field does not carry the predetermined modification. For example... Figure 4 As shown, if the judgment module 412 of the front end 410 of the loading storage cache module determines that the predetermined field of the current memory access related instruction carries a predetermined modification (for example, the relax flag indicates "1" or "true"), the current memory access request is provided to the bypass module 414 (as indicated by label 404) so ​​that the bypass module 414 can process the current memory access request.

[0047] At step 208, if the computing device 100 determines that the predetermined domain segment of the current memory access related instruction does not carry the predetermined modification, the current memory access related instruction is provided to the queuing module of the load store cache module for processing.

[0048] In some embodiments, if the judging module included in the front end of the load store cache module determines that the predetermined domain segment of the current memory access related instruction does not carry the predetermined modification, the queuing module processes memory access related instructions that carry the predetermined modification for the predetermined domain segment. As Figure 4 As shown, if the judging module 412 of the front end 410 of the load store cache module determines that the predetermined domain segment of the current memory access related instruction does not carry the predetermined modification (e.g., the relax flag indicates "0" or "false"), the current memory request is provided to the queuing module 416 (as indicated by the marker 406) for processing by the queuing module 416 for the current memory request.

[0049] As to the method of processing, by the bypass module, memory access related instructions (or memory requests) that carry the predetermined modification for the predetermined domain segment, it comprises, for example: the computing device 100 merges, via the bypass module included in the front end of the load store cache module, memory requests of multiple threads that access the same address segment, so as to output the merged memory request to the cache controller; the cache controller determines whether the merged memory request is provided to the static random access memory read-write module or to the external system of the current computing unit; in response to the cache controller determining that the merged memory request is provided to the static random access memory read-write module, the static random access memory read-write module obtains the returned mixed data based on the merged memory request, the mixed data comprising returned data corresponding to the multiple thread memory requests respectively; and via the back end of the load store cache module, the mixed data is remapped to threads corresponding to the multiple thread memory requests respectively. The method will be described in detail below in conjunction with Figure 5 The method 500 of processing, by the bypass module, memory access related instructions that carry the predetermined modification for the predetermined domain segment will not be described in detail here.

[0050] The following takes the application scenario of matrix multiplication (referred to as “MMA operation” for short) as an example to further illustrate the method 200. It should be understood that the MMA operation is divided into two parts: data carrying and matrix multiplication. The calculation task related to matrix multiplication is divided into multiple sub-tasks, each thread bundle (warp) in the thread bundle group executes different sub-tasks, so that multiple warps asynchronously and concurrently execute sub-tasks related to matrix multiplication, and then cooperatively complete the calculation task of MMA operation. For example, each warp that executes different sub-tasks is referred to as a warp specialized, the thread bundle or thread bundle group (warp / warp group) that carries data between the global memory (GLM) and the group shared memory (GSM) is referred to as a producer warp or warp group, and the warp group that performs matrix multiplication operation is referred to as a consumer warp group. Taking the warps (warp0 to warp11) that execute the MMA operation as an example, the zeroth thread bundle (warp0) is used to carry the input data A, which, for example, carries the input data A from the GLM to the GSM. The first thread bundle (warp1) is used to carry the input data B, which, for example, carries the input data B from the GLM to the GSM. The second thread bundle (warp2) is used to store the output result data, for example, the output data D (for example, the accumulated result), and the warp2, for example, carries the output data D in the GSM to the GLM. The fourth thread bundle to the seventh thread bundle (warp4-warp7) are, for example, scheduled together as a consumer warp group, which, for example, includes reading the input data A or the input data B in the GSM, performing multiplication and addition calculation, and outputting the accumulated result to the GSM. The eighth thread bundle to the eleventh thread bundle (warp8-warp11) are, for example, scheduled together as another consumer warp group, which, for example, includes reading the input data A or the input data B in the GSM, performing multiplication and addition calculation, and outputting the accumulated result to the GSM. It should be understood that reading the input data A in the GSM can use a tensor memory instruction (specifically, for example, a tensor load instruction “ldmatrix”) to carry data from the GSM to the thread register (TLR) for subsequent multiplication and addition calculation. After the calculation is completed, a tensor memory instruction (specifically, for example, a tensor storage instruction “stmatrix”) can be used to carry the output data from the TLR to the GSM. It should be understood that, for example, two consumer warp groups consume (or calculate) different data blocks respectively, when one consumer warp group is executing the stmatrix instruction, the other consumer warp group is executing the ldmatrix instruction, and the data operated by the two consumer warp groups has no dependency.The bypass module 414 can be configured to bypass the queuing module 416 for the memory access related instruction with the predetermined domain segment carrying the predetermined modifier. The bypass module 414 can be configured to process the memory access related instruction with the predetermined domain segment carrying the predetermined modifier. The two consumer thread bundle groups can be configured to perform the instruction execution in parallel in the load store cache module 332 without affecting each other via the bypass module 414. Thus, the efficiency of returning data of the memory access related instruction can be improved, and the response time of the memory access request can be reduced.

[0051] The following detailed description of the artificial intelligence chip 300 and the load store cache module 332 is made with reference to the accompanying drawings. Figure 3 and Figure 4 The artificial intelligence chip 300 and the load store cache module 332 are described in detail.

[0052] As shown in FIG. 3, the artificial intelligence chip 300 at least includes a compute core 320. In some embodiments, the compute core 320 is a stream processor cluster (SPC) which, for example, includes a plurality of compute execution units, and the SPC implements high-performance computing by task management and resource scheduling of the plurality of compute execution units. Figure 3 Regarding the compute core 320, it further includes one or more compute units.

[0053] A plurality of compute units (e.g., compute unit 330-1 to compute unit 330-N, N being a natural number) are schematically shown in FIG. 3. Figure 3

[0054] Regarding the compute unit, it is configured to perform computation. The compute unit at least includes a load store cache module 332. In some embodiments, each of the plurality of compute units (e.g., compute unit 330-1 to compute unit 330-N) includes a corresponding load store cache module. For example, the compute unit 330-1 to the compute unit 330-N respectively include the load store cache module 332-1 to the load store cache module 332-N.

[0055] As shown in FIG. 3, the load store cache module 332 at least includes a queuing module 416 and a bypass module 414. Figure 4 Regarding the queuing module 416, it is configured to, for example, process the memory access related instruction with the predetermined domain segment not carrying the predetermined modifier, the predetermined modifier indicating whether the memory access related instruction is provided to the bypass module 414 without bypassing the queuing module 416.

[0056] Regarding the bypass module 414, it is configured to, for example, process the memory access related instruction with the predetermined domain segment carrying the predetermined modifier.

[0057]

[0058] ​​It should be understood that, by making the load-store buffer module within the compute unit of the artificial intelligence chip include a queuing module and a bypass module, and making the queuing module be configured to process the memory access related instruction for the predetermined domain segment without carrying the predetermined modification indicating whether the memory access related instruction is provided to the bypass module bypassing the queuing module, and making the bypass module be configured to process the memory access related instruction for the predetermined domain segment carrying the predetermined modification, the present application can make the memory access related instruction for the predetermined domain segment carrying the predetermined modification bypass the queuing constraint of the queuing module and enter the independent bypass module for processing, so that the queuing and arbitration delay of the memory request can be significantly reduced, and the response time of the memory request can be reduced.

[0059] In some embodiments, as shown in FIG. 3, the compute core 320 further includes, for example, a cluster bus 340. The artificial intelligence chip 300 further includes, for example, a level two cache 350 and a high bandwidth memory 360. Figure 3

[0060] As shown in FIG. 4, in some embodiments, the load-store buffer module 332 further includes, for example, a judging module 412. The judging module 412 is located at the entrance of the load-store buffer module 332 and is upstream of the queuing module 416 and the bypass module 414. The judging module 412 is configured to determine whether the predetermined domain segment of the received current memory access related instruction carries the predetermined modification, and in response to determining that the predetermined domain segment of the current memory access related instruction carries the predetermined modification, provide the current memory access related instruction to the bypass module 414, and in response to determining that the predetermined domain segment of the current memory access related instruction does not carry the predetermined modification, provide the current memory access related instruction to the queuing module 416. Figure 4

[0061] In some embodiments, the load-store buffer module 332 further includes, for example, a front end 410 of the load-store buffer module. For example, the queuing module 416, the bypass module 414 and the judging module 412 are included in the front end 410 of the load-store buffer module.

[0062] As for the bypass module 414, it is further configured to, for example, merge the memory access requests of multiple threads accessing the same address segment and send the generated merged memory access request to a cache controller 420 included in the load-store buffer module 332. It should be understood that, in some embodiments, the memory access related instructions provided to the queuing module 416 and the memory access related instructions provided to the bypass module 414 are processed in parallel. Thus, the present application can further reduce the delay of the memory request.

[0063] In some embodiments, the load-store buffer module 332 further includes, for example, a cache controller 420, a static random access memory read-write module 430, and a back end 440 of the load-store buffer module. ​​

[0064] As to the cache controller 420, it is configured to determine whether the merged memory access request output by the bypass module 414 is provided to the static random-access memory read-write module 430 or the external system 342 of the current computing unit, for example.

[0065] As to the external system 342, it is a system outside the current computing unit to which the current load-store cache module 332 belongs, for example. In some embodiments, as shown in Figure 4 , the external system 342 includes a compute core bus interface (or referred to as “CBI”, such as the CBI 344 in Figure 4 , for example) and an external memory 352. In some embodiments, the CBI 344 is a cluster bus included inside the compute core 320, as indicated by the label 340 in Figure 3 , for example. The external memory 352 includes a level two cache, as indicated by the label 350 in Figure 3 , for example, and a high bandwidth memory, as indicated by the label 360 in Figure 3 , for example.

[0066] In some embodiments, as shown in Figure 4 , the load-store cache module 332 further includes a static random-access memory read-write module 430 and a backend 440 of the load-store cache module.

[0067] As to the static random-access memory (SRAM) read-write module, it is referred to as the SRAM read-write module for short. The static random-access memory read-write module 430 is configured to obtain returned mixed data based on the merged memory access request provided to the static random-access memory read-write module 430, the mixed data including returned data corresponding to the plurality of thread memory access requests respectively.

[0068] As to the backend 440 of the load-store cache module, it is configured to remap the mixed data output by the static random-access memory read-write module 430 to threads corresponding to the plurality of thread memory access requests respectively, for example.

[0069] In the above scheme, by determining the different types of the merged memory access request output by the bypass module 414 through the cache controller 420, and then providing the merged memory access request of the corresponding type to the static random-access memory read-write module 430 and the external system 342 of the current computing unit respectively, the present application can further improve the parallel processing efficiency of the plurality of thread memory access requests processed via the bypass module, thereby facilitating further reducing the response time of the memory access request.

[0070] It should be appreciated that the scheduling order of the memory access related instructions carrying the predetermined modification for the predetermined domain segment is changed by the bypass module, and the cache coherence protocol is not changed. Therefore, the memory access related instructions carrying the predetermined modification for the predetermined domain segment can still access the existing cache / coherence hardware access level (for example, the computing core 320, the second-level cache 350, and / or the high-bandwidth memory 360) of the artificial intelligence chip 300. Thus, the bypass module can bypass the memory request queue and arbitration under the premise of not changing the cache / coherence protocol, thereby reducing the response time of the memory request. In addition, the bypass module has the advantages of small modification to the existing microarchitecture of the artificial intelligence chip, low risk, easy deployment, and good compatibility.

[0071] The bypass module of the embodiment of the present application will be described below in conjunction with Figure 3 、 Figure 4 to Figure 5 The bypass module of the embodiment of the present application will be described below in conjunction with Figure 5 A flowchart of a method 500 for processing the memory access related instructions carrying the predetermined modification for the predetermined domain segment by the bypass module according to the embodiment of the present application is shown. It should be appreciated that the method 500 can be executed at the computing device 100 described above, for example. The method 500 can also include additional actions not shown and / or can omit actions shown, and the scope of the present application is not limited in this respect. Figure 1 The method 500 can also include additional actions not shown and / or can omit actions shown, and the scope of the present application is not limited in this respect.

[0072] At step 502, the bypass module included in the front end of the load store cache module merges the memory access requests associated with the threads within the same address segment, so as to output the merged memory access request to the cache controller.

[0073] As shown in Figure 4 , the front end 410 (or “LSC FrontEnd”) of the load store cache module receives the memory access request (as indicated by the mark 402 in Figure 4 ). The memory access request includes a plurality of thread memory access requests, for example. It should be appreciated that the 32 threads in the same thread bundle can access different addresses. The front end 410 of the load store cache module merges the memory access requests associated with the threads within the 128-byte range of the same address segment, so as to send the generated merged memory access request to the cache controller.

[0074] For example, the bypass module 414 included in the front end 410 of the load store cache module receives the memory access request provided by the judgment module 412 (as indicated by the mark 404 in Figure 4 ), merges the memory access requests associated with the threads within the same address segment, so as to generate a merged memory access request; then, the bypass module 414 outputs the merged memory access request to the cache controller 420 (as shown inFigure 4 (As indicated by mark 408 in the text).

[0075] At step 504, the cache controller of the computing device 100 determines whether the merged memory access request is provided to the static random access memory read / write module or to an external system of the current computing unit.

[0076] It should be understood that merged memory access requests have different types, and different types of merged memory access requests can be routed to different paths.

[0077] For example, if the cache controller determines that the current merged memory access request is a Level 1 cache access request (e.g., an L1 hit request) or a Group Shared Memory (GSM) access request, it provides the current merged memory access request to the Static Random Access Memory (SRAM) read / write module. If the cache controller determines that the current merged memory access request is an access request other than a Level 1 cache access request (e.g., an L1 hit request) or a GSM access request, it provides the current merged memory access request to external system 342.

[0078] like Figure 4 As shown, if the cache controller 420 determines that the current merged memory access request is an L1 hit request or a GSM memory access request, it provides the current merged memory access request to the static random access memory read / write module 430 (e.g., ...). Figure 4 (As indicated by marker 422 in the diagram). If the cache controller 420 determines that the current merged memory access request output by the bypass module 414 is an access request other than a L1 cache access request and a GSM access request (e.g., an L2 cache access request or an HBM access request), then the current merged memory access request is provided to the CBI 344 of the external system 342 providing the current computing unit (e.g., ...). Figure 4 (as indicated by mark 424 in the text).

[0079] In some embodiments, the CBI 344 of the external system 342 feeds back the processing result (e.g., returned data) of the received merged memory access request to the static random access memory read / write module 430 (e.g., ...). Figure 4 (As indicated by mark 426 in the document). For example, the CBI 344 of external system 342 further provides the received merged memory access request to external memory 352 (e.g., Figure 4 (as indicated by marker 428 in the diagram), thereby obtaining the return data of the merged memory access request; then, CBI 344 feeds back the obtained return data of the merged memory access request to the static random access memory read / write module 430 (e.g., Figure 4 (As indicated by mark 426 in the text). It should be understood that the external memory 352 is, for example, a... Figure 3The secondary cache is 350 or the high-bandwidth memory is 360.

[0080] At step 506, if the cache controller of the computing device 100 determines that a merged memory access request is provided to the static random access memory read / write module, the static random access memory obtains the returned mixed data based on the merged memory access request, the mixed data including return data corresponding to the multiple thread memory access requests respectively.

[0081] like Figure 4 As shown, if the cache controller 420 determines that the current merged memory access request is a Level 1 cache access request (e.g., an L1 hit request), the static random access memory read / write module 430, based on the received merged memory access request (e.g., a Level 1 cache access request), obtains mixed data including corresponding return data from multiple threads, and provides the obtained mixed data to the backend 440 of the loading cache module (e.g., ...). Figure 4 (As indicated by mark 432 in the text).

[0082] At step 508, the mixed data is remapped to the threads corresponding to the multiple thread memory access requests via the backend of the load storage cache module of the computing device 100.

[0083] For example, the backend 440 (or "LSC BackEnd") of the loading storage cache module remaps the mixed data returned by the static random access memory read / write module 430, which corresponds to the merged memory access request and includes corresponding data from multiple threads, to the threads corresponding to the multiple thread memory access requests respectively, thereby providing response output (such as...). Figure 4 (As indicated by mark 442 in the text).

[0084] In the above scheme, the cache controller determines the different types of merged memory access requests output by the bypass module, and then provides the corresponding types of merged memory access requests to the static random access memory read / write module and the external system of the current computing unit. Thus, the present invention can further improve the parallel processing efficiency of multiple thread memory access requests processed by the bypass module, thereby helping to further reduce the response time of memory access requests.

[0085] The various processes and processes described above, such as the methods 200, 500, can be performed at a computing device. The computing device includes, for example, at least one processor (at least one graphics processor and at least one central processor) and a memory communicatively connected with the at least one processor, wherein the memory stores instructions executable by the at least one processor, which are executed by the at least one processor. In some embodiments, the methods 200, 500 can be implemented as a computer software program or program product tangibly embodied in a machine-readable medium. In some embodiments, part or all of the computer program can be loaded and / or installed onto the computing device via a Read-Only Memory (ROM) and / or a communication unit. When the computer program is loaded into a Random-Access Memory (RAM) and executed by the GPU and CPU, one or more actions of the methods 200, 500 described above can be performed.

[0086] The present application can be a method, an apparatus, a system, and / or a computer program product. The computer program product can include a computer readable storage medium (or media) having computer readable program instructions thereon for performing various aspects of the present application. The computer readable storage medium can be a tangible device that can retain and store instructions for use by an instruction execution device. The computer readable storage medium can be, for example, but is not limited to, an electronic storage device, a magnetic storage device, an optical storage device, an electromagnetic storage device, a semiconductor storage device, or any suitable combination of the foregoing.

[0087] Computer readable program instructions described herein can be downloaded to respective computing / processing devices from a computer readable storage medium or to an external computer or external storage device via a network, for example, the Internet, a local area network, a wide area network and / or a wireless network. The various aspects of the application can be described herein with reference to flowchart illustrations and / or block diagrams of methods, apparatus (systems) and computer program products according to embodiments of the application. It will be understood that each block of the flowchart illustrations and / or block diagrams, and combinations of blocks in the flowchart illustrations and / or block diagrams, can be implemented by computer readable program instructions.

[0088] These computer readable program instructions can be provided to a processor of a general purpose computer, special purpose computer, or other programmable data processing apparatus to produce a machine, such that the instructions, which execute via the processor of the computer or other programmable data processing apparatus, create means for implementing the functions / acts specified in the flowchart and / or block diagram block or blocks. These computer readable program instructions can also be stored in a computer readable storage medium that can include a non-transitory computer readable storage medium that can be a computer- readable storage medium having no data, programs, program modules, e.g., instructions for operation, or digital content stored thereon or therein for a short time or not at all. The computer readable storage medium can also have other meanings inhered thereby, which will be apparent to those reading the present disclosure. The instructions stored in the computer readable storage medium, when executed by a computer or other programmable data processing apparatus, produce a machine such that the instructions, which execute via the processor of the computer or other programmable data processing apparatus, create means for implementing the functions / acts specified in the flowchart and / or block diagram block or blocks.

[0089] The flow diagrams and the block diagrams in the drawings are presented to illustrate the architecture, functionality, and operation of possible implementations of systems, methods and computer program products according to the present application. In this regard, each block in the flow diagrams and the block diagrams can represent a module, segment, or portion of code, which comprises one or more executable instructions for implementing the specified logical functions ("instructions"). In some alternative implementations, the functions noted in the blocks can occur out of the order noted in the figures. For example, two blocks shown in succession may, in fact, be executed substantially concurrently, or the blocks can sometimes be executed in the reverse order, depending upon the functionality involved. It will also be noted that each block of the block diagrams and / or flow diagrams, and combinations thereof, can be implemented by special purpose hardware-based systems that perform the specified functions or acts, or combinations of special purpose hardware and

[0090] It should be understood that various forms of flow shown above can be re-ordered, added to, or deleted from without departing from the spirit of the present disclosure. For example, the steps loaded in the present application can be executed in parallel, in sequence, or in different orders, as long as the desired results of the technical solutions disclosed in the present application can be achieved, which are not limited herein.

[0091] The specific embodiments have been shown and described for purposes of illustrating the embodiments, and not for purposes of limitation. It will be clear to those skilled in the art that various modifications, combinations, sub-combinations and alternatives can be made to the specific embodiments without departing from the spirit or scope of the disclosure.

Claims

1. An artificial intelligence chip, characterized in that, The artificial intelligence chip includes a computing core, which includes: The computing unit includes at least a load-store-cache module, which includes: A queuing module is configured to process memory access-related instructions in a predetermined domain that do not carry a predetermined modifier, wherein the predetermined modifier indicates whether the memory access-related instructions bypass the queuing module and are provided to a bypass module; and The bypass module is configured to process memory access-related instructions with predetermined modifications for a predetermined domain segment.

2. The artificial intelligence chip according to claim 1, characterized in that, The loading and storage cache module also includes: The determination module, located at the entry point of the loading storage cache module and upstream of the queuing module and the bypass module, is configured to: determine whether a predetermined field of the received current memory access related instruction carries a predetermined modification; and in response to determining that the predetermined field of the current memory access related instruction carries the predetermined modification, provide the current memory access related instruction to the bypass module; and in response to determining that the predetermined field of the current memory access related instruction does not carry the predetermined modification, provide the current memory access related instruction to the queuing module.

3. The artificial intelligence chip according to claim 2, characterized in that, The loading and storing cache module includes a front-end, and the queuing module, bypass module and judgment module are included in the front-end of the loading and storing cache module. The bypass module is configured to merge multiple thread memory access requests within the same address range and send the generated merged memory access request to the cache controller included in the loading and storing cache module. The computing core includes multiple computing units, and each of the multiple computing units includes the loading and storage cache module.

4. The artificial intelligence chip according to claim 1, characterized in that, The bypass module is also configured to merge multiple thread memory access requests within the same address range so that the merged memory access request is output to the cache controller.

5. The artificial intelligence chip according to claim 4, characterized in that, The loading and storage cache module also includes: The cache controller is configured to determine whether merged memory access requests output by the bypass module are provided to the static random access memory read / write module or to an external system of the current computing unit; and The static random access memory read / write module is configured to obtain returned mixed data based on the merged memory access request provided to the static random access memory read / write module, wherein the mixed data includes return data corresponding to the memory access requests of multiple threads respectively.

6. The artificial intelligence chip according to claim 5, characterized in that, The loading and storage cache module also includes: The backend of the loading storage cache module is configured to remap the mixed data output by the static random access memory read / write module to the threads corresponding to the multiple thread memory access requests.

7. The artificial intelligence chip according to claim 1, characterized in that, The memory access instructions provided to the queuing module and the memory access instructions provided to the bypass module are processed in parallel.

8. A method for processing memory access related instructions, characterized in that, The method includes: Configure the predefined domain segment in memory access related instructions; Determine whether a predetermined field of the current memory access instruction carries a predetermined modifier, wherein the predetermined modifier indicates whether the memory access instruction bypasses the queuing module and is provided to the bypass module; In response to determining that a predetermined field of the current memory access instruction carries a predetermined modification, the current memory access instruction is provided to a bypass module that loads the memory cache module for processing; and In response to the determination that the predetermined field of the current memory access instruction does not carry the predetermined modifier, the current memory access instruction is provided to the queuing module of the loading memory cache module for processing.

9. The method according to claim 8, characterized in that, This allows the current memory access-related instructions to be provided to the bypass module that loads the memory cache module for processing, including: By using the bypass module included in the front end of the loading storage cache module, multiple thread memory access requests within the same address range are merged so that the merged memory access request is output to the cache controller. as well as The cache controller determines whether the merged memory access request is provided to the static random access memory read / write module or to an external system of the current computing unit.

10. The method according to claim 8, characterized in that, This also includes enabling the current memory access-related instructions to be provided to the bypass module that loads the memory cache module for processing: In response to the cache controller determining that a merged memory access request is provided to the static random access memory read / write module, the static random access memory read / write module obtains the returned mixed data based on the merged memory access request, the mixed data including the returned data corresponding to the memory access requests of multiple threads respectively; as well as By loading the storage cache module's backend, the mixed data is remapped to the threads corresponding to the multiple thread memory access requests.

11. The method according to claim 8, characterized in that, Configuring predefined domains in memory access-related instructions includes: Configure a flag bit in the extended modifier bitmap of memory access related instructions to set a predetermined modifier, wherein the predetermined modifier is an additional instruction modifier.

12. The method according to claim 8, characterized in that, Determining whether a predetermined field of a memory access instruction carries a predetermined modifier includes: The front-end judgment module of the loading storage cache module decodes the memory access related instructions to determine whether the predetermined field of the memory access related instructions carries the predetermined modification.

13. A computing device, characterized in that, include: At least one processor; as well as A memory that is communicatively connected to the at least one processor; in The memory stores instructions executable by the at least one processor, which, when executed by the at least one processor, enables the at least one processor to perform the method of any one of claims 8-12.

14. A computer-readable storage medium, characterized in that, The computer-readable storage medium stores a computer program that, when executed by a machine, performs the method according to any one of claims 8-12.

15. A computer program product, characterized in that, Includes a computer program, which, when executed by a machine, performs the method according to any one of claims 8-12.

Images