Thread bundle scheduling method and apparatus, graphics processor, and computer device

Abstract

Embodiments of the present disclosure provide a thread bundle scheduling method and device, a graphics processor and a computer device. The method is applied to a graphics processor and comprises: obtaining current hardware resource usage information of the graphics processor; determining a target thread bundle from a thread bundle queue based on the hardware resource usage information and a task type of each thread bundle in the thread bundle queue; the task type of each thread bundle corresponds to hardware resource requirements of the thread bundle; and scheduling and executing the target thread bundle. In this way, by comprehensively considering the hardware resource requirements of the thread bundle and the current hardware resource usage of the graphics processor, the hardware resources in the graphics processor can be more fully and reasonably utilized, the situation of delayed execution of the scheduled target thread bundle due to waiting for resource release can be reduced, and the processing performance of the graphics processor and the overall execution efficiency of each thread bundle can be improved.

Description

TECHNICAL FIELD

[0001] The present disclosure relates to, but is not limited to, the technical field of computer, and particularly relates to a thread bundle scheduling method and device, a graphics processing unit and a computer device. BACKGROUND

[0002] A graphics processing unit (GPU) is widely used in a multi-thread task processing scenario, and is executed in a thread bundle (Warp or Wave) granularity, and each thread bundle usually contains a plurality of threads (such as 32, 64, etc.). When there are many parallel thread bundles, the GPU needs to determine which group of thread bundles can be executed in a certain scheduling manner. However, the thread bundle scheduling scheme in the related art is difficult to reasonably utilize the hardware resources in the GPU, and affects the processing performance of the GPU. SUMMARY

[0003] Therefore, the embodiments of the present disclosure provide at least a thread bundle scheduling method and device, a graphics processing unit and a computer device.

[0004] The technical scheme of the embodiments of the present disclosure is implemented as follows:

[0005] The embodiments of the present disclosure provide a thread bundle scheduling method applied to a graphics processing unit, and the method comprises the following steps:

[0006] obtaining current hardware resource usage information of the graphics processing unit;

[0007] determining a target thread bundle from a thread bundle queue based on the hardware resource usage information and a task type of each thread bundle in the thread bundle queue; the task type of each thread bundle corresponds to a hardware resource requirement of the thread bundle;

[0008] scheduling and executing the target thread bundle.

[0009] The embodiments of the present disclosure provide a thread bundle scheduling device applied to a graphics processing unit, and the device comprises the following steps:

[0010] a resource monitoring module configured to obtain current hardware resource usage information of the graphics processing unit;

[0011] a scheduling execution module configured to determine a target thread bundle from a thread bundle queue based on the hardware resource usage information and a task type of each thread bundle in the thread bundle queue; the task type of each thread bundle corresponds to a hardware resource requirement of the thread bundle; and schedule and execute the target thread bundle.

[0012] The embodiments of the present disclosure provide a graphics processing unit comprising the thread bundle scheduling device in the above embodiments.

[0013] The embodiment of the present disclosure provides a computer device, comprising the graphic processor in the above embodiment.

[0014] In the embodiment of the present disclosure, the current hardware resource usage information of the graphic processor is acquired; the target thread bundle is determined from the thread bundle queue based on the hardware resource usage information and the task type of each thread bundle in the thread bundle queue to be scheduled; the task type of each thread bundle corresponds to the hardware resource requirement of the thread bundle; and the target thread bundle is scheduled and executed. In this way, by comprehensively considering the hardware resource requirement of the thread bundle and the current hardware resource usage of the graphic processor, the hardware resources in the graphic processor can be more fully and reasonably utilized, the situation that the target thread bundle to be scheduled is delayed due to waiting for resource release can be reduced, and the processing performance of the graphic processor and the overall execution efficiency of each thread bundle can be improved.

[0015] It should be understood that the above general description and the following detailed description are only exemplary and explanatory, but not limiting the technical solutions of the present disclosure. BRIEF DESCRIPTION OF DRAWINGS

[0016] The accompanying drawings, which are incorporated into and form part of the specification, illustrate embodiments consistent with the present disclosure and, together with the specification, serve to explain the technical solutions of the present disclosure.

[0017] Figure 1 An implementation flowchart of a thread bundle scheduling method provided by the embodiment of the present disclosure is shown in the figure.

[0018] Figure 2 The component structure of a thread bundle scheduling device provided by the embodiment of the present disclosure is shown in the figure. Figure 1 ;

[0019] Figure 3 The component structure of a thread bundle scheduling device provided by the embodiment of the present disclosure is shown in the figure. Figure 1 ;

[0020] Figure 4 The component structure of a graphic processor provided by the embodiment of the present disclosure is shown in the figure.

[0021] Figure 5 The component structure of a computer device provided by the embodiment of the present disclosure is shown in the figure.

[0022] Figure 6 The implementation architecture of a thread bundle scheduling method provided by the embodiment of the present disclosure is shown in the figure.

[0023] Figure 7 The task execution timing comparison diagram of a thread bundle scheduling method provided by the embodiment of the present disclosure and the polling scheduling scheme in the related art is shown in the figure. DETAILED DESCRIPTION

[0024] To make the objectives, technical solutions, and advantages of this disclosure clearer, the technical solutions of this disclosure are further described in detail below with reference to the accompanying drawings and embodiments. The described embodiments should not be regarded as limitations on this disclosure. All other embodiments obtained by those skilled in the art without creative effort are within the scope of protection of this disclosure.

[0025] In the following description, references are made to “some embodiments,” which describe a subset of all possible embodiments. However, 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.

[0026] The terms “first / second / third” are used merely to distinguish similar objects and do not represent a specific ordering of objects. It is understood that “first / second / third” may be interchanged in a specific order or sequence where permitted, so that the embodiments of this disclosure described herein can be implemented in an order other than that illustrated or described herein.

[0027] Unless otherwise defined, all technical and scientific terms used herein have the same meaning as commonly understood by one of ordinary skill in the art to which this disclosure pertains. The terminology used herein is for descriptive purposes only and is not intended to limit the scope of this disclosure.

[0028] This disclosure provides a thread beam scheduling method. This method can be executed by a graphics processor. Figure 1 This is a schematic diagram illustrating the implementation flow of a thread beam scheduling method provided in an embodiment of this disclosure, as shown below. Figure 1 As shown, the method may include the following steps S101 to S103:

[0029] Step S101: Obtain the current hardware resource usage information of the graphics processor.

[0030] Here, hardware resource usage information may include usage information of any suitable hardware resources in the graphics processor. These hardware resources may include, but are not limited to, at least one of memory access resources, computing resources, and communication resources.

[0031] Hardware resource usage information may include, but is not limited to, at least one of the following: hardware resource utilization rate, hardware resource idle rate, hardware resource idle amount, hardware resource busy level, and hardware resource idle level.

[0032] Step S102: Based on the hardware resource usage information and the task type of each thread bundle in the queue of thread bundles to be scheduled, determine the target thread bundle from the thread bundle queue; the task type of each thread bundle corresponds to the hardware resource requirements of the thread bundle.

[0033] Here, the thread bundle queue can include a plurality of thread bundles currently to be scheduled in the graphics processor.

[0034] Each thread bundle in the thread bundle queue has a corresponding task type, which can reflect the hardware resource requirement for executing the thread bundle.

[0035] In implementation, the person skilled in the art can determine the task type corresponding to the hardware resource requirement for the thread bundle in any appropriate manner according to actual conditions, and the embodiments of the present disclosure are not limited thereto.

[0036] For example, in the case that the thread bundle requires more memory access resources for execution, the task type of the thread bundle can be memory intensive, that is, the thread bundle executes a memory intensive task.

[0037] For another example, in the case that the thread bundle requires more computation resources for execution, the task type of the thread bundle can be computation intensive, that is, the thread bundle executes a computation intensive task.

[0038] For another example, in the case that the thread bundle requires more communication resources for execution, the task type of the thread bundle can be communication type, that is, the thread bundle executes a communication type task.

[0039] In some embodiments, the task type of each thread bundle can be manually set by the developer according to the hardware resource requirement of the thread bundle. For example, the task type of each thread bundle corresponding to the task execution instruction can be set in the software interface or function for generating the task execution instruction when triggering the task execution instruction.

[0040] In some embodiments, the task type of each thread bundle can be determined by the graphics processor after analyzing the hardware resource requirement of each thread bundle.

[0041] The target thread bundle can be determined from the thread bundle queue after comprehensively considering the hardware resource requirement of the thread bundle and the current hardware resource usage of the graphics processor.

[0042] Step S103, scheduling to execute the target thread bundle.

[0043] Here, after determining the target thread bundle, the target thread bundle can be scheduled to allocate corresponding hardware resources to execute the target thread bundle.

[0044] In the embodiments of the present disclosure, current hardware resource usage information of the graphics processor is acquired; a target thread bundle is determined from a thread bundle queue based on the hardware resource usage information and a task type of each thread bundle in the thread bundle queue, the task type of each thread bundle corresponding to hardware resource requirement of the thread bundle; and the target thread bundle is scheduled for execution. In this way, by comprehensively considering the hardware resource requirement of the thread bundle and the current hardware resource usage of the graphics processor, the thread bundle is scheduled, the hardware resources in the graphics processor can be more fully and reasonably utilized, the situation that the target thread bundle to be scheduled is delayed for execution due to waiting for resource release is reduced, and thus the processing performance of the graphics processor and the execution efficiency of the overall thread bundle can be improved. In some embodiments, the task type is acquired through a software interface, the current hardware resource usage information of the graphics processor is acquired through hardware monitoring, and the target thread bundle is scheduled based on the task type acquired through the software interface and the current hardware resource usage information acquired through the hardware monitoring, so that more reasonable scheduling is achieved.

[0045] In some embodiments, the step S102 can include the following steps S111 and S112:

[0046] In step S111, a target task type to be preferentially scheduled is determined from a plurality of task types based on the hardware resource usage information.

[0047] Here, based on the hardware resource usage information of the graphics processor, the hardware resources that are relatively idle in the graphics processor can be determined, and thus the target task type to be preferentially scheduled can be determined from the plurality of task types according to the hardware resources that are relatively idle.

[0048] In some embodiments, in a case where the hardware resource usage information includes utilization rates of a plurality of hardware resources, the target task type to be preferentially scheduled can be determined based on at least one hardware resource whose utilization rate is lower than a first threshold value. The first threshold value can be determined by a person skilled in the art according to actual conditions, and the embodiments of the present disclosure are not limited thereto.

[0049] In some embodiments, in a case where the hardware resource usage information includes idle amounts of a plurality of hardware resources, the target task type to be preferentially scheduled can be determined based on at least one hardware resource whose idle amount is higher than a second threshold value. The second threshold value can be determined by a person skilled in the art according to actual conditions, and the embodiments of the present disclosure are not limited thereto.

[0050] In step S112, a target thread bundle is determined from at least one candidate thread bundle in the thread bundle queue, the task type of the candidate thread bundle being the target task type.

[0051] Here, the manner of determining the target thread bundle from the at least one candidate thread bundle can be determined according to actual conditions, and embodiments of the present disclosure do not limit this.

[0052] In some embodiments, the candidate thread bundle that is ranked first in the thread bundle queue can be determined as the target thread bundle in the order of the candidate thread bundles in the thread bundle queue.

[0053] In some embodiments, the candidate thread bundle that consumes the least hardware resources among the candidate thread bundles can be determined as the target thread bundle.

[0054] In the above embodiments, the target task type to be preferentially scheduled is determined from the plurality of task types based on the hardware resource usage information, and the target thread bundle is determined from the at least one candidate thread bundle in the thread bundle queue, and the task type of the candidate thread bundle is the target task type. In this way, the target task type to be preferentially scheduled can be determined simply and quickly.

[0055] In some embodiments, the hardware resource usage information includes utilization rates of a plurality of hardware resources.

[0056] The above step S111 can include the following step S121 and step S122:

[0057] Step S121, determining a target hardware resource with the lowest utilization rate among the plurality of hardware resources.

[0058] Step S122, determining a target task type to be preferentially scheduled from the plurality of task types based on the target hardware resource, the target task type corresponding to a hardware resource requirement matching the target hardware resource.

[0059] It can be understood that the target hardware resource is the hardware resource with the lowest utilization rate among the plurality of hardware resources, i.e., the target hardware resource is the hardware resource with the highest idle degree among the plurality of hardware resources, so that based on the target hardware resource, a task type with a hardware resource requirement matching the target hardware resource can be selected as the target task type to be preferentially scheduled from the plurality of task types.

[0060] In this way, the determined target thread bundle can be more adapted to the current hardware resource usage of the graphics processor, so that the hardware resources in the graphics processor can be more reasonably utilized, and the processing performance of the graphics processor and the execution efficiency of the overall thread bundle can be further improved.

[0061] In some embodiments, the plurality of hardware resources include memory access resources, computing resources, and communication resources, and the plurality of task types include memory-intensive, computation-intensive, and communication types.

[0062] The step S122 can include at least one of the following steps S131-S133.

[0063] In a case where the target hardware resource is the memory access resource, the step S131 determines that the target task type to be preferentially scheduled is a memory access intensive type.

[0064] In this way, since the target hardware resource is the memory access resource, it indicates that the utilization of the current memory access resource in the graphics processor is the lowest. Therefore, by determining that the target task type to be preferentially scheduled is the memory access intensive type, the graphics processor can preferentially schedule the thread bundle of the memory access intensive type which has a relatively large demand for the memory access resource, thereby improving the utilization of the memory access resource. Moreover, since the thread bundle of the memory access intensive type has a relatively small demand for other hardware resources except the memory access resource, the situation that the target thread bundle to be scheduled is delayed due to waiting for the release of other hardware resources except the memory access resource can be reduced, thereby further improving the processing performance of the graphics processor.

[0065] In a case where the target hardware resource is the operation resource, the step S132 determines that the target task type to be preferentially scheduled is a computation intensive type.

[0066] In this way, since the target hardware resource is the operation resource, it indicates that the utilization of the current operation resource in the graphics processor is the lowest. Therefore, by determining that the target task type to be preferentially scheduled is the computation intensive type, the graphics processor can preferentially schedule the thread bundle of the computation intensive type which has a relatively large demand for the operation resource, thereby improving the utilization of the operation resource. Moreover, since the thread bundle of the computation intensive type has a relatively small demand for other hardware resources except the operation resource, the situation that the target thread bundle to be scheduled is delayed due to waiting for the release of other hardware resources except the operation resource can be reduced, thereby further improving the processing performance of the graphics processor.

[0067] In a case where the target hardware resource is the communication resource, the step S133 determines that the target task type to be preferentially scheduled is a communication type.

[0068] In this way, since the target hardware resource is the communication resource, it indicates that the utilization of the current communication resource in the graphics processor is the lowest. Therefore, by determining that the target task type to be preferentially scheduled is the communication type, the graphics processor can preferentially schedule the thread bundle of the communication type which has a relatively large demand for the communication resource, thereby improving the utilization of the communication resource. Moreover, since the thread bundle of the communication type has a relatively small demand for other hardware resources except the communication resource, the situation that the target thread bundle to be scheduled is delayed due to waiting for the release of other hardware resources except the communication resource can be reduced, thereby further improving the processing performance of the graphics processor.

[0069] In some embodiments, the step S101 can include at least one of the following steps S141-S143.

[0070] The step S141, in a case where the plurality of hardware resources include a memory access resource, determining a utilization rate of the memory access resource based on a maximum access bandwidth supported by the graphics processor and a current access bandwidth usage of the graphics processor.

[0071] Here, the maximum access bandwidth refers to a maximum access bandwidth supported by a memory in the graphics processor, where the memory can be a main memory and / or other storage module in the graphics processor, which is not limited in the embodiments of the present disclosure.

[0072] The maximum access bandwidth represents a maximum number of access request bytes that the memory can receive in a single clock cycle. For example, the maximum access bandwidth can be a theoretical bandwidth of the memory, which can be an inherent hardware parameter of the memory.

[0073] The current access bandwidth usage of the graphics processor can refer to a number of access request bytes actually received by the memory in a current clock cycle, or can refer to an average number of access request bytes received by the memory in each clock cycle from a plurality of clock cycles before the current clock cycle, which is not limited in the embodiments of the present disclosure.

[0074] In some embodiments, a ratio between the current access bandwidth usage of the graphics processor and the maximum access bandwidth supported by the graphics processor can be determined as the utilization rate of the current memory access resource of the graphics processor.

[0075] In this way, the utilization rate of the current memory access resource of the graphics processor can be quickly and accurately determined.

[0076] The step S142, in a case where the plurality of hardware resources include an operation resource, determining a utilization rate of the operation resource based on a maximum operation amount supported by the graphics processor in a single clock cycle and an actual operation amount of the graphics processor in a current clock cycle.

[0077] Here, the maximum operation amount supported by the graphics processor in a single clock cycle refers to a maximum operation amount (i.e., operation times) that can be provided by an operation unit in the graphics processor in a single clock cycle.

[0078] In some embodiments, the maximum operation amount can be a hardware design parameter of the graphics processor.

[0079] The actual operation amount of the graphics processor in the current clock cycle can refer to the number of operations actually performed by the operation unit in the graphics processor in the current clock cycle, or can refer to the average number of operations performed by the operation unit in the graphics processor in each clock cycle from the current clock cycle to the previous clock cycle, and the embodiments of the present disclosure are not limited thereto.

[0080] In some embodiments, the ratio between the actual operation amount of the graphics processor in the current clock cycle and the maximum operation amount supported by the graphics processor in a single clock cycle can be determined as the utilization of the current operation resource of the graphics processor.

[0081] In this way, the utilization of the current operation resource of the graphics processor can be quickly and accurately determined.

[0082] Step S143, in the case where the plurality of hardware resources includes a communication resource, determining the utilization of the communication resource based on the maximum communication bandwidth supported by the graphics processor and the current communication bandwidth usage of the graphics processor.

[0083] Here, the maximum communication bandwidth refers to the maximum communication bandwidth supported by the communication unit in the graphics processor, wherein the communication unit can include but is not limited to at least one of a Peripheral Component Interconnect Express (PCIE) unit, a communication unit based on NVLink (a bus and communication protocol), etc., and the embodiments of the present disclosure are not limited thereto.

[0084] The maximum communication bandwidth represents the maximum number of communication request bytes that the communication unit in the graphics processor can receive in a single clock cycle. For example, the maximum communication bandwidth can be the theoretical bandwidth of the communication unit, which can be a hardware design parameter of the communication unit in the graphics processor.

[0085] The current communication bandwidth usage of the graphics processor can refer to the number of communication request bytes actually received by the communication unit in the current clock cycle, or can refer to the average number of communication request bytes received by the communication unit in each clock cycle from the current clock cycle to the previous clock cycle, and the embodiments of the present disclosure are not limited thereto.

[0086] In some embodiments, the communication resource in the graphics processor is used to support communication between the graphics processor and other processors. The other processors can include but are not limited to at least one of other graphics processors, other Central Processing Units (CPUs), etc.

[0087] In some embodiments, the ratio between the current communication bandwidth usage of the graphics processor and the maximum communication bandwidth supported by the graphics processor can be determined as the utilization of the current communication resource of the graphics processor.

[0088] In this way, the utilization of the current operation resource of the graphics processor can be determined quickly and accurately.

[0089] In some embodiments, the above method can further include steps S151 and S152.

[0090] In step S151, a task execution instruction is received, the task execution instruction including task information of a task to be executed, the task information including a task execution parameter and a preset task type.

[0091] Here, the task execution instruction can be generated and sent to the graphics processor by the host processor (such as a CPU or other GPU, etc.) in the process of executing a task processing program by calling an interface or function in the program software.

[0092] In implementation, the host processor can generate the task execution instruction by calling any suitable interface or function and send the task execution instruction to the graphics processor, and the embodiments of the present disclosure are not limited thereto. For example, the task execution instruction can be generated and sent to the graphics processor by calling a kernel function (GPU kernel) of the graphics processor under a unified computing device architecture (Compute Unified Device Architecture, CUDA).

[0093] The task execution parameter can include any suitable parameter required for executing the task, which can be reasonably set by a person skilled in the art according to the actual executed task.

[0094] The preset task type is used to represent the task type of the task to be executed, and the preset task type corresponds to the hardware resource requirement of the task to be executed.

[0095] In some embodiments, the preset task type of the task to be executed can be manually set by a developer according to the hardware resource requirement of the task to be executed. For example, the task type of the task to be executed can be set in a software interface or function for generating the task execution instruction corresponding to the task to be executed when the software interface or function is written.

[0096] In some embodiments, the task type of the task to be executed can be determined by analyzing the hardware resource requirement of each thread bundle.

[0097] In some embodiments, the preset task type can be represented by a preset type value. Among them, a corresponding type value can be set for each of a plurality of task types in advance. For example, "0" can be used to represent compute-intensive, "1" can be used to represent memory-intensive, and "2" can be used to represent communication type.

[0098] In step S152, at least one new thread bundle is created based on the task execution parameter, and the at least one new thread bundle is added to the thread bundle queue; the task type of each new thread bundle is the preset task type.

[0099] Here, at least one new thread bundle for executing the to-be-executed task can be created based on the task execution parameter of the to-be-executed task, so as to divide the task content of the to-be-executed task into each new thread bundle for execution. In addition, the preset task type of the to-be-executed task can be used as the task type of each new thread bundle.

[0100] In practice, those skilled in the art can create at least one new thread bundle based on the task execution parameter in any suitable manner according to actual conditions, and the embodiments of the present disclosure are not limited thereto.

[0101] In the above embodiments, the task execution instruction is received, the task information of the to-be-executed task is included in the task execution instruction, the task information includes the task execution parameter and the preset task type; at least one new thread bundle is created based on the task execution parameter, and the at least one new thread bundle is added to the thread bundle queue; the task type of each new thread bundle is the preset task type. In this way, the preset task type of the to-be-executed task can be carried in the task execution instruction, so as to specify the corresponding task type for at least one new thread bundle corresponding to the to-be-executed task, so that the task types of the thread bundles in the thread bundle queue can be more accurate and flexible to set.

[0102] In some embodiments, the above step S151 can include the following step S161:

[0103] In step S161, the task execution instruction sent by the main processor is received; the task execution instruction is generated by the main processor based on the task execution parameter and the preset task type, and the preset task type is determined by the main processor based on the hardware resource requirement of the to-be-executed task.

[0104] Here, the main processor can determine the preset task type based on the hardware resource requirement of the to-be-executed task, and generate the task execution instruction based on the task execution parameter of the to-be-executed task and the preset task type, and then issue the task execution instruction to the graphics processor.

[0105] In some embodiments, the host processor can generate a task execution instruction based on the task execution parameter of the to-be-executed task and the preset task type by invoking a preset software interface, and send the task execution instruction to the graphics processor. The graphics processor can receive the task execution instruction sent by the host processor through the preset software interface. For example, the preset software interface can include, but is not limited to, a kernel function of the graphics processor.

[0106] In this way, the host processor can flexibly set a suitable task type for the to-be-executed task according to the hardware resource requirement of the to-be-executed task, and generate a corresponding task execution instruction, so that the corresponding task type of at least one newly created thread bundle for the to-be-executed task can be flexibly and conveniently specified by a software manner.

[0107] The step S152 of creating at least one newly created thread bundle based on the task execution parameter can include the following steps S162 and S163:

[0108] In step S162, a task type code is determined based on the preset task type.

[0109] In some embodiments, after the graphics processor receives the task execution instruction, the graphics processor can parse the task execution instruction to obtain the task execution parameter of the to-be-executed task and the preset task type included in the task execution instruction.

[0110] After obtaining the preset task type, the graphics processor can perform encoding processing on the preset task type to obtain a task type code. The graphics processor can perform encoding processing on the preset task type in any suitable manner, which is not limited in the present disclosure.

[0111] In some embodiments, different encodings can be set for each candidate task type in the set of candidate task types, and the graphics processor can obtain the encoding corresponding to the preset task type as the task type code after obtaining the preset task type. For example, in the case of a set of candidate task types including a compute-intensive type, a memory-intensive type, and a communication type, the encoding corresponding to the compute-intensive type can be preset as "00", the encoding corresponding to the memory-intensive type can be preset as "01", and the encoding corresponding to the communication type can be preset as "10".

[0112] In some embodiments, a type value corresponding to each task type can be preset, and the type value corresponding to the task type can be converted into a binary encoding to obtain the task type code corresponding to the task type. For example, the type value 0 can be used to represent the compute-intensive type, the type value 1 can be used to represent the memory-intensive type, and the type value 2 can be used to represent the communication type. The task type code corresponding to the compute-intensive type is "00", the task type code corresponding to the memory-intensive type is "01", and the task type code corresponding to the communication type is "10".

[0113] In step S163, at least one new thread bundle is created based on the task execution parameter, and the at least one new thread bundle is associated with the task type code respectively, and the task type code is used to represent that the task type of the new thread bundle is the preset task type.

[0114] Here, after the graphics processor creates at least one new thread bundle for executing the to-be-executed task based on the task execution parameter, the graphics processor can associate each new thread bundle with the task type code. In this way, the task type code associated with each thread bundle can be referred to in the process of determining and scheduling the target thread bundle from the thread bundle queue, so that the task type of each thread bundle can be quickly and intuitively determined according to the task type code.

[0115] Embodiments of the present disclosure provide a thread bundle scheduling device, which can be applied to a graphics processor. Figure 2 The composition structure of a thread bundle scheduling device provided by an embodiment of the present disclosure is shown in Figure 1 As shown in Figure 2 The thread bundle scheduling device 200 includes a resource monitoring module 210 and a scheduling execution module 220, wherein:

[0116] The resource monitoring module 210 is configured to obtain the current hardware resource usage information of the graphics processor.

[0117] The scheduling execution module 220 is configured to determine a target thread bundle from the thread bundle queue based on the hardware resource usage information and the task type of each thread bundle in the thread bundle queue to be scheduled, wherein the task type of each thread bundle corresponds to the hardware resource requirement of the thread bundle, and the target thread bundle is scheduled to be executed.

[0118] In the embodiments of the present disclosure, by comprehensively considering the hardware resource requirement of the thread bundle and the current hardware resource usage of the graphics processor, the thread bundle is scheduled, the hardware resources in the graphics processor can be more fully and reasonably utilized, the situation that the target thread bundle to be scheduled is delayed due to waiting for resource release can be reduced, and thus the processing performance of the graphics processor and the overall execution efficiency of each thread bundle can be improved.

[0119] In some embodiments, the scheduling execution module 220 is further configured to determine a target task type to be preferentially scheduled from a plurality of task types based on the hardware resource usage information, and determine the target thread bundle from at least one candidate thread bundle in the thread bundle queue, wherein the task type of the candidate thread bundle is the target task type.

[0120] In some embodiments, the hardware resource usage information comprises utilization rates of a plurality of hardware resources; the scheduling execution module 220 is further configured to: determine a target hardware resource with the lowest utilization rate among the plurality of hardware resources; and determine, based on the target hardware resource, a target task type to be preferentially scheduled from among a plurality of the task types, the target task type corresponding to hardware resource requirements matching the target hardware resource.

[0121] In some embodiments, the plurality of hardware resources comprises memory access resources, computation resources, and communication resources; the plurality of task types comprises memory-intensive, computation-intensive, and communication types; and the scheduling execution module 220 is further configured to at least one of: in a case where the target hardware resource is the memory access resources, determine that the target task type to be preferentially scheduled is the memory-intensive; in a case where the target hardware resource is the computation resources, determine that the target task type to be preferentially scheduled is the computation-intensive; and in a case where the target hardware resource is the communication resources, determine that the target task type to be preferentially scheduled is the communication type.

[0122] In some embodiments, the resource monitoring module 210 is further configured to at least one of: in a case where the plurality of hardware resources comprises memory access resources, determine the utilization rate of the memory access resources based on a maximum memory bandwidth supported by the graphics processor and a current memory bandwidth usage of the graphics processor; in a case where the plurality of hardware resources comprises computation resources, determine the utilization rate of the computation resources based on a maximum computation amount supported by the graphics processor in a single clock cycle and an actual computation amount of the graphics processor in a current clock cycle; and in a case where the plurality of hardware resources comprises communication resources, determine the utilization rate of the communication resources based on a maximum communication bandwidth supported by the graphics processor and a current communication bandwidth usage of the graphics processor.

[0123] In some embodiments, as shown in FIG. 2, the thread bundle scheduling apparatus 200 further comprises a thread bundle creation module 230 configured to: receive a task execution instruction, the task execution instruction comprising task information of a task to be executed, the task information comprising a task execution parameter and a preset task type; create at least one new thread bundle based on the task execution parameter, and add the at least one new thread bundle to the thread bundle queue; and the task type of each new thread bundle is the preset task type. Figure 3

[0124] Embodiments of the present disclosure provide a graphics processor. As shown in FIG. 4, the graphics processor 400 comprises the above-described thread bundle scheduling apparatus 200. Figure 4

[0125] Embodiments of the present disclosure provide a computer device. As shown in FIG. 5, the computer device 500 comprises the above-described thread bundle scheduling apparatus 200. Figure 5 ​​As shown, the computer device 500 includes the aforementioned graphics processor 400.

[0126] The following describes the application of the thread beam scheduling method provided in this disclosure in a real-world scenario.

[0127] GPUs are widely used in multi-threaded task processing scenarios, and they execute tasks at the thread bundle level, typically containing several threads. When there are many parallel thread bundles, the GPU needs some kind of scheduling method to determine which thread bundle can be executed. In related technologies, the scheduling method for thread bundles in GPUs is round-robin scheduling.

[0128] However, the workload (i.e. the hardware resources required for execution) of different thread bundles may be different, and the hardware resources in the GPU (such as computing units or memory access bandwidth) have an upper limit. Therefore, the polling scheduling method may cause some hardware resources to be occupied too much, so the scheduled thread bundles cannot be executed effectively because they need to wait for the hardware resources to be released.

[0129] In the thread beam scheduling method provided in this disclosure, the type of workload for each thread beam in the GPU (corresponding to the task type in the aforementioned embodiments) can be provided through a software interface, such as compute-intensive, memory-intensive, or inter-card communication type (corresponding to the communication type in the aforementioned embodiments). During the hardware scheduling of thread beams, scheduling can be performed more effectively based on the type of workload of each thread beam in the queue to be scheduled, combined with the current hardware resource utilization in the GPU (such as main memory bandwidth utilization, computing resource utilization, communication resource utilization, etc.), so that hardware resources can be better utilized.

[0130] Figure 6 This is a schematic diagram illustrating the implementation architecture of a thread beam scheduling method provided in an embodiment of this disclosure. Figure 6 As shown, this method can be implemented collaboratively by software layer 61 and hardware 62. In software layer 61, the programmer, based on their understanding of the task to be executed, can define a preset task type (KERNEL_TYPE) through a software interface when the GPU kernels are started to execute the task. This indicates the task's requirements for different hardware resources in the GPU. The preset task type (KERNEL_TYPE) can be, for example, compute-intensive, memory-intensive, or inter-card communication-intensive. Taking CUDA calling the GPU kernel (kernel function) to generate task execution instructions as an example, the software interface called can be:

[0131] gpu_kernel<<<BLOCK_NUMBER,BLOCK_SIZE,KERNEL_TYPE>>>(ARG_A,ARG_B,ARG_C);

[0132] wherein, BLOCK_NUMBER represents the number of thread blocks for executing the corresponding to-be-executed task, BLOCK_SIZE represents the size of the thread block, KERNEL_TYPE represents a preset task type, and ARG_A, ARG_B and ARG_C are all task execution parameters of the to-be-executed task. Wherein, KERNEL_TYPE can be defined as 0, indicating that the to-be-executed task is a compute-intensive task; KERNEL_TYPE is 1, indicating that the to-be-executed task is a memory-intensive task; KERNEL_TYPE is 2, indicating that the to-be-executed task is an inter-card communication type task used for inter-card communication.

[0133] Continuing to refer to Figure 6 In the hardware 62 of the GPU, according to the preset task type KERNEL_TYPE transmitted by the software, the hardware 62 can encode the preset task type KERNEL_TYPE to obtain the task type of each thread bundle corresponding to the to-be-executed task. Taking the three types of task types currently supported (i.e., compute-intensive, memory-intensive and inter-card communication type) as an example, the hardware 62 only needs to do 2-bit encoding, for example, encoding "00" represents compute-intensive, encoding "01" represents memory-intensive, and encoding "10" represents inter-card communication type. After dividing the work of the to-be-executed task corresponding to the GPU kernel into multiple thread bundles (corresponding to the newly created thread bundle in the foregoing embodiment), each thread bundle after the division is attached with the 2-bit encoding as the task type (such as task types Type0, Type1, TypeN) for the scheduler to identify; the thread bundles after the division are added to the thread bundle queue Warps_Array, and the thread bundle queue can record the task type of each thread bundle (such as thread bundles Warp_0, Warp_1, Warp_N).

[0134] The hardware 62 of the GPU can include a scheduler 621, a selector 622, a hardware monitor 623, and hardware resources for executing tasks, including an operation unit 624, a memory access unit 625 and a communication unit 626.

[0135] The hardware monitor 623 can be used to monitor the busy degree of different hardware resources in the GPU. For the memory access resource, the bandwidth of the main memory can be monitored. For example, the number of received request bytes in a fixed number of clock cycles. Then divided by the theoretically maximum number of request bytes that can be received, the bandwidth utilization of memory access can be obtained. For example, the theoretical bandwidth of the main memory is 10 bytes / cycle (this is a hardware inherent attribute, so for a certain hardware, it is a constant value), and a total of 500 bytes of request data are received in 1000 cycles, so the hardware can calculate that the bandwidth utilization of memory access is 500 / (10*1000) = 5%. Similarly, for the operation unit / communication unit, the corresponding utilization can also be calculated. For example, for the communication unit, the utilization of the communication module is the bandwidth utilization of the PCIE communication unit (here, it is not limited to only the PCIE communication unit). For the operation unit, such as the ALU, there is also an index of the operation amount (i.e., the number of operations) that can be provided per unit time. The hardware monitor 623 can count the actual workload of a certain type of hardware resource in a certain time period, and then divide the theoretical maximum workload upper limit of the hardware resource to obtain the utilization of the hardware resource.

[0136] The scheduler 621 needs to combine the utilization of each hardware resource monitored by the hardware monitor 623 to perform weighted round-robin scheduling for the execution thread bundle. For example, there are 10 thread bundles in the current thread bundle queue, and the last two thread bundles are of the compute-intensive type (the scheduler 621 can identify according to the code representing the task type). At this time, the hardware monitor 623 feedbacks that the utilization of the operation unit 624 is 50%, and the utilization of the memory access unit 625 is 90%. Therefore, the scheduler 621 can consider that the operation unit 624 is relatively idle at this time, so the last two compute-intensive thread bundles can be preferentially selected by the selector 622.

[0137] The thread bundle scheduling method provided by the embodiments of the present disclosure can more reasonably utilize the hardware resources in the GPU. Otherwise, when the utilization of the memory access unit is very high, the thread bundle with heavy memory access (i.e., the thread bundle of the task type of memory-intensive type) is still allocated, which may cause the execution of the thread bundle to stagnate, thereby affecting the overall processing performance of the GPU. Figure 7 The task execution time sequence comparison diagram of the thread bundle scheduling method provided by the embodiments of the present disclosure and the round-robin scheduling scheme in the related art is as follows: Figure 7As shown, the thread bundle corresponding to the memory access task 1, the thread bundle corresponding to the memory access task 2 and the thread bundle corresponding to the computing task 3 are scheduled in turn in the polling scheduling scheme 71, and due to the memory bandwidth limitation, the memory access task 2 can be executed only after the memory access task 1 is completed, so that the memory access task 1, the memory access task 2 and the computing task 3 finally cost 12 ms to be executed and completed; however, in the thread bundle scheduling method 72 provided by the embodiment of the present disclosure, considering the memory bandwidth limitation when the memory access task 1 is executed, after the thread bundle corresponding to the memory access task 1 is scheduled, the thread bundle corresponding to the computing task 3 is preferentially scheduled, and then the thread bundle corresponding to the memory access task 2 is scheduled, so that the thread bundle corresponding to the computing task 3 does not need to wait for the execution of the thread bundle corresponding to the memory access task 1 to be completed and can be executed, so that the memory access task 1, the memory access task 2 and the computing task 3 can be completed in 10 ms. It can be seen that the thread bundle scheduling method provided by the embodiment of the present disclosure can more reasonably utilize the hardware resources inside the GPU, and improve the overall performance of task processing.

[0138] It should be noted that the resource monitoring module in the foregoing embodiments can be implemented by the hardware monitor described above, and the scheduling execution module can be implemented by the scheduler and the selector described above.

[0139] It should be noted that the descriptions of the various embodiments above tend to emphasize the differences between the various embodiments, and the same or similar parts can be referred to each other. The descriptions of the above device, graphics processor and computer device embodiments are similar to the descriptions of the above method embodiments, and have similar beneficial effects as the method embodiments. In some embodiments, the device, graphics processor and computer device provided by the embodiments of the present disclosure have functions or include modules that can be used to execute the methods described in the above method embodiments. For technical details not disclosed in the device embodiments of the present disclosure, please refer to the descriptions of the method embodiments of the present disclosure.

[0140] It should be understood that the "one embodiment" or "an embodiment" mentioned throughout the specification means that the specific features, structures or characteristics related to the embodiment are included in at least one embodiment of the present disclosure. Therefore, "in one embodiment" or "in an embodiment" appearing throughout the specification does not necessarily mean 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 disclosure, the size of the serial number of each step / process does not mean the execution order, and the execution order of each step / process should be determined according to its function and inherent logic, and should not constitute any limitation on the implementation process of the embodiments of the present disclosure. The serial number of the above embodiments of the present disclosure is only for description, not representing the pros and cons of the embodiments.

[0141] It should be noted that, in the present document, the terms "comprising", "containing" or any other similar term are intended to encompass non-exclusive inclusions, such that a process, method, article, or apparatus that comprises a list of elements does not necessarily include those elements only, but can include other elements not expressly listed or inherent to such process, method, article, or apparatus. An element specified by the phrase "comprising a" does not exclude the presence of additional identical elements in the process, method, article, or apparatus that includes the element.

[0142] In several embodiments provided by the present disclosure, it should be understood that the disclosed graphic processor and method can be implemented in other manners. The above described device embodiments are merely schematic, for example, the division of the units is only a logical function division, and there can be another division manner in actual implementation. For example, a plurality of units or components can be combined or integrated into another system, or some features can be ignored or not executed. In addition, the displayed or discussed coupling, or direct coupling or communication connection between the components can be indirect coupling or communication connection through some interfaces, devices, or units, and can be electrical, mechanical, or in other forms.

[0143] The units described as separate components can or can not be physically separate, and the components shown as units can or can not be physical units; they can be located in one place, or distributed on a plurality of network units; and some or all of the units can be selected according to actual needs to achieve the purpose of the embodiments. In addition, each functional unit in each embodiment of the present disclosure can be integrated into a processing unit, or each unit can be a separate unit, or two or more units can be integrated into a unit; the integrated unit can be implemented in the form of hardware, or in the form of hardware plus software functional units.

[0144] Those skilled in the art can understand that all or part of the steps of the above-mentioned method embodiments can be completed by program instruction related hardware, and the foregoing program can be stored in a computer readable storage medium, and the program executes the steps of the above-mentioned method embodiments when executed; and the foregoing storage medium includes mobile storage devices, read only memory (ROM), magnetic discs or optical discs, and various storage medium that can store program codes.

[0145] The above description is merely an implementation of the present disclosure, but the protection scope of the present disclosure is not limited thereto, and any person skilled in the art can easily think of changes or replacements within the technical range disclosed by the present disclosure, which should be covered by the protection scope of the present disclosure.

Claims

1. A thread beam scheduling method, characterized in that, Applied to a graphics processor, the method includes: Obtain the current hardware resource usage information of the graphics processor; Based on the hardware resource usage information and the task type of each thread bundle in the thread bundle queue to be scheduled, a target thread bundle is determined from the thread bundle queue; the task type of each thread bundle corresponds to the hardware resource requirements of the thread bundle; each thread bundle in the thread bundle queue has a corresponding task type, and the target thread bundle is the thread bundle in the thread bundle queue to be scheduled first. Schedule the execution of the target thread bundle.

2. The thread beam scheduling method according to claim 1, characterized in that, The step of determining the target thread bundle from the thread bundle queue based on the hardware resource usage information and the task type of each thread bundle in the queue to be scheduled includes: Based on the hardware resource usage information, determine the target task type to be prioritized for scheduling from among the various task types; From at least one candidate thread bundle in the thread bundle queue, a target thread bundle is determined, wherein the task type of the candidate thread bundle is the target task type.

3. The thread beam scheduling method according to claim 2, characterized in that, The hardware resource usage information includes the utilization rate of various hardware resources; The step of determining the target task type to be prioritized for scheduling from multiple task types based on the hardware resource usage information includes: Identify the target hardware resource with the lowest utilization rate among the various hardware resources; Based on the target hardware resources, a target task type to be prioritized for scheduling is determined from a variety of task types, and the hardware resource requirements corresponding to the target task type are matched with the target hardware resources.

4. The thread beam scheduling method according to claim 3, characterized in that, The various hardware resources include memory access resources, computing resources, and communication resources; the various task types include memory-intensive, computation-intensive, and communication-intensive tasks. The step of determining the target task type to be prioritized for scheduling from a variety of task types based on the target hardware resources includes at least one of the following: When the target hardware resource is the memory access resource, the target task type to be prioritized for scheduling is determined to be memory-intensive. When the target hardware resource is the computing resource, the target task type to be prioritized for scheduling is determined to be compute-intensive. If the target hardware resource is the communication resource, the target task type to be prioritized for scheduling is determined to be the communication type.

5. The thread beam scheduling method according to claim 3, characterized in that, The step of obtaining the current hardware resource usage information of the graphics processor includes at least one of the following: When the various hardware resources include memory access resources, the utilization rate of the memory access resources is determined based on the maximum memory access bandwidth supported by the graphics processor and the current memory access bandwidth usage of the graphics processor. When the various hardware resources include computing resources, the utilization rate of the computing resources is determined based on the maximum amount of computing power supported by the graphics processor in a single clock cycle and the actual amount of computing power of the graphics processor in the current clock cycle. When the various hardware resources include communication resources, the utilization rate of the communication resources is determined based on the maximum communication bandwidth supported by the graphics processor and the current communication bandwidth usage of the graphics processor.

6. The thread beam scheduling method according to any one of claims 1 to 5, characterized in that, The method further includes: Receive a task execution instruction, the task execution instruction including task information of the task to be executed, the task information including task execution parameters and preset task type; Based on the task execution parameters, at least one new thread bundle is created and the at least one new thread bundle is added to the thread bundle queue; the task type of each new thread bundle is the preset task type.

7. The thread beam scheduling method according to claim 6, characterized in that, The task type is obtained through a software interface, and the current hardware resource usage information of the graphics processor is obtained through hardware monitoring. The target thread bundle is scheduled based on the task type obtained through the software interface and the current hardware resource usage information obtained through hardware monitoring.

8. A thread beam scheduling device, characterized in that, The device, applied to a graphics processor, includes: The resource monitoring module is used to obtain the current hardware resource usage information of the graphics processor; The scheduling and execution module is used to: determine a target thread bundle from the thread bundle queue based on the hardware resource usage information and the task type of each thread bundle in the thread bundle queue to be scheduled; the task type of each thread bundle corresponds to the hardware resource requirements of the thread bundle; schedule the execution of the target thread bundle; each thread bundle in the thread bundle queue has a corresponding task type, and the target thread bundle is the thread bundle to be scheduled first in the thread bundle queue.

9. The thread beam scheduling device according to claim 8, characterized in that, The device further includes: A thread bundle creation module is configured to: receive a task execution instruction, the task execution instruction including task information of the task to be executed, the task information including task execution parameters and a preset task type; based on the task execution parameters, create at least one new thread bundle and add the at least one new thread bundle to the thread bundle queue; the task type of each new thread bundle is the preset task type.

10. A graphics processor, characterized in that, Includes the thread beam scheduling device as described in claim 8 or 9.

11. A computer device, characterized in that, Includes the graphics processor as described in claim 10.

Images