Reverse-engineering NVIDIA's cuda-checkpoint for faster cold starts
There’s a little known feature in the closed-source NVIDIA driver that lets you
freeze a running CUDA process, serialize its GPU state into host memory, and
later restore it to the GPU exactly as it was. We used it in an earlier post to
speed up SGLang server
startup by up to 70x.
[SGLang server
startup](/fast-sglang-starts)
The utility is called cuda-checkpoint. The feature is documented, but how it
works isn’t. One very frustrating aspect, that dogs anyone trying to use it to
checkpoint complex GPU processes, is that the checkpoint transfers come nowhere
close to saturating PCIe bandwidth. We left off our investigation in the
earlier post without a good answer for why that was the caseIn the end, we just used cooperation from the application side to work
around it..
With some tooling from our last
post, we
can find out why it costs so much, and how to make it faster without modifying
the application, or the driver.
[last
post](/what-happens-when-you-run-a-cuda-kernel)
How to checkpoint a CUDA process
How to checkpoint a CUDA process
Here’s a small CUDA program:
It binds a UDP socket, and every time a packet arrives it launches a
one-thread kernel that increments a __device__ int, reads it back, and
replies with the valueThis is NVIDIA's demo, shipped alongside the
cuda-checkpoint tool. I've
trimmed the error handling. Same 4090 and driver 590.48.01 as the
kernel-launch post.. The counter lives in GPU memory and starts at 100.
Ping it, and it says 101.
We can freeze this process — copy its GPU state out, tear its CUDA context down
to nothing, remove it from the GPU entirely — and then, some time later, bring
it back exactly where it was:
In between those two commands the process holds no GPU memory, has no CUDA
context, and does not appear in nvidia-smi. The counter, which lived only on
the device, survives anyway. This is the mechanism that a previous
post leaned on to restore a 122B-parameter server in
a few seconds — there, cuda-checkpoint was a black box called by CRIU. This
post is about how we can find out what’s inside the box.
[previous
post](/fast-sglang-starts)
Watching the process disappear
Watching the process disappear
Let’s watch the process disappear from the device. cuda-checkpoint drives a
small state machine over the target process; --action lock moves it from
running to locked, and then --action checkpoint moves it from locked to
checkpointed.
Watching the process across the two calls, with nvidia-smi, its
/proc/$P/maps, its open file descriptors, and the RssAnon line of
/proc/$P/status:
lock doesn’t change anything observable. But checkpoint does: every mapping
of a /dev/nvidia* file is gone, every file descriptor pointing at the driver
is closed, and the process vanishes from nvidia-smi. As far as the kernel
driver is concerned, this process is no longer using a GPU.
The resident anonymous memory jumped by 407,952 kB at the same moment, as the
GPU state moves into ordinary host memory, into the process’s own address
space. Can we find it?
Poking the checkpoint
The jump in RssAnon is almost exactly the size of one new anonymous mapping
that appears in /proc/$P/maps at the checkpoint. If we attach strace to the
target across the checkpoint we can
catch it being allocated:
417,739,792 bytes is about 398 MiB, against the 388 MiB of device memory
nvidia-smi had attributed to the process. So the inference is the counter’s
device footprint has been serialized into this buffer, plus about
ten megabytes of something else. MAP_POPULATE asks the kernel to fault
the whole thing in immediately rather than lazily, which matters later.
What’s in there? We know what the increment kernel compiles to — cuobjdump -sass counter gives us its SASS — so we can prove that it’s in the mapping by
taking the first few instruction words as a needle and searching the anonymous
mapping for them. We can also find the counter. Nearby, in a page that is
otherwise entirely zeros, is a single non-zero integer holding 0x67 — 103,
because I’d pinged it a couple of times before checkpointing.
With a few more of these kinds of tricks, we can see that the checkpoint buffer
structure is pretty simple:
The counter demo's checkpoint image, mapped by classifying every 4 KiB
page and locating the driver's GPU-mapped surfaces inside it. The driver state
is the constant ~10 MiB diff between the image and the nvidia-smi footprint.
It seems to be GPU-mapped host memory (the increment() SASS is in there, as
are kernel-launch parameter banks, and the channels' notifier pages).
What’s more, the checkpoint image is plain anonymous memory in a process we
own. Let’s mess with it. Rather than increment the counter through the GPU, we
can reach into the frozen image and tweak it there. /proc/$P/mem lets us
write the process’s memory directlyNeeds CAP_SYS_PTRACE or ptrace_scope=0., so we seek to the offset where we
found the counter and write 424242. Then:
The restore uploads the number back onto the GPU, the next packet runs
counter++ on it, and the process replies 424243, incrementing our injected
value.
So we know that the host-side anonymous buffer is the device memory across a
checkpoint. But how did that memory get there?
Who does the work
cuda-checkpoint, the process we invoked, cannot read the target’s device
memory. The tooling that does lives inside the target process and worse, inside
the closed-source userspace driver.
If you strace the utility, essentially everything it does to the target process is
this:
In words, it walks /proc/$P/fd/, finds a pipe that libcuda opened inside
the target process when the CUDA context was first created, and writes a
command word into it. The action lives in a single word, in exactly the order
the CLI lists them:
On the other end of that pipe, inside the target, a thread named
cuda00001400006cuda-checkpoint --get-restore-tid returns a tid that matches this thread. has been sitting in poll() since the process started.
Its kernel stack, the entire time the process is running, shows it waiting for
work:
That thread does the entire checkpoint, from inside the process that is being
checkpointed.
What the driver sees
If the service thread is doing the checkpoint, then it must be making driver
calls. So let’s try to watch those. An LD_PRELOAD shim on the target that decodes
each ioctl’s command number and parameter structThe shim wraps ioctl, matches NVIDIA's 'F' magic, and decodes the
NVOS54 (RM_CONTROL) and NVOS21/NVOS64 (RM_ALLOC) parameter structs
against the open kernel
modules. The
appendix has more details. gives us a per-phase
histogram:
[open kernel
modules](https://github.com/NVIDIA/open-gpu-kernel-modules)
There is no checkpoint ioctl. All the commands that the driver sees when a
checkpoint is performed are ordinary resource-manager ioctls, the same ones
that libcuda uses to create and destroy contexts, allocate and free memory,
and so on. It’s Checkpoint and Restore in
Userspace, but for GPU contexts.
[Checkpoint and Restore in
Userspace](https://criu.org/Main_Page)
Coming back
At checkpoint the context was torn down completely, so at restore time we start
from host buffer and a process with no GPU context.
Decoding the classes restore passes to NV_ESC_RM_ALLOC makes it clear that
the result is basically just context creation (see the previous post for some
of the
details),
run again from scratch, followed by refilling the fresh allocations from the
host image. The memory re-allocations track the anatomy above: restore
walks the buffer front to back, re-creating each allocation in the order it was
serialized.
Forcing cuda-checkpoint to be faster
Forcing cuda-checkpoint to be faster
So that’s the mechanism. When we first were doing checkpoint and restore
experiments to build out our infrastructure for fast starting inference
engines, this cuda-checkpoint blob frustrated me
in its opacity: 3s of driver time, that we couldn’t explain or explore. But
now we know what it’s doing, maybe we can force it to be fast.
[infrastructure for fast starting inference
engines](/fast-sglang-starts)
Here’s the benchmark: one process, with 8578 MiB of allocations, run through a
full checkpoint cycle:
Locking and unlocking are doing nothing here, because the process doesn’t have
anything running when it’s being lockedIn general, it looks like lock is the quiesce barrier the checkpoint
needs to see a still GPU. You can see it by running it against a process
running a kernel that spins for a known duration, launched without a sync,
lock blocks for exactly the remaining kernel time. While the process is
locked, new launches queue in libcuda and only land after unlock.. But they still take 200ms. We
could skip them, because we know that the process is quiesced. But that’s not
really sound in general. The reason they take so long is because they’re
independent invocations of cuda-checkpoint, which means they spin up contexts
and open all the device files, etc. But we know they’re actually only writing a
single command word into a pipe.
So let’s speak their language directlyWe do this with a little client that speaks the pipe protocol.
Realistically, we're probably just rebuilding what the C API does., and skip them:
With the utility out of the way, the remaining time is all driver. Swept
across device footprint with a program that just allocates N MiB and idles:
checkpoint and restore are both linear in the amount of device memory,
which makes sense since they’re copying it. They run at very different rates
though: checkpoint at about 3 GiB/s, restore at about 8. Restore, which has all
that context to rebuild, is well over twice as fast as checkpoint, which mostly
just copies memory out. Neither rate is anywhere near the PCIe ceiling: this
link moves pageable host memory at 20.6 GiB/s host-to-device and
17.4 GiB/s device-to-host, and pinned memory at about 25 either wayPCIe 4.0 x16 here, in principle 31.5 GB/s per direction..
It turns out almost all the time is going into allocating and deallocating the
staging buffer. Timing the individual syscalls inside the 8,578 MiB checkpoint,
the mmap(..., MAP_POPULATE) of the staging region alone accounts for 2.08 s
of the 2.8 s checkpoint, because MAP_POPULATE faults in every page up front
and the kernel has to zero each fresh page before handing it over. On restore,
the mirror operation is the munmap that frees those pages (at 435 ms) and
freeing is much cheaper than zeroing.
So the real story looks like this:
How can we do better?
Without going too crazy, we can just turn on Transparent Huge
Pages (THP). THP lets
the kernel automatically choose to back the anonymous mapping with 2 MiB pages
instead of 4 KiB ones, and zeroing a 2 MiB page is enormously cheaper per byte
than doing it 512 timesA small victory for humanity here, Claude can't launch processes with THP
enabled. The problem is in bun's mimalloc build, which sets
prctl(PR_SET_THP_DISABLE) for some kind of internal allocation reason, but
then (presumably accidentally) the result percolates out into all of Claude
code's subprocesses..
[Transparent Huge
Pages](https://docs.kernel.org/admin-guide/mm/transhuge.html)
With huge pages, the registration and unregistration costs collapse:
What if we don’t have or want huge pages? or what if we really want to get rid
of the allocation completely?
cuda-checkpoint allocates the staging buffer inside closed libcuda, so we
can’t work with it nicely. But, the allocation is an ordinary mmap through
libc, and it’s the only large MAP_ANONYMOUS | MAP_POPULATE mapping the
process makes, so an LD_PRELOAD shim can recognise it, and we can replace the
result with whatever we want.
If we know we’re about to perform a checkpoint, we can create a faulted, zeroed
buffer, wait for the mmap, and give it to cuda-checkpoint immediately. And if
we know our process is going to keep running after restore, we can unmap the
buffer later, when we feel like it. Putting everything together:
Checkpoint ends at 540 ms, at about 15.5 GiB/s — the pageable PCIe copy.
Restore’s floor is a little higher (12.7 GiB/s effective); it still has a
context to rebuild. The context is of fixed size, so it amortizes over larger
transfers.
Can we hit the PCIe ceiling? Not by handing libcuda a pinned buffer: when we
do, it doesn’t get used in pinned form. We could try to patch the binary, or
keep pushing on new hooks. But at some point, you’ve got to choose to be
done, and this is a reasonable place: by poking and prodding at how
cuda-checkpoint does its work, we’ve sped up the checkpoint/restore cycle
4x.
Can we productionise this stuff? THP is a kernel setting, so that’s easy
enough. cuCheckpointProcess plays our pipe tricks in a supported C API, so
you can pay the 200ms context setup cost once at startup instead of for each
verb. In exchange you get stability guarantees; you don’t so much when, like we
do here, you’re writing bytes to a pipe. The staging-buffer swap is more
fragile: it depends on catching an mmap that closed libcuda happens to make.
But before, this was NVIDIA’s playground. Now we can play too.
Appendix: the diagnostic hooks
Appendix: the diagnostic hooks
Almost everything here was read off two LD_PRELOAD shims and /proc, because
libcuda is closed and the interesting work happens inside a process we can
watch but not read the source of. As in the kernel-launch
post, the trick is to interpose
on the ordinary libc calls the driver makes and decode them against the open
kernel modules.
[kernel-launch
post](/what-happens-when-you-run-a-cuda-kernel)
Watching the driver’s ioctls
The per-phase histogram comes from wrapping ioctl, filtering for NVIDIA’s
'F' magic, and decoding the command number and parameter struct. The command
numbers live in
nv_escape.h;
the RM_CONTROL and RM_ALLOC structs (NVOS54, NVOS21) in nvos.h; and the
allocation class numbers in src/common/sdk/nvidia/inc/class/. To split the
histogram by phase, mark the shim’s log at each cuda-checkpoint invocation and
diff the offsets.
Finding the staging buffer, and the counter in it
Finding the staging buffer, and the counter in it
To catch the allocation, wrap mmap and log any large MAP_ANONYMOUS mapping.
The staging buffer is the only big one with MAP_POPULATE set. There’s
probably something more selective here. To find the contents afterwards, scan
the process’s anonymous VMAs (from /proc/$P/maps) for a needle: the first
four 128-bit instruction words of the increment kernel works. The device
global turns up as an isolated non-zero integer in a page of zeros near the
SASS. Reading and writing the image itself is pread/pwrite on
/proc/$P/mem.
The control protocol
To see the command words, wrap write in the utility rather than the target
and dump any write to a pipe fd: you’ll get the opcode-5 rendezvous, the
handshake, and the 2064-byte command struct whose first two words are the opcode
and the action index. It’s almost entirely zeros; the only non-zero fields for a
plain checkpoint are word0 = 2 and word1 = action. Passing
--device-map <uuid>=<uuid> to a restore fills a few more words — a mapping
count and inline source/destination GPU UUIDs — which is the machinery for
restoring a checkpoint onto a different physical GPU than it was taken on.
A minimal client
Putting the protocol together: the client below does the rendezvous, then
drives the state machine one 2064-byte command at a time. It never links
libcuda, so there is no init to pay; even invoked cold, once per action,
lock lands in a few milliseconds. Error handling trimmed.
Doubleword is an inference provider delivering highly efficient inference for open-source and custom models. By optimizing the full stack for throughput—from hardware to inference engine—we offer some of the lowest token costs available on the market. That means abundant, affordable tokens for background agents, data processing, batch inference, and embeddings.
Start building with free credits→
Interested in helping us make inference 100x more efficient? We’re hiring—reach us at careers@doubleword.ai.