Academia.eduAcademia.edu

Figure from:Code Generation and Runtime Techniques for Enabling Data Efficient Deep Learning Training on GPUs

See full PDFdownloadDownload figure
Figure 3.4: Inefficiency (in red) exists in all layers of existing systems.

Figure 3 4: Inefficiency (in red) exists in all layers of existing systems.

Related Figures (55)

inter-device (D2D) bandwidth, PCle bandwidth, memory bandwidth, and
inter-device (D2D) bandwidth, PCle bandwidth, memory bandwidth, and
Figure 1.2: Comparison of the memory bandwidth and FP16 throughput of Nvidia B100 SXM [13] with the arithmetic intensity of Google internal production workloads [7]. capacity, and their growth has kept up with these demands. Section 5.2.1 de- tails the reasoning. Given this, we choose to offload tensors to SSDs for LLM training to overcome GPU memory limitations. Nevertheless, SSD band- width is limited, and the gap between SSD bandwidth growth and GPU com- putational throughput growth continues to widen, as shown in Figures 1.1 and 1.3. It is essential to carefully manage data transfers to prevent training throughput from being constrained by SSD bandwidth. We propose tech- niques for selecting which tensors to offload and hiding transfer latency. To evaluate the trade-off between performance and memory savings, we com- pare different strategies, i.e., tensor recomputation, offloading, and keeping tensors in GPU memory. These are detailed in Chapter 5. Despite the importance of data efficiency. several obstacles exist to address
Figure 1.2: Comparison of the memory bandwidth and FP16 throughput of Nvidia B100 SXM [13] with the arithmetic intensity of Google internal production workloads [7]. capacity, and their growth has kept up with these demands. Section 5.2.1 de- tails the reasoning. Given this, we choose to offload tensors to SSDs for LLM training to overcome GPU memory limitations. Nevertheless, SSD band- width is limited, and the gap between SSD bandwidth growth and GPU com- putational throughput growth continues to widen, as shown in Figures 1.1 and 1.3. It is essential to carefully manage data transfers to prevent training throughput from being constrained by SSD bandwidth. We propose tech- niques for selecting which tensors to offload and hiding transfer latency. To evaluate the trade-off between performance and memory savings, we com- pare different strategies, i.e., tensor recomputation, offloading, and keeping tensors in GPU memory. These are detailed in Chapter 5. Despite the importance of data efficiency. several obstacles exist to address
for GNNs, is its assumption that both the input and output of each operator reside on the same device to which the operator is dispatched. However, the optimized transfer scheme perform PCle transfer simultaneously as an integral process, demanding the input to be in the host memory. To adopt such a data-efficient scheme and retain the PyTorch programming interface, the PyTorch runtime code has to be recompiled with the addition of a full-fledged new tensor type, its specia uses the GPU to gather node features and host memory allocator, and its set of new dispatch rules. Chapter 4 details our solution for incorporating ] OQ mos nan Al a LP Weasel Del Panvnh = EMOGI PCle transfer into PyTorch. ae enemy “bitin bk aan an sawndkh. [hewvtaa
for GNNs, is its assumption that both the input and output of each operator reside on the same device to which the operator is dispatched. However, the optimized transfer scheme perform PCle transfer simultaneously as an integral process, demanding the input to be in the host memory. To adopt such a data-efficient scheme and retain the PyTorch programming interface, the PyTorch runtime code has to be recompiled with the addition of a full-fledged new tensor type, its specia uses the GPU to gather node features and host memory allocator, and its set of new dispatch rules. Chapter 4 details our solution for incorporating ] OQ mos nan Al a LP Weasel Del Panvnh = EMOGI PCle transfer into PyTorch. ae enemy “bitin bk aan an sawndkh. [hewvtaa
Figure 2.1: Hierarchical breakdown of the GPT model. In training, dropout is applied to the output of each layer with red borders. ing, and their representation power is limited. In comparison, GNNs offer a
Figure 2.1: Hierarchical breakdown of the GPT model. In training, dropout is applied to the output of each layer with red borders. ing, and their representation power is limited. In comparison, GNNs offer a
Figure 2.2: Structure of a streaming multiprocessor (SM) in an Nvidia Volta V100 GPU [52-56]. The execution units include FP64 units, FP32 units, arithmetic logic units (ALUs), tensor cores (TCs), transcendental and data type conversion units (XUs), and load-store units (LSUs). compile time or runtime via the GPU driver.
Figure 2.2: Structure of a streaming multiprocessor (SM) in an Nvidia Volta V100 GPU [52-56]. The execution units include FP64 units, FP32 units, arithmetic logic units (ALUs), tensor cores (TCs), transcendental and data type conversion units (XUs), and load-store units (LSUs). compile time or runtime via the GPU driver.
Listing 2.1: Python code to sort tuples according to their second elements. Listing 2.2: C++ code to sort tuples by their second elements [57]
Listing 2.1: Python code to sort tuples according to their second elements. Listing 2.2: C++ code to sort tuples by their second elements [57]
VM PMs MO VUVELVUALELOLO 444 ULL VUE fit Po HHt UK Js 444110 IVIL. One of the biggest concerns of Python is the significant serialization penalty caused by the global interpreter lock (GIL) in multithreading programs. As the most widely-used Python implementation, CPython [58] uses GIL to ensure thread safety. To mitigate this, frameworks work around GIL. One method is to put performance-critical logic in the C++ framework libraries and release the GIL once the control flow goes outside the Python code [59]. Another direction is to remove t he GIL from the Python implementation. Although there are some alternate GIL-free Python implementations [60] to CPython, many frameworks re y on CPython-specific implementation de- tails, making it challenging to migrate these frameworks to such alternatives. These limitations led to the Pyt Non Enhancement Proposal (P EP) 703 to make GIL optional [61] in CPython, which has been accepted recently.
VM PMs MO VUVELVUALELOLO 444 ULL VUE fit Po HHt UK Js 444110 IVIL. One of the biggest concerns of Python is the significant serialization penalty caused by the global interpreter lock (GIL) in multithreading programs. As the most widely-used Python implementation, CPython [58] uses GIL to ensure thread safety. To mitigate this, frameworks work around GIL. One method is to put performance-critical logic in the C++ framework libraries and release the GIL once the control flow goes outside the Python code [59]. Another direction is to remove t he GIL from the Python implementation. Although there are some alternate GIL-free Python implementations [60] to CPython, many frameworks re y on CPython-specific implementation de- tails, making it challenging to migrate these frameworks to such alternatives. These limitations led to the Pyt Non Enhancement Proposal (P EP) 703 to make GIL optional [61] in CPython, which has been accepted recently.
Listing 2.3: PyTorch code to define model in Figure 2.3 and perform a training step. We denote the input hidden dimension as IN_DIM and th output hidden dimension as OUT_DIM.
Listing 2.3: PyTorch code to define model in Figure 2.3 and perform a training step. We denote the input hidden dimension as IN_DIM and th output hidden dimension as OUT_DIM.
and y_expected is the ground-truth labels. Notice that users only need to define the forward propagation logic, as shown in Figure 2.3. PyTorch’s auto-differentiation mechanism handles the computation of gradients without requiring users to manually specify the backward propagation logic. During forward propagation, PyTorch records activations, i.e., the intermediate tensors, and weights required for back- ward propagation and constructs the corresponding computational graph. n Figure 2.3, for example, y and y_expected are the activations stored for the backward propagation process of MSELoss, MSELossBackward. During backward propagation, PyTorch executes the backward-propagation compu- tational graph. PyTorch calls precompiled kernels to execute operators in orward propagation and backward propagation.
and y_expected is the ground-truth labels. Notice that users only need to define the forward propagation logic, as shown in Figure 2.3. PyTorch’s auto-differentiation mechanism handles the computation of gradients without requiring users to manually specify the backward propagation logic. During forward propagation, PyTorch records activations, i.e., the intermediate tensors, and weights required for back- ward propagation and constructs the corresponding computational graph. n Figure 2.3, for example, y and y_expected are the activations stored for the backward propagation process of MSELoss, MSELossBackward. During backward propagation, PyTorch executes the backward-propagation compu- tational graph. PyTorch calls precompiled kernels to execute operators in orward propagation and backward propagation.
Figure 3.1: The forward propagation of an RGCN layer could be divided into (1) message generation on edges and () node aggregation. We focus on paper node z in a large citation graph as an example. z only has two incoming edges, from a and 5, respectively. nC”) and h(&4) are node features. Wy rites is the weight for the type “writes”. Wo is the weight for virtual self-loops. o is the activation function. Notably, some runtime implementations may replicate data, e.g., Writes: a scalar that emphasizes the message associated with the same edge during
Figure 3.1: The forward propagation of an RGCN layer could be divided into (1) message generation on edges and () node aggregation. We focus on paper node z in a large citation graph as an example. z only has two incoming edges, from a and 5, respectively. nC”) and h(&4) are node features. Wy rites is the weight for the type “writes”. Wo is the weight for virtual self-loops. o is the activation function. Notably, some runtime implementations may replicate data, e.g., Writes: a scalar that emphasizes the message associated with the same edge during
Figure 3.3: Breakdown of inference time by Graphiler and Hector. Matrix multiply (MM) includes SpMM. We categorize PyTorch time not accounted for by kernels as “PyTorch Other Compute”.
Figure 3.3: Breakdown of inference time by Graphiler and Hector. Matrix multiply (MM) includes SpMM. We categorize PyTorch time not accounted for by kernels as “PyTorch Other Compute”.
Table 3.1: Features of Hector and prior [81-83] GNN compilers.
Table 3.1: Features of Hector and prior [81-83] GNN compilers.
Figure 3.5: Hector workflow and software architecture.
Figure 3.5: Hector workflow and software architecture.
Hector provides a decorator, @hector.compile, to take the existing PyG or DGL forward propagation logic and generate code for it, as exemplified by the input in Figure 3.5. The decorator, when applied to a method, invokes a simple transpiling pass that replaces the PyG and DGL method calls, e.g., SpMM/SDDMM, with an implementation in the inter-operator level IR, and replaces supported constructs from PyG and DGL with expressions in Hector IR. Similarly to statically-typed compilers in other Python packages [93,94], the function to compile can use most of the Python features except dynamic ones, e.g., assigning objects of different types to the same variable. We support a few types as the function arguments for heterogeneous graphs, involving Tensor and dict[str, Tensor] objects, i.e., dict objects where the keys are str objects and the values are Tensor objects. the keys are str objects and the values are Tensor objects.
Hector provides a decorator, @hector.compile, to take the existing PyG or DGL forward propagation logic and generate code for it, as exemplified by the input in Figure 3.5. The decorator, when applied to a method, invokes a simple transpiling pass that replaces the PyG and DGL method calls, e.g., SpMM/SDDMM, with an implementation in the inter-operator level IR, and replaces supported constructs from PyG and DGL with expressions in Hector IR. Similarly to statically-typed compilers in other Python packages [93,94], the function to compile can use most of the Python features except dynamic ones, e.g., assigning objects of different types to the same variable. We support a few types as the function arguments for heterogeneous graphs, involving Tensor and dict[str, Tensor] objects, i.e., dict objects where the keys are str objects and the values are Tensor objects. the keys are str objects and the values are Tensor objects.
Figure 3.6: The citation graph used as the example in Figures 3.7 and 3.8.
Figure 3.6: The citation graph used as the example in Figures 3.7 and 3.8.
n Table 3.3: Heterogeneous graph datasets [85-90] used in our evaluation The numbers reflect the default preprocessing by the OGB and DGL packages, e.g., adding inverse edges.
n Table 3.3: Heterogeneous graph datasets [85-90] used in our evaluation The numbers reflect the default preprocessing by the OGB and DGL packages, e.g., adding inverse edges.
Figure 3.9: Comparing the execution (Exec.) time of Hector best optimized code with previous work. Table 3.3 shows the datasets used
Figure 3.9: Comparing the execution (Exec.) time of Hector best optimized code with previous work. Table 3.3 shows the datasets used
optimized Hector code in Table 3.4.
optimized Hector code in Table 3.4.
*Normalized by the performance with compact materialization (C) because the unoptimized version triggers OOM errors.
*Normalized by the performance with compact materialization (C) because the unoptimized version triggers OOM errors.
Figure 3.10: Memory usage when Hector runs training and inference on H Ss Vi S GT. (b) shows the inference memory use (Infer. mem.) and training memory use (Train. mem.) of the unoptimized Hector code in MBs. (a) hows the portion of the memory use after applying compact materialization s. the unoptimized Hector code. For comparison, the number of nodes (# nodes), number of edges (# edges), and average degree of datasets are hown as dot scatters. The entity compaction ratio of each dataset is also Ss hown. Legend entries of each data series are placed next to the series’ axis.
Figure 3.10: Memory usage when Hector runs training and inference on H Ss Vi S GT. (b) shows the inference memory use (Infer. mem.) and training memory use (Train. mem.) of the unoptimized Hector code in MBs. (a) hows the portion of the memory use after applying compact materialization s. the unoptimized Hector code. For comparison, the number of nodes (# nodes), number of edges (# edges), and average degree of datasets are hown as dot scatters. The entity compaction ratio of each dataset is also Ss hown. Legend entries of each data series are placed next to the series’ axis.
Figure 3.11: Breakdown of Hector RGAT inference on two datasets. Input and output dimensions are 64. Cases with compaction (C), linear operator reordering (R), and no optimization (U) are presented. provement due to the less significant GEMM reduction.
Figure 3.11: Breakdown of Hector RGAT inference on two datasets. Input and output dimensions are 64. Cases with compaction (C), linear operator reordering (R), and no optimization (U) are presented. provement due to the less significant GEMM reduction.
Figure 3.12: Hector unoptimized performance. Each cell corresponds to one pair of dataset and model, where it is shown the time of (input dimension, output dimension) as (32, 32), (64, 64), and (128, 128) from the top to the bottom. Vacancy indicates OOM errors. Hector achieves better memory efficiency than state-of-the-art systems. the massive resources on GPUs. By eliminating unnecessary data copies,
Figure 3.12: Hector unoptimized performance. Each cell corresponds to one pair of dataset and model, where it is shown the time of (input dimension, output dimension) as (32, 32), (64, 64), and (128, 128) from the top to the bottom. Vacancy indicates OOM errors. Hector achieves better memory efficiency than state-of-the-art systems. the massive resources on GPUs. By eliminating unnecessary data copies,
Figure 4.1: An example demonstrating the GraphSAGE neighbor sampling approach for output node 9. There are two aggregation layers in the model. The left shows the data layout of node features of this graph.
Figure 4.1: An example demonstrating the GraphSAGE neighbor sampling approach for output node 9. There are two aggregation layers in the model. The left shows the data layout of node features of this graph.
system as long as it is pinned and memory-mapped into the GPU’s address space. In zero-copy access, GPU sends the request through the interconnect to get data, without explicit copying or migration in the previously mentioned two mechanisms. When accessing host memory, the GPU issues at most cacheline-sized, i.e., 128 bytes, data requests onto PCIe [17]. There are three APIs or combinations that enable zero-copy in a memory region [25], but for simplicity, we choose cudaMallocHost() in PyTorch-Direct.
system as long as it is pinned and memory-mapped into the GPU’s address space. In zero-copy access, GPU sends the request through the interconnect to get data, without explicit copying or migration in the previously mentioned two mechanisms. When accessing host memory, the GPU issues at most cacheline-sized, i.e., 128 bytes, data requests onto PCIe [17]. There are three APIs or combinations that enable zero-copy in a memory region [25], but for simplicity, we choose cudaMallocHost() in PyTorch-Direct.
Figure 4.3: CPU utilization and data loader time comparison between C} and GNN training. CPU utilization can go beyond 100% as it is multithreaded. TensorF low.
Figure 4.3: CPU utilization and data loader time comparison between C} and GNN training. CPU utilization can go beyond 100% as it is multithreaded. TensorF low.
All unified tensors are GPU-affinitive. At least one unified tensor is host-affinitive.
All unified tensors are GPU-affinitive. At least one unified tensor is host-affinitive.
Table 4.2: Rules of placements for operators involving unified tensors as operands.
Table 4.2: Rules of placements for operators involving unified tensors as operands.
This section evaluates PyTorch-Direct performance using a well-defined mi- crobenchmark and end-to-end GNN training. Using the microbenchmark, we demonstrate that (1) PyTorch-Direct is faster than the baseline PyTorch approach in accessing features from the GPU under different combinations of data sizes and systems and (2) the effectiveness of our optimized mem- ory alignment mechanism. In GNN training, we show the benefit of using PyTorch-Direct for faster training. PyTorch-Direct for faster training.
This section evaluates PyTorch-Direct performance using a well-defined mi- crobenchmark and end-to-end GNN training. Using the microbenchmark, we demonstrate that (1) PyTorch-Direct is faster than the baseline PyTorch approach in accessing features from the GPU under different combinations of data sizes and systems and (2) the effectiveness of our optimized mem- ory alignment mechanism. In GNN training, we show the benefit of using PyTorch-Direct for faster training. PyTorch-Direct for faster training.
Table 4.3: Datasets for GNN training, their characteristics, and abbreviations (abbv.) used in the text.
Table 4.3: Datasets for GNN training, their characteristics, and abbreviations (abbv.) used in the text.
Figure 4.5: Memory alignment optimization with a circular shift. The example is identical to the case in Figure 4.4. Alignment reduces the tota number of PCIe requests (req) from seven to five in this case.
Figure 4.5: Memory alignment optimization with a circular shift. The example is identical to the case in Figure 4.4. Alignment reduces the tota number of PCIe requests (req) from seven to five in this case.
Figure 4.6: Irregular host data access pattern microbenchmark comparisons between PyTorch (Py) and PyTorch-Direct (PyD) on different systems. The ideal case shows only the pure data transfer time with a peak PCle bandwidth.
Figure 4.6: Irregular host data access pattern microbenchmark comparisons between PyTorch (Py) and PyTorch-Direct (PyD) on different systems. The ideal case shows only the pure data transfer time with a peak PCle bandwidth.
Figure 4.7: Memory access alignment and its impact on PyTorch-Direct (PyD) performance. PyTorch (Py) results were added for comparison.
Figure 4.7: Memory access alignment and its impact on PyTorch-Direct (PyD) performance. PyTorch (Py) results were added for comparison.
Figure 4.8: Single epoch execution time breakdown for both PyTorch (Py) vs. PyTorch-Direct (PyD) when running (a) GraphSAGE and (b) GAT in different datasets. Training epoch time reductions are written on the bars.
Figure 4.8: Single epoch execution time breakdown for both PyTorch (Py) vs. PyTorch-Direct (PyD) when running (a) GraphSAGE and (b) GAT in different datasets. Training epoch time reductions are written on the bars.
Figure 5.1: The growth of FP16 throughput (right vertical axis) of GPUs for deep learning training is aligned with the model size of LLMs (left vertical axis), but GPU memory capacity (left vertical axis) falls behind [153]. The horizontal axis shows the release date. Points represent both Nvidia 100-level GPUs since K100 and Google TPUs. The growth rate of FP16 throughput (yellow dotted line) is more than 2x of that of the memory capacity growth rate (red dotted line).
Figure 5.1: The growth of FP16 throughput (right vertical axis) of GPUs for deep learning training is aligned with the model size of LLMs (left vertical axis), but GPU memory capacity (left vertical axis) falls behind [153]. The horizontal axis shows the release date. Points represent both Nvidia 100-level GPUs since K100 and Google TPUs. The growth rate of FP16 throughput (yellow dotted line) is more than 2x of that of the memory capacity growth rate (red dotted line).
Figure 5.4: The trend of enterprise SSD capacity [19]. For each model, only the data of the variant with maximal capacity is collected. Red lines show the growth rates predicted by quantile regression. The visualization code is adapted from Derek Jones’s work [20]. cated, causing write amplification. In turn, vendors adopt over-provisioning
Figure 5.4: The trend of enterprise SSD capacity [19]. For each model, only the data of the variant with maximal capacity is collected. Red lines show the growth rates predicted by quantile regression. The visualization code is adapted from Derek Jones’s work [20]. cated, causing write amplification. In turn, vendors adopt over-provisioning
To take a closer look at the uniqueness of SSDTrain, let us compare it with related work Stronghold [183] and ZeRO-Infinity [184]. For difference (a), SS- DTrain offloads activations. Data in LLM training can be categorized into mutually exclusive types: parameters, optimizer states, gradients, and ac- tivations. In contrast, existing work offloads other data than activations. E.g., Stronghold offloads parameters and gradients. Although ZeRO-Infinity offloads many types of data, when it comes to activations, only a subset as defined in the activation checkpoints, is optionally offloaded. Activations are the intermediate tensors produced in the forward propagation and kept for gradient computation. They are consumed in the backward propagation im- gradient computation. They are consumed in the backward propagation im-
To take a closer look at the uniqueness of SSDTrain, let us compare it with related work Stronghold [183] and ZeRO-Infinity [184]. For difference (a), SS- DTrain offloads activations. Data in LLM training can be categorized into mutually exclusive types: parameters, optimizer states, gradients, and ac- tivations. In contrast, existing work offloads other data than activations. E.g., Stronghold offloads parameters and gradients. Although ZeRO-Infinity offloads many types of data, when it comes to activations, only a subset as defined in the activation checkpoints, is optionally offloaded. Activations are the intermediate tensors produced in the forward propagation and kept for gradient computation. They are consumed in the backward propagation im- gradient computation. They are consumed in the backward propagation im-
Figure 5.5: SSDTrain timeline of a step of a two-micro-batch three-layer L) model. PyTorch hooks are used to trigger tensor cache bookkeeping, tensor offloading (@), and tensor loading (()). In the forward (F) propagation, SSDTrain records the order of scopes (@)) and switches between micro-batches at the end of the stages (@)). SSDTrain starts oading when it is switched to the backward (B) propagation (@).
Figure 5.5: SSDTrain timeline of a step of a two-micro-batch three-layer L) model. PyTorch hooks are used to trigger tensor cache bookkeeping, tensor offloading (@), and tensor loading (()). In the forward (F) propagation, SSDTrain records the order of scopes (@)) and switches between micro-batches at the end of the stages (@)). SSDTrain starts oading when it is switched to the backward (B) propagation (@).
Figure 5.6: SSDTrain software architecture. The components of SSDTrain are shown as blue blocks with white text. The CUDA malloc hook is a C++ library, while others are Python code. Figure 5.6: SSDTrain software architecture. The components of SSD Train whether they are being offloaded, the path in the file system, etc.
Figure 5.6: SSDTrain software architecture. The components of SSDTrain are shown as blue blocks with white text. The CUDA malloc hook is a C++ library, while others are Python code. Figure 5.6: SSDTrain software architecture. The components of SSD Train whether they are being offloaded, the path in the file system, etc.
Figure 5.7: ‘Tensor cache registers pack—unpack hook pair to offload tensors and reload tensors. (a) shows the PyTorch computational graph. (b) shows the hardware data path. (c) and (d) show the tensor cache state when the pack or unpack hook is triggered. During an operator ()), PyTorch calls the pack hook with tensors to be saved for backward propagation and registers the return values on the computational graph (()). Tensor cache tracks the tensors, offloads them (@)), and returns identifiers for the tensors. In an operator (@) in the backward propagation, PyTorch calls the unpack hook with the identifiers to get tensors (()). The tensor cache blocks until the requested tensors are loaded in GPU memory (©).
Figure 5.7: ‘Tensor cache registers pack—unpack hook pair to offload tensors and reload tensors. (a) shows the PyTorch computational graph. (b) shows the hardware data path. (c) and (d) show the tensor cache state when the pack or unpack hook is triggered. During an operator ()), PyTorch calls the pack hook with tensors to be saved for backward propagation and registers the return values on the computational graph (()). Tensor cache tracks the tensors, offloads them (@)), and returns identifiers for the tensors. In an operator (@) in the backward propagation, PyTorch calls the unpack hook with the identifiers to get tensors (()). The tensor cache blocks until the requested tensors are loaded in GPU memory (©).
Figure 5.8: Memory footprint of one A100 in a BI ERT training step with offloading (black) and without (blue) on Table 5.3’s system. Run with offloading incurs more allocator events because of memory release and allocation caused by tensor offoading and reloading. SSDTrain reduces memory footprint at the beginning of backward propagation by 45% and end-to-end peak memory footprint by 25%.
Figure 5.8: Memory footprint of one A100 in a BI ERT training step with offloading (black) and without (blue) on Table 5.3’s system. Run with offloading incurs more allocator events because of memory release and allocation caused by tensor offoading and reloading. SSDTrain reduces memory footprint at the beginning of backward propagation by 45% and end-to-end peak memory footprint by 25%.
Figure 5.9: The adaptive offloading algorithm uses profiling to decide modules in which the activations are to be kept in GPU memory. The model is represented as a tree where each scope is a node. On each node, the forward computation time and data transfer size are recorded during profiling. The I/O time in the forward propagation is also recorded and shown in parenthesis in the root node.
Figure 5.9: The adaptive offloading algorithm uses profiling to decide modules in which the activations are to be kept in GPU memory. The model is represented as a tree where each scope is a node. On each node, the forward computation time and data transfer size are recorded during profiling. The I/O time in the forward propagation is also recorded and shown in parenthesis in the root node.
Figure 5.10: bandwidth (le Estimate of SSD lifespan (left pink vertical axis), PCle write t blue vertical axis) and maximal activations size per GPU (right vertical axis). Lifespans longer than 5 years are shown on top of the pink bars. The horizontal axis shows the number of GPUs, the framework, and the model size [47]. ZeRO3 stands for DeepSpeed with stage-3 ZeRO, i.e., all optimizer states, gradients, and parameters are sharded across data parallel ranks.
Figure 5.10: bandwidth (le Estimate of SSD lifespan (left pink vertical axis), PCle write t blue vertical axis) and maximal activations size per GPU (right vertical axis). Lifespans longer than 5 years are shown on top of the pink bars. The horizontal axis shows the number of GPUs, the framework, and the model size [47]. ZeRO3 stands for DeepSpeed with stage-3 ZeRO, i.e., all optimizer states, gradients, and parameters are sharded across data parallel ranks.
(b) Hidden dimension as 12,288. In all cases, the required write bandwidth per GPU is smaller than the original case’s 13.8 GB/s.
(b) Hidden dimension as 12,288. In all cases, the required write bandwidth per GPU is smaller than the original case’s 13.8 GB/s.
Figure 5.13: Estimate of SSD lifespan in scenarios of Figure 5.10 with the assumption that each A100 is paired with four Samsung 980 PRO 1TB.
Figure 5.13: Estimate of SSD lifespan in scenarios of Figure 5.10 with the assumption that each A100 is paired with four Samsung 980 PRO 1TB.
One of the most important compiler research problems is the algorithm that makes choices among candidates in the design space. It remains an open high-performance computing Besides, the specific microarc problem how the data-dependent sparse operations and layout choices fit in the cost model and layout choices. Pertinently, in various applications in HPC) with multiple layout and kernel choices, researchers have developed heuristics to make the optimized choice [241-243]. hitecture of each GPU model also makes a dif- erence due to the architecture-specific features available, e.g., asynchronous oading to shared memory since Ampere [244], and different microarchitec- ture characteristics in each model. Therefore, it is meaningful to investigate their impact and incorporate them into decision-making. Ine of the most important compiler research problems is the algorithm that
One of the most important compiler research problems is the algorithm that makes choices among candidates in the design space. It remains an open high-performance computing Besides, the specific microarc problem how the data-dependent sparse operations and layout choices fit in the cost model and layout choices. Pertinently, in various applications in HPC) with multiple layout and kernel choices, researchers have developed heuristics to make the optimized choice [241-243]. hitecture of each GPU model also makes a dif- erence due to the architecture-specific features available, e.g., asynchronous oading to shared memory since Ampere [244], and different microarchitec- ture characteristics in each model. Therefore, it is meaningful to investigate their impact and incorporate them into decision-making. Ine of the most important compiler research problems is the algorithm that