Parallelization strategies for training neural networks

Abstract

A parallelization strategy for training neural networks is disclosed. Given a plurality of devices (e.g., graphics processing units) for training a deep learning (DL) model, a system and method of training a neural network to systematically find a cross-over point that indicates which parallelization strategy should be implemented for optimizing the training of the DL model on a particular system to achieve the greatest efficiency gain.

Description

BACKGROUND

[0001] The use of deep learning (DL) and other machine learning models is growing in prevalence, and the size of the data sets used to train them continues to increase, which in turn leads to longer and longer training times. As machine learning becomes more complex, models become larger and the data sets used to train these models contain more information. Thus, training machine learning models with larger data sets begins to impact training times, given the increased number of devices that need to be employed. The synchronization and communication overhead necessary to ensure accuracy between the increased number of devices further limits the overall training time. Conventional techniques for processing large data sets result in inefficiencies that become more pronounced as data sets become larger and models become more complex. BRIEF DESCRIPTION OF DRAWINGS

[0002] Various techniques will be described with reference to the drawings, in which:

[0003] Figure 1 A schematic diagram showing a system in which embodiments according to the present disclosure are implemented is shown;

[0004] Figure 2 A schematic diagram showing a stochastic batch gradient descent process representing training of a neural network with a batch of inputs according to one embodiment is shown;

[0005] Figure 3 A schematic diagram showing a synchronous stochastic gradient descent process representing data parallel training of a neural network with a batch of inputs using multiple computing devices according to one embodiment is shown.

[0006] Figure 4 A schematic diagram showing a general deep network trained in data parallel using two computing devices with a batch of inputs according to one embodiment is shown.

[0007] Figure 5 A schematic diagram showing a general deep network trained in model parallel using two computing devices with a batch of inputs according to one embodiment is shown;

[0008] Figure 6 A schematic diagram showing a general deep network trained using both data parallel training and model parallel training according to one embodiment is shown;

[0009] Figure 7 An illustrative example of a process for training a neural network according to one embodiment is shown;

[0010] Figure 8 Another illustrative example of a process for training a neural network according to one embodiment is shown;

[0011] Figure 9An example of a parallel processing unit ("PPU") is shown in accordance with one embodiment;

[0012] Figure 10 An example of a general processing cluster ("GPC") is shown in accordance with one embodiment;

[0013] Figure 11 An example of a memory partition unit is shown in accordance with one embodiment;

[0014] Figure 12 An example of a streaming multiprocessor is shown in accordance with one embodiment; and

[0015] Figure 13 A computer system in which various examples can be implemented in accordance with one embodiment is shown. DETAILED DESCRIPTION

[0016] In one embodiment, efficiency gains (e.g., improved training time) for large, complex deep learning (DL) models are achieved with systems and methods implemented in accordance with the present disclosure. One example of a DL model is a neural network. In one embodiment, improved ways of configuring neural network training using both model parallelism (MP) and data parallelism (DP) are identified and described herein. In one embodiment, DP involves using independent subsets of a training dataset (referred to as mini-batches) on independent workers (e.g., devices, processors) using copies of the model. MP involves distributing components of the neural network being trained across multiple workers (e.g., devices, processors), so each worker has a different component of the entire neural network being trained. In one embodiment, MP involves splitting the model across multiple workers when processing the same mini-batch. In one embodiment, a subset of the training dataset (e.g., a mini-batch) is a batch of inputs that is fed to each worker. In one embodiment, a collection of all mini-batches in a training step is a global batch. In one embodiment, a worker comprises one or more devices organized as a group of devices or one or more processors organized as a group of processors to form a single unit. In one embodiment, a worker is one device or one processor. In one embodiment, improved configurations for training a neural network are identified by combining the use of MP and DP. In one embodiment, a combination of MP and DP is employed to minimize end-to-end training time for a given neural network on available hardware.

[0017] In one embodiment, the amount of parallelism using DP (e.g., the number of mini-batches used by different workers / devices / processors) is increased until efficiency gains are no longer significant (e.g., as defined by a set of importance criteria). At this point, in one embodiment, MP is introduced, and the amount of introduced MP parallelism combined with the use of DP (e.g., the number of components that the neural network being trained is split into) is increased to achieve greater efficiency gains (e.g., improved training time).

[0018] In one embodiment, a dataset used to train a neural network is subdivided into subsets (e.g., mini-batches), and a collection of mini-batches used to form a global batch is processed across multiple computing devices such as graphics processing units (GPUs) (e.g., using DP). In one embodiment, multiple tensor processing units (TPUs) are used to process the subsets. As the global batch size scales / increases in number or size, the overall training time to train a given neural network begins to become inefficient. In one embodiment, as the global batch size increases, the amount of time it takes to accurately train a neural network using DP begins to take longer. As such, in one embodiment, a second type of parallelism (e.g., MP) is implemented to work in conjunction with DP to create a hybrid parallelism (e.g., DP plus MP). In one embodiment, training a neural network using this hybrid parallelism allows each copy of the model that exists on a given worker to parallelize the model, such that now, each worker includes multiple computing devices, rather than each worker being one computing device. In one embodiment, each neural network will have a unique scale at which to train the network by introducing MP to work in conjunction with DP and improve training time, which can overcome the scaling and statistical efficiency decline of DP. Thus, in one embodiment, when scaling the global batch size, hybrid parallelism training provides better efficiency gains (e.g., is more effective at minimizing end-to-end training time) than using DP alone. In one embodiment, the systems and methods described herein use measured information to deterministically determine a strategy, which information is measured to indicate the scaling inefficiency of DP alone and the speedup of MP alone, and is used to select the number of GPUs to run when the correct mechanism to use is selected for a given neural network.

[0019] In one embodiment, the techniques described herein are used to determine when to implement a mixed parallelism approach using MP with DP to overcome the inherent scaling efficiency loss and statistical efficiency loss experienced by DP training alone when scaled (e.g., when an increase in global batch size is beyond the range that DP can continue to process effectively to maintain accuracy). In one embodiment, end-to-end training times related to training a neural network using DP are monitored and analyzed to understand how the scaling efficiency loss and statistical efficiency loss impact training scalability, and the techniques described herein indicate that adding MP allows for training speedup to be achieved for a given neural network, which is critical to the overall scalability of the mixed parallelization strategy. Splitting the dataset (e.g., using DP) and splitting the network (e.g., using MP) each can provide efficiency gains separately, but for each parallelization approach, there is a point where the efficiency gains are offset by the overhead required for the combined parallel training to be worthwhile. In one embodiment, when the size of the global batch in DP grows to a point where the training speedup of DP alone is significantly diminished (e.g., training time is longer), MP can be used in conjunction with DP to improve training time beyond what is achievable with DP alone.

[0020] In the foregoing and in the following description, various techniques are described. For purposes of explanation, specific configurations and details are set forth in order to provide a thorough understanding of the techniques. However, it will also be apparent to one skilled in the art that the techniques described below can be practiced in different configurations and that the following description is not intended to limit the techniques in any way.

[0021] Figure 1A schematic diagram illustrating a system 100 in which embodiments in accordance with the present disclosure are implemented is shown. In one embodiment, the system 100 is a computing environment that includes a data store 102 or a data store service / device. The data store 102 can be configured to store data, metadata, training data, training data threads, and the like. In one embodiment, the training data includes a collection of image data or video data that, when obtained by one or more neural networks, trains the one or more neural networks for face recognition purposes. In one embodiment, the training data includes speech data or audio data that, when obtained by one or more neural networks, trains the one or more neural networks for speech recognition purposes. In one embodiment, the data store 102 stores training data 104 (e.g., training data threads) for training a neural network. In one embodiment, the neural network is also referred to as a deep network or simply as the network described herein. The training data 104 can include data structures such as arrays or matrices. A subset of the training data 104 can be referred to as a mini-batch, and a collection of mini-batches can be referred to as a global batch. In one embodiment, these mini-batches for training the neural network can be overlapping or proper subsets (non-overlapping). In one embodiment, the data store 102 is configured to satisfy a request to obtain the training data 104 by sending the training data 104 to a computing device 108 via a network 106 (via a wired or wireless network).

[0022] In one embodiment, the computing device 108 is a graphics processing unit (GPU), or in another embodiment, it represents a plurality of computing devices that include a plurality of GPUs. In one embodiment, the computing device 108 represents a neural network in which the neural network uses the training data 104 to train the network. In one embodiment, the computing device 108 is configured to receive or obtain the input training data 104 and determine a parallelization strategy to implement such that the computing device 108 is trained to maximize training efficiency. The input training data 104 can be received by the computing device 108 from the data store 102, another computing device, a computing service, a virtual machine, a virtual service (not shown in FIG. 1) external or internal to the system 100 via the network 106. Figure 1

[0023] ​In one embodiment, the computing device 108 includes a front-end service 110. The front-end service 110 can be a computer system or component thereof configured to receive or obtain the input training data 104. In one embodiment, the front-end service 110 is a service front-end that includes one or more processors and memory for storing instructions that, as a result of execution, cause the front-end service 110 to perform the various processes and methods described herein. In one embodiment, each of the one or more processors is a logical processor, where each logical processor is capable of simultaneously executing its own instruction stream. In one embodiment, the front-end service 110 is configured with executable code to submit and receive network application programming interface (API) requests. In one embodiment, the front-end service 110 is a computing service / device, such as a server computer system, a server cluster, a virtual computer system, a virtual server, a virtual runtime environment, a container environment, a serverless execution environment, a service hosting system, or any suitable computing entity associated with and used by the computing device 108 to receive the input training data 104.

[0024] In one embodiment, the front-end service 110 monitors training efficiency 150, which indicates training time, power consumption, and / or other information of the computing device 108 when training using at least one of the DP 120, the MP 130, or the combination of DP and MP 140. That is, in one embodiment, the front-end service 110 processes the input training data 104 and sends an indication or set of instructions for at least one type of parallelism to be implemented to train the computing device 108. In one embodiment, the type of parallelism applied is data parallelism (DP) 120, where the computing device 108 is trained using independent subsets of the training data 104 by using copies of the model network on independent computing devices. In one embodiment, and as described above, the subset of the training data 104 is referred to as a mini-batch. In one embodiment, another type of parallelism applied is model parallelism (MP) 130, where the model is split across multiple workers while working on the same mini-batch. In one embodiment, another type of parallelism applied is hybrid parallelism, which includes using both DP and MP to train the computing device 108.

[0025] In one embodiment, training efficiency 150 of each of these parallelism strategies is monitored and information related to training efficiency 150 is provided as feedback to the front-end service 110 or another service associated with the computing device 108 to determine, based on various training tests or predetermined information, which parallelism approach can be used to maximize training efficiency 150 of the computing device 108 given the training data 104 and the amount of training data 104 that has been received. In one embodiment, training efficiency 150 is a data store, data storage device, buffer, or message queue that stores values or information related to training time. Further, in one embodiment, training efficiency 150 is a data store that is separate from the computing device 108 but maintains communication with the computing device 108 via the network 106.

[0026] In one embodiment, the front-end service 110 is removed from the computing device 108 and a tool based on integer linear programming (e.g., DLP placer) is software that runs in conjunction with the computing device 108 such that when the tool is executed, the tool is optimally placed to maximize training acceleration (e.g., maximize training efficiency). In one embodiment, the tool is implemented to produce results that indicate training efficiency 150 when using DP 120 alone and further provide information to the computing device 108 to indicate how much of the neural network should be split (e.g., when to use MP and how much MP to use) to provide maximum efficiency when the global batch size of the training data 104 scales beyond what the DP can efficiently handle alone. In one embodiment, the tool is implemented to produce results that indicate training efficiency 150 of MP alone and further provide information to the computing device 108 to indicate whether to introduce DP and implement DP with MP.

[0027] In one embodiment, an integer linear programming based tool indicates to the computing devices 108 how to maximize resource utilization. The integer linear programming based tool can extract parallelism between operations in the network. In one embodiment, as a result of using the integer linear programming based tool, resources required to minimize communication overhead to move data between nodes can also be minimized. In one embodiment, using the tool, the neural network is expressed as a dataflow graph (DFG). In one embodiment, nodes of the DFP correspond to computational operations and unidirectional edges, where edges show dependencies between operations. In one embodiment, each node has an execution time and memory footprint expected by the operation. In one embodiment, edge weights correspond to the number of bytes exchanged between operations connected in it. In one embodiment, the node and edge weights are received by profiling the model on a computing device (e.g., GPU). In one embodiment, the node and edge weights can be computed analytically. In one embodiment, the tool reduces per-step training time by providing allocation of computing DFG operations through hardware graph (e.g., placement), schedule, and communication routing of activations, weights, and gradients. In one embodiment, the tool returns a correct solution when each operation is mapped to a single and unique device, dependencies between all operations are satisfied, non-overlapping operations that are co-located during execution, and total memory amount does not exceed device memory capacity. In one embodiment, while satisfying all the conditions listed above, the tool also assumes that two operations (which are co-located on a device, executed back-to-back) have no delay between termination of one operation and start of another, transferring a data block of size S over a link with bandwidth B and latency L takes (S / B + L) time, and tensor communication between devices overlaps with computation. In one embodiment, based on these assumptions, the tool predicts training speedup for a given MP solution, then applies MP in conjunction with DP to better train the network.

[0028] As described above, in one embodiment, the computing device 108 using the tool or frontend service 110 determines which of the DP 120, MP 130, or combination of DP and MP 140 to use. In one embodiment, as the global batch size continues to grow, the amount of processors needed to process the increase in training data 104 also grows. In one embodiment, as a result and at some point, the DP 120 can reach a tipping point at which the DP 120 begins to experience longer training times (e.g., training slows down) related to the training computing device 108. In one embodiment, this tipping point is also referred to as a crossover-point, threshold, or maximum efficiency value. At this point, in one embodiment, the computing device 108 applies another type of parallelism such that the training time is reduced or improved after the introduction of this other type of parallelism. In one embodiment, the improvement in training time is identified by comparing the training time to the training time that would be spent if the computing device 108 continued to train using the DP 120 alone and to the training time that would be spent if the computing device 108 used the combination of DP and MP 140.

[0029] In one embodiment, the combination of training DP and MP 140 is used to train the computing device 108 to reach yet another level of training efficiency. That is, in one embodiment, by identifying the tipping point of the DP (e.g., the maximum amount of training data or training data threads that the DP 120 is allowed to train the network without experiencing longer training times), it is determined when to apply hybrid parallelization (for various neural networks, hybrid parallelization is generally occurring at different scales than DP alone). In one embodiment, if the particular level of training efficiency is not met when training the computing device 108 using the combination of DP and MP 140, then the MP is reduced and the computing device 108 is trained using only the DP 120 and reduced MP 130 or only the DP 120 alone. In one embodiment, the DP 120 is reduced rather than the MP 130.

[0030] As such, in one embodiment, the training efficiency 150 is based on measurements from the training time associated with the training computing device 108. In one embodiment, a certain level of training efficiency indicates that the training time associated with training the computing device 108 using the combination of training data DP and MP 140 is equal to or lower than the training time associated with training the computing device 108 using only the DP and, as such, this information can be used to cause the frontend service 110 or the integer linear programming based tool associated with the computing device 108 to use the DP and MP 140 to apply this hybrid parallelization strategy until the training computing device 108 reaches another maximum efficiency.

[0031] Figure 2 A schematic diagram 200 is shown representing a stochastic batch gradient descent process for training a neural network using a batch of inputs according to one embodiment. In neural network training, in one embodiment, a batch of inputs is first propagated forward through the network (forward pass) 202 so that a loss calculation is performed for each input in the batch. The loss is then propagated backward through the network (backward pass) 204 so that a gradient calculation is performed. In one embodiment, the gradients of the batch are then combined and averaged and further used to update the weights 206. In one embodiment, the batch of inputs is of a size selected so that all computing resources of the particular device used for training can be fully utilized. In one embodiment, this entire process is referred to as stochastic batch gradient descent representing training a neural network. That is, in one embodiment, one forward pass 202 and back pass 204 together to update the gradient associated with the weights 206 is referred to as a training step. In one embodiment, assuming that all inputs are processed once, a single iteration of running through the training data set involves multiple steps and is referred to as an epoch. In one embodiment, the training process (e.g. Figure 2 ) runs for multiple epochs until some predetermined or optimal training accuracy is achieved.

[0032] Figure 3 A diagram 300 is shown illustrating a synchronized stochastic gradient descent process for training a neural network using DP with a batch of inputs using multiple computing devices, according to one embodiment. In one embodiment, using DP speeds up training when the full set of model parameters (e.g., weights) is replicated across multiple computing devices / workers (e.g., GPU 0, GPU 1, ..., GPU N-1). Figure 3 As shown, each worker works independently on a different batch of inputs through a forward pass 302 and a backward pass 304. In at least one embodiment, the gradients are then communicated between each worker and averaged. Thereafter, in one embodiment, each worker then applies the same set of gradient values ​​to the model weights 306, and any communication of gradients between each worker is done using all-reduce communication 308. In one embodiment, a model update is performed after each iteration (by using the average calculated over all gradients). In one embodiment, the batch of inputs for each worker is identified as a mini-batch, and the set of all mini-batches in a training step is the global batch.

[0033] Figure 4A diagram 400 of a general deep network using two compute devices (e.g., 2 GPUs) to train a network with a batch input is shown, according to one embodiment. In one embodiment, each compute device computes error and gradients for a subset of the dataset (or also known as a mini-batch input from the global batch input). As described above, the dataset can be composed of data structures such as arrays or matrices, where each element in the data structure is used to train a neural network. The training data can include labeled data, such as image data, video data, speech data, audio data, text, etc. In one embodiment, the parameters and gradients are then communicated between the two compute devices using Peripheral Component Interconnect Express (PCI-e) or any other communication protocol between compute devices.

[0034] In one embodiment, as Figure 4 diagram 400 shows only two compute devices; however, more than two compute devices are implemented to train a deep network. As described herein, as the global batch size continues to increase or scale, the individual DP implementation can start to experience inefficiencies (e.g., longer training times), and implementing additional parallelism techniques such as MP, further described in Figure 5 , can improve the efficiency gains of training a deep network.

[0035] Figure 5 A diagram 500 of a general deep network using two compute devices to train a network with a batch input using MP is shown, according to one embodiment. In one embodiment, the network is split by placing different operations of the diagram 500 on different devices. Traditionally, MP training is used for networks whose parameters do not fit into the memory of a single device. However, in one embodiment, MP provides per-step training speedup in cases where the entire network fits into one device. In one embodiment, as shown in Figure 5 , independent operations are performed simultaneously on various devices. For many networks, it is not trivial to split the dataflow graph (DFG) among multiple devices. In one embodiment, the overhead with respect to communication between compute nodes related to moving data between devices is considerable, so sometimes it can outweigh any gains provided by MP. Thus, in one embodiment, the DFG for the network is partitioned, and many considerations should be taken into account, such as the compute intensity for each of these devices, the bandwidth between devices, and even the network topology.

[0036] Figure 6A schematic diagram 600 of a general deep network being trained using both DP and MP, according to one embodiment, is shown. In one embodiment, using both DP and MP together can improve the end-to-end training time associated with training a neural network when the global batch size is scaled. In one embodiment, there are at least three factors that contribute to the change in the end-to-end training time of a neural network: the average time per step (T), the number of steps per epoch (S), and the number of epochs required to converge to a desired accuracy (E). (See above for Figure 2 The description of

[0055] discusses the definitions of training steps and epochs.) In one embodiment, the amount of training time (e.g., time to convergence (C)) is expressed as:

[0037] C=T×S×E (1)

[0038] In one embodiment, T is primarily determined by computational efficiency, i.e., given the same training settings, algorithm, and mini-batch size, T depends only on the computational power of the device. In one embodiment, hardware devices with better processing performance provide smaller values ​​of T. However, in one embodiment, S depends on the number of inputs in the training dataset and also on the global batch size. In one embodiment, each input or each item in the dataset is processed once per epoch. Therefore, in one embodiment, the number of steps per epoch (S) is equal to the number of items in the dataset, divided by the global batch size (e.g., the size of the training input data). In one embodiment, the number of epochs to converge (E) depends on the global batch size and other training hyper-parameters.

[0039] In one embodiment, the training time of using DP alone is monitored, calculated, and / or analyzed before MP is combined with DP and applied to train the network. In one embodiment, using N-way data parallelism (sU) is more efficient than training on a single device. N ) to speed up the training by using DP alone, which is expressed as:

[0040]

[0041] In one embodiment, T1 is the average training time per step when only a single device is used for training, and T N is the time per step when using N data-parallel devices (each with a constant mini-batch size). In one embodiment, T N is always greater than T1, which is mainly due to the fact that in DP, after each device has performed the forward pass and the reverse pass, the full communication is used (see Figure 3 In one embodiment, the gradient is exchanged between devices. The communication overhead represented will never be greater than one, and is typically less than one. In one embodiment, the ratio is the scaling efficiency (SE N ) of an N-way DP.

[0042] In one embodiment, S1 represents the total number of steps required per epoch when using a single device, while S N is the number of steps required per epoch when using N devices. In one embodiment, the global batch size is equal to the mini-batch size when using a single device. In one embodiment, the N-way data parallelism of each device performs independent steps with its own mini-batch data. Thus, in one embodiment, the global batch size is N times the mini-batch size for each device. Thus, in one embodiment, is also equal to N.

[0043] In one embodiment, E1 represents the number of epochs to converge when using a single device. In one embodiment, E N is the number of epochs required when using N number of devices. In one embodiment, with a larger global batch size (higher than the number N), the gradients from a larger number of training samples are averaged. In one embodiment, as a result of the averaging, it leads to a tendency of the model to over-fit (over-fitting) and get attracted to local minima or saddle points. In one embodiment, then the network will require more epochs to converge. As such, in one embodiment, is typically less than one. Thus, equation 2 can be simplified as:

[0044]

[0045] In one embodiment, when training with the number of devices (N), both SE N and decrease. In one embodiment, with a larger global batch size, hyperparameter tuning is applied in an effort to minimize the increase in the number of epochs required for convergence. However, in one embodiment, beyond a certain global batch size, for any particular network, the number of epochs required for convergence increases rapidly, regardless of the hyperparameter tuning.

[0046] In one embodiment, MP training time is also monitored, computed, and / or analyzed on a given neural network. As described above, using MP alone does not scale well and thus cannot be considered a widely applicable strategy with large numbers of devices and datasets. However, for reference purposes, training time using MP can be described herein. That is, in one embodiment, using MP enables multiple devices to perform operations on the same mini-batch simultaneously. In one embodiment, using MP reduces the training step time for a single step (e.g., the term “T” in Equation 1). In one embodiment, the speedup from M-way MP is represented by SU M In one embodiment, the measurement of M-way MP is measured by splitting the network among multiple devices and measuring the time to perform each step. In one embodiment, it should be noted that the SU M speedup already includes the communication overhead related to data movement between related operations placed across multiple devices.

[0047] As described previously, in one embodiment, when using MP alone, the global batch size remains unchanged. Thus, in one embodiment, the number of steps per epoch (e.g., the term “S” in Equation 1) and the number of epochs required for convergence (e.g., the term “E” in Equation 1) remain unchanged. In one embodiment, when improving SU M , it minimizes the convergence time by only reducing the term T in Equation 1 (while the other two terms remain unchanged). In one embodiment, MP alone is not considered a widely applicable scalable parallelization strategy. However, in one embodiment and as described throughout this disclosure, MP is combined with DP to extend training scalability beyond the limits that DP alone and MP alone can provide.

[0048] In one embodiment, a hybrid parallelism strategy (e.g., using a combination of DP and MP) is applied to achieve maximum efficiency (e.g., improve training time) in training a neural network. As described above, in one embodiment, N-way DP is used as provided in Equation 3 to obtain speedup. In one embodiment, assuming that the system is scaled up to N devices using N-way DP, and the training speedup achieved is equal to or exceeds the maximum efficiency, and if additional devices become available for training, the neural network can be trained using DP alone or using a hybrid approach. In one embodiment, the network identifies when to use DP alone, and when to combine DP with MP to obtain the maximum possible training speedup (e.g., maximize efficiency or training time). In one embodiment, by applying DP alone, the speedup using additional devices (e.g., M x N) compared to using one device is (replacing N in Equation 3 with M x N):

[0049]

[0050] In one embodiment, the speedup of MxN-way DP (Equation 4) is compared to the speedup of N-way DP (Equation 3). In one embodiment, the scaling efficiency of a system with MxN-way DP is generally lower compared to N-way, thus all-reduce communication occurs across a large number of devices and thus results in slower training time. In one embodiment, depending on the values of N, M, and system configuration, all-reduce communication can span slower inter-node links. In one embodiment, this results in an increase in all-reduce time and reduces SE M×N Furthermore, in one embodiment, due to the large global batch size at MxN devices (to maintain a constant mini-batch size), the number of steps per epoch is reduced by a factor of M compared to N-way DP. Even further, in one embodiment, the number of epochs E M×N required is greater than or equal to E N In one embodiment, these factors all provide lower efficiency as the number of devices used in DP training increases.

[0051] In one embodiment, the application of MxN devices in a hybrid parallelization strategy of N-way DP, where each worker uses M-way MP, is applied. The speedup per step of each worker is denoted by SU M . Thus, the total training time / speedup is denoted as:

[0052]

[0053] In one embodiment, a comparison is made between hybrid N-way DP with M-way model-parallel workers and N-way DP with a single GPU worker. In one embodiment, the global batch size will remain the same. In one embodiment, in the MxN device configuration, each of the M devices is grouped as a single data-parallel worker. In one embodiment, the number of steps per epoch does not change because N and the number of steps per epoch of N-way DP at N M .

[0054] In one embodiment, Equations 4 and 5 are substituted into Equation 6, and as a result, conditions are determined that indicate that using hybrid parallelization (e.g., a combination of DP and MP) is better than scaling DP alone. In one embodiment, Equation 6 indicates that the application of a hybrid parallelization (e.g., a combination of DP and MP) strategy improves network training time if the speedup obtained from MP (for a given model-parallelization step) is large enough to overcome the scaling and statistical efficiency loss due to the increase in communication, synchronization overhead, and global batch size, respectively.

[0055]

[0056] In one embodiment, in a system, the MP implementation provides 45% and 65% improvement in training efficiency when using two and four GPUs, respectively. In one embodiment, using DP only is a good strategy, which scales well for up to 32 devices, after which the improvement in training time or training speedup slows down. In one embodiment, for this example, this enables a hybrid 32-way DP and 2-way MP hybrid parallelization strategy to perform better than 64-way DP, given the scaling and statistical efficiency loss at 64 devices.

[0057] In one example, similarly, when the devices scale from 32 devices to 128 devices, a hybrid 16-way DP and 4-way MP hybrid strategy outperforms using DP only. However, in this example, the performance of this hybrid strategy is less efficient than the hybrid strategy of 32-way DP and 2-way MP. In one embodiment, the per step speedup (SU 4 ) of 4-way MP cannot overcome the tradeoff (each data parallel worker uses four machines) as effectively as the per step speedup SU 2 of 2-way MP (each data parallel worker uses two machines). Thus, in one embodiment, the relative improvement over any device count is one factor that is provided prior to selecting a parallelization strategy, which is important for the training speedup obtained when scaling to a larger number of devices. In one embodiment, most of the time, the selection depends on the properties of the neural network and system configuration parameters as described above, and thus there is no one-size-fits-all solution or a solution that fits all to efficiently scale multi-device training.

[0058] Figure 7 An illustrative example of a process 700 for training a neural network is shown, in accordance with one embodiment. That is, in one embodiment, a neural network is trained using a first number of parallel threads of training data, which results in a first level of training efficiency 702. In one embodiment, the first number of parallel threads of training data is referred to as a global batch size as described above with respect to Figure 1 For example, the neural network is trained using a global batch on multiple parallel working computing devices (e.g., DP). In one embodiment, the first level of training efficiency indicates a training time when the neural network is trained using DP. In one embodiment, the first level of training efficiency is also referred to as an intermediate level of training efficiency.

[0059] In one embodiment, the neural network is trained using DP, taking into account the configuration parameters of the network and the size of the global batch, which results in a first level of training efficiency. In one embodiment, training the neural network using DP, which results in a first level of efficiency, is an indication that maximum training efficiency has been reached. If DP is further increased, the network will result in a decrease in efficiency due to the overhead of merging the results of all devices (nodes) running in parallel. In one embodiment, instead of reaching maximum efficiency, the first level of training efficiency is reached when, given a finite number of devices (nodes) available to train the network, the efficiency cannot get better. For example, although the network can gain some training efficiency gain by using DP alone, the gain can be small, and thus the network can determine that a greater gain is obtained by introducing MP (e.g., the gain is more significant here). As described in connection with Figure 1 In one embodiment, a front-end service or software tool (e.g., DLPiacer) is run in conjunction with the neural network that determines when to train the network using DP to reach the first level of efficiency.

[0060] In one embodiment, once the first level of efficiency has been met, then the front-end service or software tool (e.g., DLPiacer) determines to train the neural network using a second number of portions of the neural network in parallel with a first number of parallel training data threads, resulting in a second level of training efficiency 704. In one embodiment, the neural network is now trained in conjunction with using a first number of parallel training data threads and multiple portions of the neural network in parallel (e.g., using both DP and MP). In one embodiment, as described above in connection with Figure 1 In one embodiment, the multiple portions of the neural network are determined and applied such that the neural network is trained to reach the second level of efficiency. In one embodiment, the determination of the optimal combination of using DP and MP to train the network is performed by analytically testing different numbers of portions of the neural network and further applying different numbers of portions. In one embodiment, although the testing is performed analytically, the testing is also performed non-analytically.

[0061] In one embodiment, as described above in connection with Figure 6 In one embodiment, if the speedup obtained from MP (for a given model parallelization step) is large enough to overcome the scaling and statistical efficiency loss due to increased communication, synchronization overhead, and global batch size increase when using DP alone, then employing a hybrid parallelization (e.g., a combination of DP and MP) strategy allows the neural network to be trained with a second level of efficiency (which indicates an improvement in training time).

[0062] As a description Figure 7For the example of the process of training a neural network, the dataset used to train the neural network is first split into four subsets, where a first level of efficiency (e.g., maximum efficiency) is obtained: T1, T2, T3, and T4. In one embodiment, subsequently, when the first level of efficiency is met, portions of the neural network are split and applied such that training reaches a second level of efficiency (e.g., another maximum efficiency). In one embodiment, the neural network is implemented using a strategy that uses a combination of MP and DP with respect to the most efficient way to train the network. In one embodiment, the neural network is efficiently trained by splitting the neural network into three portions: M1, M2, and M3. In one embodiment, the hybrid parallelism approach then dictates that each of T1, T2, T3, and T4 is used to train each of M1, M2, and M3. In this example, training using 12 processors (four training subsets times three model components) would be the most efficient, with one for each combination of training subset and model component. That is, in one embodiment, the neural network is efficiently trained by changing how the training dataset is distributed among multiple computing devices and how the components of the neural network are split and distributed among multiple computing devices, based at least in part on the efficiency of training the neural network.

[0063] Figure 8 Another illustrative example of a process 800 for training a neural network is shown in accordance with one embodiment. That is, in one embodiment, a neural network is trained to reach or above a first level of training efficiency (e.g., maximum efficiency) by first determining a first number of parallel training data threads (e.g., using DP) 802. In one embodiment, as shown in Figure 1 the training data threads are global batch sizes. In one embodiment, the process 800 describes a neural network that is trained by adjusting a first type of parallelism (e.g., DP) to train the neural network based at least in part on a first training test set indicating a first level of training efficiency until the first level of training efficiency is reached. In one embodiment, this is done until the point at which any further training of the network using increased global batch size results in slower training. In one embodiment, the overhead required to communicate results with all devices (nodes) reaches the point at which it starts to become inefficient. In one embodiment, as described with respect to Figure 7 the first level of training efficiency that the neural network is trained to reach or above is not the maximum efficiency, but a level at which the efficiency gain the network obtains is lower compared to the introduction of MP.

[0064] As Figure 8As shown, in one embodiment, the neural network is further trained to reach or above a second level of efficiency 804 by determining a second number of portions of the neural network to be trained in parallel (e.g., using MP) while using a first number of parallel training data threads. In one embodiment, the neural network is further trained by implementing a second type of parallelism (e.g., hybrid parallelism including a combination of both DP and MP) to train the neural network based at least in part on a second training test set indicative of another level of training efficiency. In one embodiment, as indicated above with respect to Figure 1 and Figure 6 As indicated, training tests are performed to determine the number of MP to be introduced and further implemented in combination with DP to improve the training time of the training network. In one embodiment, the tests are performed analytically or non-analytically to determine the number of MP to be introduced. In one embodiment, if the test results of applying MP (for a given model parallelization step) are large enough to overcome the scaling and statistical efficiency loss due to increased communication, synchronization overhead and global batch size increase when using DP alone, then a hybrid parallelization (e.g., using a combination of DP and MP) strategy can improve the network training time. In one embodiment, as described with respect to Figure 1 As described, training efficiency indicates training time and is measured and represented as a value such that when it is analyzed, it provides information to a user or a separate system as to whether the neural network is being trained efficiently.

[0065] Figure 9 A parallel processing unit ("PPU") 900 according to one embodiment is shown. In one embodiment, PPU 900 is configured with machine-readable code that, if executed by PPU, causes PPU to perform certain or all of the processes and techniques described throughout this disclosure. In one embodiment, PPU 900 is implemented on one or more integrated circuit devices and utilizes multithreading as a latency-hiding technique designed to process computer-readable instructions (also referred to as machine-readable instructions or simply instructions) in parallel across multiple threads. In one embodiment, a thread refers to an execution thread and is an instance of an instruction set configured to be executed by PPU 900. In one embodiment, PPU 900 is an image processing unit ("GPU") configured to implement a graphics rendering pipeline for processing three-dimensional ("3D") graphics data in order to generate two-dimensional ("2D") image data for display on a display device (e.g., a liquid crystal display ("LCD") device). In one embodiment, PPU 900 is used to perform computations such as linear algebraic operations and machine learning operations. Figure 9Example parallel processors are shown for illustrative purposes only and should be interpreted as non-limiting examples of processor architectures contemplated within the scope of the present disclosure, and any appropriate processor can be employed in addition to and / or in lieu of the same.

[0066] In one embodiment, one or more PPUs are configured to accelerate high performance computing (“HPC”), datacenter, and machine learning applications. In one embodiment, PPU 900 is configured to accelerate deep learning systems and applications including the following non-limiting examples: autonomous vehicle platforms, deep learning, high-precision speech, image, text recognition systems, intelligent video analytics, molecular simulations, drug discovery, disease diagnosis, weather forecasting, big data analytics, astronomy, molecular dynamics simulations, financial modeling, robotics, factory automation, real-time language translation, online search optimization, and personalized user recommendations, among others.

[0067] In one embodiment, PPU 900 includes an input / output (“I / O”) unit 906, a front-end unit 910, a scheduler unit 912, a work distribution unit 914, a hub 916, a crossbar (“Xbar”) 920, one or more general processing clusters (“GPCs”) 918, and one or more partition units 922. In one embodiment, PPU 900 connects to a host processor or other PPUs 900 via one or more high-speed GPU interconnects 908. In one embodiment, PPU 900 connects to the host processor or other peripherals via an interconnect 908. In one embodiment, PPU 900 connects to a local memory comprising one or more memory devices 904. In one embodiment, local memory includes one or more dynamic random access memory (“DRAM”) devices. In one embodiment, one or more DRAM devices are configured and / or configurable as high bandwidth memory (“HBM”) subsystems, with multiple DRAM dies stacked

[0068] High-speed GPU interconnect 908 can refer to a wired-based, multi-lane communication link that systems use to scale and includes one or more PPUs 900 in conjunction with one or more CPUs, support cache coherency between PPUs 900 and CPUs, and CPU mastering. In one embodiment, high-speed GPU interconnect 908 sends data and / or commands through hub 916 to / from other units of PPU 900 (e.g., one or more copy engines, video encoders, video decoders, power management units, and other components not explicitly shown in FIG. 9). Figure 9

[0069] In one embodiment, I / O unit 906 is configured to perform reads and writes in conjunction with the host processor (not shown) via system bus 902. In one embodiment, I / O unit 906 communicates with the host processor via Figure 9 ​The I / O unit 906 sends and receives communications (e.g., commands, data) to and from the system bus 902. In one embodiment, the I / O unit 906 communicates with the host processor directly via the system bus 902, or through one or more intermediate devices such as a memory bridge. In one embodiment, the I / O unit 906 can communicate with one or more other processors (e.g., one or more PPUs 900) via the system bus 902. In one embodiment, the I / O unit 906 implements a Peripheral Component Interconnect Express (“PCIe”) interface for communications over a PCIe bus. In one embodiment, the I / O unit 906 implements an interface for communications with external devices.

[0070] In one embodiment, the I / O unit 906 decodes packets received via the system bus 902. In one embodiment, at least some of the packets represent commands configured to cause the PPU 900 to perform various operations. In one embodiment, the I / O unit 906 sends the decoded commands to various other units of the PPU 900 as specified by the commands. In one embodiment, the commands are sent to the front-end unit 910 and / or to the hub 916 or other units of the PPU 900 such as one or more copy engines, a video encoder, a video decoder, a power management unit, etc. (not explicitly shown in Figure 9 In one embodiment, the I / O unit 906 is configured to route communications between and among various logical units of the PPU 900.

[0071] In one embodiment, a program executed by the host processor encodes a command stream in a buffer that provides a workload to the PPU 900 for processing. In one embodiment, the workload includes instructions and data to be processed by those instructions. In one embodiment, the buffer is a region of memory that is accessible (e.g., read / write) by both the host processor and the PPU 900 - the host interface unit can be configured to access the buffer in system memory connected to the system bus 902 via memory requests sent by the I / O unit 906 over the system bus 902. In one embodiment, the host processor writes the command stream to the buffer and then sends a pointer to the beginning of the command stream to the PPU 900, causing the front-end unit 910 to receive the pointer to one or more command streams and manage the one or more streams, read commands from the streams, and forward the commands to various units of the PPU 900.

[0072] In one embodiment, front-end unit 910 is coupled to a scheduler unit 912, which is configured to assign respective GPCs 918 to process tasks defined by one or more workgroups. In one embodiment, scheduler unit 912 is configured to track state information related to various tasks managed by scheduler unit 912, where state information can indicate which GPC 918 a task is assigned to, whether the task is active or inactive, priority associated with a task, etc. In one embodiment, scheduler unit 912 manages execution of a plurality of tasks on one or more GPCs 918.

[0073] In one embodiment, scheduler unit 912 is coupled to a work distribution unit 914, which is configured to dispatch tasks for execution on GPCs 918. In one embodiment, work distribution unit 914 tracks a number of scheduled tasks received from scheduler unit 912, and work distribution unit 914 manages a pending task pool and an active task pool for each of GPCs 918. In one embodiment, the pending task pool includes a number of slots (e.g., 32 slots) that hold tasks assigned to be processed by a particular GPC 918; the active task pool can include a number of slots (e.g., 4 slots) for tasks that are actively being processed by a GPC 918, such that when a GPC 918 completes processing of a task, the task is evicted from the active task pool for the GPC 918, and one of other tasks from the pending task pool is selected and scheduled for execution on the GPC 918. In one embodiment, if an active task is idle on a GPC 918, for example, while waiting for data dependency, the active task is evicted from the GPC 918 and returned to the pending task pool, while another task from the pending task pool is selected and scheduled for execution on the GPC 918.

[0074] In one embodiment, work distribution unit 914 communicates with one or more GPCs 918 via XBar 920. In one embodiment, XBar 920 is an interconnect network coupling many units of PPU 900 to other units of PPU 900, and can be configured to couple work distribution unit 914 to a particular GPC 918. Although not explicitly shown, one or more other units of PPU 900 can also be connected to XBar 920 via hub 916.

[0075] Tasks are managed by a scheduler unit 912 and dispatched to GPCs 918 by a work distribution unit 914. GPCs 918 are configured to process tasks and generate results. The results can be consumed by other tasks in the GPC 918, routed to a different GPC 918 via an XBar 920, or stored in memory 904. The results can be written to memory 904 via a partition unit 922, which implements a memory interface for reading data from and writing data to memory 904. The results can be sent to another PPU 900 or CPU via a high-speed GPU interconnect 908. In one embodiment, the PPU 900 includes a number U of partition units 922, which is equal to the number of independent and distinct memories 904 coupled to the PPU 900. This will be discussed below in conjunction with Figure 11 The partition unit 922 is described in more detail.

[0076] In one embodiment, the host processor executes a driver kernel that implements an application programming interface (API) that enables one or more applications executing on the host processor to schedule operations for execution on the PPU 900. In one embodiment, the PPU 900 executes multiple computing applications simultaneously, and the PPU 900 provides isolation, quality of service ("QoS"), and independent address spaces for the multiple computing applications. In one embodiment, the application generates instructions (e.g., in the form of API calls) that cause the driver kernel to generate one or more tasks to be executed by the PPU 900, and the driver kernel outputs the tasks to one or more streams being processed by the PPU 900. In one embodiment, each task includes one or more related groups of threads, which may be referred to as warps. In one embodiment, a warp includes multiple related threads (e.g., 32 threads) that can be executed in parallel. In one embodiment, a cooperative thread may refer to a plurality of threads that include instructions to perform tasks and exchange data through shared memory. According to one embodiment, in combination Figure 11 Describes threads and cooperative threads in more detail.

[0077] Figure 10 According to one embodiment, Figure 9 900. In one embodiment, each GPC 1000 includes multiple hardware units for processing tasks, and each GPC 1000 includes a pipeline manager 1002, a pre-raster operations unit ("PROP") 1004, a raster engine 1008, a work distribution crossbar ("WDX") 1016, a memory management unit ("MMU") 1018, one or more data processing clusters ("DPCs") 1006, and any suitable combination of the parts. It will be understood thatFigure 10 GPC 1000 can include other hardware units in place of or in addition to the units illustrated in FIG. 10. Figure 10

[0078] In one embodiment, the operation of GPC 1000 is controlled by pipeline manager 1002. Pipeline manager 1002 manages configuration of one or more DPCs 1006 for processing tasks allocated to GPC 1000. In one embodiment, pipeline manager 1002 configures at least one of one or more DPCs 1006 to implement at least a portion of a graphics rendering pipeline. In one embodiment, DPC 1006 is configured to execute vertex shader programs on the programmable streaming multi-processor (“SM”) 1014. Pipeline manager 1002 is configured to route packets received from a work distribution unit to the appropriate logic within GPC 1000, and in one embodiment, some packets can be routed to fixed function hardware units in PROP 1004 and / or raster engine 1008 while other packets can be routed to DPCs 1006 for processing by the geometry engine 1012 or SM 1014. In one embodiment, pipeline manager 1002 configures at least one of one or more DPCs 1006 to implement a neural network model and / or compute pipeline.

[0079] In one embodiment, PROP unit 1004 is configured to route data generated by raster engine 1008 and DPCs 1006 to a memory partition unit, a render output pipeline (“ROP”) unit, as described in more detail above. In one embodiment, PROP unit 1004 is configured to perform optimizations for color blending, organize pixel data, perform address translations, and / or the like. In one embodiment, raster engine 1008 includes a number of fixed function hardware units to perform various raster operations; and the raster engine 1008 includes a setting engine, a coarse raster engine, a cull engine, a clip engine, a fine raster engine, a tile merge engine, and any suitable combination thereof. In one embodiment, the setting engine receives transformed vertices and generates a plane equation associated with a geometric primitive defined by the vertices; the plane equation is sent to the coarse raster engine to produce coverage information (e.g., x, y coverage masks for tiles) of the primitive; the output of the coarse raster engine is sent to the cull engine where fragments associated with primitives that fail a z-test are culled; and fragments that pass the z-test are sent to the clip engine where fragments that are outside the view frustum range are culled. In one embodiment, fragments that survive clipping and culling are passed to the fine raster engine to generate attributes for the pixel fragments based on the plane equation generated by the setting engine. In one embodiment, the output of the raster engine 1008 includes fragments to be processed by any suitable entity, such as a fragment shader implemented within a DPC 1006. ​

[0080] In one embodiment, each DPC 1006 included in GPC 1000 includes an M-Pipe controller (“MPC”) 1010; a primitive engine 1012; one or more SMs 1014; and any suitable combination thereof. In one embodiment, MPC 1010 controls the operation of DPC 1006, routing received commands from pipeline manager 1002 to appropriate units in DPC 1006. In one embodiment, commands associated with a vertex are routed to primitive engine 1012, which is configured to fetch vertex attributes associated with the vertex from memory; in contrast, commands associated with a shader program are sent to SM 1014.

[0081] In one embodiment, SM 1014 includes a programmable streaming processor that is configured to process tasks represented by a number of threads. In one embodiment, SM 1014 is multi-threaded and configured to execute a plurality of threads (e.g., 32 threads) from a particular warping of threads concurrently and implements a SIMD (Single Instruction, Multiple Data) architecture wherein each thread of the warping of threads is configured to process a different data set based on the same set of instructions. In one embodiment, all threads in a warp execute identical instructions. In another embodiment, individual threads of a warp can diverge from execution of identical instructions based on a predicate output by a previous instruction and only threads for which the predicate is true will execute subsequent instructions. In one embodiment, the threads in the warp can be configured to execute the same instruction simultaneously. In one embodiment, SM 1014 implements a SIMT (Single Instruction, Multiple Thread) architecture wherein each thread of the warping of threads is configured to process a different data set based on the same set of instructions, but where individual threads of the warping of threads can diverge in execution depending on the results of predication instructions or other instructions. In one embodiment, program counters, call stacks and execution states for each thread are maintained such that threads within a warp can be executed in parallel with no dependencies between threads of the warp. In another embodiment, program counters, call stacks and execution states for each individual thread are maintained such that threads within a warp and between warps can be executed with dependencies between individual threads when instructions for the threads generate data that other threads can use. In one embodiment, threads execute instructions concurrently by reading instructions from a common instruction cache and by reading data from a common data cache.

[0082] In one embodiment, MMU 1018 provides an interface between GPC 1000 and memory partition unit, and MMU 1018 provides translations of virtual addresses into physical addresses, memory protection, and arbitration of memory requests. In one embodiment, MMU 1018 provides one or more translation lookaside buffers (TLBs) for handling translation of virtual addresses into physical addresses.

[0083] Figure 11A memory partition unit of a PPU is shown in accordance with one embodiment. In one embodiment, the memory partition unit 1100 includes a raster operations ("ROP") unit 1102; a level two ("L2") cache 1104; a memory interface 1106; and any suitable combination thereof. The memory interface 806 is coupled to memory. The memory interface 1106 can implement a 32-, 64-, 128-, 1024-bit data bus, etc., for high-speed data transfer. In one embodiment, the PPU includes U memory interfaces 1106, one for each pair of partition units 1100, where each pair of partition units 1100 is connected to a corresponding memory device. For example, the PPU can be connected to up to Y memory devices, such as high bandwidth memory stacks or graphics double data rate version 5 synchronous dynamic random access memory ("GDDR5 SDRAM").

[0084] In one embodiment, the memory interface 1106 implements an HBM2 memory interface, and Y is equal to half of U. In one embodiment, the HBM2 memory stacks are located on the same physical package as the PPU, saving a significant amount of power and area compared to a conventional GDDR5 SDRAM system. In one embodiment, each HBM2 stack includes four memory dies, and Y is equal to 4, while the HBM2 stack includes two 128-bit channels per die, for a total of 8 channels, and the data bus width is 1024 bits.

[0085] In one embodiment, the memory supports single error correction double error detection ("SECDED") error-correcting code ("ECC") to protect data. ECC provides higher reliability for compute applications that are sensitive to data corruption. Reliability is especially important in PPU processing very large data sets and / or large-scale cluster computing environments running applications for long periods of time.

[0086] In one embodiment, the PPU implements a multi-level memory hierarchy. In one embodiment, the memory partition unit 1000 supports a unified memory to provide a single unified virtual address space for the memory of CPUs and GPGPUs, enabling data sharing between GPGPUs and CPUs running on different processors. In one embodiment, the frequency of PPU accesses to pages residing on other processors' memories is tracked such that memory pages frequently accessed by a PPU are moved to the PPU's physical memory. In one embodiment, the high-speed GPU interconnect 908 supports address translation services, which allow a PPU to access pages resident in a CPU's memory space. In one embodiment, the address translation services provide full access to all pages, including supervisor and OS pages, thus providing the PPU with full access to the operating system and other processes.

[0087] In one embodiment, the copy engine transfers data between multiple PPUs or between a PPU and a CPU. In one embodiment, the copy engine can generate a page fault for an address that is not mapped into a page table, and then the memory partition unit 1100 services the page fault, maps the address into a page table, after which the copy engine performs the transfer. In one embodiment, the memory is pinned (i.e., not pageable) for multiple copy engine operations between multiple processors, thus significantly reducing the available memory. In one embodiment, due to the hardware page fault, the address can be passed to the copy engine without having to consider whether the memory page is resident, and the copy process is transparent.

[0088] According to one embodiment, data from Figure 9 memory or other system memory is fetched by the memory partition unit 1100 and stored in an L2 cache 1104, which is on-chip and shared among the various GPCs. In one embodiment, each memory partition unit 1100 includes at least a portion of the L2 cache 1060 associated with a corresponding memory device. In one embodiment, lower level caches are implemented in the various units within a GPC. In one embodiment, each SM 1140 can implement a level one (“Ll”) cache, where the Ll cache is private per SM 1140 and is a private memory for a particular SM 1140 and fetches data

[0089] In one embodiment, the ROP unit 1102 performs graphics raster operations including, for example, color compression, pixel blending, etc. In one embodiment, the ROP unit 1102 incorporates depth testing alongside the raster engine 1125, receiving depth values for sample positions associated with pixel fragments from a culling engine of the raster engine 1125. In one embodiment, the depth for a sample position associated with a fragment is tested against a corresponding depth in a depth buffer. In one embodiment, if the fragment passes the depth test for the sample position, the ROP unit 1102 updates the depth buffer and sends the results of the depth test to the raster engine 1125. It will be appreciated that the number of partition units 1100 can be different from the number of GPCs, and thus, in one embodiment, each ROP unit 1102 can be coupled to each GPC. In one embodiment, the ROP unit 1102 tracks the packets received from different GPCs and determines to which GPC to route results generated by the ROP unit 1102 via the Xbar.

[0090] Figure 12 FIG. 1 1 shows a graphics processor 1100 according to one embodiment, such as graphics processor 1100 of FIG. 10, in accordance with one embodiment of the application.Figure 10 Stream Multiprocessors. In one embodiment, the SM 1200 comprises: an instruction cache 1202; one or more scheduler units 1204; a register file 1208; one or more processor cores 1210; one or more special-function units ("SFUs") 1212; one or more load / store units ("LSUs") 1214; an interconnect network 1216; shared memory / L1 cache 1218; and any suitable combination thereof. In one embodiment, a work distribution unit dispatches tasks for execution on GPCs of a PPU, and each task is assigned a particular DPC within a GPC and, if the task is associated with a shader program, the task is assigned to a SM 1200. In one embodiment, the scheduler unit 1204 receives tasks from the work distribution unit and manages instruction scheduling for one or more thread blocks assigned to the SM 1200. In one embodiment, the scheduler unit 1204 schedules thread blocks for execution as thread warps of parallel threads, with each thread block being assigned at least one thread warp. In one embodiment, each thread warp executes a thread. In one embodiment, the scheduler unit 1204 manages a plurality of different thread blocks, assigns warps to different thread blocks, and then dispatches instructions from different ones of the plurality of different cooperative groups to various functional units (e.g., the cores 1210, the SFUs 1212, and the LSUs 1214) during each clock cycle.

[0091] A cooperative group can refer to a programming model for organizing groups of communicating threads that allows developers to express the granularity of threads that are communicating, enabling richer expressiveness, more efficient parallel decomposition. In one embodiment, a cooperative launch API supports synchronization between thread blocks to execute parallel algorithms. In one embodiment, an application of a conventional programming model provides a single simple construct for synchronizing cooperating threads: a barrier across all threads of a thread block (e.g., a syncthreads() function). However, programmers often want to define thread groups at a granularity smaller than a thread block and synchronize within the defined groups to achieve higher performance, design flexibility, and reuse of software in the form of collective group-wide functional interfaces. Cooperative groups enable programmers to explicitly define thread groups at sub-block (i.e., down to a single thread) and multi-block granularity and to perform collective operations such as synchronization on threads in a cooperative group. The programming model supports clean composition across software boundaries, so library and utility functions can safely synchronize in their local context without having to make assumptions about convergence. Cooperative group primitives enable new patterns of cooperative parallelism, including producer-consumer parallelism, opportunistic parallelism, and global synchronization across a grid of thread blocks.

[0092] In one embodiment, the dispatch unit 1206 is configured to send instructions to one or more functional units, and the scheduler unit 1204 includes two dispatch units 1206 that enable two different instructions from the same thread bundle to be dispatched during each clock cycle. In one embodiment, each scheduler unit 1204 includes a single dispatch unit 1206 or an additional dispatch unit 1206.

[0093] In one embodiment, each SM 1200 includes a register file 1208 that provides a set of registers for the functional units of the SM 1200. In one embodiment, the register file 1208 is partitioned between functional units, with each functional unit being allocated a dedicated portion of the register file 1208. In one embodiment, the register file 1208 is partitioned between different thread bundles being executed by the SM 1200 with each thread bundle being allocated dedicated storage for

[0094] According to one embodiment, the tensor cores are configured to perform matrix operations. In one embodiment, one or more tensor cores are included in the cores 1210. In one embodiment, the tensor cores are configured to perform deep learning matrix arithmetic, such as convolution operations for neural network training and inference. In one embodiment, each tensor core operates on 4x4 matrix operations, and performs a matrix multiply and accumulate operation D = AxB + C, where A, B, C, and D are 4x4 matrices.

[0095] In one embodiment, the matrix multiply inputs A and B are 16-bit floating point matrices, and the accumulate matrices C and D are 16-bit floating point or 32-bit floating point matrices. In one embodiment, a tensor core performs 32-bit floating point accumulation operations on 16-bit floating point input data. In one embodiment, the 16-bit floating point multiplication requires 64 operations and produces a full precision product which is then accumulated using 32-bit floating point addition with other intermediate products for a 4x4x4 matrix multiply. In one embodiment, tensor cores are used to perform larger two-dimensional or higher dimensional matrix operations composed of these smaller elements. In one embodiment, an API such as CUDA 9 C++ API exposes specialized matrix load, matrix multiply and accumulate, and matrix store operations to efficiently use tensor cores from a CUDA-C++ program. In one embodiment, at the CUDA level, the warp-level interface assumes 16x16 size matrices across all 32 threads of a warp.

[0096] In one embodiment, each SM 1200 includes M SFUs 1212 that perform special functions, for example, certain mathematical functions, bit and byte operations, etc. In one embodiment, SFU 1212 includes tree traversal units configured to traverse a hierarchical tree data structure. In one embodiment, SFU 1212 includes a texture unit configured to perform texture map filtering operations. In one embodiment, the texture unit is configured to load a texture map (e.g., a 2D array of texels) from memory and sample the texture map to produce sampled texture values for use in a shader program executed by SM 1200. In one embodiment, the texture map is stored in shared memory / L1 cache. According to one embodiment, the texture unit implements texture operations, for example, filtering operations using mipmaps (e.g., texture maps of different levels of detail). In one embodiment, each SM 1200 includes two texture units.

[0097] In one embodiment, each SM 1200 includes N LSUs 1214 that implement load and store operations between shared memory / L1 cache 1218 and register file 1208. In one embodiment, each SM 1200 includes an interconnect network 1216 that connects each of the functional units to register file 1208 and LSUs 1214 to register file 1208, shared memory / L1 cache 1218. In one embodiment, interconnect network 1216 is a cross-bar switch that can be configured to connect any of the functional units to any of the registers in register file 1208 and connect LSUs 1214 to registers in register file and memory locations in shared memory / L1 cache 1218.

[0098] In one embodiment, shared memory / L1 cache 1218 is an array of on-chip memory that allows data storage and communication between threads within the SM 1200 and between the SM 1200 and the graphics processing engine. In one embodiment, shared memory / L1 cache 1218 includes 128 KB of storage capacity and is located on the path from the SM 1200 to the partition units. In one embodiment, shared memory / L1 cache 1218 is used for caching reads and writes. One or more of shared memory / L1 cache 1218, L2 cache, and memory are backed up by main memory.

[0099] In one embodiment, combining data cache and shared memory functionality into a single memory block provides improved performance over using separate memory blocks for the two functions. In one embodiment, this capacity is used as either a cache or not, depending on whether shared memory is configured, for example if shared memory is configured to use half the capacity, textures and load / store operations can use the remaining capacity. According to one embodiment, the integration within shared memory / L1 cache 1218 enables shared memory / L1 cache 1218 to be used as a high-throughput pipe for streaming data to the SM 1200, while also providing high bandwidth and low latency access to frequently reused data. When configured for general parallel compute, simpler configurations can be used compared to graphics processing. In one embodiment, fixed function graphics processing units are bypassed, creating a simpler programming model. In one embodiment, within general parallel compute configurations, the work distribution unit allocates and dispatches blocks of threads directly to DPCs. According to one embodiment, threads in a block execute the same program, with each thread using a unique thread ID to perform computations, the SM 1200 executes programs and performs computations, the shared memory / L1 cache 1218 is used for communication between threads, and the LSU 1214 reads and writes from / to global memory through shared memory / L1 cache 1218 and memory partition units. In one embodiment, when configured for general parallel compute, the SM 1200 writes commands that the scheduler unit can use to launch new work on DPCs.

[0100] In one embodiment, the PPU is included in or coupled to a desktop computer, a laptop computer, a tablet computer, a server, a supercomputer, a smart-phone (e.g., a wireless, hand-held device), a personal digital assistant (“PDA”), a digital camera, a vehicle, a head-mounted display, a hand-held electronic device, etc. In one embodiment, the PPU is embodied on a single semiconductor substrate. In one embodiment, the PPU is included in a system-on-a-chip (“SoC”) along with one or more other devices (e.g., additional PPUs, memory, reduced instruction set computer (“RISC”) CPUs, memory management units (“MMUs”), digital-to-analog converters (“DACs”), etc.).

[0101] In one embodiment, the PPU can be included on a graphics card that includes one or more memory devices. The graphics card can be configured to interface with a PCIe slot on a motherboard of a desktop computer. In another embodiment, the PPU can be an integrated graphics processing unit (“iGPU”) included in a chipset of the motherboard.

[0102] Figure 13 A computer system 1300 according to one embodiment in which various architectures and / or functionalities can be implemented is shown. In one embodiment, the computer system 1300 is configured to implement the various processes and methods described throughout this disclosure.

[0103] In one embodiment, the computer system 1300 includes at least one central processing unit 1302 connected to a communication bus 1310, which employs any suitable protocol, for example, PCI (Peripheral Component Interconnect), PCI-Express, AGP (Accelerated Graphics Port), HyperTransport, or any other bus or point-to-point communication protocol. In one embodiment, the computer system 1300 includes a main memory 1304 and control logic (e.g., implemented in hardware, software, or a combination thereof) and data are stored in the main memory 1304, which can take the form of random access memory (“RAM”). In one embodiment, a network interface subsystem 1322 provides an interface to other computing devices and networks for receiving data from or transmitting data to other systems from the computer system 1300.

[0104] In one embodiment, computer system 1300 includes input device 1308, parallel processing system 1312, and display device 1306, which can be implemented using a conventional CRT (cathode ray tube), LCD (liquid crystal display), LED (light emitting diode), plasma display or other suitable display technologies. In one embodiment, user input is received from input device 1308 such as a keyboard, mouse, touchpad, microphone, etc. In one embodiment, each of the aforementioned modules can be located on a single semiconductor platform to form a processing system.

[0105] In this specification, a single semiconductor platform can refer to a sole unitary semiconductor-based integrated circuit or chip as standing alone (singular form) or in inclusion with other parts (plural form). It should be noted that the term semiconductor platform is also meant to encompass a multi-chip module with increased connectivity which simulate a single semiconductor platform with respect to operation and functionality, and is also meant to encompass a system on a chip (SoC) with increased connectivity which simulates a single semiconductor platform with respect to operation and functionality.

[0106] In one embodiment, computer programs in the form of machine-readable executable code or computer-control logic algorithms are stored in main memory 1304 and / or secondary memory. The computer programs, if executed by one or more processors, enable the system 1300 to perform various functions. The memory, storage, and / or any other storage is a possible example of computer-readable media. Secondary memory can refer to any suitable storage device or system including a hard disk drive and / or a removable storage drive, representing a floppy disk drive, a magnetic tape drive, an optical storage drive, a digital versatile disk ("DVD") drive, a recording device, a universal serial bus ("USB") flash drive, among others.

[0107] In one embodiment, the architecture and / or functionality of the various previous figures are implemented as: central processing unit 1302; parallel processing system 1312; integrated circuit(s) having at least a portion of the functionality of both a central processing unit 1302 and parallel processing system 1312; a chip set including central processing unit 1302 and parallel processing system 1312 each designed to work as a single unit and sold as such; and any suitable combination thereof.

[0108] In one embodiment, the architecture and / or functionality of the various previous figures is implemented as a universal computer system, a circuit board system, a game console system specifically designed for entertainment purposes, a dedicated system, etc. In one embodiment, computer system 1300 can take the form of a desktop computer, laptop computer, tablet computer, server computer, super computer, smart-phone (e.g., a wireless handheld device), personal digital assistant ("PDA"), digital cameras, vehicles, head mounted display, hand-held electronic device, mobile telephone device, television, workstation, game console, embedded system, and / or any other type of logic.

[0109] In one embodiment, parallel processing system 1312 includes a plurality of PPUs 1314 and associated memories 1316. In one embodiment, the PPUs are connected to a host processor or other peripheral devices via interconnect 1318 and switch 1320 or multiplexer. In one embodiment, parallel processing system 1312 distributes compute tasks across the PPUs 1314 that are parallelizable— e.g., as part of distributing a compute task across a plurality of GPU thread blocks. In one embodiment, although shared memory is shared and accessed among some or all of the PPUs 1314 (e.g., for read and / or write access), such shared memory can result in performance loss relative to using local memory and registers resident to the PPU. In one embodiment, the operation of the PPUs 1314 are synchronized by use of commands such as __syncthreads(), which require all threads in a block (e.g., executing across multiple PPUs 1314) to reach a certain point of execution in the code before moving on.

[0110] Therefore, the specification and drawings should be considered as illustrative only and not in a limiting sense. For example, while the disclosed technology is described in the context of a particular implementation, it is clear that various modifications and changes can be made thereto without departing from the broader spirit and scope of the invention as set forth in the claims.

[0111] Other variations are within the spirit of the present disclosure. Thus, while the disclosed technology is susceptible to various modifications and alternative constructions, certain illustrated embodiments thereof are shown in the drawings and have been described above in detail. It should be understood, however, that there is no intention to limit the invention to the specific form or forms disclosed, but on the contrary, the intention is to cover all modifications, alternative constructions, and equivalents falling within the spirit and scope of the invention, as defined in the appended claims.

[0112] The use of the terms "a" and "an" and "the" and similar referents in the context of describing the disclosed embodiments (especially in the context of the following claims) are to be construed to cover both the singular and the plural, unless otherwise indicated herein or clearly contradicted by context. The terms "comprising," "having," "including," and "containing" are to be construed as open-ended terms (i.e., meaning "including, but not limited to,") unless otherwise noted. The term "connected" as used in the context of the present disclosure, is to be construed as partially or wholly contained within, connected or joined to, or otherwise being adjacent to, even if there is anything intervening. Unless specifically stated otherwise, the description herein of a numerical range is meant to include each individual number within that range, and is meant to be shorthand for "a range from the lower limit of the range to the upper limit of the range." Unless otherwise indicated herein, the use of the term "set" (e.g., "set of items") or "subset" is to be construed as a non-empty set having one or more members. Also, unless otherwise indicated herein or otherwise apparent from context, the term "subset" of a respective set does not necessarily denote a proper subset of the respective set, but rather the subset and the respective set can be equal.

[0113] Unless specifically stated otherwise, or otherwise clear from context, combinational language such as "at least one of A, B, and C" or "at least one of A, B, and C" is to be construed to cover all cases, e.g., wherein A is actually X, B is actually Y, C is actually Z, A is actually X, B is actually X, C is actually Y, A is actually X, B is actually Y, C is actually Y, A is actually X, B is actually Y, C is actually Z, and so on. For example, in the case of an illustrative example having a set of three members, the conjunction phrase "at least one of A, B, and C" and "at least one of A, B, and C" refers to any of the following sets: {A}, {B}, {C}, {A, B}, {A, C}, {B, C}, {A, B, C}. Thus, this conjunctive language is generally not intended to imply certain embodiments require at least one of A, at least one of B, and at least one of C to each be present. Further, unless otherwise noted, the term "plurality" denotes a state of plurality (e.g., "a plurality of items" denotes a plurality of items). The number of items in the plurality is at least two, but can be more when explicitly or by context indicated. Also, unless otherwise noted or clearly contradicted by context, the phrase "based on" means "based, at least in part, on" rather than "based solely on."

[0114] The operations of a process described herein can be performed in any suitable order unless otherwise indicated herein or otherwise clearly contradicted by context. In one embodiment, a process such as those described herein (or variations and / or combinations thereof) is performed under the control of one or more computer systems configured with executable instructions to perform the operations of the process. In one embodiment, executable instructions are stored on a computer-readable storage medium that is non-transitory in nature, such as a hard drive, read only memory, etc. In one embodiment, the computer system is configured with the executable instructions to perform the operations of the process, such as those described herein. In one embodiment, the computer system is configured with the executable instructions to perform the operations of the process, such as those described herein, collectively or individually. In one embodiment, the computer system is configured with the executable instructions to perform the operations of the process, such as those described herein, collectively or individually, in conjunction with hardware and / or software of the computer system.

[0115] In one embodiment, accordingly, a computer system is configured to implement one or more services that individually or collectively perform the operations of the processes described herein, and such a computer system is configured with the applicable hardware and / or software to enable the operations to be performed. Moreover, a computer system implementing an embodiment of the present disclosure is a single device, and in another embodiment it is a distributed computer system comprising multiple devices operating in different ways such that the distributed computer system performs the operations described herein and such that no single device is capable of performing all of the operations.

[0116] The use of any and all examples, or exemplary language (e.g., "such as") provided herein, is intended merely to better illuminate embodiments of the application and does not pose a limitation on the scope of the application unless otherwise claimed. No language in the specification should be construed as indicating any non-claimed element as essential to the practice of the application.

[0117] Embodiments of the present disclosure are described herein, including the best mode known to the inventors for practicing the application. Variations of those embodiments will become apparent to those of ordinary skill in the art upon reading the foregoing description. The inventors expect skilled artisans to employ such variations as appropriate, and the inventors intend for the scope of the present disclosure to include all such modifications and equivalents of the subject matter recited in the claims appended hereto as permitted by applicable law. Moreover, additional

[0118] All references cited herein, including publications, patent applications, and patents, are incorporated by reference as if each were individually and specifically indicated to be incorporated by reference. The disclosure herein of any

[0119] In the description and claims, the terms "coupled" and "connected," along with derivatives thereof, can be used. It should be understood that these terms are not intended as synonyms for each other. Rather, in particular embodiments, "connected" or "coupled" is used to indicate that two or more elements are in direct or indirect physical or electrical contact with each other. "Coupled" can also mean that two or more elements are not in direct contact with each other, but yet still co-operate or interact with each other.

[0120] Unless specifically stated otherwise, it can be appreciated that throughout the specification terms such as "processing," "computing," "calculating," "determining," or the like, refer to the action and / or processes of a computer or computing system, or similar electronic computing device, that manipulates and / or transforms data represented as physical quantities (e.g., electronic quantities) within the computer's or computing system's registers and / or memories into other data similarly represented as physical quantities within the computer's or computing system's memories, registers, or other such information storage, transmission, or display devices.

[0121] In a similar manner, the term“processor” can refer to any device or portion of a device that processes electronic data from registers and / or memory to transform that electronic data into other electronic data that can be stored in registers and / or memory. As a non-limiting example,“processor” can refer to a central processing unit (CPU), or a graphics processing unit (GPU). A“computing platform” can include one or more processors. As used herein, a“software” process can include, for example, software and / or hardware entities such as tasks, threads, and intelligent agents that perform work over time. Also, each process can refer to multiple processes that execute instructions either continuously or intermittently, either sequentially or in parallel. The terms“system” and“method” are used interchangeably herein, as long as the system can embody one or more methods and the method can be considered a system.

[0122] In this document, references to obtaining, acquiring, receiving, or inputting analog or digital data can refer to obtaining, acquiring, receiving, or inputting the analog or digital data into a subsystem, computer system, or computer-implemented machine. The process of obtaining, acquiring, receiving, or inputting analog and digital data can be accomplished in a variety of ways, such as by receiving the data as a parameter of a function call or a call to an application programming interface, in some implementations, the process of obtaining, acquiring, receiving, or inputting analog or digital data can be accomplished by transmitting the data via a serial or parallel interface. In another implementation, the process of obtaining, acquiring, receiving, or inputting analog or digital data can be accomplished by transmitting the data from a providing entity to an acquiring entity via a computer network. References to providing, outputting, transmitting, sending, or presenting analog or digital data can also be made. In various examples, the process of providing, outputting, transmitting, sending, or presenting analog or digital data can be accomplished by transmitting the data as an input or output parameter of a function call, a parameter of an application programming interface, or an interprocess communication mechanism.

[0123] Although the above discussion discusses example implementations of the described technology, other architectures can be used to implement the described functionality and are intended to be within the scope of this disclosure. Moreover, although specific

[0124] Furthermore, although the subject matter has been described in language specific to structural features and / or methodological acts, it is to be understood that the subject matter defined in the appended claims is not necessarily limited to the specific features or acts described. Rather, the specific features and acts are disclosed as illustrative forms of implementing the claims.

Claims

1. One or more processors comprising: circuitry to train a machine learning model by iteratively switching between applying a machine learning model parallelism technique and applying a data parallelism technique until a target level of training efficiency is reached based on comparing a training time associated with applying the machine learning model parallelism technique and a training time associated with applying the data parallelism technique.

2. The one or more processors of claim 1, wherein the machine learning model is trained by increasing the data parallelism technique until an intermediate level of training efficiency is reached and increasing the model parallelism technique until an intermediate level of training efficiency is reached.

3. The one or more processors of claim 1, wherein the intermediate level of training efficiency is measured based at least in part on a training time associated with using the data parallelism technique and the target level of training efficiency is measured based at least in part on a training time associated with using a combination of the data parallelism technique and the machine learning model parallelism technique.

4. The one or more processors of claim 1, wherein the machine learning model is trained by: comparing a training time associated with using the data parallelism technique and a training time associated with using a combination of the data parallelism technique and the machine learning model parallelism technique; and based on the comparison, using the combination of the data parallelism technique and the machine learning model parallelism technique.

5. The one or more processors of claim 1, wherein the machine learning model is trained by: using a combination of the data parallelism technique and the machine learning model parallelism technique within one or more layers of the machine learning model to reach the target level of training efficiency.

6. The one or more processors of claim 1, wherein the circuitry is to use the machine learning model to infer information, wherein the machine learning model is trained by using a data parallelism technique followed by the machine learning model parallelism technique.

7. A system comprising: one or more computers having one or more processors to train a machine learning model by iteratively switching between applying a machine learning model parallelism technique and applying a data parallelism technique until a target level of training efficiency is reached based on comparing a training time associated with applying the machine learning model parallelism technique and a training time associated with applying the data parallelism technique.

8. The system of claim 7, wherein a first number of threads of training data are split into subsets to be distributed among the one or more processors to train the machine learning model.

9. The system of claim 7, wherein a second number of portions of the machine learning model being trained are split into components to be distributed among the one or more processors to train the machine learning model.

10. The system of claim 7, wherein the one or more computers with the one or more processors further train the machine learning model by increasing a first number of training data threads until the target level of training efficiency is reached.

11. The system of claim 7, wherein the one or more computers with the one or more processors further train the machine learning model by: comparing a training time associated with using a first number of training data threads to a training time associated with using a second number of portions of the machine learning model trained in parallel using the first number of training data threads; and based on the comparison, using the second number of portions of the machine learning model trained in parallel using the first number of training data threads.

12. The system of claim 7, wherein a second number of portions of the machine learning model are trained in parallel using a first number of training data threads.

13. The system of claim 7, wherein the target level of training efficiency is based at least in part on information for power consumption of the system.

14. A non-transitory machine-readable medium having stored thereon a set of instructions which, if executed by one or more processors, cause the one or more processors to at least: train a machine learning model to be iteratively adjusted to continually improve performance of the machine learning model by iteratively switching between applying machine learning model parallelism techniques and applying data parallelism techniques based on comparing a training time associated with applying machine learning model parallelism techniques to a training time associated with applying data parallelism techniques until a target level of training efficiency is reached.

15. The non-transitory machine-readable medium of claim 14, wherein a first level of training efficiency is based at least in part on a training time associated with training the machine learning model using a first number of parallel training data threads.

16. The non-transitory machine-readable medium of claim 14, wherein a second level of training efficiency indicates a training time associated with training the machine learning model in parallel using a first number of parallel training data threads in combination with a second number of portions of the machine learning model.

17. The non-transitory machine-readable medium of claim 14, wherein the set of instructions further cause the one or more processors to at least train the machine learning model by adjusting a first number of parallel training data threads until the target level of training efficiency is reached.

18. The non-transitory machine-readable medium of claim 14, wherein the set of instructions further cause the one or more processors to at least train the machine learning model by: comparing a training time associated with using a first number of parallel training data threads to a training time associated with using a second number of portions of the machine learning model in parallel with the first number of parallel training data threads; and training the machine learning model using the second number of portions of the machine learning model in parallel with the first number of parallel training data threads based on the comparison.

19. The non-transitory machine-readable medium of claim 14, wherein the set of instructions further cause the one or more processors to train the machine learning model using at least the first number of parallel training data threads prior to training the machine learning model using the second number of portions of the machine learning model in parallel with the first number of parallel training data threads.

20. A computer-implemented method comprising: training a machine learning model by iteratively switching between applying a machine learning model parallelism technique and applying a data parallelism technique until a target level of training efficiency is reached based on comparing a training time associated with applying the machine learning model parallelism technique to a training time associated with applying the data parallelism technique.

21. The method of claim 20, wherein a first number of parallel training data threads is split into subsets to be distributed among one or more processors of a computer system to train the machine learning model.

22. The method of claim 20, wherein a second number of portions of the machine learning model to be trained in parallel is split into components to be distributed among one or more processors of a computer system to train the machine learning model.

23. The method of claim 20, wherein a first number of parallel training data threads to train the machine learning model is increased until the machine learning model is trained to reach or exceed a first level of training efficiency, wherein the first level of training efficiency is based at least in part on a training speedup associated with training the machine learning model using the first number of parallel training data threads.

24. The method of claim 23, wherein responsive to the machine learning model being trained to reach or exceed the first level of training efficiency, a second number of portions of the machine learning model to be trained in parallel using the first number of parallel training data threads is determined.

25. The method of claim 20, wherein a second number of portions of the machine learning model to be trained in parallel using a first number of parallel training data threads is based at least in part on a training speedup associated with training the machine learning model using the second number of portions.

Images