From a3303c75569177788544bed303409969a0bf176c Mon Sep 17 00:00:00 2001
From: Jay Leaton <58495321+jayleaton@users.noreply.github.com>
Date: Sat, 10 Oct 2026 10:28:21 +0700
Subject: [PATCH 1/9] CUDA keyed graph cache and device arena, both off until a
 family installs them.

graph_cache.Cache captures one executable graph per caller-named key and replays it when the key comes back; ranks agree on every capture, eviction and fallback through an Agree hook, and a memory floor turns misses eager.
arena.Arena sub-allocates DeviceBuffer.alloc from a few large chunks (VMM or plain cuMemAlloc) through memory.setHook; borrowed spans and pointers it does not own are freed as before.
---
 zig/src/cuda/abi.zig         |   2 +
 zig/src/cuda/arena.zig       | 331 +++++++++++++++++++++++++++++
 zig/src/cuda/graph_cache.zig | 392 +++++++++++++++++++++++++++++++++++
 zig/src/cuda/memory.zig      |  52 ++++-
 zig/src/cuda/root.zig        |   5 +
 5 files changed, 779 insertions(+), 3 deletions(-)
 create mode 100644 zig/src/cuda/arena.zig
 create mode 100644 zig/src/cuda/graph_cache.zig

diff --git a/zig/src/cuda/abi.zig b/zig/src/cuda/abi.zig
index 485edaabde..a099686c67 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 = 126,
 };
 
 pub const FunctionAttribute = enum(c_int) {
diff --git a/zig/src/cuda/arena.zig b/zig/src/cuda/arena.zig
new file mode 100644
index 0000000000..2c80b7581b
--- /dev/null
+++ b/zig/src/cuda/arena.zig
@@ -0,0 +1,331 @@
+//! 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
+    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);
+        if (a.free_small.getPtr(n)) |l| if (l.pop()) |p| {
+            try a.live.put(a.gpa, 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;
+        try a.live.put(a.gpa, 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| {
+            a.be.release(a.be.ctx, p, e.value);
+            return true;
+        }
+        const e = a.live.fetchRemove(p) orelse return false;
+        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;
+    }
+
+    // -- 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));
+}
diff --git a/zig/src/cuda/graph_cache.zig b/zig/src/cuda/graph_cache.zig
new file mode 100644
index 0000000000..a2fb8e18bb
--- /dev/null
+++ b/zig/src/cuda/graph_cache.zig
@@ -0,0 +1,392 @@
+//! 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, so ranks capture, evict and fall back at the same steps; never called in a capture.
+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
+};
+
+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 } = .{},
+
+        pub fn init(gpa: std.mem.Allocator, engine: Engine, settings: Settings) Self {
+            return .{ .gpa = gpa, .engine = engine, .settings = settings, .cap = settings.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 launched, captured first on a miss; a new `fingerprint` drops every graph first.
+        pub fn run(c: *Self, key: Key, stream: Stream, fingerprint: u64, body: Body) !Outcome {
+            if (!c.settings.on or c.failed) 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, else shed the least recent quarter and lower the cap (once below it, eager)
+            const room_here = if (c.room) |f| f() else true;
+            if (!try c.agreed(room_here and c.cap > 0)) {
+                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));
+                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;
+            };
+            if (!try c.agreed(exec != null)) {
+                if (exec) |x| c.engine.free(c.engine.ctx, x);
+                c.failed = true;
+                c.dropAll();
+                return c.eager(stream, body);
+            }
+            try c.entries.put(c.gpa, key, .{ .exec = exec.?, .used = c.tick });
+            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.
+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: fingerprint" {
+    try std.testing.expect(fingerprintOf(&.{ 1, 2 }) != fingerprintOf(&.{ 2, 1 }));
+}
diff --git a/zig/src/cuda/memory.zig b/zig/src/cuda/memory.zig
index 2c469bae2c..6d802980c7 100644
--- a/zig/src/cuda/memory.zig
+++ b/zig/src/cuda/memory.zig
@@ -49,16 +49,37 @@ 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 {
         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");
+        }
         held(&device_bytes, len);
         return .{ .d = d, .ptr = p, .len = len };
     }
@@ -73,7 +94,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;
@@ -209,6 +233,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 3053c09a6b..22f9454a83 100644
--- a/zig/src/cuda/root.zig
+++ b/zig/src/cuda/root.zig
@@ -6,6 +6,7 @@ pub const Error = @import("driver.zig").Error;
 pub const Context = @import("context.zig").Context;
 pub const DeviceBuffer = @import("memory.zig").DeviceBuffer;
 pub const HostBuffer = @import("memory.zig").HostBuffer;
+pub const memory = @import("memory.zig");
 pub const usage = @import("memory.zig").usage;
 pub const Usage = @import("memory.zig").Usage;
 pub const carveout = @import("carveout.zig");
@@ -36,6 +37,8 @@ pub const embed_rows = @import("embed_rows.zig");
 pub const experts = @import("experts.zig");
 pub const sample = @import("sample.zig");
 pub const modelopt = @import("modelopt.zig");
+pub const graph_cache = @import("graph_cache.zig");
+pub const arena = @import("arena.zig");
 
 test {
     _ = @import("memory.zig");
@@ -54,4 +57,6 @@ test {
     _ = carveout;
     _ = @import("nvfp4.zig");
     _ = @import("modelopt.zig");
+    _ = graph_cache;
+    _ = arena;
 }

From 69e8a655adf0a37426f8b8f760c422734f140359 Mon Sep 17 00:00:00 2001
From: Jay Leaton <58495321+jayleaton@users.noreply.github.com>
Date: Sat, 10 Oct 2026 14:24:18 +0700
Subject: [PATCH 2/9] A failed CUDA kernel launch names its kernel and launch
 configuration; CUresult names fall back to a table when the driver cannot
 answer.

cuFuncGetName is optional (CUDA 12.3+), so an older driver still opens. During
the exit's teardown cuGetErrorName fails too, which printed
CUDA_ERROR_DEINITIALIZED as "CUDA_ERROR_UNKNOWN_CODE (4)"; code 4 now says the
process is exiting.
---
 zig/src/cuda/driver.zig | 67 ++++++++++++++++++++++++++++++++++++++---
 zig/src/cuda/launch.zig | 36 ++++++++++++++++++++--
 2 files changed, 96 insertions(+), 7 deletions(-)

diff --git a/zig/src/cuda/driver.zig b/zig/src/cuda/driver.zig
index 0a860eac92..88ee9050dc 100644
--- a/zig/src/cuda/driver.zig
+++ b/zig/src/cuda/driver.zig
@@ -8,6 +8,10 @@ pub const Error = error{ DriverUnavailable, MissingSymbol, CudaFailed, OutOfDevi
 pub const Driver = struct {
     lib: std.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("libcuda.so.1");
@@ -25,7 +29,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;
     }
@@ -38,18 +42,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 {
@@ -82,3 +101,43 @@ pub const Driver = struct {
         return @intCast(10 * major + minor);
     }
 };
+
+/// 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/launch.zig b/zig/src/cuda/launch.zig
index 56bba5272a..fe161fabf1 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,36 @@ 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;
+    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 {

From d7e89dff4fa4b3f4ac699a6cb2de51fe0e111bed Mon Sep 17 00:00:00 2001
From: Jonathan Regan-Beasley
 <197533392+jregan-beasley-smc@users.noreply.github.com>
Date: Sun, 11 Oct 2026 10:23:24 +1100
Subject: [PATCH 3/9] fix(cuda): correct RDMA VMM capability attribute

Use CUDA device attribute 110 for GPUDirect RDMA with VMM support. Attribute 126 reports the memory sync domain count and can incorrectly enable the RDMA allocation flag. Add a regression for the external ABI contract.
---
 zig/src/cuda/abi.zig | 9 ++++++++-
 1 file changed, 8 insertions(+), 1 deletion(-)

diff --git a/zig/src/cuda/abi.zig b/zig/src/cuda/abi.zig
index a099686c67..5418ce0109 100644
--- a/zig/src/cuda/abi.zig
+++ b/zig/src/cuda/abi.zig
@@ -40,7 +40,7 @@ pub const DeviceAttribute = enum(c_int) {
     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 = 126,
+    gpu_direct_rdma_with_cuda_vmm_supported = 110,
 };
 
 pub const FunctionAttribute = enum(c_int) {
@@ -243,3 +243,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,
+    ));
+}

From 5826f89bd31ce0f4d875051b61dec78cfd764ae7 Mon Sep 17 00:00:00 2001
From: Jay Leaton <58495321+jayleaton@users.noreply.github.com>
Date: Sun, 11 Oct 2026 17:29:31 +0700
Subject: [PATCH 4/9] Graph cache: a capture is freed until the map holds it,
 the keep floor stays under the max, a max of 0 is eager

- A capture whose agreement throws, or whose insert runs out of memory, no longer leaks its executable: the map's
  slot is reserved before the capture, and the executable is owned here until the insert, which cannot fail.
- A host map that cannot grow counts as no room in the existing room vote, so every rank declines that miss together.
- `min_keep` is clamped to `max`, so under the memory floor the cap never rises above the configured max. A max of 0
  runs every step eagerly, with no collective.
- A cap lowered to 0 under the floor goes back to max once every rank has room, instead of staying eager for good.
- Hits make no collective: with `Agree` set, keys, settings and fingerprint changes must already match on every rank.
  The docs now say so: the fingerprint changes at the same step on every rank, and `fingerprintOf` is for one rank.

Co-authored-by: Jonathan Regan-Beasley <197533392+jregan-beasley-smc@users.noreply.github.com>
---
 zig/src/cuda/graph_cache.zig                 |  31 ++--
 zig/src/cuda/graph_cache_regression_test.zig | 181 +++++++++++++++++++
 2 files changed, 201 insertions(+), 11 deletions(-)
 create mode 100644 zig/src/cuda/graph_cache_regression_test.zig

diff --git a/zig/src/cuda/graph_cache.zig b/zig/src/cuda/graph_cache.zig
index a2fb8e18bb..693f8ce25a 100644
--- a/zig/src/cuda/graph_cache.zig
+++ b/zig/src/cuda/graph_cache.zig
@@ -12,7 +12,7 @@ pub const Body = struct {
     run: *const fn (ctx: *anyopaque, stream: Stream) anyerror!void,
 };
 
-/// Every rank's verdict AND-ed, so ranks capture, evict and fall back at the same steps; never called in a capture.
+/// Every rank's verdict AND-ed on a miss (hits never call it), so ranks capture, evict and fall back together.
 pub const Agree = struct {
     ctx: *anyopaque,
     all: *const fn (ctx: *anyopaque, ok: bool) anyerror!bool,
@@ -53,7 +53,9 @@ pub fn Cache(comptime Key: type) type {
         stats: struct { replayed: u64 = 0, captured: u64 = 0, eager: u64 = 0, evicted: u64 = 0, dropped: u64 = 0, floor_hits: u64 = 0 } = .{},
 
         pub fn init(gpa: std.mem.Allocator, engine: Engine, settings: Settings) Self {
-            return .{ .gpa = gpa, .engine = engine, .settings = settings, .cap = settings.max };
+            var s = settings;
+            s.min_keep = @min(s.min_keep, s.max); // the floor never holds more than the configured max
+            return .{ .gpa = gpa, .engine = engine, .settings = s, .cap = s.max };
         }
 
         pub fn deinit(c: *Self) void {
@@ -98,9 +100,9 @@ pub fn Cache(comptime Key: type) type {
             return .eager;
         }
 
-        /// One step: the key's graph launched, captured first on a miss; a new `fingerprint` drops every graph first.
+        /// 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) return c.eager(stream, body);
+            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;
@@ -110,9 +112,10 @@ pub fn Cache(comptime Key: type) type {
                 c.stats.replayed += 1;
                 return .replayed;
             }
-            // a miss: room on every rank, else shed the least recent quarter and lower the cap (once below it, eager)
-            const room_here = if (c.room) |f| f() else true;
-            if (!try c.agreed(room_here and c.cap > 0)) {
+            // 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;
@@ -123,7 +126,7 @@ pub fn Cache(comptime Key: type) type {
                     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));
+                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);
@@ -138,13 +141,15 @@ pub fn Cache(comptime Key: type) type {
                 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)) {
-                if (exec) |x| c.engine.free(c.engine.ctx, x);
                 c.failed = true;
                 c.dropAll();
                 return c.eager(stream, body);
             }
-            try c.entries.put(c.gpa, key, .{ .exec = exec.?, .used = c.tick });
+            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;
@@ -196,7 +201,7 @@ pub const CudaEngine = struct {
     }
 };
 
-/// FNV-1a over the addresses a body captures (buffer bases, sizes): the fingerprint `run` compares.
+/// 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| {
@@ -390,3 +395,7 @@ test "graph cache: off runs every step eagerly" {
 test "graph cache: fingerprint" {
     try std.testing.expect(fingerprintOf(&.{ 1, 2 }) != fingerprintOf(&.{ 2, 1 }));
 }
+
+test {
+    _ = @import("graph_cache_regression_test.zig");
+}
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);
+}

From 5806e551cdcdcee9bff7cb90c596379826681d0e Mon Sep 17 00:00:00 2001
From: Jay Leaton <58495321+jayleaton@users.noreply.github.com>
Date: Sun, 11 Oct 2026 18:33:36 +0700
Subject: [PATCH 5/9] CUDA arena: the address maps are rehashed before
 tombstones fill them

The arena's `live` (small ranges) and `large` (own chunks) maps drop an entry on every free and gain one on every
allocation for the life of the process. std's tombstones never clear, so in time they take every free slot, and each
allocation, free and miss then probes the whole map. Each map now counts its removals and is rehashed in place once
they reach a tenth of its capacity.

Test: 20,000 random allocations and frees (small ranges and large chunks) keep a tenth of each map free at every step;
without the rehash it fails.
---
 zig/src/cuda/arena.zig | 43 ++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 43 insertions(+)

diff --git a/zig/src/cuda/arena.zig b/zig/src/cuda/arena.zig
index 2c80b7581b..3886c81403 100644
--- a/zig/src/cuda/arena.zig
+++ b/zig/src/cuda/arena.zig
@@ -36,6 +36,7 @@ pub const Arena = struct {
     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 {
@@ -114,16 +115,26 @@ pub const Arena = struct {
         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 {
@@ -329,3 +340,35 @@ test "arena: under DeviceBuffer.alloc through memory.setHook" {
     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());
+        }
+    }
+}

From d335520663539b54ee1cada6b32a599c3e2a2bee Mon Sep 17 00:00:00 2001
From: Jay Leaton <58495321+jayleaton@users.noreply.github.com>
Date: Wed, 14 Oct 2026 21:05:29 +0700
Subject: [PATCH 6/9] Graph cache: TF_GRAPH_CACHE=0 or off runs every step
 eagerly in any family that installs it, without a rebuild

Cache.init applies the operator's value through Settings.operator, so a deployment that measures graphs slower
can turn them off at serve time. The value is read once a cache, and is set alike on every rank like the
settings themselves: no agreement is added, and hits still make no collective.
---
 RUNBOOK.md                   |  2 ++
 zig/src/cuda/graph_cache.zig | 25 ++++++++++++++++++++++++-
 2 files changed, 26 insertions(+), 1 deletion(-)

diff --git a/RUNBOOK.md b/RUNBOOK.md
index 408f2d91e8..525cdf49e5 100644
--- a/RUNBOOK.md
+++ b/RUNBOOK.md
@@ -157,6 +157,8 @@ budget free for the next states.
 `TF_HEAT_HIGH` and `TF_HEAT_LOW`, in degrees Celsius and set together, make each prompt chunk wait while the hottest
 thermal zone (under `/sys/class/thermal`, or `TF_HEAT_ROOT`) is above the high band, until it is at or under the low one.
 The server logs `heat_wait_s`, a cancel ends the wait, and the reply's tokens do not change.
+`TF_GRAPH_CACHE=0` (or `off`) 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.
 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/graph_cache.zig b/zig/src/cuda/graph_cache.zig
index 693f8ce25a..9c62e4fc7d 100644
--- a/zig/src/cuda/graph_cache.zig
+++ b/zig/src/cuda/graph_cache.zig
@@ -31,6 +31,14 @@ pub const Settings = struct {
     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;
+    }
 };
 
 pub const Outcome = enum { replayed, captured, eager };
@@ -52,8 +60,10 @@ pub fn Cache(comptime Key: type) type {
         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`, unless TF_GRAPH_CACHE turns it off (set it alike on every rank, as `settings`).
         pub fn init(gpa: std.mem.Allocator, engine: Engine, settings: Settings) Self {
-            var s = settings;
+            var s = settings.operator(if (std.c.getenv("TF_GRAPH_CACHE")) |v| std.mem.span(v) else null);
+            if (!quiet and settings.on and !s.on) std.log.info("graph cache: off (TF_GRAPH_CACHE), every step eager", .{});
             s.min_keep = @min(s.min_keep, s.max); // the floor never holds more than the configured max
             return .{ .gpa = gpa, .engine = engine, .settings = s, .cap = s.max };
         }
@@ -392,6 +402,19 @@ test "graph cache: off runs every step eagerly" {
     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 }));
 }

From 864fe253277eef5909680c5d5ed42e1aacce6874 Mon Sep 17 00:00:00 2001
From: Jay Leaton <58495321+jayleaton@users.noreply.github.com>
Date: Fri, 16 Oct 2026 11:26:38 +0700
Subject: [PATCH 7/9] Graph cache: effective settings have one digest for ranks
 to compare at start; RUNBOOK says what a mismatch does

An off cache never calls Agree, so a rank left on waits at its first miss for a vote that never comes. Settings.effective
applies TF_GRAPH_CACHE, clamps min_keep and maps every off cache (0, off, OFF, a max of 0) to one value; Cache.init runs
with it, and Settings.digest gives ranks those bytes to compare before building a cache. Hits still make no collective.
---
 RUNBOOK.md                              |   6 +-
 zig/src/cuda/graph_cache.zig            |  34 +++++-
 zig/src/cuda/graph_cache_ranks_test.zig | 138 ++++++++++++++++++++++++
 3 files changed, 171 insertions(+), 7 deletions(-)
 create mode 100644 zig/src/cuda/graph_cache_ranks_test.zig

diff --git a/RUNBOOK.md b/RUNBOOK.md
index 525cdf49e5..4fd649e9be 100644
--- a/RUNBOOK.md
+++ b/RUNBOOK.md
@@ -157,8 +157,10 @@ budget free for the next states.
 `TF_HEAT_HIGH` and `TF_HEAT_LOW`, in degrees Celsius and set together, make each prompt chunk wait while the hottest
 thermal zone (under `/sys/class/thermal`, or `TF_HEAT_ROOT`) is above the high band, until it is at or under the low one.
 The server logs `heat_wait_s`, a cancel ends the wait, and the reply's tokens do not change.
-`TF_GRAPH_CACHE=0` (or `off`) 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.
+`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/graph_cache.zig b/zig/src/cuda/graph_cache.zig
index 9c62e4fc7d..1916b22b57 100644
--- a/zig/src/cuda/graph_cache.zig
+++ b/zig/src/cuda/graph_cache.zig
@@ -12,7 +12,7 @@ pub const Body = struct {
     run: *const fn (ctx: *anyopaque, stream: Stream) anyerror!void,
 };
 
-/// Every rank's verdict AND-ed on a miss (hits never call it), so ranks capture, evict and fall back together.
+/// 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,
@@ -39,6 +39,30 @@ pub const Settings = struct {
         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 };
@@ -60,11 +84,10 @@ pub fn Cache(comptime Key: type) type {
         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`, unless TF_GRAPH_CACHE turns it off (set it alike on every rank, as `settings`).
+        /// 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 {
-            var s = settings.operator(if (std.c.getenv("TF_GRAPH_CACHE")) |v| std.mem.span(v) else null);
-            if (!quiet and settings.on and !s.on) std.log.info("graph cache: off (TF_GRAPH_CACHE), every step eager", .{});
-            s.min_keep = @min(s.min_keep, s.max); // the floor never holds more than the configured max
+            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 };
         }
 
@@ -421,4 +444,5 @@ test "graph cache: fingerprint" {
 
 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..608ebe87a1
--- /dev/null
+++ b/zig/src/cuda/graph_cache_ranks_test.zig
@@ -0,0 +1,138 @@
+//! 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 {
+        var ts: std.c.timespec = undefined;
+        _ = std.c.clock_gettime(.MONOTONIC, &ts);
+        return @as(u64, @intCast(ts.sec)) * std.time.ns_per_s + @as(u64, @intCast(ts.nsec));
+    }
+};
+
+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);
+}

From 3aeebb97b18f3799c83069d48fd6a3f36f150824 Mon Sep 17 00:00:00 2001
From: Jay Leaton <58495321+jayleaton@users.noreply.github.com>
Date: Sat, 17 Oct 2026 09:19:52 +0700
Subject: [PATCH 8/9] CUDA arena: a reuse whose address-map insertion fails
 keeps its range free

allocLocked reserves room in `live` before it takes a range. A failed insertion used to pop a freed range
and then lose it: the range was in neither the free list nor the live map. Both small-path inserts now cannot fail.
A failure leaves the free list, the cursor and the chunks as they were, and the retry gets the same range.

The new test spends the map's spare room so that the reuse's insertion must grow it, fails that allocation,
and checks range, cursor and chunk conservation, the retry, and a clean deinit with every chunk released.

Not in this change: tying memory.setHook / currentHook installation to the family that owns the arena.
Jonathan Regan-Beasley's capture-ownership fix (5826f89b) is unchanged.

Spark receipt: lane job doug-635-arena-01 on dgx-02 (aarch64, Zig 0.17.0), 2026-10-17T02:10:06Z, VERDICT PASS,
candidate tree 663ff7bc2e2d0fa2eb404d5bc9d31eb55696bc68. zig fmt --check clean; lean 517/517 against merge-base
46a08c18, 0 new; zig build 12/12; test test-kernel-contract test-golden fuzz-cortheon 94/94 steps, 1430/1437
(7 skipped); test-cuda-host 39/39. With only the new test on abcad7d4, it fails (the freed range left its free list).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
---
 zig/src/cuda/arena.zig | 35 +++++++++++++++++++++++++++++++++--
 1 file changed, 33 insertions(+), 2 deletions(-)

diff --git a/zig/src/cuda/arena.zig b/zig/src/cuda/arena.zig
index 3886c81403..de3da5718c 100644
--- a/zig/src/cuda/arena.zig
+++ b/zig/src/cuda/arena.zig
@@ -86,8 +86,9 @@ pub const Arena = struct {
             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| {
-            try a.live.put(a.gpa, p, n);
+            a.live.putAssumeCapacityNoClobber(p, n);
             a.stats.reused += 1;
             return p;
         };
@@ -104,7 +105,7 @@ pub const Arena = struct {
             a.stats.chunk_bytes += cl;
         }
         const p = a.cursor;
-        try a.live.put(a.gpa, p, n);
+        a.live.putAssumeCapacityNoClobber(p, n);
         a.cursor += n;
         a.stats.small += 1;
         return p;
@@ -372,3 +373,33 @@ test "arena: a long run of allocations and frees leaves both address maps free s
         }
     }
 }
+
+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
+}

From adbdb8ab6ef873b75aea7cacddcf0f9e1f9d98f7 Mon Sep 17 00:00:00 2001
From: Jay Leaton <58495321+jayleaton@users.noreply.github.com>
Date: Sun, 18 Oct 2026 18:00:11 +0700
Subject: [PATCH 9/9] CUDA launches: a not-ready launch returns NotReady
 unlogged, as Driver.check does

launched() named the failing kernel for every non-success code, so a launch that main's tests answer with
CUDA_ERROR_NOT_READY (qwen_cuda_test, experts3_host_test) reached errorName on a fake driver and crashed.
It now returns error.NotReady first, like Driver.check. The ranks test reads the monotonic clock through
std.Io, so it also compiles on Windows hosts (zig-windows).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
---
 zig/src/cuda/graph_cache_ranks_test.zig | 4 +---
 zig/src/cuda/launch.zig                 | 1 +
 2 files changed, 2 insertions(+), 3 deletions(-)

diff --git a/zig/src/cuda/graph_cache_ranks_test.zig b/zig/src/cuda/graph_cache_ranks_test.zig
index 608ebe87a1..18b8a87e9d 100644
--- a/zig/src/cuda/graph_cache_ranks_test.zig
+++ b/zig/src/cuda/graph_cache_ranks_test.zig
@@ -53,9 +53,7 @@ const Board = struct {
     }
 
     fn now() u64 {
-        var ts: std.c.timespec = undefined;
-        _ = std.c.clock_gettime(.MONOTONIC, &ts);
-        return @as(u64, @intCast(ts.sec)) * std.time.ns_per_s + @as(u64, @intCast(ts.nsec));
+        return @intCast(std.Io.Clock.awake.now(std.testing.io).toNanoseconds());
     }
 };
 
diff --git a/zig/src/cuda/launch.zig b/zig/src/cuda/launch.zig
index fe161fabf1..47cbf720b6 100644
--- a/zig/src/cuda/launch.zig
+++ b/zig/src/cuda/launch.zig
@@ -111,6 +111,7 @@ pub fn launch(f: Function, cfg: Config, stream: Stream, args: *Args) Error!void
 /// 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) });