Instruction execution method and device, equipment and storage medium

Abstract

The embodiment of the invention discloses an instruction execution method and device, equipment and a storage medium, and the method comprises the steps: carrying out the parallel processing of one or more paths of data through a working group; wherein the working group comprises a plurality of threads, multiple paths of data are mutually independent, and the threads in one path of data have correlation; the thread corresponding to the path of data is dormant through the corresponding control parameter, and / or the thread in the dormant state is synchronously awakened. According to the instruction execution method, the graphics processor can more efficiently process the multi-thread synchronization problem.

Description

Technical Field

[0001] The present application relates to the field of computer technology, and in particular to an instruction execution method, device, equipment and storage medium. Background Art

[0002] In multithreaded programming, barrier instructions are widely used to synchronize the execution of threads. Barrier instructions can be understood as a synchronization point that prevents threads from continuing to execute until all threads participating in the barrier reach this point. In traditional barrier implementations, each thread will be blocked when encountering a barrier until the slowest thread arrives. This implementation method leads to a waste of threads because other threads may have completed their respective work but still need to wait for the slowest thread to complete. Summary of the invention

[0003] In view of this, embodiments of the present application provide at least one instruction execution method, apparatus, device, and storage medium.

[0004] The technical solution of the embodiment of the present application is implemented as follows: On the one hand, an embodiment of the present application provides an instruction execution method, which includes: processing one or more data in parallel through a work group; wherein the work group includes multiple threads, the multiple data are independent of each other, and there is a correlation between the threads within one data channel; the thread corresponding to the one data channel is put into sleep through corresponding control parameters, and / or the synchronous awakening of the thread in sleep state.

[0005] In some embodiments, the control parameter includes an arrival count value corresponding to the one-way data; the arrival count value is used to represent the number of threads that have not yet executed to the waiting point among the at least two threads corresponding to the one-way data.

[0006] In some embodiments, the instruction execution method further includes: in response to the thread corresponding to the one-way data executing to the waiting point, updating the arrival count value corresponding to the one-way data.

[0007] In some embodiments, the instruction execution method further includes: in response to the arrival count value indicating that there is no thread among the at least two threads that has not yet executed to the waiting point, waking up a thread in a dormant state among the at least two threads.

[0008] In some embodiments, the control parameter includes a sleep indication parameter; the sleep indication parameter is used to indicate a thread in a sleep state among the at least two threads corresponding to the one channel of data.

[0009] In some embodiments, the instruction execution method further includes: in response to the thread corresponding to the one channel of data entering a sleep state, updating the sleep indication parameter.

[0010] In some embodiments, the instruction execution method further includes: in response to reaching a wake-up condition, based on the sleep indication parameter, waking up the thread in a sleep state corresponding to the data path; the wake-up condition includes that there is no thread among the at least two threads that has not yet executed to the waiting point.

[0011] In some embodiments, the control parameters include an arrival count value and a sleep control parameter; the instruction execution method includes: updating the arrival count value after one or more threads corresponding to the data have executed the arrival instruction; when the arrival count value indicates that there is no thread that has not yet executed the arrival instruction among the at least two threads corresponding to the data, waking up the thread in a sleep state based on the sleep control parameter.

[0012] In some embodiments, the instruction execution method also includes: for a first thread that executes a reach instruction, in response to the first thread executing a wait instruction, controlling the first thread to enter a sleep state and updating the sleep indication parameter; the first thread is a thread that needs to use shared memory data among the at least two threads corresponding to the one path of data.

[0013] In some embodiments, the instruction execution method further includes: for a second thread that executes to a reach instruction and does not need to use shared memory data, after the second thread executes to the reach instruction, continuing to execute subsequent instructions.

[0014] In some embodiments, the instruction execution method further includes: executing an initialization instruction corresponding to the data path through any thread in the workgroup; the initialization instruction is used to set the initial value of the arrival count value corresponding to the data path to the total number of threads corresponding to the data path.

[0015] In some embodiments, the instruction execution method further includes: in response to all threads in the work group executing the initialization barrier instruction, controlling each thread in the work group to continue executing subsequent instructions.

[0016] In some embodiments, after one or more threads corresponding to the one path of data execute the arrival instruction, updating the arrival count value includes: in response to each thread corresponding to the one path of data executing the arrival instruction, reducing the arrival count value corresponding to the one path of data by one.

[0017] In some embodiments, updating the sleep indication parameter includes: based on the thread identifier of the first thread, updating the sleep indication bit corresponding to the thread identifier of the first thread in the sleep indication parameter corresponding to the one channel of data from a first value to a second value.

[0018] In some embodiments, waking up a thread in a dormant state based on a sleep control parameter includes: determining a thread identifier of the thread in a dormant state based on a sleep indication bit in the sleep control parameter that is at the second value; and waking up the thread in a dormant state through the thread identifier of the thread in a dormant state.

[0019] In some embodiments, the instruction execution method further includes: in response to a thread in a dormant state being awakened, based on the thread identifier of the thread, updating the dormant indication bit corresponding to the thread identifier of the thread in the dormant indication parameter corresponding to the data path from the second value to the first value.

[0020] In some embodiments, the instruction execution method is applied to a graphics processor, and an arrival count cache and a sleep control cache are provided in the on-chip storage of the graphics processor; wherein the arrival count cache is used to store the arrival count value corresponding to each data path in the work group; and the sleep control cache is used to store the sleep control parameters corresponding to each data path in the work group.

[0021] In some embodiments, the instruction execution method also includes: in response to updating the arrival count value corresponding to the one path of data, controlling the first thread to execute an independent instruction; the independent instruction is an instruction before the waiting instruction and does not require the use of shared memory data; the first thread is a thread that needs to use shared memory data among the at least two threads corresponding to the one path of data.

[0022] On the other hand, an embodiment of the present application provides an instruction execution device, which includes: an instruction parallel execution module, which is used to parallelly process one or more data through a work group; wherein the work group includes multiple threads, the multiple data are independent of each other, and there is a correlation between the threads within a data channel; the thread corresponding to the data channel is put into sleep through corresponding control parameters, and / or the synchronous awakening of the thread in the sleep state.

[0023] On the other hand, an embodiment of the present application provides a computer device, including a memory and a processor, wherein the memory stores a computer program that can be executed on the processor, and when the processor executes the program, some or all of the steps in the above method are implemented.

[0024] On the other hand, an embodiment of the present application provides a computer-readable storage medium having a computer program stored thereon, which implements part or all of the steps in the above method when executed by a processor.

[0025] Based on the above embodiments provided by the present application, by processing one or more data channels in parallel with a work group, the computing resources of a multi-core processor can be fully utilized to improve the efficiency of data processing; at the same time, for threads in one data channel, sleep is achieved through corresponding control parameters, and when it is necessary to continue processing, the threads in sleep state are synchronously awakened, so that instruction synchronization of one or more data channels of the work group can be achieved. The above embodiments enable the graphics processor to handle multi-thread synchronization issues more efficiently.

[0026] It should be understood that the above general description and the following detailed description are merely exemplary and explanatory, and are not intended to limit the technical solutions of the present application. BRIEF DESCRIPTION OF THE DRAWINGS

[0027] The drawings herein are incorporated into the specification and constitute a part of the specification. These drawings illustrate embodiments consistent with the present application and are used together with the specification to illustrate the technical solution of the present application.

[0028] Figure 1 A schematic diagram of the relationship between a kernel, a work group, a wavefront and a thread provided in an embodiment of the present application; Figure 2 A schematic diagram of a data relationship of data dependency between different threads provided in an embodiment of the present application; Figure 3 A schematic diagram of data storage of an arrival count cache provided in an embodiment of the present application; Figure 4 A schematic diagram of a thread sleep and wake-up solution based on barrier instructions provided in an embodiment of the present application; Figure 5 A schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application Figure 1 ; Figure 6 A schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application Figure 2 ; Figure 7 A schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application Figure 3 ; Figure 8 A schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application Figure 4 ; Fig. 9 A schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application Figure 5 ; Fig.10 A schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application Figure 6 ; Fig.11A schematic diagram of initialization record values ​​of an arrival count cache and a sleep control cache provided in an embodiment of the present application; Fig.12 A schematic diagram of a thread asynchronous waiting solution based on an arrival instruction and a wait instruction provided in an embodiment of the present application; Fig.13 A schematic diagram of the structure of an instruction execution device provided in an embodiment of the present application; Fig.14 A hardware entity diagram of a computer device provided in an embodiment of the present application. DETAILED DESCRIPTION

[0029] In order to make the purpose, technical solutions and advantages of the present application clearer, the technical solutions of the present application are further elaborated in detail below in conjunction with the drawings and embodiments. The described embodiments should not be regarded as limiting the present application. All other embodiments obtained by ordinary technicians in the field without making creative work are within the scope of protection of the present application.

[0030] In the following description, reference is made to "some embodiments", which describe a subset of all possible embodiments, but it is understood that "some embodiments" may be the same subset or different subsets of all possible embodiments, and may be combined with each other without conflict. The terms "first / second / third" involved are merely used to distinguish similar objects and do not represent a specific order for the objects. It is understood that "first / second / third" may be interchanged in a specific order or sequence where permitted, so that the embodiments of the present application described herein can be implemented in an order other than that illustrated or described herein.

[0031] Unless otherwise defined, all technical and scientific terms used herein have the same meaning as those commonly understood by those skilled in the art to which this application belongs. The terms used herein are only for the purpose of describing this application and are not intended to limit this application.

[0032] (1) Graphics Processing Unit (GPU): A microprocessor that performs image and graphics-related computing tasks on personal computers, workstations, game consoles, and some mobile devices (such as tablets and smartphones). Graphics processors are also called display cores, visual processors, and display chips.

[0033] (2) Kernel: A program defined by general computing standard frameworks such as OpenCL / DirectX / Vulkan and executed on the GPU.

[0034] (3) Thread (work-item): The smallest unit of program execution. The kernel is a three-dimensional space composed of a specified number of work-items. The work-items execute the program independently and in parallel.

[0035] (4) Work-group: The kernel is divided into three-dimensional work-item sets of equal size. The same work-group can only be executed in the same instruction processor. The work-items within the work-group can share internal memory (local memory) and are synchronized through work-group barriers.

[0036] (5) Instructions: The program consists of instructions that instruct the GPU to perform various operations, such as arithmetic budgeting, data reading and writing, synchronization, etc.

[0037] (6) Work-group barrier: used for inter-process synchronization between work-items within the same work-group.

[0038] (7) Wavefront: A thread bundle or thread group is a way for GPUs to implement parallel computing. A wavefront is a group of threads that start executing at the same time and are executed in parallel on the GPU's stream multiprocessor (SM). Each SM can only execute one wavefront at a time. A wavefront usually contains 16 or 32 work-items, which are executed in parallel. The execution process of a wavefront usually requires multiple clock cycles.

[0039] See also Figure 1 , Figure 1 A schematic diagram of the relationship between a kernel, a workgroup, a wavefront, and a thread is provided for an embodiment of the present application. In which, a kernel 10 may include multiple workgroups 11, and the smallest unit of a workgroup 11 is a thread 12. During parallel computing, multiple threads 12 executed simultaneously in a workgroup 11 form a wavefront 13.

[0040] In parallel computing programs, a large amount of data is divided into three-dimensional threads (work-items) that execute the same instruction operations in parallel to achieve data parallel processing. However, in actual applications, the data between threads within a work-group inevitably has certain correlations.

[0041] For example, please refer to the following code implementation for data dependency scenarios: __kernel void sample_kernel(__global int*a, __global int*b, __localint*local_mem) { int global_id = get_global_id(0); / / thread id in kernel int local_id = get_local_id(0); / / thread id in the current workgroup int local_size = get_local_size(0); / / Number of threads in the working group local_mem[local_id] = a[global_id]; / / Use the id of the current thread in the kernel as the offset to get the value of cache A and store it at the corresponding offset in the internal memory.

[0042] barrier(CLK_LOCAL_MEM_FENCE); / / Barrier instruction ensures that all threads in the current working group execute the barrier instruction before executing the following instructions.

[0043] b[global_id] = local_mem [(local_id+1)%local_size]; / / Use the id of the adjacent thread in the kernel as the offset to obtain the value of the internal memory, and store it in the B cache with the current thread as the offset position. In other words, the current thread has data dependency with the adjacent thread.

[0044] } The above program code describes the data dependencies between different threads. Figure 2 , Figure 2 A data relationship diagram of data dependency between different threads provided in an embodiment of the present application. The program copies the contents of the A cache 21 to the B cache 24 with an offset of one bit (a work-item offset), and temporarily stores the intermediate results through the internal memory 22. The barrier instruction 23 is used to ensure that all threads of the current work group can execute subsequent instructions only after executing the barrier instruction 23, so that when assigning a value to the B cache 24, it is ensured that the corresponding contents of the A cache 21 have been written into the internal memory 22.

[0045] Figure 3 A schematic diagram of data storage of an arrival count cache provided in an embodiment of the present application, and implementation of barrier instructions by existing technical solutions Figure 3As shown, it is generally achieved by counting the number of threads corresponding to the execution of the barrier instruction. The value is stored in the arrival counter buffer 31. Each working group has a corresponding record value (entry), such as Figure 3 Includes the number of threads corresponding to each of the N+1 working groups.

[0046] In some embodiments, if the program uses a barrier instruction, before the program is started, the hardware will initialize the arrival count cache 31, and the initialization value is fixed to the number of all threads in the work group (the number of work group threads), that is, the total number of threads that the current work group expects to execute the barrier instruction. When the thread executes the barrier instruction, the record value corresponding to the current work group in the arrival count cache 31 will be atomically reduced by one, and the current thread will be descheduled, so that it cannot continue to execute subsequent instructions.

[0047] When the record value corresponding to the current working group is reduced to 0, it means that all threads of the current working group have executed the barrier instruction, the synchronization has been completed, and all threads of the current working group are immediately awakened (re-scheduled) to continue executing the instructions after the barrier instruction.

[0048] Figure 4 A schematic diagram of a thread sleep and wake-up scheme based on barrier instructions provided in an embodiment of the present application.

[0049] The current work group includes six threads, from thread 0 to thread 5. Before the program starts, the arrival count cache is initialized to the number of threads in the work group, that is, the record value 41 of the arrival count cache is initialized to 6. After the program starts, each thread will execute the corresponding instruction (the instruction before the barrier instruction). After a thread executes the barrier instruction 42, it will enter the sleep state and decrement the record value 41. Figure 4 Thread 5 in the example has no data dependency and can directly execute barrier instruction 42 and enter the sleep state. At the same time, it decrements the record value 41, and the record value 41 is 5. After that, thread 2 also executes barrier instruction 42 and enters the sleep state while decrementing the record value 41. At this time, the record value 41 is 4. This process continues in this way until thread 3 executes barrier instruction 42. At this time, the record value 41 has been decremented to 0, and thread 3 does not need to enter the sleep state. At the same time, it executes wake-up operation 43, and threads 0 to 5 are all awakened (thread 3 does not enter the sleep state) and continue to execute subsequent instructions.

[0050] It can be seen that the above scheme has the following problems: (1) The barrier instruction is strongly synchronized. The condition for the wake-up operation to be executed is that all threads in the same work group have executed the barrier instruction. In other words, all threads will be blocked by the thread that executes the barrier instruction the slowest (e.g. Figure 4 Thread 3 in the process). (2) The record value of the arrival count cache is initialized to the number of threads in the work group. Some threads without data dependencies also need to participate in the barrier instruction and are blocked (such as Figure 4 Thread 5 in the thread), reducing execution efficiency. (3) Only one barrier instruction is supported. For multi-channel data scenarios (data within each channel is correlated, but data between channels is not correlated), barrier instructions must be used serially, which affects efficiency.

[0051] In order to solve the above problems, the embodiment of the present application provides an instruction execution method. The computer device may refer to a server, a laptop, a tablet computer, a desktop computer, a smart TV, a set-top box, a mobile device (such as a mobile phone, a portable video player, a personal digital assistant, a dedicated messaging device, a portable gaming device) and other devices with data processing capabilities.

[0052] Figure 5 A schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application Figure 1 ,like Figure 5 As shown, the method includes the following step S101: Step S101, processing one or more data in parallel through a working group; wherein the working group includes multiple threads, the multiple data are independent of each other, and there is correlation between threads within one data; the thread corresponding to the one data is put into sleep through corresponding control parameters, and / or synchronously awakens the thread in sleep state.

[0053] The work group is a basic organizational unit in GPU parallel computing, consisting of multiple threads (Work-Items), where the multiple threads are executed on the same instruction processor and can share internal memory. In related technologies, threads within a work group can be synchronized through barrier instructions (Work-Group Barrier).

[0054] In some embodiments, the work group may process multiple paths of data in parallel, that is, the multiple threads included in the work group may be divided into multiple parts, and a part of the threads is used to process one path of data.

[0055] In some embodiments, when a work group processes one path of data in parallel, all threads in the work group may participate in processing this path of data; or some threads (at least one thread) in the work group may participate in processing this path of data.

[0056] In some embodiments, when a work group processes multiple data paths in parallel, all threads in the work group may participate in processing this path of data. In this case, all threads may be divided into multiple parts, and some of the threads are used to process one path of data. Alternatively, some threads (at least two threads) in the work group may participate in processing this path of data. In this case, some of the threads may be divided into multiple parts, and some of the threads are used to process one path of data.

[0057] It should be noted that one thread can only process one path of data, and different paths of data correspond to different threads.

[0058] Among them, multiple data paths are independent of each other, that is, there is no data dependency between different data paths, so they can be processed in parallel without interfering with each other; accordingly, due to the independence between different data paths, different data paths can be assigned to different threads in the process group for processing, thereby speeding up the calculation speed.

[0059] There is correlation between threads in one channel of data, that is, there may be data dependency between at least two threads processing the same channel of data, and thus synchronization between threads is required to ensure the correctness and consistency of data processing.

[0060] It can be understood that in the embodiments of the present application, a synchronization method can be provided for each channel of data in multiple channels. The following embodiments will be described with "one channel of data", where one channel of data can be each channel of data in multiple channels, or any channel of data in multiple channels.

[0061] In the embodiment of the present application, the thread corresponding to one path of data is controlled by the control parameters corresponding to the one path of data, and the control parameters corresponding to different paths of data are different.

[0062] For example, if there are A-channel data and B-channel data, corresponding control parameters can be set for A-channel data and B-channel data respectively. The A-channel data is controlled by the control parameters corresponding to the A-channel data; the B-channel data is controlled by the control parameters corresponding to the B-channel data, and the control parameters corresponding to the A-channel data and the control parameters corresponding to the B-channel data are two data.

[0063] In some embodiments, the control parameters corresponding to the one-way thread can control all threads corresponding to the one-way data to sleep; in other embodiments, the control parameters corresponding to the one-way thread can control the synchronous awakening of all threads corresponding to the one-way data that are in a sleeping state; in yet other embodiments, the control parameters corresponding to the one-way thread can control all threads corresponding to the one-way data to sleep, and control the synchronous awakening of all threads corresponding to the one-way data that are in a sleeping state.

[0064] Based on the above embodiments provided by the present application, by processing one or more data channels in parallel with a work group, the computing resources of a multi-core processor can be fully utilized to improve the efficiency of data processing; at the same time, for threads in one data channel, sleep is achieved through corresponding control parameters, and when it is necessary to continue processing, the threads in sleep state are synchronously awakened, so that instruction synchronization of one or more data channels of the work group can be achieved. The above embodiments enable the graphics processor to handle multi-thread synchronization issues more efficiently.

[0065] In some embodiments, the control parameter includes an arrival count value corresponding to the one-way data; the arrival count value is used to represent the number of threads that have not yet executed to the waiting point among the at least two threads corresponding to the one-way data.

[0066] Among them, a corresponding arrival count value can be set for each channel of data, wherein the arrival count value can change with the execution of instructions by each thread corresponding to the channel of data. At any time, the arrival count value represents the number of threads that have not yet executed to the waiting point among at least two threads corresponding to the channel of data.

[0067] In some embodiments, the initial value of the arrival count value is used to indicate the total number of threads participating in the barrier instruction. It is understandable that the initial values ​​of the arrival count values ​​corresponding to different paths of data may be the same or different, and may be customized based on the actual operating environment.

[0068] Among them, the above-mentioned waiting point is the point where all threads corresponding to one path of data need to be synchronized during the execution of instructions; in some embodiments, the thread corresponding to one path of data needs to sleep (wait) after executing to the waiting point until all threads (at least two threads) corresponding to one path of data have executed to the waiting point, and all threads in the sleeping state continue to execute.

[0069] Based on the above-mentioned embodiments provided in the present application, by setting a corresponding arrival count value for each path of data and making it change dynamically as the thread executes instructions, the number of threads that have not yet executed to the waiting point among at least two threads can be reflected in real time, thereby providing accurate data support for thread synchronization; at the same time, the initial value of the arrival count value can be customized according to the actual operating environment, and can be flexibly adapted to the needs of different scenarios.

[0070] In some embodiments, the instruction execution method further includes: in response to the thread corresponding to the one-way data executing to the waiting point, updating the arrival count value corresponding to the one-way data.

[0071] Among them, since the arrival count value is used to represent the number of threads that have not yet executed to the waiting point among the at least two threads corresponding to one path of data, after the thread corresponding to the one path of data executes to the waiting point, the arrival count value corresponding to the one path of data needs to be updated.

[0072] Exemplarily, the thread group includes threads 1 to 8, wherein threads 1 to 4 correspond to A-way data, and threads 5 to 8 correspond to B-way data; for A-way data, the corresponding arrival count value can be initialized to 4. During the process of threads 1 to 4 executing instructions in parallel, if thread 3 executes to the waiting point, the arrival count value is updated to 3; if thread 1 executes to the waiting point, the arrival count value is updated to 2, and so on, until threads 1 to 4 all execute to the waiting point, at which time the arrival count value is 0, indicating that the number of threads whose A-way data has not yet executed to the waiting point is 0.

[0073] Based on the above embodiments provided by the present application, since the arrival count value is used to characterize the number of threads that have not yet executed to the waiting point, the gradual decrease of the arrival count value as the thread executes can intuitively display the progress of thread synchronization of one data path.

[0074] In some embodiments, the instruction execution method further includes: in response to the arrival count value indicating that there is no thread among the at least two threads that has not yet executed to the waiting point, waking up a thread in a dormant state among the at least two threads.

[0075] In some embodiments, in response to the thread executing to the waiting point, while updating the arrival count value corresponding to the one-way data, the thread also needs to be set to a sleep state.

[0076] In the above embodiment, in response to the arrival count value indicating that there is no thread among the at least two threads that has not yet executed to the waiting point, that is, the arrival count value is 0, a thread in a dormant state among the at least two threads is awakened.

[0077] In some embodiments, when the arrival count value indicates that there is no thread in the at least two threads that has not yet executed to the waiting point, the number of threads in the dormant state may be the number of all threads corresponding to the data path. For example, taking the data path A as an example, when the corresponding arrival count value is 0, all threads corresponding to the data path A are in the dormant state, that is, the number of threads in the dormant state is 4.

[0078] In other embodiments, when the arrival count value indicates that there is no thread that has not yet executed to the waiting point among the at least two threads, the number of threads in the dormant state may also be less than the number of all threads corresponding to the one path of data. In the current scenario, among the at least two threads corresponding to one path of data, there may be one or more threads that do not need to use shared memory data, that is, the one or more threads do not need to participate in the instruction synchronization process of the current waiting point. Therefore, in response to the thread that does not need to use the shared memory data executing to the waiting point, while updating the arrival count value corresponding to the one path of data, the thread that does not need to use the shared memory data does not need to enter the dormant state and continues to execute instructions after the waiting point. Exemplarily, taking the above-mentioned A-way data as an example, its corresponding arrival count value can be initialized to 4, where thread 1 is a thread that does not need to use shared memory data. In the process of threads 1 to 4 executing instructions in parallel, if thread 3 executes to the waiting point, the arrival count value is updated to 3, and thread 3 is controlled to enter a sleep state; if thread 1 executes to the waiting point, the arrival count value is updated to 2. Since thread 1 is a thread that does not need to use shared memory data, thread 1 does not need to sleep and continues to execute subsequent instructions; and so on, until threads 1 to 4 all execute to the waiting point, the arrival count value is 0, and the threads in the sleep state include thread 1, thread 2 and thread 4, which is less than the number of all threads corresponding to the A-way data.

[0079] Based on the above embodiments provided in the present application, by waking up a dormant thread in response to a count value indicating that there is no thread among at least two threads that has not yet executed to the waiting point, accurate synchronization between threads can be achieved, so that the thread corresponding to one path of data continues to execute after completing the synchronization.

[0080] In some embodiments, the control parameter includes a sleep indication parameter; the sleep indication parameter is used to indicate a thread in a sleep state among the at least two threads corresponding to the one channel of data.

[0081] Each path of data corresponds to a sleep indication parameter, and the sleep indication parameter is used to indicate a thread in a sleep state among at least two threads corresponding to the path of data.

[0082] In some embodiments, a corresponding thread identifier may be set for each thread, and the sleep indication parameter may include the thread identifier of each thread in a sleep state. For example, continuing to take the above-mentioned A-way data as an example, when the sleep indication parameter is "2, 3", it means that among the threads corresponding to the A-way data, thread 2 and thread 3 are in a sleep state.

[0083] In other embodiments, the sleep indication parameter may include a sleep indication bit for each thread in one channel of data, and the sleep indication bit is used to indicate whether the corresponding thread is in a sleep state. Exemplarily, continuing with the above-mentioned A channel data as an example, when the sleep indication parameter is "1010", it means that among the threads corresponding to the A channel data, thread 1 and thread 3 are in a sleep state. Here, when the sleep indication bit is "1", it indicates that the corresponding thread is in a sleep state, and when the sleep indication bit is "0", it indicates that the corresponding thread is not in a sleep state. Of course, other values ​​can also be set in the sleep indication bit to indicate whether the thread is in a sleep state, such as "Y / N", "T / F", etc.

[0084] In some embodiments, the instruction execution method further includes: in response to the thread corresponding to the one channel of data entering a sleep state, updating the sleep indication parameter.

[0085] For this channel of data, each time a thread enters the sleep state, the sleep indication parameter is updated once.

[0086] Exemplarily, continuing with the above-mentioned A-way data as an example, in the process of threads 1 to 4 executing instructions in parallel, if thread 3 executes to the waiting point, the sleep indication parameter is updated to "3", and thread 3 is controlled to enter the sleep state; if thread 1 executes to the waiting point, the sleep indication parameter is updated to "1, 3", and thread 3 is controlled to enter the sleep state; and so on, until threads 1 to 4 all execute to the waiting point and enter the sleep state, at which time the sleep indication parameter is "1, 2, 3, 4".

[0087] As another example, continuing to take the above-mentioned A-way data as an example, in the process of parallel execution of instructions by threads 1 to 4, if thread 3 executes to the waiting point, the sleep indication parameter is updated to "0010", and thread 3 is controlled to enter the sleep state; if thread 1 executes to the waiting point, the sleep indication parameter is updated to "1010", and thread 3 is controlled to enter the sleep state; and so on, until threads 1 to 4 all execute to the waiting point and enter the sleep state, at which time the sleep indication parameter is "1111"; here, when the sleep indication bit is "1", it indicates that the corresponding thread is in the sleep state, and when the sleep indication bit is "0", it indicates that the corresponding thread is not in the sleep state. Of course, other values ​​can also be set in the sleep indication bit to indicate whether the thread is in the sleep state, such as "Y / N", "T / F", etc.

[0088] In some embodiments, the instruction execution method further includes: in response to reaching a wake-up condition, waking up the thread in a dormant state corresponding to the data channel based on the dormancy indication parameter.

[0089] The wake-up condition includes that there is no thread that has not yet executed to the waiting point among the at least two threads. If there is no thread that has not yet executed to the waiting point among the threads corresponding to one line of data, that is, all threads have executed to the waiting point, it means that all related threads have completed their respective instructions to reach the waiting point and reached a common execution state. At this time, the wake-up condition is met, and the thread that is in a dormant state due to waiting for other threads can be awakened based on the state of the sleep indication parameter.

[0090] In some embodiments, considering that the arrival count value is used to represent the number of threads that have not yet executed to the waiting point among the at least two threads corresponding to the one path of data, the above wake-up condition can be updated as: the arrival count value is 0.

[0091] In some embodiments, since the sleep indication parameter is used to indicate the thread in the sleep state among the at least two threads, after the wake-up condition is met, the thread in the sleep state can be determined among the at least two threads based on the sleep indication parameter, and the thread in the sleep state can be awakened.

[0092] Exemplarily, continuing to take the above-mentioned A-way data as an example, in response to reaching the wake-up condition, if the sleep indication parameter is "0011", it means that the threads in the sleep state at this time include thread 3 and thread 4, and thread 3 and thread 4 are woken up.

[0093] Based on the above-mentioned embodiments provided by the present application, by waking up a thread in a dormant state in response to reaching a wake-up condition (i.e., there is no thread in at least two threads that has not yet executed to the waiting point, or the count value reaches 0), accurate synchronization between threads can be achieved, so that the thread corresponding to one path of data continues to execute after completing the synchronization; at the same time, the sleep indication parameter is used to indicate which threads are in a dormant state, so that when the wake-up condition is met, these threads can be accurately determined and woken up, avoiding unnecessary waiting and waste of resources.

[0094] Figure 6 This is a schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application. Figure 2 The control parameters include the arrival count value and the sleep control parameter; the process of processing one or more data channels in parallel through the working group includes step S201 and step S202, combining Figure 6 The steps shown are explained.

[0095] Step S201: After one or more threads corresponding to the one channel of data execute an arrival instruction, the arrival count value is updated.

[0096] The barrier instruction is used to control a group of threads to wait at a waiting point until all participating threads reach the waiting point, and then let the group of threads continue to execute subsequent instructions together. In order to avoid the phenomenon in the related art that all threads are blocked by the slowest thread executing the barrier instruction, the embodiment of the present application provides an arrival instruction to achieve the same function as the barrier instruction in the related art. The arrival instruction is used to indicate that the process has been executed to a position where synchronization is required.

[0097] In the above step S201, each thread corresponding to one data channel will update the arrival count value corresponding to the one data channel when executing the arrival instruction. It can be understood that the arrival count value at the current moment reflects the number of threads that have not executed the arrival instruction at the current moment among all threads corresponding to one data channel.

[0098] Step S202: When the arrival count value indicates that there is no thread that has not yet executed the arrival instruction among the at least two threads corresponding to the one path of data, waking up the thread in the dormant state based on the dormancy control parameter.

[0099] Among them, if there is no thread that has not yet executed the arrival instruction in the threads corresponding to one data channel, that is, all threads have executed the arrival instruction, it means that all related threads have completed their respective instructions to reach the arrival instruction and reached a common execution state. At this time, the thread that is in a dormant state due to waiting for other threads can be awakened based on the state of the sleep indication parameter.

[0100] Based on the above embodiments provided in the present application, synchronization between threads is achieved by introducing arrival instructions and arrival count values. Furthermore, through flexible sleep and wake-up control, the utilization of thread resources is optimized and unnecessary waiting time is reduced.

[0101] In some embodiments, the instruction execution method also includes: for a first thread that executes a reach instruction, in response to the first thread executing a wait instruction, controlling the first thread to enter a sleep state and updating the sleep indication parameter; the first thread is a thread that needs to use shared memory data among the at least two threads corresponding to the one path of data.

[0102] In some embodiments, within the same workgroup, or in other words, among all threads participating in the same barrier instruction, all threads can be divided into a first thread and a second thread based on whether the thread needs to use shared memory data. The first thread is a thread that needs to use shared memory data, and the second thread is a thread that does not need to use shared memory data. Exemplarily, the first thread often needs to use shared memory data as data input, and after processing, output the data to an external memory; the second thread does not need shared memory data, but needs to output data to the shared memory for use by the first thread. The shared memory data is data in shared memory, and the shared memory is an on-chip memory located in the processor.

[0103] In some embodiments, in order to achieve the same function as the barrier instruction in the related art, the embodiment of the present application provides a wait instruction on the basis of providing an arrive instruction; compared with the solution in the related art that the thread enters a sleep state after executing to the wait point, the embodiment of the present application fully implements the function of the barrier instruction in the related art through the two instructions of the arrive instruction and the wait instruction, wherein the arrive instruction is used to indicate that the process has executed to the wait point, and the wait instruction is used to control the process to enter the sleep state.

[0104] In the above embodiment, the reach instruction is located before the wait instruction. In some embodiments, after the process executes the reach instruction, it can continue to execute the wait instruction to enter the sleep state; in other embodiments, after the process executes the reach instruction, it may not execute the wait instruction, that is, it does not need to execute the sleep state.

[0105] In some embodiments, for the thread (first process) that needs to use shared memory data among at least two threads corresponding to one path of data, since the first process needs to use shared memory data, that is, there is a data dependency relationship with other processes corresponding to the one path of data, it needs to enter a sleep state to wait for other threads to execute the arrival instruction.

[0106] That is to say, for the first thread, after executing the reach instruction, the enlarge instruction will certainly be executed.

[0107] In some embodiments, in response to the first thread executing a wait instruction, the first thread is controlled to enter a sleep state, and the sleep indication parameter is updated.

[0108] In some embodiments, a corresponding thread identifier may be set for each thread, and the sleep indication parameter may include the thread identifier of each thread in a sleep state. In other embodiments, the sleep indication parameter may include a sleep indication bit of each thread in a data path, and the sleep indication bit is used to indicate whether the corresponding thread is in a sleep state. For the data path, the sleep indication parameter is updated once each time a thread enters a sleep state.

[0109] It should be noted that, here, after the first thread executes the wait instruction, the specific implementation method of "controlling the first thread to enter the sleep state and updating the sleep indication parameter" is the same as that in the above embodiment, and specific examples can refer to the above embodiment.

[0110] In some embodiments, the instruction execution method further includes: for a second thread that executes to a reach instruction and does not need to use shared memory data, after the second thread executes to the reach instruction, continuing to execute subsequent instructions.

[0111] In some embodiments, the above-mentioned arrival instruction is executed for each thread of the data path. After executing the arrival instruction, if the thread is the first thread, it needs to execute the wait instruction to perform inter-process synchronization; if the thread is the second thread, it does not need to execute the wait instruction.

[0112] In the embodiment of the present application, after the second thread runs to the waiting instruction, since it does not need to execute the waiting instruction, compared with the first thread, the second thread will not enter the dormant state, and the second thread can directly execute the subsequent instructions. The subsequent instructions are instructions after the arrival instruction.

[0113] In some embodiments, each thread corresponding to one path of data runs the same program, but the data processed by the program is different. Therefore, among all the threads for the same path of data, there may be some threads (first threads) that must wait for all threads to run until the instruction is reached in order to obtain the shared memory data output by all threads; there may be another part of the threads (second threads) that do not need to wait for all threads to run until the instruction is reached, that is, do not need the shared memory data output by all threads, and can execute subsequent instructions.

[0114] Therefore, for the second thread that does not need to use the shared memory data corresponding to one data path, after the second thread executes the arrival instruction, the subsequent instructions can be executed immediately. Since these threads do not need to use the shared memory data, the second thread can continue to execute when the first thread executes the wait instruction, thereby improving the parallelism and efficiency of the program.

[0115] In an embodiment of the present application, by dividing the thread that does not need to use shared memory data into a second thread, and distinguishing the second thread from other first threads that need to use shared memory data, more efficient thread management and synchronization can be achieved, thereby improving the performance and response speed of the program; at the same time, the present application controls the second thread to continue executing subsequent instructions after executing the arrival instruction, without waiting for all threads to execute the arrival instruction, thereby achieving the effect of an asynchronous barrier instruction.

[0116] In some embodiments, the above-mentioned instruction execution method is applied to a graphics processor, and an arrival count cache and a sleep control cache are set in the on-chip storage of the graphics processor; wherein the arrival count cache is used to store the arrival count value corresponding to each data path in the work group; and the sleep control cache is used to store the sleep control parameters corresponding to each data path in the work group.

[0117] In the embodiment of the present application, an arrival count cache and a sleep control cache are provided in the on-chip storage of the graphics processor; the arrival count value is stored in the arrival count cache, and the sleep control parameter is stored in the sleep control cache.

[0118] Exemplarily, taking the case where there are two working groups, each of which contains two paths of data, please refer to Table 1 for a storage schematic table of arrival count values ​​and sleep control parameters.

[0119] Table 1

[0120] It can be seen that the arrival count buffer stores the arrival count value corresponding to each data path in each working group; the sleep control buffer stores the sleep control parameter corresponding to each data path in each working group.

[0121] In the embodiment of the present application, the arrival count value corresponding to each data path in the work group is stored in the arrival count cache, and the graphics processor can quickly update and obtain the arrival count value of each data path, thereby accurately controlling the execution of the thread; at the same time, the sleep control parameter corresponding to each data path in the work group is stored in the sleep control cache, and the graphics processor can quickly determine which threads are in a sleep state and wake up these threads accurately when needed, thereby optimizing the execution efficiency of the thread. The above embodiment enables the graphics processor to handle multi-thread synchronization problems more efficiently, thereby improving the efficiency and performance of the program.

[0122] Figure 7 This is a schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application. Figure 3 The instruction execution method further comprises step S301, combining Figure 7 The steps shown are explained.

[0123] Step S301, executing an initialization instruction corresponding to one channel of data through any thread in the working group; the initialization instruction is used to set an initial value of an arrival count value corresponding to one channel of data to the total number of threads corresponding to one channel of data.

[0124] In some embodiments, the initialization instruction may be executed by any thread in the work group. For example, when the work group includes threads 1 to 8, the initialization instruction may be executed by any thread in threads 1 to 8. For example, it may be executed by thread 1.

[0125] In some embodiments, each data path has an initialization instruction corresponding to the data path, and the initialization instruction corresponding to the data path is used to set the initial value of the arrival count value corresponding to the data path to the total number of threads corresponding to the data path. Exemplarily, in the case where the work group includes threads 1 to 8, if there are A-path data and B-path data, wherein threads 1 to 4 correspond to A-path data, and threads 5 to 8 correspond to B-path data. Then two initialization instructions need to be executed, the first initialization instruction is used to set the initial value of the arrival count value corresponding to A-path data to 4; the second initialization instruction is used to set the initial value of the arrival count value corresponding to B-path data to 4.

[0126] In some embodiments, the initialization instruction is a barrier instruction in the related art.

[0127] It should be noted that for at least one initialization instruction (barrier instruction) corresponding to a work group, the at least one initialization instruction is executed by the same thread. Based on the above example, when the work group includes threads 1 to 8, if there are two paths of data, the initialization instruction corresponding to each of the two paths of data can be executed by thread 1 in sequence.

[0128] In some embodiments, the initialization instruction is further used to set the initial value of the sleep control parameter corresponding to one channel of data to null, that is, to set the initial value of the sleep control parameter corresponding to one channel of data to a state where no thread identifier of any thread is stored.

[0129] In the embodiment of the present application, the initial value of the arrival count value corresponding to one path of data is set to the total number of threads corresponding to one path of data through an initialization instruction, so that the arrival count value can accurately represent the number of threads that have not executed the arrival instruction.

[0130] In some embodiments, in order to prevent each thread from executing instructions before the initialization instruction, after step S301, the instruction execution method further includes step S302.

[0131] Step S302: In response to all threads in the work group executing the initialization barrier instruction, control each thread in the work group to continue executing subsequent instructions.

[0132] In the embodiment of the present application, the above-mentioned initialization barrier instruction is a barrier instruction defined in the relevant technology, and has not been updated to an arrival instruction and a waiting instruction. The initialization barrier instruction is intended to ensure that all threads in the work group have their control parameters corresponding to each data path initialized before executing subsequent instructions; here, the subsequent instructions of the initialization barrier instruction include the instructions between the initialization barrier instruction and the arrival instruction.

[0133] In an embodiment of the present application, by starting to execute subsequent instructions only after all threads have reached the initialization barrier instruction, it can be ensured that all threads start executing subsequent instructions in the same initial state, thereby avoiding the situation where the thread has executed the arrival instruction before the control parameters (arrival count value / sleep indication parameter) are initialized.

[0134] Figure 8 This is a schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application. Figure 4 .based on Figure 8 The instruction execution method further includes step S401, combining Figure 8 The steps shown are explained.

[0135] Step S401 : in response to each thread corresponding to the one-way data executing the arrival instruction, subtract one from the arrival count value corresponding to the one-way data.

[0136] In the embodiment of the present application, whenever a thread of the data of one path executes the arrival instruction, an atomic subtraction operation is performed on the arrival count value corresponding to the data of one path.

[0137] Among them, since the initial value of the arrival count value corresponding to the one-way data is the total number of threads of the one-way data, as the number of threads that execute to the arrival instruction increases, the arrival count value corresponding to the one-way data will decrease accordingly until it reaches zero. At this time, each thread representing the one-way data executes to the arrival instruction, and the thread in the dormant state can be awakened.

[0138] In the embodiment of the present application, by performing an atomic operation on the arrival count value of one data path at the arrival instruction, the consistency and correctness of the data in a multi-threaded environment can be ensured. At the same time, by reducing the arrival count value of one data path by one, the arrival count value can accurately represent the number of threads that have not executed the arrival instruction.

[0139] Fig. 9 This is a schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application. Figure 5The instruction execution method further comprises step S501, combining Fig. 9 The steps shown are explained.

[0140] Step S501: Based on the thread identifier of the first thread, in the sleep control parameter corresponding to the one channel of data, a sleep indication bit corresponding to the thread identifier of the first thread is updated from a first value to a second value.

[0141] Wherein, while the sleep indication bit corresponding to the thread identifier of the first thread is updated from the first value to the second value, the first thread is controlled to enter the sleep state.

[0142] In the embodiment of the present application, since the first thread is a thread that needs to use shared memory data, after executing the above-mentioned arrival instruction, it is also necessary to execute the wait instruction. In this way, it can be ensured that all first threads continue to execute subsequent instructions only after executing the arrival instruction.

[0143] In the above embodiment, after the first thread executes the wait instruction, it needs to perform two operations at the same time: first, based on the thread identifier of the first thread, it needs to update the sleep indication bit corresponding to the thread identifier of the first thread in the sleep control parameter corresponding to the data path from the first value to the second value; second, it needs to control the first thread to enter the sleep state.

[0144] The thread ID is a unique identifier for each thread, which is usually provided by an operating system or a thread library. The thread ID can be an integer, or a unique string or handle. In some embodiments, the thread ID of the first thread is the thread program number of the first thread in the work group. Exemplarily, when the work group includes threads 1 to 8, the thread ID of thread 1 is "1", the thread ID of thread 2 is "2", and so on.

[0145] Correspondingly, the above-mentioned updating of the sleep indication bit corresponding to the thread identifier of the first thread from the first value to the second value in the sleep control parameter corresponding to the one data path includes: updating the sleep indication bit corresponding to the thread program number of the first thread from the first value to the second value in the sleep control parameter corresponding to the one data path.

[0146] The sleep control parameter may include at least one sleep indication bit arranged in sequence, and each sleep indication bit corresponds to a thread.

[0147] In some embodiments, the order of the sleep indication bits in the sleep control parameter is the same as the thread program number of the thread corresponding to the sleep indication bit in the work group. For example, when the work group includes threads 1 to 8, the sleep indication bit corresponding to thread 1 may be the first bit in the sleep control parameter, the sleep indication bit corresponding to thread 2 may be the second bit in the sleep control parameter, and so on, the sleep indication bit corresponding to thread 8 may be the eighth bit in the sleep control parameter.

[0148] In some embodiments, the sleep indication bit can be set to a first value or a second value. When the sleep indication bit of a thread is set to the first value, it indicates that the thread is not in a sleep state; when the sleep indication bit of the thread is set to the second value, it indicates that the thread is in a sleep state.

[0149] In an embodiment of the present application, while controlling the first thread to enter the sleep state, it is necessary to update the sleep indication bit corresponding to the first thread from a first value indicating that the first thread is not in the sleep state to a second value indicating that the first thread is in the sleep state. Exemplarily, the first value can be set to 0, and the second value can be set to 1.

[0150] To facilitate understanding of the above embodiment, a working group including 8 threads from thread 1 to thread 8 is taken as an example for explanation, wherein these 8 threads all belong to the first thread, and thread 1 to thread 3 belong to A-way data, and thread 4 to thread 8 belong to B-way data.

[0151] After all eight threads have executed the initialization barrier instruction, for the A-way data, if thread 2 executes to the reach instruction first and then executes to the wait instruction, the sleep control parameter corresponding to the B-way data can be updated from "000b" (also represented as "0b") to "010b" (also represented as "10b"), that is, the second bit of the sleep control parameter is updated from the first value "0" to the second value "1". Afterwards, if thread 1 executes to the reach instruction and then executes to the wait instruction, the sleep control parameter corresponding to the A-way data can be updated from "010b" (also represented as "10b") to "011b" (also represented as "11b"), that is, the first bit of the sleep control parameter is updated from the first value "0" to the second value "1".

[0152] After all eight threads have executed the initialization barrier instruction, for the B-way data, if thread 8 executes to the reach instruction first and then executes to the wait instruction, the sleep control parameter corresponding to the B-way data can be updated from "00000000b" (also represented as "0b") to "10000000b", that is, the eighth bit of the sleep control parameter is updated from the first value "0" to the second value "1". Afterwards, if thread 4 executes to the reach instruction and then executes to the wait instruction, the sleep control parameter corresponding to the B-way data can be updated from "10000000b" to "10001000b", that is, the fourth bit of the sleep control parameter is updated from the first value "0" to the second value "1".

[0153] In some embodiments, while in the sleep state, the thread does not perform any operations and does not consume processor resources.

[0154] In an embodiment of the present application, for a first thread that needs to use shared memory data for one path of data, when executing a wait instruction, the sleep indication bit corresponding to the thread identifier is updated from a first value to a second value, thereby marking the first thread as being in a sleep state. In this way, it can be ensured that which threads are in a sleep state are correctly recorded to facilitate subsequent wake-up operations.

[0155] Accordingly, the instruction execution method may further include step S502 to step S503.

[0156] Step S502: Determine the thread identifier of the thread in the dormant state based on the dormancy indication bit in the dormancy control parameter that is at the second value.

[0157] In some embodiments, in response to the arrival count value corresponding to the one-way data being zero, the thread identifier of the thread in the dormant state is determined based on the dormant indication bit in the dormant control parameter being at the second value. Here, the dormant thread is the first thread in the dormant state.

[0158] In some embodiments, after the arrival count value is initialized to the total number of threads, in response to the arrival count value corresponding to the data being reduced to zero, the sleep control parameter corresponding to the data at the current moment is queried. At this time, the sleep control parameter stores the thread identifier of the thread (first thread) that is in sleep state at the current moment.

[0159] Based on the implementation of step S401 above, here, the thread identifier of the first thread currently in the sleep state can be determined by determining the order of the sleep indication bits in the sleep control parameter that are in the second value in the sleep control parameter.

[0160] Exemplarily, the above work group includes 8 threads from thread 1 to thread 8 as an example for explanation, wherein these 8 threads all belong to the first thread, and thread 1 to thread 3 belong to A-way data, and thread 4 to thread 8 belong to B-way data. At the moment when the arrival count value corresponding to the A-way data is reduced to zero, if the sleep control parameter corresponding to the A-way data is obtained as "011b", it can be determined that the sleep indication bit at the second value "1" includes the "1" bit and the "2" bit, and accordingly, it can be determined that the threads in the sleep state in the work group are thread 1 and thread 2. At the moment when the arrival count value corresponding to the B-way data is reduced to zero, if the sleep control parameter corresponding to the B-way data is obtained as "10001000b", it can be determined that the sleep indication bit at the second value "1" includes the "5" bit and the "8" bit, and accordingly, it can be determined that the first threads in the sleep state in the work group are thread 5 and thread 8.

[0161] Step S503: wake up the dormant thread by using the thread identifier of the dormant thread.

[0162] In an embodiment of the present application, after obtaining the thread identifier of a thread in a dormant state, the thread corresponding to the thread identifier can be awakened by a wake-up function provided by the thread library, whereby the awakened thread can regain processor resources and continue to execute subsequent instructions.

[0163] It should be noted that, when there are multiple threads in the dormant state, each of the threads in the dormant state needs to be woken up synchronously.

[0164] In the embodiment of the present application, the thread in the dormant state is awakened by the thread identifier of the dormant thread, so that the correct resumption of execution of the thread in the multi-threaded environment can be achieved.

[0165] In some embodiments, the instruction execution method further includes step S504.

[0166] Step S504: in response to the thread in the dormant state being awakened, based on the thread identifier of the thread, in the dormancy control parameter corresponding to the one channel of data, a dormancy indication bit corresponding to the thread identifier of the thread is updated from the second value to the first value.

[0167] In an embodiment of the present application, when threads in a dormant state are awakened simultaneously, the sleep indication bit corresponding to the thread identifier of the thread in the sleep control parameter corresponding to one channel of data is updated from the second value to the first value. In fact, it can be understood that all sleep indication bits of the sleep control parameter corresponding to the one channel of data are set to the first value.

[0168] Exemplarily, the above work group includes 8 threads from thread 1 to thread 8 as an example for explanation, wherein these 8 threads all belong to threads, and thread 1 to thread 3 belong to the first A-way data, and thread 4 to thread 8 belong to B-way data. After the threads (thread 1 and thread 2) in the dormant state are awakened based on the dormant control parameter "011b" corresponding to the A-way data, the second value "1" of the dormant control parameter "011b" corresponding to the A-way data is reset to the first value "0", that is, "000b" (which can also be expressed as "0b"). After the threads (thread 5 and thread 8) in the dormant state are awakened based on the dormant control parameter "10001000b" corresponding to the B-way data, the second value "1" of the dormant control parameter "10001000b" corresponding to the B-way data is reset to the first value "0", that is, "00000000b" (which can also be expressed as "0b").

[0169] In an embodiment of the present application, when a thread in a dormant state is awakened, the dormant indication bit corresponding to the thread identifier of the thread in the dormant control parameter is updated from the second value to the first value. In this way, at any time, the thread in a dormant state can be accurately determined through the dormant indication bit in all instructions corresponding to one line of data.

[0170] Fig.10 This is a schematic diagram of an implementation process of an instruction execution method provided in an embodiment of the present application. Figure 6 The instruction execution method further comprises step S601, combining Fig.10 The steps shown are explained.

[0171] Step S601, in response to updating the arrival count value corresponding to the one path of data, controlling the first thread to execute an independent instruction; the independent instruction is an instruction before the waiting instruction and does not require the use of shared memory data; the first thread is a thread that needs to use shared memory data among the at least two threads corresponding to the one path of data.

[0172] In an embodiment of the present application, for a first thread that needs to use shared memory data, during the execution of an instruction, between the arrival instruction and the wait instruction, an independent instruction may also be executed. The independent instruction is an instruction in the program run by the first thread that does not need to use the shared memory data output by all threads. Exemplarily, the instruction may be an instruction that uses other non-shared memory data, or an instruction that does not use input data, etc.

[0173] In an embodiment of the present application, in the first thread that needs to use shared memory data, after the execution is completed and before the waiting instruction is executed, instructions that do not need to use the shared memory data output by all threads can be executed first. In this way, before waiting for synchronization between processes, some instructions that do not require shared memory data can be executed first, thereby improving the overall program running efficiency.

[0174] The following describes the application of the instruction execution method provided in the embodiment of the present application in a practical scenario.

[0175] In parallel computing programs, a large amount of data is divided into three-dimensional threads (work-items) that execute the same instruction operations in parallel to achieve data parallel processing. However, in actual applications, the data between threads within a work-group inevitably has certain correlations.

[0176] The existing technical solution is generally implemented by counting the number of threads corresponding to the execution of barrier instructions, see Figure 3 and Figure 4 It can be seen that the above scheme has the following problems: (1) The barrier instruction is strongly synchronized. The condition for the execution of the wake-up operation is that all threads in the same work group have executed the barrier instruction. In other words, all threads will be blocked by the thread that executes the barrier instruction the slowest (such as Figure 4 Thread 3 in the process). (2) The record value of the arrival count cache is initialized to the number of threads in the work group. Some threads without data dependencies also need to participate in the barrier instruction and are blocked (such as Figure 4 Thread 5 in the thread), reducing execution efficiency. (3) Only one barrier instruction is supported. For multi-channel data scenarios (data within each channel is correlated, but data between channels is not correlated), barrier instructions must be used serially, which affects efficiency.

[0177] Considering the defects of the above scheme, the present application proposes an asynchronous barrier instruction, splits the barrier instruction into two stages: arrival and wait, increases the number of barrier instructions, introduces an initialization arrival instruction (init arrive), and supports configurable initialization values ​​of arrival count values ​​of arrival count cache. Among them, the present scheme is still implemented by counting the total number of threads participating in the barrier instruction, and the arrival count cache stores the number of threads participating in the barrier instruction.

[0178] Different from the related technical solutions, this solution supports multiple barrier instructions in the same work group, and the specific number is specified by the user. Therefore, each barrier instruction in each work group occupies a record value (entry) in the arrival count cache. In addition, the newly added arrival instruction supports the user to specify the total number of threads participating in each barrier instruction (expected count). The general usage is that only one thread is required to call the arrival instruction to initialize the total number of threads for each barrier instruction. In addition, a new sleep control cache (de-schedule mask buffer) is added to store the current sleeping thread id for subsequent wake-up operations. Like the arrival count cache, the sleep control cache also occupies one record value for each barrier instruction in each work group, which is subsequently called the sleep control parameter and is initialized to 0.

[0179] See also Fig.11 , Fig.11 A schematic diagram of the initialization record values ​​of an arrival count cache and a sleep control cache provided in an embodiment of the present application. Among them, the arrival count cache 51 stores the arrival count value corresponding to each barrier instruction in each work group, which is the total number of threads participating in each barrier instruction in each work group, such as the barrier instruction 0 of work group 0, and the arrival count value stored in the arrival count cache 51 is the total number of threads participating in the barrier instruction 0 of the work group 0; the sleep control cache 52 also stores the sleep control parameters corresponding to each barrier instruction in each work group, which are initialized to 0.

[0180] In the embodiment of the present application, the barrier instruction is divided into two instructions: a reach instruction and a wait instruction: When the thread executes the arrival instruction, it only atomically decrements the arrival count value of the current working group and the current barrier instruction in the arrival count cache, and then it can continue to execute instructions without data dependencies. When the thread executes the wait instruction, the id of the current thread is stored in the sleep control cache, and the thread then sleeps. When the arrival count value is reduced to 0, the sleep control parameter is used to wake up the thread that has executed the wait instruction, and the sleep control parameter is cleared to 0. Compared with related solutions, the wait state is asynchronous, and there is no need for all threads to execute to the wait state before waking up.

[0181] For example, the procedure of the above scheme is as follows: __kernel void sample_async_bar_kernel(__global int*a, __global int*b,__global int*c, __global int*d,__local int*local_mem_a, __local int* local_mem_b){ int global_id = get_global_id(0); int local_id = get_local_id(0); int local_size = get_local_size(0); if(local_id == 0){ init_arrive(bar0, 5); / / Initialize the arrival count value of a barrier instruction bar0 in the working group init_arrive(bar1, 3); / / Initialize the arrival count value of another barrier instruction bar1 in the working group } barrier(CLK_LOCAL_MEM_FENCE); / / Ensure that all threads in the current working group execute instructions after initialization reaches the count value if(local_id <5){ / / Threads 0 to 4 participate in the barrier instruction bar0 local_mem_a[local_id] = a[global_id]; bar0.arrive(); if(local_id == 0 || local_id == 3){ a[global_id] +=1; / / The 0th and 3rd threads are producers. Since they only need to write the generated data to the internal memory, they do not need to use the data in the internal memory and do not need to wait. } else{ bar0.wait(); b[global_id] = local_mem_a [(local_id+1)%local_size]; } } else if(local_id<8){ / / Threads 5 to 7 participate in barrier instruction bar1 local_mem_b[local_id] = c[global_id]; bar1.arrive(); if(local_id == 7){ c[global_id] +=1; / / The 7th thread is the producer, no need to wait} else{ bar1.wait(); d[global_id] = local_mem_b [(local_id+1)%local_size]; } } else{ / / The 8th work-item does not participate in any barrier instructions a[global_id] +=2; } } See also Fig.12 , Fig.12 A schematic diagram of a thread asynchronous waiting scheme based on an arrival instruction and a wait instruction provided in an embodiment of the present application. The current work group includes nine threads, thread 0 to thread 8, among which threads 0 to thread 4 are the first barrier instruction bar0 supported by the current work group, threads 5 to thread 7 are the second barrier instruction bar1 supported by the current work group, and thread 8 does not participate in any barrier instruction (except the barrier instruction used for initialization).

[0182] Before executing the instruction, it is necessary to initialize the arrival count value and sleep control parameter corresponding to each barrier instruction through any thread. Exemplarily, for the above-mentioned first barrier instruction bar0, the first initialization instruction "Initialize 0" can be called through thread 0 in the figure, and the arrival count value 61 and sleep control parameter 62 of the first barrier instruction bar0 are set to "5, 0b" respectively; for the second barrier instruction bar1, the second initialization instruction "Initialize 1" can be called through thread 0 in the figure, and the arrival count value 66 and sleep control parameter 67 of the second barrier instruction bar1 are set to "3, 0b" respectively. It should be noted that after the initialization instruction of all barrier instructions in the current working group, an initialization barrier instruction 65 is required to ensure that all threads of the current working group execute instructions after the arrival count value is initialized.

[0183] After each thread executes the initialization barrier instruction 65, it needs to execute subsequent instructions. After the threads corresponding to different barrier instructions execute the arrival instruction corresponding to the current barrier instruction, the arrival count value corresponding to the current barrier instruction will be atomically subtracted by one; at the same time, it will determine whether to execute the waiting instruction based on whether the current instruction is a producer.

[0184] The following embodiments will be described by taking the first barrier instruction bar0 as an example. Fig.12 As shown, in chronological order, the execution process of each thread corresponding to the first barrier instruction bar0 includes: (1) Thread 3 first executes the arrival instruction 63 and decrements the arrival count value 61 by one. At this time, the arrival count value 61 is 4.

[0185] (2) Thread 0 executes instruction 63 and decrements the arrival count value 61. At this time, the arrival count value 61 is 3.

[0186] For the above-mentioned thread 0 and thread 3, since thread 0 and thread 3 are producers, there is no need to execute the wait instruction 64, and the sleep control parameter 62 remains unchanged.

[0187] (3) Thread 2 executes the arrival instruction 63 and decrements the arrival count value 61 by one. At this time, the arrival count value 61 is 2.

[0188] (4) Thread 1 executes the arrival instruction 63 and decrements the arrival count value 61 by one. At this time, the arrival count value 61 is 1.

[0189] (5) Since thread 2 is not a producer, it needs to execute the wait instruction 64. Therefore, when thread 2 executes the wait instruction 64, the sleep control parameter 62 is updated to 100b and enters the sleep state.

[0190] (6) Since thread 1 is not a producer, it needs to execute the wait instruction 64. Therefore, when thread 1 executes the wait instruction 64, the sleep control parameter 62 is updated to 110b and enters the sleep state.

[0191] (7) Thread 4 executes to the arrival instruction 63 and decrements the arrival count value 61. At this time, the arrival count value 61 is 0. It should be noted that since the arrival count value 61 is 0, it is necessary to wake up the dormant threads, that is, it is necessary to wake up the dormant threads 1 and 2. Threads 1 and 2 continue to execute subsequent instructions.

[0192] (8) Since thread 4 is not a producer, it needs to execute the wait instruction 64. Therefore, when thread 4 executes the wait instruction 64, the sleep control parameter 62 is updated to 10000b and it enters the sleep state. Since the arrival count value 61 is 0, thread 4 will be awakened immediately and execute subsequent instructions.

[0193] The following embodiments will be described by taking the second barrier instruction bar1 as an example. Fig.12 As shown, in chronological order, the execution process of each thread corresponding to the second barrier instruction bar1 includes: (1) Thread 7 first executes the arrival instruction 68 and decrements the arrival count value 66. At this time, the arrival count value 66 is 2. Since thread 7 is the producer, it does not need to execute the wait instruction 69. Therefore, the sleep control parameter 67 remains unchanged.

[0194] (2) Thread 5 executes the arrival instruction 68 and decrements the arrival count value 66 by one. At this time, the arrival count value 66 is 1.

[0195] (3) Thread 6 executes the arrival instruction 68 and decrements the arrival count value 66 by one. At this time, the arrival count value 66 is 0.

[0196] (4) Since thread 5 is not a producer, it needs to execute the wait instruction 69. When thread 5 executes the wait instruction 69, the sleep control parameter 67 is updated to 100000b, and it enters the sleep state. At the same time, since the arrival count value 66 is 0, thread 5 is immediately awakened and executes subsequent instructions. It can be understood that after thread 5 is awakened, the sleep control parameter 67 is updated to 0b (not shown in the figure).

[0197] (5) Since thread 6 is not a producer, it needs to execute the wait instruction 69. When thread 6 executes the wait instruction 69, the sleep control parameter 67 is updated to 1000000b, and it enters the sleep state. At the same time, since the arrival count value 66 is 0, thread 6 is immediately awakened and executes subsequent instructions. It can be understood that after thread 6 is awakened, the sleep control parameter 67 is updated to 0b (not shown in the figure).

[0198] Based on the above embodiment, the barrier instruction is split into an arrival stage and a waiting stage, that is, into an arrival instruction and a waiting instruction, wherein the arrival instruction only performs an atomic subtraction operation on the arrival count value, and the waiting instruction is used to control the thread to perform a sleep operation to enter a sleep state. When the arrival count value is reduced to 0, the thread that has performed the sleep operation can be awakened without waiting for all threads to execute the wait instruction, thereby realizing an asynchronous barrier instruction; at the same time, the user is supported to specify the number of threads participating in the barrier instruction by adding an initialization arrival instruction. Threads without data dependencies will not be blocked, thereby improving program execution efficiency; at the same time, the sleep control cache is also used to store the thread ID that needs to be waited for, so as to wake up the correct thread later; the same workgroup supports multiple barrier instructions, and the specific number is specified by the user, which can greatly improve the parallelism of multi-channel data scenarios.

[0199] Based on the foregoing embodiments, the embodiments of the present application provide an instruction execution device, which includes the units included and the modules included in the units, and can be implemented by a processor in a computer device; of course, it can also be implemented by a specific logic circuit; in the implementation process, the processor can be a central processing unit (CPU), a microprocessor (MPU), a digital signal processor (DSP) or a field programmable gate array (FPGA), etc.

[0200] Fig.13 A schematic diagram of the structure of an instruction execution device provided in an embodiment of the present application is shown in FIG. Fig.13 As shown, the instruction execution device 1300 includes: an instruction parallel execution module 1310, wherein: Instruction parallel execution module 1310, used for processing one or more data in parallel through a working group; The working group includes multiple threads, multiple data paths are independent of each other, and threads within one data path are correlated; the thread corresponding to one data path realizes sleep and / or synchronous awakening of threads in sleep state through corresponding control parameters.

[0201] In some embodiments, the control parameter includes an arrival count value corresponding to the one-way data; the arrival count value is used to represent the number of threads that have not yet executed to the waiting point among the at least two threads corresponding to the one-way data.

[0202] In some embodiments, the instruction parallel execution module 1310 is further configured to update the arrival count value corresponding to the one-way data in response to the thread corresponding to the one-way data executing to the waiting point.

[0203] In some embodiments, the instruction parallel execution module 1310 is further configured to wake up a dormant thread among the at least two threads in response to the arrival count value indicating that there is no thread among the at least two threads that has not yet executed to the waiting point.

[0204] In some embodiments, the control parameter includes a sleep indication parameter; the sleep indication parameter is used to indicate a thread in a sleep state among the at least two threads corresponding to the one channel of data.

[0205] In some embodiments, the instruction parallel execution module 1310 is further configured to update the sleep indication parameter in response to the thread corresponding to the one channel of data entering a sleep state.

[0206] In some embodiments, the instruction parallel execution module 1310 is further configured to wake up the thread corresponding to the one channel of data in a dormant state based on the dormancy indication parameter in response to reaching the wake-up condition; The wake-up condition includes that there is no thread among the at least two threads that has not yet executed to the waiting point.

[0207] In some embodiments, the control parameters include an arrival count value and a sleep control parameter; the instruction parallel execution module 1310 is also used to update the arrival count value after one or more threads corresponding to the one path of data execute the arrival instruction; when the arrival count value indicates that there is no thread that has not yet executed the arrival instruction among the at least two threads corresponding to the one path of data, the thread in the sleep state is woken up based on the sleep control parameter.

[0208] In some embodiments, the instruction parallel execution module 1310 is also used to control the first thread that executes the reach instruction to enter a sleep state and update the sleep indication parameter in response to the first thread executing the wait instruction; the first thread is a thread that needs to use shared memory data among the at least two threads corresponding to the one path of data.

[0209] In some embodiments, the instruction parallel execution module 1310 is further used for continuing to execute subsequent instructions for a second thread that executes to an arrival instruction and does not need to use shared memory data after the second thread executes to the arrival instruction.

[0210] In some embodiments, the instruction parallel execution module 1310 is also used to execute the initialization instruction corresponding to the one path of data through any thread in the work group; the initialization instruction is used to set the initial value of the arrival count value corresponding to the one path of data to the total number of threads corresponding to the one path of data.

[0211] In some embodiments, the instruction parallel execution module 1310 is further configured to control each thread in the work group to continue to execute subsequent instructions in response to all threads in the work group executing the initialization barrier instruction.

[0212] In some embodiments, the instruction parallel execution module 1310 is further configured to, in response to each thread corresponding to the one path of data executing the arrival instruction, reduce the arrival count value corresponding to the one path of data by one.

[0213] In some embodiments, the instruction parallel execution module 1310 is further used to update the sleep indication bit corresponding to the thread identifier of the first thread in the sleep indication parameter corresponding to the one path of data from a first value to a second value based on the thread identifier of the first thread.

[0214] In some embodiments, the instruction parallel execution module 1310 is also used to determine the thread identifier of the thread in the dormant state based on the sleep indication bit in the sleep control parameter being at the second value; and wake up the thread in the dormant state through the thread identifier of the thread in the dormant state.

[0215] In some embodiments, the instruction parallel execution module 1310 is also used to respond to a thread in a dormant state being awakened, and based on the thread identifier of the thread, update the sleep indication bit corresponding to the thread identifier of the thread in the sleep indication parameter corresponding to the data path from the second value to the first value.

[0216] In some embodiments, the instruction execution method is applied to a graphics processor, and an arrival count cache and a sleep control cache are provided in the on-chip storage of the graphics processor; wherein the arrival count cache is used to store the arrival count value corresponding to each data path in the work group; and the sleep control cache is used to store the sleep control parameters corresponding to each data path in the work group.

[0217] In some embodiments, the instruction parallel execution module 1310 is also used to control the first thread to execute an independent instruction in response to updating the arrival count value corresponding to the one path of data; the independent instruction is an instruction before the waiting instruction and does not require the use of shared memory data; the first thread is a thread that needs to use shared memory data among the at least two threads corresponding to the one path of data.

[0218] The description of the above device embodiment is similar to the description of the above method embodiment, and has similar beneficial effects as the method embodiment. In some embodiments, the functions or modules included in the device provided in the embodiment of the present application can be used to execute the method described in the above method embodiment. For technical details not disclosed in the device embodiment of the present application, please refer to the description of the method embodiment of the present application for understanding.

[0219] It should be noted that in the embodiments of the present application, if the above-mentioned instruction execution method is implemented in the form of a software function module and sold or used as an independent product, it can also be stored in a computer-readable storage medium. Based on this understanding, the technical solution of the embodiment of the present application can be essentially or partly reflected in the form of a software product that contributes to the relevant technology. The software product is stored in a storage medium, including several instructions to enable a computer device (which can be a personal computer, a server, or a network device, etc.) to execute all or part of the methods described in each embodiment of the present application. The aforementioned storage medium includes: various media that can store program codes, such as a U disk, a mobile hard disk, a read-only memory (ROM), a disk or an optical disk. In this way, the embodiments of the present application are not limited to any specific hardware, software or firmware, or any combination of hardware, software, and firmware.

[0220] An embodiment of the present application provides a computer device, including a memory and a processor, wherein the memory stores a computer program that can be run on the processor, and when the processor executes the program, some or all of the steps in the above method are implemented.

[0221] The embodiment of the present application provides a computer-readable storage medium on which a computer program is stored, and when the computer program is executed by a processor, some or all of the steps in the above method are implemented. The computer-readable storage medium can be transient or non-transient.

[0222] An embodiment of the present application provides a computer program, including a computer-readable code. When the computer-readable code is run in a computer device, a processor in the computer device executes some or all of the steps for implementing the above method.

[0223] The embodiment of the present application provides a computer program product, which includes a non-transitory computer-readable storage medium storing a computer program, and when the computer program is read and executed by a computer, some or all of the steps in the above method are implemented. The computer program product can be implemented in hardware, software, or a combination thereof. In some embodiments, the computer program product is specifically embodied as a computer storage medium, and in other embodiments, the computer program product is specifically embodied as a software product, such as a software development kit (SDK), etc.

[0224] It should be noted here that the description of the various embodiments above tends to emphasize the differences between the various embodiments, and the same or similar aspects can be referenced to each other. The description of the above device, storage medium, computer program and computer program product embodiments is similar to the description of the above method embodiment, and has similar beneficial effects as the method embodiment. For technical details not disclosed in the embodiments of the device, storage medium, computer program and computer program product of this application, please refer to the description of the method embodiment of this application for understanding.

[0225] Fig.14 A hardware entity diagram of a computer device provided in an embodiment of the present application is shown in FIG. Fig.14 As shown, the hardware entity of the computer device 1400 includes: a processor 1401 and a memory 1402, wherein the memory 1402 stores a computer program that can be run on the processor 1401, and the processor 1401 implements the steps in the method of any of the above embodiments when executing the program.

[0226] The memory 1402 stores a computer program that can be run on the processor. The memory 1402 is configured to store instructions and applications executable by the processor 1401. It can also cache data to be processed or processed by the processor 1401 and various modules in the computer device 1400 (for example, image data, audio data, voice communication data, and video communication data). This can be achieved through flash memory (FLASH) or random access memory (Random Access Memory, RAM).

[0227] When the processor 1401 executes the program, the steps of any of the above instruction execution methods are implemented. The processor 1401 generally controls the overall operation of the computer device 1400.

[0228] An embodiment of the present application provides a computer storage medium, which stores one or more programs. The one or more programs can be executed by one or more processors to implement the steps of the instruction execution method of any of the above embodiments.

[0229] It should be noted here that the description of the above storage medium and device embodiments is similar to the description of the above method embodiments, and has similar beneficial effects as the method embodiments. For technical details not disclosed in the storage medium and device embodiments of this application, please refer to the description of the method embodiments of this application for understanding.

[0230] The processor may be at least one of an Application Specific Integrated Circuit (ASIC), a Digital Signal Processor (DSP), a Digital Signal Processing Device (DSPD), a Programmable Logic Device (PLD), a Field Programmable Gate Array (FPGA), a Central Processing Unit (CPU), a controller, a microcontroller, and a microprocessor. It is understandable that the electronic device that implements the functions of the processor may also be other, and the embodiments of the present application do not specifically limit this.

[0231] The above-mentioned computer storage medium / memory can be a read-only memory (ROM), a programmable read-only memory (PROM), an erasable programmable read-only memory (EPROM), an electrically erasable programmable read-only memory (EEPROM), a magnetic random access memory (FRAM), a flash memory (Flash Memory), a magnetic surface memory, an optical disk, or a compact disc read-only memory (CD-ROM) and the like; it can also be various terminals including one or any combination of the above-mentioned memories, such as a mobile phone, a computer, a tablet device, a personal digital assistant, etc.

[0232] It should be understood that "one embodiment" or "an embodiment" mentioned throughout the specification means that specific features, structures or characteristics related to the embodiment are included in at least one embodiment of the present application. Therefore, "in one embodiment" or "in an embodiment" appearing throughout the specification does not necessarily refer to the same embodiment. In addition, these specific features, structures or characteristics can be combined in one or more embodiments in any suitable manner. It should be understood that in various embodiments of the present application, the size of the serial number of each step / process mentioned above does not mean the order of execution, and the execution order of each step / process should be determined by its function and internal logic, and should not constitute any limitation on the implementation process of the embodiment of the present application. The serial numbers of the embodiments of the present application mentioned above are for description only and do not represent the advantages and disadvantages of the embodiments.

[0233] It should be noted that, in this article, the terms "include", "comprises" or any other variations thereof are intended to cover non-exclusive inclusion, so that a process, method, article or device including a series of elements includes not only those elements, but also other elements not explicitly listed, or also includes elements inherent to such process, method, article or device. In the absence of further restrictions, an element defined by the sentence "comprises a ..." does not exclude the existence of other identical elements in the process, method, article or device including the element.

[0234] In the several embodiments provided in the present application, it should be understood that the disclosed devices and methods can be implemented in other ways. The device embodiments described above are only schematic. For example, the division of the units is only a logical function division. There may be other division methods in actual implementation, such as: multiple units or components can be combined, or can be integrated into another system, or some features can be ignored or not executed. In addition, the coupling, direct coupling, or communication connection between the components shown or discussed can be through some interfaces, and the indirect coupling or communication connection of the devices or units can be electrical, mechanical or other forms.

[0235] The units described above as separate components may or may not be physically separated, and the components displayed as units may or may not be physical units; they may be located in one place or distributed on multiple network units; some or all of the units may be selected according to actual needs to achieve the purpose of the present embodiment.

[0236] In addition, all functional units in the embodiments of the present application can be integrated into one processing unit, or each unit can be a separate unit, or two or more units can be integrated into one unit; the above integrated unit can be implemented in the form of hardware or in the form of hardware plus software functional units. A person of ordinary skill in the art can understand that all or part of the steps of the above method embodiments can be completed by hardware related to program instructions, and the above program can be stored in a computer-readable storage medium. When the program is executed, the steps of the above method embodiments are executed; and the above storage medium includes: mobile storage devices, read-only memory (ROM), disks or optical disks, etc. Various media that can store program codes.

[0237] Alternatively, if the above-mentioned integrated unit of the present application is implemented in the form of a software function module and sold or used as an independent product, it can also be stored in a computer-readable storage medium. Based on this understanding, the technical solution of the present application can essentially or in other words, the part that contributes to the relevant technology can be embodied in the form of a software product, which is stored in a storage medium and includes a number of instructions for a computer device (which can be a personal computer, server, or network device, etc.) to execute all or part of the methods described in each embodiment of the present application. The aforementioned storage medium includes: various media that can store program codes, such as mobile storage devices, ROMs, magnetic disks, or optical disks.

[0238] The above is only an implementation method of the present application, but the protection scope of the present application is not limited thereto. Any technician familiar with the technical field can easily think of changes or substitutions within the technical scope disclosed in the present application, which should be included in the protection scope of the present application.

Claims

1. A method for executing an instruction, characterized in that: The method comprises: Process one or more data channels in parallel through workgroups; The working group includes multiple threads, multiple data paths are independent of each other, and threads within one data path are correlated; the thread corresponding to one data path realizes sleep and / or synchronous awakening of threads in sleep state through corresponding control parameters.

2. The instruction execution method according to claim 1, characterized in that: The control parameter includes an arrival count value corresponding to the one-way data; the arrival count value is used to represent the number of threads that have not yet executed to the waiting point among the at least two threads corresponding to the one-way data.

3. The instruction execution method according to claim 2, characterized in that: The method further comprises: In response to the thread corresponding to the one-way data executing to the waiting point, the arrival count value corresponding to the one-way data is updated.

4. The instruction execution method according to claim 2, characterized in that: The method further comprises: In response to the arrival count value indicating that there is no thread among the at least two threads that has not yet executed to the waiting point, a thread in a dormant state among the at least two threads is awakened.

5. The instruction execution method according to claim 1, characterized in that: The control parameter includes a sleep indication parameter; the sleep indication parameter is used to indicate a thread in a sleep state among the at least two threads corresponding to the one channel of data.

6. The instruction execution method according to claim 5, characterized in that: The method further comprises: In response to the thread corresponding to the one channel of data entering a sleep state, the sleep indication parameter is updated.

7. The instruction execution method according to claim 5, characterized in that: The method further comprises: In response to reaching a wake-up condition, based on the sleep indication parameter, waking up the thread in the sleep state corresponding to the data channel; The wake-up condition includes that there is no thread among the at least two threads that has not yet executed to the waiting point.

8. The instruction execution method according to any one of claims 1 to 7, characterized in that: The control parameters include an arrival count value and a sleep control parameter; the method includes: After one or more threads corresponding to the one channel of data execute an arrival instruction, updating the arrival count value; When the arrival count value indicates that there is no thread that has not yet executed the arrival instruction among the at least two threads corresponding to the one path of data, the thread in the dormant state is awakened based on the dormancy control parameter.

9. The instruction execution method according to claim 8, characterized in that: The method further comprises: For the first thread that executes the reach instruction, in response to the first thread executing the wait instruction, the first thread is controlled to enter a sleep state and a sleep indication parameter is updated; the first thread is a thread that needs to use shared memory data among the at least two threads corresponding to the one path of data.

10. The instruction execution method according to claim 8, characterized in that: The method further comprises: For the second thread that does not need to use the shared memory data after executing to the arrival instruction, after the second thread executes to the arrival instruction, the subsequent instructions continue to be executed.

11. The instruction execution method according to claim 8, characterized in that: The method further comprises: The initialization instruction corresponding to one path of data is executed by any thread in the working group; the initialization instruction is used to set the initial value of the arrival count value corresponding to one path of data to the total number of threads corresponding to one path of data.

12. The instruction execution method according to claim 11, characterized in that: The method further comprises: In response to all threads in the work group executing the initialization barrier instruction, each thread in the work group is controlled to continue executing subsequent instructions.

13. The instruction execution method according to claim 8, characterized in that: After one or more threads corresponding to the one path of data execute the arrival instruction, updating the arrival count value includes: In response to each thread corresponding to the one-way data executing the arrival instruction, the arrival count value corresponding to the one-way data is reduced by one.

14. The instruction execution method according to claim 9, characterized in that: The updating of the sleep indication parameter includes: Based on the thread identifier of the first thread, a sleep indication bit corresponding to the thread identifier of the first thread in the sleep indication parameter corresponding to the one channel of data is updated from a first value to a second value.

15. The instruction execution method according to claim 14, characterized in that: The step of waking up a thread in a dormant state based on a dormancy control parameter includes: Determine a thread identifier of a thread in a dormant state based on the dormancy indication bit in the dormancy control parameter that is at the second value; The thread in the dormant state is awakened by using the thread identifier of the thread in the dormant state.

16. The instruction execution method according to claim 15, characterized in that: The method further comprises: In response to a thread in a dormant state being awakened, based on a thread identifier of the thread, a dormancy indication bit corresponding to the thread identifier of the thread in a dormancy indication parameter corresponding to the one channel of data is updated from a second value to a first value.

17. The instruction execution method according to claim 8, characterized in that: The method is applied to a graphics processor, wherein an arrival count cache and a sleep control cache are provided in the on-chip storage of the graphics processor; wherein the arrival count cache is used to store an arrival count value corresponding to each data path in the work group; and the sleep control cache is used to store a sleep control parameter corresponding to each data path in the work group.

18. The instruction execution method according to claim 8, characterized in that: The method further comprises: In response to updating the arrival count value corresponding to the one path of data, the first thread is controlled to execute an independent instruction; the independent instruction is an instruction before the waiting instruction and does not require the use of shared memory data; the first thread is a thread that needs to use shared memory data among the at least two threads corresponding to the one path of data.

19. An instruction execution device, characterized in that: The instruction execution device comprises: An instruction parallel execution module is used to process one or more data in parallel through a working group; The working group includes multiple threads, multiple data paths are independent of each other, and threads within one data path are correlated; the thread corresponding to one data path realizes sleep and / or synchronous awakening of threads in sleep state through corresponding control parameters.

20. A computer device comprising a memory and a processor, wherein the memory stores a computer program executable on the processor, wherein: When the processor executes the program, the steps in the instruction execution method according to any one of claims 1 to 18 are implemented.

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

Images