diff --git a/RUNBOOK.md b/RUNBOOK.md
index 67d8ecdc2a..2b85fbbdce 100644
--- a/RUNBOOK.md
+++ b/RUNBOOK.md
@@ -168,6 +168,10 @@ The server logs `heat_wait_s`, a cancel ends the wait, and the reply's tokens do
 On three GB10s serving GLM-5.3-Flash NVFP4 at TP=3, a 1,035,295-token prompt without the bands kept one host's hottest
 zone at 94 C or more for most of the prompt (97.7 C at the top); `TF_HEAT_HIGH=92` and `TF_HEAT_LOW=88` held all three
 at 93.0 C or under, and the first token came 10% later.
+`TF_GRAPH_CACHE=0` (or `off`, any case) runs every step eagerly in any family that installs the keyed CUDA graph cache,
+without a rebuild. Set it alike on every tensor-parallel rank: an off rank never joins the cache's per-miss vote, so a
+rank left on would wait at its first miss, and only the off rank logs the switch. Ranks that share a cache compare
+`Settings.digest` of their effective settings at start, before building it, and refuse to start when they differ.
 Archives whose names contain `host-only` are CPU verification artifacts and cannot serve CUDA inference.
 
 ## Flash Next and paired Metal serving
diff --git a/zig/src/cuda/abi.zig b/zig/src/cuda/abi.zig
index b784479964..fe21a66fc4 100644
--- a/zig/src/cuda/abi.zig
+++ b/zig/src/cuda/abi.zig
@@ -39,6 +39,8 @@ pub const DeviceAttribute = enum(c_int) {
     compute_capability_minor = 76,
     max_shared_memory_per_multiprocessor = 81,
     max_shared_memory_per_block_optin = 97,
+    virtual_memory_management_supported = 102,
+    gpu_direct_rdma_with_cuda_vmm_supported = 110,
 };
 
 pub const FunctionAttribute = enum(c_int) {
@@ -242,3 +244,10 @@ comptime {
     std.debug.assert(@sizeOf(KernelNodeParams) == 72 and @offsetOf(KernelNodeParams, "params") == 40);
     std.debug.assert(@sizeOf(ExecUpdateResultInfo) == 24);
 }
+
+test "CUDA device attribute matches the RDMA VMM capability ABI" {
+    // CUDA's CUdevice_attribute assigns 110 to RDMA VMM support; 126 is the memory sync domain count.
+    try @import("std").testing.expectEqual(@as(c_int, 110), @backingInt(
+        DeviceAttribute.gpu_direct_rdma_with_cuda_vmm_supported,
+    ));
+}
diff --git a/zig/src/cuda/arena.zig b/zig/src/cuda/arena.zig
new file mode 100644
index 0000000000..de3da5718c
--- /dev/null
+++ b/zig/src/cuda/arena.zig
@@ -0,0 +1,405 @@
+//! A device sub-allocator under DeviceBuffer.alloc: buffers carved from a few large chunks, not one cuMemAlloc each.
+
+const std = @import("std");
+const abi = @import("abi.zig");
+const Driver = @import("driver.zig").Driver;
+const memory = @import("memory.zig");
+
+pub const Error = error{ DriverUnavailable, MissingSymbol, Unsupported, CudaFailed, OutOfMemory };
+
+/// Where chunks come from: `chunk` maps `len` bytes (a multiple of `granularity`), `release` gives a chunk back.
+pub const Backend = struct {
+    ctx: *anyopaque,
+    granularity: usize,
+    chunk: *const fn (ctx: *anyopaque, len: usize) ?u64,
+    release: *const fn (ctx: *anyopaque, base: u64, len: usize) void,
+    name: []const u8,
+};
+
+pub const Options = struct {
+    small_max: usize = 4 << 20, // requests up to this many bytes share chunks
+    small_chunk: usize = 64 << 20, // a shared chunk's bytes, rounded up to the granularity
+    alignment: usize = 512, // every range starts at a multiple of this
+};
+
+pub const Stats = struct { small: u64 = 0, large: u64 = 0, reused: u64 = 0, fallback: u64 = 0, chunk_bytes: u64 = 0 };
+
+/// Small requests bump 512-byte aligned through shared chunks and are reused by size; large ones get their own chunk.
+pub const Arena = struct {
+    gpa: std.mem.Allocator,
+    be: Backend,
+    o: Options,
+    mutex: std.atomic.Mutex = .unlocked,
+    small: std.ArrayList([2]u64) = .empty, // shared chunks (base, len)
+    cursor: u64 = 0, // the current shared chunk's next free byte
+    limit: u64 = 0,
+    free_small: std.AutoHashMapUnmanaged(u64, std.ArrayList(u64)) = .empty, // freed small ranges by rounded size
+    live: std.AutoHashMapUnmanaged(u64, u64) = .empty, // live small ranges: address to rounded size
+    large: std.AutoHashMapUnmanaged(u64, u64) = .empty, // large chunks: address to mapped size
+    dead: [2]u32 = .{ 0, 0 }, // removals from `live` and `large` since each was last rehashed
+    stats: Stats = .{},
+
+    pub fn init(gpa: std.mem.Allocator, be: Backend, o: Options) Arena {
+        return .{ .gpa = gpa, .be = be, .o = o };
+    }
+
+    pub fn deinit(a: *Arena) void {
+        var it = a.large.iterator();
+        while (it.next()) |e| a.be.release(a.be.ctx, e.key_ptr.*, e.value_ptr.*);
+        for (a.small.items) |c| a.be.release(a.be.ctx, c[0], c[1]);
+        var f = a.free_small.valueIterator();
+        while (f.next()) |l| l.deinit(a.gpa);
+        a.free_small.deinit(a.gpa);
+        a.live.deinit(a.gpa);
+        a.large.deinit(a.gpa);
+        a.small.deinit(a.gpa);
+        a.* = undefined;
+    }
+
+    fn lock(a: *Arena) void {
+        while (!a.mutex.tryLock()) std.Thread.yield() catch {};
+    }
+
+    fn roundUp(x: usize, m: usize) usize {
+        return std.mem.alignForward(usize, x, m);
+    }
+
+    /// An address for `len` bytes, or null (the caller then allocates as before).
+    pub fn alloc(a: *Arena, len: usize) ?u64 {
+        if (len == 0) return null;
+        a.lock();
+        defer a.mutex.unlock();
+        return a.allocLocked(len) catch null;
+    }
+
+    fn allocLocked(a: *Arena, len: usize) !?u64 {
+        if (len > a.o.small_max) {
+            const n = roundUp(len, a.be.granularity);
+            const p = a.be.chunk(a.be.ctx, n) orelse {
+                a.stats.fallback += 1;
+                return null;
+            };
+            errdefer a.be.release(a.be.ctx, p, n);
+            try a.large.put(a.gpa, p, n);
+            a.stats.large += 1;
+            a.stats.chunk_bytes += n;
+            return p;
+        }
+        const n: u64 = roundUp(len, a.o.alignment);
+        try a.live.ensureUnusedCapacity(a.gpa, 1); // before a range is taken: a failure leaves free list and cursor alone
+        if (a.free_small.getPtr(n)) |l| if (l.pop()) |p| {
+            a.live.putAssumeCapacityNoClobber(p, n);
+            a.stats.reused += 1;
+            return p;
+        };
+        if (a.limit - a.cursor < n) {
+            const cl = roundUp(@max(a.o.small_chunk, n), a.be.granularity);
+            const base = a.be.chunk(a.be.ctx, cl) orelse {
+                a.stats.fallback += 1;
+                return null;
+            };
+            errdefer a.be.release(a.be.ctx, base, cl);
+            try a.small.append(a.gpa, .{ base, cl });
+            a.cursor = base;
+            a.limit = base + cl;
+            a.stats.chunk_bytes += cl;
+        }
+        const p = a.cursor;
+        a.live.putAssumeCapacityNoClobber(p, n);
+        a.cursor += n;
+        a.stats.small += 1;
+        return p;
+    }
+
+    /// True when `p` was the arena's (and is now free); false: not ours, the caller frees it as before.
+    pub fn free(a: *Arena, p: u64) bool {
+        a.lock();
+        defer a.mutex.unlock();
+        if (a.large.fetchRemove(p)) |e| {
+            removed(&a.large, &a.dead[1]);
+            a.be.release(a.be.ctx, p, e.value);
+            return true;
+        }
+        const e = a.live.fetchRemove(p) orelse return false;
+        removed(&a.live, &a.dead[0]);
+        const g = a.free_small.getOrPut(a.gpa, e.value) catch return true; // leaked range, still ours
+        if (!g.found_existing) g.value_ptr.* = .empty;
+        g.value_ptr.append(a.gpa, p) catch {};
+        return true;
+    }
+
+    /// std's tombstones never clear and keys fill 80 %: rehash at a tenth, so a tenth of the map always stays free.
+    fn removed(map: *std.AutoHashMapUnmanaged(u64, u64), dead: *u32) void {
+        dead.* += 1;
+        if (@as(u64, dead.*) * 10 < map.capacity()) return;
+        map.rehash(std.hash_map.AutoContext(u64){});
+        dead.* = 0;
+    }
+
+    // -- memory.zig's hook ----------------------------------------------------------------------------------------
+
+    fn hookAlloc(ctx: *anyopaque, len: usize) ?u64 {
+        const a: *Arena = @ptrCast(@alignCast(ctx));
+        return a.alloc(len);
+    }
+
+    fn hookFree(ctx: *anyopaque, p: u64) bool {
+        const a: *Arena = @ptrCast(@alignCast(ctx));
+        return a.free(p);
+    }
+
+    /// Every later DeviceBuffer.alloc comes from the arena, which must outlive its buffers; only addresses change.
+    pub fn install(a: *Arena) void {
+        memory.setHook(.{ .ctx = a, .alloc = hookAlloc, .free = hookFree });
+    }
+};
+
+// -- backends ---------------------------------------------------------------------------------------------------------
+
+/// One cuMemAlloc a chunk (the driver's own 2 MiB-aligned allocations).
+pub const Plain = struct {
+    d: *const Driver,
+
+    pub fn backend(p: *Plain) Backend {
+        return .{ .ctx = p, .granularity = 2 << 20, .chunk = chunk, .release = release, .name = "plain" };
+    }
+
+    fn chunk(ctx: *anyopaque, len: usize) ?u64 {
+        const p: *Plain = @ptrCast(@alignCast(ctx));
+        var q: abi.DevicePtr = 0;
+        if (p.d.api.cuMemAlloc_v2(&q, len) != abi.success) return null;
+        return q;
+    }
+
+    fn release(ctx: *anyopaque, base: u64, _: usize) void {
+        const p: *Plain = @ptrCast(@alignCast(ctx));
+        _ = p.d.api.cuMemFree_v2(base);
+    }
+};
+
+/// CUDA's virtual memory management: one reserved virtual range, chunks of the recommended granularity mapped in order.
+pub const Vmm = struct {
+    lib: std.DynLib,
+    device: abi.Device,
+    prop: Prop,
+    granularity: usize,
+    va: u64 = 0,
+    va_len: usize = 0,
+    next: u64 = 0,
+    api: Api,
+
+    const Api = struct {
+        reserve: *const fn (*u64, usize, usize, u64, c_ulonglong) callconv(.c) abi.Result,
+        addressFree: *const fn (u64, usize) callconv(.c) abi.Result,
+        create: *const fn (*u64, usize, *const Prop, c_ulonglong) callconv(.c) abi.Result,
+        release: *const fn (u64) callconv(.c) abi.Result,
+        map: *const fn (u64, usize, usize, u64, c_ulonglong) callconv(.c) abi.Result,
+        unmap: *const fn (u64, usize) callconv(.c) abi.Result,
+        setAccess: *const fn (u64, usize, *const AccessDesc, usize) callconv(.c) abi.Result,
+        granularity: *const fn (*usize, *const Prop, c_int) callconv(.c) abi.Result,
+    };
+
+    pub const Prop = extern struct { // CUmemAllocationProp
+        type: c_int = 1, // CU_MEM_ALLOCATION_TYPE_PINNED
+        handle_types: c_int = 0,
+        loc_type: c_int = 1, // CU_MEM_LOCATION_TYPE_DEVICE
+        loc_id: c_int = 0,
+        win32: ?*anyopaque = null,
+        compression: u8 = 0,
+        rdma: u8 = 0,
+        usage: u16 = 0,
+        reserved: [4]u8 = .{ 0, 0, 0, 0 },
+    };
+    const AccessDesc = extern struct { loc_type: c_int = 1, loc_id: c_int, flags: c_int = 3 }; // CUmemAccessDesc, read-write
+
+    comptime {
+        std.debug.assert(@sizeOf(Prop) == 32 and @offsetOf(Prop, "compression") == 24 and @sizeOf(AccessDesc) == 12);
+    }
+
+    /// `va_bytes` of virtual addresses reserved (cheap: no memory behind them until a chunk maps).
+    pub fn open(d: *const Driver, device: abi.Device, va_bytes: usize) Error!Vmm {
+        var vmm_ok: c_int = 0;
+        if (d.api.cuDeviceGetAttribute(&vmm_ok, .virtual_memory_management_supported, device) != abi.success or vmm_ok == 0) return error.Unsupported; // VIRTUAL_MEMORY_MANAGEMENT_SUPPORTED
+        var lib = std.DynLib.open("libcuda.so.1") catch return error.DriverUnavailable;
+        errdefer lib.close();
+        var api: Api = undefined;
+        const names = .{ .{ "reserve", "cuMemAddressReserve" }, .{ "addressFree", "cuMemAddressFree" }, .{ "create", "cuMemCreate" }, .{ "release", "cuMemRelease" }, .{ "map", "cuMemMap" }, .{ "unmap", "cuMemUnmap" }, .{ "setAccess", "cuMemSetAccess" }, .{ "granularity", "cuMemGetAllocationGranularity" } };
+        inline for (names) |n| @field(api, n[0]) = lib.lookup(@FieldType(Api, n[0]), n[1]) orelse return error.MissingSymbol;
+        var prop: Prop = .{ .loc_id = device };
+        var rdma: c_int = 0; // GPU_DIRECT_RDMA_WITH_CUDA_VMM_SUPPORTED: keep the memory registrable as cuMemAlloc's is
+        if (d.api.cuDeviceGetAttribute(&rdma, .gpu_direct_rdma_with_cuda_vmm_supported, device) == abi.success and rdma != 0) prop.rdma = 1;
+        var g: usize = 0;
+        if (api.granularity(&g, &prop, 1) != abi.success or g == 0) return error.Unsupported; // RECOMMENDED
+        g = @max(g, 2 << 20);
+        const len = std.mem.alignForward(usize, va_bytes, g);
+        var va: u64 = 0;
+        if (api.reserve(&va, len, g, 0, 0) != abi.success) return error.CudaFailed;
+        return .{ .lib = lib, .device = device, .prop = prop, .granularity = g, .va = va, .va_len = len, .next = va, .api = api };
+    }
+
+    pub fn close(v: *Vmm) void {
+        _ = v.api.addressFree(v.va, v.va_len);
+        v.lib.close();
+    }
+
+    pub fn backend(v: *Vmm) Backend {
+        return .{ .ctx = v, .granularity = v.granularity, .chunk = chunk, .release = release, .name = "vmm" };
+    }
+
+    fn chunk(ctx: *anyopaque, len: usize) ?u64 {
+        const v: *Vmm = @ptrCast(@alignCast(ctx));
+        if (v.va + v.va_len - v.next < len) return null;
+        var h: u64 = 0;
+        if (v.api.create(&h, len, &v.prop, 0) != abi.success) return null;
+        defer _ = v.api.release(h); // the mapping holds the memory until it is unmapped
+        const at = v.next;
+        if (v.api.map(at, len, 0, h, 0) != abi.success) return null;
+        const desc: AccessDesc = .{ .loc_id = v.device };
+        if (v.api.setAccess(at, len, &desc, 1) != abi.success) {
+            _ = v.api.unmap(at, len);
+            return null;
+        }
+        v.next += len;
+        return at;
+    }
+
+    fn release(ctx: *anyopaque, base: u64, len: usize) void {
+        const v: *Vmm = @ptrCast(@alignCast(ctx));
+        _ = v.api.unmap(base, len); // the virtual range is not reused: the reservation is sized for the process
+    }
+};
+
+// -- tests (a fake backend: addresses only) ---------------------------------------------------------------------------
+
+const Fake = struct {
+    next: u64 = 1 << 40,
+    chunks: u32 = 0,
+    released: u32 = 0,
+    fail: bool = false,
+
+    fn backend(f: *Fake, g: usize) Backend {
+        return .{ .ctx = f, .granularity = g, .chunk = chunk, .release = release, .name = "fake" };
+    }
+    fn chunk(ctx: *anyopaque, len: usize) ?u64 {
+        const f: *Fake = @ptrCast(@alignCast(ctx));
+        if (f.fail) return null;
+        const p = f.next;
+        f.next += len;
+        f.chunks += 1;
+        return p;
+    }
+    fn release(ctx: *anyopaque, _: u64, _: usize) void {
+        const f: *Fake = @ptrCast(@alignCast(ctx));
+        f.released += 1;
+    }
+};
+
+test "arena: small requests pack 512-aligned into shared chunks, are reused by size, large ones get their own" {
+    const t = std.testing;
+    var f: Fake = .{};
+    var a = Arena.init(t.allocator, f.backend(2 << 20), .{ .small_max = 1 << 20, .small_chunk = 4 << 20 });
+    defer a.deinit();
+    const p1 = a.alloc(20480).?; // a norm weight
+    const p2 = a.alloc(983040).?; // an mHC fn16
+    const p3 = a.alloc(100).?;
+    try t.expectEqual(p1 + 20480, p2);
+    try t.expectEqual(p2 + 983040, p3);
+    for ([_]u64{ p1, p2, p3 }) |p| try t.expectEqual(@as(u64, 0), p % 512);
+    try t.expectEqual(@as(u32, 1), f.chunks);
+    // a freed range comes back for the same rounded size
+    try t.expect(a.free(p3));
+    try t.expectEqual(p3, a.alloc(17).?);
+    try t.expectEqual(@as(u64, 1), a.stats.reused);
+    // three more 1 MiB ranges fit the 4 MiB chunk, the fourth opens a new shared chunk
+    for (0..3) |_| _ = a.alloc(1 << 20).?;
+    try t.expectEqual(@as(u32, 1), f.chunks);
+    _ = a.alloc(1 << 20).?;
+    try t.expectEqual(@as(u32, 2), f.chunks);
+    // large: its own chunk, rounded to the granularity, released when freed
+    const big = a.alloc(5 << 20).?;
+    try t.expectEqual(@as(u32, 3), f.chunks);
+    try t.expectEqual(@as(u64, 6 << 20), a.large.get(big).?);
+    try t.expect(a.free(big));
+    try t.expectEqual(@as(u32, 1), f.released);
+    // not ours
+    try t.expect(!a.free(12345));
+    // a backend that cannot map: null, the caller allocates as before
+    f.fail = true;
+    try t.expectEqual(@as(?u64, null), a.alloc(64 << 20));
+    try t.expectEqual(@as(u64, 1), a.stats.fallback);
+}
+
+test "arena: under DeviceBuffer.alloc through memory.setHook" {
+    const t = std.testing;
+    var f: Fake = .{};
+    var a = Arena.init(t.allocator, f.backend(2 << 20), .{});
+    defer a.deinit();
+    a.install();
+    defer memory.setHook(null);
+    const h = memory.currentHook().?;
+    const p = h.alloc(h.ctx, 4096).?;
+    try t.expect(h.free(h.ctx, p));
+    try t.expect(!h.free(h.ctx, 7));
+}
+
+/// Slots of `m` that hold neither a key nor a tombstone.
+fn freeSlots(m: *const std.AutoHashMapUnmanaged(u64, u64)) usize {
+    const meta = m.metadata orelse return 0;
+    var n: usize = 0;
+    for (meta[0..m.capacity()]) |x| n += @intFromBool(x.isFree());
+    return n;
+}
+
+test "arena: a long run of allocations and frees leaves both address maps free slots" {
+    const t = std.testing;
+    var f: Fake = .{};
+    var a = Arena.init(t.allocator, f.backend(2 << 20), .{ .small_max = 1 << 20, .small_chunk = 4 << 20 });
+    defer a.deinit();
+    var held: std.ArrayList(u64) = .empty;
+    defer held.deinit(t.allocator);
+    var rng = std.Random.DefaultPrng.init(3);
+    for (0..20000) |_| {
+        // a serving process's buffers come and go (prompt workspaces, grown buffers): small and large alike
+        if (held.items.len > 0 and rng.random().boolean()) {
+            const k = rng.random().uintLessThan(usize, held.items.len);
+            try t.expect(a.free(held.swapRemove(k)));
+        } else {
+            const big = rng.random().uintLessThan(u32, 4) == 0;
+            try held.append(t.allocator, a.alloc(if (big) 3 << 20 else 512 * (1 + rng.random().uintLessThan(u64, 64))).?);
+        }
+        for ([_]*const std.AutoHashMapUnmanaged(u64, u64){ &a.live, &a.large }) |m| {
+            if (m.capacity() == 0) continue;
+            try t.expect(freeSlots(m) * 10 >= m.capacity());
+        }
+    }
+}
+
+test "arena: a reuse whose address-map insertion fails keeps its range free, and the retry gets it" {
+    const t = std.testing;
+    var f: Fake = .{};
+    var fa = t.FailingAllocator.init(t.allocator, .{});
+    {
+        var a = Arena.init(fa.allocator(), f.backend(2 << 20), .{ .small_max = 1 << 20, .small_chunk = 4 << 20 });
+        defer a.deinit(); // testing.allocator underneath reports any leaked host bytes
+        const p = a.alloc(4096).?;
+        try t.expect(a.free(p));
+        // other sizes take back the room the free gave: inserting the reused range must grow the map
+        while (a.live.available > 0) _ = a.alloc(512).?;
+        const live = a.live.count();
+        const cursor = a.cursor;
+        const chunks = f.chunks;
+        fa.fail_index = fa.alloc_index; // the next allocation fails, and every one after it until healed
+        try t.expectEqual(@as(?u64, null), a.alloc(4096));
+        try t.expectEqual(live, a.live.count());
+        try t.expectEqual(cursor, a.cursor);
+        try t.expectEqual(chunks, f.chunks);
+        try t.expectEqualSlices(u64, &.{p}, a.free_small.get(4096).?.items);
+        try t.expectEqual(@as(u64, 0), a.stats.reused);
+        fa.fail_index = std.math.maxInt(usize);
+        try t.expectEqual(p, a.alloc(4096).?);
+        try t.expectEqual(@as(u64, 1), a.stats.reused);
+        try t.expectEqual(@as(usize, 0), a.free_small.get(4096).?.items.len);
+        try t.expect(a.free(p));
+    }
+    try t.expectEqual(f.chunks, f.released); // every chunk went back to the backend
+}
diff --git a/zig/src/cuda/driver.zig b/zig/src/cuda/driver.zig
index 7c7fb7f068..38a9b930d7 100644
--- a/zig/src/cuda/driver.zig
+++ b/zig/src/cuda/driver.zig
@@ -9,6 +9,10 @@ pub const Error = error{ CopySourceLimit, StageBusy, StageUnprepared, UnownedCap
 pub const Driver = struct {
     lib: platform.DynLib,
     api: abi.Api,
+    /// cuFuncGetName (CUDA 12.3+; optional: an older driver opens without it): a failed launch names its kernel
+    func_name: ?FuncGetName = null,
+
+    pub const FuncGetName = *const fn (*?[*:0]const u8, abi.Function) callconv(.c) abi.Result;
 
     pub fn open() Error!Driver {
         return openPath(platform.cuda_driver);
@@ -26,7 +30,7 @@ pub const Driver = struct {
                 return error.MissingSymbol;
             };
         }
-        const d: Driver = .{ .lib = lib, .api = api };
+        const d: Driver = .{ .lib = lib, .api = api, .func_name = lib.lookup(FuncGetName, "cuFuncGetName") };
         try d.check(api.cuInit(0), "cuInit");
         return d;
     }
@@ -43,18 +47,33 @@ pub const Driver = struct {
     pub fn check(self: *const Driver, res: abi.Result, what: []const u8) Error!void {
         if (res == abi.success) return;
         if (res == abi.error_not_ready) return error.NotReady;
-        std.log.err("{s}: {s} ({d}) {s}", .{ what, self.errorName(res), res, self.errorText(res) });
+        std.log.err("{s}: {s} ({d}) {s}{s}", .{ what, self.errorName(res), res, self.errorText(res), shutdownNote(res) });
+        return errorOf(res);
+    }
+
+    /// The error a failed call returns (`check` without its log line, for callers that log more).
+    pub fn errorOf(res: abi.Result) Error {
         return switch (res) {
             2 => error.OutOfDeviceMemory,
             abi.error_not_found => error.NotFound,
+            abi.error_not_ready => error.NotReady,
             else => error.CudaFailed,
         };
     }
 
+    /// The driver's name for the code, else knownName's: during the exit's teardown cuGetErrorName fails too.
     pub fn errorName(self: *const Driver, res: abi.Result) []const u8 {
         var s: ?[*:0]const u8 = null;
-        if (self.api.cuGetErrorName(res, &s) != abi.success) return "CUDA_ERROR_UNKNOWN_CODE";
-        return if (s) |p| std.mem.span(p) else "CUDA_ERROR_UNKNOWN_CODE";
+        if (self.api.cuGetErrorName(res, &s) == abi.success) if (s) |p| return std.mem.span(p);
+        return knownName(res) orelse "CUDA_ERROR_UNKNOWN_CODE";
+    }
+
+    /// A kernel's symbol name (cuFuncGetName), or "?" when the driver cannot say.
+    pub fn functionName(self: *const Driver, f: abi.Function) []const u8 {
+        const get = self.func_name orelse return "?";
+        var s: ?[*:0]const u8 = null;
+        if (get(&s, f) != abi.success) return "?";
+        return if (s) |p| std.mem.span(p) else "?";
     }
 
     pub fn errorText(self: *const Driver, res: abi.Result) []const u8 {
@@ -96,3 +115,43 @@ pub const Driver = struct {
         return std.mem.sliceTo(buf, 0);
     }
 };
+
+/// CUresult names (cuda.h) for when cuGetErrorName cannot answer.
+pub fn knownName(res: abi.Result) ?[]const u8 {
+    return switch (res) {
+        1 => "CUDA_ERROR_INVALID_VALUE",
+        2 => "CUDA_ERROR_OUT_OF_MEMORY",
+        3 => "CUDA_ERROR_NOT_INITIALIZED",
+        4 => "CUDA_ERROR_DEINITIALIZED",
+        101 => "CUDA_ERROR_INVALID_DEVICE",
+        200 => "CUDA_ERROR_INVALID_IMAGE",
+        201 => "CUDA_ERROR_INVALID_CONTEXT",
+        209 => "CUDA_ERROR_NO_BINARY_FOR_GPU",
+        400 => "CUDA_ERROR_INVALID_HANDLE",
+        500 => "CUDA_ERROR_NOT_FOUND",
+        600 => "CUDA_ERROR_NOT_READY",
+        700 => "CUDA_ERROR_ILLEGAL_ADDRESS",
+        701 => "CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES",
+        702 => "CUDA_ERROR_LAUNCH_TIMEOUT",
+        709 => "CUDA_ERROR_CONTEXT_IS_DESTROYED",
+        710 => "CUDA_ERROR_ASSERT",
+        714 => "CUDA_ERROR_HARDWARE_STACK_ERROR",
+        715 => "CUDA_ERROR_ILLEGAL_INSTRUCTION",
+        716 => "CUDA_ERROR_MISALIGNED_ADDRESS",
+        719 => "CUDA_ERROR_LAUNCH_FAILED",
+        999 => "CUDA_ERROR_UNKNOWN",
+        else => null,
+    };
+}
+
+/// CUDA_ERROR_DEINITIALIZED: the process is exiting (a stop signal, then exit) while a thread still calls the driver.
+pub fn shutdownNote(res: abi.Result) []const u8 {
+    return if (res == 4) " (the CUDA driver is shutting down: the process is exiting)" else "";
+}
+
+test "CUresult names without the driver: DEINITIALIZED is 4, unknown codes stay unknown" {
+    try std.testing.expectEqualStrings("CUDA_ERROR_DEINITIALIZED", knownName(4).?);
+    try std.testing.expectEqualStrings("CUDA_ERROR_ILLEGAL_ADDRESS", knownName(700).?);
+    try std.testing.expect(knownName(12345) == null);
+    try std.testing.expect(shutdownNote(4).len > 0 and shutdownNote(1).len == 0);
+}
diff --git a/zig/src/cuda/graph_cache.zig b/zig/src/cuda/graph_cache.zig
new file mode 100644
index 0000000000..1916b22b57
--- /dev/null
+++ b/zig/src/cuda/graph_cache.zig
@@ -0,0 +1,448 @@
+//! Keyed CUDA graphs for any family: one captured executable per key, launched again when the key comes back.
+
+const std = @import("std");
+const graph = @import("graph.zig");
+const Driver = @import("driver.zig").Driver;
+const Stream = @import("stream.zig").Stream;
+const quiet = @import("builtin").is_test; // a test's stderr is the build runner's channel
+
+/// One step's launches on `stream`; it must not synchronize, allocate device memory or read the device on the host.
+pub const Body = struct {
+    ctx: *anyopaque,
+    run: *const fn (ctx: *anyopaque, stream: Stream) anyerror!void,
+};
+
+/// Every rank's verdict AND-ed on a miss, never a hit; an off cache never votes, so ranks match `Settings.digest`.
+pub const Agree = struct {
+    ctx: *anyopaque,
+    all: *const fn (ctx: *anyopaque, ok: bool) anyerror!bool,
+};
+
+/// Capture a body into an executable graph, launch it, free it: a fake engine runs the policy in host tests.
+pub const Engine = struct {
+    ctx: *anyopaque,
+    capture: *const fn (ctx: *anyopaque, stream: Stream, body: Body) anyerror!*anyopaque,
+    launch: *const fn (ctx: *anyopaque, exec: *anyopaque, stream: Stream) anyerror!void,
+    free: *const fn (ctx: *anyopaque, exec: *anyopaque) void,
+};
+
+pub const Settings = struct {
+    on: bool = true, // off: every step runs eagerly
+    max: u32 = 48, // graphs held at most
+    min_keep: u32 = 8, // under the memory floor the cache evicts down to this many, then stops capturing
+    hold: bool = false, // under the floor a miss runs eagerly with no eviction: held graphs cost next to no memory
+
+    /// These settings under the operator's TF_GRAPH_CACHE value: `0` or `off` runs every step eagerly; else unchanged.
+    pub fn operator(s: Settings, value: ?[]const u8) Settings {
+        var out = s;
+        const v = value orelse return out;
+        if (std.mem.eql(u8, v, "0") or std.ascii.eqlIgnoreCase(v, "off")) out.on = false;
+        return out;
+    }
+
+    /// What a cache runs with: the operator's value, min_keep under max (0 with hold), and every off cache one value.
+    pub fn effective(s: Settings, value: ?[]const u8) Settings {
+        var out = s.operator(value);
+        if (!out.on or out.max == 0) return .{ .on = false, .max = 0, .min_keep = 0, .hold = false };
+        out.min_keep = if (out.hold) 0 else @min(out.min_keep, out.max);
+        return out;
+    }
+
+    /// `effective` under this process's TF_GRAPH_CACHE.
+    pub fn fromEnv(s: Settings) Settings {
+        return s.effective(if (std.c.getenv("TF_GRAPH_CACHE")) |v| std.mem.span(v) else null);
+    }
+
+    /// The effective settings as bytes for ranks to compare at start (a family's bootstrap agreement), never per step.
+    pub fn digest(s: Settings) [10]u8 {
+        const e = s.effective(null);
+        var out: [10]u8 = undefined;
+        out[0] = @intFromBool(e.on);
+        out[1] = @intFromBool(e.hold);
+        std.mem.writeInt(u32, out[2..6], e.max, .little);
+        std.mem.writeInt(u32, out[6..10], e.min_keep, .little);
+        return out;
+    }
+};
+
+pub const Outcome = enum { replayed, captured, eager };
+
+pub fn Cache(comptime Key: type) type {
+    return struct {
+        const Self = @This();
+        const Entry = struct { exec: *anyopaque, used: u64 };
+
+        gpa: std.mem.Allocator,
+        engine: Engine,
+        settings: Settings,
+        agree: ?Agree = null,
+        room: ?*const fn () bool = null, // false when this rank is below its memory floor; null: always room
+        entries: std.AutoArrayHashMapUnmanaged(Key, Entry) = .empty,
+        cap: u32,
+        tick: u64 = 0,
+        fingerprint: ?u64 = null,
+        failed: bool = false, // a capture failed on some rank: eager from here on
+        stats: struct { replayed: u64 = 0, captured: u64 = 0, eager: u64 = 0, evicted: u64 = 0, dropped: u64 = 0, floor_hits: u64 = 0 } = .{},
+
+        /// The cache with `settings.fromEnv()`: every rank sharing an `agree` must hold an equal `Settings.digest`.
+        pub fn init(gpa: std.mem.Allocator, engine: Engine, settings: Settings) Self {
+            const s = settings.fromEnv();
+            if (!quiet and settings.on and settings.max > 0 and !s.on) std.log.info("graph cache: off (TF_GRAPH_CACHE), every step eager", .{});
+            return .{ .gpa = gpa, .engine = engine, .settings = s, .cap = s.max };
+        }
+
+        pub fn deinit(c: *Self) void {
+            c.dropAll();
+            c.entries.deinit(c.gpa);
+        }
+
+        pub fn count(c: *const Self) usize {
+            return c.entries.count();
+        }
+
+        pub fn has(c: *const Self, key: Key) bool {
+            return c.entries.contains(key);
+        }
+
+        /// Every graph freed (a fingerprint change, a shutdown).
+        pub fn dropAll(c: *Self) void {
+            for (c.entries.values()) |e| c.engine.free(c.engine.ctx, e.exec);
+            c.stats.dropped += c.entries.count();
+            c.entries.clearRetainingCapacity();
+        }
+
+        fn agreed(c: *Self, ok: bool) !bool {
+            const a = c.agree orelse return ok;
+            return a.all(a.ctx, ok);
+        }
+
+        fn evictOldest(c: *Self) void {
+            var at: ?usize = null;
+            for (c.entries.values(), 0..) |e, i| {
+                if (at == null or e.used < c.entries.values()[at.?].used) at = i;
+            }
+            const i = at orelse return;
+            c.engine.free(c.engine.ctx, c.entries.values()[i].exec);
+            c.entries.swapRemoveAt(i);
+            c.stats.evicted += 1;
+        }
+
+        fn eager(c: *Self, stream: Stream, body: Body) !Outcome {
+            try body.run(body.ctx, stream);
+            c.stats.eager += 1;
+            return .eager;
+        }
+
+        /// One step: the key's graph, captured first on a miss; a new `fingerprint` (same step on all ranks) drops all.
+        pub fn run(c: *Self, key: Key, stream: Stream, fingerprint: u64, body: Body) !Outcome {
+            if (!c.settings.on or c.failed or c.settings.max == 0) return c.eager(stream, body);
+            if (c.fingerprint != null and c.fingerprint.? != fingerprint) c.dropAll();
+            c.fingerprint = fingerprint;
+            c.tick += 1;
+            if (c.entries.getPtr(key)) |e| {
+                e.used = c.tick;
+                try c.engine.launch(c.engine.ctx, e.exec, stream);
+                c.stats.replayed += 1;
+                return .replayed;
+            }
+            // a miss: room on every rank (a host map that cannot grow is none), else shed a quarter and lower the cap
+            const reserved = if (c.entries.ensureUnusedCapacity(c.gpa, 1)) |_| true else |_| false;
+            const room_here = reserved and (if (c.room) |f| f() else true);
+            if (!try c.agreed(room_here)) {
+                if (c.settings.hold) {
+                    if (!quiet and c.stats.floor_hits == 0) std.log.warn("graph cache: under the memory floor on {s} rank: this key eager, {d} graphs held (hold)", .{ if (room_here) "another" else "this", c.entries.count() });
+                    c.stats.floor_hits += 1;
+                    return c.eager(stream, body);
+                }
+                if (c.entries.count() > c.settings.min_keep) {
+                    const drop = @max(1, c.entries.count() / 4);
+                    for (0..drop) |_| if (c.entries.count() > c.settings.min_keep) c.evictOldest();
+                }
+                const was = c.cap;
+                c.cap = @intCast(@max(c.entries.count(), c.settings.min_keep)); // min_keep <= max (init)
+                c.stats.floor_hits += 1;
+                if (!quiet and c.cap != was) std.log.warn("graph cache: under the memory floor on {s} rank: cap {d} -> {d} ({d} held, {d} evicted, {d} eager so far)", .{ if (room_here) "another" else "this", was, c.cap, c.entries.count(), c.stats.evicted, c.stats.eager + 1 });
+                return c.eager(stream, body);
+            }
+            if (c.cap < c.settings.max) {
+                if (!quiet) std.log.info("graph cache: room again on every rank: cap {d} -> {d}", .{ c.cap, c.settings.max });
+                c.cap = c.settings.max;
+            }
+            while (c.entries.count() >= c.cap) c.evictOldest();
+            const exec = c.engine.capture(c.engine.ctx, stream, body) catch |err| blk: {
+                // the only trace of why graphs went off (the step then runs eagerly, which may fail the same way)
+                if (!quiet) std.log.warn("graph capture failed: {s} (key {any}, {d} graphs held)", .{ @errorName(err), key, c.entries.count() });
+                break :blk null;
+            };
+            var owned = exec; // freed here until the map holds it, so a failed agreement cannot leak it
+            defer if (owned) |x| c.engine.free(c.engine.ctx, x);
+            if (!try c.agreed(exec != null)) {
+                c.failed = true;
+                c.dropAll();
+                return c.eager(stream, body);
+            }
+            c.entries.putAssumeCapacityNoClobber(key, .{ .exec = exec.?, .used = c.tick });
+            owned = null;
+            try c.engine.launch(c.engine.ctx, exec.?, stream);
+            c.stats.captured += 1;
+            return .captured;
+        }
+    };
+}
+
+/// The driver's engine: thread-local stream capture (other threads' CUDA calls stay out), instantiate, upload, launch.
+pub const CudaEngine = struct {
+    d: *const Driver,
+
+    pub fn engine(e: *CudaEngine) Engine {
+        return .{ .ctx = e, .capture = capture, .launch = launch, .free = free };
+    }
+
+    const Held = struct { exec: graph.Exec };
+
+    fn capture(ctx: *anyopaque, stream: Stream, body: Body) anyerror!*anyopaque {
+        const e: *CudaEngine = @ptrCast(@alignCast(ctx));
+        try graph.beginCapture(stream, .thread_local);
+        body.run(body.ctx, stream) catch |err| {
+            // end the capture so the stream is usable again; the partial graph is dropped
+            if (graph.endCapture(stream)) |g| {
+                var gg = g;
+                gg.deinit();
+            } else |_| {}
+            return err;
+        };
+        var g = try graph.endCapture(stream);
+        defer g.deinit();
+        const h = try std.heap.c_allocator.create(Held);
+        errdefer std.heap.c_allocator.destroy(h);
+        h.exec = try g.instantiate();
+        errdefer h.exec.deinit();
+        try h.exec.upload(stream);
+        _ = e;
+        return h;
+    }
+
+    fn launch(_: *anyopaque, exec: *anyopaque, stream: Stream) anyerror!void {
+        const h: *Held = @ptrCast(@alignCast(exec));
+        try h.exec.launchOn(stream);
+    }
+
+    fn free(_: *anyopaque, exec: *anyopaque) void {
+        const h: *Held = @ptrCast(@alignCast(exec));
+        h.exec.deinit();
+        std.heap.c_allocator.destroy(h);
+    }
+};
+
+/// FNV-1a over the addresses a body captures (buffer bases, sizes): the fingerprint `run` compares, on one rank only.
+pub fn fingerprintOf(words: []const u64) u64 {
+    var h: u64 = 0xcbf29ce484222325;
+    for (words) |w| {
+        var x = w;
+        for (0..8) |_| {
+            h = (h ^ (x & 0xff)) *% 0x100000001b3;
+            x >>= 8;
+        }
+    }
+    return h;
+}
+
+// -- host tests: the policy with a fake engine --------------------------------------------------------------------
+
+const Fake = struct {
+    captures: u32 = 0,
+    launches: u32 = 0,
+    live: u32 = 0,
+    fail_capture: bool = false,
+    slots: [64]u8 = undefined,
+
+    fn engine(f: *Fake) Engine {
+        return .{ .ctx = f, .capture = cap, .launch = lau, .free = fre };
+    }
+    fn cap(ctx: *anyopaque, stream: Stream, body: Body) anyerror!*anyopaque {
+        const f: *Fake = @ptrCast(@alignCast(ctx));
+        if (f.fail_capture) return error.CaptureFailed;
+        _ = stream;
+        _ = body; // a capture records the body; nothing runs
+        f.captures += 1;
+        f.live += 1;
+        return &f.slots[f.captures % f.slots.len];
+    }
+    fn lau(ctx: *anyopaque, _: *anyopaque, _: Stream) anyerror!void {
+        const f: *Fake = @ptrCast(@alignCast(ctx));
+        f.launches += 1;
+    }
+    fn fre(ctx: *anyopaque, _: *anyopaque) void {
+        const f: *Fake = @ptrCast(@alignCast(ctx));
+        f.live -= 1;
+    }
+};
+
+const Ran = struct {
+    n: u32 = 0,
+    fn body(r: *Ran) Body {
+        return .{ .ctx = r, .run = go };
+    }
+    fn go(ctx: *anyopaque, _: Stream) anyerror!void {
+        const r: *Ran = @ptrCast(@alignCast(ctx));
+        r.n += 1;
+    }
+};
+
+const no_stream: Stream = .{ .d = undefined, .handle = null };
+
+test "graph cache: capture on first use, replay after, LRU at the cap" {
+    var f: Fake = .{};
+    var r: Ran = .{};
+    var c = Cache(u32).init(std.testing.allocator, f.engine(), .{ .max = 2 });
+    defer c.deinit();
+    try std.testing.expectEqual(Outcome.captured, try c.run(1, no_stream, 7, r.body()));
+    try std.testing.expectEqual(Outcome.replayed, try c.run(1, no_stream, 7, r.body()));
+    try std.testing.expectEqual(Outcome.captured, try c.run(2, no_stream, 7, r.body()));
+    _ = try c.run(1, no_stream, 7, r.body()); // 1 is now the most recent
+    try std.testing.expectEqual(Outcome.captured, try c.run(3, no_stream, 7, r.body())); // evicts 2
+    try std.testing.expect(c.has(1) and c.has(3) and !c.has(2));
+    try std.testing.expectEqual(@as(u32, 0), r.n); // the body never ran eagerly
+    try std.testing.expectEqual(@as(u32, 2), f.live);
+    try std.testing.expectEqual(@as(u32, 5), f.launches);
+}
+
+test "graph cache: a new fingerprint drops every graph" {
+    var f: Fake = .{};
+    var r: Ran = .{};
+    var c = Cache(u32).init(std.testing.allocator, f.engine(), .{});
+    defer c.deinit();
+    _ = try c.run(1, no_stream, 7, r.body());
+    _ = try c.run(2, no_stream, 7, r.body());
+    try std.testing.expectEqual(Outcome.captured, try c.run(1, no_stream, 8, r.body()));
+    try std.testing.expectEqual(@as(usize, 1), c.count());
+    try std.testing.expectEqual(@as(u32, 1), f.live);
+}
+
+const Peer = struct {
+    theirs: []const bool, // the other rank's verdicts, in call order
+    at: usize = 0,
+    fn agree(p: *Peer) Agree {
+        return .{ .ctx = p, .all = all };
+    }
+    fn all(ctx: *anyopaque, ok: bool) anyerror!bool {
+        const p: *Peer = @ptrCast(@alignCast(ctx));
+        const t = p.theirs[p.at];
+        p.at += 1;
+        return ok and t;
+    }
+};
+
+test "graph cache: a capture failing on another rank turns graphs off here too" {
+    var f: Fake = .{};
+    var r: Ran = .{};
+    // per miss: room, then the capture's verdict; the peer's capture fails on the second key
+    var peer: Peer = .{ .theirs = &.{ true, true, true, false } };
+    var c = Cache(u32).init(std.testing.allocator, f.engine(), .{});
+    c.agree = peer.agree();
+    defer c.deinit();
+    try std.testing.expectEqual(Outcome.captured, try c.run(1, no_stream, 7, r.body()));
+    try std.testing.expectEqual(Outcome.eager, try c.run(2, no_stream, 7, r.body()));
+    try std.testing.expect(c.failed);
+    try std.testing.expectEqual(@as(u32, 0), f.live); // this rank's good graphs are freed as well
+    try std.testing.expectEqual(Outcome.eager, try c.run(1, no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(u32, 2), r.n);
+}
+
+var room_flag = true;
+fn roomFn() bool {
+    return room_flag;
+}
+
+test "graph cache: below the memory floor the least recent quarter goes, the cap follows, the step runs eagerly" {
+    var f: Fake = .{};
+    var r: Ran = .{};
+    var c = Cache(u32).init(std.testing.allocator, f.engine(), .{ .max = 48, .min_keep = 2 });
+    c.room = roomFn;
+    defer c.deinit();
+    room_flag = true;
+    for (0..8) |k| _ = try c.run(@intCast(k), no_stream, 7, r.body());
+    room_flag = false;
+    try std.testing.expectEqual(Outcome.eager, try c.run(100, no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(usize, 6), c.count());
+    try std.testing.expectEqual(@as(u32, 6), c.cap);
+    try std.testing.expect(!c.has(0) and !c.has(1) and c.has(7));
+    room_flag = true;
+    // room again: the cap is back to max, the miss captures without evicting
+    try std.testing.expectEqual(Outcome.captured, try c.run(100, no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(usize, 7), c.count());
+    try std.testing.expectEqual(@as(u32, 48), c.cap);
+    try std.testing.expect(c.has(7) and c.has(2));
+    try std.testing.expectEqual(@as(u64, 1), c.stats.floor_hits);
+}
+
+test "graph cache: a dip under the floor does not leave later keys churning once memory is back" {
+    var f: Fake = .{};
+    var r: Ran = .{};
+    var c = Cache(u32).init(std.testing.allocator, f.engine(), .{ .max = 16, .min_keep = 2 });
+    c.room = roomFn;
+    defer c.deinit();
+    room_flag = true;
+    for (0..12) |k| _ = try c.run(@intCast(k), no_stream, 7, r.body());
+    room_flag = false; // a long prompt's workspace held
+    _ = try c.run(200, no_stream, 7, r.body());
+    room_flag = true; // freed: the decode after it meets 12 new keys (a new context bucket), then replays them
+    const before = c.stats.evicted;
+    for (0..12) |k| _ = try c.run(@intCast(300 + k), no_stream, 7, r.body());
+    for (0..12) |k| try std.testing.expectEqual(Outcome.replayed, try c.run(@intCast(300 + k), no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(u32, 16), c.cap);
+    // only what the cap of 16 forces: 9 left after the dip + 12 new = 21 -> 5 evictions, none during the replays
+    try std.testing.expectEqual(before + 5, c.stats.evicted);
+}
+
+test "graph cache: hold - below the floor a miss runs eagerly, nothing is evicted, the held keys replay" {
+    var f: Fake = .{};
+    var r: Ran = .{};
+    var c = Cache(u32).init(std.testing.allocator, f.engine(), .{ .max = 48, .min_keep = 2, .hold = true });
+    c.room = roomFn;
+    defer c.deinit();
+    room_flag = true;
+    for (0..8) |k| _ = try c.run(@intCast(k), no_stream, 7, r.body());
+    room_flag = false;
+    try std.testing.expectEqual(Outcome.eager, try c.run(100, no_stream, 7, r.body()));
+    try std.testing.expectEqual(Outcome.eager, try c.run(101, no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(usize, 8), c.count());
+    try std.testing.expectEqual(@as(u32, 48), c.cap);
+    for (0..8) |k| try std.testing.expectEqual(Outcome.replayed, try c.run(@intCast(k), no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(u64, 0), c.stats.evicted);
+    try std.testing.expectEqual(@as(u64, 2), c.stats.floor_hits);
+    room_flag = true;
+    try std.testing.expectEqual(Outcome.captured, try c.run(100, no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(usize, 9), c.count());
+}
+
+test "graph cache: off runs every step eagerly" {
+    var f: Fake = .{};
+    var r: Ran = .{};
+    var c = Cache(u32).init(std.testing.allocator, f.engine(), .{ .on = false });
+    defer c.deinit();
+    try std.testing.expectEqual(Outcome.eager, try c.run(1, no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(u32, 0), f.captures);
+}
+
+test "graph cache: TF_GRAPH_CACHE=0 or off turns a family's cache off without a rebuild; anything else leaves it" {
+    const on: Settings = .{};
+    for ([_][]const u8{ "0", "off", "OFF" }) |v| try std.testing.expect(!on.operator(v).on);
+    for ([_]?[]const u8{ null, "1", "on", "" }) |v| try std.testing.expect(on.operator(v).on);
+    try std.testing.expect(!(Settings{ .on = false }).operator("1").on); // the operator only turns a cache off
+    var f: Fake = .{};
+    var r: Ran = .{};
+    var c = Cache(u32).init(std.testing.allocator, f.engine(), on.operator("off"));
+    defer c.deinit();
+    for (0..3) |_| try std.testing.expectEqual(Outcome.eager, try c.run(1, no_stream, 7, r.body()));
+    try std.testing.expectEqual(@as(u32, 0), f.captures);
+}
+
+test "graph cache: fingerprint" {
+    try std.testing.expect(fingerprintOf(&.{ 1, 2 }) != fingerprintOf(&.{ 2, 1 }));
+}
+
+test {
+    _ = @import("graph_cache_regression_test.zig");
+    _ = @import("graph_cache_ranks_test.zig");
+}
diff --git a/zig/src/cuda/graph_cache_ranks_test.zig b/zig/src/cuda/graph_cache_ranks_test.zig
new file mode 100644
index 0000000000..18b8a87e9d
--- /dev/null
+++ b/zig/src/cuda/graph_cache_ranks_test.zig
@@ -0,0 +1,136 @@
+//! Two in-process ranks with fake engines: equal effective settings vote on misses only; a mismatch is a stall.
+
+const std = @import("std");
+const cache = @import("graph_cache.zig");
+const Stream = @import("stream.zig").Stream;
+const no_stream: Stream = .{ .d = undefined, .handle = null };
+const equal = std.testing.expectEqual;
+const Settings = cache.Settings;
+
+const Engine = struct {
+    captures: usize = 0,
+    live: usize = 0,
+    handles: [8]u8 = undefined,
+
+    fn engine(e: *Engine) cache.Engine {
+        return .{ .ctx = e, .capture = capture, .launch = launch, .free = free };
+    }
+    fn capture(p: *anyopaque, _: Stream, _: cache.Body) !*anyopaque {
+        const e: *Engine = @ptrCast(@alignCast(p));
+        e.captures += 1;
+        e.live += 1;
+        return &e.handles[e.captures % e.handles.len];
+    }
+    fn launch(_: *anyopaque, _: *anyopaque, _: Stream) !void {}
+    fn free(p: *anyopaque, _: *anyopaque) void {
+        const e: *Engine = @ptrCast(@alignCast(p));
+        e.live -= 1;
+    }
+    fn body(e: *Engine) cache.Body {
+        return .{ .ctx = e, .run = run };
+    }
+    fn run(_: *anyopaque, _: Stream) !void {}
+};
+
+/// The two ranks' votes in call order; a rank waits at most `wait_ns` for the other's vote of the same round.
+const Board = struct {
+    votes: [2][16]bool = undefined,
+    cast: [2]std.atomic.Value(u32) = .{ .init(0), .init(0) },
+    wait_ns: u64,
+
+    fn vote(b: *Board, rank: u1, ok: bool) !bool {
+        const r = b.cast[rank].load(.monotonic);
+        if (r == b.votes[rank].len) return error.TooManyVotes;
+        b.votes[rank][r] = ok;
+        b.cast[rank].store(r + 1, .release);
+        const other = 1 - rank;
+        const deadline = now() + b.wait_ns;
+        while (b.cast[other].load(.acquire) <= r) {
+            if (now() > deadline) return error.PeerNeverVoted;
+            std.Thread.yield() catch {};
+        }
+        return ok and b.votes[other][r];
+    }
+
+    fn now() u64 {
+        return @intCast(std.Io.Clock.awake.now(std.testing.io).toNanoseconds());
+    }
+};
+
+const Rank = struct {
+    board: *Board,
+    id: u1,
+    engine: Engine = .{},
+    outcomes: [8]cache.Outcome = undefined,
+    err: ?anyerror = null,
+
+    fn all(ctx: *anyopaque, ok: bool) anyerror!bool {
+        const r: *Rank = @ptrCast(@alignCast(ctx));
+        return r.board.vote(r.id, ok);
+    }
+
+    /// One rank's steps over `keys` with `settings` exactly as its process resolved them.
+    fn steps(r: *Rank, settings: Settings, ks: []const u32) void {
+        var c = cache.Cache(u32).init(std.testing.allocator, r.engine.engine(), settings);
+        defer c.deinit();
+        c.agree = .{ .ctx = r, .all = all };
+        for (ks, 0..) |k, i| r.outcomes[i] = c.run(k, no_stream, 7, r.engine.body()) catch |e| {
+            r.err = e;
+            return;
+        };
+    }
+};
+
+const keys = [_]u32{ 1, 2, 1, 2, 1 };
+
+test "graph cache ranks: equal enabled settings capture together and vote on each miss, never on a hit" {
+    var board: Board = .{ .wait_ns = 5 * std.time.ns_per_s };
+    var ranks = [2]Rank{ .{ .board = &board, .id = 0 }, .{ .board = &board, .id = 1 } };
+    const s = (Settings{ .max = 4 }).effective(null);
+    const t = try std.Thread.spawn(.{}, Rank.steps, .{ &ranks[1], s, &keys });
+    ranks[0].steps(s, &keys);
+    t.join();
+    const want = [_]cache.Outcome{ .captured, .captured, .replayed, .replayed, .replayed };
+    for (&ranks) |*r| {
+        try equal(@as(?anyerror, null), r.err);
+        try std.testing.expectEqualSlices(cache.Outcome, &want, r.outcomes[0..keys.len]);
+        try equal(@as(u32, 4), board.cast[r.id].load(.acquire)); // two misses: room and capture each, hits none
+        try equal(@as(usize, 0), r.engine.live);
+    }
+}
+
+test "graph cache ranks: TF_GRAPH_CACHE 0, off and OFF are one effective setting, and an off cache never votes" {
+    const d: Settings = .{};
+    const off = d.effective("0").digest();
+    try std.testing.expectEqualSlices(u8, &off, &d.effective("off").digest());
+    try std.testing.expectEqualSlices(u8, &off, &d.effective("OFF").digest());
+    try std.testing.expectEqualSlices(u8, &off, &(Settings{ .max = 0, .hold = true }).digest());
+    try std.testing.expect(!std.mem.eql(u8, &off, &d.effective(null).digest()));
+    const held: Settings = .{ .hold = true, .min_keep = 2 };
+    try std.testing.expectEqualSlices(u8, &held.digest(), &(Settings{ .hold = true, .min_keep = 30 }).digest());
+    try std.testing.expectEqualSlices(u8, &d.effective("1").digest(), &d.effective(null).digest());
+    var board: Board = .{ .wait_ns = 0 };
+    var ranks = [2]Rank{ .{ .board = &board, .id = 0 }, .{ .board = &board, .id = 1 } };
+    ranks[0].steps(d.effective("0"), &keys);
+    ranks[1].steps(d.effective("OFF"), &keys);
+    for (&ranks) |*r| {
+        try equal(@as(?anyerror, null), r.err);
+        for (r.outcomes[0..keys.len]) |o| try equal(cache.Outcome.eager, o);
+        try equal(@as(u32, 0), board.cast[r.id].load(.acquire));
+        try equal(@as(usize, 0), r.engine.captures);
+    }
+}
+
+test "graph cache ranks: one rank switched off leaves the other waiting at its first miss; digests show it first" {
+    const d: Settings = .{};
+    try std.testing.expect(!std.mem.eql(u8, &d.effective(null).digest(), &d.effective("off").digest()));
+    var board: Board = .{ .wait_ns = 50 * std.time.ns_per_ms };
+    var ranks = [2]Rank{ .{ .board = &board, .id = 0 }, .{ .board = &board, .id = 1 } };
+    ranks[1].steps(d.effective("off"), &keys);
+    try equal(@as(?anyerror, null), ranks[1].err);
+    try equal(@as(u32, 0), board.cast[1].load(.acquire)); // the off rank ran every step eagerly without a vote
+    ranks[0].steps(d.effective(null), &keys);
+    try equal(@as(?anyerror, error.PeerNeverVoted), ranks[0].err); // in production this wait has no deadline
+    try equal(@as(u32, 1), board.cast[0].load(.acquire));
+    try equal(@as(usize, 0), ranks[0].engine.live);
+}
diff --git a/zig/src/cuda/graph_cache_regression_test.zig b/zig/src/cuda/graph_cache_regression_test.zig
new file mode 100644
index 0000000000..7f2da1063a
--- /dev/null
+++ b/zig/src/cuda/graph_cache_regression_test.zig
@@ -0,0 +1,181 @@
+//! Failure ownership, rank verdict counts and bounded-cap regressions for the graph cache.
+
+const std = @import("std");
+const cache = @import("graph_cache.zig");
+const Stream = @import("stream.zig").Stream;
+const no_stream: Stream = .{ .d = undefined, .handle = null };
+const expect = std.testing.expect;
+const equal = std.testing.expectEqual;
+
+const Engine = struct {
+    captures: usize = 0,
+    live: usize = 0,
+    launches: usize = 0,
+    bodies: usize = 0,
+    handles: [32]u8 = undefined,
+
+    fn engine(e: *Engine) cache.Engine {
+        return .{ .ctx = e, .capture = capture, .launch = launch, .free = free };
+    }
+    fn capture(p: *anyopaque, _: Stream, _: cache.Body) !*anyopaque {
+        const e: *Engine = @ptrCast(@alignCast(p));
+        e.captures += 1;
+        e.live += 1;
+        return &e.handles[e.captures % e.handles.len];
+    }
+    fn launch(p: *anyopaque, _: *anyopaque, _: Stream) !void {
+        const e: *Engine = @ptrCast(@alignCast(p));
+        e.launches += 1;
+    }
+    fn free(p: *anyopaque, _: *anyopaque) void {
+        const e: *Engine = @ptrCast(@alignCast(p));
+        std.debug.assert(e.live > 0);
+        e.live -= 1;
+    }
+    fn body(e: *Engine) cache.Body {
+        return .{ .ctx = e, .run = run };
+    }
+    fn run(p: *anyopaque, _: Stream) !void {
+        const e: *Engine = @ptrCast(@alignCast(p));
+        e.bodies += 1;
+    }
+};
+
+const Peer = struct {
+    verdicts: []const bool,
+    at: usize = 0,
+    throw_at: ?usize = null,
+    fn agree(p: *Peer) cache.Agree {
+        return .{ .ctx = p, .all = all };
+    }
+    fn all(ctx: *anyopaque, ok: bool) !bool {
+        const p: *Peer = @ptrCast(@alignCast(ctx));
+        const at = p.at;
+        p.at += 1;
+        if (p.throw_at == at) return error.CollectiveFailed;
+        if (at >= p.verdicts.len) return error.UnexpectedCollective;
+        return ok and p.verdicts[at];
+    }
+};
+
+test "graph cache regression: capture agreement throw frees the not-yet-inserted executable" {
+    var e: Engine = .{};
+    // per miss: room, then the capture's verdict
+    var p: Peer = .{ .verdicts = &.{ true, true }, .throw_at = 1 };
+    var c = cache.Cache(u32).init(std.testing.allocator, e.engine(), .{});
+    defer c.deinit();
+    c.agree = p.agree();
+    try std.testing.expectError(error.CollectiveFailed, c.run(1, no_stream, 7, e.body()));
+    try equal(@as(usize, 1), e.captures);
+    try equal(@as(usize, 0), e.live);
+    try equal(@as(usize, 0), c.count());
+    try equal(@as(usize, 2), p.at);
+}
+
+test "graph cache regression: local map OOM declines before any capture and can recover" {
+    var e: Engine = .{};
+    var failing = std.testing.FailingAllocator.init(std.testing.allocator, .{ .fail_index = 0 });
+    var c = cache.Cache(u32).init(failing.allocator(), e.engine(), .{});
+    defer c.deinit();
+    try equal(cache.Outcome.eager, try c.run(1, no_stream, 7, e.body()));
+    try expect(failing.has_induced_failure);
+    try equal(@as(usize, 0), e.captures);
+    try equal(@as(usize, 0), e.live);
+    try expect(!c.failed);
+    failing.fail_index = std.math.maxInt(usize);
+    try equal(cache.Outcome.captured, try c.run(1, no_stream, 7, e.body()));
+    try equal(@as(usize, 1), e.live);
+}
+
+test "graph cache regression: peer map OOM makes a locally prepared rank eager before capture" {
+    var e: Engine = .{};
+    var p: Peer = .{ .verdicts = &.{false} };
+    var c = cache.Cache(u32).init(std.testing.allocator, e.engine(), .{});
+    defer c.deinit();
+    c.agree = p.agree();
+    try equal(cache.Outcome.eager, try c.run(1, no_stream, 7, e.body()));
+    try equal(@as(usize, 1), p.at);
+    try equal(@as(usize, 0), e.captures);
+    try equal(@as(usize, 1), e.bodies);
+}
+
+test "graph cache regression: local OOM still contributes the room verdict" {
+    var e: Engine = .{};
+    var failing = std.testing.FailingAllocator.init(std.testing.allocator, .{ .fail_index = 0 });
+    var p: Peer = .{ .verdicts = &.{true} };
+    var c = cache.Cache(u32).init(failing.allocator(), e.engine(), .{});
+    defer c.deinit();
+    c.agree = p.agree();
+    try equal(cache.Outcome.eager, try c.run(1, no_stream, 7, e.body()));
+    try equal(@as(usize, 1), p.at);
+    try equal(@as(usize, 0), e.captures);
+}
+
+test "graph cache regression: hits make no collective; a fingerprint change shared by every rank recaptures" {
+    var e: Engine = .{};
+    var p: Peer = .{ .verdicts = &.{ true, true, true, true, true, true } };
+    var c = cache.Cache(u32).init(std.testing.allocator, e.engine(), .{});
+    defer c.deinit();
+    c.agree = p.agree();
+    _ = try c.run(1, no_stream, 7, e.body());
+    _ = try c.run(2, no_stream, 7, e.body());
+    try equal(@as(usize, 4), p.at);
+    for (0..4) |_| try equal(cache.Outcome.replayed, try c.run(1, no_stream, 7, e.body()));
+    try equal(@as(usize, 4), p.at);
+    // a new layout generation on every rank: both graphs go, the next miss agrees twice as any other
+    try equal(cache.Outcome.captured, try c.run(1, no_stream, 8, e.body()));
+    try equal(@as(usize, 6), p.at);
+    try equal(@as(u64, 2), c.stats.dropped);
+    try equal(@as(usize, 1), c.count());
+    try equal(@as(usize, 1), e.live);
+}
+
+var room = true;
+fn available() bool {
+    return room;
+}
+
+test "graph cache regression: zero maximum remains eager across pressure recovery, with no collective" {
+    var e: Engine = .{};
+    var p: Peer = .{ .verdicts = &.{} };
+    var c = cache.Cache(u32).init(std.testing.allocator, e.engine(), .{ .max = 0 });
+    defer c.deinit();
+    c.room = available;
+    c.agree = p.agree();
+    for ([_]bool{ false, true, false, true }) |ok| {
+        room = ok;
+        try equal(cache.Outcome.eager, try c.run(1, no_stream, 7, e.body()));
+        try equal(@as(u32, 0), c.cap);
+        try equal(@as(usize, 0), c.count());
+    }
+    try equal(@as(usize, 0), e.captures);
+    try equal(@as(usize, 0), p.at);
+}
+
+test "graph cache regression: default floor cannot enlarge a two-entry maximum" {
+    var e: Engine = .{};
+    var c = cache.Cache(u32).init(std.testing.allocator, e.engine(), .{ .max = 2 });
+    defer c.deinit();
+    c.room = available;
+    try equal(@as(u32, 2), c.settings.min_keep);
+    room = false;
+    _ = try c.run(1, no_stream, 7, e.body());
+    room = true;
+    for (0..8) |key| {
+        _ = try c.run(@intCast(key), no_stream, 7, e.body());
+        try expect(c.count() <= 2 and c.cap <= 2 and e.live <= 2);
+    }
+}
+
+test "graph cache regression: zero keep floor recovers after reducing the cap to zero" {
+    var e: Engine = .{};
+    var c = cache.Cache(u32).init(std.testing.allocator, e.engine(), .{ .max = 2, .min_keep = 0 });
+    defer c.deinit();
+    c.room = available;
+    room = false;
+    try equal(cache.Outcome.eager, try c.run(1, no_stream, 7, e.body()));
+    try equal(@as(u32, 0), c.cap);
+    room = true;
+    try equal(cache.Outcome.captured, try c.run(1, no_stream, 7, e.body()));
+    try equal(@as(u32, 2), c.cap);
+}
diff --git a/zig/src/cuda/launch.zig b/zig/src/cuda/launch.zig
index 56bba5272a..47cbf720b6 100644
--- a/zig/src/cuda/launch.zig
+++ b/zig/src/cuda/launch.zig
@@ -4,7 +4,8 @@ const std = @import("std");
 const abi = @import("abi.zig");
 const Function = @import("module.zig").Function;
 const Stream = @import("stream.zig").Stream;
-const Error = @import("driver.zig").Error;
+const driver = @import("driver.zig");
+const Error = driver.Error;
 
 pub const Dim3 = abi.Dim3;
 
@@ -88,7 +89,7 @@ pub fn launch(f: Function, cfg: Config, stream: Stream, args: *Args) Error!void
     const g = cfg.grid;
     const b = cfg.block;
     if (cfg.plain()) {
-        return d.check(d.api.cuLaunchKernel(f.handle, g.x, g.y, g.z, b.x, b.y, b.z, cfg.shared, stream.handle, args.pointers(), null), "cuLaunchKernel");
+        return launched(f, cfg, d.api.cuLaunchKernel(f.handle, g.x, g.y, g.z, b.x, b.y, b.z, cfg.shared, stream.handle, args.pointers(), null), "cuLaunchKernel");
     }
     var attrs: [4]abi.LaunchAttribute = undefined;
     const n = fillAttributes(cfg, &attrs);
@@ -104,7 +105,37 @@ pub fn launch(f: Function, cfg: Config, stream: Stream, args: *Args) Error!void
         .attrs = &attrs,
         .num_attrs = @intCast(n),
     };
-    try d.check(d.api.cuLaunchKernelEx(&lc, f.handle, args.pointers(), null), "cuLaunchKernelEx");
+    try launched(f, cfg, d.api.cuLaunchKernelEx(&lc, f.handle, args.pointers(), null), "cuLaunchKernelEx");
+}
+
+/// A launch's result; a failure names the kernel and its configuration with the code.
+fn launched(f: Function, cfg: Config, res: abi.Result, what: []const u8) Error!void {
+    if (res == abi.success) return;
+    if (res == abi.error_not_ready) return error.NotReady; // as Driver.check: not a failure to log
+    const d = f.d;
+    var buf: [160]u8 = undefined;
+    std.log.err("{s}: {s} ({d}) {s}; {s}{s}", .{ what, d.errorName(res), res, d.errorText(res), describe(&buf, d.functionName(f.handle), cfg), driver.shutdownNote(res) });
+    return driver.Driver.errorOf(res);
+}
+
+/// "kernel <name>, grid XxYxZ, block XxYxZ, shared N B" (+ cluster / PDL / cooperative when set).
+pub fn describe(buf: []u8, name: []const u8, cfg: Config) []const u8 {
+    const g = cfg.grid;
+    const b = cfg.block;
+    var w = std.Io.Writer.fixed(buf);
+    w.print("kernel {s}, grid {d}x{d}x{d}, block {d}x{d}x{d}, shared {d} B", .{ name, g.x, g.y, g.z, b.x, b.y, b.z, cfg.shared }) catch return w.buffered();
+    if (cfg.cluster) |c| w.print(", cluster {d}x{d}x{d}", .{ c.x, c.y, c.z }) catch return w.buffered();
+    if (cfg.pdl) w.writeAll(", pdl") catch return w.buffered();
+    if (cfg.cooperative) w.writeAll(", cooperative") catch return w.buffered();
+    return w.buffered();
+}
+
+test "a failed launch's description: the kernel, its grid, block and shared bytes, bounded" {
+    var buf: [160]u8 = undefined;
+    const cfg: Config = .{ .grid = .{ .x = 70000, .y = 2, .z = 1 }, .block = .{ .x = 128, .y = 1, .z = 1 }, .shared = 49152, .pdl = true };
+    try std.testing.expectEqualStrings("kernel _stream_pf, grid 70000x2x1, block 128x1x1, shared 49152 B, pdl", describe(&buf, "_stream_pf", cfg));
+    var small: [24]u8 = undefined;
+    try std.testing.expect(describe(&small, "a_very_long_kernel_symbol_name", cfg).len <= small.len);
 }
 
 fn attribute(id: abi.LaunchAttributeId) abi.LaunchAttribute {
diff --git a/zig/src/cuda/memory.zig b/zig/src/cuda/memory.zig
index 28c5b6244f..0f8f407806 100644
--- a/zig/src/cuda/memory.zig
+++ b/zig/src/cuda/memory.zig
@@ -100,18 +100,39 @@ fn lockCounts() void {
     while (!counts_mutex.tryLock()) std.Thread.yield() catch {};
 }
 
+/// A sub-allocator under DeviceBuffer (arena.zig): `alloc` gives an address or null for cuMemAlloc, `free` owns it.
+pub const Hook = struct {
+    ctx: *anyopaque,
+    alloc: *const fn (ctx: *anyopaque, len: usize) ?abi.DevicePtr,
+    free: *const fn (ctx: *anyopaque, p: abi.DevicePtr) bool, // true when the address was the hook's
+};
+
+var alloc_hook: ?Hook = null;
+
+/// Installs, or with null removes, the sub-allocator: once at boot, before the buffers it should hold.
+pub fn setHook(h: ?Hook) void {
+    alloc_hook = h;
+}
+
+pub fn currentHook() ?Hook {
+    return alloc_hook;
+}
+
 pub const DeviceBuffer = struct {
     d: *const Driver,
     ptr: abi.DevicePtr,
     len: usize,
     borrowed: bool = false, // a span of memory another value owns (a carveout's): free leaves it alone
 
-    /// `len` bytes, 256-byte aligned by the driver; zero bytes allocate nothing and hold address 0.
+    /// `len` bytes, 256-byte aligned (512 from an installed arena); zero bytes allocate nothing and hold address 0.
     pub fn alloc(d: *const Driver, len: usize) Error!DeviceBuffer {
         try reserveAllocation(len);
         errdefer finishAllocation(null, len);
         var p: abi.DevicePtr = 0;
-        if (len > 0) try d.check(d.api.cuMemAlloc_v2(&p, len), "cuMemAlloc");
+        if (len > 0) {
+            if (alloc_hook) |h| p = h.alloc(h.ctx, len) orelse 0;
+            if (p == 0) try d.check(d.api.cuMemAlloc_v2(&p, len), "cuMemAlloc");
+        }
         finishAllocation(&device_bytes, len);
         return .{ .d = d, .ptr = p, .len = len };
     }
@@ -129,7 +150,10 @@ pub const DeviceBuffer = struct {
 
     pub fn free(self: *DeviceBuffer) void {
         if (!self.borrowed) {
-            if (self.ptr != 0) _ = self.d.api.cuMemFree_v2(self.ptr);
+            if (self.ptr != 0) {
+                const mine = if (alloc_hook) |h| h.free(h.ctx, self.ptr) else false;
+                if (!mine) _ = self.d.api.cuMemFree_v2(self.ptr);
+            }
             freed(&device_bytes, self.len);
         }
         self.* = undefined;
@@ -266,6 +290,28 @@ test "freeing a borrowed span leaves the driver and the counts alone" {
     try std.testing.expectEqual(before.device, usage(false).device);
 }
 
+test "a buffer the hook owns is freed by the hook, never by the driver" {
+    const Owner = struct {
+        var frees: u32 = 0;
+        fn alloc(_: *anyopaque, _: usize) ?abi.DevicePtr {
+            return null;
+        }
+        fn free(_: *anyopaque, p: abi.DevicePtr) bool {
+            frees += 1;
+            return p == 0x2000;
+        }
+    };
+    var ctx: u8 = 0;
+    setHook(.{ .ctx = &ctx, .alloc = Owner.alloc, .free = Owner.free });
+    defer setHook(null);
+    held(&device_bytes, 4096);
+    var b: DeviceBuffer = .{ .d = undefined, .ptr = 0x2000, .len = 4096 };
+    b.free();
+    var c: DeviceBuffer = .{ .d = undefined, .ptr = 0x3000, .len = 64, .borrowed = true };
+    c.free();
+    try std.testing.expectEqual(@as(u32, 1), Owner.frees);
+}
+
 test "peak reset preserves an allocation between its read and store" {
     const before = usage(true);
     held(&device_bytes, 100);
diff --git a/zig/src/cuda/root.zig b/zig/src/cuda/root.zig
index bd3ffa6845..0a52db2a17 100644
--- a/zig/src/cuda/root.zig
+++ b/zig/src/cuda/root.zig
@@ -54,6 +54,8 @@ pub const tp_group = @import("tp_group.zig");
 pub const tp_watch = @import("tp_watch.zig");
 pub const ck = @import("ck.zig");
 pub const modelopt = @import("modelopt.zig");
+pub const graph_cache = @import("graph_cache.zig");
+pub const arena = @import("arena.zig");
 pub const tree_attention = @import("tree_attention.zig");
 pub const block_attention = @import("block_attention.zig");
 pub const window_attention = @import("window_attention.zig");
@@ -100,6 +102,8 @@ test {
     _ = @import("modelopt.zig");
     _ = @import("ck.zig");
     _ = @import("tp_watch.zig");
+    _ = graph_cache;
+    _ = arena;
     _ = @import("shared_test.zig");
     _ = @import("tree_attention.zig");
     _ = @import("block_attention.zig");