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.

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.

How to checkpoint a CUDA process

Here’s a small CUDA program:

__device__ int counter = 100;
__global__ void increment() { counter++; }

int main(void) {
    cudaFree(0);                                   // force context creation
    int sock = socket(PF_INET, SOCK_DGRAM, IPPROTO_UDP);
    sockaddr_in addr = {AF_INET, htons(10000), inet_addr("127.0.0.1")};
    bind(sock, (sockaddr *)&addr, sizeof addr);

    while (true) {
        char buffer[16] = {0};
        sockaddr_in peer = {0}; socklen_t n = sizeof peer;
        recvfrom(sock, buffer, sizeof buffer, 0, (sockaddr *)&peer, &n);

        increment<<<1,1>>>();                       // one thread, counter++
        int h = 0;
        cudaMemcpyFromSymbol(&h, counter, sizeof counter);

        size_t bytes = sprintf(buffer, "%d\n", h);
        sendto(sock, buffer, bytes, 0, (sockaddr *)&peer, n);
    }
}

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.

$ ./counter &
$ echo -n ping | nc -u -w1 127.0.0.1 10000 # send the packet, and view the response
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:

$ P=$(pgrep -xn counter)
$ cuda-checkpoint --action checkpoint --pid $P     # (lock first; see below)
$ cuda-checkpoint --action restore   --pid $P
$ echo -n ping | nc -u -w1 127.0.0.1 10000
102

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.

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.

cuda-checkpoint --action lock       --pid $P
cuda-checkpoint --action checkpoint --pid $P

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:

$ strace -f -p $P -e trace=mmap
...
mmap(NULL, 417739792, PROT_READ|PROT_WRITE,
     MAP_PRIVATE|MAP_ANONYMOUS|MAP_POPULATE, -1, 0) = 0x...

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:

$ cuda-checkpoint --action restore --pid $P
$ cuda-checkpoint --action unlock  --pid $P
$ echo -n ping | nc -u -w1 127.0.0.1 10000
424243

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:

$ strace -f -e trace=openat,read,write cuda-checkpoint --action checkpoint --pid $P
...
openat(AT_FDCWD, "/proc/$P/task", O_RDONLY|O_DIRECTORY) = 35
openat(AT_FDCWD, "/proc/$P/task/2863110/comm", O_RDONLY) = 36
read(36, "cuda00001400006\n", 1024)          = 16    # the CUDA service thread
openat(AT_FDCWD, "/proc/$P/fd/6",  O_RDONLY) = 35    # its reply pipe
openat(AT_FDCWD, "/proc/$P/fd/5",  O_WRONLY) = 36    # its command pipe
write(36, "\5\0\0\0", 4)                     = 4     # "where do I talk to you?"
read(35, "-\0\0\0\20\0\0\0", 8)              = 8     # -> use fds 45 and 16
openat(AT_FDCWD, "/proc/$P/fd/45", O_RDONLY) = 37
openat(AT_FDCWD, "/proc/$P/fd/16", O_WRONLY) = 38
write(38, "\6\0\0\0\0\0\0\0"..., 2064)       = 2064  # handshake
read(37, "\0\0\0\0\1\0\0\0", 8)              = 8
write(38, "\2\0\0\0\1\0\0\0"..., 2064)       = 2064  # opcode 2, action 1: checkpoint
read(37, "\0\0\0\0\2\0\0\0", 8)              = 8     # status
...

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:

$ sudo cat /proc/$P/task/<tid>/stack
[<0>] do_poll.constprop.0+0x315/0x3c0
[<0>] do_sys_poll+0x1ef/0x290
[<0>] __x64_sys_poll+0x4e/0x150

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:

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.

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

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.

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..

$ echo always | sudo tee /sys/kernel/mm/transparent_hugepage/enabled

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

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.

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

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.

// Speak cuda-checkpoint's pipe protocol directly, no libcuda needed.
// usage: ckpt_ctl <pid> <action>...    (lock | checkpoint | restore | unlock)
#define _GNU_SOURCE
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <stdint.h>
#include <fcntl.h>
#include <unistd.h>
#include <poll.h>
#include <sys/ioctl.h>
#include <dirent.h>

static int openfd(int pid, int fd, int flags) {
    char p[64];
    snprintf(p, sizeof p, "/proc/%d/fd/%d", pid, fd);
    return open(p, flags);
}

// The channel is a loopback pipe: the target reads commands and writes replies
// on the same pipe, so never read until the whole reply has been queued.
static void reply(int fd, uint32_t status[2]) {
    int avail = 0;
    while (ioctl(fd, FIONREAD, &avail), avail < 8)
        poll(&(struct pollfd){fd, POLLIN, 0}, 1, 1);
    read(fd, status, 8);
}

int main(int argc, char **argv) {
    int pid = atoi(argv[1]);

    // find the target's pipe fds
    char path[64];
    snprintf(path, sizeof path, "/proc/%d/fd", pid);
    DIR *dir = opendir(path);
    int pipes[64], np = 0;
    for (struct dirent *de; (de = readdir(dir)) && np < 64;) {
        char lnk[64], tgt[64] = {0};
        snprintf(lnk, sizeof lnk, "/proc/%d/fd/%s", pid, de->d_name);
        if (readlink(lnk, tgt, sizeof tgt) > 4 && !strncmp(tgt, "pipe:", 5))
            pipes[np++] = atoi(de->d_name);
    }

    // rendezvous: opcode 5 into the lowest pipe; some pipe answers with
    // 8 bytes naming the command channel's fds
    uint32_t op5 = 5, chan[2];
    write(openfd(pid, pipes[0], O_WRONLY), &op5, 4);
    for (int i = 0, avail = 0; ; i = (i + 1) % np) {
        int p = openfd(pid, pipes[i], O_RDONLY);
        if (ioctl(p, FIONREAD, &avail), avail >= 8) { read(p, chan, 8); break; }
        close(p);
        usleep(1000);
    }
    int rep = openfd(pid, chan[0], O_RDONLY);
    int cmd = openfd(pid, chan[1], O_WRONLY);

    // each action: 2064-byte handshake (word0=6), then the command itself
    // (word0=2, word1=action), each acknowledged by an 8-byte status
    const char *names[] = {"lock", "checkpoint", "restore", "unlock"};
    for (int i = 2; i < argc; i++) {
        uint32_t buf[516] = {6}, status[2];
        write(cmd, buf, sizeof buf);
        reply(rep, status);
        memset(buf, 0, sizeof buf);
        buf[0] = 2;
        for (uint32_t a = 0; a < 4; a++) if (!strcmp(argv[i], names[a])) buf[1] = a;
        write(cmd, buf, sizeof buf);
        reply(rep, status);
        printf("%s: status=%u state=%u\n", argv[i], status[0], status[1]);
        if (status[0]) return 1;
    }
}