From 0211138b88ecab27e646509c39d080b6d51e58e6 Mon Sep 17 00:00:00 2001
From: ashhart <71890694+ashhart@users.noreply.github.com>
Date: Sun, 1 Nov 2026 18:59:46 +0000
Subject: [PATCH 1/3] Keep Kolibri CUDA prompts in shared KV pages

Use the existing TF_KV_SHARE adapter for global page references and
pool-owned sliding snapshots. Stage tree tables in active-chain order,
bind captured steps to stable slot tables, and keep the original linear
path when the switch is unset.

Add 128-dimensional prompt page addressing to the original row/replay
template, with replacement/reference/control cases and pending resource
records. Cover resumed and concurrent cache oracles, copy-draft bindings,
partial acceptance, cancellation and eviction with host regressions and
production-source faults. NVCC and checkpoint qualification remain pending.

Refs #1657
---
 docs/recipes/kolibri-1.md                     |   5 +
 zig/build/kolibri.zig                         |   1 +
 zig/build/kolibri_shared_host.zig             |   9 +
 zig/kernels/cuda/window_attention.cu          |  46 +-
 zig/kolibri_page_driver.zig                   |  90 ++++
 zig/kolibri_shared_host.zig                   | 425 ++++++++++++++++++
 .../core/contract/kolibri_pages_pending.zig   |  10 +
 zig/src/core/contract/registry.zig            |   2 +-
 zig/src/cuda/window_attention.zig             |  24 +-
 zig/src/families/kolibri1/backend.zig         |  86 ++--
 zig/src/families/kolibri1/forward.zig         |  77 +++-
 zig/src/families/kolibri1/kernels.zig         |   2 +-
 zig/src/families/kolibri1/native.zig          |  53 ++-
 zig/src/families/kolibri1/pages.zig           | 129 ++++++
 zig/src/families/kolibri1/shared.zig          | 169 +++++++
 zig/src/families/kolibri1/snap.zig            | 123 +++--
 zig/tests/cuda/contract_window.zig            |  53 ++-
 17 files changed, 1195 insertions(+), 109 deletions(-)
 create mode 100644 zig/build/kolibri_shared_host.zig
 create mode 100644 zig/kolibri_page_driver.zig
 create mode 100644 zig/kolibri_shared_host.zig
 create mode 100644 zig/src/core/contract/kolibri_pages_pending.zig
 create mode 100644 zig/src/families/kolibri1/pages.zig
 create mode 100644 zig/src/families/kolibri1/shared.zig

diff --git a/docs/recipes/kolibri-1.md b/docs/recipes/kolibri-1.md
index 9601290198..fb27e8b825 100644
--- a/docs/recipes/kolibri-1.md
+++ b/docs/recipes/kolibri-1.md
@@ -48,6 +48,11 @@ rounds run eager. Both give a row the same bits.
 - Refused at load: other quantizations, `norm_topk_prob`, other layer kinds, a config with a vision tower, shapes the
   shared kernels do not take, `--vision`, `--precision` and a captured kernel set (`TENSORFOLD_CUDA_KERNELS`).
 
+`TF_KV_SHARE=N` selects the experimental shared prompt store, with at most N kept states; unset or zero keeps the
+existing linear-cache path. Full-layer K/V is held as core-pool pages, while each snapshot keeps the sliding layers'
+last 512 K/V rows in pages of the same pool. Shared keeps and resumes use 256-token page edges. Copy drafts use
+the existing lane core; this does not add EAGLE-3 or a draft head. The new paged prompt kernels and complete shared-cache checkpoint path still require CUDA qualification.
+
 ## Checks
 
 ```bash
diff --git a/zig/build/kolibri.zig b/zig/build/kolibri.zig
index c868c5ba01..ba865154fb 100644
--- a/zig/build/kolibri.zig
+++ b/zig/build/kolibri.zig
@@ -9,6 +9,7 @@ pub fn host(b: *std.Build, m: *std.Build.Module, api: *std.Build.Module, all: *s
     const run = b.addRunArtifact(b.addTest(.{ .root_module = m }));
     b.step("test-kolibri-host", "Kolibri config, ownership and copy-draft metadata without a GPU").dependOn(&run.step);
     all.dependOn(&run.step);
+    @import("kolibri_shared_host.zig").host(b, m.import_table.get("cuda").?, m.import_table.get("core").?, m.import_table.get("lanes").?, api, m, all);
     b.step("test-kolibri-context-host", "Logical Kolibri context and physical verification reserve").dependOn(
         &b.addRunArtifact(b.addTest(.{ .root_module = m, .filters = &.{"Kolibri context:"} })).step,
     );
diff --git a/zig/build/kolibri_shared_host.zig b/zig/build/kolibri_shared_host.zig
new file mode 100644
index 0000000000..62ec1a0967
--- /dev/null
+++ b/zig/build/kolibri_shared_host.zig
@@ -0,0 +1,9 @@
+//! Real native prompt adapter, pool and Kolibri state copies under an inert CUDA driver.
+const std = @import("std");
+pub fn host(b: *std.Build, cuda: *std.Build.Module, core: *std.Build.Module, lanes: *std.Build.Module, api: *std.Build.Module, kolibri: *std.Build.Module, all: *std.Build.Step) void {
+    const adapter = b.createModule(.{ .root_source_file = b.path("zig/src/native/cuda_prompt_pages.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{ .{ .name = "core", .module = core }, .{ .name = "engine_api", .module = api } } });
+    const root = b.createModule(.{ .root_source_file = b.path("zig/kolibri_shared_host.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{ .{ .name = "cuda", .module = cuda }, .{ .name = "core", .module = core }, .{ .name = "lanes", .module = lanes }, .{ .name = "kolibri1", .module = kolibri }, .{ .name = "prompt_pages", .module = adapter } } });
+    const run = b.addRunArtifact(b.addTest(.{ .root_module = root, .filters = &.{"Kolibri shared:"} }));
+    b.step("test-kolibri-shared-host", "Kolibri paged resumes, shared streams and copy-draft cache bindings without a GPU").dependOn(&run.step);
+    all.dependOn(&run.step);
+}
diff --git a/zig/kernels/cuda/window_attention.cu b/zig/kernels/cuda/window_attention.cu
index a8754df06c..191f70bf06 100644
--- a/zig/kernels/cuda/window_attention.cu
+++ b/zig/kernels/cuda/window_attention.cu
@@ -8,6 +8,7 @@
 #include "attn_odd.cuh"
 #include "common/qmm_frag.cuh"
 #include "common/cp_async.cuh"
+#include "common/kv_pages.cuh"
 #include "window_split.cuh"
 
 // Prompt rows: prefill_attention.cu's tensor-core kernel with a window (keys (pos - win, pos]) and a ring of keys.
@@ -57,10 +58,10 @@ __device__ __forceinline__ float rowsum(const float (&c)[8]) {
 }
 
 // ODD: every NaN and infinite value leaves the mma, and the rows that see them add them back at float64's weights.
-template <int D, int WARPS, int HPC, int NS, bool ODD>
+template <int D, int WARPS, int HPC, int NS, bool ODD, typename... P>
 __device__ __forceinline__ bool wattn_rows(const __nv_bfloat16* __restrict__ q, const __nv_bfloat16* __restrict__ kc,
                                            const __nv_bfloat16* __restrict__ vc, __nv_bfloat16* __restrict__ out,
-                                           int p0, int W, int H, int HK, int G, float scale, int win, int ring) {
+                                           int p0, int W, int H, int HK, int G, float scale, int win, int ring, const P&... pg) {
     constexpr int C = D / 8, K16 = D / 16, NB = D / 8, RB = WARPS / HPC, THREADS = 32 * WARPS;
     constexpr bool QREG = D <= 128;               // the queries' fragments stay in registers, else in shared memory
     extern __shared__ uint4 wsmem[];
@@ -81,7 +82,13 @@ __device__ __forceinline__ bool wattn_rows(const __nv_bfloat16* __restrict__ q,
             const int key = key0 + i / C;
             const bool in = key < limit;             // past the cache's filled keys: zeros, never read
             const int row = in ? (ring > 0 ? key % ring : key) : 0; // a ring keeps key p at row p % ring
-            cp16z(dst + sw<D>(i / C, i % C), src + ((int64_t)row * HK + hk) * D + (i % C) * 8, in);
+            if constexpr (sizeof...(P) != 0) {
+                const tf_common::Pg& t = (pg, ...);
+                const int k = in ? key : 0;
+                const int64_t at = (int64_t)t.tab[k / t.tok] * t.page + ((int64_t)(k % t.tok) * HK + hk) * D;
+                cp16z(dst + sw<D>(i / C, i % C), src + at + (i % C) * 8, in);
+            } else
+                cp16z(dst + sw<D>(i / C, i % C), src + ((int64_t)row * HK + hk) * D + (i % C) * 8, in);
         }
     };
     uint32_t qa[QREG ? K16 : 1][4];
@@ -234,8 +241,14 @@ __device__ __forceinline__ bool wattn_rows(const __nv_bfloat16* __restrict__ q,
     if constexpr (!ODD) { // a NaN or infinite output: the block runs again as ODD (finite values: the same bits)
         if (__syncthreads_or(attn_odd::nonfinite_out<NB>(o, r0 < W, r1 < W))) return true;
     } else {
-        const auto row = [&](int k) { return vc + ((int64_t)(ring > 0 ? k % ring : k) * HK + hk) * D; };
-        const auto krow = [&](int k) { return kc + ((int64_t)(ring > 0 ? k % ring : k) * HK + hk) * D; };
+        const auto offset = [&](int k) -> int64_t {
+            if constexpr (sizeof...(P) != 0) {
+                const tf_common::Pg& t = (pg, ...);
+                return (int64_t)t.tab[k / t.tok] * t.page + ((int64_t)(k % t.tok) * HK + hk) * D;
+            } else return ((int64_t)(ring > 0 ? k % ring : k) * HK + hk) * D;
+        };
+        const auto row = [&](int k) { return vc + offset(k); };
+        const auto krow = [&](int k) { return kc + offset(k); };
         const auto lo = [&](int r) { return win > 0 ? max(pw + r - win + 1, 0) : 0; };
         const auto vis = [&](int r, int k) { return wrow + r < W && k >= lo(r) && k <= pw + r; };
         const auto each = [&](int r, auto f) { for (int k = lo(r); k <= pw + r; ++k) f(krow(k)); }; // row r's keys
@@ -256,11 +269,11 @@ __device__ __forceinline__ bool wattn_rows(const __nv_bfloat16* __restrict__ q,
 }
 
 // A block with a NaN or infinite output, run again with ODD on (out of the common path).
-template <int D, int WARPS, int HPC, int NS>
+template <int D, int WARPS, int HPC, int NS, typename... P>
 __device__ __noinline__ void wattn_odd(const __nv_bfloat16* q, const __nv_bfloat16* kc, const __nv_bfloat16* vc,
                                        __nv_bfloat16* out, int p0, int W, int H, int HK, int G, float scale, int win,
-                                       int ring) {
-    wattn_rows<D, WARPS, HPC, NS, true>(q, kc, vc, out, p0, W, H, HK, G, scale, win, ring);
+                                       int ring, const P&... pg) {
+    wattn_rows<D, WARPS, HPC, NS, true>(q, kc, vc, out, p0, W, H, HK, G, scale, win, ring, pg...);
 }
 
 template <int D, int WARPS, int HPC, int NS>
@@ -533,3 +546,20 @@ template __global__ void tf_window_prompt::wattn_kernel<128, 8, 4, 8>(const __nv
 template __global__ void tf_window_prompt::wattn_kernel<128, 8, 8, 8>(const __nv_bfloat16* __restrict__,
     const __nv_bfloat16* __restrict__, const __nv_bfloat16* __restrict__, __nv_bfloat16* __restrict__, int, int, int,
     int, int, float, int, int);
+
+// The same prompt rows and nonfinite replay, with only their cache addresses taken from pages.
+extern "C" __global__ void __launch_bounds__(256, 1) tf_window_prompt_128_h4_pg(
+    const __nv_bfloat16* q, const __nv_bfloat16* kc, const __nv_bfloat16* vc, __nv_bfloat16* out,
+    int p0, int W, int H, int HK, int G, float scale, int win, int ring, const int* tab, long long page, int tok) {
+    const tf_common::Pg pg{tab, page, tok, 0};
+    if (tf_window_prompt::wattn_rows<128, 8, 4, 8, false>(q, kc, vc, out, p0, W, H, HK, G, scale, win, ring, pg))
+        tf_window_prompt::wattn_odd<128, 8, 4, 8>(q, kc, vc, out, p0, W, H, HK, G, scale, win, ring, pg);
+}
+
+extern "C" __global__ void __launch_bounds__(256, 1) tf_window_prompt_128_h8_pg(
+    const __nv_bfloat16* q, const __nv_bfloat16* kc, const __nv_bfloat16* vc, __nv_bfloat16* out,
+    int p0, int W, int H, int HK, int G, float scale, int win, int ring, const int* tab, long long page, int tok) {
+    const tf_common::Pg pg{tab, page, tok, 0};
+    if (tf_window_prompt::wattn_rows<128, 8, 8, 8, false>(q, kc, vc, out, p0, W, H, HK, G, scale, win, ring, pg))
+        tf_window_prompt::wattn_odd<128, 8, 8, 8>(q, kc, vc, out, p0, W, H, HK, G, scale, win, ring, pg);
+}
diff --git a/zig/kolibri_page_driver.zig b/zig/kolibri_page_driver.zig
new file mode 100644
index 0000000000..a9fb649cb6
--- /dev/null
+++ b/zig/kolibri_page_driver.zig
@@ -0,0 +1,90 @@
+//! Inert CUDA allocations and ordered copies; no GPU or kernel arithmetic is executed.
+const std = @import("std");
+const cuda = @import("cuda");
+const a = std.testing.allocator;
+const Entry = struct { ptr: u64 = 0, len: usize = 0 };
+var blocks: [512]Entry = @splat(.{});
+pub var skip_copy = false;
+pub var allocations: usize = 0;
+pub var graphs_dropped: usize = 0;
+
+fn alloc(out: *u64, n: usize) callconv(.c) cuda.abi.Result {
+    const mem = a.alignedAlloc(u8, .fromByteUnits(256), n) catch return 2;
+    @memset(mem, 0);
+    for (&blocks) |*b| if (b.ptr == 0) {
+        b.* = .{ .ptr = @intFromPtr(mem.ptr), .len = n };
+        allocations += 1;
+        out.* = b.ptr;
+        return 0;
+    };
+    a.free(mem);
+    return 2;
+}
+fn free(ptr: u64) callconv(.c) cuda.abi.Result {
+    for (&blocks) |*b| if (b.ptr == ptr) {
+        const mem: [*]align(256) u8 = @ptrFromInt(ptr);
+        a.free(mem[0..b.len]);
+        b.* = .{};
+        return 0;
+    };
+    return 1;
+}
+fn h2d(dst: u64, src: ?*const anyopaque, n: usize, _: cuda.abi.Stream) callconv(.c) cuda.abi.Result {
+    const out: [*]u8 = @ptrFromInt(dst);
+    const in: [*]const u8 = @ptrCast(src.?);
+    @memcpy(out[0..n], in[0..n]);
+    return 0;
+}
+fn d2d(dst: u64, src: u64, n: usize, _: cuda.abi.Stream) callconv(.c) cuda.abi.Result {
+    if (!skip_copy) {
+        const out: [*]u8 = @ptrFromInt(dst);
+        const in: [*]const u8 = @ptrFromInt(src);
+        @memcpy(out[0..n], in[0..n]);
+    }
+    return 0;
+}
+fn fill(dst: u64, value: u8, n: usize, _: cuda.abi.Stream) callconv(.c) cuda.abi.Result {
+    const out: [*]u8 = @ptrFromInt(dst);
+    @memset(out[0..n], value);
+    return 0;
+}
+fn sync(_: cuda.abi.Stream) callconv(.c) cuda.abi.Result {
+    return 0;
+}
+fn contextSync() callconv(.c) cuda.abi.Result {
+    return 0;
+}
+fn dropGraph(_: cuda.abi.GraphExec) callconv(.c) cuda.abi.Result {
+    graphs_dropped += 1;
+    return 0;
+}
+
+pub fn driver() cuda.Driver {
+    var api: cuda.abi.Api = undefined;
+    api.cuMemAlloc_v2 = alloc;
+    api.cuMemFree_v2 = free;
+    api.cuMemcpyHtoDAsync_v2 = h2d;
+    api.cuMemcpyDtoDAsync_v2 = d2d;
+    api.cuMemsetD8Async = fill;
+    api.cuStreamSynchronize = sync;
+    api.cuCtxSynchronize = contextSync;
+    api.cuGraphExecDestroy = dropGraph;
+    api.cuLaunchKernel = launch;
+    return .{ .lib = undefined, .api = api };
+}
+
+pub fn empty() bool {
+    for (blocks) |b| if (b.ptr != 0) return false;
+    return true;
+}
+
+pub var last_function: usize = 0;
+pub var last_table: u64 = 0;
+fn launch(fun: cuda.abi.Function, _: c_uint, _: c_uint, _: c_uint, _: c_uint, _: c_uint, _: c_uint, _: c_uint, _: cuda.abi.Stream, args: ?[*]?*anyopaque, _: ?[*]?*anyopaque) callconv(.c) cuda.abi.Result {
+    last_function = @intFromPtr(fun);
+    if (last_function >= 3) {
+        const ptr: *const u64 = @ptrCast(@alignCast(args.?[12].?));
+        last_table = ptr.*;
+    } else last_table = 0;
+    return 0;
+}
diff --git a/zig/kolibri_shared_host.zig b/zig/kolibri_shared_host.zig
new file mode 100644
index 0000000000..b4e02b0b14
--- /dev/null
+++ b/zig/kolibri_shared_host.zig
@@ -0,0 +1,425 @@
+//! Real Kolibri leases, snapshots and prompt adapter with a deterministic row oracle, not checkpoint arithmetic.
+const std = @import("std");
+const k1 = @import("kolibri1");
+const cuda = @import("cuda");
+const lanes = @import("lanes");
+const adapter = @import("prompt_pages");
+const fake = @import("kolibri_page_driver.zig");
+const a = std.testing.allocator;
+const t = std.testing;
+const fwd = k1.forward;
+const be = k1.backend;
+const marks = [_]u32{ 8192, 8448, 8704, 8960 };
+const config: k1.Config = blk: {
+    var c: k1.Config = .{ .layers = 5, .hidden = 128, .heads = 12, .kv_heads = 1, .head_dim = 128, .vocab = 256, .experts = 2, .top_k = 1, .moe_width = 128, .shared_width = 128, .window = 513, .rope_theta = 1e4, .eps = 1e-6 };
+    c.full[4] = true;
+    break :blk c;
+};
+
+const Fixture = struct {
+    d: cuda.Driver,
+    m: fwd.Model,
+    g: be.Kolibri,
+    k: k1.kernels.Kernels = undefined,
+    store: ?*adapter.Adapter = null,
+
+    fn init(f: *Fixture, paged: bool, share: bool) !void {
+        return f.initWith(paged, share);
+    }
+    fn initWith(f: *Fixture, paged: bool, share: bool) !void {
+        f.store = null;
+        f.d = fake.driver();
+        const queue: cuda.Stream = .{ .d = &f.d, .handle = @ptrFromInt(1) };
+        f.m = .{ .gpa = a, .d = &f.d, .w = undefined, .c = config, .context = 16384, .slots = if (paged) 1 else 2, .ring = fwd.ringSize(fwd.prompt_rows, config.window), .max_rows = fwd.prompt_rows, .groups = 33, .k = &.{}, .v = &.{}, .norms = &.{}, .cos = 0, .sin = 0, .origin = 0, .s = undefined, .paged = paged };
+        errdefer f.m.deinit();
+        f.m.k = try a.alloc(u64, 5);
+        f.m.v = try a.alloc(u64, 5);
+        try fwd.pages.build(&f.m, queue);
+        f.g = try be.Kolibri.init(a, &f.d, &f.m, &f.k, queue);
+        errdefer f.g.deinit();
+        if (paged) {
+            const dropped = fake.graphs_dropped;
+            f.g.steps[0] = .{ .exec = .{ .d = &f.d, .handle = @ptrFromInt(123) }, .slot = 0, .flat = 0, .logits = 0 };
+            try be.shared.attach(&f.g, 1, 128);
+            try t.expectEqual(dropped + 1, fake.graphs_dropped);
+            var prng = std.Random.DefaultPrng.init(0x7a92);
+            const pool = &f.g.pool.?;
+            prng.random().shuffle(u32, pool.free[0..pool.nfree]);
+        }
+        if (share) f.store = (try adapter.Adapter.open(k1.native, a, &f.g, fwd.pages.tokens, 8)).?;
+    }
+    fn deinit(f: *Fixture) void {
+        if (f.store) |x| x.close();
+        f.g.deinit();
+        f.m.deinit();
+    }
+    fn release(f: *Fixture, s: *lanes.Stream) void {
+        const b = f.g.backend();
+        b.vtable.release(b.ptr, s);
+    }
+};
+
+fn stream(ids: []const u32) !lanes.Stream {
+    return lanes.Stream.init(a, .{ .id = "fixture", .prompt = ids, .max_new = 16, .drafts = false, .shared_prefixes = &marks });
+}
+
+fn prompt() ![]u32 {
+    const out = try a.alloc(u32, 9003);
+    for (out, 0..) |*x, i| x.* = @intCast(i % 97 + 3);
+    return out;
+}
+
+fn address(m: *const fwd.Model, cache: []const u64, layer: usize, slot: usize, position: usize) u64 {
+    const rb = m.c.kv_heads * m.c.head_dim * 2;
+    const row = if (m.c.full[layer] and m.paged) @as(usize, m.ids[slot * fwd.pages.perSlot(m) + position / fwd.pages.tokens]) * fwd.pages.tokens + position % fwd.pages.tokens else slot * (if (m.c.full[layer]) m.context else m.ring) + position % (if (m.c.full[layer]) m.context else m.ring);
+    return cache[layer] + row * rb;
+}
+
+fn value(token: u32, position: usize, layer: usize, v: bool) u16 {
+    return @truncate(token *% 73 +% @as(u32, @truncate(position)) *% 17 +% @as(u32, @intCast(layer)) *% 101 +% @as(u32, @intFromBool(v)) *% 19);
+}
+
+/// Emulates only row writes at the device boundary; the production store performs every keep, map and restore.
+fn fill(f: *Fixture, s: *lanes.Stream) !u64 {
+    return finish(f, s, try f.g.opening(s));
+}
+
+fn finish(f: *Fixture, s: *lanes.Stream, l: anytype) !u64 {
+    while (l.fill) |start| {
+        var end = @min(s.prompt().len, (start / fwd.prompt_rows + 1) * fwd.prompt_rows);
+        for (l.marks) |mark| if (mark > start and mark < end) {
+            end = mark;
+        };
+        for (start..end) |pos| for (0..f.m.c.layers) |i| for ([_][]u64{ f.m.k, f.m.v }, 0..) |cache, v| {
+            const row: [*]u16 = @ptrFromInt(address(&f.m, cache, i, l.slot, pos));
+            @memset(row[0 .. f.m.c.kv_heads * f.m.c.head_dim], value(s.prompt()[pos], pos, i, v == 1));
+        };
+        l.fill = if (end == s.prompt().len) null else end;
+        l.len = end;
+        s.fill_at.store(@intCast(end), .monotonic);
+        f.g.marked(s, end);
+    }
+    return digest(&f.m, l.slot, s.prompt().len);
+}
+
+/// A reply oracle depends on every global row and every live sliding-tail byte, separately for each slot.
+fn digest(m: *const fwd.Model, slot: usize, at: usize) u64 {
+    var hash = std.hash.Wyhash.init(0);
+    for (0..m.c.layers) |i| {
+        const first = if (m.c.full[i]) 0 else at - @min(at, m.c.window - 1);
+        for (first..at) |pos| for ([_][]u64{ m.k, m.v }) |cache| {
+            const row: [*]const u8 = @ptrFromInt(address(m, cache, i, slot, pos));
+            hash.update(row[0 .. m.c.kv_heads * m.c.head_dim * 2]);
+        };
+    }
+    return hash.final();
+}
+
+fn reply(hash: u64) [16]u32 {
+    var prng = std.Random.DefaultPrng.init(hash);
+    var out: [16]u32 = undefined;
+    for (&out) |*x| x.* = prng.random().uintLessThan(u32, config.vocab);
+    return out;
+}
+
+fn idle(f: *Fixture) !void {
+    if (f.g.pool) |*p| try t.expect(p.idle());
+    try t.expect(f.g.lanes.count() == 0);
+}
+
+const Resume = struct { fresh: [16]u32, resumed: [16]u32, cached: u64 };
+fn resumed(skip_tail: bool) !Resume {
+    var f: Fixture = undefined;
+    try f.init(true, true);
+    defer f.deinit();
+    const ids = try prompt();
+    defer a.free(ids);
+    var teach = try stream(ids);
+    defer teach.deinit(a);
+    const cold = reply(try fill(&f, &teach));
+    f.release(&teach);
+    var blocker = try stream(&.{ 99, 98, 97 });
+    defer blocker.deinit(a);
+    _ = try fill(&f, &blocker); // slot 0 holds another request while the resumed pass takes slot 1
+    defer f.release(&blocker);
+    var s = try stream(ids);
+    defer s.deinit(a);
+    fake.skip_copy = skip_tail;
+    _ = try f.g.opening(&s);
+    fake.skip_copy = false;
+    defer fake.skip_copy = false;
+    const got = reply(try finish(&f, &s, f.g.lanes.getPtr(&s).?));
+    f.release(&s);
+    return .{ .fresh = cold, .resumed = got, .cached = s.cached };
+}
+
+test "Kolibri shared: resumed reply restores wrapped sliding tails with shared global pages" {
+    const out = try resumed(false);
+    try t.expectEqual(@as(u64, 8960), out.cached);
+    try t.expectEqualSlices(u32, &out.fresh, &out.resumed);
+    const bad = try resumed(true);
+    try t.expectEqual(@as(u64, 8960), bad.cached);
+    try t.expect(!std.mem.eql(u32, &bad.fresh, &bad.resumed));
+    try t.expect(fake.empty());
+}
+
+test "Kolibri shared: concurrent continuations preserve solo replies and a swapped-slot control fails" {
+    const ids = try prompt();
+    defer a.free(ids);
+    const other = try a.dupe(u32, ids);
+    defer a.free(other);
+    other[8990] = 77;
+    var plain: Fixture = undefined;
+    try plain.init(false, false);
+    var sa = try stream(ids);
+    defer sa.deinit(a);
+    const oracle_a = reply(try fill(&plain, &sa));
+    plain.release(&sa);
+    var sb = try stream(other);
+    defer sb.deinit(a);
+    const oracle_b = reply(try fill(&plain, &sb));
+    plain.release(&sb);
+    plain.deinit();
+    var f: Fixture = undefined;
+    try f.init(true, true);
+    defer f.deinit();
+    var a1 = try stream(ids);
+    defer a1.deinit(a);
+    const r1 = reply(try fill(&f, &a1));
+    var b1 = try stream(other);
+    defer b1.deinit(a);
+    const r2 = reply(try fill(&f, &b1));
+    try t.expectEqualSlices(u32, &oracle_a, &r1);
+    try t.expectEqualSlices(u32, &oracle_b, &r2);
+    try t.expect(b1.cached > 0);
+    const la = f.g.lanes.getPtr(&a1).?;
+    const lb = f.g.lanes.getPtr(&b1).?;
+    const page = 8990 / fwd.pages.tokens;
+    const wrong = f.m.ids[la.slot * fwd.pages.perSlot(&f.m) + page];
+    try fwd.pages.map(&f.m, f.g.s, lb.slot, &.{wrong}, page);
+    const bad = reply(digest(&f.m, lb.slot, other.len));
+    try t.expect(!std.mem.eql(u32, &oracle_b, &bad));
+    f.release(&a1);
+    f.release(&b1);
+    try idle(&f);
+    f.store.?.close();
+    f.store = null;
+    try t.expectEqual(@as(usize, 128), f.g.pool.?.available());
+}
+
+/// Read the same chain-order metadata and page bindings passed to the real tree operator.
+fn treeReply(f: *Fixture, chains: []const fwd.Chain, graph_slot: ?usize) ![16]u32 {
+    const windows = try a.alloc(k1.tree.Window, chains.len);
+    defer a.free(windows);
+    for (chains, windows) |ch, *w| w.* = .{ .committed = ch.p0, .rows = ch.tokens.len };
+    const plan = try k1.tree.build(a, windows);
+    defer a.free(plan.flat);
+    const flat = try f.m.alloc(plan.flat.len * 4);
+    try f.m.upload(f.g.s, flat, std.mem.sliceAsBytes(plan.flat));
+    const offs = try f.m.alloc(f.m.c.layers * chains.len * 16);
+    try f.m.offsets(f.g.s, chains, offs);
+    try fwd.pages.chains(&f.m, f.g.s, chains);
+    const rows = f.m.treeRows(flat, plan, 0, 0, offs, graph_slot);
+    if (f.m.paged) try t.expect(rows.pages != null);
+    var hash = std.hash.Wyhash.init(0);
+    for (chains, 0..) |ch, j| for (0..f.m.c.layers) |i| {
+        const first = if (f.m.c.full[i]) 0 else ch.p0 - @min(ch.p0, f.m.c.window - 1);
+        for (first..ch.p0) |pos| for ([_][]u64{ f.m.k, f.m.v }, 0..) |cache, v| {
+            const ptr = if (f.m.c.full[i]) blk: {
+                const off: [*]const i64 = @ptrFromInt(rows.offs);
+                const base: u64 = @bitCast(@as(i64, @intCast(rows.origin)) + off[(i * chains.len + j) * 2 + v] * 2);
+                const rb = f.m.c.kv_heads * f.m.c.head_dim * 2;
+                if (rows.pages) |pg| {
+                    const table: [*]const u32 = @ptrFromInt(pg.table);
+                    break :blk base + (@as(usize, table[j * pg.stride + pos / pg.tok]) * pg.page + pos % pg.tok * rb / 2) * 2;
+                }
+                break :blk base + pos * rb;
+            } else address(&f.m, cache, i, ch.slot, pos);
+            const bytes: [*]const u8 = @ptrFromInt(ptr);
+            hash.update(bytes[0 .. f.m.c.kv_heads * f.m.c.head_dim * 2]);
+        };
+        // Each chain's staged path is private; its proposed rows never replace another chain's keys.
+        for (ch.tokens, 0..) |token, r| for (0..2) |v| {
+            var data: [128]u16 = undefined;
+            @memset(&data, value(token, ch.p0 + r, i, v == 1));
+            hash.update(std.mem.sliceAsBytes(&data));
+        };
+    };
+    return reply(hash.final());
+}
+
+fn extend(f: *Fixture, slot: usize, from: usize, ids: []const u32) void {
+    for (ids, from..) |token, pos| for (0..f.m.c.layers) |i| for ([_][]u64{ f.m.k, f.m.v }, 0..) |cache, v| {
+        const row: [*]u16 = @ptrFromInt(address(&f.m, cache, i, slot, pos));
+        @memset(row[0..128], value(token, pos, i, v == 1));
+    };
+}
+
+test "Kolibri shared: copy-draft tree bindings preserve plain replies in reversed chain order and after rejection" {
+    const ids = try prompt();
+    defer a.free(ids);
+    const other = try a.dupe(u32, ids);
+    defer a.free(other);
+    other[8990] = 77;
+    var f: Fixture = undefined;
+    try f.init(true, true);
+    defer f.deinit();
+    var s = try stream(ids);
+    defer s.deinit(a);
+    _ = try fill(&f, &s);
+    var peer = try stream(other);
+    defer peer.deinit(a);
+    _ = try fill(&f, &peer);
+    defer f.release(&peer);
+    const l = f.g.lanes.getPtr(&s).?;
+    const p = f.g.lanes.getPtr(&peer).?;
+    const proposed = [_]u32{ 4, 5, 88, 99 };
+    const chains = [_]fwd.Chain{ .{ .slot = p.slot, .p0 = ids.len, .tokens = &proposed }, .{ .slot = l.slot, .p0 = ids.len, .tokens = &.{ 6, 7 } } };
+    const got = try treeReply(&f, &chains, null);
+    const solo = try treeReply(&f, chains[0..1], p.slot);
+    try t.expectEqual(f.m.table + p.slot * fwd.pages.perSlot(&f.m) * 4, fwd.pages.treeView(&f.m, p.slot).?.table);
+    var plain: Fixture = undefined;
+    try plain.init(false, false);
+    defer plain.deinit();
+    var ps = try stream(ids);
+    defer ps.deinit(a);
+    _ = try fill(&plain, &ps);
+    var pp = try stream(other);
+    defer pp.deinit(a);
+    _ = try fill(&plain, &pp);
+    const fresh_chains = [_]fwd.Chain{ .{ .slot = 1, .p0 = ids.len, .tokens = &proposed }, .{ .slot = 0, .p0 = ids.len, .tokens = &.{ 6, 7 } } };
+    try t.expectEqualSlices(u32, &try treeReply(&plain, &fresh_chains, null), &got);
+    try t.expectEqualSlices(u32, &try treeReply(&plain, fresh_chains[0..1], 1), &solo);
+    plain.release(&ps);
+    plain.release(&pp);
+    // The real keep hook discards two speculative rows. Subsequent writes replace exactly those positions.
+    l.base = ids.len;
+    try l.rows.appendSlice(a, &proposed);
+    extend(&f, l.slot, ids.len, &proposed);
+    l.len = ids.len + proposed.len;
+    const backend = f.g.backend();
+    const positions = [_]u64{ ids.len + 1, ids.len + 2, ids.len + 3, ids.len + 4 };
+    const win: lanes.backend.Window = .{ .stream = &s, .pending = proposed[0], .tokens = proposed[1..], .positions = &positions, .held = 0, .parents = null };
+    try backend.vtable.keep(backend.ptr, &.{win}, &.{&.{ 0, 1 }});
+    try t.expectEqual(@as(u64, ids.len + 2), l.len);
+    extend(&f, l.slot, l.len, &.{8});
+    extend(&plain, 0, ids.len, &.{ 4, 5, 8 });
+    try t.expectEqual(reply(digest(&plain.m, 0, ids.len + 3)), reply(digest(&f.m, l.slot, ids.len + 3)));
+    f.release(&s);
+}
+
+fn cancelled(_: *anyopaque) bool {
+    return true;
+}
+
+test "Kolibri shared: cancellation returns leases and page coverage fails before launch" {
+    var f: Fixture = undefined;
+    try f.init(true, true);
+    defer f.deinit();
+    try t.expectEqual(@as(u32, 64), f.g.facts().max_streams);
+    var flag: u8 = 0;
+    var s = try lanes.Stream.init(a, .{ .id = "cancel", .prompt = &.{ 1, 2, 3 }, .max_new = 16, .cancel_check = .{ .ptr = &flag, .check = cancelled } });
+    defer s.deinit(a);
+    const backend = f.g.backend();
+    try t.expectError(error.Cancelled, backend.vtable.prefill_step.?(backend.ptr, &s, 8192));
+    try idle(&f);
+    try t.expectEqual(@as(usize, 128), f.g.pool.?.available());
+    try t.expectError(error.PastPages, f.m.forward(.{ .k = &f.k, .s = f.g.s }, &.{.{ .slot = 0, .p0 = 0, .tokens = &.{1} }}, true, null, 0));
+}
+
+test "Kolibri shared: snapshots and fifty-layer geometry are charged to the pool" {
+    var f: Fixture = undefined;
+    try f.init(true, true);
+    defer f.deinit();
+    try t.expectEqual(@as(u64, 1048576), k1.snap.copyBytes(&f.m, 8960, 8960));
+    try t.expectEqual(@as(usize, 8), k1.snap.blobPages(&f.m, 8960, 8960));
+    const ids = try prompt();
+    defer a.free(ids);
+    var s = try stream(ids);
+    defer s.deinit(a);
+    const before = fake.allocations;
+    _ = try fill(&f, &s);
+    try t.expectEqual(before, fake.allocations);
+    f.release(&s);
+    try idle(&f);
+    f.store.?.close();
+    f.store = null;
+    try t.expectEqual(@as(usize, 128), f.g.pool.?.available());
+    var full = f.m;
+    full.c.layers = 50;
+    full.c.kv_heads = 4;
+    for (0..50) |i| full.c.full[i] = i % 5 == 4;
+    full.context = 33792;
+    try t.expectEqual(@as(u64, 5242880), fwd.pages.pageBytes(&full));
+    try t.expectEqual(@as(u64, 718275088), fwd.pages.layout(&full).fixed_bytes);
+    try t.expectEqual(@as(u64, 41943040), k1.snap.copyBytes(&full, 8960, 8960));
+    try t.expectEqual(@as(usize, 8), k1.snap.blobPages(&full, 8960, 8960));
+}
+
+test "Kolibri shared: 128-dimensional prompt dispatch keeps linear functions and binds paged functions" {
+    const d = fake.driver();
+    const queue: cuda.Stream = .{ .d = &d, .handle = @ptrFromInt(1) };
+    var kernel: cuda.window_attention.WindowAttention = undefined;
+    kernel.prompt_fns = .{ .{ .d = &d, .handle = @ptrFromInt(1) }, .{ .d = &d, .handle = @ptrFromInt(2) } };
+    kernel.prompt_pages128 = .{ .{ .d = &d, .handle = @ptrFromInt(3) }, .{ .d = &d, .handle = @ptrFromInt(4) } };
+    var w: cuda.window_attention.Window = .{ .h = 48, .hk = 4, .d = 128, .win = 0, .ring = 0, .scale = 1 };
+    try kernel.prompt(queue, 0, 0, 0, 0, 0, 32, w);
+    try t.expectEqual(@as(usize, 1), fake.last_function);
+    w.pg = .{ .table = 42, .page_elems = 256 * 4 * 128, .tokens = 256 };
+    try kernel.prompt(queue, 0, 0, 0, 0, 0, 32, w);
+    try t.expectEqual(@as(usize, 3), fake.last_function);
+    try t.expectEqual(@as(u64, 42), fake.last_table);
+    w.h = 32;
+    try kernel.prompt(queue, 0, 0, 0, 0, 0, 32, w);
+    try t.expectEqual(@as(usize, 4), fake.last_function);
+    w.pg = null;
+    try kernel.prompt(queue, 0, 0, 0, 0, 0, 32, w);
+    try t.expectEqual(@as(usize, 2), fake.last_function);
+}
+
+test "Kolibri shared: ring restore precedes forced snapshot eviction in reservation" {
+    var f: Fixture = undefined;
+    try f.init(true, true);
+    defer f.deinit();
+    const ids = try prompt();
+    defer a.free(ids);
+    var teach = try stream(ids);
+    defer teach.deinit(a);
+    const expected = reply(try fill(&f, &teach));
+    f.release(&teach);
+    f.store.?.store.limit(.{ .pages = 128, .snaps = 1 });
+    try t.expectEqual(@as(u32, 8960), f.store.?.store.find(ids).?.end);
+    var blocker = try stream(&.{ 99, 98, 97 });
+    defer blocker.deinit(a);
+    _ = try fill(&f, &blocker);
+    defer f.release(&blocker);
+    const pool = &f.g.pool.?;
+    const held = try a.alloc(u32, pool.available());
+    defer a.free(held);
+    try t.expect(pool.hold(held));
+    defer for (held) |id| pool.release(id);
+    const Reclaim = struct {
+        f: *Fixture,
+        ids: []const u32,
+        restored_first: bool = false,
+        fn call(ptr: *anyopaque, want: usize) bool {
+            const x: *@This() = @ptrCast(@alignCast(ptr));
+            x.restored_first = x.f.g.restored_at == 8960;
+            x.f.store.?.store.forget(x.ids, 8960);
+            return x.f.g.pool.?.available() >= want;
+        }
+    };
+    var reclaim: Reclaim = .{ .f = &f, .ids = ids };
+    const original = pool.reclaim;
+    pool.reclaim = .{ .ctx = &reclaim, .call = Reclaim.call };
+    defer pool.reclaim = original;
+    var s = try stream(ids);
+    defer s.deinit(a);
+    const l = try f.g.opening(&s);
+    defer f.release(&s);
+    try t.expect(reclaim.restored_first);
+    try t.expectEqual(@as(u64, 8960), s.cached);
+    try t.expect(f.store.?.store.find(ids) == null); // its blob was reclaimed by the new lane's promise
+    const got = reply(try finish(&f, &s, l));
+    try t.expectEqualSlices(u32, &expected, &got);
+}
diff --git a/zig/src/core/contract/kolibri_pages_pending.zig b/zig/src/core/contract/kolibri_pages_pending.zig
new file mode 100644
index 0000000000..ae4c358842
--- /dev/null
+++ b/zig/src/core/contract/kolibri_pages_pending.zig
@@ -0,0 +1,10 @@
+//! Paged 128-dimensional prompt launches; resource measurements remain pending for every compiler and SM.
+const registry = @import("registry.zig");
+pub const records = blk: {
+    var out: [2]registry.CudaRecord = undefined;
+    for ([_]u32{ 4, 8 }, 0..) |heads, i| {
+        const symbol = @import("std").fmt.comptimePrint("tf_window_prompt_128_h{d}_pg", .{heads});
+        out[i] = .{ .launch = .{ .name = symbol, .symbol = symbol, .image = "window_attention", .threads = 256, .dynamic = 65536, .opt_in = true }, .usage = registry.none };
+    }
+    break :blk out;
+};
diff --git a/zig/src/core/contract/registry.zig b/zig/src/core/contract/registry.zig
index e6bb1f89ed..8305d6289f 100644
--- a/zig/src/core/contract/registry.zig
+++ b/zig/src/core/contract/registry.zig
@@ -575,7 +575,7 @@ const learning = @import("registry_learning.zig");
 pub const gradient_pending = learning.gradient_pending;
 pub const qwen_learning_pending = learning.qwen_learning_pending;
 pub const glm_pending = @import("glm_sampling_pending.zig");
-pub const landed_cuda = cuda ++ @import("dory_cuda_records.zig").cuda ++ gradient_pending ++ qwen_learning_pending ++ paging.extra ++ @import("qwen_served_records.zig").records ++ cuda_pending ++ @import("qwen_decode_pending.zig").records ++ glm_pending.sampling ++ @import("glm_kv_pending.zig").records ++ @import("exl3_grouped_records.zig").records ++ @import("proposal_sampling_records.zig").records ++ @import("flashnext_pages_pending.zig").records ++ @import("glm_f32w_pending.zig").records ++ @import("penalty_sampling_records.zig").records;
+pub const landed_cuda = @import("kolibri_pages_pending.zig").records ++ cuda ++ @import("dory_cuda_records.zig").cuda ++ gradient_pending ++ qwen_learning_pending ++ paging.extra ++ @import("qwen_served_records.zig").records ++ cuda_pending ++ @import("qwen_decode_pending.zig").records ++ glm_pending.sampling ++ @import("glm_kv_pending.zig").records ++ @import("exl3_grouped_records.zig").records ++ @import("proposal_sampling_records.zig").records ++ @import("flashnext_pages_pending.zig").records ++ @import("glm_f32w_pending.zig").records ++ @import("penalty_sampling_records.zig").records;
 test {
     _ = learning;
 }
diff --git a/zig/src/cuda/window_attention.zig b/zig/src/cuda/window_attention.zig
index 1d4061b99c..b368f6abf3 100644
--- a/zig/src/cuda/window_attention.zig
+++ b/zig/src/cuda/window_attention.zig
@@ -10,7 +10,7 @@ const Module = module.Module;
 const Function = module.Function;
 const Stream = stream_.Stream;
 
-/// A Window's `pg` (win 0, head dims 256, 512): caches name a layer in page 0; `table` rows of ids, `stride` a slot.
+/// Global cache pages: layer K/V in page 0, table rows of IDs, stride a stream (head dims 128, 256, 512).
 pub const Pages = struct { table: u64, page_elems: usize, tokens: usize, stride: usize = 0 };
 
 /// A row's heads and keys (pos - win, pos] (win 0: all) at cache row slot * ring + key % ring; p32: wide P in halves.
@@ -22,22 +22,22 @@ pub const page_unit = chunk_keys;
 /// Keys a prompt tile (window_rows.cuh's BN): a paged tile takes one page base, so prompt pages hold whole tiles.
 pub const prompt_tile = 32;
 
-/// Pages any paged kernel reads or `put` writes: no window, a wide head dim, a position count, i64 page steps.
+/// Pages any paged kernel reads or `put` writes: no window, a served head dim, a position count, i64 page steps.
 fn pageShape(w: Window) bool {
     const p = w.pg orelse return true;
-    return w.win == 0 and wide(w.d) and p.tokens > 0 and p.page_elems <= std.math.maxInt(i63);
+    return w.win == 0 and (w.d == 128 or wide(w.d)) and p.tokens > 0 and p.page_elems <= std.math.maxInt(i63);
 }
 
 /// Head dim 256 decode pages hold whole chunks; 512's decode looks up each key's page (prompt rows: promptShape).
 fn pagedShape(w: Window) bool {
     const p = w.pg orelse return true;
-    return pageShape(w) and (w.d != 256 or p.tokens % page_unit == 0);
+    return wide(w.d) and pageShape(w) and (w.d != 256 or p.tokens % page_unit == 0);
 }
 
 /// Prompt rows at either head dim: a 32-key tile takes one page base, so pages hold whole tiles (prompt_tile).
 fn promptShape(w: Window) bool {
     const p = w.pg orelse return true;
-    return pageShape(w) and p.tokens % prompt_tile == 0;
+    return wide(w.d) and pageShape(w) and p.tokens % prompt_tile == 0;
 }
 
 /// Head dims window_rows.cu serves, with any query heads a key head up to 16.
@@ -69,6 +69,7 @@ pub const WindowAttention = struct {
     tc_fn: Function,
     prompt_fns: [2]Function, // 4, then 8 query heads a block
     rows: Module, // window_rows.cu
+    prompt_pages128: [2]?Function = .{ null, null },
     pages: ?Module = null, // window_pages.cu (loadPages): window_rows.cuh over paged caches
     put_fn: ?Function = null, // tf_wrows_put: rows into paged caches
     wide_fns: [rows_smem.len]Wide,
@@ -103,7 +104,7 @@ pub const WindowAttention = struct {
         }, .prompt_fns = p, .rows = rows, .wide_fns = wf };
     }
 
-    /// Loads window_pages.cu: decode, prompt rows (bf16 P and P32) and `put` over caches in pages (head dims 256, 512).
+    /// Page functions: 128-dimensional prompt rows, wide decode/prompt rows, and the shared cache writer.
     pub fn loadPages(g: *WindowAttention, d: *const Driver) !void {
         if (g.pages != null) return;
         var pages = try Module.load(d, kernels.window_pages);
@@ -113,6 +114,8 @@ pub const WindowAttention = struct {
         const p512: Paged = .{ .part = try pages.function("tf_wrows_part_512_pg"), .merge = try pages.function("tf_wrows_merge_512_pg"), .prompt = .{ try pages.function("tf_wrows_prompt_512_pg"), try pages.function("tf_wrows_prompt_512_p32_pg") } };
         try p512.part.allowDynamicShared(staged_smem);
         for (p512.prompt) |pf| try pf.?.allowDynamicShared(rows_smem[1]);
+        g.prompt_pages128 = .{ try g.mod.function("tf_window_prompt_128_h4_pg"), try g.mod.function("tf_window_prompt_128_h8_pg") };
+        for (g.prompt_pages128) |pf| try pf.?.allowDynamicShared(prompt_smem);
         g.put_fn = try pages.function("tf_wrows_put");
         g.wide_fns[0].paged = p;
         g.wide_fns[1].paged = p512;
@@ -157,6 +160,7 @@ pub const WindowAttention = struct {
 
     /// n rows q [n, h, d] at pos/slot (int32 [n]) -> out [n, h, d]; `scratch`: scratchBytes(w, n) (d != 128).
     pub fn decode(g: *const WindowAttention, s: Stream, q: u64, kc: u64, vc: u64, out: u64, pos: u64, slot: ?u64, w: Window, scratch: u64, n: usize) !void {
+        if (w.d == 128 and w.pg != null) return error.WindowShape;
         if (wide(w.d)) return decodeWith(g.wideFns(w.d), s, q, kc, vc, out, pos, slot, w, scratch, n);
         if (w.hk == 0 or w.h % w.hk != 0 or w.h / w.hk > 16 or w.d > 128 or w.ring == 0) return error.WindowShape;
         if (w.win == 0 and w.ring > 4096) return error.WindowShape; // an unwindowed row's tiles come from the ring
@@ -262,9 +266,15 @@ pub const WindowAttention = struct {
         a.add(w.scale);
         a.add(@as(i32, @intCast(w.win)));
         a.add(@as(i32, @intCast(w.ring)));
+        if (w.pg != null) {
+            if (w.win != 0 or !pageShape(w)) return error.WindowShape;
+            addPages(&a, w.pg.?, false);
+        }
         const rows = 16 * (8 / hpc);
         const cfg: launch_.Config = .{ .grid = .{ .x = @intCast((n + rows - 1) / rows), .y = @intCast(w.hk * (grp / hpc)), .z = 1 }, .block = .{ .x = 256 }, .shared = prompt_smem };
-        try launch_.launch(g.prompt_fns[if (hpc == 8) 1 else 0], cfg, s, &a);
+        const which: usize = if (hpc == 8) 1 else 0;
+        const fun = if (w.pg != null) g.prompt_pages128[which] orelse return error.WindowShape else g.prompt_fns[which];
+        try launch_.launch(fun, cfg, s, &a);
     }
 
     /// Wide prompt rows, G up to 16; row i's last key hi[i] (int32, in [p0 + i, p0 + n)) or its position (null).
diff --git a/zig/src/families/kolibri1/backend.zig b/zig/src/families/kolibri1/backend.zig
index bff15c4987..f1f6fdbac8 100644
--- a/zig/src/families/kolibri1/backend.zig
+++ b/zig/src/families/kolibri1/backend.zig
@@ -6,6 +6,7 @@ const fwd = @import("forward.zig");
 const kern = @import("kernels.zig");
 const snap = @import("snap.zig");
 const limits = @import("context.zig");
+pub const shared = @import("shared.zig");
 
 const be = lanes.backend;
 const Stream = lanes.Stream;
@@ -21,6 +22,10 @@ const ring = 1024;
 
 const Lane = struct {
     slot: usize,
+    lease: @import("core").kv_pool.Lease = .{},
+    marks: []const u32 = &.{},
+    shared: bool = false,
+    kept: u32 = 0,
     len: u64 = 0, // the cache length: positions [0, len) hold this stream's keys
     base: u64 = 0, // the cache length the last verify started from
     rows: std.ArrayList(u32) = .empty, // the last verify's row tokens
@@ -44,6 +49,10 @@ pub const Kolibri = struct {
     meta: cuda.DeviceBuffer, // a draw's first position minus one, a window each
     slot_bufs: []cuda.DeviceBuffer, // a slot's last-row logits, then its rule: two a slot
     next: u64 = 0,
+    pool: ?@import("core").kv_pool.Pool = null,
+    prompts: ?@import("core").kv_prompts.Prompts = null,
+    restore_slot: ?usize = null,
+    restored_at: u32 = 0,
     steps: []?fwd.Model.Step, // a slot's one-row step, captured at its first use (null: not yet)
     graphed: bool, // TENSORFOLD_KOLIBRI_GRAPHS=0 runs every step eager
 
@@ -54,7 +63,8 @@ pub const Kolibri = struct {
 
     pub fn init(gpa: std.mem.Allocator, d: *const cuda.Driver, m: *fwd.Model, k: *const kern.Kernels, s: cuda.Stream) !Kolibri {
         const v = m.c.vocab;
-        const free = try gpa.alloc(bool, m.slots);
+        const slots = if (m.paged) round_rows else m.slots;
+        const free = try gpa.alloc(bool, slots);
         errdefer gpa.free(free);
         @memset(free, true);
         var logits = try cuda.DeviceBuffer.alloc(d, round_rows * v * 4);
@@ -65,7 +75,7 @@ pub const Kolibri = struct {
         errdefer tok.free();
         var meta = try cuda.DeviceBuffer.alloc(d, round_rows * 16);
         errdefer meta.free();
-        const bufs = try gpa.alloc(cuda.DeviceBuffer, 2 * m.slots);
+        const bufs = try gpa.alloc(cuda.DeviceBuffer, 2 * slots);
         errdefer gpa.free(bufs);
         var made: usize = 0;
         errdefer for (bufs[0..made]) |*b| b.free();
@@ -73,18 +83,23 @@ pub const Kolibri = struct {
             b.* = try cuda.DeviceBuffer.alloc(d, if (i % 2 == 0) v * 4 else @sizeOf(cuda.sample.Rule));
             made += 1;
         }
-        const steps = try gpa.alloc(?fwd.Model.Step, m.slots);
+        const steps = try gpa.alloc(?fwd.Model.Step, slots);
         @memset(steps, null);
         return .{ .gpa = gpa, .d = d, .m = m, .k = k, .s = s, .free = free, .logits = logits, .logits16 = logits16, .tok = tok, .meta = meta, .slot_bufs = bufs, .steps = steps, .graphed = !envOff("TENSORFOLD_KOLIBRI_GRAPHS") };
     }
 
     pub fn deinit(x: *Kolibri) void {
         var it = x.lanes.valueIterator();
-        while (it.next()) |l| l.rows.deinit(x.gpa);
+        while (it.next()) |l| {
+            if (x.pool) |*p| p.finish(&l.lease);
+            x.gpa.free(l.marks);
+            l.rows.deinit(x.gpa);
+        }
+        if (x.pool) |*p| p.deinit();
         x.lanes.deinit(x.gpa);
         for (x.slot_bufs) |*b| b.free();
         x.gpa.free(x.slot_bufs);
-        for (x.steps) |*g| if (g.*) |*e| e.exec.deinit();
+        x.dropGraphs();
         x.gpa.free(x.steps);
         x.gpa.free(x.free);
         x.logits.free();
@@ -93,6 +108,13 @@ pub const Kolibri = struct {
         x.meta.free();
     }
 
+    pub fn dropGraphs(x: *Kolibri) void {
+        for (x.steps) |*g| if (g.*) |*e| {
+            e.exec.deinit();
+            g.* = null;
+        };
+    }
+
     pub fn backend(x: *Kolibri) be.Backend {
         return .{ .ptr = x, .vtable = &vtable };
     }
@@ -101,7 +123,7 @@ pub const Kolibri = struct {
 
     /// The round loop's facts: no draft head; host copies verified in windows up to 16 rows, streams sharing a forward.
     pub fn facts(x: *const Kolibri) lanes.Model {
-        return .{ .exact_width = max_window, .gpu_tokens = false, .mtp = false, .speculate = false, .hidden_rows = true, .batch_rows = round_rows, .max_streams = @intCast(@min(x.m.slots, round_rows)) };
+        return .{ .exact_width = max_window, .gpu_tokens = false, .mtp = false, .speculate = false, .hidden_rows = true, .batch_rows = round_rows, .max_streams = @intCast(@min(if (x.m.paged) round_rows else x.m.slots, round_rows)) };
     }
 
     fn envOff(name: [*:0]const u8) bool {
@@ -117,7 +139,7 @@ pub const Kolibri = struct {
         return .{ .k = x.k, .s = x.s };
     }
 
-    fn upload(x: *Kolibri, dst: u64, bytes: []const u8) !void {
+    pub fn upload(x: *Kolibri, dst: u64, bytes: []const u8) !void {
         try x.d.check(x.d.api.cuMemcpyHtoDAsync_v2(dst, bytes.ptr, bytes.len, x.s.handle), "cuMemcpyHtoDAsync");
     }
 
@@ -139,11 +161,12 @@ pub const Kolibri = struct {
         const s: *const Stream = @ptrCast(@alignCast(owner orelse return error.NoOwner));
         const l = x.lanes.getPtr(s) orelse return error.UnknownStream;
         if ((l.fill orelse l.len) != at) return error.NotAtMark; // a state is copied only where the pass stands
-        return snap.save(x.gpa, x.m, x.s, l.slot, at);
+        return if (x.prompts != null) shared.save(x, l.slot, at) else snap.save(x.gpa, x.m, x.s, l.slot, at);
     }
 
     /// The stream's slot made to hold `k`'s state (a restore outside its prompt pass).
     pub fn restoreState(x: *Kolibri, owner: ?*anyopaque, k: *const snap.Saved) !void {
+        if (x.prompts != null) return shared.restore(x, k);
         const s: *const Stream = @ptrCast(@alignCast(owner orelse return error.NoOwner));
         const l = x.lanes.getPtr(s) orelse return error.UnknownStream;
         try snap.restore(x.m, x.s, l.slot, k);
@@ -151,7 +174,7 @@ pub const Kolibri = struct {
     }
 
     pub fn dropState(x: *Kolibri, k: *snap.Saved) void {
-        snap.drop(x.gpa, x.s, k);
+        shared.drop(x, k);
     }
 
     /// The lowest free slot.
@@ -160,34 +183,12 @@ pub const Kolibri = struct {
     }
 
     /// A prompt pass begins: a slot, the stream's rule, a kept state restored; where its first chunk starts.
-    fn opening(x: *Kolibri, s: *Stream) !*Lane {
-        const p = s.prompt();
-        if (p.len == 0) return error.EmptyPrompt;
-        if (!limits.fits(x.m.context, p.len, @as(usize, s.max_new))) return error.PastContext;
-        const got = try x.lanes.getOrPut(x.gpa, s);
-        if (!got.found_existing) {
-            const slot = x.freeSlot() orelse {
-                x.lanes.removeByPtr(got.key_ptr);
-                return error.NoFreeSlot;
-            };
-            x.free[slot] = false;
-            got.value_ptr.* = .{ .slot = slot, .last = x.slot_bufs[2 * slot].ptr, .rule = x.slot_bufs[2 * slot + 1].ptr };
-        }
-        const l = got.value_ptr;
-        l.sampled = false;
-        if (s.sampling) |sp| if (sp.temperature > 0) {
-            const r = cuda.sample.Rule.from(lanes.gpu_rule.rule(sp));
-            try x.upload(l.rule, std.mem.asBytes(&r));
-            l.sampled = true;
-        };
-        s.cached = 0;
-        l.fill = x.resumeAt(s, l, p.len);
-        s.fill_at.store(@intCast(l.fill.?), .monotonic);
-        return l;
+    pub fn opening(x: *Kolibri, s: *Stream) !*Lane {
+        return if (x.pool != null) shared.opening(Lane, x, s) else shared.legacyOpening(Lane, x, s);
     }
 
     /// The kept state's position, restored into the slot (0: none, or not this prompt's); any row resumes exactly.
-    fn resumeAt(x: *Kolibri, s: *Stream, l: *const Lane, len: usize) usize {
+    pub fn resumeAt(x: *Kolibri, s: *Stream, l: *const Lane, len: usize) usize {
         const saved = s.reuse.saved orelse return 0;
         const k: *const snap.Saved = @ptrCast(@alignCast(saved));
         if (k.at != s.reuse.at or k.at == 0 or k.at >= len) {
@@ -211,7 +212,7 @@ pub const Kolibri = struct {
             return error.Cancelled;
         }
         var end = @min(p.len, (at / prompt_rows + 1) * prompt_rows);
-        for (s.reuse.marks) |mk| if (mk > at and mk < end) {
+        for (if (l.shared) l.marks else s.reuse.marks) |mk| if (mk > at and mk < end) {
             end = mk;
         };
         try x.m.forward(x.ops(), &.{.{ .slot = l.slot, .p0 = at, .tokens = p[at..end] }}, true, null, l.last);
@@ -219,10 +220,18 @@ pub const Kolibri = struct {
         const done = end == p.len;
         l.fill = if (done) null else end;
         if (done) l.len = p.len;
-        if (s.reuse.hook) |hk| if (std.mem.indexOfScalar(u32, s.reuse.marks, @intCast(end)) != null) hk.at(hk.ptr, s, @intCast(end));
+        x.marked(s, end);
         return done;
     }
 
+    pub fn marked(x: *Kolibri, s: *Stream, at: usize) void {
+        if (x.lanes.getPtr(s)) |l| if (l.shared) {
+            shared.marked(x, s, l, at) catch |err| std.log.warn("kolibri1: prompt pages not kept: {s}", .{@errorName(err)});
+            return;
+        };
+        if (s.reuse.hook) |hk| if (std.mem.indexOfScalar(u32, s.reuse.marks, @intCast(at)) != null) hk.at(hk.ptr, s, @intCast(at));
+    }
+
     fn prefill(ptr: *anyopaque, s: *Stream) anyerror!void {
         const x = self(ptr);
         const l = try x.opening(s);
@@ -270,6 +279,7 @@ pub const Kolibri = struct {
         const l = x.lanes.getPtr(s) orelse return error.UnknownStream;
         const t = try x.value(feed);
         if (position != l.len + 1) return error.PositionMismatch;
+        try shared.reach(x, l, l.len + 1);
         const logits = try x.step(l, t);
         l.len += 1;
         const h = x.next;
@@ -314,6 +324,7 @@ pub const Kolibri = struct {
             try l.rows.append(x.gpa, w.pending);
             try l.rows.appendSlice(x.gpa, w.tokens);
             for (w.positions, 0..) |pos, r| if (pos != l.base + r + 1) return error.PositionMismatch;
+            try shared.reach(x, l, l.base + l.rows.items.len);
             c.* = .{ .slot = l.slot, .p0 = l.base, .tokens = l.rows.items };
             total += l.rows.items.len;
         }
@@ -367,6 +378,9 @@ pub const Kolibri = struct {
         if (x.lanes.fetchRemove(s)) |kv| {
             var l = kv.value;
             x.free[l.slot] = true;
+            if (x.pool) |*p| p.finish(&l.lease);
+            if (x.m.paged) x.m.caps[l.slot] = 0;
+            x.gpa.free(l.marks);
             l.rows.deinit(x.gpa);
         }
     }
diff --git a/zig/src/families/kolibri1/forward.zig b/zig/src/families/kolibri1/forward.zig
index 1e89b606e3..48c51f2860 100644
--- a/zig/src/families/kolibri1/forward.zig
+++ b/zig/src/families/kolibri1/forward.zig
@@ -5,6 +5,7 @@ const Config = @import("config.zig").Config;
 const wts = @import("weights.zig");
 const kern = @import("kernels.zig");
 const tree = @import("tree.zig");
+pub const pages = @import("pages.zig");
 
 const DeviceBuffer = cuda.DeviceBuffer;
 const grouped = cuda.grouped;
@@ -37,6 +38,12 @@ pub const Model = struct {
     groups: usize, // tree attention's key groups at the context and a full chain
     k: []u64, // per layer: the cache [slots, len, HK, D] bf16 (len: context for full layers, the ring for sliding)
     v: []u64,
+    paged: bool = false,
+    pages: usize = 0,
+    table: u64 = 0,
+    tree_table: u64 = 0,
+    ids: []u32 = &.{},
+    caps: []usize = &.{},
     norms: []u64, // per layer: q_norm then k_norm, fp32 [2, D]
     cos: u64,
     sin: u64,
@@ -47,7 +54,7 @@ pub const Model = struct {
 
     const Scratch = struct { ids: u64, pos: u64, slot: u64, res: u64, x: u64, qkv: u64, q: u64, kn: u64, vn: u64, att: u64, o: u64, part: u64, logits_e: u64, pick: u64, wts: u64, plan: grouped.Plan, act: u64, y: u64, mo: u64, flat: u64, paths: u64, depths: u64, offs: u64, sel: u64, win: u64, tree: u64 };
 
-    fn alloc(m: *Model, bytes: usize) !u64 {
+    pub fn alloc(m: *Model, bytes: usize) !u64 {
         if (m.planned) |*p| {
             p.* += @max(bytes, 256);
             return 0;
@@ -61,7 +68,7 @@ pub const Model = struct {
         return b.ptr;
     }
 
-    fn zeros(m: *Model, s: cuda.Stream, bytes: usize) !u64 {
+    pub fn zeros(m: *Model, s: cuda.Stream, bytes: usize) !u64 {
         const p = try m.alloc(bytes);
         if (m.planned != null) return p;
         try m.d.check(m.d.api.cuMemsetD8Async(p, 0, @max(bytes, 256), s.handle), "cuMemsetD8Async");
@@ -77,7 +84,12 @@ pub const Model = struct {
 
     /// Caches for `slots` streams of `context` positions, RoPE tables, and scratch for `max_rows` rows a forward.
     pub fn init(gpa: std.mem.Allocator, d: *const cuda.Driver, s: cuda.Stream, w: *const wts.Weights, k: *const kern.Kernels, context: usize, slots: usize, max_rows: usize) !Model {
-        var m = blank(gpa, d, w, context, slots, max_rows);
+        return initPaged(gpa, d, s, w, k, context, slots, max_rows, false);
+    }
+
+    pub fn initPaged(gpa: std.mem.Allocator, d: *const cuda.Driver, s: cuda.Stream, w: *const wts.Weights, k: *const kern.Kernels, context: usize, slots: usize, max_rows: usize, paged: bool) !Model {
+        var m = blank(gpa, d, w, context, if (paged) 1 else slots, max_rows);
+        m.paged = paged;
         errdefer m.deinit();
         try m.build(s, k);
         return m;
@@ -85,7 +97,12 @@ pub const Model = struct {
 
     /// The device bytes `init` allocates for these sizes, counted by the same path without allocating.
     pub fn openBytes(gpa: std.mem.Allocator, d: *const cuda.Driver, s: cuda.Stream, w: *const wts.Weights, k: *const kern.Kernels, context: usize, slots: usize, max_rows: usize) !usize {
-        var m = blank(gpa, d, w, context, slots, max_rows);
+        return openBytesPaged(gpa, d, s, w, k, context, slots, max_rows, false);
+    }
+
+    pub fn openBytesPaged(gpa: std.mem.Allocator, d: *const cuda.Driver, s: cuda.Stream, w: *const wts.Weights, k: *const kern.Kernels, context: usize, slots: usize, max_rows: usize, paged: bool) !usize {
+        var m = blank(gpa, d, w, context, if (paged) 1 else slots, max_rows);
+        m.paged = paged;
         defer m.deinit();
         m.planned = 0;
         try m.build(s, k);
@@ -103,17 +120,18 @@ pub const Model = struct {
         const d = m.d;
         const c = m.c;
         const w = m.w;
-        const slots = m.slots;
-        const context = m.context;
         const dh = c.head_dim;
         m.k = try gpa.alloc(u64, c.layers);
         m.v = try gpa.alloc(u64, c.layers);
         m.norms = try gpa.alloc(u64, c.layers);
+        if (m.paged) try pages.build(m, s);
         for (0..c.layers) |i| {
-            const len = if (c.full[i]) context else m.ring;
-            const bytes = slots * len * c.kv_heads * dh * 2;
-            m.k[i] = try m.zeros(s, bytes);
-            m.v[i] = try m.zeros(s, bytes);
+            if (!m.paged) {
+                const len = if (c.full[i]) m.context else m.ring;
+                const bytes = m.slots * len * c.kv_heads * dh * 2;
+                m.k[i] = try m.zeros(s, bytes);
+                m.v[i] = try m.zeros(s, bytes);
+            }
             m.norms[i] = try m.alloc(2 * dh * 4);
             if (m.planned != null) continue;
             const L = w.layers[i];
@@ -132,6 +150,8 @@ pub const Model = struct {
         m.gpa.free(m.k);
         m.gpa.free(m.v);
         m.gpa.free(m.norms);
+        m.gpa.free(m.ids);
+        m.gpa.free(m.caps);
     }
 
     /// RoPE tables: angle p / theta^(2j / D) in float64, cos and sin rounded to fp32, the halves repeated (neox).
@@ -211,10 +231,10 @@ pub const Model = struct {
             .act = try m.alloc(pairs * c.moe_width * 2),
             .y = try m.alloc(pairs * c.hidden * 4),
             .mo = try m.alloc(r * c.hidden * 2),
-            .flat = try m.alloc((2 * cap + 4 * m.slots) * 4),
+            .flat = try m.alloc((2 * cap + 4 * (if (m.paged) round_rows else m.slots)) * 4),
             .paths = try m.alloc(cap * tree.max_rows * 4),
             .depths = try m.alloc(cap * 4),
-            .offs = try m.alloc(c.layers * m.slots * 2 * 8),
+            .offs = try m.alloc(c.layers * (if (m.paged) round_rows else m.slots) * 2 * 8),
             .sel = try m.alloc(r * c.hidden * 2),
             .tree = try m.alloc(TreeAttention.scratchBytes(cap, c.heads, c.kv_heads, m.groups, dh)),
             .win = try m.alloc(cuda.window_attention.WindowAttention.scratchBytes(m.window(), cap)),
@@ -243,6 +263,7 @@ pub const Model = struct {
             if (ch.tokens.len > m.ring - m.c.window) return error.ChainOverrunsRing;
             if (ch.p0 + ch.tokens.len > m.context) return error.PastContext;
             if (ch.slot >= m.slots) return error.NoSuchSlot;
+            if (m.paged and ch.p0 + ch.tokens.len > m.caps[ch.slot]) return error.PastPages;
             n += ch.tokens.len;
         }
         if (n > m.max_rows or (!prompt and n > round_rows)) return error.TooManyRows;
@@ -272,7 +293,8 @@ pub const Model = struct {
             }
             plan = try tree.build(gpa, windows);
             try m.upload(s, m.s.flat, std.mem.sliceAsBytes(plan.?.flat));
-            t = m.treeRows(m.s.flat, plan.?, m.s.paths, m.s.depths, m.s.offs);
+            try pages.chains(m, s, chains);
+            t = m.treeRows(m.s.flat, plan.?, m.s.paths, m.s.depths, m.s.offs, null);
             try o.k.tree.paths(s, m.s.flat + plan.?.parentsAt(), m.s.paths, m.s.depths, n);
             try m.offsets(s, chains, m.s.offs);
             ng = TreeAttention.groups(most);
@@ -291,19 +313,22 @@ pub const Model = struct {
         try m.launch(o, n, if (prompt) chains[0] else null, t, chains.len, ng, @min(m.max_rows, round_rows), pick, out);
     }
 
-    fn treeRows(m: *const Model, base: u64, p: tree.Plan, paths: u64, depths: u64, offs: u64) cuda.tree_attention.Rows {
-        return .{ .row_stream = base, .streams = base + p.streamsAt(), .offs = offs, .origin = m.origin, .paths = paths, .depths = depths };
+    pub fn treeRows(m: *const Model, base: u64, p: tree.Plan, paths: u64, depths: u64, offs: u64, slot: ?usize) cuda.tree_attention.Rows {
+        return .{ .row_stream = base, .streams = base + p.streamsAt(), .offs = offs, .origin = m.origin, .paths = paths, .depths = depths, .pages = pages.treeView(m, slot) };
     }
 
     /// Each full layer's cache offsets a chain, K then V, in bf16 elements from `origin`: [layer][chain][2].
-    fn offsets(m: *Model, s: cuda.Stream, chains: []const Chain, dst: u64) !void {
+    pub fn offsets(m: *Model, s: cuda.Stream, chains: []const Chain, dst: u64) !void {
         const c = m.c;
         const per = m.context * c.kv_heads * c.head_dim * 2;
         const offs = try m.gpa.alloc(i64, c.layers * 2 * chains.len);
         defer m.gpa.free(offs);
         for (0..c.layers) |i| for (chains, 0..) |ch, j| {
-            offs[(i * chains.len + j) * 2] = @intCast((m.k[i] + ch.slot * per - m.origin) / 2);
-            offs[(i * chains.len + j) * 2 + 1] = @intCast((m.v[i] + ch.slot * per - m.origin) / 2);
+            const shift = if (m.paged and c.full[i]) 0 else ch.slot * per;
+            const key = m.k[i] + shift;
+            const val = m.v[i] + shift;
+            offs[(i * chains.len + j) * 2] = if (m.paged) @divExact(@as(i64, @bitCast(key -% m.origin)), 2) else @intCast((key - m.origin) / 2);
+            offs[(i * chains.len + j) * 2 + 1] = if (m.paged) @divExact(@as(i64, @bitCast(val -% m.origin)), 2) else @intCast((val - m.origin) / 2);
         };
         try m.upload(s, dst, std.mem.sliceAsBytes(offs));
     }
@@ -328,13 +353,22 @@ pub const Model = struct {
         for (w.layers, 0..) |L, i| {
             try (if (prompt) lin.prompt(s, L.qkv, .{ .x = m.s.x, .part = m.s.part }, m.s.qkv, n) else lin.decode(s, L.qkv, .{ .x = m.s.x, .part = m.s.part }, m.s.qkv, n));
             const len = if (c.full[i]) m.context else m.ring;
-            try o.qkv(.{ .rows = m.s.qkv, .norms = m.norms[i], .cos = m.cos, .sin = m.sin, .pos = m.s.pos, .slots = m.s.slot, .q = m.s.q, .k = m.s.kn, .v = m.s.vn, .kc = m.k[i], .vc = m.v[i], .len = len }, n, h, hk, dh, c.eps, !c.full[i]);
+            const pg = if (c.full[i]) pages.view(m, null) else null;
+            try o.qkv(.{ .rows = m.s.qkv, .norms = m.norms[i], .cos = m.cos, .sin = m.sin, .pos = m.s.pos, .slots = m.s.slot, .q = m.s.q, .k = m.s.kn, .v = m.s.vn, .kc = if (pg != null) null else m.k[i], .vc = if (pg != null) null else m.v[i], .len = len }, n, h, hk, dh, c.eps, !c.full[i]);
+            if (pg != null) {
+                var win = m.window();
+                win.win = 0;
+                win.ring = m.context;
+                win.pg = pg;
+                try o.k.win.put(s, m.s.kn, m.s.vn, m.k[i], m.v[i], m.s.pos, m.s.slot, win, n);
+            }
             if (prompt_chain) |ch| {
-                const per = len * hk * dh * 2; // a stream's rows of this layer's cache
+                const per = if (pg != null) 0 else len * hk * dh * 2; // a stream's rows of this layer's cache
                 var win = m.window();
                 if (c.full[i]) {
                     win.win = 0; // every key before the row
                     win.ring = m.context;
+                    win.pg = pages.view(m, ch.slot);
                 }
                 try o.promptRows(m.s.q, m.k[i] + ch.slot * per, m.v[i] + ch.slot * per, m.s.att, ch.p0, n, win);
             } else if (!c.full[i]) {
@@ -385,7 +419,7 @@ pub const Model = struct {
         const logits = try m.alloc(c.vocab * 4);
         try s.synchronize();
         try cuda.graph.beginCapture(s, .thread_local);
-        m.launch(o, 1, null, m.treeRows(flat, p, paths, depths, offs), 1, m.groups, 1, &.{0}, logits) catch |e| {
+        m.launch(o, 1, null, m.treeRows(flat, p, paths, depths, offs, slot), 1, m.groups, 1, &.{0}, logits) catch |e| {
             if (cuda.graph.endCapture(s)) |g| {
                 var gg = g;
                 gg.deinit();
@@ -401,6 +435,7 @@ pub const Model = struct {
     pub fn step(m: *Model, o: kern.Ops, g: *const Step, token: u32, p0: usize) !void {
         const s = o.s;
         if (p0 + 1 > m.context) return error.PastContext;
+        if (m.paged and p0 + 1 > m.caps[g.slot]) return error.PastPages;
         const staged = [3]i32{ @intCast(token), @intCast(p0), @intCast(g.slot) };
         try m.upload(s, m.s.ids, std.mem.asBytes(&staged[0]));
         try m.upload(s, m.s.pos, std.mem.asBytes(&staged[1]));
diff --git a/zig/src/families/kolibri1/kernels.zig b/zig/src/families/kolibri1/kernels.zig
index de488c55f7..70c2484c62 100644
--- a/zig/src/families/kolibri1/kernels.zig
+++ b/zig/src/families/kolibri1/kernels.zig
@@ -134,7 +134,7 @@ pub const Ops = struct {
         try o.k.norms.addRms(o.s, x, .bf16, res, w1, w2, out, .bf16, rows, n, eps);
     }
 
-    pub const Qkv = struct { rows: u64, norms: u64, cos: u64, sin: u64, pos: u64, slots: u64, q: u64, k: u64, v: u64, kc: u64, vc: u64, len: usize };
+    pub const Qkv = struct { rows: u64, norms: u64, cos: u64, sin: u64, pos: u64, slots: u64, q: u64, k: u64, v: u64, kc: ?u64, vc: ?u64, len: usize };
 
     /// q and k RMS-normed (RoPE'd on sliding layers), v as is; k and v also into the cache at pos % len.
     pub fn qkv(o: Ops, a: Qkv, n: usize, h: usize, hk: usize, d: usize, eps: f32, rope: bool) !void {
diff --git a/zig/src/families/kolibri1/native.zig b/zig/src/families/kolibri1/native.zig
index 34d2132632..7af79be24d 100644
--- a/zig/src/families/kolibri1/native.zig
+++ b/zig/src/families/kolibri1/native.zig
@@ -8,6 +8,8 @@ const fwd = @import("forward.zig");
 const be = @import("backend.zig");
 const snap = @import("snap.zig");
 const limits = @import("context.zig");
+const kv = @import("core").kv_pool;
+const shared = @import("shared.zig");
 
 pub const model_type = "kolibri1";
 pub const formats: []const []const u8 = &.{"fp8-block"};
@@ -16,6 +18,7 @@ pub const max_segments: u32 = 1;
 pub const prompt_rows: u32 = be.prompt_rows;
 /// A kept state is the caches through its position and nothing past it (no draft head reads a token on).
 pub const snap_lookahead: u32 = 0;
+pub const window_rows: u32 = be.max_window;
 /// No drafter yet: --drafter is refused before the load.
 pub const takes_drafter = false;
 pub const drafter_refusal = "--drafter: Kolibri-1 on CUDA loads no drafter yet; serve it without --drafter";
@@ -41,6 +44,8 @@ pub const Loaded = struct {
     deinit: *const fn (*anyopaque) void,
     lone: ?LoneRun = null,
     learn_bytes: usize = 0,
+    layout: ?kv.Layout = null,
+    engine_paged: bool = false,
     target: *be.Kolibri, // the prompt cache keeps this backend's states (snap_*)
 };
 
@@ -62,10 +67,11 @@ const Plan = struct {
     own: *const Owned,
     d: *const cuda.Driver,
     context: usize,
+    paged: bool,
 
     fn bytes(p: Plan, n: usize) !u64 {
         const o = p.own;
-        return try fwd.Model.openBytes(o.gpa, p.d, o.s, &o.w, &o.k, p.context, n, be.prompt_rows) + be.Kolibri.openBytes(o.w.config.vocab, n);
+        return try fwd.Model.openBytesPaged(o.gpa, p.d, o.s, &o.w, &o.k, p.context, n, be.prompt_rows, p.paged) + be.Kolibri.openBytes(o.w.config.vocab, if (p.paged) be.round_rows else n);
     }
 };
 
@@ -79,6 +85,7 @@ pub fn open(gpa: std.mem.Allocator, io: std.Io, ctx: *const cuda.Context, dir: [
     if (kernels != null) return error.CapturedKernels;
     if (o.lanes_fixed and o.lanes > be.round_rows) return error.ParallelPastRound; // before the weights load
     const context = try positions(o.context);
+    const paged = try @import("core").kv_prompts.sharing() > 0;
     const own = try gpa.create(Owned);
     errdefer gpa.destroy(own);
     own.gpa = gpa;
@@ -86,19 +93,20 @@ pub fn open(gpa: std.mem.Allocator, io: std.Io, ctx: *const cuda.Context, dir: [
     errdefer own.w.deinit();
     own.k = try kern.Kernels.load(ctx);
     errdefer own.k.deinit();
+    if (paged) try own.k.win.loadPages(ctx.d);
     own.s = try cuda.Stream.init(ctx.d, false);
     errdefer own.s.deinit();
-    // every stream's caches are allocated here and none later: plan them inside the room the weights leave
-    const plan: Plan = .{ .own = own, .d = ctx.d, .context = context };
+    // Price the original caches, or the first paged slot and metadata, inside the room the weights leave.
+    const plan: Plan = .{ .own = own, .d = ctx.d, .context = context, .paged = paged };
     const room: u64 = if (o.room) |r| r.read(r.ctx) orelse 0 else std.math.maxInt(u64);
-    const n = fit(room, o.lanes, o.lanes_fixed, plan) catch |e| {
+    const n = (if (paged) fitFirst(room, plan) else fit(room, o.lanes, o.lanes_fixed, plan)) catch |e| {
         if (e == error.ParallelPastMemory) {
             const most = fit(room, o.lanes, false, plan) catch 0;
             std.log.err("--parallel {d} needs {d:.2} GiB at a {d}-token window and the CUDA memory budget leaves {d:.2} GiB after the weights: {d} stream{s} fit", .{ o.lanes, gib(try plan.bytes(o.lanes)), context, gib(room), most, if (most == 1) "" else "s" });
         }
         return e;
     };
-    own.m = try fwd.Model.init(gpa, ctx.d, own.s, &own.w, &own.k, context, n, be.prompt_rows);
+    own.m = try fwd.Model.initPaged(gpa, ctx.d, own.s, &own.w, &own.k, context, n, be.prompt_rows, paged);
     errdefer own.m.deinit();
     own.kb = try be.Kolibri.init(gpa, ctx.d, &own.m, &own.k, own.s);
     return .{
@@ -110,6 +118,8 @@ pub fn open(gpa: std.mem.Allocator, io: std.Io, ctx: *const cuda.Context, dir: [
         .ctx = own,
         .deinit = release,
         .target = &own.kb,
+        .layout = if (paged) fwd.pages.layout(&own.m) else null,
+        .engine_paged = paged,
     };
 }
 
@@ -121,9 +131,39 @@ fn kolibri(ptr: *anyopaque) *be.Kolibri {
     return @ptrCast(@alignCast(ptr));
 }
 
+fn fitFirst(room: u64, plan: Plan) !usize {
+    if (try plan.bytes(1) > room) return error.NoStreamFits;
+    return 1;
+}
+
+pub fn attach(p: *anyopaque, extra: u32, count: usize) !void {
+    const own: *Owned = @ptrCast(@alignCast(p));
+    try shared.attach(&own.kb, extra, count);
+}
+
+pub fn fits(p: *anyopaque, prompt: u32, max_new: u32) bool {
+    const own: *Owned = @ptrCast(@alignCast(p));
+    return shared.admits(&own.kb, prompt, max_new);
+}
+
+pub fn share_pool(p: *anyopaque) ?*kv.Pool {
+    const x = kolibri(p);
+    return if (x.pool) |*pool| pool else null;
+}
+
+pub fn share_attach(p: *anyopaque, prompts: ?@import("core").kv_prompts.Prompts) void {
+    kolibri(p).prompts = prompts;
+}
+
+pub fn kvPool(p: *anyopaque) ?kv.Source {
+    const own: *Owned = @ptrCast(@alignCast(p));
+    return if (own.kb.pool) |*pool| .{ .pool = pool, .page_tokens = fwd.pages.tokens } else null;
+}
+
 /// Bytes a state at `at` prompt tokens takes.
 pub fn snap_bytes(ptr: *anyopaque, at: u32) u64 {
-    return snap.bytes(kolibri(ptr).m, at);
+    const x = kolibri(ptr);
+    return if (x.prompts != null) snap.blobPages(x.m, at, at) * fwd.pages.pageBytes(x.m) else snap.bytes(x.m, at);
 }
 
 /// The prompt pass stands at `at` (a chunk start): copy the stream's caches there.
@@ -142,6 +182,7 @@ pub fn snap_drop(ptr: *anyopaque, saved: *anyopaque) void {
 /// A request or load this engine refuses, in words; null: none of its own.
 pub fn explain(_: ?*anyopaque, err: anyerror) ?[]const u8 {
     return switch (err) {
+        error.BadKvShare => "TF_KV_SHARE takes a nonnegative count of kept prompt states (0 or unset: copies)",
         error.PastContext => "the prompt and its reply exceed this server's context window: shorten it or lower max_tokens",
         error.NoFreeSlot => "every stream's caches are in use: retry once another request ends",
         error.NoDraftHead => "Kolibri-1 on CUDA has no draft head yet: drafts come from the prompt's own text only",
diff --git a/zig/src/families/kolibri1/pages.zig b/zig/src/families/kolibri1/pages.zig
new file mode 100644
index 0000000000..b8da1d4eb0
--- /dev/null
+++ b/zig/src/families/kolibri1/pages.zig
@@ -0,0 +1,129 @@
+//! Global K/V as core-pool pages; sliding layers retain their existing per-slot rings.
+const std = @import("std");
+const cuda = @import("cuda");
+const kv = @import("core").kv_pool;
+const fwd = @import("forward.zig");
+
+pub const tokens = cuda.window_attention.page_unit;
+
+pub fn perSlot(m: anytype) usize {
+    return (m.context + tokens - 1) / tokens;
+}
+
+pub fn pageElems(m: anytype) usize {
+    return tokens * m.c.kv_heads * m.c.head_dim;
+}
+
+pub fn globalLayers(m: anytype) usize {
+    var n: usize = 0;
+    for (m.c.full[0..m.c.layers]) |full| n += @intFromBool(full);
+    return n;
+}
+
+pub fn pageBytes(m: anytype) u64 {
+    return globalLayers(m) * 2 * pageElems(m) * 2;
+}
+
+pub fn layout(m: anytype) kv.Layout {
+    var fixed: u64 = perSlot(m) * 4;
+    for (0..m.c.layers) |i| if (!m.c.full[i]) {
+        fixed += 2 * m.ring * m.c.kv_heads * m.c.head_dim * 2;
+    };
+    return .{ .fixed_bytes = fixed, .token_bytes = pageBytes(m) / tokens, .page_tokens = tokens, .snap_bytes = 0 };
+}
+
+pub fn view(m: anytype, slot: ?usize) ?cuda.window_attention.Pages {
+    if (!m.paged) return null;
+    return .{ .table = m.table + (slot orelse 0) * perSlot(m) * 4, .page_elems = pageElems(m), .tokens = tokens, .stride = perSlot(m) };
+}
+
+/// One warmup page before attach; otherwise the unchanged linear or sliding cache allocation.
+pub fn build(m: anytype, s: cuda.Stream) !void {
+    for (0..m.c.layers) |i| {
+        const len = if (m.c.full[i]) m.context else m.ring;
+        const n = if (m.paged and m.c.full[i]) pageElems(m) else m.slots * len * m.c.kv_heads * m.c.head_dim;
+        m.k[i] = try m.zeros(s, n * 2);
+        m.v[i] = try m.zeros(s, n * 2);
+    }
+    if (m.paged) {
+        try tables(m, s);
+        m.tree_table = try m.alloc(fwd.round_rows * perSlot(m) * 4);
+        if (m.planned == null) m.caps[0] = tokens;
+    }
+}
+
+fn tables(m: anytype, s: cuda.Stream) !void {
+    m.table = try m.alloc(m.slots * perSlot(m) * 4);
+    if (m.planned != null) return;
+    m.ids = try m.gpa.alloc(u32, m.slots * perSlot(m));
+    @memset(m.ids, 0);
+    m.caps = try m.gpa.alloc(usize, m.slots);
+    @memset(m.caps, 0);
+    try m.upload(s, m.table, std.mem.sliceAsBytes(m.ids));
+}
+
+fn forget(m: anytype, ptr: u64) void {
+    for (m.buffers.items, 0..) |*b, i| if (b.ptr == ptr) {
+        b.free();
+        _ = m.buffers.swapRemove(i);
+        return;
+    };
+}
+
+/// The native host has priced every slot and page; no forward may run until attach succeeds.
+pub fn attach(m: anytype, s: cuda.Stream, extra: usize, count: usize) !void {
+    if (!m.paged or m.pages != 0 or count == 0 or extra >= fwd.round_rows or globalLayers(m) == 0) return error.BadAttach;
+    try s.synchronize();
+    for (m.k, m.v) |k, v| {
+        forget(m, k);
+        forget(m, v);
+    }
+    forget(m, m.table);
+    m.gpa.free(m.ids);
+    m.gpa.free(m.caps);
+    m.ids = &.{};
+    m.caps = &.{};
+    m.slots = extra + 1;
+    m.pages = count;
+    const region = count * pageElems(m) * 2;
+    const arena = try m.zeros(s, globalLayers(m) * 2 * region);
+    var block: usize = 0;
+    for (0..m.c.layers) |i| {
+        if (m.c.full[i]) {
+            m.k[i] = arena + block * region;
+            m.v[i] = m.k[i] + region;
+            block += 2;
+        } else {
+            const bytes = m.slots * m.ring * m.c.kv_heads * m.c.head_dim * 2;
+            m.k[i] = try m.zeros(s, bytes);
+            m.v[i] = try m.zeros(s, bytes);
+        }
+    }
+    try tables(m, s);
+    try s.synchronize();
+}
+
+pub fn map(m: anytype, s: cuda.Stream, slot: usize, ids: []const u32, from: usize) !void {
+    const per = perSlot(m);
+    if (!m.paged or slot >= m.slots or from > per or ids.len > per - from) return error.BadChain;
+    for (ids) |id| if (id >= m.pages) return error.BadChain;
+    const at = slot * per + from;
+    @memcpy(m.ids[at..][0..ids.len], ids);
+    m.caps[slot] = (from + ids.len) * tokens;
+    try m.upload(s, m.table + at * 4, std.mem.sliceAsBytes(m.ids[at..][0..ids.len]));
+}
+
+/// Tree rows index chains rather than physical slots; a graph keeps its slot's stable table directly.
+pub fn treeView(m: anytype, slot: ?usize) ?cuda.tree_attention.Pages {
+    if (!m.paged) return null;
+    return .{ .table = if (slot) |i| m.table + i * perSlot(m) * 4 else m.tree_table, .page = pageElems(m), .tok = tokens, .stride = perSlot(m) };
+}
+
+/// Stage active slots in chain order; Q/K writes continue to use the physical-slot table.
+pub fn chains(m: anytype, s: cuda.Stream, rows: []const fwd.Chain) !void {
+    if (!m.paged) return;
+    const bytes = perSlot(m) * 4;
+    for (rows, 0..) |ch, j| {
+        try m.d.check(m.d.api.cuMemcpyDtoDAsync_v2(m.tree_table + j * bytes, m.table + ch.slot * bytes, bytes, s.handle), "cuMemcpyDtoDAsync");
+    }
+}
diff --git a/zig/src/families/kolibri1/shared.zig b/zig/src/families/kolibri1/shared.zig
new file mode 100644
index 0000000000..a6d2cbfbd0
--- /dev/null
+++ b/zig/src/families/kolibri1/shared.zig
@@ -0,0 +1,169 @@
+//! Kolibri's paged leases and prompt store; snapshots hold sliding tails in pages of the same pool.
+const std = @import("std");
+const cuda = @import("cuda");
+const core = @import("core");
+const lanes = @import("lanes");
+const fwd = @import("forward.zig");
+const snap = @import("snap.zig");
+const Stream = lanes.Stream;
+const max_window = @import("backend.zig").max_window;
+
+pub fn pagesFor(prompt: usize, max_new: usize) usize {
+    return pagesOf(prompt + max_new + max_window);
+}
+
+fn pagesOf(n: usize) usize {
+    return (n + fwd.pages.tokens - 1) / fwd.pages.tokens;
+}
+
+pub fn attach(x: anytype, extra: usize, count: usize) !void {
+    try x.s.synchronize();
+    x.dropGraphs();
+    try fwd.pages.attach(x.m, x.s, extra, count);
+    x.pool = try core.kv_pool.Pool.init(x.gpa, count, fwd.pages.pageBytes(x.m));
+}
+
+pub fn admits(x: anytype, prompt: usize, max_new: usize) bool {
+    if (!@import("context.zig").fits(x.m.context, prompt, max_new)) return true;
+    if (std.mem.indexOfScalar(bool, x.free[0..x.m.slots], true) == null) return false;
+    return if (x.pool) |*p| p.fits(pagesFor(prompt, max_new)) else true;
+}
+
+pub fn reach(x: anytype, l: anytype, end: usize) !void {
+    const p = if (x.pool) |*p| p else return;
+    const n = pagesOf(end);
+    if (n <= l.lease.ids.items.len) return;
+    const Map = struct {
+        x: @TypeOf(x),
+        slot: usize,
+        fn write(ctx: *anyopaque, ids: []const u32, from: usize) anyerror!void {
+            const m: *@This() = @ptrCast(@alignCast(ctx));
+            try fwd.pages.map(m.x.m, m.x.s, m.slot, ids[from..], from);
+        }
+    };
+    var map: Map = .{ .x = x, .slot = l.slot };
+    try p.cover(&l.lease, n, .{ .ctx = &map, .write = Map.write });
+}
+
+/// Restore rings while begin protects its found state, before a promise can reclaim the snapshot.
+pub fn restore(x: anytype, k: *const snap.Saved) !void {
+    const slot = x.restore_slot orelse return error.NoOwner;
+    if (k.from != k.at or k.at == 0) return error.KeptStateMismatch;
+    try snap.restore(x.m, x.s, slot, k);
+    x.restored_at = k.at;
+}
+
+pub fn save(x: anytype, slot: usize, at: u32) !*snap.Saved {
+    const p = if (x.pool) |*p| p else return error.NoPool;
+    const ids = try x.gpa.alloc(u32, snap.blobPages(x.m, at, at));
+    if (!p.hold(ids)) {
+        x.gpa.free(ids);
+        return error.NoCachePages;
+    }
+    return snap.saveShared(x.gpa, x.m, x.s, slot, at, at, ids) catch |err| {
+        cuda.completion.drain(x.d, x.s, "Kolibri shared snapshot save failed before pool pages could be returned");
+        for (ids) |id| p.release(id);
+        x.gpa.free(ids);
+        return err;
+    };
+}
+
+pub fn drop(x: anytype, k: *snap.Saved) void {
+    if (k.blob.len > 0) {
+        cuda.completion.drain(x.d, x.s, "Kolibri shared snapshot drop did not drain before releasing its pool pages");
+        for (k.blob) |id| x.pool.?.release(id);
+        x.gpa.free(k.blob);
+        k.blob = &.{};
+    }
+    snap.drop(x.gpa, x.s, k);
+}
+
+pub fn opening(comptime Lane: type, x: anytype, s: *Stream) !*Lane {
+    const prompt = s.prompt();
+    if (prompt.len == 0) return error.EmptyPrompt;
+    if (!@import("context.zig").fits(x.m.context, prompt.len, s.max_new)) return error.PastContext;
+    if (x.lanes.contains(s)) return error.StreamAlreadyOpen;
+    const slot = std.mem.indexOfScalar(bool, x.free[0..x.m.slots], true) orelse return error.NoFreeSlot;
+    var adopt: std.ArrayList(u32) = .empty;
+    defer adopt.deinit(x.gpa);
+    x.restore_slot = slot;
+    defer x.restore_slot = null;
+    x.restored_at = 0;
+    var adopted = false;
+    errdefer if (!adopted) for (adopt.items) |id| x.pool.?.release(id);
+    const pr = x.prompts;
+    const plan: core.kv_prompts.Plan = if (pr) |p| try p.begin(prompt, s.history_len, s.shared_prefixes, s, &adopt) else .{};
+    var l: Lane = .{ .slot = slot, .last = x.slot_bufs[2 * slot].ptr, .rule = x.slot_bufs[2 * slot + 1].ptr, .marks = plan.marks, .shared = pr != null };
+    errdefer {
+        x.m.caps[slot] = 0;
+        x.pool.?.finish(&l.lease);
+        x.gpa.free(l.marks);
+    }
+    try x.pool.?.promise(&l.lease, pagesFor(prompt.len, s.max_new) - adopt.items.len);
+    adopted = true; // adopt releases its input references itself on failure
+    try x.pool.?.adopt(&l.lease, adopt.items);
+    if (adopt.items.len > 0) try fwd.pages.map(x.m, x.s, slot, adopt.items, 0);
+    try reach(x, &l, prompt.len);
+    s.cached = 0;
+    if (pr != null) {
+        if (plan.from != x.restored_at) return error.KeptStateMismatch;
+        s.cached = plan.from;
+        l.fill = plan.from;
+    } else l.fill = x.resumeAt(s, &l, prompt.len);
+    if (s.sampling) |sp| if (sp.temperature > 0) {
+        const rule = cuda.sample.Rule.from(lanes.gpu_rule.rule(sp));
+        try x.upload(l.rule, std.mem.asBytes(&rule));
+        l.sampled = true;
+    };
+    try x.lanes.put(x.gpa, s, l);
+    x.free[slot] = false;
+    s.fill_at.store(@intCast(l.fill.?), .monotonic);
+    return x.lanes.getPtr(s).?;
+}
+
+pub fn marked(x: anytype, s: *Stream, l: anytype, at: usize) !void {
+    const pr = x.prompts orelse return;
+    if (!l.shared) return;
+    if (std.mem.indexOfScalar(u32, l.marks, @intCast(at)) != null) {
+        const n = at / fwd.pages.tokens;
+        const path = try x.gpa.alloc(u32, n);
+        defer x.gpa.free(path);
+        const mine = try x.gpa.dupe(u32, l.lease.ids.items[0..n]);
+        defer x.gpa.free(mine);
+        const shared = pr.keep(s.prompt(), @intCast(at), s, mine, path);
+        if (shared == n) l.kept += 1;
+        if (!std.mem.eql(u32, mine[0..shared], path[0..shared])) {
+            try fwd.pages.map(x.m, x.s, l.slot, path[0..shared], 0);
+            for (path[0..shared], 0..) |id, i| x.pool.?.share(&l.lease, i, id);
+            // Remapping a prefix must not shorten the already covered writable suffix.
+            x.m.caps[l.slot] = l.lease.ids.items.len * fwd.pages.tokens;
+        }
+    }
+    if (at == s.prompt().len) pr.report(at, s.cached, l.kept);
+}
+
+pub fn legacyOpening(comptime Lane: type, x: anytype, s: *Stream) !*Lane {
+    const p = s.prompt();
+    if (p.len == 0) return error.EmptyPrompt;
+    if (!@import("context.zig").fits(x.m.context, p.len, @as(usize, s.max_new))) return error.PastContext;
+    const got = try x.lanes.getOrPut(x.gpa, s);
+    if (!got.found_existing) {
+        const slot = std.mem.indexOfScalar(bool, x.free[0..x.m.slots], true) orelse {
+            x.lanes.removeByPtr(got.key_ptr);
+            return error.NoFreeSlot;
+        };
+        x.free[slot] = false;
+        got.value_ptr.* = .{ .slot = slot, .last = x.slot_bufs[2 * slot].ptr, .rule = x.slot_bufs[2 * slot + 1].ptr };
+    }
+    const l = got.value_ptr;
+    l.sampled = false;
+    if (s.sampling) |sp| if (sp.temperature > 0) {
+        const r = cuda.sample.Rule.from(lanes.gpu_rule.rule(sp));
+        try x.upload(l.rule, std.mem.asBytes(&r));
+        l.sampled = true;
+    };
+    s.cached = 0;
+    l.fill = x.resumeAt(s, l, p.len);
+    s.fill_at.store(@intCast(l.fill.?), .monotonic);
+    return l;
+}
diff --git a/zig/src/families/kolibri1/snap.zig b/zig/src/families/kolibri1/snap.zig
index 0c7961455d..2e7cd00800 100644
--- a/zig/src/families/kolibri1/snap.zig
+++ b/zig/src/families/kolibri1/snap.zig
@@ -1,36 +1,59 @@
-//! A stream's cache state at a prompt position, for the core prompt cache: full layers' keys, sliding layers' window.
+//! A stream's caches at a prompt position: full layers' [0, at), each sliding layer's last window - 1 positions.
 const std = @import("std");
 const cuda = @import("cuda");
 const fwd = @import("forward.zig");
+const Blob = @import("core").kv_prompts.Blob;
 
-/// A kept state: the keys of positions [0, at) of full layers, then the window's last keys of each sliding layer.
-pub const Saved = struct { at: u32, buf: cuda.DeviceBuffer };
+/// A kept state: per layer K then V, full [from, at), sliding last positions; [0, from) is the store's pages.
+pub const Saved = struct {
+    at: u32,
+    from: u32 = 0,
+    blob: []u32 = &.{}, // the copied rows in pages the store holds (TF_KV_SHARE), else in `buf`
+    buf: cuda.DeviceBuffer,
 
-/// Positions a sliding layer keeps at `at`: the next row reads (at - window, at], so window - 1 before it.
+    fn copyTo(k: *const Saved, m: *const fwd.Model, s: cuda.Stream, out: bool) Copy {
+        if (k.blob.len > 0) return .{ .m = m, .s = s, .blob = .{ .ids = k.blob, .chunk = chunk(m), .regions = regions(m) }, .out = out };
+        return .{ .m = m, .s = s, .buf = k.buf.ptr, .out = out };
+    }
+};
+
+/// Positions a sliding layer keeps at `at`: the keys the next row's window reaches back to.
 fn tail(m: *const fwd.Model, at: usize) usize {
     return @min(at, m.c.window - 1);
 }
 
-fn rowBytes(m: *const fwd.Model) usize {
+fn rowBytes(m: *const fwd.Model, _: usize) usize {
     return m.c.kv_heads * m.c.head_dim * 2;
 }
 
-/// Bytes a state at `at` takes.
-pub fn bytes(m: *const fwd.Model, at: usize) u64 {
+/// Bytes a state at `at` copies, its full layers' positions below `from` held as pages.
+pub fn copyBytes(m: *const fwd.Model, at: usize, from: usize) u64 {
     var n: usize = 0;
-    for (m.c.full[0..m.c.layers]) |full| n += 2 * (if (full) at else tail(m, at));
-    return n * rowBytes(m);
+    for (0..m.c.layers) |i| n += 2 * (if (m.c.full[i]) at - from else tail(m, at)) * rowBytes(m, i);
+    return n;
 }
 
-/// Each layer's cache, K then V, and positions [first, at) of `slot`: one copy a run of ring rows, in order.
-fn each(m: *const fwd.Model, slot: usize, at: usize, ctx: anytype, comptime f: fn (@TypeOf(ctx), dev: u64, off: usize, len: usize) anyerror!void) !void {
-    const rb = rowBytes(m);
+/// Each layer's cache, K then V, positions [first, at) of `slot` (full: from `from`), one copy a run, in order.
+fn each(m: *const fwd.Model, slot: usize, at: usize, from: usize, ctx: anytype, comptime f: fn (@TypeOf(ctx), dev: u64, off: usize, len: usize) anyerror!void) !void {
+    if (from > 0 and (!m.paged or from % fwd.pages.tokens != 0 or from > at)) return error.KeptPagesUnpaged;
     var off: usize = 0;
     for (0..m.c.layers) |i| for ([_][]u64{ m.k, m.v }) |cache| {
-        const full = m.c.full[i];
-        const len = if (full) m.context else m.ring;
+        const rb = rowBytes(m, i);
+        if (m.c.full[i] and m.paged) { // pages: p in page p / page_tokens of the slot's row, a copy a run within one
+            var p: usize = from;
+            while (p < at) {
+                const into = p % fwd.pages.tokens;
+                const run = @min(at - p, fwd.pages.tokens - into);
+                const row = @as(usize, m.ids[slot * fwd.pages.perSlot(m) + p / fwd.pages.tokens]) * fwd.pages.tokens + into;
+                try f(ctx, cache[i] + row * rb, off, run * rb);
+                off += run * rb;
+                p += run;
+            }
+            continue;
+        }
+        const len = if (m.c.full[i]) m.context else m.ring;
         const base = cache[i] + slot * len * rb;
-        var p = if (full) 0 else at - tail(m, at);
+        var p = if (m.c.full[i]) 0 else at - tail(m, at);
         while (p < at) {
             const row = p % len;
             const run = @min(at - p, len - row);
@@ -41,29 +64,70 @@ fn each(m: *const fwd.Model, slot: usize, at: usize, ctx: anytype, comptime f: f
     };
 }
 
-const Copy = struct { d: *const cuda.Driver, s: cuda.Stream, buf: u64, out: bool };
+const Copy = struct { m: *const fwd.Model, s: cuda.Stream, buf: u64 = 0, blob: ?Blob = null, out: bool };
 
 fn copy(c: Copy, dev: u64, off: usize, len: usize) anyerror!void {
-    const src = if (c.out) dev else c.buf + off;
-    const dst = if (c.out) c.buf + off else dev;
-    try c.d.check(c.d.api.cuMemcpyDtoDAsync_v2(dst, src, len, c.s.handle), "cuMemcpyDtoDAsync");
+    if (c.blob) |b| return b.each(off, len, Piece{ .c = c, .dev = dev }, Piece.copy);
+    try move(c, dev, c.buf + off, len);
 }
 
-/// Copy `slot`'s state at `at` into new storage, queued on `s` behind the chunk that ends there.
-pub fn save(gpa: std.mem.Allocator, m: *const fwd.Model, s: cuda.Stream, slot: usize, at: u32) !*Saved {
+fn move(c: Copy, dev: u64, kept: u64, len: u64) !void {
+    const src = if (c.out) dev else kept;
+    const dst = if (c.out) kept else dev;
+    try c.m.d.check(c.m.d.api.cuMemcpyDtoDAsync_v2(dst, src, len, c.s.handle), "cuMemcpyDtoDAsync");
+}
+
+/// A run's piece in one blob chunk: region r is full layer r / 2's K (even) or V (odd) region, page id's rows there.
+const Piece = struct {
+    c: Copy,
+    dev: u64,
+    fn copy(p: Piece, region: usize, id: u32, within: u64, from: u64, n: u64) anyerror!void {
+        const i = fullLayer(p.c.m, region / 2);
+        const base = (if (region % 2 == 0) p.c.m.k else p.c.m.v)[i];
+        try move(p.c, p.dev + from, base + id * chunk(p.c.m) + within, n);
+    }
+};
+
+/// The bytes a page holds in one full layer's K or V region.
+fn chunk(m: *const fwd.Model) u64 {
+    return fwd.pages.pageElems(m) * 2;
+}
+
+fn regions(m: *const fwd.Model) usize {
+    return @intCast(fwd.pages.pageBytes(m) / chunk(m));
+}
+
+fn fullLayer(m: *const fwd.Model, j: usize) usize {
+    var n: usize = 0;
+    for (0..m.c.layers) |i| if (m.c.full[i]) {
+        if (n == j) return i;
+        n += 1;
+    };
+    unreachable;
+}
+
+/// Pages a state at `at` (full layers from `from`) takes when its copied rows are kept in pool pages.
+pub fn blobPages(m: *const fwd.Model, at: usize, from: usize) usize {
+    return Blob.pagesFor(copyBytes(m, at, from), chunk(m), regions(m));
+}
+
+/// Copy `slot`'s state at `at` (full layers from `from`) to new storage or `blob`'s pages, queued after its chunk.
+pub fn saveShared(gpa: std.mem.Allocator, m: *const fwd.Model, s: cuda.Stream, slot: usize, at: u32, from: u32, blob: []u32) !*Saved {
     if (at == 0 or at > m.context) return error.NotAtMark;
+    if (from > at or (from > 0 and (!m.paged or from % fwd.pages.tokens != 0))) return error.KeptPagesUnpaged;
     const k = try gpa.create(Saved);
     errdefer gpa.destroy(k);
-    k.* = .{ .at = at, .buf = try cuda.DeviceBuffer.alloc(m.d, bytes(m, at)) };
+    k.* = .{ .at = at, .from = from, .blob = blob, .buf = .{ .d = m.d, .ptr = 0, .len = 0, .borrowed = true } };
+    if (blob.len == 0) k.buf = try cuda.DeviceBuffer.alloc(m.d, copyBytes(m, at, from));
     errdefer k.buf.free();
-    try each(m, slot, at, Copy{ .d = m.d, .s = s, .buf = k.buf.ptr, .out = true }, copy);
+    try each(m, slot, at, from, k.copyTo(m, s, true), copy);
     return k;
 }
 
 /// Make `slot`'s caches hold `k`'s state, queued on `s` before the next chunk.
 pub fn restore(m: *const fwd.Model, s: cuda.Stream, slot: usize, k: *const Saved) !void {
     if (k.at > m.context) return error.PastContext;
-    try each(m, slot, k.at, Copy{ .d = m.d, .s = s, .buf = k.buf.ptr, .out = false }, copy);
+    try each(m, slot, k.at, k.from, k.copyTo(m, s, false), copy);
 }
 
 /// Free a state once the stream's queued copies of it are done.
@@ -73,6 +137,14 @@ pub fn drop(gpa: std.mem.Allocator, s: cuda.Stream, k: *Saved) void {
     gpa.destroy(k);
 }
 
+pub fn bytes(m: *const fwd.Model, at: usize) u64 {
+    return copyBytes(m, at, 0);
+}
+
+pub fn save(gpa: std.mem.Allocator, m: *const fwd.Model, s: cuda.Stream, slot: usize, at: u32) !*Saved {
+    return saveShared(gpa, m, s, slot, at, 0, &.{});
+}
+
 const testing = std.testing;
 
 const Count = struct {
@@ -90,6 +162,7 @@ test "a state holds full layers' positions and each sliding ring's window, wrapp
     var v = [_]u64{ 3 << 20, 4 << 20 };
     var m: fwd.Model = undefined;
     m.c = c;
+    m.paged = false;
     m.context = 64;
     m.ring = 8;
     m.k = &k;
@@ -98,7 +171,7 @@ test "a state holds full layers' positions and each sliding ring's window, wrapp
     try testing.expectEqual(@as(u64, 2 * (4 + 10) * rb), bytes(&m, 10));
     var runs: std.ArrayList([3]usize) = .empty;
     defer runs.deinit(testing.allocator);
-    try each(&m, 1, 10, Count{ .runs = &runs }, Count.add);
+    try each(&m, 1, 10, 0, Count{ .runs = &runs }, Count.add);
     // sliding K: positions 6..9 at ring rows 6, 7 then 0, 1 of slot 1; then V; then full K and V, positions 0..9
     const want = [_][3]usize{
         .{ (1 << 20) + (8 + 6) * rb, 0, 2 * rb },      .{ (1 << 20) + 8 * rb, 2 * rb, 2 * rb },
diff --git a/zig/tests/cuda/contract_window.zig b/zig/tests/cuda/contract_window.zig
index 46e855cefb..390ae0955d 100644
--- a/zig/tests/cuda/contract_window.zig
+++ b/zig/tests/cuda/contract_window.zig
@@ -14,9 +14,12 @@ const wa = cuda.window_attention;
 const split_why = "a score's absolute rounding error is its weight's relative error: gates 1 and 3 hold outputs to attention_ref's bound";
 
 /// A decode or prompt call's fixed heads, window and ring (the variant), and which kernels it runs.
-pub const Cfg = struct { name: []const u8, h: usize, hk: usize, d: usize, win: usize, ring: usize, streams: usize = 1, prompt: bool = false, served: []const c.Shape };
+pub const Cfg = struct { name: []const u8, h: usize, hk: usize, d: usize, win: usize, ring: usize, streams: usize = 1, prompt: bool = false, pages: usize = 0, swap: bool = false, served: []const c.Shape };
 
 pub const cfgs = [_]Cfg{
+    .{ .name = "cuda window prompt 128 h4 over shuffled pages, same bytes as linear", .h = 48, .hk = 4, .d = 128, .win = 0, .ring = 0, .prompt = true, .pages = 256, .served = &.{ .{ .rows = 128, .n = 48, .k = 0 }, .{ .rows = 96, .n = 48, .k = 8177 }, .{ .rows = 31, .n = 48, .k = 513 } } },
+    .{ .name = "cuda window prompt 128 h8 over shuffled pages, same bytes as linear", .h = 16, .hk = 2, .d = 128, .win = 0, .ring = 0, .prompt = true, .pages = 256, .served = &.{ .{ .rows = 64, .n = 16, .k = 100 }, .{ .rows = 48, .n = 16, .k = 511 } } },
+    .{ .name = "control: window prompt 128 reads an unheld page", .h = 48, .hk = 4, .d = 128, .win = 0, .ring = 0, .prompt = true, .pages = 256, .swap = true, .served = &.{.{ .rows = 64, .n = 48, .k = 513 }} },
     // Kolibri's sliding layers (48 heads over 4, window 513): decode on tensor cores, prompt rows 4 heads a block
     .{ .name = "cuda window decode tf_window_decode_tc, 48/4 heads, window 513, ring 640", .h = 48, .hk = 4, .d = 128, .win = 513, .ring = 640, .streams = 2, .served = &.{ .{ .rows = 16, .n = 48, .k = 3000 }, .{ .rows = 64, .n = 48, .k = 1000 } } },
     .{ .name = "cuda window decode tf_window_decode_tc, 16/1 heads, no window, ring 300", .h = 16, .hk = 1, .d = 128, .win = 0, .ring = 300, .served = &.{.{ .rows = 16, .n = 16, .k = 300 }} },
@@ -60,6 +63,10 @@ pub const Window = struct {
     dev: [5]u64 = @splat(0),
     scratch: u64 = 0,
     out: u64 = 0,
+    pg: ?wa.Pages = null,
+    paged_k: u64 = 0,
+    paged_v: u64 = 0,
+    plain: u64 = 0,
 
     pub fn deinit(self: *Window) void {
         const a = self.rig.gpu.gpa;
@@ -69,7 +76,7 @@ pub const Window = struct {
     }
 
     pub fn kernel(self: *Window) g.Kernel {
-        return .{ .spec = spec(self.f), .ctx = self, .vt = &vt };
+        return .{ .spec = spec(self.f), .ctx = self, .vt = &vt, .control = if (self.f.swap) .shapes else null };
     }
 
     const vt = g.hooks(Window);
@@ -133,10 +140,38 @@ pub const Window = struct {
         }
         for ([_][]const u8{ std.mem.sliceAsBytes(self.q.items), std.mem.sliceAsBytes(self.kc.items), std.mem.sliceAsBytes(self.vc.items), std.mem.sliceAsBytes(self.pos.items), std.mem.sliceAsBytes(self.slot.items) }, &self.dev) |bytes, *p| p.* = try self.rig.input(bytes);
         self.scratch = if (f.prompt) 0 else try self.rig.scratch(wa.WindowAttention.scratchBytes(self.win(), n));
+        self.pg = null;
+        if (f.pages > 0) try self.place(rng);
         self.out = try self.rig.output(n * f.h * f.d * 2);
+        if (f.pages > 0) self.plain = try self.rig.scratch(n * f.h * f.d * 2);
         return true;
     }
 
+    fn place(self: *Window, rng: std.Random) !void {
+        const a = self.rig.gpu.gpa;
+        const row = self.f.hk * self.f.d;
+        const tok = self.f.pages;
+        const n = self.kc.items.len / row;
+        const per = (n + tok - 1) / tok;
+        const ids = try a.alloc(u32, per + 1);
+        defer a.free(ids);
+        for (ids, 0..) |*id, i| id.* = @intCast(i);
+        rng.shuffle(u32, ids);
+        const pool = try a.alloc(u16, (per + 1) * tok * row);
+        defer a.free(pool);
+        for ([_][]const u16{ self.kc.items, self.vc.items }, 0..) |src, part| {
+            @memset(pool, c.toBf16(std.math.nan(f32)));
+            for (0..n) |i| {
+                const dst = (@as(usize, ids[i / tok]) * tok + i % tok) * row;
+                @memcpy(pool[dst..][0..row], src[i * row ..][0..row]);
+            }
+            const ptr = try self.rig.input(std.mem.sliceAsBytes(pool));
+            if (part == 0) self.paged_k = ptr else self.paged_v = ptr;
+        }
+        if (self.f.swap) ids[0] = ids[per];
+        self.pg = .{ .table = try self.rig.input(std.mem.sliceAsBytes(ids[0..per])), .page_elems = tok * row, .tokens = tok, .stride = per };
+    }
+
     fn cell(self: *Window, x: []u16, r: usize, key: usize, kh: usize, e: usize) *u16 {
         const w = self.at(r);
         return &x[(self.cacheRow(w.base, key) * self.f.hk + kh) * self.f.d + e];
@@ -226,7 +261,12 @@ pub const Window = struct {
     pub fn run(self: *Window) !void {
         const n = self.shape.rows;
         if (self.f.prompt) {
-            try self.g.prompt(self.rig.s, self.dev[0], self.dev[1], self.dev[2], self.out, self.shape.k, n, self.win());
+            if (self.pg) |pg| {
+                var w = self.win();
+                w.pg = pg;
+                try self.g.prompt(self.rig.s, self.dev[0], self.paged_k, self.paged_v, self.out, self.shape.k, n, w);
+                try self.g.prompt(self.rig.s, self.dev[0], self.dev[1], self.dev[2], self.plain, self.shape.k, n, self.win());
+            } else try self.g.prompt(self.rig.s, self.dev[0], self.dev[1], self.dev[2], self.out, self.shape.k, n, self.win());
         } else try self.g.decode(self.rig.s, self.dev[0], self.dev[1], self.dev[2], self.out, self.dev[3], self.dev[4], self.win(), self.scratch, n);
         try self.rig.s.synchronize();
     }
@@ -262,7 +302,12 @@ pub const Window = struct {
     }
 
     pub fn check(self: *Window, gate: g.Gate) !?[]const u8 {
-        return self.r.?.check(self, self.heads(), self.scale(), try self.rig.outputBytes(), gate);
+        const got = try self.rig.outputBytes();
+        if (self.pg != null) {
+            const plain = try self.rig.bytes(self.plain, got.len);
+            if (!std.mem.eql(u8, got, plain)) return "paged prompt differs from its original linear kernel";
+        }
+        return self.r.?.check(self, self.heads(), self.scale(), got, gate);
     }
 
     pub fn reference(self: *Window, want: []f64, mag: []f64) !void {

From 663cbd66d96c9ce002861971368a15cb70bcda39 Mon Sep 17 00:00:00 2001
From: ashhart <71890694+ashhart@users.noreply.github.com>
Date: Sun, 1 Nov 2026 21:30:44 +0000
Subject: [PATCH 2/3] Keep narrow prompt page loads explicit

---
 zig/build/spark25.zig                |   9 +++
 zig/spark25_load_host.zig            | 117 +++++++++++++++++++++++++++
 zig/src/cuda/window_attention.zig    |  13 ++-
 zig/src/families/kolibri1/native.zig |   2 +-
 zig/tests/cuda/contract_window.zig   |   2 +-
 5 files changed, 138 insertions(+), 5 deletions(-)
 create mode 100644 zig/spark25_load_host.zig

diff --git a/zig/build/spark25.zig b/zig/build/spark25.zig
index ac977b30e1..495d757ca5 100644
--- a/zig/build/spark25.zig
+++ b/zig/build/spark25.zig
@@ -9,4 +9,13 @@ pub fn host(b: *std.Build, m: *std.Build.Module, all: *std.Build.Step) void {
     const run = b.addRunArtifact(b.addTest(.{ .root_module = m }));
     b.step("test-spark25-host", "Spark2.5 config admission, RoPE tables, cache sizes, snapshot runs and refusals without a GPU").dependOn(&run.step);
     all.dependOn(&run.step);
+    const cb = @import("cuda.zig");
+    const inert = b.addWriteFiles().add("inert-load-only.fatbin", "INERT DRIVER FIXTURE, NEVER A GPU IMAGE");
+    const images: [cb.kernels.len]?std.Build.LazyPath = @splat(inert);
+    const fake_cuda = cb.runtime(b, b.graph.host, .debug, &images);
+    const fake = module(b, b.graph.host, .debug, fake_cuda, m.import_table.get("core").?, m.import_table.get("lanes").?);
+    const root = b.createModule(.{ .root_source_file = b.path("zig/spark25_load_host.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{ .{ .name = "cuda", .module = fake_cuda }, .{ .name = "spark25", .module = fake } } });
+    const load = b.addRunArtifact(b.addTest(.{ .root_module = root, .filters = &.{"Spark load:"} }));
+    b.step("test-spark25-load-host", "Actual Spark kernel loading over inert images with every driver lookup observed").dependOn(&load.step);
+    all.dependOn(&load.step);
 }
diff --git a/zig/spark25_load_host.zig b/zig/spark25_load_host.zig
new file mode 100644
index 0000000000..55e5c8f90c
--- /dev/null
+++ b/zig/spark25_load_host.zig
@@ -0,0 +1,117 @@
+//! The actual Spark kernel constructor over inert images and a driver that records every load boundary.
+const std = @import("std");
+const cuda = @import("cuda");
+const spark = @import("spark25");
+var modules: usize = 0;
+var closed: usize = 0;
+var functions: usize = 0;
+var attributes: usize = 0;
+var prompt128: usize = 0;
+var trace: std.crypto.hash.sha2.Sha256 = undefined;
+fn reset() void {
+    modules = 0;
+    closed = 0;
+    functions = 0;
+    attributes = 0;
+    prompt128 = 0;
+    trace = .init(.{});
+}
+fn number(n: usize) void {
+    const v: u64 = n;
+    trace.update(std.mem.asBytes(&v));
+}
+fn moduleLoad(out: *cuda.abi.Module, _: ?*const anyopaque) callconv(.c) cuda.abi.Result {
+    modules += 1;
+    out.* = @ptrFromInt(modules);
+    trace.update("L");
+    number(modules);
+    return 0;
+}
+fn moduleClose(m: cuda.abi.Module) callconv(.c) cuda.abi.Result {
+    closed += 1;
+    trace.update("U");
+    number(@intFromPtr(m));
+    return 0;
+}
+fn functionGet(out: *cuda.abi.Function, m: cuda.abi.Module, name: [*:0]const u8) callconv(.c) cuda.abi.Result {
+    const s = std.mem.span(name);
+    if (std.mem.eql(u8, s, "tf_window_prompt_128_h4_pg") or std.mem.eql(u8, s, "tf_window_prompt_128_h8_pg")) prompt128 += 1;
+    functions += 1;
+    out.* = @ptrFromInt(1000 + functions);
+    trace.update("F");
+    number(@intFromPtr(m));
+    trace.update(s);
+    trace.update(&.{0});
+    return 0;
+}
+fn attrSet(f: cuda.abi.Function, a: cuda.abi.FunctionAttribute, value: c_int) callconv(.c) cuda.abi.Result {
+    attributes += 1;
+    trace.update("A");
+    number(@intFromPtr(f));
+    number(@intCast(@backingInt(a)));
+    number(@intCast(value));
+    return 0;
+}
+fn occupancy(out: *c_int, _: cuda.abi.Function, _: c_int, _: usize) callconv(.c) cuda.abi.Result {
+    out.* = 1;
+    return 0;
+}
+fn deviceAttr(out: *c_int, a: cuda.abi.DeviceAttribute, _: cuda.abi.Device) callconv(.c) cuda.abi.Result {
+    out.* = switch (a) {
+        .multiprocessor_count => 32,
+        .compute_capability_major => 12,
+        .compute_capability_minor => 1,
+        else => 0,
+    };
+    return 0;
+}
+fn currentDevice(out: *cuda.abi.Device) callconv(.c) cuda.abi.Result {
+    out.* = 0;
+    return 0;
+}
+fn driver() cuda.Driver {
+    var d: cuda.Driver = .{ .lib = undefined, .api = undefined };
+    d.api.cuModuleLoadData = moduleLoad;
+    d.api.cuModuleUnload = moduleClose;
+    d.api.cuModuleGetFunction = functionGet;
+    d.api.cuFuncSetAttribute = attrSet;
+    d.api.cuOccupancyMaxActiveBlocksPerMultiprocessor = occupancy;
+    d.api.cuDeviceGetAttribute = deviceAttr;
+    d.api.cuCtxGetDevice = currentDevice;
+    return d;
+}
+test "Spark load: default constructor never touches the new 128 page kernels" {
+    reset();
+    var d = driver();
+    const ctx: cuda.Context = .{ .d = &d, .device = 0, .handle = null };
+    var k = try spark.kernels.Kernels.load(&ctx);
+    errdefer k.deinit();
+    try std.testing.expect(k.win.pages != null);
+    try std.testing.expectEqual(@as(usize, 0), prompt128);
+    k.deinit();
+    try std.testing.expectEqual(modules, closed);
+    std.debug.print("SPARK_DEFAULT_LOAD modules={d} functions={d} attributes={d} closed={d} trace={x}\n", .{ modules, functions, attributes, closed, trace.finalResult() });
+}
+
+test "Spark load: explicit 128 page opt-in resolves both new functions exactly once" {
+    reset();
+    var d = driver();
+    const ctx: cuda.Context = .{ .d = &d, .device = 0, .handle = null };
+    var k = try spark.kernels.Kernels.load(&ctx);
+    defer k.deinit();
+    try std.testing.expectEqual(@as(usize, 0), prompt128);
+    const previous_functions = functions;
+    const previous_modules = modules;
+    const previous_attributes = attributes;
+    try k.win.loadPromptPages128(&d);
+    try std.testing.expectEqual(@as(usize, 2), prompt128);
+    try std.testing.expectEqual(previous_functions + 2, functions);
+    try std.testing.expectEqual(previous_attributes + 2, attributes);
+    try std.testing.expectEqual(previous_modules, modules);
+    try std.testing.expect(k.win.prompt_pages128[0] != null and k.win.prompt_pages128[1] != null);
+    try k.win.loadPromptPages128(&d);
+    try k.win.loadPages(&d);
+    try std.testing.expectEqual(previous_functions + 2, functions);
+    try std.testing.expectEqual(previous_attributes + 2, attributes);
+    try std.testing.expectEqual(previous_modules, modules);
+}
diff --git a/zig/src/cuda/window_attention.zig b/zig/src/cuda/window_attention.zig
index b368f6abf3..5b7f5aa989 100644
--- a/zig/src/cuda/window_attention.zig
+++ b/zig/src/cuda/window_attention.zig
@@ -104,7 +104,7 @@ pub const WindowAttention = struct {
         }, .prompt_fns = p, .rows = rows, .wide_fns = wf };
     }
 
-    /// Page functions: 128-dimensional prompt rows, wide decode/prompt rows, and the shared cache writer.
+    /// Wide page functions and the shared cache writer; 128-dimensional prompt pages are an explicit opt-in.
     pub fn loadPages(g: *WindowAttention, d: *const Driver) !void {
         if (g.pages != null) return;
         var pages = try Module.load(d, kernels.window_pages);
@@ -114,14 +114,21 @@ pub const WindowAttention = struct {
         const p512: Paged = .{ .part = try pages.function("tf_wrows_part_512_pg"), .merge = try pages.function("tf_wrows_merge_512_pg"), .prompt = .{ try pages.function("tf_wrows_prompt_512_pg"), try pages.function("tf_wrows_prompt_512_p32_pg") } };
         try p512.part.allowDynamicShared(staged_smem);
         for (p512.prompt) |pf| try pf.?.allowDynamicShared(rows_smem[1]);
-        g.prompt_pages128 = .{ try g.mod.function("tf_window_prompt_128_h4_pg"), try g.mod.function("tf_window_prompt_128_h8_pg") };
-        for (g.prompt_pages128) |pf| try pf.?.allowDynamicShared(prompt_smem);
         g.put_fn = try pages.function("tf_wrows_put");
         g.wide_fns[0].paged = p;
         g.wide_fns[1].paged = p512;
         g.pages = pages;
     }
 
+    /// Kolibri's shared prompts opt in separately, leaving Spark's existing wide-page load unchanged.
+    pub fn loadPromptPages128(g: *WindowAttention, d: *const Driver) !void {
+        if (g.prompt_pages128[0] != null and g.prompt_pages128[1] != null) return;
+        try g.loadPages(d);
+        const narrow_functions = [_]?Function{ try g.mod.function("tf_window_prompt_128_h4_pg"), try g.mod.function("tf_window_prompt_128_h8_pg") };
+        for (narrow_functions) |pf| try pf.?.allowDynamicShared(prompt_smem);
+        g.prompt_pages128 = narrow_functions;
+    }
+
     pub fn unload(g: *WindowAttention) void {
         if (g.pages) |*p| p.unload();
         g.rows.unload();
diff --git a/zig/src/families/kolibri1/native.zig b/zig/src/families/kolibri1/native.zig
index 7af79be24d..b72b8c4a4c 100644
--- a/zig/src/families/kolibri1/native.zig
+++ b/zig/src/families/kolibri1/native.zig
@@ -93,7 +93,7 @@ pub fn open(gpa: std.mem.Allocator, io: std.Io, ctx: *const cuda.Context, dir: [
     errdefer own.w.deinit();
     own.k = try kern.Kernels.load(ctx);
     errdefer own.k.deinit();
-    if (paged) try own.k.win.loadPages(ctx.d);
+    if (paged) try own.k.win.loadPromptPages128(ctx.d);
     own.s = try cuda.Stream.init(ctx.d, false);
     errdefer own.s.deinit();
     // Price the original caches, or the first paged slot and metadata, inside the room the weights leave.
diff --git a/zig/tests/cuda/contract_window.zig b/zig/tests/cuda/contract_window.zig
index 390ae0955d..95f3efaaa6 100644
--- a/zig/tests/cuda/contract_window.zig
+++ b/zig/tests/cuda/contract_window.zig
@@ -45,7 +45,7 @@ pub fn spec(f: Cfg) c.Spec {
 pub fn load(d: *const cuda.Driver) !wa.WindowAttention {
     var w = try wa.WindowAttention.load(d);
     errdefer w.unload();
-    try w.loadPages(d);
+    try w.loadPromptPages128(d);
     return w;
 }
 

From c6a4a51379a34b325c5fc8dab55b6d87c83928c4 Mon Sep 17 00:00:00 2001
From: ashhart <71890694+ashhart@users.noreply.github.com>
Date: Mon, 2 Nov 2026 03:30:29 +0000
Subject: [PATCH 3/3] changelog: Kolibri-1 on CUDA: prompts in the shared KV
 store with TF_KV_SHARE (off by default)

---
 CHANGELOG.md | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/CHANGELOG.md b/CHANGELOG.md
index 35d9f087dd..b751ccb97a 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -4,6 +4,8 @@ Each release's page on GitHub has its notes and measurements. See [the 1.0.0 rel
 
 ## Unreleased
 
+- `TF_KV_SHARE=N` now keeps Kolibri-1's prompts in CUDA's shared paged KV store too, as for Gemma 4. It is off by default (#1657).
+
 - GLM, GLM-DSA and Flash Next on Metal drop their checkpoint from the file cache once loaded, keeping the SSD expert and PLE files that serving still reads (#1868).
 
 - On Macs, the q27, Nemotron, LFM2 and Qwen3.5 loaders drop their checkpoint from the file cache once the weights are loaded, so unified memory doesn't hold them twice: on an M5 Ultra a loaded Qwen3.8-27B frees about 14 GiB and Nemotron 3.5 Lightning about 17 GiB, with the same replies (part of #1868).