A data scheduling method and computing device

Abstract

A data scheduling method and a computing device, the method is applied to a main processing unit, the main processing unit is in communication connection with a coprocessing unit, the coprocessing unit includes a plurality of processing cores, and the method includes: aligning target data of a target data amount according to a preset alignment boundary; determining a target level according to the aligned target data amount; determining a data splitting parameter of the target data according to the target data amount and the target level; and sending the data splitting parameter to the coprocessing unit, so that the coprocessing unit schedules the processing cores to perform calculation according to the data splitting parameter. Thus, by designing a targeted data splitting scheme to determine the data splitting parameter for target data of different sizes of target data amount, reasonable scheduling and allocation of the plurality of processing cores can be achieved, and the rate of executing target binocular operator calculation of the system is improved.

Description

Technical Field

[0001] This application relates to the field of artificial intelligence technology, and in particular to a data scheduling method and computing device. Background Technology

[0002] A neural processing unit (NPU) is a processor specifically designed for performing artificial intelligence (AI) and deep learning tasks. It aims to improve the computational efficiency and performance of AI tasks, and one of its core components is multiple AI cores. AI cores feature efficient on- and off-core parallel computing capabilities, flexibility supporting mixed precision, and a high-speed interconnect and shared memory architecture, significantly improving the processing performance of AI tasks and providing robust hardware support for the training and inference of deep learning models.

[0003] The AI ​​Core comprises core components such as computing units, storage units, and data transfer units, supporting accelerated large-scale neural network computations. The computing units include three basic computing resources: Cube computing units, Vector computing units, and Scalar computing units. The storage units are the internal storage of the AI ​​Core, collectively referred to as Local Memory. In contrast, the external storage of the AI ​​Core is called Global Memory. The Direct Memory Access (DMA) transfer unit is responsible for transferring data between Global Memory and Local Memory.

[0004] In AI applications such as digital vision image processing, task operators developed using programming languages ​​such as C / C++ and Python run in parallel on multiple AI Cores of the NPU. Because image processing (such as high-resolution image processing) involves massive amounts of computational data, these operator computations are typically data-intensive.

[0005] However, AI Core has limited memory storage resources (usually 192K Byte or 256K Byte), which means that computational data cannot be loaded, computed, and unloaded all at once. Therefore, it is necessary to design reasonable data partitioning parameters to complete operator computation tasks in blocks and efficiently. Summary of the Invention

[0006] This application provides a data scheduling method and computing device that can efficiently complete operator computation tasks.

[0007] In a first aspect, embodiments of this application provide a data scheduling method. The method is applied to a main processing unit, which is communicatively connected to a coprocessing unit. The coprocessing unit includes multiple processing cores. The method includes: aligning target data of a target data volume according to a preset alignment boundary; the target data is any input data for performing target binocular operator calculations; the target data volume is the number of target data; determining a target level based on the aligned target data volume; determining data segmentation parameters for the target data based on the target data volume and the target level; the data segmentation parameters include the target number of processing cores used, the number of processing times for each of the target number of processing cores, and the single data volume for each processing; and sending the data segmentation parameters to the coprocessing unit, so that it schedules the processing cores to perform calculations according to the data segmentation parameters.

[0008] Therefore, designing targeted data segmentation schemes and determining data segmentation parameters for target data of different sizes can enable reasonable scheduling and allocation of multiple processing cores, reduce communication requirements and overhead between multiple processing cores, improve their parallel processing performance, and thus improve the system's speed of executing target binocular operator calculations.

[0009] In one possible implementation, each processing core includes a memory space of the same size; the target level includes a first level; the lower boundary parameter of the first level is 1, and the upper boundary parameter of the first level is the product of the total number of processing cores and the first preset number of target data processed by a single processing core each time.

[0010] In one possible implementation, the target level also includes a second level; the lower boundary parameter of the second level is greater than the upper boundary parameter of the first level, and the upper boundary parameter of the second level is the product of the total number of processing cores and the second preset number of target data processed by a single processing core each time; wherein, the second preset number is at least related to the memory size used by a single processing core to perform one target binocular operator calculation using the in-place calculation method.

[0011] In one possible implementation, if the target binocular operator supports in-place computation, the second preset number is the quotient of the memory space size of a single processing core divided by the memory size used; the memory size used is the memory size used by a single processing core to perform one computation of the target binocular operator using the in-place computation method.

[0012] In one possible implementation, if the coprocessor unit also supports in-core pipelined parallelism, the second preset number is the memory space size of a single processing core, which is the quotient of the product of the memory size used and the maximum parallelism of the in-core pipelined parallelism.

[0013] Therefore, reasonable upper and lower boundary parameters are set for multiple levels based on the capabilities supported by the coprocessor unit.

[0014] In one possible implementation, the target level also includes a third level; the lower boundary parameter of the third level is greater than the upper boundary parameter of the second level.

[0015] Therefore, multiple levels are preset for target data volumes of different sizes.

[0016] In one possible implementation, for the target data corresponding to the target data volume of the first level, the target number is the value of the target data volume divided by the first preset number and rounded up; the number of processing cores for the target number is 1; the amount of data processed by each processing core that is 1 less than the target number is the first preset number, and the amount of data processed by the remaining 1 processing core in the target number is the difference between the target data volume and the sum of the amount of data processed by each processing core that is 1 less than the target number.

[0017] In one possible implementation, for the target data corresponding to the target data volume of the second level, the target number is the total number of multiple processing cores: the processing cores of the target number each process once; the data volume processed by each processing core that is one less than the target number is the quotient of the target data volume divided by the target number, and the data volume processed by the remaining one processing core of the target number is the difference between the target data volume and the sum of the data volumes processed by each processing core that is one less than the target number.

[0018] In one possible implementation, for the target data corresponding to the target data volume of the third level, the target number is the total number of multiple processing cores; the total data volume processed by each processing core that is one less than the target number is the quotient of the target data volume divided by the target number; and the total data volume processed by the remaining one processing core in the target number is the difference between the target data volume and the sum of the total data volume processed by each processing core that is one less than the target number.

[0019] In one possible implementation, for each of the target number of processing cores, the number of processing times of the processing core is the value of its corresponding total data volume divided by a second preset number and rounded up; wherein, for processing cores other than the last processing, the amount of data processed by the processing core in each processing session is the second preset number, and for the last processing session, the amount of data processed by the processing core in each processing session is the difference between its corresponding total data volume and the sum of the amounts of data processed in each processing session other than the last processing session.

[0020] Therefore, targeted data partitioning schemes are designed for target data of different sizes to guide the coprocessing unit in rationally scheduling multiple processing cores.

[0021] Secondly, embodiments of this application provide a computing device, including: a memory, a main processor, and a coprocessor;

[0022] The memory is coupled to the main processor and coprocessors.

[0023] Memory, used to store programs;

[0024] The main processor is used to call the program to execute the method described in any possible implementation of the first aspect, so that the coprocessor can perform the scheduling.

[0025] Thirdly, embodiments of this application provide a computer-readable storage medium storing computer-executable instructions, which, when executed by a computer, implement the method as described in the first aspect or any one of the first aspects.

[0026] Fourthly, embodiments of this application provide a program product including a computer program that, when executed by a processor, implements the method as described in the first aspect or any one of the first aspects.

[0027] It is understood that the beneficial effects of the second to fourth aspects mentioned above can be found in the relevant descriptions in the first aspect mentioned above, and will not be repeated here. Attached Figure Description

[0028] To more clearly illustrate the technical solutions of the embodiments of this application, the drawings used in the description of the embodiments will be briefly introduced below. Obviously, the drawings described below are only some embodiments of this application. For those skilled in the art, other drawings can be obtained based on these drawings without creative effort.

[0029] Figure 1 A flowchart of a custom operator development process provided for an embodiment of this application;

[0030] Figure 2 A flowchart of a data scheduling method provided in an embodiment of this application;

[0031] Figure 3 This is a schematic diagram of data segmentation provided in an embodiment of this application;

[0032] Figure 4 This is a schematic diagram of data segmentation provided in an embodiment of this application;

[0033] Figure 5 This is a schematic diagram of data segmentation parameters provided in an embodiment of this application;

[0034] Figure 6 This is a schematic diagram of the structure of a computing device provided in an embodiment of this application. Detailed Implementation

[0035] To make the objectives, technical solutions, and advantages of the embodiments of this application clearer, the technical solutions in the embodiments of this application will be described below with reference to the accompanying drawings.

[0036] In the description of the embodiments in this application, any embodiment or design that is “exemplary,” “for example,” or “by way of example” should not be construed as being more preferred or advantageous than other embodiments or designs. Specifically, the use of terms such as “exemplary,” “for example,” or “by way of example” is intended to present the relevant concepts in a concrete manner.

[0037] Furthermore, the terms "first" and "second" are used for descriptive purposes only and should not be construed as indicating or implying relative importance or implicitly specifying the indicated technical features. Thus, a feature defined as "first" or "second" may explicitly or implicitly include one or more of that feature. The terms "comprising," "including," "having," and variations thereof all mean "including but not limited to," unless otherwise specifically emphasized. "At least one" can refer to one or more, where multiple can refer to two or more.

[0038] In fields such as deep learning and computer vision, operators play a crucial role in neural network models, corresponding to the computational logic of layers or nodes within the network. General operators such as "tanh", "ReLU", and "Conv2D" are primarily used for analyzing and processing single-view images. In stereo vision systems, stereo operators, such as "Add", are used to process two images from different perspectives of the same scene. One image corresponds to one view in the stereo operator. By comparing and analyzing these two images, the depth information of objects in the scene is calculated. The natural two-dimensional structure of image data allows these operators to be processed in parallel using multiple AI Cores to improve computational speed.

[0039] Currently, for computing systems with multiple AI Cores, the common method for data partitioning between cores is load balancing. This distributes the computational data for operator calculations evenly across each AI Core to optimize resource utilization and improve overall performance. However, this data partitioning method cannot dynamically adjust the allocation of computational data based on the characteristics of the operators involved in the current load and task. In certain specific scenarios, such as when there are large differences in the amount of computation or the execution time of operators, simple load balancing is not effective enough.

[0040] Therefore, this scheme first determines the data partitioning parameters of the target data in the main processing unit based on the target data volume and target level. The target data refers to any input data from the target binocular operator, and the target data volume is the number of target data items. Secondly, the main processing unit transmits the data partitioning parameters to the coprocessing unit it communicates with. The coprocessing unit then schedules multiple processing cores according to these parameters to execute the target binocular operator calculation. Thus, by designing targeted data partitioning schemes for target data of different sizes, this scheme enables reasonable scheduling and allocation of multiple processing cores, reduces communication requirements and overhead between cores, improves their parallel processing performance, and ultimately enhances the system's performance in executing the target binocular operator calculation.

[0041] In addition, the boundary parameters used at the target level in this scheme are at least related to the memory size used by a single processing core to perform a target binocular operator calculation once using the in-place calculation method, thereby improving the rationality of the target level division; and, the in-place calculation method used in this scheme to perform target binocular operator calculation can improve the memory utilization efficiency of a single processing core, reduce data transfer time, and thus improve the system's speed of performing target binocular operator calculation.

[0042] For example, Figure 1 The diagram shows a flowchart of a custom operator development process provided in an embodiment of this application.

[0043] like Figure 1 As shown, in fields such as computer vision or image processing, developers complete efficient and reliable operators through operator analysis and operator creation engineering design based on specific image processing tasks.

[0044] The operator analysis step is used to analyze and evaluate the operators used in image processing algorithms. Operator analysis helps determine the operator types and parameters best suited for a specific task.

[0045] The operator creation project step is used to create a new project or program during the software development process, based on a defined operator type and parameters, specifically for developing and implementing image processing operators. This may include setting up the development environment, configuring dependencies, defining the project structure, etc.

[0046] In binocular vision systems, custom binocular operators can be developed using programming languages ​​such as C / C++ and Python. For example, "Add" is used to combine the features or data of two images for processing.

[0047] like Figure 1As shown, the host side is the environment where the main processing unit resides, and its code implementation is used to execute high-level control logic and coordinate computational tasks. The main processing unit can be a Central Processing Unit (CPU) or a Graphics Processing Unit (GPU), etc. In this scheme, the main processing unit determines the data segmentation parameters of the target data according to the target level of the target data volume, and then passes the data segmentation parameters to the core side.

[0048] The core side consists of coprocessor units responsible for executing parallel computing tasks. Their code implementation is used to perform these tasks. Corresponding coprocessor units are assigned to different types of main processing units. For example, when the main processing unit is a GPU, the coprocessor unit can be a Tensor Processing Unit (TPU) or an NPU; when the main processing unit is a CPU, the coprocessor unit can be a GPU, TPU, or NPU. These coprocessor units include multiple processing cores and perform excellently when executing highly parallelized tasks. In this scheme, the coprocessor unit schedules multiple processing cores based on the dataset partitioning parameters sent by the main processing unit, transferring data from the dataset to perform computations on the target binocular operator.

[0049] It is understandable that the main processing unit and the coprocessor unit establish a communication connection to work collaboratively, and this heterogeneous computing model can fully leverage the advantages of each. Depending on the application scenario, some systems may integrate the coprocessor unit into a system-on-chip (SoC), providing an integrated solution together with the main processing unit, memory, and I / O controllers. Alternatively, the coprocessor unit can also exist as a standalone device, for example, as a pluggable accelerator card, connected to the host system including the main processing unit via a high-speed interface (such as PCIe). This form factor allows it to be used in existing systems that require additional high-speed parallel computing capabilities. Both of these methods are communication connection methods supported by this solution.

[0050] The following section will use the Neural Processing Unit (NPU) as a coprocessor and its application in digital visual image processing as examples to provide a detailed explanation of this solution.

[0051] For example, Figure 2 The diagram illustrates a data scheduling method according to an embodiment of this application. This method can run on any computing unit, server, device, or main processing unit of a device cluster with computing and processing capabilities. These hardware devices also integrate coprocessing units or have communication connections with coprocessing units via high-speed interfaces. The coprocessing unit includes multiple processing cores.

[0052] like Figure 2 As shown, the method mainly includes the following steps:

[0053] Step S201: Align the target data according to a preset alignment boundary. The target data is any input data used to perform the target binocular operator calculation. The target data volume is the number of target data items.

[0054] In one embodiment, the required alignment boundary, such as 32 bytes or 64 bytes, is determined based on the access performance of the storage medium storing the target data.

[0055] Before performing calculations using the target data, aligning the target data volume in the storage medium according to a preset alignment boundary can reduce storage space access latency and improve data loading and storage efficiency.

[0056] For example, target data is data collected in the application domain used to perform target binocular operator calculations. A target binocular operator is a type of binocular operator, meaning that performing the calculation requires two input data streams, X and Y; one input data stream is called "one binocular". Target data is either of the two input data streams, X and Y. The two input data streams X and Y used to perform the target operator calculation have the same data type and the same target data volume.

[0057] The target data size refers to the number of target data items, that is, the number of input data items used to execute the target binocular operator. The target data size is independent of the data type. The larger the target data size, the more target data items there are, and vice versa.

[0058] It is understandable that when coprocessing units are applied in AI fields such as digital vision image processing (e.g., high-resolution image processing), the computational data involved is intensive, resulting in target data with large volumes. Conversely, when coprocessing units are applied in situations where digital vision images do not change significantly or in low-resolution image processing, the target data has small or medium volumes. Therefore, different data partitioning schemes can be designed for target data of different sizes to improve the efficiency of performing target binocular operator computations.

[0059] Step S202: Determine the target level based on the aligned target data volume.

[0060] In one embodiment, at least three levels are preset for different target data volumes. The first, second, and third levels represent target data volumes of small, medium, and large, respectively. Upper and lower boundary parameters are set for each level to determine the range of data volumes included in each level.

[0061] After acquiring the target data volume through data collection and aligning it to preset boundaries, the target level is determined based on the target data volume. It can be understood that the target level could be any of the first, second, or third levels.

[0062] In one implementation, each processing core of the coprocessing unit includes a memory space of the same size. The lower boundary parameter of the first level is set to 1, and the upper boundary parameter of the first level is the product of the total number of processing cores and the first preset number of target data processed by a single processing core each time.

[0063] For example, for the NPU coprocessor, the range of the number of target data corresponding to the first level is set to [1, totalIBlockDim * Length]. Here, totalBlockDim is the total number of AI Cores in the NPU. Length is a first preset number, which can be a constant set empirically, indicating the amount of data that can be processed in a single operation on a single AI Core, for example, the optimal number of data points that a single AI Core can process. The storage space occupied by the target data of the first preset number Length is less than the memory space of a single AI Core.

[0064] The first preset number can be adjusted according to the computational complexity of different target binocular operators. For example, the more complex the computation of the target binocular operator, the smaller the first preset number should be, and vice versa. Storage space refers to the storage space occupied by the target data in the storage medium, which is the AI ​​Core's Global Memory.

[0065] In one implementation, based on the range of the number of target data corresponding to the first level, the lower boundary parameter of the second level is set to be greater than the upper boundary parameter of the first level. The upper boundary parameter of the second level is the product of the total number of processing cores and the second preset number of target data processed by a single processing core each time. Wherein, when the target binocular operator supports in-place computation, the second preset number is at least related to the memory size used by a single processing core to perform one target binocular operator computation using the in-place computation method.

[0066] For example, for the NPU (Network Processing Unit), the range of the number of target data corresponding to the second level is set to (totalBlockDim * Length, totalBlockDim * ubLength). Here, totalBlockDim is the total number of AI Cores in the NPU. ubLength is a second preset number, indicating the maximum amount of data a single AI Core can process in a single operation, such as the maximum number of data items a single AI Core can handle. The storage space occupied by the target data of the second preset number ubLength is no greater than the memory space of a single AI Core.

[0067] Optionally, if the target binocular operator supports in-place computation, the second preset number is the quotient of the memory space size of a single processing core divided by the memory size used; the memory size used is the memory size used by a single processing core to perform one target binocular operator computation using the in-place computation method.

[0068] For example, for the NPU coprocessor, the second preset number ubLength is set to the memory space size of a single AI Core. The quotient is calculated by dividing the memory size used for one target binocular operator computation by the memory size used for a single AI Core using an in-place computation method. This can be expressed as: ubLength = ubSize / (num * sizeof(datatype)).

[0069] Here, `ubSize` represents the memory size of a single AI Core. `num` represents the number of input and output variables required by the target binocular operator, and `sizeof(datatype)` represents the memory space occupied by the data type `datatype` of the target data. For example, for the data type `float16`, `sizeof(float16) = 2`, meaning that a single `float16` data type requires 2 bytes of memory. `num * sizeof(datatype)` is the memory size used by a single AI Core to execute the target binocular operator once. Next, we need to determine the value of `num` in the target binocular operator.

[0070] Taking a typical binocular operator `Add` with the expression `z = X + y` as an example, executing this binocular operator calculation once requires two input data points, X and Y, with one input data point being the "one-way" input. The two input data points X and Y correspond to variables x and y, respectively. Therefore, it is necessary to move the data corresponding to the two input variables x and y from the AI ​​Core's Global Memory to the AI ​​Core's Local Memory. Storing these two variables in the AI ​​Core's Local Memory requires memory space for two data types. Additionally, memory space for the result variable z also needs to be allocated for a different data type. Therefore, executing a single binocular operator `Add` calculation requires memory space for three data types.

[0071] However, when developing custom target binary operators using programming languages, the algorithm of the binary operator can be designed to support in-place computation by updating the data structure in place or by using techniques such as swapping values ​​instead of copying values, for example, by using pointers and references to avoid unnecessary data copying.

[0072] Taking the ordinary binary operator Add as an example, its expression can be replaced with y+=x during programming development, resulting in a target binary operator that supports in-place computation, expressed as y+=x. As can be seen from the expression y+=x, the calculation can be performed directly on the input variable y, and the result can be stored in the original memory space of the input variable y, without needing to allocate additional memory space for the intermediate result z. Therefore, completing one calculation of the target binary operator that supports in-place computation only requires memory space for two data types, meaning the number of input and output variables required for the target binary operator to support in-place computation is num=2.

[0073] Therefore, this solution, by developing a target binocular operator that supports in-place computation and scheduling the processing core to execute the computation of the target binocular operator using the in-place computation method, can reduce the overhead of data copying and movement, thereby improving the memory utilization of a single processing core and reducing data transfer time.

[0074] Optionally, when the target binocular operator supports in-place computation, if the coprocessor unit also supports in-core pipelined parallelism, the second preset number is the quotient of the product of the memory space size of a single processing core and the maximum parallelism of the in-core pipelined parallelism function. The memory size used is the memory size used by a single AI Core to perform one target binocular operator computation using the in-place computation method.

[0075] For example, when the NPU supports in-core pipelined parallelism, the second preset number is the memory space size of a single AI Core. This is the quotient of the product of the memory size used to perform one target binocular operator computation on a single AI Core using in-place computation and the maximum parallelism of the in-core pipelined parallelism. For instance, if the parallelism of the in-core pipelined parallelism in the NPU, `doubleBuffer`, can be set to 1 or 2, then the maximum parallelism of the in-core pipelined parallelism is 2. Setting `doubleBuffer = 1` disables the in-core pipelined parallelism, and correspondingly, setting `doubleBuffer` greater than 1 enables it.

[0076] At this point, the second preset number ubLength is set to: ubLength = ubSize / (doubleBuffer * num * sizeof(datatype)). Here, doubleBuffer = 2 represents the maximum parallelism of the in-core pipelined parallelism function. Since enabling in-core pipelined parallelism allows for pipelined processing of instructions or tasks within a single AI Core, thereby improving the computational efficiency and throughput of the AI ​​Core, it is more reasonable to set the second preset number based on the maximum parallelism of the in-core pipelined parallelism function for application scenarios with frequent large data volumes, especially when the NPU coprocessor supports in-core pipelined parallelism.

[0077] Optionally, when the target binocular operator supports in-place computation, if the coprocessor unit also supports in-core pipelined parallelism, and the difference between the number of target data read by a single processing core per call to the level 0 interface and the quotient of the product of the memory space size of a single processing core, the memory size used, and the maximum parallelism of the in-core pipelined parallelism function, is within a preset range, then the second preset number is the number of target data read by a single processing core per call to the level 0 interface. The memory size used is the memory size used by a single AI Core to perform one target binocular operator computation using the in-place computation method.

[0078] In computing, interface levels (such as level 0, level 1, level 2, or level 3 interfaces) typically refer to different levels of abstraction for interacting with hardware. Higher interface levels are generally easier to use and integrate into applications. Lower-level interfaces may offer better control over hardware performance but require more programming work. Accordingly, level 0 interfaces are the closest to the hardware, providing direct control over hardware resources and allowing users to perform low-level operations such as memory access and register manipulation. Level 3 interfaces are the most abstract, providing more advanced functions and services, such as application frameworks, database management systems, and advanced graphical user interface libraries.

[0079] In the NPU coprocessor, compared to level 1, 2, or 3 interfaces, the level 0 interface allows the AI ​​Core to directly access and control hardware resources, such as CPU registers, memory addresses, or dedicated hardware accelerators, thus improving the real-time performance of target binocular operator computation.

[0080] Taking the aforementioned binocular operator Add as an example, if the Add operator supports in-place computation and the NPU supports in-core pipelined parallelism, then the quotient of the memory space size of a single AI Core divided by the product of the memory size used and the maximum parallelism of the in-core pipelined parallelism function is ubSize / (doubleBuffer*num*sizeof(datatype)), hereinafter referred to as the first quotient, representing the number of target data processed by a single AI Core in this case. The memory size used is the memory size used by a single AI Core to perform one target binocular operator computation using the in-place computation method.

[0081] For an AI Core with a memory size of 256KB and target data with a datatype of float16, the memory space occupied by the target data in the first quotient is: sizeof(datatype)*(ubSize / (doubleBuffer*num*sizeof(datatype)))=256K / (2*2)=65536Byte, which is one-quarter of the memory space of a single AI Core. Therefore, the first quotient is 65536Byte / sizeof(float16).

[0082] When the NPU coprocessor uses a level 0 interface, the amount of data that a single AI Core can directly access and control hardware resources is 255 * 256 bytes = 65280 bytes. Here, 255 refers to the number of times data can be read in one call to the level 0 interface, and 256 bytes refers to the storage space occupied by each data read. Therefore, the number of target data items that can be read in each call to the level 0 interface is 65280 bytes / sizeof(float16).

[0083] From the formula (65536Byte - 65280Byte) = 256Byte, we can see that the memory space occupied by the target data of the first quotient is slightly larger than the storage space occupied by the data read by a single AI Core in a single call to the level 0 interface; the two are quite close. From the formula 65536Byte / sizeof(float16) - 65280Byte / sizeof(float16) = (65536Byte - 65280Byte) / sizeof(float16) = 256Byte / 2, we can see that the number of target data items that a single AI Core can read in a single call to the level 0 interface differs very little from the first quotient, falling within a very small preset range. Therefore, we can approximately consider that the amount of data read per call by the Add level 0 interface is the same as the first quotient, and the hardware and software coordination requirements of the system are basically the same, achieving consistent computational performance and metrics.

[0084] Furthermore, each time the AI ​​Core processes the target data of the first quotient, it can achieve 100% memory utilization. However, since the memory space occupied by the target data of the first quotient is slightly larger than the storage space occupied by the data read by a single AI Core in a single call to the level 0 interface, it is necessary to call the level 0 interface twice and execute the level 0 interface scheduling twice.

[0085] When the AI ​​Core processes the target data read from the level 0 interface once, although the AI ​​Core memory utilization will decrease slightly, only one level 0 interface scheduling is required.

[0086] Compared to the slight reduction in AI Core memory utilization, setting the second preset number to the number of target data that a single AI Core can read each time it calls the level 0 interface can avoid the situation where the AI ​​Core calls the level 0 interface twice for each data transfer, thereby reducing the time spent on executing the target binocular operator calculation, which is more reasonable.

[0087] In one implementation, based on the range of the number of target data corresponding to the second level, the lower boundary parameter of the third level is set to be greater than the upper boundary parameter of the second level, and the upper boundary parameter of the third level is infinity.

[0088] For example, for the NPU coprocessor, the range of the number of target data corresponding to the third level is set to (totalBlockDim*ubLength, +∞). Where totalBlockDim is the total number of AI Cores in the NPU coprocessor, and ubLength is the second preset number.

[0089] Therefore, this solution, at least based on the in-situ computation characteristics of the target binocular operator, sets the upper and lower boundary parameters used at the target level, which can make the boundary parameter settings more reasonable. Furthermore, this solution's in-situ computation of the target binocular operator can also improve the memory utilization efficiency of each processing core, reduce data transfer time, and thus improve the system's performance in calculating the target binocular operator.

[0090] Step S203: Determine the data segmentation parameters for the target data based on the target data volume and target level. The data segmentation parameters include the target number of processing cores, the number of processing cycles per core for the target number of data segments, and the data volume per processing cycle.

[0091] In one embodiment, this scheme designs targeted data segmentation schemes for target data of different sizes and determines the data segmentation parameters of the target data. The target number is the number of processing cores scheduled by the coprocessor unit for parallel processing of the target binocular operator computation for the target data of the target data size. The target number is not greater than the total number of processing cores included in the coprocessor unit.

[0092] For example, for the target data corresponding to the target data volume of the first level, the data segmentation scheme includes: the target number is the target data volume divided by a first preset number and rounded up. Each processing core of the target number processes the data once. The data volume processed by each processing core that is one less than the target number is the first preset number, and the data volume processed by the remaining processing core of the target number is the difference between the target data volume and the sum of the data volumes processed by each processing core that is one less than the target number.

[0093] For example, for the NPU (Neural Processing Unit), the target number of AI Cores to be scheduled is blockDim = (X + Length - 1) / Length. This formula represents dividing the target data volume by a first preset number and rounding up. Here, X is the target data volume, and Length is the first preset number. Each AI Core of the target number is processed once, and the data volume per processing iteration for each AI Core of the target number can be determined by a single core partitioning operation. Specifically, the data volume per iteration processed by the first (blockDim - 1) AI Cores (hereinafter referred to as large cores) is formLength = Length, and the data volume per iteration processed by the last AI Core (hereinafter referred to as small core) is tailLength = X - (blockDim - 1) * formLength.

[0094] For example, for the target data corresponding to the target data volume of the second level, the data segmentation scheme includes: the target number is the total number of multiple processing cores; each of the processing cores of the target number processes data once; the data volume processed by each processing core that is one less than the target number is the quotient of the target data volume divided by the target number; and the data volume processed by the remaining one processing core in the target number is the difference between the target data volume and the sum of the data volumes processed by each processing core that is one less than the target number.

[0095] For example, for a coprocessor unit (NPU), the target number of AI Cores to be scheduled is blockDim = totalBlockDim, where totalBlockDim is the total number of AI Cores in the NPU. Each AI Core of the target number is processed once, and the amount of data processed per cycle by each AI Core can be determined by a single core partitioning. Specifically, the amount of data processed per cycle by the first (blockDim-1) AI Cores (large cores) is formerLength = X / blockDim, and the amount of data processed per cycle by the last AI Core (small core) is tailLength = X - (blockDim-1) * formerLength.

[0096] For example, for the target data corresponding to the target data volume of the third level, the data segmentation scheme includes: the target number is the total number of the plurality of processing cores. The total data volume processed by each processing core that is one less than the target number is the quotient of the target data volume divided by the target number, and the total data volume processed by the remaining one processing core out of the target number is the difference between the target data volume and the sum of the total data volume processed by each processing core that is one less than the target number.

[0097] For each processing core in the target number of processing cores, the number of processing times of the processing core is the value of dividing the corresponding total data volume by the second preset number and rounding up; wherein, for processing cores other than the last processing, the amount of data processed by the processing core in each processing session is the second preset number, and for the last processing, the amount of data processed by the processing core in each processing session is the difference between the corresponding total data volume and the sum of the amounts of data processed in each processing session other than the last processing session.

[0098] For example, for a coprocessor unit (NPU), the target number of AI Cores to be scheduled is blockDim = totalIBlockDim, where totalBlockDim is the total number of AI Cores in the NPU. Since each of the target number of AI Cores processes more than once, one inter-core partitioning and one intra-core partitioning can be performed to determine the amount of data processed per cycle by each of the target number of AI Cores.

[0099] First, an inter-core partitioning is performed to determine the total data volume to be processed by each of the target number of AI Cores. The first (blockDim-1) AI Cores (large cores) process a total data volume `formerLength = XIblockDim`, and the last AI Core (small core) processes a total data volume `tailLength = X - (blockDim-1) * formerLength`. Since the total data volume allocated to a single AI Core exceeds the memory size of that single AI Core, the data transfer and computation cannot be completed in one round. Therefore, in addition to the inter-core partitioning, an intra-core partitioning is required.

[0100] Secondly, an intra-core partitioning is performed to determine the amount of data processed by each AI Core in each processing iteration for the target number of cores. The intra-core partitioning strategy is consistent for both large and small cores. Taking the large core as an example, the processing iterations for a single AI Core of the target number are calculated as follows: `formerNum = (formerLength + ubLength - 1) / ubLength`. This formula represents dividing the total amount of data processed by the core by a second preset number and rounding up. The amount of data processed by a single AI Core of the target number in the first (formerNum - 1) iterations is `formerTilelLength = ubLength`, and the amount of data processed by a single AI Core in the last iteration is `formerTilelLength = formerLength - (formerNum - 1) * formerTilelLength`.

[0101] Therefore, the target number of AI Cores each process multiple times until the target data amount allocated to each of the target number of AI Cores is scheduled. During the calculation process, consideration is given to improving the memory utilization of each AI Core while also improving the parallel processing performance of the target number of AI Cores, trying to achieve a balance between the two, so as to improve the overall speed of using multi-core processors to move the target data amount to execute the target binocular operator.

[0102] Step S204: Send the data segmentation parameters to the coprocessing unit so that it can schedule the processing core to perform calculations according to the data segmentation parameters.

[0103] In one embodiment, the main processing unit transmits data segmentation parameters to guide the coprocessing unit to schedule multiple processing cores according to the data segmentation parameters, and to move the target data volume of target data to perform target binocular operator calculation.

[0104] Therefore, this solution designs targeted data partitioning schemes for target data of different sizes, determines data partitioning parameters based on the data partitioning schemes, and guides the coprocessing unit to perform reasonable scheduling and allocation of multiple processing cores, reducing communication requirements and overhead between multiple processing cores, improving their parallel processing performance, and thus improving the system's computing performance under different computing loads.

[0105] The following example will use the target binocular operator as the aforementioned binocular operator Add that supports in-place computation, the NPU coprocessor unit comprising a single AI Core with a memory size of 256KB, and the target data data type as float16. Figures 3-5 The data partitioning scheme used in this solution is described in detail.

[0106] For example, Figure 3 The diagram shows a data segmentation scheme provided in an embodiment of this application.

[0107] like Figure 3 As shown, the total number of AI Cores in the NPU is totalBlockDim = N. For the target data corresponding to the target data volume of the first level, the number of AI Cores scheduled and used by the NPU can be less than N. For the target number of AI Cores, each AI Core is moved once and processed once, so only one inter-core splitting is needed.

[0108] For the large cores in the target number of AI Cores, each AI Core processes a first preset number of data points (Length), and occupies a memory space of Length * sizeof(float16). The number of small cores in the target number of AI Cores is 1. For a small core, the amount of data processed in a single run is the difference between the target data amount and the sum of the data amounts processed in a single run by all the large cores. This difference is typically less than the first preset number of data points (Length). Therefore, for the target data corresponding to the first level of target data amount, the concurrency rate of the AI ​​Cores can be less than 100%, and the memory utilization rate of each AI Core can also be less than 100%.

[0109] For the target data corresponding to the second level of target data volume, the number of AI Cores scheduled and used by the coprocessor unit (NPU) is equal to N. For each AI Core with the target data volume, each AI Core performs one data transfer and one computation, thus requiring only one inter-core splitting.

[0110] For the large cores in the target number of AI Cores, each AI Core processes a first preset number of data points (ubLength), occupying a memory space of ubLength * sizeof(float16). The number of small cores in the target number is 1. For a small core, the amount of data processed in a single run is the difference between the target data amount and the sum of the data amounts processed in a single run by all the large cores. This difference is typically less than the second preset number (ubLength). Therefore, for the target data corresponding to the second level of target data amount, the concurrency rate of the AI ​​Cores is 100%, and the memory utilization rate of each AI Core is typically less than 100%.

[0111] For the target data corresponding to the target data volume of the third level, the number of AI Cores scheduled and used by the coprocessing unit (NPU) is equal to N. For the number of AI Cores corresponding to the target data volume, each AI Core moves data m times and processes data m times, where m is greater than 1. Therefore, in addition to performing one inter-core partitioning, another intra-core partitioning is required.

[0112] The first step is to perform inter-core partitioning. For the large cores in the target number of AI Cores, the total amount of data processed by each AI Core is the quotient of the target amount of data divided by the target number. The number of small cores in the target number is 1. For a small core, the total amount of data processed by that AI Core is the difference between the target amount of data and the sum of the total amount of data processed by all the large cores, which is usually less than the total amount of data processed by a single large core.

[0113] The second step is to perform kernel segmentation. The same segmentation strategy is applied to both large and small cores. Taking large cores as an example, the number of computations *m* for a single AI Core is first determined as the total data volume divided by a second preset number and rounded up. For the first (m-1) processing iterations, the data volume per iteration is the second preset number *ubLength*. For the *m*th processing iteration, the data volume per iteration is the difference between the total data volume and the sum of the data volumes per iteration that are not the last processing iteration.

[0114] Therefore, for the target data corresponding to the target data volume of the third level, the concurrency rate of AI Core is 100%. The memory utilization rate of each AI Core may be 100% in the first (m-1) times out of m times, while it is usually less than 100% in the mth time. The memory utilization rate of the first (m-1) times is usually greater than that in the mth time.

[0115] As mentioned earlier, if the NPU (Non-Processing Unit) supports in-core pipelined parallelism, the second preset number `ubLength` will typically be smaller compared to when it doesn't support in-core pipelined parallelism. The difference is roughly one-tenth of the maximum parallelism of the in-core pipelined parallelism function. Therefore, for target data corresponding to the third level of target data volume, since each AI Core processes the target data more than once, enabling in-core pipelined parallelism and setting `doubleBuffer` to the maximum parallelism can further improve the computation speed of the target binocular operator. For target data corresponding to the first or second level of target data volume, since each AI Core processes the target data only once, setting `doubleBuffer` to 1 can disable in-core pipelined parallelism. These settings optimize the performance of multiple AI Cores executing the target binocular operator.

[0116] For example, Figure 4 The diagram shows a data segmentation scheme provided in an embodiment of this application.

[0117] Among them, the target binocular operator supports in-situ computation, the NPU supports in-core pipelined parallelism, and the AI ​​Core in the MPU uses a level 0 interface to read the target data.

[0118] like Figure 4 As shown, since the target binocular operator uses a level 0 interface to read target data, as mentioned earlier, the second preset number ubLength is slightly less than the maximum number of target data that a single AI Core can handle at one time. Therefore, for the target data corresponding to the third level of target data volume, the concurrency rate of a single AI Core is 100%, and the memory utilization rate of a single AI Core is less than 100% for m rounds. The memory utilization rate in the first (m-1) rounds is usually greater than that in the mth round. Other content is related to... Figure 3 Similarly, I will not elaborate further.

[0119] For example, Figure 5 The diagram shows a data segmentation parameter provided in an embodiment of this application.

[0120] like Figure 5 As shown, the marking symbols involved in the embodiments are as follows:

[0121] Total number of AI Cores in the NPU coprocessor: totalBlockDim

[0122] Total data volume for large-core computation: formerLength

[0123] Total data volume calculated by the small core: tailLength

[0124] Largest number of cores: formerBlock.

[0125] The first, second, and third levels represent the target data volume as small, medium, and large, respectively.

[0126] The target data is aligned in the storage medium according to the preset alignment boundary of 32 bytes.

[0127] For the target data corresponding to the first level of target data volume, the target data volume is in the range [1, totalBlockDim*Length]. Among them, the first preset number Length = 16*256 / sizeof(float16) = 128*16, which is equivalent to calling the Add 0 level interface to read the target data, but the number of times the target data can be read by calling the 0 level interface is set to 16; the in-core pipeline parallel function is disabled, that is, doubleBuffer = 1 is set.

[0128] For the target data corresponding to the second level of target data volume, the target data volume is within the range (totalBlockDim*Length, totalBlockDim*ubLength). Among them, the first preset number Length = 128*16; the second preset number ubLength = 255*256 / sizeof(float16) = 128*255; similarly, the in-kernel pipeline parallel function is disabled, that is, doubleBuffer = 1 is set.

[0129] For the target data corresponding to the third level of target data volume, the target data volume is within the range (totalBlockDim*ubLength, +∞). Where, ubLength = 255*256 / sizeof(float16) = 128*255; enable the kernel pipeline parallel function, that is, set doubleBuffer = 2 as the maximum parallel amount.

[0130] For target data corresponding to different levels of target data volume, the data segmentation parameters determined according to the data segmentation scheme corresponding to each level include:

[0131] For the target data corresponding to the first level of target data volume, the number of AI Cores called by the NPU is blockDim = (X + Length - 1) / Length, where Length = 128 * 16. One inter-core partitioning is performed. In the first (blockDim - 1) large cores, each large core is allocated a single data volume of formerLength = Length. In the last small core, the single data volume allocated to the small core is tailLength = X - (blockDim - 1) * formerLength; typically, tailLength <= formerLength.

[0132] For the target data corresponding to the second level of target data volume, the number of AI Cores called by the NPU is blockDim = totalBlockDim. One inter-core partitioning is performed. In the first (blockDim-1) large cores, each large core is allocated a single data processing volume `formerLength = X / bIockDim`. In the last small core, the single data processing volume `tailLength = X - (blockDim-1) * formerLength`. Where `tailLength <= formerLength`.

[0133] For the target data corresponding to the target data volume of the third level, the target number of AI Cores called by the NPU is blockDim = totalBlockDim. First, one inter-core partitioning is performed. In the first (blockDim-1) large cores, the total data volume processed by each large core is formerLength = X / blockDim. In the last small core, the total data volume processed by the small core is tailLength = X-(blockDim-1)*formerLength; where tailLength <= formerLength.

[0134] Secondly, for each large core and one small core, perform one kernel partitioning operation using the same strategy:

[0135] Set `formerTileLength = tailTileLength = 128 * 255`. In the first (blockDim - 1) large cores, the number of processing iterations for each large core is `formerNum = (formerLength + formerTileLength - 1) / formerTileLength`. The amount of data processed per iteration in the first (formerNum - 1) iterations is `formerTileLength`. The amount of data processed per iteration in the last iteration is `formerTailLength = formerLength - (formerNum - 1) * formerTileLength`. Where `formerTailLength <= formerTileLength`.

[0136] In the last small core, the number of processing times is tailNum = (tailLength + tailTileLength - 1) / / tailTileLength. In the first (tailNum - 1) processing times, the amount of data processed per time is tailTileLength. The amount of data processed per time in the last processing time is tailTailLength = tailLength - (tailNum - 1) * tailTileLength. Where tailNum <= formerNum and tailTailLength <= tailTileLength.

[0137] It is understood that the sequence number of each step in the above embodiments does not imply the order of execution. The execution order of each process should be determined by its function and internal logic, and should not constitute any limitation on the implementation process of the embodiments of this application. Furthermore, in some possible implementations, each step in the above embodiments may be selectively executed according to actual circumstances; it may be partially or fully executed, without limitation here. Additionally, all or part of any feature in the above embodiments can be freely and arbitrarily combined without contradiction. The combined technical solutions are also within the scope of this application.

[0138] This application also provides a computing device 1000. For example... Figure 6 As shown, the computing device 1000 includes: a bus 1002, a main processor 1004, a coprocessor 1006, a memory 1008, and a communication interface 1010. The main processor 1004, coprocessor 1006, memory 1008, and communication interface 1010 communicate with each other via the bus 1002. It should be understood that this application does not limit the number of main processors, coprocessors, and memories in the computing device 1000.

[0139] Bus 1002 can be a Peripheral Component Interconnect (PCI) bus or an Extended Industry Standard Architecture (EIA) bus, etc. Buses can be categorized as address buses, data buses, control buses, etc. For ease of representation, Figure 6 The bus 1002 may be represented by a single line, but this does not mean that there is only one bus or one type of bus. The bus 1002 may include a path for transmitting information between various components of the computing device 1000 (e.g., main processor 1004, coprocessor 1006, memory 1008, communication interface 1010).

[0140] The main processor 1004 may include any one or more processors such as a central processing unit (CPU), a graphics processing unit (GPU), a microprocessor (MP), or a digital signal processor (DSP).

[0141] The coprocessor 1006 may include any one or more processors such as a graphics processing unit (GPU), a neural network processing unit (NPU), a tensor processing unit (TPU), or an image processing unit (IPU). For example, when the main processor 1004 is a central processing unit (CPU), the coprocessor 1006 may be a GPU, an NPU, a TPU, or an IPU. Similarly, when the main processor 1004 is a GPU, the coprocessor 1006 may be an NPU, a TPU, or an IPU.

[0142] The memory 1008 may include volatile memory, such as random access memory (RAM). The memory 1008 may also include non-volatile memory, such as read-only memory (ROM), flash memory, hard disk drive (HDD), or solid state drive (SSD).

[0143] The memory 1008 stores executable program code / instructions, and the main processor 1004 executes the executable program code / instructions to implement... Figure 2The method flow shown above implements all or part of the steps of the method in the above embodiments, so that the coprocessor 1006 performs calculations. In other words, the memory 1008 stores a program / instruction for performing all or part of the steps of the method in the above embodiments.

[0144] The communication interface 1010 uses transceiver modules such as, but not limited to, network interface cards and transceivers to enable communication between the computing device 1000 or a cluster of multiple computing devices 1000 and other devices or communication networks.

[0145] This application provides a computing device, including: a memory, a main processor, and a coprocessor; the memory is coupled to the main processor and the coprocessor; the memory is used to store a program; the main processor is used to execute the program stored in the memory. When the program stored in the memory is executed, the main processor is used to execute the method described in the above embodiment, so that the coprocessor performs scheduling. In this application embodiment, the computing device may be a server or a computer, etc.

[0146] Based on the methods in the above embodiments, this application provides a computer-readable storage medium storing a computer program that, when run on a processor, causes the processor to execute the methods in the above embodiments.

[0147] Based on the methods in the above embodiments, this application provides a computer program product that, when run on a processor, causes the processor to execute the methods in the above embodiments.

[0148] The method steps in the embodiments of this application can be implemented in hardware or by a processor executing software instructions. The software instructions can consist of corresponding software modules, which can be stored in random access memory (RAM), flash memory, read-only memory (ROM), programmable read-only memory (PROM), erasable programmable read-only memory (EPRCM), electrically erasable programmable read-only memory (EEPROM), registers, hard disks, portable hard disks, CD-ROMs, or any other form of storage medium known in the art. An exemplary storage medium is coupled to the processor, enabling the processor to read information from and write information to the storage medium. Of course, the storage medium can also be a component of the processor. The processor and the storage medium can reside in an AS1C.

[0149] In the above embodiments, implementation can be achieved, in whole or in part, through software, hardware, firmware, or any combination thereof. When implemented in software, it can be implemented, in whole or in part, as a computer program product. A computer program product includes one or more computer instructions. When the computer program instructions are loaded and executed on a computer, all or part of the flow or function according to the embodiments of this application is generated. The computer can be a general-purpose computer, a special-purpose computer, a computer network, or other programmable device. The computer instructions can be stored in or transmitted through a computer-readable storage medium. The computer instructions can be transmitted from one website, computer, server, or data center to another website, computer, server, or data center via wired (e.g., coaxial cable, fiber optic, digital subscriber line (DSL)) or wireless (e.g., infrared, wireless, microwave, etc.) means. The computer-readable storage medium can be any available medium that a computer can access or a data storage device such as a server or data center that integrates one or more available media. The available medium can be a magnetic medium (e.g., floppy disk, hard disk, magnetic tape), an optical medium (e.g., DVD), or a semiconductor medium (e.g., solid-state disk (SSD)).

[0150] It is understood that the various numerical designations used in the embodiments of this application are merely for descriptive convenience and are not intended to limit the scope of the embodiments of this application.

Claims

1. A data scheduling method applied to digital visual image processing, characterized in that, The method is applied to a main processing unit, which is communicatively connected to a coprocessor unit, the coprocessor unit including multiple processing cores, and the method includes: The target data of the target data volume is aligned according to a preset alignment boundary; the target data is any input data of the target binocular operator; the target data volume is the number of target data. The target level is determined based on the aligned target data volume; The data segmentation parameters of the target data are determined based on the target data volume and the target level; the data segmentation parameters include the target number of processing cores used, the number of processing times of each of the target number of processing cores, and the data volume per processing time. The data segmentation parameters are sent to the coprocessing unit, which then schedules the processing core to perform calculations based on the data segmentation parameters. Each of the aforementioned processing cores includes the same size of memory space; The target level includes the first level; The lower boundary parameter of the first level is 1, and the upper boundary parameter of the first level is the product of the total number of processing cores and the first preset number of target data processed by a single processing core each time. The target level also includes a second level; The lower boundary parameter of the second level is greater than the upper boundary parameter of the first level. The upper boundary parameter of the second level is the product of the total number of processing cores and the second preset number of target data processed by a single processing core each time. For the target data corresponding to the target data volume of the first level, the target number is the value of dividing the target data volume by the first preset number and then rounding up; Each of the processing cores for the target number is processed once; Each processing core, which is one less than the target number, processes a single data volume of the first preset number, and The amount of data processed in a single run by the remaining processing core out of the target number is the difference between the target data amount and the sum of the data amounts processed in a single run by each of the processing cores that are one less than the target number. For the target data corresponding to the target data volume of the second level, the target number is the total number of the plurality of processing cores; Each of the processing cores for the target number is processed once; The amount of data processed per cycle by each processing core that is one less than the target number is the quotient of the target data amount divided by the target number, and The amount of data processed in a single run by the remaining processing core out of the target number is the difference between the target data amount and the sum of the data amounts processed in a single run by each of the processing cores that are one less than the target number. For the target data corresponding to the target data volume of the third level, the number of targets is the total number of the plurality of processing cores; The total amount of data processed by each of the processing cores that is one less than the target number is the quotient of the target data amount divided by the target number, and The total amount of data processed by the remaining processing core out of the target number is the difference between the target data amount and the sum of the total data amounts processed by each of the processing cores that are one less than the target number.

2. The method according to claim 1, characterized in that, If the target binocular operator supports in-place computation The second preset number is the memory space size of a single processing core, which is the quotient of the memory size used; the memory size used is the memory size used by a single processing core to perform the target binocular operator calculation once using the in-place calculation method.

3. The method according to claim 2, characterized in that, If the coprocessor unit also supports in-core pipelined parallelism. The second preset number is the memory space size of a single processing core, which is the quotient of the product of the memory size used and the maximum parallelism of the pipelined parallel function within the core.

4. The method according to any one of claims 1-3, characterized in that, The target level also includes a third level; The lower boundary parameter of the third level is greater than the upper boundary parameter of the second level.

5. The method according to claim 1, characterized in that, For each processing core in the target number of processing cores, the number of processing times for that processing core is the value of dividing the corresponding total data volume by the second preset number and then rounding up; Specifically, for data not processed in the last iteration, the processing core processes a single data set of the second preset number of data items each time. For the last processing, the amount of data processed by the processing core in each iteration is the difference between the total amount of data processed and the sum of the amounts of data processed in each iteration that were not the last processing iteration.

6. A computing device, characterized in that, include: Memory, main processor, and coprocessor; The memory is coupled with the main processor and coprocessors; Memory, used to store programs; The main processor is used to invoke a program to execute the method as described in any one of claims 1-5, so that the coprocessor performs the scheduling.

Images