#  ParallelGPUOS: A Concurrent OS-level GPU Checkpoint and Restore System
using Validated Speculation

Zhuobin Huang111Work done while Zhuobin Huang and Yinyi Hao were at Institute
of Parallel and Distributed Systems, SEIEE, Shanghai Jiao Tong University.
Institute of Parallel and Distributed Systems, SEIEE, Shanghai Jiao Tong
University  Xingda Wei Xingda Wei is the corresponding author
(wxdwfc@sjtu.edu.cn). Institute of Parallel and Distributed Systems, SEIEE,
Shanghai Jiao Tong University  Yingyi Hao  Institute of Parallel and
Distributed Systems, SEIEE, Shanghai Jiao Tong University  Rong Chen
Institute of Parallel and Distributed Systems, SEIEE, Shanghai Jiao Tong
University  Mingcong Han  Institute of Parallel and Distributed Systems,
SEIEE, Shanghai Jiao Tong University  Jinyu Gu  Institute of Parallel and
Distributed Systems, SEIEE, Shanghai Jiao Tong University  Haibo Chen
Institute of Parallel and Distributed Systems, SEIEE, Shanghai Jiao Tong
University

## Abstract

Checkpointing (C) and restoring (R) are key components for GPU tasks. POS is
an OS-level GPU C/R system: It can transparently checkpoint or restore
processes that use the GPU, without requiring any cooperation from the
application, a key feature required by modern systems like the cloud.
Moreover, POS is the _first_ OS-level C/R system that can concurrently execute
C/R with the application execution: a critical feature that can be trivially
achieved when the processes only running on the CPU, but becomes challenging
when the processes use GPU. The problem is how to ensure consistency during
concurrent execution with the lack of application semantics due to
transparency. CPU processes can leverage OS and hardware paging to fix
inconsistency without application semantics. Unfortunately, GPU bypasses OS
and paging for high performance. POS fills the semantic gap by speculatively
extracting buffer access information of GPU kernels during runtime. Thanks to
the simple and well-structured nature of GPU kernels, our speculative
extraction (with runtime validation) achieves 100% accuracy on applications
from training to inference whose domains span from vision, large language
models, and reinforcement learning. Based on the extracted semantics, we
systematically overlap C/R with application execution, and achieves orders of
magnitude higher performance under various tasks compared with the state-of-
the-art OS-level GPU C/R, including training fault tolerance, live GPU process
migration, and cold starts acceleration in GPU-based serverless computing.

##  Introduction

The advancement in Machine Learning (ML) in fields like computer vision [34],
large language models [54], and reinforcement learning [68], has led to
increasing deployment of processes that use GPU. Checkpoint and restore (C/R)
is a key building block in GPU applications, which provides fault tolerance
[67], live migration [67, 78], and auto-scaling [31, 12] (detailed in §2.1)
capabilities. Specifically, checkpoint stops the tasks and dumps their states
(e.g., CPU and GPU memory) into a condensed binary (_image_). Restore
recreates the task from the image.

C/R can be implemented at the operating system level (OS-level), i.e., the OS
directly C/Rs the process running the task [10, 2, 67, 4] or at the task-
level, i.e., leveraging the ML frameworks to perform C/R [61, 43]. A key
advantage of OS-level C/R is _transparency_ , which means it doesn’t require
any _explicit cooperation from the applications_ , such as using a specific
C/R framework or being aware of the C/R. Transparency is particularly
important when the system that performs the C/R is not the task, e.g., the
cloud infrastructure. Consider a cloud scenario where the vendor wants to
migrate tasks between machines for better resource utilization [67, 78]. Since
the tasks can be arbitrary user applications, it is impossible to do so at
task-level, unless the cloud poses an impractical restriction on the tasks
that can run, i.e., the users can only submit applications using a specific
version of the framework. Meanwhile, transparency can improve efficiency,
e.g., the task-level C/R is not implemented efficiently, and improve user
experience, e.g., the user does not need to explicitly use C/R for fault
tolerance, a selling point of GPU cloud services [67].

Reducing the downtime caused by C/R is critical to applications. One canonical
way is by conducting the C/R concurrently with the application’s execution.
However, ensuring consistency under concurrent C/R is challenging in OS-level
C/R due to the _lack of application semantic_ at the OS-level. For example,
the OS may checkpoint a memory page while the application is concurrently
modifying it. Similarly, the OS may restore a memory page while the
application concurrently reads it. Without application semantics, the OS
cannot coordinate the application execution with C/R correctly.

For processes that only run on the CPU, the OS can cooperate with the hardware
paging without application semantics. For example, we can use copy-on-write
(CoW) to isolate concurrent reads and writes to the same page [65, 77], or
leverage dirty-bit to track the memory page (dirty page) modified by the
application during checkpoint [9, 22, 65], or utilize on-demand paging to
ensure the application reads a complete page during concurrent restore [1, 27,
26, 76].

Unfortunately, being designed for high performance, GPU programs bypass the OS
during (most) runtime, and the GPU hardware has limited paging capabilities,
e.g., no dirty bit and copy-on-write support. Therefore, existing systems
cannot concurrently execute GPU C/R with applications [2, 15, 58, 67, 18, 4].
As the CPU part of the C/R can be concurrently executed, the GPU part of the
C/R can dominate the downtime of the applications (§2.3). A foundational
question answered by this paper is thus _how to correctly execute concurrent
GPU C/R with the application execution?_

Key insight and methodology.  Our key insight is that the OS can extract
sufficient semantics—fine-grained GPU buffer access patterns of GPU programs
at runtime with high accuracy. This bridges the semantic gap between the OS
and the applications for efficient C/R. First, unlike CPU, GPU computation are
abstracted as fine-grained computation units called _kernels_ [23] and memory
move operations222Example: cudaMemcpy. For brevity, we also term GPU memory
move operations as _kernel_ s in this paper without losing generality., each
operates on a few GPU buffers (e.g., 2 buffers over a total of 1,000 buffers).
More importantly, we can quickly and accurately extract the fine-grained GPU
buffers accessed by each kernel during runtime, because they have a relatively
simple execution pattern, e.g., conduct a matrix multiplication on three GPU
buffers.

With the extracted semantics, we can imitate features required for efficient
C/R unsupported by GPU hardware in software. For example, before executing a
kernel, we can check whether the kernel may modify a buffer that is being
checkpointed. This essentially realizes a “soft dirty bit” for GPU buffers. We
can also realize features like “soft copy-on-write (CoW)” and “soft on-demand
load” similarly to avoid stalling the GPU application during C/R.

ParallelGPUOS (POS). We present POS, an efficient OS-level GPU C/R system.
Unlike existing designs, POS can execute GPU kernels concurrently with the
C/R, while the C/R is coordinated with “Soft CoW”, “Soft dirty bit” and “Soft
on-demand load”. To support the above features, we propose _kernel DAG_
(directed acyclic graph)—a dataflow constructed during runtime that tracks the
fine-grained application semantics of the GPU execution. Specifically, the
nodes are kernels (executing or waiting to be executed) and buffers, and edges
are dependencies between them. The DAG is inspired by the following
requirements for coordinating applications with C/R. First, it has to
incorporate which buffer a kernel writes for “soft CoW” and “soft dirty bit”.
Second, it must specify which buffer a kernel reads for “soft on-demand load”.
Finally, it has to record dependencies between un-executed kernels such that
we can prioritize buffer loading during the on-demand load. The DAG is
transparently constructed and maintained by POS with GPU virtualization [46,
13, 74].

A key challenge is how to infer the buffer accessed by a kernel with high
accuracy. The observation is that the buffers accessed by a GPU kernel can be
inferred from the arguments of the API that launches the kernel. They are
either specified by the driver documentation, e.g., for kernels provided by
the driver [49], or can be speculatively inferred with empirically 100%
accuracy. Interestingly, we found the first type of kernels dominate the
workloads (up to 95%) due to their efficiency on the vendor’s GPU.
Nevertheless, we have to correctly handle the second type of kernels. To this
end, we use binary instrumentation [73, 41] to validate the speculation at
runtime (§4.2). We fall back to a slow path for correctness under validation
fails. The instrumentation is done offline, and we only call the instrumented
kernel during C/R to prevent hurting normal execution performance. We don’t
experience any validation failures in emerging GPU applications.

Besides supporting ‘Soft copy-on-write”, “Soft dirty bit” and “Soft on-demand
load” with kernel DAG, our kernel DAG further decouples GPU execution from the
underlying low-level GPU driver states, such that we can avoid stalls caused
by recreating the GPU execution context during restoration—another source of
stalls during GPU restoring (up to 66.1% of the restore time in our evaluated
workloads). Besides, it also enables optimizations like CPU–GPU memory
deduplication for reduced checkpointed image size and time for several
applications (§5.2).

We have implemented POS on Linux and it can run unmodified GPU applications
that use NVIDIA GPUs. On an A100 GPU, when compared to existing stop-the-world
C/R, POS brings 1.3–3.6 $\times$ throughput improvements when using periodical
checkpointing for fault tolerance, can shorten the downtime of live migrating
a llama 2 (13 B) from 7.9 s to 0.4 s, and can reduce the cold start time for
various GPU processes by up to 82%.

Contributions. Our contributions are highlighted as follows:

  * •

We propose Kernel DAG to bridge the semantics between OS-level C/R and GPU
execution.

  * •

We present the first OS-level GPU C/R system that can concurrently run C/R
with GPU execution in various scenarios.

  * •

Our prototype system with extensive evaluations demonstrated the efficacy of
our approach.

We will open source POS upon publication.

##  Background and Motivation

Figure 1:  An illustration of the overhead of C/R for (a) C-mostly:
periodically checkpointing for fault tolerance in training tasks, (b) C+R: the
timeline of migration of an inference task and (c) R-mostly: a breakdown of
restore inference tasks in a serverless manner. Note that the GPU memory used
is much larger than the model parameter size marked in the graph.

###  The case of OS-level C/R for GPU applications

Use cases.  C/R is commonly used in ML tasks that use GPU, which all require
high C/R performance:

  * •

Fault tolerance (C-mostly)333Restore is also necessary in fault tolerance for
failure recoveries. We omit it as its execution frequencies are relatively low
compared to C.. ML tasks, such as training, are vulnerable to GPU failures
[20, 16, 67]. C/R provides fault tolerance by periodically checkpointing and
persisting the images [70, 77, 67]. Upon failure and recovery, the OS simply
restores the task using the checkpointed image, creating an illusion that the
tasks never failed. The checkpointed image could be stale, i.e., not
checkpointing the latest state on the GPU. But the frequency should be high to
minimize computation loss from too stale checkpoints [75].

  * •

Live migration (C+R). Process migration is widely used by cloud vendors to
improve resource utilization or providing elasticity [67, 78, 30, 82]: the OS
first stops and checkpoints the running process, transfers the image to
another machine, and then resumes the process. Since the process is stalled
during migration, it is critical to minimize the checkpoint, transfer, and
restore time to achieve live migration. Unlike the C-mostly case, migration
requires the latest state to be checkpointed.

  * •

Fast process startup (R-mostly). In emerging cloud computing paradigms like
serverless computing, the GPU process can be created on demand to serve
requests [32]. OS-level restoring is a key technique in accelerating process
startup [76, 12, 72, 31] because it skips the language runtime and framework
initialization overhead commonly found costly in ML frameworks like Pytorch
[11, 76].

Why OS-level C/R.  Though task framework can also implement the C/R, e.g.,
GEMINI [75] and Nebula [43] uses task-level C/R for fault tolerance;
TorchElastic [61] for elasticity, the key benefit of OS-level C/R is that it
_does not require any application cooperation_. (1) There exists cases when
C/R is conducted outside of the tasks. Consider a cloud that runs applications
from different users. It is desirable for the cloud provider to migrate
applications to improve cluster utilizations [67, 78, 30]. From the cloud
provider’s perspective, the applications are black boxes, which means
providers can only use OS-level C/R unless imposing impractical restrictions
on the applications. For example, using an application-level solution like
TorchElastic requires all users to use TorchElastic, which is clearly
impossible. Meanwhile, TorchElastic has additional compatibility issues even
for applications written in Pytorch [60]. (2) User-level C/R may be
implemented inefficiently, e.g., our experiments in §6.1 show that a carefully
designed OS-level C/R (without developing efforts) is 12.4–62.3% faster than a
user-level C/R (with extra coding and configuration). (3) The OS-level C/R can
provide better user experiences. For instance, developers do not need to
handle failure recovery manually since the OS can manage the restore
transparently, which is a selling point for vendors that provide GPU services
[67].

Figure 2: The CPU and GPU objects to be checkpointed in (a) traditional GPU
checkpoint [15, 58, 67, 18, 4] and (b) POS.

###  Basic OS-level C/R for GPU processes

GPU and kernel basics.  GPU is an accelerator with high computational capacity
commonly used in ML tasks. Although various GPU vendors exist, they all follow
the same _command–launch_ execution model: to execute a job (a.k.a _kernel_)
on the GPU, the CPU executes the driver API that sends a command to the GPU to
launch the kernel on it. The command contains the kernel function name and its
arguments. The arguments typically indicate the GPU memory (abstracted as
fine-grained buffers) that the kernel accesses. GPU executes kernels
asynchronously, i.e., the return of the API to the caller thread on the CPU
does not necessarily indicate that the kernel has finished on the GPU. The CPU
can use APIs such as cudaDeviceSynchronize to wait for the in-flight kernels
to complete.

Before kernel execution, the GPU driver will first initialize a context, which
includes buffers for arguments, command queues (streams), and various
dynamically loaded modules on the GPU. Kernels sent to the same stream execute
in a FIFO order, while kernels in different streams can execute concurrently.

The kernel binaries can be intercepted by the OS before launching them onto
the GPU (e.g., by intercepting the loaded CUDA modules), allowing us to
instrument these binaries [73] or conduct static analysis on them [41] before
execution. Since most kernel binaries (except the JIT-ed ones) are packaged
within the process’s binaries, we can perform the instrumentation and analysis
offline to avoid runtime costs.

C/R a GPU process [2, 15, 58, 67, 18, 4]. For the stop-the-world C/R, GPU does
not make any high-level differences compared to C/R a pure-CPU process. For
checkpointing, the OS first halts the execution on both the CPU and GPU. To
halt GPU execution, the OS stops submitting kernels to the GPU, stores them
locally, and utilizes cudaDeviceSynchronize to drain unfinished kernels off
the GPU. Afterward, it dumps the CPU and GPU memory, along with necessary data
structures such as the GPU context, streams, and buffered unexecuted kernels
into the checkpointed image (see Figure 2 (a)). For restoring, the OS creates
the required GPU context and streams, copies the buffers to the GPU, and
finally re-submits unexecuted kernels to resume execution.

###  Analysis of costs of existing GPU C/R

GPU C/R has non-trivial overheads to various tasks.  Figure 1 presents the
analysis of existing stop-the-world C/R on three representative use cases of
C/R on common ML tasks444Detailed setup is in §6. We compare the training
iteration time without checkpointing to two baselines: checkpoint with driver
utilities, i.e., NVIDIA/cuda-checkpoint [4, 52] (Driver-C), and the approach
described in Singularity [67], the state-of-the-art stop-of-the-world C/R
system [67] (Base-C).

For fault tolerance (a), adding checkpoints at an interval close to twice the
training iteration time in most training tasks with Base-C, incurring
1.62–4.55 $\times$ overhead for various training tasks with Base-C. This is
because the checkpointing will stop the execution, and the checkpoint time is
close to the training iteration time. Checkpointing with driver-C incurs
larger overhead as it further clears and restores the entire GPU state
including context and memory during C/R. Since driver-C is close-sourced it
performs worse than Base-C, we focus on Base-C in this paper.

For migration (b), C/R using Base-C adds an about 8-second service downtime to
llama 2 inference, causing hundreds of tokens to be stalled. The disruption is
significant as stalled time caused by C/R is 2.1$\times$ longer than the
average request handling time [69]. Finally, when using restore to speed up
serverless function startups (c), the restore of GPU tasks is 1.25–3.9$\times$
longer than the task execution time.

GPU memory copy dominates the checkpointing.  Figure 1 (b) breaks down the C/R
time into CPU and GPU parts. The memory copy dominates because (1) the size of
the copied memory is non-trivial (50 GB) and (2) other costs like CPU-side can
be hidden with existing techniques, e.g., using pre-copy [9] to concurrently
execute the CPU memory copy with the application. Unfortunately, we cannot do
similarly on the GPU.

The first challenge is how to _overlap GPU memory copy with GPU execution_ for
checkpointing?

GPU memory copy contributes significantly to the restoring.  Similarly, as
shown in Figure 1 (b) and (c), memory copy contributes a 50% restore time
especially for relatively large models (e.g., llama 2 13 B, 50 GB). CPU-side
has less impact because its execution only requires a small amount of memory
footprint, which can be effectively hidden via om demand paging [1, 27, 26,
76].

The second challenge is how to _overlap GPU memory copy with GPU execution_
for restoring?

GPU context creation is slow during restoring.  Finally, context creation
takes 50–83% of the restoring time, and the time is dominated by the GPU part.
Creating an execution context on a GPU is more costly than on a CPU due to
extensive GPU hardware configuration and driver state loading. In comparison,
creating a CPU execution context (a Linux process) only requires initializing
a few Linux kernel data structures.

The third challenge is how to _avoid costly context creation for restoring GPU
execution_?

##  Overview of POS

Figure 3: The architecture of POS.

Figure 4: An overview of the POS C/R procedure under (a) fault tolerance, (b)
task migration, and (c) fast process startup. For brevity, we omit C/R CPU
memory.

POS is a system service for efficient C/R for GPU applications. The system
administrator can use command line instructions to checkpoint, restore or
migrate a GPU process. Unlike existing GPU C/R systems, POS can effectively
overlap GPU execution with both checkpointing and restoring, thereby
minimizing or eliminating stalls that interrupt the execution of GPU
processes.

Figure 3 shows the internal architecture of POS. POS is an OS service that
maintains a kernel DAG (§4)—our key data structure to fill the semantic gap
between C/R and application execution. If the application submits a GPU kernel
(by calling a driver API) (❶), we will intercept it (❷), update the DAG (❸),
and then forward the kernel to the driver (❹) for execution (❺).

Besides the kernel DAG, POS has an engine for rehearsing the C/R. For the
application’s CPU part, POS leverages CRIU [10]—the state-of-the-art CPU C/R
library to checkpoint and restore its state. For the GPU part, POS itself will
C/R the GPU buffer (buffer allocated by the application, recorded as the
buffer handles in POS) with the help of kernel DAG (§5). Finally, POS
transfers the checkpointed image to a wide range of supported targets,
including host CPU memory, remote CPU memory and SSD.

To achieve efficient C/R, POS needs to answer two questions: (1) how to
concurrently execute memory copies in C/R with minimal interference in normal
GPU execution (challenges 1 and 2 summarized in §2.3), and (2) how to avoid
costly GPU context creation during restoration (challenge 3). We use two
system designs accordingly.

Kernel DAG. We observe that a kernel DAG is sufficient to track the runtime
state on the GPU, therefore giving the OS sufficient semantics to coordinate
C/R with concurrent GPU execution. Specifically, in a kernel DAG, each node
represents a GPU kernel or a memory buffer manipulated by the kernel. The edge
between two nodes represents the dependencies and data flow between them. As
shown in Figure 2 (b), the flow $1\rightarrow K0\rightarrow(1,2)$ indicates
that a GPU computation (or memory copy) kernel ($k0$) will read the buffer 1
and write buffers 1 and 2. The buffer can be either the CPU memory or the GPU
buffer. Including the CPU buffer in the DAG is critical for deduplication
during a checkpoint (§5.2).

POS maintains the kernel DAG through GPU virtualization. When the
application’s GPU process calls a driver API, we intercept and analyze it, and
update the DAG according to predefined rules (see §4.1). The update either
appends a node to the DAG, e.g., upon calling cudaLaunchKernel that launches a
kernel on the GPU, or removes nodes if they are done.

In addition to facilitating C/R coordination, another key advantage of the
kernel DAG is that _it decouples GPU execution from the GPU execution
context._ Specifically, a DAG can be viewed as a high-level abstraction of the
application semantic that transforms memory buffers from one state to another.
Thus, a DAG can be executed as long as there’s an available GPU context for
launching kernels. Based on this observation, POS further leverages the
context pool to avoid expensive GPU context creation during restoration.

Execution context pool.  Specifically, as long as the kernels in the DAG are
correctly recorded according to the dependencies, we can re-execute them
regardless of the GPU context. Therefore, POS pre-initializes a pool of GPU
contexts for executing restoration kernels. When restoring a context, we
simply return a handler from the pool. The context is returned to the pool
when the application exits. As a result, restoring doesn’t need to create new
contexts. Maintaining a context pool is possible because POS is a long-running
system service.

Overview of C/R execution flow. Figure 4 illustrates the overall C/R procedure
in various C/R cases (§2.1). In the C-mostly scenarios like fault tolerance
(a), the goal is to take a snapshot of the application state periodically. We
will first stop the CPU and GPU execution, marking the CPU execution as copy-
on-write (CoW), and marking the GPU as a soft CoW implemented based on the
kernel DAG (❶). This stop is short (within 1 ms). Afterward, POS will resume
application execution, and start copying the CPU and GPU memory from GPU
buffers to the image (❷). If the CPU issues a GPU kernel during the copy, we
will copy its accessed buffers (CoW) if this kernel would affect the buffers
being checkpointed (❸). Once all the buffers have been copied, the checkpoint
is done.

In the cases of C+R like live migration (b), the goal is to checkpoint the
latest process state to a remote machine (GPU). We will first stop the CPU and
GPU execution and enable the soft dirty bit mode (❶). We then start copying
the CPU and GPU memory to the target and resume application execution (❷).
During the copy, we track the dirty buffers modified by the execution. After
all the buffers have been copied, we will stop the execution again, retransmit
both the dirty CPU pages and GPU buffers (❸), and resume the execution on the
remote GPU. Note that if the number of dirty buffers exceeds a threshold, we
will use re-execution for reduced memory transfer by delaying the execution,
recording them in the kernel DAG and copying the kernel DAG to the remote
machine for restoring (❹).

Finally, for R-mostly cases like starting a process with restoring, e.g., in
serverless, the goal is to start the GPU process from the checkpointed image
as soon as possible. Thus, POS designs an on-demand GPU buffers restore
procedure (c). We will first restore the kernel DAG if it exists in the
checkpoint (❶) and load buffers that are required by it (❷). The order to load
the buffers is prioritized by the kernel execution order defined in the DAG.
We will launch a kernel as long as its input buffer is ready (❸). Note that
loading buffer and launching kernel can directly use context in the pool. For
new kernels issued by the application, we will also load the missing buffers
on demand (❹) and execute it (❺). Finally, even without kernel execution, POS
will concurrently load unloaded buffers from the image (❻).

##  Runtime kernel DAG

###  Runtime kernel DAG management

POS leverages a rule-based method to adjust the DAG according to GPU APIs
called from the applications. Since all the driver APIs are known and remain
unchanged during the application execution, we statically analyze them to
generate a rule table and incorporate the table with POS for runtime DAG
adjustment.

The rule for a specific API should consider whether and how to adjust the DAG.
For APIs that don’t read or write GPU execution states (e.g., cudaGetDevice),
the rule simply skips changing the DAG. Other APIs can generally be
categorized into two types:

  1. 1.

APIs that perform computation or memory movement on the GPU.

  2. 2.

APIs that query the device state of the GPU, i.e., checking whether the kernel
launched by the previous APIs have been completed

For the first type of API, we first speculatively get which buffers are read
and write given the kernel it launches using its argument (§4.2). Then, we
construct a data flow and add it to the current DAG. The addition should
consider the data dependencies as well as implicit dependencies defined by the
API. For example, APIs submitted to the same GPU stream should be executed in
a FIFO order. Therefore, we first look up the last node belonging to the
stream (if exists) and then adds it to the node accordingly.

The rule for the second type of API is simple, e.g., upon the finish of
cudaDeviceSynchronize, we just cleared the current DAG. Note that before
clearing the DAG, POS may run a pass on it to analyze how the DAG affects the
C/R (see §5.1).

We generate the rules for APIs manually since the rule depends on the API
behaviors. The human effort is acceptable because (1) the process is
mechanical (described above) and (2) applications only utilize a few number of
GPU APIs during runtime. For example, generating rules for up to 50 out of
2,449 total CUDA 12.3 APIs [51] is sufficient to run all our applications, and
we take 2 human days to generate them (several 10–20 LoC for each API).

###  Argument-based kernel dataflow analyze

Basic method.  POS analyzes the arguments of APIs that launch GPU kernels to
infer their accessed buffers. Since GPU APIs usually have few arguments (the
medium is 5), the overhead of analysis is negligible. Generally, the APIs that
launch GPU kernels can be further categorized based on the types of kernels
they launch:

  * •

Memory move kernels. These kernels transfer data from one GPU (or CPU) buffer
to another. Example: cudaMemcpy (void* dst, const void* src, sizet count, ...)
.

  * •

Computation kernels with known dataflow.  Computation kernels read GPU
buffers, do computations on them and write other GPU buffers. For kernels that
are provided by the driver (e.g., those packaged in cuBLAS), their data flow
is known so we don’t need speculation. Example: cublasSgemm(...,const float
*A,const float *B,const float *beta,float *C,...).

  * •

Computation kernels with opaque dataflow (Opaque kernel).  As GPUs are
general-purpose, kernels can have opaque data flow. Example: cudaLaunchKernel
(const void* func, ... , void** args, ...) .

Analyzing the dataflow of the first two types of APIs does not need
speculation, because the dataflow is already in the arguments of the APIs
[49]. For instance, the (src,count) is the input buffer and (dst,count) is the
output buffer in memory move kernels, and in cublasSgemm, A and B are input
buffers and C is the output buffer. Interestingly, we found both types of
kernels dominate the runtime kernels invoked in the runtime in various
applications, as they are highly-optimized kernels provided by the vendors for
better utilizing vendors’ GPUs. As a result, applications would invoke them
for high performance.

Figure 5: An illustration of inferring vecadd’s buffer accesses buffers (a and
b) using only the arguments of launch kernel API.

Speculative opaque kernel data flow analyze. The challenge is how to infer the
accessed buffers for the opaque kernels. Our observation is that even for
kernels with opaque data flows, we can still guess the accessed buffers from
the opaque arguments (args), because the arguments typically include the
buffer addresses. Specifically, we can treat all the 8 B arguments555The
argument size can be acquired via CUDA’s expory table. However, the type of
the argument is still unknown.  as tentative buffer addresses. Then we compare
these tentative addresses with the buffers allocated on the GPU to filter out
non-address inputs (e.g., offsets or lengths). Figure 5 shows a concrete
example. The vec_add functions on buffers a and b. To infer from the argument
(line 6), we first treat the first two arguments as tentative buffer
addresses; and compare them with the memory buffers allocated for the
application (line 5). If the addresses fall within the range of an allocated
buffer, we will treat the buffer as accessed by the kernel (both read and
write).

By analyzing traces from kernels used by various applications (see Table 1),
we found that simple argument-based speculation can accurately infer all the
buffers of these kernels. (We have run a validation test as well as manually
checked all 162 opaque kernels of various libraries using Pytorch: they only
access the buffers in the arguments.) However, the speculation can still be
incorrect. First, it has false positives, e.g., treating a non-addressed
argument as a buffer address. Fortunately, most false positives can be ruled
out since the GPU uses a specific address range for the buffers (e.g., CUDA
allocates GPU buffers from the high address space) not commonly used in other
argument values (e.g., offsets or lengths). A more subtle issue is
incompleteness, i.e., the speculation may miss some buffers. For example, if
the kernels access memory indirectly like *(a[1024]) = 73, it is impossible to
infer the real buffer accessed by the kernel. Fortunately, GPU kernels rarely
access memory indirectly, simply because it is inefficient to do so [50]. One
possible solution is to conduct a static analysis of how the kernel accesses
buffers based on its arguments, but static analysis is notoriously hard to be
complete due to the possibility of unbounded loops. Existing solutions like
Honeycomb [41] require application developers to annotate the kernel, which
breaks the transparency of OS-level C/R.

Figure 6: An illustration of the instrumentation-based runtime validation in
POS.

Speculative analyze with runtime validation. To ensure correctness, POS
validates the speculation results for opaque dataflow kernels during runtime.
If the validation fails, we will fallback to a slow path for correctness
(§5.1). For transparency, we instrument the kernel binary (PTX ISA [53]) for
the validation [41]. The instrumentation only happens the first time POS meets
a kernel, so it is offline. Furthermore, we only run the instrumented kernels
during C/R to minimize their impact on normal execution.

Figure 6 shows the high-level validation procedure. During offline, we will
instrument the opaque kernel binaries to add a runtime validator that can
return the buffers not accurately inferred by POS (❶). During execution with
C/R, POS will first guess the buffers (❷) and then execute the instrumented
kernel (❸). After the kernel finishes, we will check the validation results
and fallback if necessary (❹). The handle of failed validation is
asynchronous, which minimizes the impact on the application (§5.1). Note that
during runtime, POS will not invoke the instrumented kernel (❺).

More specifically, our instrumented code will validate whether an accessed
buffer is within the inferred buffers. If not, we will record it and notify
POS. For checkpoint, we only instrument the write access of the kernel, as it
affects the checkpoint correctness. For restore, we also instrument the read
access. The detailed instrumentation is as follows. First, we use cuobjdump to
extract PTX code from the kernel. Then, before each global read/write
instruction, we insert a series of comparison instructions to check whether
the read/write address falls within the speculatively inferred buffers. The
inferred buffers are passed as additional arguments of the kernel. If the
validation fails, we will record the buffers causing the failure. To quickly
find the required buffer, we will pass the allocated buffers as a list to the
kernel (via an argument). Each buffer is accompanied by a validate bit. The
kernel will mark the validate bit if the read/write address causing the
failures is within the allocated buffer. Finally, to report the validation
results to the POS, we also pass a flag to the instrumented kernel such that
POS can check whether the verification is ok.

Empirically, we found kernel instrumentation has little overhead for two
reasons. First, the check is lightweight if the validation passes, e.g., no
need to search and mark which buffer caused the speculation failure. Note that
we have met no runtime violations in various applications. Second, the
instrumented kernels are invoked less often, e.g., less than 20% in our
evaluated workloads (§6.4).

##  DAG-based Checkpoint and Restore

###  Concurrent C/R execution

POS concurrently executes the C/R with GPU execution to hide the C/R costs. To
ensure correctness, we retrofit the following classic CPU techniques in
software for the GPU with the help of the kernel DAG. This section first
describes the common cases when the DAG generated with speculation is correct,
and then elaborate on the fallback path.

Figure 7: An illustration of how POS uses the DAG information to correctly and
efficiently overlap GPU kernel execution with buffer checkpointing.

Soft CoW.  Copy-on-write (CoW) pessimistically isolates the application’s
accesses to buffers being checkpointed, i.e., before modifying the buffer, we
first copy it to a new buffer, and execute the kernel with the new buffer. It
is typically used in checkpoint mostly cases, e.g., for fault tolerance [77,
66]. POS leverages the dataflow information in the DAG to copy the buffer in
time. As shown in Figure 7, before launching a kernel ($K0$), we will check if
its access buffer conflicts with the buffer being checkpointed (➀). If so, we
will utilize an on-device GPU buffer managed by the POS to first copy the
conflicted buffer to it, and then launch the kernel. Otherwise ($K1$), if the
buffer (buffer 1) has been checkpointed, we can launch the kernel directly
(➃).

A tricky case is how to let the kernels with opaque dataflow using the CoW
buffer: we cannot modify its arguments due to the lack of knowledge on how it
uses the arguments. Thus, we instead let the checkpoint update the copied
buffer using the CoW buffer (➁). This may cause extra copies, as the buffer
has to be copied three times, once before the kernel execution666This copy is
efficient since it is on-device. and twice for the checkpoints. To reduce
extra copies, we logically break a buffer into small trunks to checkpoint it
in a fine-grained way, As a result, we can break the original copy as soon as
possible and copy the new buffer. Since the buffers are copied fine-grained,
it allows us to estimate the time to finish the checkpointing of the buffer,
so we can delay the kernel execution if we will finish checkpointing the
conflict buffer soon (e.g., $k2$ ➂).

Soft dirty-bit.  Besides CoW, we can optimistically execute the checkpointing
with the applications and re-copy buffers being concurrently modified (dirty
buffers). The challenge is that there is no dirty bit support in existing
GPUs, so we implement it in software with kernel DAG (Figure 7). During
execution, once part of the DAG is cleared due to finished kernel execution
(❶), we will record output buffers in the DAG as dirty buffers. After all the
buffers have been checkpointed, we will stop execution and wait for all
executing kernels to finish to record the final dirty buffers. Finally, we
will re-copy all the dirty buffers to finish the checkpoint (❷).

Optimization: checkpoint the kernel DAG instead of dirty buffers. One issue is
that the GPU can modify a large amount of buffers during execution, e.g., in
the case of training. POS leverages the kernel DAG to reduce the number of
dirty buffers with re-execution. Specifically, if the number of dirty buffers
exceeds a pre-defined threshold during concurrent execution, we will stop
launching kernels on the GPU, record all future kernels sent by the CPU in the
DAG ($K1$ in the example in Figure 7). Meanwhile, we will concurrently copy
all existing dirty buffers to the image. After all the dirty buffers have been
checkpointed, we will checkpoint the DAG (❸) such that POS can re-execute the
kernels for restoring.

Soft on demand load.  We use on-demand buffer load to concurrently execute
kernels without requiring all the buffers to be loaded from the image. After
the kernel DAG has been restored from the image, we will first sort the
buffers in it using a topological sort based on the DAG. Afterward, we will
copy the buffers in this order, and launch a new kernel as long as all its
speculated buffer has been copied. The new kernel can either be unexecuted
kernels in the restored DAG or the new kernel sent by the application.

Figure 8: An illustration of the correctness problem in concurrent
checkpointing (a) and restoring (b) when the speculatively inferred buffer
access of a kernel mismatches the real accessed buffer.

Table 1:  Emerging GPU applications evaluated. PPO does not have an inference
task. Single GPU is unable to run the llama-2 training.

| ResNET (152M) | GPT-2 (1.56B) | BERT (340M) | PPO (335.5M) | llama-2 (7B)
---|---|---|---|---|---
Workloads | Train | Inference | Train | Inference | Train | Inference | Train | Inference | Train | Inference
Application library |  torchvision [62] |  HuggingFace [17] |  HuggingFace [17] |  OpenAI Gym [55] |  Meta Llama 2 [42]
Total GPU memory | 1.3 GB | 354.4 MB | 30.8 GB | 6.5 GB | 15.6 GB | 5.8 GB | 5.6 GB | N/A | N/A | 51.7GB
#GPU buffers | 224 | 52 | 1,044 | 249 | 409 | 271 | 97 | N/A | N/A | 328
Active GPU kernels | 3,562 | 1,221 | 125,482 | 72,710 | 14,754 | 3,025 | 628,886 | N/A | N/A | 825,627

Fallback handlers. Though we don’t experience any validation fails during
runtime, we should provide a functional fallback handler for correctness. The
fallback strategy for the techniques described above is different. For
concurrent checkpointing, an incorrect guess of the write-accessed buffers
could cause a concurrently checkpointed buffer to be incorrect. For example,
in Figure 8 (a), the checkpoint and the kernel $k0$ have a race condition on
buffer 0, which could not happen if the speculation is correct. Under soft
dirty bit, we can trivially mark buffer 0 as dirty and use retransmission to
fix the problem. Under soft CoW, it is more tricky: we are unable to recover
buffer 0 to the correct version, as the buffer has not been copied on write.
Thus, we can only restart the entire checkpoint. To prevent re-checkpointing
caused by the kernel again, one approach is to record the kernel that causes
the validation failure and fall back to a stop-the-world checkpoint once we
encounter it again. Alternatively, we can use the soft dirty bit for
concurrent checkpointing.

For concurrent restoring (Figure 8 (b)), a missing GPU buffer (0) will cause
$k0$ reporting a failure—either due to unallocated buffers or validation
failures. To fall back, an intuitive approach is to re-execute $k0$ when the
missing buffer has been loaded. However, this is incorrect as the kernel can
update another buffer (buffer 1) using the wrong buffer content, e.g.,
partially loaded buffer 0. Hence, we will first reload all the original
buffers of the kernel in the image before re-executing it. The reload is done
recursively based on the DAG if there are multiple ongoing kernels.

Figure 9: Coordinated CPU and GPU checkpointing.

###  Coordinating CPU and GPU Checkpointing

Coordinated CPU and GPU checkpoint.  In general, CPU memory also needs to be
checkpointed with the GPU memory. The problem is that both checkpoints compete
for the bandwidth used for C/R (e.g., the network bandwidth in case of
migration), so checkpointing the CPU affects the performance of the GPU (and
vice versa), see Figure 9 (a). The contention slows down the individual
process, causing extra dirty memory in the case of delta copies or more CoW
buffers.

POS adopts a coordinated design that sequentially checkpoints the GPU and CPU
memory. Figure 9 (b) shows an example using the soft dirty bit, but the same
methodology also applies to soft CoW, which we omitted due to space
limitations. With coordinated checkpoints, we can speed up the checkpoint of
each device’s memory.

Checkpointed memory deduplication.  For GPU processes, the memory of the CPU
and GPU may contain duplicate content, e.g., the model parameters. Identifying
and deduplicating them can dramatically reduce the checkpoint time due to the
reduced image size (§6.2). POS leverages the kernel DAG for deduplication.
Specifically, once a memory buffer is copied from the host to the GPU ($K2$ in
Figure 2), we track the upstream CPU buffer at POS. If the host and GPU memory
don’t change up to the checkpoint time, we can treat them as identical and
avoid checkpointing the GPU buffer.

There are two approaches to tracking whether a buffer has changed. First, we
can use dirty bits (both in hardware at the CPU side and in software at the
GPU). Second, we can use the checksum. POS adopts a hybrid solution: we use
hardware dirty bit to track whether the copied buffers have been changed at
the CPU (since the copy), and use the checksum to verify that the GPU buffer
is the same as the CPU buffer, so it transitively implies that the GPU buffer
has not been changed. The design choice is based on the fact that using soft
dirty bits require running the instrumented kernels during normal execution,
which has a performance overhead.

More specifically, we apply the cyclic redundancy check (CRC) checksum to
check whether the GPU buffer is identical to the CPU buffer. After copying a
CPU buffer to the GPU, we clear the CPU-side dirty bits relating to the
buffer; and calculate its CRC checksum on the GPU before executing kernels on
it. The checksum calculation is orders of magnitude faster than the memory
copy, so the overhead is negligible both during copy and during checkpoint.
For example, in the BERT-340M inference workload, it takes 4 ms to calculate
the checksums for 1.12 GB of HBM which are potentially duplicate, while taking
206 ms for transferring 4.6 GB of remaining non-duplicate HBM. However, some
workloads contain a large proportion of duplicate memories, introducing non-
negligible checksum calculation overhead. For instance, in llama2-13b
inference, it takes 155 ms to calculate the checksums for 50.6 GB of duplicate
HBM, while taking only 46 ms to transfer 1.04 GB of non-duplicate HBM. Note
that the checksum calculation can be overlapped with the kernel execution as
well as the checkpoint to minimize its overhead.

Priority-based GPU memory copy coordination. One issue we encounter during
overlapped checkpointing is that the memory transfer in checkpointing competes
for memory transfer bandwidth as well as memory copy engines in GPUs. Though
the PCIe bandwidth can be shared between memory copies, GPUs have a few memory
copy engines [3]. As a result, blindly checkpointing buffers at full speed
would block host-to-device memory transfers required by the application, which
would make the application execution stall. Since the application memory
transfers are less often, we prioritize the application memory transfer over
the checkpointing.

##  Evaluation

We have implemented POS from scratch in 2,4885 C++ LoC, excluding the GPU
virtualization framework, communication libraries and integration with CRIU.
We have carefully tuned our virtualization frameworks following existing
methods [74, 18] such that virtualization has no overhead to our evaluated
applications.

Testbed.  All our experiments are conducted on two local servers, each with
eight A100 (80 GB HBM, SXM4) GPUs, two Intel Xeon Gold 6348 CPUs (total 56
cores), and 1 TB of host memory. Two servers are connected via a 200 Gbps
network (ConnectX-6 InfiniBand). Except for the live migration experiments
that requires 2 GPUs (so we use two A100s), other experiments only use a
single A100. Both servers install CUDA 12.3 [51].

Baselines.  To the best of our knowledge, no available open-source OS-level
C/R systems exist that can work on recent NVIDIA GPUs like A100. Therefore,
we’ve implemented the state-of-the-art system—Singularity [67] in our codebase
as the baseline (Base). Specifically, we’ve implemented the stop-the-world GPU
C/R described in the original paper. We’ve carefully optimized its
performance, e.g., leveraging pin memory to saturate the PCIe bandwidth during
C/R. We also apply our CPU-GPU coordinated checkpoint described in §5.2 to it.
Meanwhile, we also compare POS with Pytorch’s default checkpoint method, a
task-level checkpointing mechanism that supports fault tolerance (Task). For
task-level restore, we directly compare with starting the application’s GPU
process from scratch. For task-level migration, we omit its comparison
because, to the best of our knowledge, no existing task-level C/R system can
lively migrate GPU applications like llama-2.

Evaluated applications.  We evaluated various ML tasks (see Table 1), ranging
from training to inference, whose domains span vision (ResNET), language
(BERT), large foundational models (llama-2 and GPT-2), and reinforcement
learning (PPO). The tasks are chosen from different applications that use
various task libraries, aiming to show that POS, generalizes to various
applications.

###  Checkpoint performance

Figure 10:  The runtime overhead of POS checkpointing. The time is normalized
to a setup when no checkpoint is taken (non-fault-tolerant).

Figure 11:  Time breakdown of how application execution time (exe.) is
affected by the OS-level checkpointing.
Figure 12:  Migration timelines for different workloads. (blue line for Base,
red line for POS)

Experiment setup.  To evaluate the effectiveness of POS in checkpointing, we
use workloads containing training ML models from vision to RL, and evaluate
the performance when using POS and other baselines to provide fault tolerance.
For POS and Base, we implement a manager that periodically takes checkpoints
on the training process. Since training tasks benefit from a high checkpoint
frequency to reduce wasted computing time [43, 75], we chose two checkpoint
frequencies, every 3 and 10 iterations. As POS and Base don’t have the
applications’ iteration semantics, for a fair comparison with Task, we set the
interval to be the same time required to run 3 iterations when no checkpoint
is needed. Finally, as the image source can significantly impact the
performance of checkpointing, we follow prior work by using the host CPU
memory as buffers for accelerating checkpoints [75], where the in-memory
buffer is flushed to the SSD in the background.

Overall performance.  Figure 10 shows the checkpointing overhead of POS and
other baselines for different workloads. We run the workload 1,000 iterations
and report the time per iteration. We can see that POS is 5.5–72.2% and
12.4–62.3% faster than Base and Task, respectively. Compared to Base, the
performance improvements mainly come from executing applications concurrently
with checkpointing, as Base must stop the application during this period. Our
next paragraph will break down the time in detail. Compared to Task, the
performance improvements come from the fact that the default checkpointing
mechanism provided by the library has not adopted optimization like in-memory
buffer [43, 75]. As an OS-level C/R system, POS can transparently integrating
so without requiring applications to change their code.

Breakdown of checkpointing.  Figure 11 shows the breakdown of the GPU
checkpointing time for different workloads. We omit the CPU checkpointing time
because it is negligible in our applications: (1) CPU memory rarely changes
and (2) the CPU-side can use incremental checkpointing so the real
checkpointed data is negligible. Therefore, we focus on the GPU part
checkpoint time to understand how POS behaves. From GPT-2, we can see that POS
has orders of magnitude smaller stop time (1 ms vs. 726 ms) compared with
Base, since the application can run concurrently with the checkpointing. This
contributes most to the low overhead of POS on the training applications. On
the other hand, POS does incur extra overhead during application due to the
CoW and instrumented kernels (e.g., GPT-2). Nevertheless, such overhead is
negligible compared to the overall time saved by concurrent execution.

###  Migration performance

Table 2:  The checkpointed image size of our evaluated applications.

| ResNET (152M) | GPT-2 (1.56B) | BERT (340M) | PPO (335.5M) | llama-2 (7B)
---|---|---|---|---|---
Workloads | Train | Inference | Train | Inference | Train | Inference | Train | Inference | Train | Inference
Base | 1.2 GB | 330.1 MB | 28.7 GB | 6.1 GB | 14.5 GB | 5.4 GB | 5.2 GB | N/A | N/A | 48.1 GB
POS | 1.2 GB | 144.5 MB | 28.7 GB | 709.0 MB | 14.5 GB | 4.3 GB | 5.2 GB | N/A | N/A | 970.7 MB

Figure 13:  A breakdown of the live GPU process migration timeline for the
llama 2 inference and BERT training with POS C/R.

Figure 14:  The runtime overhead added by the instrumentation in (a) training
and (b) inference workloads. (c) The ratio of the instrumented kernel in the
overall workloads.

Experiment setup.  To demonstrate the effectiveness of live migration enabled
by POS, we measure the performance using different approaches to migrate a GPU
process from one machine to another. For both POS and Base, we will use RDMA
to checkpoint the GPU process directly from the source node to the target
node’s GPU memory in a peer-to-peer manner, and restoring execution
accordingly. POS uses soft dirty bit to enable the application to run during
migration (being live), while Base must stop the application to perform
checkpointing. We don’t evaluate Task as no existing framework can migrate our
applications.

Overall performance.  As shown in Figure 12, POS can migrate the GPU process
58.0–95.2% faster than Base. Note that both approaches start the migration at
the same time. POS is faster because: first, it can concurrently execute the
checkpoint in the migration with the application. Second, it reduces the
checkpoint image size (and thus, the transfer time) for inference workloads
with a high duplication ratio between CPU and GPU memory (see Table 2). By
profiling the application semantics, we found the duplication is due to the
model parameters, as expected. This also explains why training has no
duplicated memory, because the GPU part of the model parameters needs to be
extensively updated. Finally, POS reduces the restore time of GPU context
creation with context pool. §6.3 will discuss the restore performance in
detail.

Breakdown of checkpointing and restoring during migration.  Figure 13 shows
the breakdown timeline of the live migrating a llama 2 inference workload and
BERT training workloads. Other workloads share similar patterns so we omit
them. The training and inference workloads share the benefit of POS that the
application can run concurrently with the checkpointing in migration.
Nevertheless, they do exhibit different behaviors. For the inference
workloads, first, the checkpointing time on GPU can be effectively minimized
with deduplication, as most GPU buffers are used for storing model parameters,
which have already been checkpointed when checkpointing the CPU memory. For
the training workloads, POS has to additionally checkpoint the GPU memory as
the GPU extensively modifies it. Moreover, training has a relatively long
dirty buffer copy phase (compared to the inference) because many dirty buffers
are generated. Therefore, we will stop application execution before all the
buffer has been copied to reduce extra buffer copied, and use re-execution to
recover them.

###  Restoring performance

Figure 15:  Restore latency of Task, Base and POS

Experiment setup.  Restoring from a well-formed, checkpointed image can
improve application startup, which is critical in scenarios like serverless
computing [76, 12]. This section demonstrates (1) how OS-level restores can
enhance GPU application startup and (2) how soft on-demand load and execution
pools of POS can improve OS-level restoration on a GPU. Since startup
performance is negligible in training workloads, we focus on the startup
performance of inference workloads, which is important in scenarios like
serverless computing [19].

Overall performance.  Figure 15 shows the overall startup performance of
various inference tasks. First, OS-level restoring can have orders of
magnitude higher performance than task-level startup: POS and Base is
6.8–111.9 $\times$ and 2.1–5.7 $\times$ faster than Task—cold start using
Docker, respectively. This is aligned with existing works on optimization
container startup with OS-level restoring [76, 12]. Note that these works
don’t support GPU but POS do. POS is further 3.2–19.5 $\times$ faster than
Base thanks to skipping most GPU context creation and using soft on-demand
load to concurrently execute restore with the application.

Figure 16:  Restore latency breakdown of Task, Base and POS

Breakdown of restoring. Figure 16 further shows a breakdown of the restoring
time of various workloads. Since the Task is so slow, we skip it in the
discussion. For workloads with relatively small GPU footprint, i.e., ResNET,
BERT, and GPT-2, the memory loading is small (e.g., 354.4 MB–6.5 GB), so the
restoring time is bottlenecked by the GPU context creation. POS eliminates so
with the context pool. For llama 2, as the GPU memory footprint is large (51.7
GB), loading memory to it contributes a significant portion of restoring time.
This causes the stalls in Base. With soft on-demand loading, POS can
effectively overlap the memory loading with the kernel execution, and reduces
most of the stall time.

###  Analysis of the instrumentation overhead

Figure 14 (a) and (b) show the runtime overhead added by the instrumentation.
Note that POS only runs the instrumented code during C/R such that _these
overheads do not exist during normal application execution_. We observe a
related small slowdown 2–21% for training workloads, and 1–19% for inference
workloads. Although the instrumentation—add a bunch of checks for each global
store (for C) and global load and store (for R)—is not trivial, the observable
overhead is small because: (1) the kernels won’t intensively interact with the
global memory for high computation intensity and (2) the instrumented kernels
are only a small portion of the overall kernels in the workloads, as shown in
Figure 14 (c). The workload with the biggest overhead in training is RL (21%).
This is because it has many kernels to be instrumented: 54%. On the other
hand, the biggest overhead in inference is llama 2 (19%). This is because its
instrumented kernels interact with global memory more intensively due to its
larger model size, e.g., we found a 3.2 $\times$ overhead for an instrumented
unrolled elementwise kernel on llama 2, while it is only 1.7 $\times$ in
GPT-2.

###  Discussion

Multi-GPU applications. POS currently only supports single-GPU applications,
but our methodology generalizes to multiple-GPU setups, e.g., by virtualizing
group communication primitives and speculatively inferring their buffer access
semantics. We plan to extend POS to support processes that use multiple GPUs
in the future. Nevertheless, we should note that this does not pose a
limitation on the applicability of POS because single-GPU applications still
dominate in modern datacenters. First, a single GPU is sufficient to run many
workloads, such as small model training (and finetuning [39]) and inference of
small to medium-scale models (e.g., llama-2 13 B). Second, a majority of GPU
jobs deployed on the cloud only use one GPU, as observed by recent studies
[28, 79].

Applications with few kernels or buffers.  Our fine-grained coordination works
best when applications use kernels in a fine-grained manner (i.e., invoking
numerous kernels), and allocate GPU buffers in a fine-grained manner (i.e.,
allocating many buffers). Fortunately, for emerging ML applications, it is
commonly observed that tens and hundreds of kernels are invoked during runtime
[23]. Meanwhile, existing frameworks like PyTorch allocate GPU buffers in a
fine-grained tensor granularity [59], i.e., one for each tensor. As each AI
model may have tens of thousands of tensors (e.g., tens per layer), emerging
GPU applications naturally allocate buffers fine-grainedly.

##  Related work

GPU and CPU C/R. POS continues the line of research on C/R for different
tasks, with a particular focus on OS-level C/R for GPU applications and on
efficient and concurrent C/R execution. Though concurrent OS-level C/R has
been extensively explored for processes that only run on the CPU [14, 40, 24,
66, 71, 10, 35, 25, 77], to the best of our knowledge, existing C/R systems
that support GPU all leverage a stop-the-world design [48, 47, 21, 15, 67],
which we have shown has notable performance degradation or application
downtime in various scenarios. Meanwhile, there are also many task-level C/R
designs, e.g., for ML tasks [44, 75, 16], HPC tasks [57] or leverage NVM for
acceleration [56]. They are not transparent and only support a limited
category of applications. They also don’t support C/R executed outside of the
task (e.g., from the cloud vendors).

DAG-based computing graph. The DAG-based abstraction, originated from
distributed and concurrent programming literatures [7, 29], has been widely
used in various systems for abstracting and coordinating concurrent execution
[64, 81, 80, 45, 63]. POS retrofits the DAG and uses it to abstract the
computation on the GPU. As a result, we can coordinate application execution
with C/R for better performance.

Analyzing GPU kernels.  Many existing works analyze GPU programs (kernels),
either through static analysis or with runtime analysis [8, 37, 5, 6, 33, 36,
38]. The closest to POS’s design is Honeycomb [41], a GPU software TEE that
validates the runtime analysis of kernels to avoid invalid buffer accesses.
The key difference is that Honeycomb requires users to annotate the buffers
accessed by each kernel for validation, while POS does this transparently
through speculation. POS can do so because a failed speculation won’t affect
application execution correctness (only affecting C/R performance), while it
can break the security guarantee in Honeycomb.

##  Conclusion

We present ParallelGPUOS, the first transparent checkpoint and restore (C/R)
system that can concurrently run C/R with the GPU execution. We show it is
possible to achieve so by bridging the semantic gap between OS and GPU
execution through abstracting GPU execution as a fine-grained dataflow graph
(DAG). The graph is speculatively constructed during runtime for efficiency
without compromising transparency, and the execution is validated during C/R
for correctness. Thanks to the DAG, POS can concurrently and correctly execute
C/R with the application execution, thereby bringing orders of magnitude
higher performance in various application scenarios.

## References

  * [1] Al-Kiswany, S., Subhraveti, D., Sarkar, P., and Ripeanu, M. Vmflock: virtual machine co-migration for the cloud.  In Proceedings of the 20th ACM International Symposium on High Performance Distributed Computing, HPDC 2011, San Jose, CA, USA, June 8-11, 2011 (2011), A. B. Maccabe and D. Thain, Eds., ACM, pp. 159–170.
  * [2] AMD.  amdgpuplugin.c.  https://github.com/checkpoint-restore/criu/blob/criu-dev/plugins/amdgpu/amdgpu_plugin.c, 2024\.
  * [3] Bakita, J., and Anderson, J. H. Demystifying nvidia gpu internals to enable reliable gpu management.
  * [4] Bernie Wu, Y. T. Achieving k8s and public cloud operational efficiency using a new checkpoint/restart feature for gpus.  https://www.nvidia.com/gtc/posters/#/session/1705106137731001cNAN, 2024\.
  * [5] Betts, A., Chong, N., Donaldson, A. F., Ketema, J., Qadeer, S., Thomson, P., and Wickerson, J. The design and implementation of a verification technique for GPU kernels.  ACM Trans. Program. Lang. Syst. 37, 3 (2015), 10:1–10:49.
  * [6] Betts, A., Chong, N., Donaldson, A. F., Qadeer, S., and Thomson, P. Gpuverify: a verifier for GPU kernels.  In Proceedings of the 27th Annual ACM SIGPLAN Conference on Object-Oriented Programming, Systems, Languages, and Applications, OOPSLA 2012, part of SPLASH 2012, Tucson, AZ, USA, October 21-25, 2012 (2012), G. T. Leavens and M. B. Dwyer, Eds., ACM, pp. 113–132.
  * [7] Blumofe, R. D., and Leiserson, C. E. Space-efficient scheduling of multithreaded computations.  In Proceedings of the Twenty-Fifth Annual ACM Symposium on Theory of Computing, May 16-18, 1993, San Diego, CA, USA (1993), S. R. Kosaraju, D. S. Johnson, and A. Aggarwal, Eds., ACM, pp. 362–371.
  * [8] Chiang, W., Gopalakrishnan, G., Li, G., and Rakamaric, Z. Formal analysis of GPU programs with atomics via conflict-directed delay-bounding.  In NASA Formal Methods, 5th International Symposium, NFM 2013, Moffett Field, CA, USA, May 14-16, 2013. Proceedings (2013), G. Brat, N. Rungta, and A. Venet, Eds., vol. 7871 of Lecture Notes in Computer Science, Springer, pp. 213–228.
  * [9] Clark, C., Fraser, K., Hand, S., Hansen, J. G., Jul, E., Limpach, C., Pratt, I., and Warfield, A. Live migration of virtual machines.  In 2nd Symposium on Networked Systems Design and Implementation NSDI (2005), May 2-4, 2005, Boston, Massachusetts, USA, Proceedings (2005), A. Vahdat and D. Wetherall, Eds., USENIX.
  * [10] CRIU.  CRIU.  https://criu.org/Main_Page, 2024.
  * [11] Du, D., Liu, Q., Jiang, X., Xia, Y., Zang, B., and Chen, H. Serverless computing on heterogeneous computers.  In ASPLOS ’22: 27th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Lausanne, Switzerland, 28 February 2022 - 4 March 2022 (2022), B. Falsafi, M. Ferdman, S. Lu, and T. F. Wenisch, Eds., ACM, pp. 797–813.
  * [12] Du, D., Yu, T., Xia, Y., Zang, B., Yan, G., Qin, C., Wu, Q., and Chen, H. Catalyzer: Sub-millisecond startup for serverless computing with initialization-less booting.  In ASPLOS ’20: Architectural Support for Programming Languages and Operating Systems, Lausanne, Switzerland, March 16-20, 2020 (2020), J. R. Larus, L. Ceze, and K. Strauss, Eds., ACM, pp. 467–481.
  * [13] Duato, J., Peña, A. J., Silla, F., Mayo, R., and Quintana-Ortí, E. S. rcuda: Reducing the number of gpu-based accelerators in high performance clusters.  In Proceedings of the 2010 International Conference on High Performance Computing & Simulation, HPCS 2010, June 28 - July 2, 2010, Caen, France (2010), W. W. Smari and J. P. McIntire, Eds., IEEE, pp. 224–231.
  * [14] Egwutuoha, I. P., Levy, D., Selic, B., and Chen, S. A survey of fault tolerance mechanisms and checkpoint/restart implementations for high performance computing systems.  The Journal of Supercomputing 65, 3 (2013), 1302–1326.
  * [15] Eiling, N., Baude, J., Lankes, S., and Monti, A. Cricket: A virtualization layer for distributed execution of CUDA applications with checkpoint/restart support.  Concurr. Comput. Pract. Exp. 34, 14 (2022).
  * [16] Eisenman, A., Matam, K. K., Ingram, S., Mudigere, D., Krishnamoorthi, R., Nair, K., Smelyanskiy, M., and Annavaram, M. Check-n-run: a checkpointing system for training deep learning recommendation models.  In 19th USENIX Symposium on Networked Systems Design and Implementation, NSDI 2022, Renton, WA, USA, April 4-6, 2022 (2022), A. Phanishayee and V. Sekar, Eds., USENIX Association, pp. 929–943.
  * [17] Face, H. The ai community building the future.  https://huggingface.co, 2024.
  * [18] Fingler, H., Zhu, Z., Yoon, E., Jia, Z., Witchel, E., and Rossbach, C. J. DGSF: disaggregated gpus for serverless functions.  In 2022 IEEE International Parallel and Distributed Processing Symposium, IPDPS 2022, Lyon, France, May 30 - June 3, 2022 (2022), IEEE, pp. 739–750.
  * [19] Fu, Y., Xue, L., Huang, Y., Brabete, A., Ustiugov, D., Patel, Y., and Mai, L. Serverlessllm: Locality-enhanced serverless inference for large language models.  CoRR abs/2401.14351 (2024).
  * [20] Gao, Y., Shi, X., Lin, H., Zhang, H., Wu, H., Li, R., and Yang, M. An empirical study on quality issues of deep learning platform.  In 45th IEEE/ACM International Conference on Software Engineering: Software Engineering in Practice, SEIP@ICSE 2023, Melbourne, Australia, May 14-20, 2023 (2023), IEEE, pp. 455–466.
  * [21] Garg, R., Mohan, A., Sullivan, M. B., and Cooperman, G. CRUM: checkpoint-restart support for cuda’s unified memory.  In IEEE International Conference on Cluster Computing, CLUSTER 2018, Belfast, UK, September 10-13, 2018 (2018), IEEE Computer Society, pp. 302–313.
  * [22] Gu, J., Hua, Z., Xia, Y., Chen, H., Zang, B., Guan, H., and Li, J. Secure live migration of SGX enclaves on untrusted cloud.  In 47th Annual IEEE/IFIP International Conference on Dependable Systems and Networks, DSN 2017, Denver, CO, USA, June 26-29, 2017 (2017), IEEE Computer Society, pp. 225–236.
  * [23] Han, M., Zhang, H., Chen, R., and Chen, H. Microsecond-scale preemption for concurrent GPU-accelerated DNN inferences.  In 16th USENIX Symposium on Operating Systems Design and Implementation (OSDI 22) (Carlsbad, CA, July 2022), USENIX Association, pp. 539–558.
  * [24] Hardy, N. Keykos architecture.  SIGOPS Oper. Syst. Rev. 19, 4 (oct 1985), 8–25.
  * [25] Hargrove, P. H., and Duell, J. C. Berkeley lab checkpoint/restart (blcr) for linux clusters.  In Journal of Physics: Conference Series (2006), vol. 46, IOP Publishing, p. 067.
  * [26] Hines, M. R., Deshpande, U., and Gopalan, K. Post-copy live migration of virtual machines.  ACM SIGOPS Oper. Syst. Rev. 43, 3 (2009), 14–26.
  * [27] Hines, M. R., and Gopalan, K. Post-copy based live virtual machine migration using adaptive pre-paging and dynamic self-ballooning.  In Proceedings of the 5th International Conference on Virtual Execution Environments, VEE 2009, Washington, DC, USA, March 11-13, 2009 (2009), A. L. Hosking, D. F. Bacon, and O. Krieger, Eds., ACM, pp. 51–60.
  * [28] Hu, Q., Sun, P., Yan, S., Wen, Y., and Zhang, T. Characterization and prediction of deep learning workloads in large-scale GPU datacenters.  In International Conference for High Performance Computing, Networking, Storage and Analysis, SC 2021, St. Louis, Missouri, USA, November 14-19, 2021 (2021), B. R. de Supinski, M. W. Hall, and T. Gamblin, Eds., ACM, p. 104.
  * [29] Isard, M., Budiu, M., Yu, Y., Birrell, A., and Fetterly, D. Dryad: distributed data-parallel programs from sequential building blocks.  In Proceedings of the 2007 EuroSys Conference, Lisbon, Portugal, March 21-23, 2007 (2007), P. Ferreira, T. R. Gross, and L. Veiga, Eds., ACM, pp. 59–72.
  * [30] Jeon, M., Venkataraman, S., Phanishayee, A., Qian, J., Xiao, W., and Yang, F. Analysis of Large-Scale Multi-Tenant GPU clusters for DNN training workloads.  In 2019 USENIX Annual Technical Conference (USENIX ATC 19) (Renton, WA, July 2019), USENIX Association, pp. 947–960.
  * [31] Johnson, E. Starting up faster with aws lambda snapstart.  https://aws.amazon.com/cn/blogs/compute/starting-up-faster-with-aws-lambda-snapstart/, 2024\.
  * [32] Jonas, E., Schleier-Smith, J., Sreekanti, V., Tsai, C., Khandelwal, A., Pu, Q., Shankar, V., Carreira, J., Krauth, K., Yadwadkar, N. J., Gonzalez, J. E., Popa, R. A., Stoica, I., and Patterson, D. A. Cloud programming simplified: A berkeley view on serverless computing.  CoRR abs/1902.03383 (2019).
  * [33] Kamath, A. K., and Basu, A. iguard: In-gpu advanced race detection.  In SOSP ’21: ACM SIGOPS 28th Symposium on Operating Systems Principles, Virtual Event / Koblenz, Germany, October 26-29, 2021 (2021), R. van Renesse and N. Zeldovich, Eds., ACM, pp. 49–65.
  * [34] Krizhevsky, A., Sutskever, I., and Hinton, G. E. Imagenet classification with deep convolutional neural networks.  In Advances in Neural Information Processing Systems 25: 26th Annual Conference on Neural Information Processing Systems 2012. Proceedings of a meeting held December 3-6, 2012, Lake Tahoe, Nevada, United States (2012), P. L. Bartlett, F. C. N. Pereira, C. J. C. Burges, L. Bottou, and K. Q. Weinberger, Eds., pp. 1106–1114.
  * [35] Laadan, O., and Hallyn, S. E. Linux-cr: Transparent application checkpoint-restart in linux.  In Linux Symposium (2010), vol. 159, Citeseer.
  * [36] Leung, A., Gupta, M., Agarwal, Y., Gupta, R., Jhala, R., and Lerner, S. Verifying GPU kernels by test amplification.  In ACM SIGPLAN Conference on Programming Language Design and Implementation, PLDI ’12, Beijing, China - June 11 - 16, 2012 (2012), J. Vitek, H. Lin, and F. Tip, Eds., ACM, pp. 383–394.
  * [37] Li, G., and Gopalakrishnan, G. Scalable smt-based verification of GPU kernel functions.  In Proceedings of the 18th ACM SIGSOFT International Symposium on Foundations of Software Engineering, 2010, Santa Fe, NM, USA, November 7-11, 2010 (2010), G. Roman and A. van der Hoek, Eds., ACM, pp. 187–196.
  * [38] Li, G., Li, P., Sawaya, G., Gopalakrishnan, G., Ghosh, I., and Rajan, S. P. GKLEE: concolic verification and test generation for gpus.  In Proceedings of the 17th ACM SIGPLAN Symposium on Principles and Practice of Parallel Programming, PPOPP 2012, New Orleans, LA, USA, February 25-29, 2012 (2012), J. Ramanujam and P. Sadayappan, Eds., ACM, pp. 215–224.
  * [39] Liao, C., Sun, M., Yang, Z., Chen, K., Yuan, B., Wu, F., and Wang, Z. Adding nvme ssds to enable and accelerate 100b model fine-tuning on a single GPU.  CoRR abs/2403.06504 (2024).
  * [40] Litzkow, M., Tannenbaum, T., Basney, J., and Livny, M. Checkpoint and migration of unix processes in the condor distributed processing system.  Tech. rep., University of Wisconsin-Madison Department of Computer Sciences, 1997.
  * [41] Mai, H., Zhao, J., Zheng, H., Zhao, Y., Liu, Z., Gao, M., Wang, C., Cui, H., Feng, X., and Kozyrakis, C. Honeycomb: Secure and efficient GPU executions via static validation.  In 17th USENIX Symposium on Operating Systems Design and Implementation (OSDI 23) (Boston, MA, 2023), USENIX Association, pp. 155–172.
  * [42] Meta.  Llama 2.  https://github.com/meta-llama/llama, 2024.
  * [43] Microsoft.  Boost checkpoint speed and reduce cost with nebula.  https://learn.microsoft.com/en-us/azure/machine-learning/reference-checkpoint-performance-for-large-models?view=azureml-api-2&tabs=PYTORCH, 2024\.
  * [44] Mohan, J., Phanishayee, A., and Chidambaram, V. Checkfreq: Frequent, fine-grained DNN checkpointing.  In 19th USENIX Conference on File and Storage Technologies, FAST 2021, February 23-25, 2021 (2021), M. K. Aguilera and G. Yadgar, Eds., USENIX Association, pp. 203–216.
  * [45] Murray, D. G., McSherry, F., Isaacs, R., Isard, M., Barham, P., and Abadi, M. Naiad: a timely dataflow system.  In ACM SIGOPS 24th Symposium on Operating Systems Principles, SOSP ’13, Farmington, PA, USA, November 3-6, 2013 (2013), M. Kaminsky and M. Dahlin, Eds., ACM, pp. 439–455.
  * [46] Ng, K. K. W., Demoulin, H. M., and Liu, V. Paella: Low-latency model serving with software-defined GPU scheduling.  In Proceedings of the 29th Symposium on Operating Systems Principles, SOSP 2023, Koblenz, Germany, October 23-26, 2023 (2023), J. Flinn, M. I. Seltzer, P. Druschel, A. Kaufmann, and J. Mace, Eds., ACM, pp. 595–610.
  * [47] Nukada, A., Suzuki, T., and Matsuoka, S. Efficient checkpoint/restart of CUDA applications.  Parallel Comput. 116 (2023), 103018.
  * [48] Nukada, A., Takizawa, H., and Matsuoka, S. NVCR: A transparent checkpoint-restart library for NVIDIA CUDA.  In 25th IEEE International Symposium on Parallel and Distributed Processing, IPDPS 2011, Anchorage, Alaska, USA, 16-20 May 2011 \- Workshop Proceedings (2011), IEEE, pp. 104–113.
  * [49] NVIDIA.  Basic linear algebra on nvidia gpus.  https://developer.nvidia.com/cublas, 2024.
  * [50] NVIDIA.  Cuda c++ programming guide.  https://docs.nvidia.com/cuda/cuda-c-programming-guide/, 2024.
  * [51] NVIDIA.  Cuda toolkit 12.3 downloads.  https://developer.nvidia.com/cuda-12-3-0-download-archive, 2024\.
  * [52] NVIDIA.  Nvidia/cuda-checkpoint.  https://github.com/NVIDIA/cuda-checkpoint, 2024.
  * [53] NVIDIA.  Parallel thread execution isa version 8.4.  https://docs.nvidia.com/cuda/parallel-thread-execution/index.html, 2024\.
  * [54] OpenAI.  Chatgpt.  https://chat.openai.com, 2024.
  * [55] OpenAI.  Openai gym.  https://github.com/openai/gym, 2024.
  * [56] Pandey, S., Kamath, A. K., and Basu, A. GPM: leveraging persistent memory from a GPU.  In ASPLOS ’22: 27th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Lausanne, Switzerland, 28 February 2022 - 4 March 2022 (2022), B. Falsafi, M. Ferdman, S. Lu, and T. F. Wenisch, Eds., ACM, pp. 142–156.
  * [57] Parasyris, K., Keller, K., Bautista-Gomez, L., and Unsal, O. S. Checkpoint restart support for heterogeneous HPC applications.  In 20th IEEE/ACM International Symposium on Cluster, Cloud and Internet Computing, CCGRID 2020, Melbourne, Australia, May 11-14, 2020 (2020), IEEE, pp. 242–251.
  * [58] Prades, J., and Silla, F. Gpu-job migration: The rcuda case.  IEEE Trans. Parallel Distributed Syst. 30, 12 (2019), 2718–2729.
  * [59] pytorch.  Cuda semantics.  https://pytorch.org/docs/stable/notes/cuda.html#cuda-memory-management, 2024\.
  * [60] Pytorch.  [rfc] upstream torchelastic to pytorch #50621.  https://github.com/pytorch/pytorch/issues/50621, 2024.
  * [61] PyTorch.  Torchelastic.  https://pytorch.org/elastic/latest/, 2024.
  * [62] pytorch.  torchvision.  https://github.com/pytorch/vision, 2024.
  * [63] Qu, H., Mashayekhi, O., Shah, C., and Levis, P. A. Decoupling the control plane from program control flow for flexibility and performance in cloud computing.  In Proceedings of the Thirteenth EuroSys Conference, EuroSys 2018, Porto, Portugal, April 23-26, 2018 (2018), R. Oliveira, P. Felber, and Y. C. Hu, Eds., ACM, pp. 1:1–1:13.
  * [64] Rossbach, C. J., Currey, J., Silberstein, M., Ray, B., and Witchel, E. Ptask: operating system abstractions to manage gpus as compute devices.  In Proceedings of the 23rd ACM Symposium on Operating Systems Principles 2011, SOSP 2011, Cascais, Portugal, October 23-26, 2011 (2011), T. Wobber and P. Druschel, Eds., ACM, pp. 233–248.
  * [65] Shapiro, J. S., Smith, J. M., and Farber, D. J. EROS: a fast capability system.  In Proceedings of the 17th ACM Symposium on Operating System Principles, SOSP 1999, Kiawah Island Resort, near Charleston, South Carolina, USA, December 12-15, 1999 (1999), D. Kotz and J. Wilkes, Eds., ACM, pp. 170–185.
  * [66] Shapiro, J. S., Smith, J. M., and Farber, D. J. EROS: a fast capability system.  In Proceedings of the 17th ACM Symposium on Operating System Principles, SOSP 1999, Kiawah Island Resort, near Charleston, South Carolina, USA, December 12-15, 1999 (1999), D. Kotz and J. Wilkes, Eds., ACM, pp. 170–185.
  * [67] Shukla, D., Sivathanu, M., Viswanatha, S., Gulavani, B. S., Nehme, R., Agrawal, A., Chen, C., Kwatra, N., Ramjee, R., Sharma, P., Katiyar, A., Modi, V., Sharma, V., Singh, A., Singhal, S., Welankar, K., Xun, L., Anupindi, R., Elangovan, K., Rahman, H., Lin, Z., Seetharaman, R., Xu, C., Ailijiang, E., Krishnappa, S., and Russinovich, M. Singularity: Planet-scale, preemptive and elastic scheduling of AI workloads.  CoRR abs/2202.07848 (2022).
  * [68] Silver, D., Huang, A., Maddison, C. J., Guez, A., Sifre, L., van den Driessche, G., Schrittwieser, J., Antonoglou, I., Panneershelvam, V., Lanctot, M., Dieleman, S., Grewe, D., Nham, J., Kalchbrenner, N., Sutskever, I., Lillicrap, T. P., Leach, M., Kavukcuoglu, K., Graepel, T., and Hassabis, D. Mastering the game of go with deep neural networks and tree search.  Nat. 529, 7587 (2016), 484–489.
  * [69] Touvron, H., Martin, L., Stone, K., Albert, P., Almahairi, A., Babaei, Y., Bashlykov, N., Batra, S., Bhargava, P., Bhosale, S., Bikel, D., Blecher, L., Canton-Ferrer, C., Chen, M., Cucurull, G., Esiobu, D., Fernandes, J., Fu, J., Fu, W., Fuller, B., Gao, C., Goswami, V., Goyal, N., Hartshorn, A., Hosseini, S., Hou, R., Inan, H., Kardas, M., Kerkez, V., Khabsa, M., Kloumann, I., Korenev, A., Koura, P. S., Lachaux, M., Lavril, T., Lee, J., Liskovich, D., Lu, Y., Mao, Y., Martinet, X., Mihaylov, T., Mishra, P., Molybog, I., Nie, Y., Poulton, A., Reizenstein, J., Rungta, R., Saladi, K., Schelten, A., Silva, R., Smith, E. M., Subramanian, R., Tan, X. E., Tang, B., Taylor, R., Williams, A., Kuan, J. X., Xu, P., Yan, Z., Zarov, I., Zhang, Y., Fan, A., Kambadur, M., Narang, S., Rodriguez, A., Stojnic, R., Edunov, S., and Scialom, T. Llama 2: Open foundation and fine-tuned chat models.  CoRR abs/2307.09288 (2023).
  * [70] Tsalapatis, E., Hancock, R., Barnes, T., and Mashtizadeh, A. J. The aurora single level store operating system.  In SOSP ’21: ACM SIGOPS 28th Symposium on Operating Systems Principles, Virtual Event / Koblenz, Germany, October 26-29, 2021 (2021), R. van Renesse and N. Zeldovich, Eds., ACM, pp. 788–803.
  * [71] Tsalapatis, E., Hancock, R., Barnes, T., and Mashtizadeh, A. J. The aurora single level store operating system.  In Proceedings of the ACM SIGOPS 28th Symposium on Operating Systems Principles (New York, NY, USA, 2021), SOSP ’21, Association for Computing Machinery, p. 788–803.
  * [72] Ustiugov, D., Petrov, P., Kogias, M., Bugnion, E., and Grot, B. Benchmarking, analysis, and optimization of serverless function snapshots.  In Proceedings of the 26th ACM International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS’21) (2021), ACM.
  * [73] Villa, O., Stephenson, M., Nellans, D. W., and Keckler, S. W. Nvbit: A dynamic binary instrumentation framework for NVIDIA gpus.  In Proceedings of the 52nd Annual IEEE/ACM International Symposium on Microarchitecture, MICRO 2019, Columbus, OH, USA, October 12-16, 2019 (2019), ACM, pp. 372–383.
  * [74] Wang, T., Chen, Z., Wei, X., Gu, J., Chen, R., and Chen, H. Characterizing network requirements for GPU API remoting in AI applications.  CoRR abs/2401.13354 (2024).
  * [75] Wang, Z., Jia, Z., Zheng, S., Zhang, Z., Fu, X., Ng, T. S. E., and Wang, Y. GEMINI: fast failure recovery in distributed training with in-memory checkpoints.  In Proceedings of the 29th Symposium on Operating Systems Principles, SOSP 2023, Koblenz, Germany, October 23-26, 2023 (2023), J. Flinn, M. I. Seltzer, P. Druschel, A. Kaufmann, and J. Mace, Eds., ACM, pp. 364–381.
  * [76] Wei, X., Lu, F., Wang, T., Gu, J., Yang, Y., Chen, R., and Chen, H. No provisioned concurrency: Fast rdma-codesigned remote fork for serverless computing.  In 17th USENIX Symposium on Operating Systems Design and Implementation, OSDI 2023, Boston, MA, USA, July 10-12, 2023 (2023), R. Geambasu and E. Nightingale, Eds., USENIX Association, pp. 497–517.
  * [77] Wu, F., Dong, M., Mo, G., and Chen, H. Treesls: A whole-system persistent microkernel with tree-structured state checkpoint on NVM.  In Proceedings of the 29th Symposium on Operating Systems Principles, SOSP 2023, Koblenz, Germany, October 23-26, 2023 (2023), J. Flinn, M. I. Seltzer, P. Druschel, A. Kaufmann, and J. Mace, Eds., ACM, pp. 1–16.
  * [78] Xiao, W., Bhardwaj, R., Ramjee, R., Sivathanu, M., Kwatra, N., Han, Z., Patel, P., Peng, X., Zhao, H., Zhang, Q., Yang, F., and Zhou, L. Gandiva: Introspective cluster scheduling for deep learning.  In 13th USENIX Symposium on Operating Systems Design and Implementation, OSDI 2018, Carlsbad, CA, USA, October 8-10, 2018 (2018), A. C. Arpaci-Dusseau and G. Voelker, Eds., USENIX Association, pp. 595–610.
  * [79] Yang, Z., Ye, Z., Fu, T., Luo, J., Wei, X., Luo, Y., Wang, X., Wang, Z., and Zhang, T. Tear up the bubble boom: Lessons learned from a deep learning research and development cluster.  In IEEE 40th International Conference on Computer Design, ICCD 2022, Olympic Valley, CA, USA, October 23-26, 2022 (2022), IEEE, pp. 672–680.
  * [80] Zaharia, M., Chowdhury, M., Das, T., Dave, A., Ma, J., McCauly, M., Franklin, M. J., Shenker, S., and Stoica, I. Resilient distributed datasets: A fault-tolerant abstraction for in-memory cluster computing.  In Proceedings of the 9th USENIX Symposium on Networked Systems Design and Implementation, NSDI 2012, San Jose, CA, USA, April 25-27, 2012 (2012), S. D. Gribble and D. Katabi, Eds., USENIX Association, pp. 15–28.
  * [81] Zhuang, S., Wang, S., Liang, E., Cheng, Y., and Stoica, I. Exoflow: A universal workflow system for exactly-once dags.  In 17th USENIX Symposium on Operating Systems Design and Implementation, OSDI 2023, Boston, MA, USA, July 10-12, 2023 (2023), R. Geambasu and E. Nightingale, Eds., USENIX Association, pp. 269–286.
  * [82] Zu, Y., Ghaffarkhah, A., Dang, H., Towles, B., Hand, S., Huda, S., Bello, A., Kolbasov, A., Rezaei, A., Du, D., Lacy, S., Wang, H., Wisner, A., Lewis, C., and Bahini, H. Resiliency at scale: Managing google’s tpuv4 machine learning supercomputer.  In 21st USENIX Symposium on Networked Systems Design and Implementation, NSDI 2024, Santa Clara, CA, April 15-17, 2024 (2024), L. Vanbever and I. Zhang, Eds., USENIX Association, pp. 761–774.