# Reverse-engineering NVIDIA's cuda-checkpoint for faster cold starts > Source: > Published: 2026-07-09 00:00:00+00:00 # 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](/fast-sglang-starts) 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.. Now we do: why it costs so much, and how to make it faster, without modifying the application, or the driver. [How to checkpoint a CUDA process](#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](https://github.com/NVIDIA/cuda-checkpoint) tool. I've trimmed the error handling. Same 4090 and driver 590.48.01 as the [kernel-launch post](/what-happens-when-you-run-a-cuda-kernel).. The counter lives in GPU memory and starts at 100. Ping it, and it says 101. ``` bash $ ./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: ``` bash $ 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](/fast-sglang-starts) 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](#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` : | running | locked | checkpointed | | |---|---|---|---| `RssAnon` | 12,860 kB | 12,860 kB | 420,812 kB | NVIDIA VMAs in `/proc/$P/maps` | 26 | 26 | 0 | NVIDIA fds in `/proc/$P/fd/` | 24 | 24 | 0 | visible in `nvidia-smi` | yes | yes | no | `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](#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: ``` bash $ 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: ``` bash $ 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](#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: ``` bash $ 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: `word1` | action | |---|---| | 0 | `lock` | | 1 | `checkpoint` | | 2 | `restore` | | 3 | `unlock` | On the other end of that pipe, inside the target, a thread named `cuda00001400006` `cuda-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: ``` bash $ sudo cat /proc/$P/task//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](#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](https://github.com/NVIDIA/open-gpu-kernel-modules). The [appendix](#appendix-the-diagnostic-hooks) has more details. gives us a per-phase histogram: | phase | RM_CONTROL | RM_ALLOC | other `NV_ESC_*` | UVM | total ioctls | |---|---|---|---|---|---| `lock` | — | — | — | — | 2 | `checkpoint` | 133 | 25 | 97 | 56 | 793 | `restore` | 213 | 124 | 117 | 108 | 1197 | `unlock` | — | — | — | — | 2 | 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](https://criu.org/Main_Page), but for GPU contexts. [Coming back](#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](/what-happens-when-you-run-a-cuda-kernel)), 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 ](#forcing-cuda-checkpoint-to-be-faster)`cuda-checkpoint` to be faster `cuda-checkpoint` to be fasterSo that’s the mechanism. When we first were doing checkpoint and restore experiments to build out our [infrastructure for fast starting inference engines](/fast-sglang-starts), 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: | GPU MiB | lock | checkpoint | restore | unlock | |---|---|---|---|---| | 450 | 0.3 ms | 231 ms | 197 ms | 1 ms | | 642 | 0.3 ms | 287 ms | 203 ms | 1 ms | | 1410 | 0.4 ms | 517 ms | 284 ms | 2 ms | | 4482 | 0.3 ms | 1407 ms | 571 ms | 4 ms | | 8578 | 0.3 ms | 2749 ms | 1056 ms | 4 ms | `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](https://docs.kernel.org/admin-guide/mm/transhuge.html) (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.. ``` bash $ 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 do some kind of `ioctl` hook. 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 you productionise any of this stuff? THP is a kernel setting, so that’s easy enough. `cuCheckpointProcess` in the C API 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; not so much when, like we do here, you’re writing bytes to a pipe. But our client doesn’t need a CUDA context at all, which is pretty nice. The staging-buffer swap is more fragile: it depends on catching an mmap that closed libcuda happens to make, and there are no ABI guarantees of any kind there. 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](/what-happens-when-you-run-a-cuda-kernel), 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](#watching-the-drivers-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](https://github.com/NVIDIA/open-gpu-kernel-modules/blob/590.48.01/src/nvidia/arch/nvalloc/unix/include/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](#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 =` 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](#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 ... (lock | checkpoint | restore | unlock) #define _GNU_SOURCE #include #include #include #include #include #include #include #include #include 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; } } @misc{doubleword-what-happens-when-you-checkpoint-a-cuda-process, title = {Reverse-engineering NVIDIA's cuda-checkpoint for faster cold starts}, author = {Fergus Finn}, year = {2026}, howpublished = {Doubleword Blog}, url = {https://blog.doubleword.ai/what-happens-when-you-checkpoint-a-cuda-process}, } ```