From 36ca33ea76c1230248f041f14b146e0851457cee Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Mon, 2 Nov 2026 23:10:18 +0000
Subject: [PATCH 01/10] feat(hip_serve): an Offer value, an image helper and a
 cancel helper

Families that serve images need one way to hand the host their image
offer, prompt cache and host settings, a way to turn a typed tower
into an api.Vision, and a cancellation checkpoint on the lanes' own
CancelCheck instead of new hook types. These helpers add that,
defaulted so a family without them changes nothing.
---
 zig/src/hip_serve/cancel.zig | 67 +++++++++++++++++++++++++++++++++
 zig/src/hip_serve/offer.zig  | 27 +++++++++++++
 zig/src/hip_serve/serve.zig  |  6 +++
 zig/src/hip_serve/vision.zig | 73 ++++++++++++++++++++++++++++++++++++
 4 files changed, 173 insertions(+)
 create mode 100644 zig/src/hip_serve/cancel.zig
 create mode 100644 zig/src/hip_serve/offer.zig
 create mode 100644 zig/src/hip_serve/vision.zig

diff --git a/zig/src/hip_serve/cancel.zig b/zig/src/hip_serve/cancel.zig
new file mode 100644
index 0000000000..3eee8856f9
--- /dev/null
+++ b/zig/src/hip_serve/cancel.zig
@@ -0,0 +1,67 @@
+//! A request's cancellation hook at a tower's checkpoints, on the lanes' own CancelCheck.
+const std = @import("std");
+const CancelCheck = @import("lanes").stream.CancelCheck;
+
+/// error.Cancelled when the hook says the request left; a null hook never cancels.
+pub fn poll(hook: ?CancelCheck) error{Cancelled}!void {
+    if (hook) |h| if (h.check(h.ptr)) return error.Cancelled;
+}
+
+/// With a hook: finish the stream's queued work, then poll, so a cancelled request frees idle buffers; null: neither.
+pub fn drain(stream: anytype, hook: ?CancelCheck) !void {
+    if (hook == null) return;
+    try stream.synchronize();
+    try poll(hook);
+}
+
+const Mock = struct {
+    queued: bool = true,
+    synchronizations: usize = 0,
+    polls: usize = 0,
+    polled_before_drain: bool = false,
+    fail_sync: bool = false,
+    left: bool = true,
+
+    pub fn synchronize(m: *Mock) !void {
+        if (m.fail_sync) return error.HipFailed;
+        m.synchronizations += 1;
+        m.queued = false;
+    }
+
+    fn check(ptr: *anyopaque) bool {
+        const m: *Mock = @ptrCast(@alignCast(ptr));
+        m.polls += 1;
+        m.polled_before_drain = m.polled_before_drain or m.queued;
+        return m.left;
+    }
+
+    fn hook(m: *Mock) CancelCheck {
+        return .{ .ptr = m, .check = check };
+    }
+};
+
+test "a null hook neither synchronizes nor polls, so a tower without one stays asynchronous" {
+    var m: Mock = .{};
+    try drain(&m, null);
+    try poll(null);
+    try std.testing.expectEqual(@as(usize, 0), m.synchronizations);
+    try std.testing.expect(m.queued);
+}
+
+test "a hook drains the stream before it polls, and a left request cancels" {
+    var m: Mock = .{};
+    try std.testing.expectError(error.Cancelled, drain(&m, m.hook()));
+    try std.testing.expectEqual(@as(usize, 1), m.synchronizations);
+    try std.testing.expectEqual(@as(usize, 1), m.polls);
+    try std.testing.expect(!m.queued and !m.polled_before_drain);
+    var stays: Mock = .{ .left = false };
+    try drain(&stays, stays.hook());
+    try poll(stays.hook());
+    try std.testing.expectEqual(@as(usize, 2), stays.polls);
+}
+
+test "a failed synchronize is the error, and the hook is not polled" {
+    var m: Mock = .{ .fail_sync = true };
+    try std.testing.expectError(error.HipFailed, drain(&m, m.hook()));
+    try std.testing.expectEqual(@as(usize, 0), m.polls);
+}
diff --git a/zig/src/hip_serve/offer.zig b/zig/src/hip_serve/offer.zig
new file mode 100644
index 0000000000..f0ac40d5c6
--- /dev/null
+++ b/zig/src/hip_serve/offer.zig
@@ -0,0 +1,27 @@
+//! What a HIP family offers the host past its lanes: images, its prompt cache and its host settings, all defaulted.
+const api = @import("engine_api");
+
+/// A family's `Loaded.offer`; a family without one serves text with the host's defaults.
+pub const Offer = struct {
+    /// Built by the family; null: text only.
+    vision: ?api.Vision = null,
+    /// The family's own store; null whenever its cache is off (no budget or no slots).
+    cache: ?*api.prompt_cache.Store = null,
+    cache_capacity: ?api.PromptCacheCapacity = null,
+    warm_turns: bool = false,
+    /// The family's frame word for an image prompt, in its cache key (LaneHost.frame_key).
+    frame_key: ?*const fn (images: []const api.Image) u32 = null,
+    /// Null: the host's default (true).
+    thinking_budget: ?bool = null,
+    /// Runs first on the lane thread; a failure closes admission (lane_thread.zig).
+    on_thread: ?api.LaneHost.ThreadInit = null,
+    /// The family's effective policy line when it runs other settings than the host's defaults.
+    policy: ?[]const u8 = null,
+};
+
+test "an empty offer is text only with the host's defaults" {
+    const std = @import("std");
+    const o: Offer = .{};
+    try std.testing.expect(o.vision == null and o.cache == null and o.cache_capacity == null and !o.warm_turns);
+    try std.testing.expect(o.frame_key == null and o.thinking_budget == null and o.on_thread == null and o.policy == null);
+}
diff --git a/zig/src/hip_serve/serve.zig b/zig/src/hip_serve/serve.zig
index c95ea83e24..ee1bc99a64 100644
--- a/zig/src/hip_serve/serve.zig
+++ b/zig/src/hip_serve/serve.zig
@@ -8,6 +8,9 @@ pub const prefix = @import("prefix.zig");
 pub const backend = @import("backend.zig");
 pub const costs = @import("costs.zig");
 pub const Options = @import("options.zig").Options;
+pub const Offer = @import("offer.zig").Offer;
+pub const vision = @import("vision.zig");
+pub const cancel = @import("cancel.zig");
 
 test {
     _ = draw;
@@ -19,4 +22,7 @@ test {
     _ = backend;
     _ = costs;
     _ = @import("options.zig");
+    _ = @import("offer.zig");
+    _ = vision;
+    _ = cancel;
 }
diff --git a/zig/src/hip_serve/vision.zig b/zig/src/hip_serve/vision.zig
new file mode 100644
index 0000000000..49f9647a1a
--- /dev/null
+++ b/zig/src/hip_serve/vision.zig
@@ -0,0 +1,73 @@
+//! A family-side helper: a typed image tower (`c.image_token`, `limits`, `prepare`, `explain`) as an api.Vision.
+const std = @import("std");
+const api = @import("engine_api");
+
+/// `tower` is an optional pointer to the tower; null gives null (text only). The host never calls this.
+pub fn offer(tower: anytype) ?api.Vision {
+    const x = tower orelse return null;
+    const T = @typeInfo(@TypeOf(x)).pointer.child;
+    const prepare = struct {
+        fn f(ctx: *anyopaque, a: std.mem.Allocator, bytes: []const u8, max_tokens: u32) anyerror!api.PreparedImage {
+            const t: *T = @ptrCast(@alignCast(ctx));
+            const p = try t.prepare(a, bytes, max_tokens);
+            return .{ .pixels = p.pixels, .gh = p.gh, .gw = p.gw, .tokens = p.tokens, .hash = p.hash };
+        }
+    }.f;
+    var out: api.Vision = .{
+        .ctx = x,
+        .prepare = prepare,
+        .image_token = x.c.image_token,
+        .image_tokens = x.limits.image_tokens,
+        .max_images = x.limits.max_images,
+        .explain = T.explain,
+    };
+    if (comptime @hasField(T, "markers")) {
+        out.markers = &x.markers;
+    } else {
+        // A marker-free template still reserves its image placeholder against literal request text.
+        out.markers = @as(*const [1]u32, @ptrCast(&x.c.image_token));
+    }
+    if (comptime @hasDecl(T, "part_type")) out.part_type = T.part_type;
+    return out;
+}
+
+test "a marker-free image offer reserves its placeholder without adding wrapper tokens" {
+    const Tower = struct {
+        pub const part_type = "image";
+        c: struct { image_token: u32 } = .{ .image_token = 57 },
+        limits: struct { image_tokens: u32, max_images: u32 } = .{ .image_tokens = 64, .max_images = 2 },
+        pub fn prepare(_: *@This(), _: std.mem.Allocator, _: []const u8, _: u32) !api.PreparedImage {
+            return error.Unused;
+        }
+        pub fn explain(_: anyerror) ?[]const u8 {
+            return null;
+        }
+    };
+    var tower: Tower = .{};
+    const v = offer(@as(?*Tower, &tower)).?;
+    try std.testing.expectEqualSlices(u32, &.{57}, v.markers);
+    try std.testing.expectEqual(@as(u32, 57), v.image_token);
+    try std.testing.expectEqualStrings("image", v.part_type);
+    try std.testing.expect(v.wrap == null);
+    try std.testing.expect(v.markers.ptr == @as([*]const u32, @ptrCast(&tower.c.image_token)));
+    try std.testing.expect(offer(@as(?*Tower, null)) == null);
+}
+
+test "a tower with markers offers them, and its limits reach the host" {
+    const Tower = struct {
+        c: struct { image_token: u32 } = .{ .image_token = 9 },
+        limits: struct { image_tokens: u32, max_images: u32 } = .{ .image_tokens = 256, .max_images = 3 },
+        markers: [3]u32 = .{ 8, 9, 10 },
+        pub fn prepare(_: *@This(), _: std.mem.Allocator, _: []const u8, _: u32) !api.PreparedImage {
+            return error.Unused;
+        }
+        pub fn explain(_: anyerror) ?[]const u8 {
+            return null;
+        }
+    };
+    var tower: Tower = .{};
+    const v = offer(@as(?*Tower, &tower)).?;
+    try std.testing.expectEqualSlices(u32, &.{ 8, 9, 10 }, v.markers);
+    try std.testing.expectEqual(@as(u32, 256), v.image_tokens);
+    try std.testing.expectEqual(@as(u32, 3), v.max_images);
+}

From c904aafe1ecb5f3772c8974ae1840545228fb923 Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Tue, 3 Nov 2026 12:47:32 +0000
Subject: [PATCH 02/10] feat(native-hip): read family offers, refusals and
 model-type aliases

The HIP host now reads one optional offer from a family's Loaded and
copies its values into the lane host, takes a family refusal hook
before any load or group join, lets a family read extra model types,
and starts text-only with a log line when --vision meets a model
without images. GLM-OCR moves its image tower, prompt cache and cache
capacity into that offer and its --tp and --mtp-drafts refusals into
the hook, so the host keeps one path; the text families declare none
of these and get the empty offer and the same host values as before.
---
 zig/src/families/glm_ocr/hip_native.zig |  17 +++-
 zig/src/families/qwen3/native.zig       |   5 +-
 zig/src/native/hip.zig                  |  87 +++++++++++++--------
 zig/src/native/hip_families.zig         | 100 ++++++++++++++++++++++++
 4 files changed, 171 insertions(+), 38 deletions(-)
 create mode 100644 zig/src/native/hip_families.zig

diff --git a/zig/src/families/glm_ocr/hip_native.zig b/zig/src/families/glm_ocr/hip_native.zig
index deb940b660..d406870e22 100644
--- a/zig/src/families/glm_ocr/hip_native.zig
+++ b/zig/src/families/glm_ocr/hip_native.zig
@@ -16,6 +16,12 @@ pub const slide_refusal = "--slide: GLM-OCR on HIP has no Sliding Weights learne
 pub const confidence_refusal = "--mtp-confidence: GLM-OCR uses shared copy drafts; its MTP head is not qualified";
 pub const mtp_refusal = "--mtp-drafts: GLM-OCR's MTP head is not qualified; shared copy drafts remain available";
 pub const tp_refusal = "--tp: GLM-OCR currently runs on one HIP rank";
+/// GLM-OCR's refusals before any load: more than one rank, or MTP drafts.
+pub fn refusal(o: api.Open) ?[]const u8 {
+    if (o.tp != 1) return tp_refusal;
+    if (o.mtp_drafts) |n| if (n > 0) return mtp_refusal;
+    return null;
+}
 /// GLM-OCR writes no reasoning block: replies are all answer content; the prompt still follows enable_thinking.
 pub const reasoning = false;
 
@@ -33,9 +39,8 @@ pub const Loaded = struct {
     deinit: *const fn (*anyopaque) void,
     streams: usize,
     startup: []const u8,
-    vision: ?api.Vision,
-    cache: ?*api.prompt_cache.Store,
-    cache_capacity: ?api.PromptCacheCapacity,
+    /// Its image tower, prompt cache and that cache's capacity, for the host.
+    offer: serve.Offer,
 };
 fn refuse(a: std.mem.Allocator, problem: *[]const u8, comptime fmt: []const u8, args: anytype) error{Refused} {
     problem.* = std.fmt.allocPrint(a, fmt, args) catch fmt;
@@ -71,7 +76,11 @@ pub fn open(a: std.mem.Allocator, gpa: std.mem.Allocator, io: std.Io, dev: hip.D
     const startup = try std.fmt.allocPrint(gpa, "GLM-OCR BF16 on HIP {s}, device {d}; {d} shared lanes at context {d}; exact verification through 32 rows; full-shape prompt cache {d} bytes in at most {d} slots", .{ hip.kernels.arch, dev.index, e.streams, e.window, e.cache_bytes, e.cache_slots });
     own.* = .{ .gpa = gpa, .e = e, .h = h, .startup = startup };
     const vision: ?api.Vision = if (o.vision) |v| .{ .ctx = own, .prepare = prepare, .image_token = c.image_token, .markers = &markers, .part_type = "image", .image_tokens = if (v.image_tokens_set) v.image_tokens else 6144, .max_images = v.max_images, .explain = @import("sys").optional_libraries.explain } else null;
-    return .{ .backend = h.backend(), .facts = h.facts(), .ctx = own, .deinit = release, .streams = e.streams, .startup = startup, .vision = vision, .cache = if (e.cache_bytes > 0) &h.store else null, .cache_capacity = .{ .budget_bytes = e.cache_bytes, .minimum_state_bytes = @import("hip_snapshot.zig").Snapshot.bytes(1) } };
+    return .{ .backend = h.backend(), .facts = h.facts(), .ctx = own, .deinit = release, .streams = e.streams, .startup = startup, .offer = .{
+        .vision = vision,
+        .cache = if (e.cache_bytes > 0) &h.store else null,
+        .cache_capacity = .{ .budget_bytes = e.cache_bytes, .minimum_state_bytes = @import("hip_snapshot.zig").Snapshot.bytes(1) },
+    } };
 }
 pub fn follow(a: std.mem.Allocator, _: std.mem.Allocator, _: std.Io, _: hip.Device, _: []const u8, _: serve.Options, problem: *[]const u8) !void {
     return refuse(a, problem, "GLM-OCR currently runs on one HIP rank", .{});
diff --git a/zig/src/families/qwen3/native.zig b/zig/src/families/qwen3/native.zig
index 0f9d79fc03..b9658f3181 100644
--- a/zig/src/families/qwen3/native.zig
+++ b/zig/src/families/qwen3/native.zig
@@ -23,8 +23,11 @@ pub const confidence_refusal = "--mtp-confidence: Qwen3 dense checkpoints have n
 pub const reasoning = false;
 /// Every product is a BF16 dense one, which gfx12's WMMA tile runs: matrix=on is admitted on gfx12.
 pub const matrix_wmma12 = true;
-/// Checked by native/hip.zig before any rank joins a group.
 pub const tp_refusal = "--tp: Qwen3 dense checkpoints run on one HIP GPU; serve with --tp 1";
+/// Checked by native/hip.zig before any rank joins a group (hip_families.refused).
+pub fn refusal(o: api.Open) ?[]const u8 {
+    return if (o.tp != 1) tp_refusal else null;
+}
 pub const open = openEngine;
 pub const follow = followRank;
 
diff --git a/zig/src/native/hip.zig b/zig/src/native/hip.zig
index 472c36a42b..c850e1279b 100644
--- a/zig/src/native/hip.zig
+++ b/zig/src/native/hip.zig
@@ -11,6 +11,7 @@ const serve = @import("hip_serve");
 const glm_ocr = @import("glm_ocr");
 const qwen3 = @import("qwen3");
 const dots = @import("dots_mocr");
+const fam = @import("hip_families.zig");
 const Allocator = std.mem.Allocator;
 const admission = hip.admission;
 const modelContext = admission.modelContext;
@@ -23,12 +24,7 @@ pub const backends: []const []const u8 = &.{"hip"};
 pub const std_options: std.Options = .{ .signal_stack_size = null };
 /// The serve flags past the common ones this backend takes (server/cli.zig lists and parses them).
 pub const serves: []const []const u8 = &.{ "--tp", "--policy", "--mtp-drafts", "--mtp-confidence", "--checkpoint-slots", "--backend=rocm" };
-pub const families: []const api.Family = blk: {
-    var out: [registry.len]api.Family = undefined;
-    for (registry, 0..) |F, i| out[i] = .{ .model_type = F.model_type, .formats = F.formats };
-    const final = out;
-    break :blk &final;
-};
+pub const families: []const api.Family = fam.list(registry);
 
 /// The GPU family gate entries name ("rdna2" for gfx103x, "rdna3" for gfx11 / gfx12); null without a usable GPU.
 pub fn chip(a: Allocator) ?[]const u8 {
@@ -206,16 +202,8 @@ fn reasoning(comptime F: type) bool {
     return if (@hasDecl(F, "reasoning")) F.reasoning else true;
 }
 
-/// The family's refusal of --slide or --mtp-confidence, when it names one.
-fn flagRefused(comptime F: type, o: api.Open) ?[]const u8 {
-    if (o.tp != 1 and @hasDecl(F, "tp_refusal")) return F.tp_refusal;
-    if (o.mtp_drafts) |n| if (n > 0 and @hasDecl(F, "mtp_refusal")) return F.mtp_refusal;
-    if (@hasDecl(F, "refusal")) if (F.refusal(o)) |why| return why;
-    if (o.rope_yarn != null) return "--rope-yarn: the native HIP engine has no YaRN rope; serve without --rope-yarn, or on CUDA";
-    if (o.slide and @hasDecl(F, "slide_refusal")) return F.slide_refusal;
-    if (o.mtp_confidence != null and @hasDecl(F, "confidence_refusal")) return F.confidence_refusal;
-    return null;
-}
+/// The family's refusals before it loads or joins a group (hip_families.zig).
+const flagRefused = fam.refused;
 
 /// One loaded model behind the lane host: everything the engine thread reads lives here.
 const Host = struct {
@@ -247,13 +235,9 @@ fn readMemory(_: ?*anyopaque, reset_peak: bool) ?api.Memory {
     return .{ .active = u.device, .cache = 0, .peak = u.peak };
 }
 
-fn visionOffer(loaded: anytype) ?api.Vision {
-    return if (@hasField(@TypeOf(loaded), "vision")) loaded.vision else null;
-}
-
-test "HIP optional family offers: image metadata comes from the loaded owner" {
-    try std.testing.expect(visionOffer(.{}) == null);
-    try std.testing.expect(visionOffer(.{ .vision = @as(?api.Vision, null) }) == null);
+test "HIP optional family offers: image metadata comes from the loaded owner's offer" {
+    try std.testing.expect(fam.offerOf(.{}).vision == null);
+    try std.testing.expect(fam.offerOf(.{ .offer = serve.Offer{} }).vision == null);
     const Stub = struct {
         fn prepare(_: *anyopaque, _: Allocator, _: []const u8, _: u32) !api.PreparedImage {
             return error.Unused;
@@ -261,7 +245,7 @@ test "HIP optional family offers: image metadata comes from the loaded owner" {
     };
     var owner: u8 = 0;
     const v: api.Vision = .{ .ctx = &owner, .prepare = Stub.prepare, .image_token = 9, .image_tokens = 6144, .max_images = 2 };
-    const offered = visionOffer(.{ .vision = @as(?api.Vision, v) }).?;
+    const offered = fam.offerOf(.{ .offer = serve.Offer{ .vision = v } }).vision.?;
     try std.testing.expectEqual(v.ctx, offered.ctx);
     try std.testing.expectEqual(v.image_token, offered.image_token);
     try std.testing.expectEqual(v.image_tokens, offered.image_tokens);
@@ -280,7 +264,7 @@ pub fn open(a: Allocator, gpa: Allocator, io: std.Io, o: api.Open, problem: *[]c
         return null;
     }
     inline for (registry) |F| {
-        if (std.mem.eql(u8, o.model_type, F.model_type)) return openWith(F, a, gpa, io, o, problem);
+        if (fam.reads(F, o.model_type)) return openWith(F, a, gpa, io, o, problem);
     }
     problem.* = try std.fmt.allocPrint(a, "the native HIP engine has no backend for {s} checkpoints yet; serve with --engine python", .{o.model_type});
     return null;
@@ -303,6 +287,9 @@ fn openWith(comptime F: type, a: Allocator, gpa: Allocator, io: std.Io, o: api.O
         else => |x| return x,
     };
     errdefer loaded.deinit(loaded.ctx);
+    const offer = fam.offerOf(loaded);
+    if (fam.textOnly(o, offer))
+        std.debug.print("[tensorfold] --vision: the HIP engine has no image tower for {s}; serving text only, and image requests are refused\n", .{o.model_type});
     const h = try gpa.create(Host);
     errdefer gpa.destroy(h);
     h.gpa = gpa;
@@ -316,13 +303,32 @@ fn openWith(comptime F: type, a: Allocator, gpa: Allocator, io: std.Io, o: api.O
     h.core = lanes.Engine.init(gpa, &h.cfg, loaded.backend, h.clock.clock());
     errdefer h.core.deinit();
     const what = Serves.of(loaded, r.options.window);
-    h.host = api.LaneHost.init(gpa, io, &h.core, .{ .benchmark_ranks = if (r.dev.group) |g| @intCast(g.world) else 1, .lanes = @intCast(loaded.streams), .context_window = @intCast(what.window), .prefill_step = F.prefill_step, .decode_share = o.decode_share, .policy = h.policy_line, .startup = loaded.startup, .structures = what.text_generation, .text_generation = what.text_generation, .embeddings = what.embeddings, .decides = what.decides, .rerank = what.rerank, .readouts = what.readouts, .vision = visionOffer(loaded), .reasoning = reasoning(F) });
+    h.host = api.LaneHost.init(gpa, io, &h.core, .{
+        .benchmark_ranks = if (r.dev.group) |g| @intCast(g.world) else 1,
+        .lanes = @intCast(loaded.streams),
+        .context_window = @intCast(what.window),
+        .prefill_step = F.prefill_step,
+        .decode_share = o.decode_share,
+        .policy = offer.policy orelse h.policy_line,
+        .startup = loaded.startup,
+        .structures = what.text_generation and loaded.backend.vtable.first_masked != null,
+        .text_generation = what.text_generation,
+        .embeddings = what.embeddings,
+        .decides = what.decides,
+        .rerank = what.rerank,
+        .readouts = what.readouts,
+        .penalties = loaded.backend.penalties,
+        .reasoning = reasoning(F),
+        .vision = offer.vision,
+        .prompt_cache = offer.cache != null,
+        .prompt_cache_capacity = offer.cache_capacity,
+        .warm_turns = offer.warm_turns,
+        .thinking_budget = offer.thinking_budget orelse true,
+    });
+    h.host.cache = offer.cache;
+    h.host.frame_key = offer.frame_key;
+    h.host.on_thread = offer.on_thread;
     h.host.memory = .{ .read = readMemory };
-    if (@hasField(@TypeOf(loaded), "cache")) {
-        h.host.cache = loaded.cache;
-        h.host.info_.prompt_cache = loaded.cache != null;
-        h.host.info_.prompt_cache_capacity = loaded.cache_capacity;
-    }
     if (@hasDecl(F, "attach")) F.attach(h.family, &h.host, o);
     try h.host.start();
     served = true;
@@ -366,7 +372,7 @@ pub fn follow(a: Allocator, gpa: Allocator, io: std.Io, o: api.Open, problem: *[
         return false;
     }
     inline for (registry) |F| {
-        if (std.mem.eql(u8, o.model_type, F.model_type)) {
+        if (fam.reads(F, o.model_type)) {
             if (flagRefused(F, o)) |why| {
                 problem.* = why;
                 return false;
@@ -394,6 +400,19 @@ test "a HIP family declares replies without reasoning with `pub const reasoning
     }));
 }
 
+test "the text families carry no offer, so the host gives them the empty one; GLM-OCR offers through its Loaded" {
+    inline for (.{ qwen35.native, qwen35.native_moe, flashnext.native, nemotron.native }) |F| {
+        const Loaded = @typeInfo(@typeInfo(@TypeOf(F.open)).@"fn".return_type.?).error_union.payload;
+        try std.testing.expect(!@hasField(Loaded, "offer"));
+    }
+    try std.testing.expect(@FieldType(glm_ocr.native.Loaded, "offer") == serve.Offer);
+    inline for (.{ "vision", "cache", "cache_capacity" }) |old| try std.testing.expect(!@hasField(glm_ocr.native.Loaded, old));
+}
+
+test {
+    _ = fam;
+}
+
 test "every registered family is listed for capabilities" {
     try std.testing.expectEqual(@as(usize, registry.len), families.len);
     try std.testing.expectEqualStrings("qwen3_5", families[0].model_type);
@@ -425,7 +444,9 @@ test "Nemotron refuses --slide and --mtp-confidence on HIP before it loads, and
 
 test "a one-GPU family refuses --tp before the group forms; the others are not refused for it" {
     const OneGpu = struct {
-        pub const tp_refusal = "--tp: one GPU only";
+        pub fn refusal(o: api.Open) ?[]const u8 {
+            return if (o.tp != 1) "--tp: one GPU only" else null;
+        }
     };
     try std.testing.expectEqualStrings("--tp: one GPU only", flagRefused(OneGpu, .{ .dir = "d", .model_type = "m", .tp = 2 }).?);
     try std.testing.expect(flagRefused(OneGpu, .{ .dir = "d", .model_type = "m" }) == null);
diff --git a/zig/src/native/hip_families.zig b/zig/src/native/hip_families.zig
new file mode 100644
index 0000000000..5ec7c48307
--- /dev/null
+++ b/zig/src/native/hip_families.zig
@@ -0,0 +1,100 @@
+//! What the HIP host reads from a family's namespace and its Loaded: names, refusals and its offer.
+const std = @import("std");
+const api = @import("engine_api");
+const Offer = @import("hip_serve").Offer;
+
+/// Whether `F` reads `model_type`: its own, or one of its `also_model_types`.
+pub fn reads(comptime F: type, model_type: []const u8) bool {
+    if (std.mem.eql(u8, model_type, F.model_type)) return true;
+    if (comptime @hasDecl(F, "also_model_types")) for (F.also_model_types) |alias| {
+        if (std.mem.eql(u8, model_type, alias)) return true;
+    };
+    return false;
+}
+
+/// One capability entry per model_type a family reads, aliases after its own, with its `raw_text` when it names one.
+pub fn list(comptime registry: anytype) []const api.Family {
+    comptime {
+        var count: usize = 0;
+        for (registry) |F| count += 1 + if (@hasDecl(F, "also_model_types")) F.also_model_types.len else 0;
+        var out: [count]api.Family = undefined;
+        var i: usize = 0;
+        for (registry) |F| {
+            const raw = if (@hasDecl(F, "raw_text")) F.raw_text else true;
+            out[i] = .{ .model_type = F.model_type, .formats = F.formats, .raw_text = raw };
+            i += 1;
+            if (@hasDecl(F, "also_model_types")) for (F.also_model_types) |alias| {
+                out[i] = .{ .model_type = alias, .formats = F.formats, .raw_text = raw };
+                i += 1;
+            };
+        }
+        const final = out;
+        return &final;
+    }
+}
+
+/// Why `F` cannot take these flags, before any load or group join: its own `refusal(o)` first, then the shared ones.
+pub fn refused(comptime F: type, o: api.Open) ?[]const u8 {
+    if (comptime @hasDecl(F, "refusal")) if (F.refusal(o)) |why| return why;
+    if (o.rope_yarn != null) return "--rope-yarn: the native HIP engine has no YaRN rope; serve without --rope-yarn, or on CUDA";
+    if (o.slide and @hasDecl(F, "slide_refusal")) return F.slide_refusal;
+    if (o.mtp_confidence != null and @hasDecl(F, "confidence_refusal")) return F.confidence_refusal;
+    return null;
+}
+
+/// The family's offer, or the empty one (text, the host's defaults) when its Loaded carries none.
+pub fn offerOf(loaded: anytype) Offer {
+    return if (comptime @hasField(@TypeOf(loaded), "offer")) loaded.offer else .{};
+}
+
+/// --vision asked of a model this family serves without images: it starts text-only and logs so.
+pub fn textOnly(o: api.Open, offer: Offer) bool {
+    return o.vision != null and offer.vision == null;
+}
+
+const Plain = struct {
+    pub const model_type = "plain";
+    pub const formats: []const []const u8 = &.{"mlx"};
+};
+const Aliased = struct {
+    pub const model_type = "first";
+    pub const also_model_types: []const []const u8 = &.{"second"};
+    pub const formats: []const []const u8 = &.{"bf16"};
+    pub const raw_text = false;
+    pub fn refusal(o: api.Open) ?[]const u8 {
+        return if (o.tp > 1) "--tp: this family serves one rank" else null;
+    }
+};
+
+test "a family reads its own model_type and its aliases; the list has one entry each, with raw_text" {
+    try std.testing.expect(reads(Plain, "plain") and !reads(Plain, "first"));
+    try std.testing.expect(reads(Aliased, "first") and reads(Aliased, "second") and !reads(Aliased, "plain"));
+    const families = comptime list(.{ Plain, Aliased });
+    try std.testing.expectEqual(@as(usize, 3), families.len);
+    try std.testing.expectEqualStrings("plain", families[0].model_type);
+    try std.testing.expect(families[0].raw_text);
+    try std.testing.expectEqualStrings("second", families[2].model_type);
+    try std.testing.expectEqualStrings("bf16", families[2].formats[0]);
+    try std.testing.expect(!families[1].raw_text and !families[2].raw_text);
+}
+
+test "a family's own refusal comes first; --vision alone refuses nothing" {
+    const many: api.Open = .{ .dir = "d", .model_type = "first", .tp = 2, .rope_yarn = 2.0 };
+    try std.testing.expectEqualStrings("--tp: this family serves one rank", refused(Aliased, many).?);
+    try std.testing.expect(std.mem.startsWith(u8, refused(Plain, many).?, "--rope-yarn"));
+    const seen: api.Open = .{ .dir = "d", .model_type = "plain", .vision = .{} };
+    try std.testing.expect(refused(Plain, seen) == null and refused(Aliased, seen) == null);
+}
+
+test "a Loaded without an offer gives the empty one; --vision with no offered images is text-only" {
+    const Bare = struct { streams: usize = 1 };
+    const Offering = struct { offer: Offer };
+    const empty = offerOf(Bare{});
+    try std.testing.expect(empty.vision == null and empty.cache == null and empty.policy == null);
+    const warm = offerOf(Offering{ .offer = .{ .warm_turns = true, .policy = "fixed" } });
+    try std.testing.expect(warm.warm_turns);
+    try std.testing.expectEqualStrings("fixed", warm.policy.?);
+    const seen: api.Open = .{ .dir = "d", .model_type = "plain", .vision = .{} };
+    try std.testing.expect(textOnly(seen, empty));
+    try std.testing.expect(!textOnly(.{ .dir = "d", .model_type = "plain" }, empty));
+}

From 136a9af0227d0f62ca8ea637bd0a85f9fc75ccfe Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Mon, 2 Nov 2026 23:04:35 +0000
Subject: [PATCH 03/10] build(hip): derive the stub count and allow strict
 per-source flags

Two families that each add a kernel image no longer edit the same
literal: the stub's image count now comes from the group list, and
each group is a source with a `strict` flag that adds -O2
-fno-fast-math -fdenormal-fp-math=ieee after the common flags, so a
family image can ask for IEEE arithmetic without touching the others.
Every existing group keeps strict off, so its hipcc command is
unchanged, and `compile` keeps its signature for the controls.
---
 zig/build/hip.zig | 33 ++++++++++++++++++++++++++-------
 1 file changed, 26 insertions(+), 7 deletions(-)

diff --git a/zig/build/hip.zig b/zig/build/hip.zig
index f1fa189e86..55110e210d 100644
--- a/zig/build/hip.zig
+++ b/zig/build/hip.zig
@@ -380,11 +380,23 @@ pub fn server(b: *std.Build, target: std.Build.ResolvedTarget, e: Engine, build_
     return exe;
 }
 
+/// A group's source and whether it compiles with IEEE arithmetic (`strict`: -O2 -fno-fast-math, IEEE denormals).
+const Source = struct { path: []const u8, strict: bool = false };
+
 /// The model-free kernels' source groups, in zig/src/hip/kernels.zig's Group order.
-const groups = [_][]const u8{
-    "ops/ops.hip",           "ops/act.hip",             "attention/attention.hip", "recurrence/gated_delta.hip",
-    "attention/prefill.hip", "decode/decode.hip",       "decode/plan.hip",         "tiles/dot2_tiles.hip",
-    "tiles/dot2.hip",        "flashnext/flashnext.hip", "recurrence/mamba.hip",    "glm_ocr/pointwise.hip",
+const groups = [_]Source{
+    .{ .path = "ops/ops.hip" },
+    .{ .path = "ops/act.hip" },
+    .{ .path = "attention/attention.hip" },
+    .{ .path = "recurrence/gated_delta.hip" },
+    .{ .path = "attention/prefill.hip" },
+    .{ .path = "decode/decode.hip" },
+    .{ .path = "decode/plan.hip" },
+    .{ .path = "tiles/dot2_tiles.hip" },
+    .{ .path = "tiles/dot2.hip" },
+    .{ .path = "flashnext/flashnext.hip" },
+    .{ .path = "recurrence/mamba.hip" },
+    .{ .path = "glm_ocr/pointwise.hip" },
 };
 
 /// What the group sources include, so an edit to one rebuilds them.
@@ -404,7 +416,7 @@ const headers = [_][]const u8{
 /// The flags of the kernels' first build: no contraction, wave32 on RDNA, C++20.
 const flags = [_][]const u8{ "-D__HIP_PLATFORM_AMD__=1", "-DUSE_ROCM=1", "-DHIPBLAS_V2", "-fPIC", "-DCUDA_HAS_FP16=1", "-DHIP_ENABLE_WARP_SYNC_BUILTINS=1", "-std=c++20", "-fno-gpu-rdc", "-mno-wavefrontsize64", "-ffp-contract=off" };
 
-const stub = "pub const arch = \"\";\npub const images: [12][]align(8) const u8 = @splat(&.{});\n";
+const stub = std.fmt.comptimePrint("pub const arch = \"\";\npub const images: [{d}][]align(8) const u8 = @splat(&.{{}});\n", .{groups.len});
 
 /// One code object a group for `arch`, built with the caps table's instruction switches, as the `hip_kernels` module.
 pub fn kernelObjects(b: *std.Build, hipcc: []const u8, hipcc_resolved: []const u8, device_lib: ?[]const u8, relax_branches: bool, arch: []const u8) *std.Build.Module {
@@ -424,19 +436,26 @@ pub const Hipcc = struct { name: []const u8, resolved: []const u8, device_lib: ?
 /// The groups' code objects for `arch` in one directory, group0.hsaco onward.
 pub fn objectFiles(b: *std.Build, cc: Hipcc, arch: []const u8) *std.Build.Step.WriteFile {
     const files = b.addWriteFiles();
-    for (groups, 0..) |source, i| {
+    for (groups, 0..) |g, i| {
         const name = b.fmt("group{d}.hsaco", .{i});
-        _ = files.addCopyFile(compile(b, cc, arch, b.path(b.fmt("zig/kernels/hip/{s}", .{source})), name), name);
+        const source = b.path(b.fmt("zig/kernels/hip/{s}", .{g.path}));
+        _ = files.addCopyFile(compileSource(b, cc, arch, source, name, g.strict), name);
     }
     return files;
 }
 
 /// One source's code object for `arch`, with the kernels' flags and the caps table's instruction switches.
 pub fn compile(b: *std.Build, cc: Hipcc, arch: []const u8, source: std.Build.LazyPath, out: []const u8) LazyPath {
+    return compileSource(b, cc, arch, source, out, false);
+}
+
+/// `compile`, with `-O2 -fno-fast-math -fdenormal-fp-math=ieee` after the common flags when `strict`.
+fn compileSource(b: *std.Build, cc: Hipcc, arch: []const u8, source: LazyPath, out: []const u8, strict: bool) LazyPath {
     const c = caps.Caps.of(arch).?;
     const root = std.fs.path.dirname(std.fs.path.dirname(cc.resolved) orelse ".") orelse ".";
     const run = b.addSystemCommand(&.{ cc.name, "--genco" });
     run.addArgs(&flags);
+    if (strict) run.addArgs(&.{ "-O2", "-fno-fast-math", "-fdenormal-fp-math=ieee" });
     run.addArg(b.fmt("-DTF_WAVE={d}", .{c.wave}));
     run.addArg(b.fmt("-DTF_DOT2_F16={d}", .{@intFromBool(c.dot2_f16)}));
     run.addArg(b.fmt("-DTF_DOT2_BF16={d}", .{@intFromBool(c.dot2_bf16)}));

From 95ca0a5e1bac4d46954a4d38957e302a94fa802f Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Mon, 2 Nov 2026 23:21:23 +0000
Subject: [PATCH 04/10] test(hip-contract): a family registration list and an
 image selector

Family kernels need gates 1-4 with their own fixtures, so the runner
gets a registration list (contract_families.zig, empty today, one
entry a group) and `image <group> [smoke|full]`, which resolves the
group's entry and its gate-5 records before any GPU access, runs the
entry, then gate 5 on those records only, and exits 1 on a failed
kernel, a missed control, a limit problem or a run with no kernel;
the no-argument run also runs every entry and merges its tally. A
pure selector (contract_select.zig) holds that logic with host
tests, family records join hip_records.records from one list file,
and the limits checks gain a variant over a given record slice.
---
 zig/build/hip_contract.zig                    |  10 ++
 zig/src/core/contract.zig                     |   1 +
 zig/src/core/contract/hip_records.zig         |   3 +-
 .../core/contract/hip_records_families.zig    |  17 ++
 zig/tests/hip/contract_entry.zig              |  14 ++
 zig/tests/hip/contract_families.zig           |  11 ++
 zig/tests/hip/contract_limits.zig             |  23 ++-
 zig/tests/hip/contract_main.zig               |  64 +++++++-
 zig/tests/hip/contract_select.zig             | 150 ++++++++++++++++++
 9 files changed, 283 insertions(+), 10 deletions(-)
 create mode 100644 zig/src/core/contract/hip_records_families.zig
 create mode 100644 zig/tests/hip/contract_entry.zig
 create mode 100644 zig/tests/hip/contract_families.zig
 create mode 100644 zig/tests/hip/contract_select.zig

diff --git a/zig/build/hip_contract.zig b/zig/build/hip_contract.zig
index b1981b4a11..faca0ab700 100644
--- a/zig/build/hip_contract.zig
+++ b/zig/build/hip_contract.zig
@@ -26,6 +26,16 @@ pub fn steps(b: *std.Build, o: Options) void {
     pair_host.has_side_effects = true;
     o.test_step.dependOn(&pair_host.step);
     b.step("test-hip-pair-contract-host", "Paired projection simulation without uninitialized magnitude reads or GPU calls").dependOn(&pair_host.step);
+    const select_host = b.addRunArtifact(b.addTest(.{ .root_module = b.createModule(.{
+        .root_source_file = b.path("zig/tests/hip/contract_select.zig"),
+        .target = b.graph.host,
+        .optimize = .Debug,
+        .link_libc = true,
+        .imports = &.{ .{ .name = "hip", .module = o.host }, .{ .name = "core", .module = o.host.import_table.get("core").? } },
+    }) }));
+    o.test_step.dependOn(&select_host.step);
+    const select_what = "Which entry and records `image <group>` runs, and its exit status, on fixtures without a GPU";
+    b.step("test-hip-contract-select-host", select_what).dependOn(&select_host.step);
     const source = b.path("zig/kernels/hip/contract_controls.hip");
     const own = hip.compile(b, o.hipcc, o.arch, source, "contract_controls.hsaco");
     const exe = runner(b, o.target, o.optimize, o.gpu, controls(b, o.arch, own));
diff --git a/zig/src/core/contract.zig b/zig/src/core/contract.zig
index 3dba7a2fa9..64fc87b1eb 100644
--- a/zig/src/core/contract.zig
+++ b/zig/src/core/contract.zig
@@ -42,6 +42,7 @@ test {
     _ = hsaco;
     _ = hip_records;
     _ = @import("contract/hip_records_qwen3.zig");
+    _ = @import("contract/hip_records_families.zig");
     _ = @import("contract/fakes.zig");
 }
 
diff --git a/zig/src/core/contract/hip_records.zig b/zig/src/core/contract/hip_records.zig
index bf3e899ffc..9099a5a5b0 100644
--- a/zig/src/core/contract/hip_records.zig
+++ b/zig/src/core/contract/hip_records.zig
@@ -498,7 +498,8 @@ const own_records = [_]Record{
 
 pub const records = own_records ++ @import("hip_records_stream.zig").records ++ @import("hip_records_router.zig").records ++
     @import("hip_records_glm_ocr.zig").records ++ @import("hip_records_item.zig").records ++ @import("hip_records_qwen3.zig").records ++
-    @import("hip_records_columns.zig").records ++ @import("hip_records_vision.zig").records ++ @import("hip_records_dense_text.zig").records;
+    @import("hip_records_columns.zig").records ++ @import("hip_records_vision.zig").records ++ @import("hip_records_dense_text.zig").records ++
+    @import("hip_records_families.zig").all;
 
 /// Gate 5's planted HIP controls: a launch past its __launch_bounds__, and 72 KiB of LDS where a workgroup has 64.
 pub const controls = [_]Record{
diff --git a/zig/src/core/contract/hip_records_families.zig b/zig/src/core/contract/hip_records_families.zig
new file mode 100644
index 0000000000..81b3e95e51
--- /dev/null
+++ b/zig/src/core/contract/hip_records_families.zig
@@ -0,0 +1,17 @@
+//! The HIP families' gate-5 records, one array a family (`hip_records_<family>.zig`); hip_records.zig joins them.
+const std = @import("std");
+const Record = @import("hip_records.zig").Record;
+
+pub const records = .{};
+
+/// Every family's records as one array.
+pub const all = join: {
+    var joined: []const Record = &.{};
+    for (records) |list| joined = joined ++ list;
+    break :join joined[0..joined.len].*;
+};
+
+test "the shared HIP records stay main's; families only add theirs" {
+    const every = @import("hip_records.zig").records;
+    try std.testing.expectEqual(@as(usize, 173), every.len - all.len);
+}
diff --git a/zig/tests/hip/contract_entry.zig b/zig/tests/hip/contract_entry.zig
new file mode 100644
index 0000000000..acb6c8bf15
--- /dev/null
+++ b/zig/tests/hip/contract_entry.zig
@@ -0,0 +1,14 @@
+//! What a HIP family's contract entry is, apart from the list of them, so the pure selector needs no family code.
+const std = @import("std");
+const hip = @import("hip");
+const core = @import("core");
+
+/// How much of its gates 1-4 an entry runs: a quick check or every shape.
+pub const Profile = enum { smoke, full };
+
+/// A family's image, the gate-5 records it covers by launch name (none: every record of `group`), and its gates 1-4.
+pub const Entry = struct {
+    group: hip.kernels.Group,
+    records: []const []const u8 = &.{},
+    run: *const fn (std.mem.Allocator, std.Io, Profile) anyerror!core.contract.gates.Tally,
+};
diff --git a/zig/tests/hip/contract_families.zig b/zig/tests/hip/contract_families.zig
new file mode 100644
index 0000000000..380b16a62e
--- /dev/null
+++ b/zig/tests/hip/contract_families.zig
@@ -0,0 +1,11 @@
+//! The HIP families' contract entries, one line a family: `tf-hip-kernel-contract image <group>` runs one.
+pub const Profile = @import("contract_entry.zig").Profile;
+pub const Entry = @import("contract_entry.zig").Entry;
+
+pub const entries = [_]Entry{};
+
+comptime { // one entry a group, so `image <group>` never has to choose
+    for (entries, 0..) |e, i| for (entries[0..i]) |f| {
+        if (e.group == f.group) @compileError("two contract entries for group " ++ @tagName(e.group));
+    };
+}
diff --git a/zig/tests/hip/contract_limits.zig b/zig/tests/hip/contract_limits.zig
index 56ae3c123b..b6a13a4c69 100644
--- a/zig/tests/hip/contract_limits.zig
+++ b/zig/tests/hip/contract_limits.zig
@@ -34,11 +34,7 @@ fn file(a: std.mem.Allocator, dir: []const u8, chip: []const u8, name: []const u
 
 /// Gate 5 from the records alone: every launch on every chip, and each control red; returns the problems.
 pub fn table(a: std.mem.Allocator) !usize {
-    var n = try records.check(a, &records.records, false);
-    const word = if (n == 0) "PASS" else "FAIL";
-    const chips = "gfx1030/1100/1151/1201";
-    const say = "{s} limits: {d} registered HIP launches against {s}, {d} problems\n";
-    std.debug.print(say, .{ word, records.records.len, chips, n });
+    var n = try tableOf(a, &records.records);
     for (records.controls) |r| {
         const missed = try records.controlProblems(a, r, false);
         n += missed;
@@ -49,8 +45,23 @@ pub fn table(a: std.mem.Allocator) !usize {
     return n;
 }
 
+/// Gate 5 from `list` alone: each launch on every chip; returns the problems.
+pub fn tableOf(a: std.mem.Allocator, list: []const records.Record) !usize {
+    const n = try records.check(a, list, false);
+    const word = if (n == 0) "PASS" else "FAIL";
+    const chips = "gfx1030/1100/1151/1201";
+    const say = "{s} limits: {d} registered HIP launches against {s}, {d} problems\n";
+    std.debug.print(say, .{ word, list.len, chips, n });
+    return n;
+}
+
 /// Gate 5 on this GPU: the device's limits against its chip's row, each loaded kernel against its record.
 pub fn live(a: std.mem.Allocator, r: *const hip.Runtime, ctx: hip.Context, arch: []const u8) !usize {
+    return liveOf(a, r, ctx, arch, &all);
+}
+
+/// `live` over `list`: each record checked in its own code object, whatever image the selection spans.
+pub fn liveOf(a: std.mem.Allocator, r: *const hip.Runtime, ctx: hip.Context, arch: []const u8, list: []const records.Record) !usize {
     const at = for (limits.rdna, 0..) |t, i| {
         if (std.mem.eql(u8, t.name, arch)) break i;
     } else {
@@ -71,7 +82,7 @@ pub fn live(a: std.mem.Allocator, r: *const hip.Runtime, ctx: hip.Context, arch:
         const say = "{s} limits live {s} {s}: the device says {d}, the table {d}\n";
         std.debug.print(say, .{ word, arch, d[0], got, d[2] });
     }
-    for (all) |rec| {
+    for (list) |rec| {
         // a planted control is checked on every chip: only a registered launch may be unserved somewhere
         if (!rec.usage[at].served and !isControl(rec)) {
             std.debug.print("N/A  limits live {s}: not served on {s}\n", .{ rec.launch.name, arch });
diff --git a/zig/tests/hip/contract_main.zig b/zig/tests/hip/contract_main.zig
index e068b6f585..fd46fa7162 100644
--- a/zig/tests/hip/contract_main.zig
+++ b/zig/tests/hip/contract_main.zig
@@ -1,9 +1,11 @@
-//! tf-hip-kernel-contract: HIP kernels through gates 1-4, the controls and gate 5; `limits DIR` audits code objects.
+//! tf-hip-kernel-contract: HIP kernels through gates 1-4, controls and gate 5; `image`: one family; `limits DIR`.
 const std = @import("std");
 const hip = @import("hip");
 const core = @import("core");
 const controls = @import("hip_controls");
 const gate5 = @import("contract_limits.zig");
+const families = @import("contract_families.zig");
+const select = @import("contract_select.zig");
 const rows_k = @import("contract_rows.zig");
 const wide_k = @import("contract_rms_wide.zig");
 const stream_k = @import("contract_stream.zig");
@@ -40,8 +42,10 @@ pub fn main(init: std.process.Init) !u8 {
     const a = init.gpa;
     const args = try init.minimal.args.toSlice(init.arena.allocator());
     if (args.len == 1) return if (try run(a, init.io) > 0) 1 else 0;
+    const image_args = args.len == 3 or args.len == 4;
+    if (image_args and std.mem.eql(u8, args[1], "image")) return image(a, init.io, args[2], if (args.len == 4) args[3] else "full");
     if (args.len != 3 or !std.mem.eql(u8, args[1], "limits")) {
-        std.debug.print("usage: tf-hip-kernel-contract [limits DIR]\n", .{});
+        std.debug.print("usage: tf-hip-kernel-contract [image GROUP [smoke|full] | limits DIR]\n", .{});
         return 2;
     }
     const n = try gate5.objects(a, init.io, args[2]);
@@ -54,11 +58,64 @@ pub fn main(init: std.process.Init) !u8 {
     return 0;
 }
 
+/// `image GROUP [smoke|full]`: the entry's own gates 1-4, then gate 5 on its records; returns the exit status.
+fn image(a: std.mem.Allocator, io: std.Io, group_name: []const u8, profile_name: []const u8) !u8 {
+    const profile = std.meta.stringToEnum(families.Profile, profile_name) orelse {
+        std.debug.print("FAIL image {s}: the profile is smoke or full, not {s}\n", .{ group_name, profile_name });
+        return 2;
+    };
+    var bad: []const u8 = "";
+    const sel = select.resolve(a, &families.entries, &core.contract.hip_records.records, group_name, &bad) catch |err| {
+        const why = switch (err) {
+            error.OutOfMemory => return err,
+            error.UnknownGroup => "no kernel group has that name",
+            error.NoEntry => "the group has no entry in contract_families.zig",
+            error.MissingRecord => "the entry names a record that is not registered",
+            error.AmbiguousRecord => "the entry names a record that is registered more than once",
+            error.EmptySelection => "the entry selects no gate-5 record",
+        };
+        std.debug.print("FAIL image {s}: {s} ({s})\n", .{ group_name, why, bad });
+        if (err == error.UnknownGroup) {
+            std.debug.print("groups:", .{});
+            for (@typeInfo(hip.kernels.Group).@"enum".field_names) |name| std.debug.print(" {s}", .{name});
+            std.debug.print("\n", .{});
+        }
+        return 2;
+    };
+    defer a.free(sel.records);
+    var built: usize = hip.kernels.images[@backingInt(sel.entry.group)].len;
+    _ = &built; // runtime-known, so a stub build still compiles all that follows
+    if (built == 0) {
+        std.debug.print("FAIL image {s}: built without hipcc, so no gate is proven\n", .{group_name});
+        return 1;
+    }
+    var r = try hip.Runtime.open();
+    defer r.close();
+    var ctx = try hip.Context.init(&r, 0);
+    defer ctx.deinit();
+    var arch_buffer: [256]u8 = undefined;
+    const arch = try hip.Device.archName(&r, 0, &arch_buffer);
+    if (!std.mem.eql(u8, arch, hip.kernels.arch)) {
+        const say = "FAIL image {s}: this GPU is {s}, the kernels are built for {s}\n";
+        std.debug.print(say, .{ group_name, arch, hip.kernels.arch });
+        return 1;
+    }
+    const t0 = ms(io);
+    const t = try sel.entry.run(a, io, profile);
+    const t1 = ms(io);
+    const lim = try gate5.tableOf(a, sel.records) + try gate5.liveOf(a, &r, ctx, arch, sel.records);
+    const t2 = ms(io);
+    const tally = .{ t.kernels, t.passed, t.failed, t.known, t.controls, t.caught, t.missed, lim, t1 - t0, t2 - t1 };
+    std.debug.print("SUMMARY {s} HIP {t}: {d} kernels: {d} pass gates 1-4, {d} fail, {d} known findings; {d} controls: " ++
+        "{d} red as planted, {d} missed; {d} limit problems; gates {d} ms, limits {d} ms\n", .{ group_name, profile } ++ tally);
+    return select.verdict(t, lim);
+}
+
 fn ms(io: std.Io) i96 {
     return @divTrunc(std.Io.Clock.awake.now(io).toNanoseconds(), std.time.ns_per_ms);
 }
 
-/// Gates 1-4 on every registered HIP kernel and each control at its own gate, then gate 5; returns the problems.
+/// Gates 1-4 on every registered HIP kernel, each control and every family entry, then gate 5; returns the problems.
 fn run(a: std.mem.Allocator, io: std.Io) !usize {
     var built: [2]usize = .{ hip.kernels.images[0].len, controls.bytes.len };
     _ = &built; // runtime-known, so a stub build still compiles all that follows
@@ -233,6 +290,7 @@ fn run(a: std.mem.Allocator, io: std.Io) !usize {
     for (&norms) |*op| try kernels.append(a, op.kernel());
     const t0 = ms(io);
     try g.runAll(a, kernels.items, 0x6b63, &t);
+    t = select.merge(t, try select.runEntries(a, io, &families.entries, .full)); // each family's own gates 1-4
     const t1 = ms(io);
     const lim = try gate5.table(a) + try gate5.live(a, &r, ctx, arch);
     const t2 = ms(io);
diff --git a/zig/tests/hip/contract_select.zig b/zig/tests/hip/contract_select.zig
new file mode 100644
index 0000000000..359f373d86
--- /dev/null
+++ b/zig/tests/hip/contract_select.zig
@@ -0,0 +1,150 @@
+//! `tf-hip-kernel-contract image <group>`: the entry and gate-5 records it runs, and its exit status; no HIP calls.
+const std = @import("std");
+const hip = @import("hip");
+const core = @import("core");
+const families = @import("contract_entry.zig"); // the entry types only: no family code on the host
+
+const Entry = families.Entry;
+const Record = core.contract.hip_records.Record;
+const Tally = core.contract.gates.Tally;
+
+pub const Error = error{ UnknownGroup, NoEntry, MissingRecord, AmbiguousRecord, EmptySelection, OutOfMemory };
+
+/// An image's entry and the records its gate 5 checks; `records` belongs to the allocator `resolve` was given.
+pub const Selection = struct { entry: Entry, records: []Record };
+
+/// `group_name`'s entry with each of its record names matched once (none: the group's records); `bad` names a fault.
+pub fn resolve(
+    a: std.mem.Allocator,
+    entries: []const Entry,
+    records: []const Record,
+    group_name: []const u8,
+    bad: *[]const u8,
+) Error!Selection {
+    bad.* = group_name;
+    const group = std.meta.stringToEnum(hip.kernels.Group, group_name) orelse return error.UnknownGroup;
+    const entry = for (entries) |e| {
+        if (e.group == group) break e;
+    } else return error.NoEntry;
+    var out: std.ArrayList(Record) = .empty;
+    errdefer out.deinit(a);
+    if (entry.records.len == 0) for (records) |r| {
+        if (std.mem.eql(u8, r.launch.image, group_name)) try out.append(a, r);
+    };
+    for (entry.records) |name| {
+        bad.* = name;
+        var found: ?Record = null;
+        for (records) |r| if (std.mem.eql(u8, r.launch.name, name)) {
+            if (found != null) return error.AmbiguousRecord;
+            found = r;
+        };
+        try out.append(a, found orelse return error.MissingRecord);
+    }
+    bad.* = group_name;
+    if (out.items.len == 0) return error.EmptySelection;
+    return .{ .entry = entry, .records = try out.toOwnedSlice(a) };
+}
+
+/// The exit status: 1 when a kernel failed, a control was missed, gate 5 found a problem or no kernel ran, else 0.
+pub fn verdict(t: Tally, gate5_problems: usize) u8 {
+    return if (t.problems() > 0 or gate5_problems > 0 or t.kernels == 0) 1 else 0;
+}
+
+/// Two runs' tallies as one, field by field.
+pub fn merge(x: Tally, y: Tally) Tally {
+    var t = x;
+    inline for (@typeInfo(Tally).@"struct".field_names) |name| @field(t, name) += @field(y, name);
+    return t;
+}
+
+/// Every entry through its own gates 1-4 at `profile`, one tally for all of them (the no-argument run).
+pub fn runEntries(a: std.mem.Allocator, io: std.Io, entries: []const Entry, profile: families.Profile) !Tally {
+    var t: Tally = .{};
+    for (entries) |e| t = merge(t, try e.run(a, io, profile));
+    return t;
+}
+
+fn noRun(_: std.mem.Allocator, _: std.Io, _: families.Profile) anyerror!Tally {
+    return .{};
+}
+
+fn cleanRun(_: std.mem.Allocator, _: std.Io, p: families.Profile) anyerror!Tally {
+    return if (p == .full) .{ .kernels = 2, .passed = 2, .controls = 1, .caught = 1 } else .{};
+}
+
+fn missRun(_: std.mem.Allocator, _: std.Io, p: families.Profile) anyerror!Tally {
+    return if (p == .full) .{ .kernels = 1, .passed = 1, .controls = 1, .missed = 1 } else .{};
+}
+
+fn rec(name: []const u8, image: []const u8) Record {
+    return .{ .launch = .{ .name = name, .symbol = name, .image = image, .threads = 64 }, .usage = @splat(.{ .vgprs = 16 }) };
+}
+
+const fixture = [_]Record{ rec("a", "ops"), rec("b", "ops"), rec("fa", "prefill"), rec("twin", "act"), rec("twin", "act") };
+
+fn pick(entries: []const Entry, group_name: []const u8, bad: *[]const u8) Error!Selection {
+    return resolve(std.testing.allocator, entries, &fixture, group_name, bad);
+}
+
+test "contract select: an entry's names resolve across images, and no names take its group's records" {
+    var bad: []const u8 = "";
+    const across = try pick(&.{.{ .group = .ops, .records = &.{ "a", "fa" }, .run = noRun }}, "ops", &bad);
+    defer std.testing.allocator.free(across.records);
+    try std.testing.expectEqual(hip.kernels.Group.ops, across.entry.group);
+    try std.testing.expectEqual(@as(usize, 2), across.records.len);
+    try std.testing.expectEqualStrings("ops", across.records[0].launch.image);
+    try std.testing.expectEqualStrings("prefill", across.records[1].launch.image);
+    const own = try pick(&.{.{ .group = .ops, .run = noRun }}, "ops", &bad);
+    defer std.testing.allocator.free(own.records);
+    try std.testing.expectEqualStrings("a", own.records[0].launch.name);
+    try std.testing.expectEqualStrings("b", own.records[1].launch.name);
+    try std.testing.expectEqual(@as(usize, 2), own.records.len);
+}
+
+test "contract select: an unknown group, a group without an entry and an empty selection are errors" {
+    var bad: []const u8 = "";
+    const entries = [_]Entry{ .{ .group = .ops, .run = noRun }, .{ .group = .decode, .run = noRun } };
+    try std.testing.expectError(error.UnknownGroup, pick(&entries, "opz", &bad));
+    try std.testing.expectEqualStrings("opz", bad);
+    try std.testing.expectError(error.NoEntry, pick(&entries, "mamba", &bad));
+    try std.testing.expectEqualStrings("mamba", bad);
+    try std.testing.expectError(error.EmptySelection, pick(&entries, "decode", &bad));
+    try std.testing.expectEqualStrings("decode", bad);
+}
+
+test "contract select: every explicit name must match exactly one record" {
+    var bad: []const u8 = "";
+    const mixed = [_]Entry{.{ .group = .ops, .records = &.{ "a", "misspelt", "fa" }, .run = noRun }};
+    try std.testing.expectError(error.MissingRecord, pick(&mixed, "ops", &bad));
+    try std.testing.expectEqualStrings("misspelt", bad);
+    const twin = [_]Entry{.{ .group = .act, .records = &.{ "a", "twin" }, .run = noRun }};
+    try std.testing.expectError(error.AmbiguousRecord, pick(&twin, "act", &bad));
+    try std.testing.expectEqualStrings("twin", bad);
+}
+
+test "contract select: the verdict fails a failed kernel, a missed control, a gate-5 problem and an empty run" {
+    const clean: Tally = .{ .kernels = 2, .passed = 2, .controls = 1, .caught = 1, .known = 1 };
+    try std.testing.expectEqual(@as(u8, 0), verdict(clean, 0));
+    try std.testing.expectEqual(@as(u8, 1), verdict(.{ .kernels = 1, .failed = 1 }, 0));
+    try std.testing.expectEqual(@as(u8, 1), verdict(.{ .kernels = 1, .passed = 1, .controls = 1, .missed = 1 }, 0));
+    try std.testing.expectEqual(@as(u8, 1), verdict(clean, 3));
+    try std.testing.expectEqual(@as(u8, 1), verdict(.{}, 0));
+    try std.testing.expectEqual(@as(u8, 1), verdict(.{ .controls = 1, .caught = 1 }, 0));
+    const missed: Tally = .{ .kernels = 1, .passed = 1, .controls = 1, .missed = 1 };
+    try std.testing.expectEqual(@as(u8, 0), verdict(merge(clean, clean), 0));
+    try std.testing.expectEqual(@as(u8, 1), verdict(merge(clean, missed), 0));
+    try std.testing.expectEqual(@as(usize, 3), merge(clean, missed).kernels);
+}
+
+test "contract select: the no-argument run counts every entry's kernels and controls" {
+    const two = [_]Entry{ .{ .group = .ops, .run = cleanRun }, .{ .group = .act, .run = missRun } };
+    const t = try runEntries(std.testing.allocator, std.testing.io, &two, .full);
+    try std.testing.expectEqual(@as(usize, 3), t.kernels);
+    try std.testing.expectEqual(@as(usize, 3), t.passed);
+    try std.testing.expectEqual(@as(usize, 2), t.controls);
+    try std.testing.expectEqual(@as(usize, 1), t.caught);
+    try std.testing.expectEqual(@as(usize, 1), t.missed);
+    try std.testing.expectEqual(@as(u8, 1), verdict(t, 0));
+    const none = try runEntries(std.testing.allocator, std.testing.io, &.{}, .full);
+    try std.testing.expectEqual(Tally{}, none);
+}

From 8a950998ca7018057e3aa1d8ffc92122d952de6a Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Mon, 2 Nov 2026 23:27:15 +0000
Subject: [PATCH 05/10] fix(hip-contract): run registered families in the full
 contract

The full run merged the family entries' tallies but judged them only
as one sum, so an entry that ran no kernel could hide beside a clean
one. Each entry now prints an ENTRY line and is judged on its own
tally (failed kernels, missed controls, and one problem for an empty
run); the full run adds those problems to today's and to gate 5, and
with no entry registered its output and exit are unchanged.
---
 zig/tests/hip/contract_main.zig   |  6 +++--
 zig/tests/hip/contract_select.zig | 45 ++++++++++++++++++++++++++-----
 2 files changed, 42 insertions(+), 9 deletions(-)

diff --git a/zig/tests/hip/contract_main.zig b/zig/tests/hip/contract_main.zig
index fd46fa7162..46225a23e9 100644
--- a/zig/tests/hip/contract_main.zig
+++ b/zig/tests/hip/contract_main.zig
@@ -290,12 +290,14 @@ fn run(a: std.mem.Allocator, io: std.Io) !usize {
     for (&norms) |*op| try kernels.append(a, op.kernel());
     const t0 = ms(io);
     try g.runAll(a, kernels.items, 0x6b63, &t);
-    t = select.merge(t, try select.runEntries(a, io, &families.entries, .full)); // each family's own gates 1-4
+    const own = t.problems();
+    const fam = try select.runEntries(a, io, &families.entries, .full); // each family's own gates 1-4
+    t = select.merge(t, fam.tally);
     const t1 = ms(io);
     const lim = try gate5.table(a) + try gate5.live(a, &r, ctx, arch);
     const t2 = ms(io);
     const tally = .{ t.kernels, t.passed, t.failed, t.known, t.controls, t.caught, t.missed, lim, t1 - t0, t2 - t1 };
     std.debug.print("SUMMARY {s} {d} kernels: {d} pass gates 1-4, {d} fail, {d} known findings; {d} controls: " ++
         "{d} red as planted, {d} missed; {d} limit problems; gates {d} ms, limits {d} ms\n", .{arch} ++ tally);
-    return t.problems() + lim;
+    return own + fam.problems + lim;
 }
diff --git a/zig/tests/hip/contract_select.zig b/zig/tests/hip/contract_select.zig
index 359f373d86..965fab1d22 100644
--- a/zig/tests/hip/contract_select.zig
+++ b/zig/tests/hip/contract_select.zig
@@ -57,11 +57,24 @@ pub fn merge(x: Tally, y: Tally) Tally {
     return t;
 }
 
-/// Every entry through its own gates 1-4 at `profile`, one tally for all of them (the no-argument run).
-pub fn runEntries(a: std.mem.Allocator, io: std.Io, entries: []const Entry, profile: families.Profile) !Tally {
-    var t: Tally = .{};
-    for (entries) |e| t = merge(t, try e.run(a, io, profile));
-    return t;
+/// One entry's problems on its own tally: failed kernels, missed controls, and one when it ran no kernel.
+pub fn entryProblems(t: Tally) usize {
+    return t.problems() + @intFromBool(t.kernels == 0);
+}
+
+/// The entries' tallies summed, and their problems each judged on its own entry's tally.
+pub const Entries = struct { tally: Tally = .{}, problems: usize = 0 };
+
+/// Every entry through its own gates 1-4 at `profile`, one ENTRY line each (the no-argument run).
+pub fn runEntries(a: std.mem.Allocator, io: std.Io, entries: []const Entry, profile: families.Profile) !Entries {
+    var out: Entries = .{};
+    for (entries) |e| {
+        const t = try e.run(a, io, profile);
+        std.debug.print("ENTRY {t}: {d} kernels, {d} fail, {d} missed\n", .{ e.group, t.kernels, t.failed, t.missed });
+        out.tally = merge(out.tally, t);
+        out.problems += entryProblems(t);
+    }
+    return out;
 }
 
 fn noRun(_: std.mem.Allocator, _: std.Io, _: families.Profile) anyerror!Tally {
@@ -72,6 +85,10 @@ fn cleanRun(_: std.mem.Allocator, _: std.Io, p: families.Profile) anyerror!Tally
     return if (p == .full) .{ .kernels = 2, .passed = 2, .controls = 1, .caught = 1 } else .{};
 }
 
+fn failRun(_: std.mem.Allocator, _: std.Io, p: families.Profile) anyerror!Tally {
+    return if (p == .full) .{ .kernels = 3, .passed = 1, .failed = 2, .known = 1 } else .{};
+}
+
 fn missRun(_: std.mem.Allocator, _: std.Io, p: families.Profile) anyerror!Tally {
     return if (p == .full) .{ .kernels = 1, .passed = 1, .controls = 1, .missed = 1 } else .{};
 }
@@ -138,7 +155,9 @@ test "contract select: the verdict fails a failed kernel, a missed control, a ga
 
 test "contract select: the no-argument run counts every entry's kernels and controls" {
     const two = [_]Entry{ .{ .group = .ops, .run = cleanRun }, .{ .group = .act, .run = missRun } };
-    const t = try runEntries(std.testing.allocator, std.testing.io, &two, .full);
+    const run = try runEntries(std.testing.allocator, std.testing.io, &two, .full);
+    const t = run.tally;
+    try std.testing.expectEqual(@as(usize, 1), run.problems);
     try std.testing.expectEqual(@as(usize, 3), t.kernels);
     try std.testing.expectEqual(@as(usize, 3), t.passed);
     try std.testing.expectEqual(@as(usize, 2), t.controls);
@@ -146,5 +165,17 @@ test "contract select: the no-argument run counts every entry's kernels and cont
     try std.testing.expectEqual(@as(usize, 1), t.missed);
     try std.testing.expectEqual(@as(u8, 1), verdict(t, 0));
     const none = try runEntries(std.testing.allocator, std.testing.io, &.{}, .full);
-    try std.testing.expectEqual(Tally{}, none);
+    try std.testing.expectEqual(Tally{}, none.tally);
+    try std.testing.expectEqual(@as(usize, 0), none.problems);
+}
+
+test "contract select: each entry is judged on its own tally, so an empty run is a problem beside a clean one" {
+    try std.testing.expectEqual(@as(usize, 0), entryProblems(try cleanRun(std.testing.allocator, std.testing.io, .full)));
+    try std.testing.expectEqual(@as(usize, 2), entryProblems(try failRun(std.testing.allocator, std.testing.io, .full)));
+    try std.testing.expectEqual(@as(usize, 1), entryProblems(try noRun(std.testing.allocator, std.testing.io, .full)));
+    const three = [_]Entry{ .{ .group = .ops, .run = cleanRun }, .{ .group = .act, .run = failRun }, .{ .group = .mamba, .run = noRun } };
+    const run = try runEntries(std.testing.allocator, std.testing.io, &three, .full);
+    try std.testing.expectEqual(@as(usize, 3), run.problems); // 0 + 2 failed + 1 for the empty run
+    try std.testing.expectEqual(@as(usize, 5), run.tally.kernels); // the merged tally alone would miss the empty run
+    try std.testing.expectEqual(@as(usize, 2), run.tally.problems());
 }

From 3027d53ac2affa839102f1efe03721527d022809 Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Mon, 2 Nov 2026 23:31:17 +0000
Subject: [PATCH 06/10] fix(hip-contract): check every family's records before
 the full run

The full contract ran each registered family's gates 1-4 without
resolving its gate-5 record selection, so a family with missing or
empty records could pass the default run while `image` refused it.
The full run now validates every entry first, before any GPU access,
and exits 2 with the reason, as `image` does.
---
 zig/tests/hip/contract_main.zig   | 21 +++++++++++----------
 zig/tests/hip/contract_select.zig | 30 ++++++++++++++++++++++++++++++
 2 files changed, 41 insertions(+), 10 deletions(-)

diff --git a/zig/tests/hip/contract_main.zig b/zig/tests/hip/contract_main.zig
index 46225a23e9..26aa0cd3d0 100644
--- a/zig/tests/hip/contract_main.zig
+++ b/zig/tests/hip/contract_main.zig
@@ -41,7 +41,15 @@ const g = core.contract.gates;
 pub fn main(init: std.process.Init) !u8 {
     const a = init.gpa;
     const args = try init.minimal.args.toSlice(init.arena.allocator());
-    if (args.len == 1) return if (try run(a, init.io) > 0) 1 else 0;
+    if (args.len == 1) {
+        var bad: []const u8 = "";
+        select.validateAll(a, &families.entries, &core.contract.hip_records.records, &bad) catch |err| {
+            if (err == error.OutOfMemory) return err;
+            std.debug.print("FAIL kernel-contract: {s} ({s})\n", .{ select.explain(err), bad });
+            return 2;
+        };
+        return if (try run(a, init.io) > 0) 1 else 0;
+    }
     const image_args = args.len == 3 or args.len == 4;
     if (image_args and std.mem.eql(u8, args[1], "image")) return image(a, init.io, args[2], if (args.len == 4) args[3] else "full");
     if (args.len != 3 or !std.mem.eql(u8, args[1], "limits")) {
@@ -66,15 +74,8 @@ fn image(a: std.mem.Allocator, io: std.Io, group_name: []const u8, profile_name:
     };
     var bad: []const u8 = "";
     const sel = select.resolve(a, &families.entries, &core.contract.hip_records.records, group_name, &bad) catch |err| {
-        const why = switch (err) {
-            error.OutOfMemory => return err,
-            error.UnknownGroup => "no kernel group has that name",
-            error.NoEntry => "the group has no entry in contract_families.zig",
-            error.MissingRecord => "the entry names a record that is not registered",
-            error.AmbiguousRecord => "the entry names a record that is registered more than once",
-            error.EmptySelection => "the entry selects no gate-5 record",
-        };
-        std.debug.print("FAIL image {s}: {s} ({s})\n", .{ group_name, why, bad });
+        if (err == error.OutOfMemory) return err;
+        std.debug.print("FAIL image {s}: {s} ({s})\n", .{ group_name, select.explain(err), bad });
         if (err == error.UnknownGroup) {
             std.debug.print("groups:", .{});
             for (@typeInfo(hip.kernels.Group).@"enum".field_names) |name| std.debug.print(" {s}", .{name});
diff --git a/zig/tests/hip/contract_select.zig b/zig/tests/hip/contract_select.zig
index 965fab1d22..9076511d37 100644
--- a/zig/tests/hip/contract_select.zig
+++ b/zig/tests/hip/contract_select.zig
@@ -57,6 +57,23 @@ pub fn merge(x: Tally, y: Tally) Tally {
     return t;
 }
 
+/// Every entry's selection resolves, so no run checks a family without its gate-5 records; `bad` names the fault.
+pub fn validateAll(a: std.mem.Allocator, entries: []const Entry, records: []const Record, bad: *[]const u8) Error!void {
+    for (entries) |e| a.free((try resolve(a, entries, records, @tagName(e.group), bad)).records);
+}
+
+/// Why a selection is refused, in words.
+pub fn explain(err: Error) []const u8 {
+    return switch (err) {
+        error.OutOfMemory => "out of memory",
+        error.UnknownGroup => "no kernel group has that name",
+        error.NoEntry => "the group has no entry in contract_families.zig",
+        error.MissingRecord => "the entry names a record that is not registered",
+        error.AmbiguousRecord => "the entry names a record that is registered more than once",
+        error.EmptySelection => "the entry selects no gate-5 record",
+    };
+}
+
 /// One entry's problems on its own tally: failed kernels, missed controls, and one when it ran no kernel.
 pub fn entryProblems(t: Tally) usize {
     return t.problems() + @intFromBool(t.kernels == 0);
@@ -179,3 +196,16 @@ test "contract select: each entry is judged on its own tally, so an empty run is
     try std.testing.expectEqual(@as(usize, 5), run.tally.kernels); // the merged tally alone would miss the empty run
     try std.testing.expectEqual(@as(usize, 2), run.tally.problems());
 }
+
+test "contract select: the full run validates every entry first; an empty registry is valid" {
+    var bad: []const u8 = "";
+    try validateAll(std.testing.allocator, &.{}, &fixture, &bad);
+    const good = [_]Entry{ .{ .group = .ops, .records = &.{ "a", "fa" }, .run = cleanRun }, .{ .group = .prefill, .run = cleanRun } };
+    try validateAll(std.testing.allocator, &good, &fixture, &bad);
+    const empty = [_]Entry{ .{ .group = .ops, .run = cleanRun }, .{ .group = .decode, .run = cleanRun } };
+    try std.testing.expectError(error.EmptySelection, validateAll(std.testing.allocator, &empty, &fixture, &bad));
+    try std.testing.expectEqualStrings("decode", bad);
+    const missing = [_]Entry{.{ .group = .ops, .records = &.{"misspelt"}, .run = cleanRun }};
+    try std.testing.expectError(error.MissingRecord, validateAll(std.testing.allocator, &missing, &fixture, &bad));
+    try std.testing.expectEqualStrings("misspelt", bad);
+}

From ab9266597352b28c05b8c91355fa337bd973348a Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Tue, 3 Nov 2026 13:09:21 +0000
Subject: [PATCH 07/10] test(hip-contract): register GLM-OCR as the first
 contract family

GLM-OCR's pointwise and norm kernels ran inline in the full contract
run, so `image glm_ocr` had no entry and the family list was empty.
They now form GLM-OCR's entry, which opens its own runtime and runs
the same 18 gate cases with the same seed, and its 12 gate-5 records
move from the shared list into the family list. The full run still
covers every kernel and all 167 records; `image glm_ocr` checks the
family alone, and a test pins the 167 total and 155 shared records.
---
 zig/src/core/contract/hip_records.zig         |  5 ++-
 .../core/contract/hip_records_families.zig    |  9 +++--
 zig/tests/hip/contract_families.zig           |  4 ++-
 zig/tests/hip/contract_glm_ocr.zig            | 35 +++++++++++++++++++
 zig/tests/hip/contract_main.zig               | 13 -------
 5 files changed, 46 insertions(+), 20 deletions(-)
 create mode 100644 zig/tests/hip/contract_glm_ocr.zig

diff --git a/zig/src/core/contract/hip_records.zig b/zig/src/core/contract/hip_records.zig
index 9099a5a5b0..4dd38d02d4 100644
--- a/zig/src/core/contract/hip_records.zig
+++ b/zig/src/core/contract/hip_records.zig
@@ -497,9 +497,8 @@ const own_records = [_]Record{
     @import("hip_records_score.zig").records ++ @import("hip_records_embed.zig").records;
 
 pub const records = own_records ++ @import("hip_records_stream.zig").records ++ @import("hip_records_router.zig").records ++
-    @import("hip_records_glm_ocr.zig").records ++ @import("hip_records_item.zig").records ++ @import("hip_records_qwen3.zig").records ++
-    @import("hip_records_columns.zig").records ++ @import("hip_records_vision.zig").records ++ @import("hip_records_dense_text.zig").records ++
-    @import("hip_records_families.zig").all;
+    @import("hip_records_item.zig").records ++ @import("hip_records_qwen3.zig").records ++ @import("hip_records_columns.zig").records ++
+    @import("hip_records_vision.zig").records ++ @import("hip_records_dense_text.zig").records ++ @import("hip_records_families.zig").all;
 
 /// Gate 5's planted HIP controls: a launch past its __launch_bounds__, and 72 KiB of LDS where a workgroup has 64.
 pub const controls = [_]Record{
diff --git a/zig/src/core/contract/hip_records_families.zig b/zig/src/core/contract/hip_records_families.zig
index 81b3e95e51..43fffd7901 100644
--- a/zig/src/core/contract/hip_records_families.zig
+++ b/zig/src/core/contract/hip_records_families.zig
@@ -2,7 +2,9 @@
 const std = @import("std");
 const Record = @import("hip_records.zig").Record;
 
-pub const records = .{};
+pub const records = .{
+    @import("hip_records_glm_ocr.zig").records,
+};
 
 /// Every family's records as one array.
 pub const all = join: {
@@ -11,7 +13,8 @@ pub const all = join: {
     break :join joined[0..joined.len].*;
 };
 
-test "the shared HIP records stay main's; families only add theirs" {
+test "the HIP records stay main's; the family lists only regroup GLM-OCR's" {
     const every = @import("hip_records.zig").records;
-    try std.testing.expectEqual(@as(usize, 173), every.len - all.len);
+    try std.testing.expectEqual(@as(usize, 182), every.len); // main's records, GLM-OCR's 12 among them
+    try std.testing.expectEqual(@as(usize, 170), every.len - all.len);
 }
diff --git a/zig/tests/hip/contract_families.zig b/zig/tests/hip/contract_families.zig
index 380b16a62e..b881aa950a 100644
--- a/zig/tests/hip/contract_families.zig
+++ b/zig/tests/hip/contract_families.zig
@@ -2,7 +2,9 @@
 pub const Profile = @import("contract_entry.zig").Profile;
 pub const Entry = @import("contract_entry.zig").Entry;
 
-pub const entries = [_]Entry{};
+pub const entries = [_]Entry{
+    .{ .group = .glm_ocr, .run = @import("contract_glm_ocr.zig").run },
+};
 
 comptime { // one entry a group, so `image <group>` never has to choose
     for (entries, 0..) |e, i| for (entries[0..i]) |f| {
diff --git a/zig/tests/hip/contract_glm_ocr.zig b/zig/tests/hip/contract_glm_ocr.zig
new file mode 100644
index 0000000000..422a929145
--- /dev/null
+++ b/zig/tests/hip/contract_glm_ocr.zig
@@ -0,0 +1,35 @@
+//! GLM-OCR's contract entry: its pointwise and norm kernels from the glm_ocr image through gates 1-4.
+const std = @import("std");
+const hip = @import("hip");
+const core = @import("core");
+const Profile = @import("contract_entry.zig").Profile;
+const ocr_ops = @import("glm_ocr_ops").Ops;
+const ocr_points = @import("contract_glm_ocr_ops.zig");
+const ocr_norms = @import("contract_glm_ocr_norm.zig");
+
+const g = core.contract.gates;
+
+/// Every GLM-OCR kernel at its fixed shapes; both profiles run the same cases.
+pub fn run(a: std.mem.Allocator, _: std.Io, _: Profile) anyerror!g.Tally {
+    var r = try hip.Runtime.open();
+    defer r.close();
+    var ctx = try hip.Context.init(&r, 0);
+    defer ctx.deinit();
+    var s = try hip.Stream.init(&r);
+    defer s.deinit();
+    var ocr = try ocr_ops.load(&r);
+    defer ocr.unload();
+    var kernels: std.ArrayList(g.Kernel) = .empty;
+    defer kernels.deinit(a);
+    var pointwise: [std.enums.values(ocr_points.Which).len]ocr_points.Pointwise = undefined;
+    for (&pointwise, std.enums.values(ocr_points.Which)) |*op, which_op| op.* = .init(a, &r, s, &ocr, which_op);
+    defer for (&pointwise) |*op| op.deinit();
+    for (&pointwise) |*op| try kernels.append(a, op.kernel());
+    var norms: [std.enums.values(ocr_norms.Which).len]ocr_norms.Norm = undefined;
+    for (&norms, std.enums.values(ocr_norms.Which)) |*op, which_op| op.* = .init(a, &r, s, &ocr, which_op);
+    defer for (&norms) |*op| op.deinit();
+    for (&norms) |*op| try kernels.append(a, op.kernel());
+    var t: g.Tally = .{};
+    try g.runAll(a, kernels.items, 0x6b63, &t);
+    return t;
+}
diff --git a/zig/tests/hip/contract_main.zig b/zig/tests/hip/contract_main.zig
index 26aa0cd3d0..a2252abf96 100644
--- a/zig/tests/hip/contract_main.zig
+++ b/zig/tests/hip/contract_main.zig
@@ -26,9 +26,6 @@ const dense12_k = @import("contract_dense_wmma12.zig");
 const gate_k = @import("contract_moe_gate.zig");
 const router_k = @import("contract_router.zig");
 const kit = @import("contract_kit.zig");
-const ocr_ops = @import("glm_ocr_ops").Ops;
-const ocr_points = @import("contract_glm_ocr_ops.zig");
-const ocr_norms = @import("contract_glm_ocr_norm.zig");
 const item_k = @import("contract_wmma_item.zig");
 const embed_k = @import("contract_embed.zig");
 const columns_k = @import("contract_columns.zig");
@@ -279,16 +276,6 @@ fn run(a: std.mem.Allocator, io: std.Io) !usize {
     defer for (&text) |*v| v.deinit();
     for (&text) |*v| try kernels.append(a, v.kernel());
     var t: g.Tally = .{};
-    var ocr = try ocr_ops.load(&r);
-    defer ocr.unload();
-    var pointwise: [std.enums.values(ocr_points.Which).len]ocr_points.Pointwise = undefined;
-    for (&pointwise, std.enums.values(ocr_points.Which)) |*op, which_op| op.* = .init(a, &r, s, &ocr, which_op);
-    defer for (&pointwise) |*op| op.deinit();
-    for (&pointwise) |*op| try kernels.append(a, op.kernel());
-    var norms: [std.enums.values(ocr_norms.Which).len]ocr_norms.Norm = undefined;
-    for (&norms, std.enums.values(ocr_norms.Which)) |*op, which_op| op.* = .init(a, &r, s, &ocr, which_op);
-    defer for (&norms) |*op| op.deinit();
-    for (&norms) |*op| try kernels.append(a, op.kernel());
     const t0 = ms(io);
     try g.runAll(a, kernels.items, 0x6b63, &t);
     const own = t.problems();

From c6a09fe786ebe6e287452234f8e9fc71a3f0b56e Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Wed, 4 Nov 2026 09:34:10 +0000
Subject: [PATCH 08/10] refactor(native-hip): fold the served routes into the
 family offer

Main's Serves read text_generation, embeddings, decides, rerank,
readouts and window from a family's Loaded field by field, beside the
Offer this branch reads for images and caches, so the host had two
probes for one question. The six fields move into hip_serve.Offer
with the same defaults; Qwen3 returns its routes and window in its
offer, Qwen3.5 its readouts, and hip_families.generates keeps main's
rule that a family answering readouts generates no text. The Qwen3
module imports hip_serve, structures needs text generation and
first_masked, and every family on main gets the values it had.
---
 zig/build/hip.zig                             |  1 +
 zig/src/families/qwen3/native.zig             | 26 +++---
 .../qwen3_5/backend/hip/engine/native.zig     |  9 ++-
 zig/src/hip_serve/offer.zig                   | 16 +++-
 zig/src/native/hip.zig                        | 79 ++++++-------------
 zig/src/native/hip_families.zig               | 10 +++
 6 files changed, 64 insertions(+), 77 deletions(-)

diff --git a/zig/build/hip.zig b/zig/build/hip.zig
index 55110e210d..53a45b96be 100644
--- a/zig/build/hip.zig
+++ b/zig/build/hip.zig
@@ -325,6 +325,7 @@ pub fn engineModules(b: *std.Build, target: std.Build.ResolvedTarget, optimize:
             .{ .name = "core", .module = core },
             .{ .name = "lanes", .module = lanes },
             .{ .name = "engine_api", .module = api },
+            .{ .name = "hip_serve", .module = serve },
         },
     });
     const dots = @import("dots_mocr.zig").hipModule(b, hip);
diff --git a/zig/src/families/qwen3/native.zig b/zig/src/families/qwen3/native.zig
index b9658f3181..1833047d02 100644
--- a/zig/src/families/qwen3/native.zig
+++ b/zig/src/families/qwen3/native.zig
@@ -3,6 +3,7 @@ const std = @import("std");
 const hip = @import("hip");
 const lanes = @import("lanes");
 const api = @import("engine_api");
+const serve = @import("hip_serve");
 const engine = @import("engine.zig");
 const prompt = @import("prompt.zig");
 const score = @import("score.zig");
@@ -31,7 +32,7 @@ pub fn refusal(o: api.Open) ?[]const u8 {
 pub const open = openEngine;
 pub const follow = followRank;
 
-/// What the server drives, as the other HIP families' `Loaded`, plus what this engine serves.
+/// What the server drives, as the other HIP families' `Loaded`; its offer says what this engine serves.
 pub const Loaded = struct {
     backend: lanes.backend.Backend,
     facts: lanes.Model,
@@ -40,15 +41,8 @@ pub const Loaded = struct {
     deinit: *const fn (*anyopaque) void,
     streams: usize,
     startup: []const u8,
-    /// No chat or completions: the server refuses the generation routes.
-    text_generation: bool = false,
-    embeddings: ?api.EmbeddingInfo = null,
-    /// The window the engine was sized to (the asked one, or less where the budget shrank it).
-    window: ?usize = null,
-    /// A reranker: the engine scores labels (the score hook) and serves no embeddings.
-    decides: bool = false,
-    /// The reranker's label tokens, from `1_LogitScore/config.json`.
-    rerank: ?api.Rerank = null,
+    /// No text generation; embeddings or a reranker's labels; the window the budget sized (openEngine).
+    offer: serve.Offer,
 };
 
 pub const Refused = error{Refused};
@@ -224,11 +218,13 @@ fn openEngine(a: Allocator, gpa: Allocator, io: std.Io, dev: hip.Device, dir: []
         .deinit = release,
         .streams = 1,
         .startup = own.startup,
-        .text_generation = false,
-        .embeddings = if (head == .rerank) null else .{ .dimensions = @intCast(spec.hidden), .vocabulary = @intCast(spec.vocab), .max_batch_tokens = @intCast(window), .max_inputs = @intCast(engine.max_inputs), .mrl_min = @min(32, @as(u32, @intCast(spec.hidden))) },
-        .window = window,
-        .decides = head == .rerank,
-        .rerank = labels,
+        .offer = .{
+            .text_generation = false,
+            .embeddings = if (head == .rerank) null else .{ .dimensions = @intCast(spec.hidden), .vocabulary = @intCast(spec.vocab), .max_batch_tokens = @intCast(window), .max_inputs = @intCast(engine.max_inputs), .mrl_min = @min(32, @as(u32, @intCast(spec.hidden))) },
+            .window = window,
+            .decides = head == .rerank,
+            .rerank = labels,
+        },
     };
 }
 
diff --git a/zig/src/families/qwen3_5/backend/hip/engine/native.zig b/zig/src/families/qwen3_5/backend/hip/engine/native.zig
index 35adfbc734..2b7f1b266b 100644
--- a/zig/src/families/qwen3_5/backend/hip/engine/native.zig
+++ b/zig/src/families/qwen3_5/backend/hip/engine/native.zig
@@ -32,7 +32,8 @@ pub fn Native(comptime type_name: []const u8) type {
 }
 
 /// What the serve flags ask of the engine, as the host read them.
-pub const Options = @import("hip_serve").Options;
+const serve = @import("hip_serve");
+pub const Options = serve.Options;
 
 /// What the native server drives: the lane backend, the facts its round loop reads, and how to free it.
 pub const Loaded = struct {
@@ -44,8 +45,8 @@ pub const Loaded = struct {
     /// The streams admitted at once, and the server's startup line (the memory plan).
     streams: usize,
     startup: []const u8,
-    /// The readouts the engine answers (null: none; it generates text instead).
-    readouts: ?lanes.readout.Info = null,
+    /// The readouts the engine answers, if any (then it generates no text); the host's defaults otherwise.
+    offer: serve.Offer = .{},
 };
 
 /// A flag or a resource the engine cannot take: `problem` says why.
@@ -280,7 +281,7 @@ fn openEngine(a: std.mem.Allocator, gpa: std.mem.Allocator, io: std.Io, dev: hip
         .deinit = release,
         .streams = p.streams,
         .startup = p.startup,
-        .readouts = own.backend.readouts(),
+        .offer = .{ .readouts = own.backend.readouts() },
     };
 }
 
diff --git a/zig/src/hip_serve/offer.zig b/zig/src/hip_serve/offer.zig
index f0ac40d5c6..06c9f54fe6 100644
--- a/zig/src/hip_serve/offer.zig
+++ b/zig/src/hip_serve/offer.zig
@@ -1,8 +1,21 @@
-//! What a HIP family offers the host past its lanes: images, its prompt cache and its host settings, all defaulted.
+//! What a HIP family offers the host past its lanes: its routes, images, prompt cache and host settings, all defaulted.
 const api = @import("engine_api");
+const lanes = @import("lanes");
 
 /// A family's `Loaded.offer`; a family without one serves text with the host's defaults.
 pub const Offer = struct {
+    /// False: no chat or completions; the server refuses the generation routes.
+    text_generation: bool = true,
+    /// The embedding route's shape; null: no embeddings.
+    embeddings: ?api.EmbeddingInfo = null,
+    /// A reranker: the engine scores labels (its backend's score hook).
+    decides: bool = false,
+    /// The reranker's label tokens.
+    rerank: ?api.Rerank = null,
+    /// The readouts the engine answers; set: it answers them and generates no text (hip_families.generates).
+    readouts: ?lanes.readout.Info = null,
+    /// The window the engine was sized to; null: the window asked.
+    window: ?usize = null,
     /// Built by the family; null: text only.
     vision: ?api.Vision = null,
     /// The family's own store; null whenever its cache is off (no budget or no slots).
@@ -22,6 +35,7 @@ pub const Offer = struct {
 test "an empty offer is text only with the host's defaults" {
     const std = @import("std");
     const o: Offer = .{};
+    try std.testing.expect(o.text_generation and o.embeddings == null and !o.decides and o.rerank == null and o.readouts == null and o.window == null);
     try std.testing.expect(o.vision == null and o.cache == null and o.cache_capacity == null and !o.warm_turns);
     try std.testing.expect(o.frame_key == null and o.thinking_budget == null and o.on_thread == null and o.policy == null);
 }
diff --git a/zig/src/native/hip.zig b/zig/src/native/hip.zig
index c850e1279b..aedd086cdf 100644
--- a/zig/src/native/hip.zig
+++ b/zig/src/native/hip.zig
@@ -288,6 +288,7 @@ fn openWith(comptime F: type, a: Allocator, gpa: Allocator, io: std.Io, o: api.O
     };
     errdefer loaded.deinit(loaded.ctx);
     const offer = fam.offerOf(loaded);
+    const text = fam.generates(offer);
     if (fam.textOnly(o, offer))
         std.debug.print("[tensorfold] --vision: the HIP engine has no image tower for {s}; serving text only, and image requests are refused\n", .{o.model_type});
     const h = try gpa.create(Host);
@@ -302,21 +303,20 @@ fn openWith(comptime F: type, a: Allocator, gpa: Allocator, io: std.Io, o: api.O
     h.clock = .{ .io = io };
     h.core = lanes.Engine.init(gpa, &h.cfg, loaded.backend, h.clock.clock());
     errdefer h.core.deinit();
-    const what = Serves.of(loaded, r.options.window);
     h.host = api.LaneHost.init(gpa, io, &h.core, .{
         .benchmark_ranks = if (r.dev.group) |g| @intCast(g.world) else 1,
         .lanes = @intCast(loaded.streams),
-        .context_window = @intCast(what.window),
+        .context_window = @intCast(offer.window orelse r.options.window),
         .prefill_step = F.prefill_step,
         .decode_share = o.decode_share,
         .policy = offer.policy orelse h.policy_line,
         .startup = loaded.startup,
-        .structures = what.text_generation and loaded.backend.vtable.first_masked != null,
-        .text_generation = what.text_generation,
-        .embeddings = what.embeddings,
-        .decides = what.decides,
-        .rerank = what.rerank,
-        .readouts = what.readouts,
+        .structures = text and loaded.backend.vtable.first_masked != null,
+        .text_generation = text,
+        .embeddings = offer.embeddings,
+        .decides = offer.decides,
+        .rerank = offer.rerank,
+        .readouts = offer.readouts,
         .penalties = loaded.backend.penalties,
         .reasoning = reasoning(F),
         .vision = offer.vision,
@@ -335,34 +335,6 @@ fn openWith(comptime F: type, a: Allocator, gpa: Allocator, io: std.Io, o: api.O
     return .{ .engine = h.host.engine(), .close = Host.close, .ctx = h };
 }
 
-/// What a family's `Loaded` says it serves; a family that names none generates text in the window asked.
-const Serves = struct {
-    text_generation: bool = true,
-    embeddings: ?api.EmbeddingInfo = null,
-    decides: bool = false,
-    rerank: ?api.Rerank = null,
-    /// A family whose `Loaded` reports readouts answers them and generates nothing.
-    readouts: ?lanes.readout.Info = null,
-    window: usize,
-
-    fn of(loaded: anytype, asked: usize) Serves {
-        const L = @TypeOf(loaded);
-        var s: Serves = .{ .window = asked };
-        if (@hasField(L, "text_generation")) s.text_generation = loaded.text_generation;
-        if (@hasField(L, "embeddings")) s.embeddings = loaded.embeddings;
-        if (@hasField(L, "decides")) s.decides = loaded.decides;
-        if (@hasField(L, "rerank")) s.rerank = loaded.rerank;
-        if (@hasField(L, "readouts")) if (loaded.readouts) |info| {
-            s.readouts = info;
-            s.text_generation = false;
-        };
-        if (@hasField(L, "window")) if (loaded.window) |w| {
-            s.window = w;
-        };
-        return s;
-    }
-};
-
 /// A rank above 0: holds its share and runs rank 0's steps until rank 0 stops; false with `problem` if it cannot start.
 pub fn follow(a: Allocator, gpa: Allocator, io: std.Io, o: api.Open, problem: *[]const u8) !bool {
     const report = @import("sys").optional_libraries.atBoundary();
@@ -400,13 +372,17 @@ test "a HIP family declares replies without reasoning with `pub const reasoning
     }));
 }
 
-test "the text families carry no offer, so the host gives them the empty one; GLM-OCR offers through its Loaded" {
-    inline for (.{ qwen35.native, qwen35.native_moe, flashnext.native, nemotron.native }) |F| {
+test "the text families carry no offer, so the host gives them the empty one; GLM-OCR and Qwen3 offer through their Loaded" {
+    inline for (.{ flashnext.native, nemotron.native }) |F| {
         const Loaded = @typeInfo(@typeInfo(@TypeOf(F.open)).@"fn".return_type.?).error_union.payload;
         try std.testing.expect(!@hasField(Loaded, "offer"));
     }
     try std.testing.expect(@FieldType(glm_ocr.native.Loaded, "offer") == serve.Offer);
     inline for (.{ "vision", "cache", "cache_capacity" }) |old| try std.testing.expect(!@hasField(glm_ocr.native.Loaded, old));
+    try std.testing.expect(@FieldType(qwen3.native.Loaded, "offer") == serve.Offer);
+    inline for (.{ "text_generation", "embeddings", "window", "decides", "rerank" }) |old| try std.testing.expect(!@hasField(qwen3.native.Loaded, old));
+    // Qwen3.5's readouts reach the Info through its offer (both of its model types share one Loaded)
+    try std.testing.expect(@FieldType(qwen35.native_engine.Loaded, "offer") == serve.Offer and !@hasField(qwen35.native_engine.Loaded, "readouts"));
 }
 
 test {
@@ -581,26 +557,15 @@ test "slide admission: HIP followers refuse before joining ranks" {
     }
 }
 
-test "a family's Loaded names what it serves; one without the fields generates text in the asked window" {
-    const Plain = struct { streams: usize };
-    const p = Serves.of(Plain{ .streams = 1 }, 4096);
-    try std.testing.expect(p.text_generation and p.embeddings == null and p.window == 4096);
-    const Embed = struct { text_generation: bool, embeddings: ?api.EmbeddingInfo, window: ?usize };
-    const e = Serves.of(Embed{ .text_generation = false, .embeddings = .{ .dimensions = 8, .vocabulary = 4, .max_batch_tokens = 2048 }, .window = 2048 }, 32768);
-    try std.testing.expect(!e.text_generation and e.embeddings.?.dimensions == 8 and e.window == 2048);
-    const Unsized = struct { window: ?usize };
-    try std.testing.expectEqual(@as(usize, 32768), Serves.of(Unsized{ .window = null }, 32768).window);
-    const Rerank = struct { text_generation: bool, decides: bool, rerank: ?api.Rerank, window: ?usize };
-    const r = Serves.of(Rerank{ .text_generation = false, .decides = true, .rerank = .{ .yes = 9693, .no = 2152 }, .window = 32768 }, 40960);
+test "a family's offer names what it serves; one without an offer generates text in the asked window" {
+    const p = fam.offerOf(.{ .streams = @as(usize, 1) });
+    try std.testing.expect(p.text_generation and p.embeddings == null and p.window == null and !p.decides and p.rerank == null);
+    const e = fam.offerOf(.{ .offer = serve.Offer{ .text_generation = false, .embeddings = .{ .dimensions = 8, .vocabulary = 4, .max_batch_tokens = 2048 }, .window = 2048 } });
+    try std.testing.expect(!e.text_generation and e.embeddings.?.dimensions == 8 and e.window == 2048 and e.rerank == null);
+    const r = fam.offerOf(.{ .offer = serve.Offer{ .text_generation = false, .decides = true, .rerank = .{ .yes = 9693, .no = 2152 }, .window = 32768 } });
     try std.testing.expect(!r.text_generation and r.decides and r.embeddings == null and r.window == 32768);
     try std.testing.expectEqual(@as(u32, 9693), r.rerank.?.yes);
-    try std.testing.expect(p.rerank == null and e.rerank == null);
-    try std.testing.expect(!p.decides and !e.decides);
-    const Readout = struct { readouts: ?lanes.readout.Info };
-    const d = Serves.of(Readout{ .readouts = .{ .max_rows = 16384, .max_questions = 64, .max_options = 1024 } }, 16384);
-    try std.testing.expect(!d.text_generation and d.embeddings == null and d.rerank == null);
+    const d = fam.offerOf(.{ .offer = serve.Offer{ .readouts = .{ .max_rows = 16384, .max_questions = 64, .max_options = 1024 } } });
+    try std.testing.expect(!fam.generates(d) and d.embeddings == null and d.rerank == null and fam.generates(p));
     try std.testing.expectEqual(@as(u32, 1024), d.readouts.?.max_options);
-    try std.testing.expect(Serves.of(Readout{ .readouts = null }, 4096).text_generation and p.readouts == null);
-    // Qwen3.5's `Loaded` must carry the field, or its readouts would never reach the Info
-    try std.testing.expect(@hasField(@typeInfo(@typeInfo(@TypeOf(qwen35.native.open)).@"fn".return_type.?).error_union.payload, "readouts"));
 }
diff --git a/zig/src/native/hip_families.zig b/zig/src/native/hip_families.zig
index 5ec7c48307..39ac1f322b 100644
--- a/zig/src/native/hip_families.zig
+++ b/zig/src/native/hip_families.zig
@@ -47,6 +47,11 @@ pub fn offerOf(loaded: anytype) Offer {
     return if (comptime @hasField(@TypeOf(loaded), "offer")) loaded.offer else .{};
 }
 
+/// Whether the family generates text: its offer says so and it answers no readouts.
+pub fn generates(offer: Offer) bool {
+    return offer.text_generation and offer.readouts == null;
+}
+
 /// --vision asked of a model this family serves without images: it starts text-only and logs so.
 pub fn textOnly(o: api.Open, offer: Offer) bool {
     return o.vision != null and offer.vision == null;
@@ -86,6 +91,11 @@ test "a family's own refusal comes first; --vision alone refuses nothing" {
     try std.testing.expect(refused(Plain, seen) == null and refused(Aliased, seen) == null);
 }
 
+test "a family that answers readouts generates no text" {
+    try std.testing.expect(generates(.{}) and !generates(.{ .text_generation = false }));
+    try std.testing.expect(!generates(.{ .readouts = .{ .max_rows = 16384, .max_questions = 64, .max_options = 1024 } }));
+}
+
 test "a Loaded without an offer gives the empty one; --vision with no offered images is text-only" {
     const Bare = struct { streams: usize = 1 };
     const Offering = struct { offer: Offer };

From 8c7acd84dfa043e96391d9fd61e378ad7e6b75b5 Mon Sep 17 00:00:00 2001
From: laudney <106273+laudney@users.noreply.github.com>
Date: Wed, 4 Nov 2026 15:01:02 +0000
Subject: [PATCH 09/10] refactor(dots): serve dots.mocr through the family
 offer

dots.mocr landed with two host hooks of its own, openVision and
attach, so main had two paths again. Its open now reads --vision
from the shared options and returns what attach set in Loaded.offer:
the prompt cache and its capacity, warm turns, images, policy line,
thinking budget, frame key and thread hook. The host drops both
probes, and the device owner test passes hip_serve.Options. With the
cache off, dots now hands the host no store, as GLM-OCR does.
---
 zig/build/hip.zig                         |  2 +-
 zig/src/families/dots_mocr/hip_native.zig | 44 ++++++++++-------------
 zig/src/native/hip.zig                    |  8 ++---
 zig/tests/hip/dots_owner_test.zig         |  7 ++--
 4 files changed, 28 insertions(+), 33 deletions(-)

diff --git a/zig/build/hip.zig b/zig/build/hip.zig
index 53a45b96be..228cb4c0a3 100644
--- a/zig/build/hip.zig
+++ b/zig/build/hip.zig
@@ -211,7 +211,7 @@ pub fn steps(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.buil
         .target = target,
         .optimize = .ReleaseSafe,
         .link_libc = true,
-        .imports = &.{ .{ .name = "hip", .module = gpu.hip }, .{ .name = "dots_mocr", .module = gpu.native.import_table.get("dots_mocr").? } },
+        .imports = &.{ .{ .name = "hip", .module = gpu.hip }, .{ .name = "dots_mocr", .module = gpu.native.import_table.get("dots_mocr").? }, .{ .name = "hip_serve", .module = gpu.native.import_table.get("hip_serve").? } },
     }), .filters = &.{"dots HIP device owner:"} });
     b.step("build-dots-hip-owner-test", "Build real owner close and failure cleanup checks without GPU execution").dependOn(&b.addInstallArtifact(dots_owner, .{}).step);
     const native = server(b, target, gpu, build_options);
diff --git a/zig/src/families/dots_mocr/hip_native.zig b/zig/src/families/dots_mocr/hip_native.zig
index bbd5ab8f08..bb13a6cc90 100644
--- a/zig/src/families/dots_mocr/hip_native.zig
+++ b/zig/src/families/dots_mocr/hip_native.zig
@@ -42,12 +42,13 @@ pub const Loaded = struct {
     ctx: *anyopaque,
     deinit: *const fn (*anyopaque) void,
     startup: []const u8,
+    /// Its prompt cache (none when the budget or the slots are zero), images under --vision, and its host settings.
+    offer: serve.Offer,
 };
 
 const Owner = struct {
     e: *engine.Engine,
     back: *backend_.Backend,
-    vision: ?api.Vision,
     startup: []const u8,
 
     fn close(ptr: *anyopaque) void {
@@ -60,15 +61,11 @@ const Owner = struct {
     }
 };
 
-pub fn open(a: std.mem.Allocator, gpa: std.mem.Allocator, io: std.Io, dev: hip.Device, dir: []const u8, o: anytype, problem: *[]const u8) !Loaded {
-    return openVision(a, gpa, io, dev, dir, o, null, problem);
-}
-
-pub fn openVision(a: std.mem.Allocator, gpa: std.mem.Allocator, io: std.Io, dev: hip.Device, dir: []const u8, o: anytype, vision: ?api.VisionOpen, problem: *[]const u8) !Loaded {
+pub fn open(a: std.mem.Allocator, gpa: std.mem.Allocator, io: std.Io, dev: hip.Device, dir: []const u8, o: serve.Options, problem: *[]const u8) !Loaded {
     if (policyRefusal(dev.policy)) |why| return refuse(a, problem, "{s}", .{why});
     const c = cfg.Config.read(a, io, dir) catch |err| return refuse(a, problem, "dots.mocr config refused: {s}", .{@errorName(err)});
-    const image_tokens: usize = if (vision) |v| (if (v.image_tokens_set) v.image_tokens else 14400) else 0;
-    if (vision) |v| if (image_tokens == 0 or v.max_images == 0) return refuse(a, problem, "dots.mocr image limits must be positive", .{});
+    const image_tokens: usize = if (o.vision) |v| (if (v.image_tokens_set) v.image_tokens else 14400) else 0;
+    if (o.vision) |v| if (image_tokens == 0 or v.max_images == 0) return refuse(a, problem, "dots.mocr image limits must be positive", .{});
     const rows = engine.plan.max_window;
     const drawer = serve.draw.Drawer.deviceBytes(rows);
     const e = engine.Engine.load(gpa, io, dir, dev.index, .{
@@ -87,8 +84,9 @@ pub fn openVision(a: std.mem.Allocator, gpa: std.mem.Allocator, io: std.Io, dev:
     const own = try gpa.create(Owner);
     errdefer gpa.destroy(own);
     const startup = try std.fmt.allocPrint(gpa, "dots.mocr BF16: {d} lanes, context {d}, scratch {d} bytes, prompt-cache budget {d} bytes", .{ e.p.streams, e.p.window, e.p.scratch_bytes, e.p.cache_bytes });
-    own.* = .{ .e = e, .back = back, .startup = startup, .vision = null };
-    if (vision) |v| own.vision = .{
+    own.* = .{ .e = e, .back = back, .startup = startup };
+    const cached = back.store.budget > 0 and back.store.most > 0;
+    const vision: ?api.Vision = if (o.vision) |v| .{
         .ctx = own,
         .prepare = prepare,
         .image_token = c.image_token,
@@ -96,21 +94,17 @@ pub fn openVision(a: std.mem.Allocator, gpa: std.mem.Allocator, io: std.Io, dev:
         .part_type = "image",
         .image_tokens = @intCast(image_tokens),
         .max_images = v.max_images,
-    };
-    return .{ .backend = back.backend(), .facts = .{ .exact_width = @intCast(rows), .drafts = 0, .mtp = false, .copy_drafts = true, .streams_exact = true }, .rows = @intCast(e.p.rows), .streams = e.p.streams, .ctx = own, .deinit = Owner.close, .startup = startup };
-}
-
-pub fn attach(ptr: *anyopaque, host: *api.LaneHost, _: api.Open) void {
-    const own: *Owner = @ptrCast(@alignCast(ptr));
-    host.cache = &own.back.store;
-    host.info_.prompt_cache = own.back.store.budget > 0 and own.back.store.most > 0;
-    host.info_.warm_turns = host.info_.prompt_cache;
-    host.info_.prompt_cache_capacity = .{ .budget_bytes = own.e.p.cache_bytes, .minimum_state_bytes = (text.Cache.bytes(own.e.c, @min(own.e.p.window, engine.plan.prompt_rows)) catch 0) + @as(u64, own.e.c.vocab) * 2 };
-    host.info_.vision = own.vision;
-    host.info_.policy = policy_line;
-    host.info_.thinking_budget = false;
-    host.frame_key = frameKey;
-    host.on_thread = .{ .ctx = own.e, .run = bind };
+    } else null;
+    return .{ .backend = back.backend(), .facts = .{ .exact_width = @intCast(rows), .drafts = 0, .mtp = false, .copy_drafts = true, .streams_exact = true }, .rows = @intCast(e.p.rows), .streams = e.p.streams, .ctx = own, .deinit = Owner.close, .startup = startup, .offer = .{
+        .vision = vision,
+        .cache = if (cached) &back.store else null,
+        .cache_capacity = .{ .budget_bytes = e.p.cache_bytes, .minimum_state_bytes = (text.Cache.bytes(e.c, @min(e.p.window, engine.plan.prompt_rows)) catch 0) + @as(u64, e.c.vocab) * 2 },
+        .warm_turns = cached,
+        .frame_key = frameKey,
+        .thinking_budget = false,
+        .on_thread = .{ .ctx = e, .run = bind },
+        .policy = policy_line,
+    } };
 }
 
 fn bind(ptr: *anyopaque) anyerror!void {
diff --git a/zig/src/native/hip.zig b/zig/src/native/hip.zig
index aedd086cdf..9457fde272 100644
--- a/zig/src/native/hip.zig
+++ b/zig/src/native/hip.zig
@@ -281,8 +281,7 @@ fn openWith(comptime F: type, a: Allocator, gpa: Allocator, io: std.Io, o: api.O
     };
     var served = false;
     defer if (!served) r.deinit(gpa);
-    const result = if (@hasDecl(F, "openVision")) F.openVision(a, gpa, io, r.dev, o.dir, r.options, o.vision, problem) else F.open(a, gpa, io, r.dev, o.dir, r.options, problem);
-    const loaded = result catch |err| switch (err) {
+    const loaded = F.open(a, gpa, io, r.dev, o.dir, r.options, problem) catch |err| switch (err) {
         error.Refused => return null,
         else => |x| return x,
     };
@@ -329,7 +328,6 @@ fn openWith(comptime F: type, a: Allocator, gpa: Allocator, io: std.Io, o: api.O
     h.host.frame_key = offer.frame_key;
     h.host.on_thread = offer.on_thread;
     h.host.memory = .{ .read = readMemory };
-    if (@hasDecl(F, "attach")) F.attach(h.family, &h.host, o);
     try h.host.start();
     served = true;
     return .{ .engine = h.host.engine(), .close = Host.close, .ctx = h };
@@ -372,7 +370,7 @@ test "a HIP family declares replies without reasoning with `pub const reasoning
     }));
 }
 
-test "the text families carry no offer, so the host gives them the empty one; GLM-OCR and Qwen3 offer through their Loaded" {
+test "the text families carry no offer, so the host gives them the empty one; GLM-OCR, Qwen3, Qwen3.5 and dots.mocr offer through their Loaded" {
     inline for (.{ flashnext.native, nemotron.native }) |F| {
         const Loaded = @typeInfo(@typeInfo(@TypeOf(F.open)).@"fn".return_type.?).error_union.payload;
         try std.testing.expect(!@hasField(Loaded, "offer"));
@@ -383,6 +381,8 @@ test "the text families carry no offer, so the host gives them the empty one; GL
     inline for (.{ "text_generation", "embeddings", "window", "decides", "rerank" }) |old| try std.testing.expect(!@hasField(qwen3.native.Loaded, old));
     // Qwen3.5's readouts reach the Info through its offer (both of its model types share one Loaded)
     try std.testing.expect(@FieldType(qwen35.native_engine.Loaded, "offer") == serve.Offer and !@hasField(qwen35.native_engine.Loaded, "readouts"));
+    try std.testing.expect(@FieldType(dots.native.Loaded, "offer") == serve.Offer);
+    try std.testing.expect(!@hasDecl(dots.native, "attach") and !@hasDecl(dots.native, "openVision"));
 }
 
 test {
diff --git a/zig/tests/hip/dots_owner_test.zig b/zig/tests/hip/dots_owner_test.zig
index e39747a4c7..344c232cc0 100644
--- a/zig/tests/hip/dots_owner_test.zig
+++ b/zig/tests/hip/dots_owner_test.zig
@@ -2,6 +2,7 @@
 const std = @import("std");
 const hip = @import("hip");
 const family = @import("dots_mocr");
+const serve = @import("hip_serve");
 pub const std_options: std.Options = .{ .signal_stack_size = null };
 
 test "dots HIP device owner: close and late allocation failures free every owned byte" {
@@ -16,10 +17,10 @@ test "dots HIP device owner: close and late allocation failures free every owned
     defer ctx.deinit();
     if ((try ctx.memInfo()).free < 10 << 30) return error.DotsMemoryBudget;
     const before = hip.usage(false);
-    const options = .{ .window = @as(usize, 512), .streams = @as(usize, 2), .fixed = true, .cache_gib = @as(?f64, 0) };
+    const options: serve.Options = .{ .window = 512, .streams = 2, .fixed = true, .cache_gib = 0, .vision = .{ .image_tokens = 64, .image_tokens_set = true } };
     var problem: []const u8 = "";
     var counted = std.testing.FailingAllocator.init(a, .{});
-    const loaded = try family.native.openVision(arena.allocator(), counted.allocator(), std.testing.io, dev, std.mem.span(path), options, .{ .image_tokens = 64, .image_tokens_set = true }, &problem);
+    const loaded = try family.native.open(arena.allocator(), counted.allocator(), std.testing.io, dev, std.mem.span(path), options, &problem);
     const allocations = counted.alloc_index;
     try std.testing.expect(hip.usage(false).device > before.device + 6078358528);
     loaded.deinit(loaded.ctx);
@@ -29,7 +30,7 @@ test "dots HIP device owner: close and late allocation failures free every owned
     // The final two host allocations own the adapter and its startup text. Both follow the device bank and drawer.
     for ([_]usize{ allocations - 2, allocations - 1 }) |fail_at| {
         var failing = std.testing.FailingAllocator.init(a, .{ .fail_index = fail_at });
-        try std.testing.expectError(error.OutOfMemory, family.native.openVision(arena.allocator(), failing.allocator(), std.testing.io, dev, std.mem.span(path), options, .{ .image_tokens = 64, .image_tokens_set = true }, &problem));
+        try std.testing.expectError(error.OutOfMemory, family.native.open(arena.allocator(), failing.allocator(), std.testing.io, dev, std.mem.span(path), options, &problem));
         try std.testing.expect(failing.has_induced_failure);
         try std.testing.expectEqual(failing.allocated_bytes, failing.freed_bytes);
         try std.testing.expectEqual(before.device, hip.usage(false).device);

From f378511c9ff1dff60ffd378221b565ccc121fd47 Mon Sep 17 00:00:00 2001
From: ashhart <71890694+ashhart@users.noreply.github.com>
Date: Wed, 4 Nov 2026 22:21:15 +0000
Subject: [PATCH 10/10] changelog: #2040 under Unreleased

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

diff --git a/CHANGELOG.md b/CHANGELOG.md
index d1f6488e5a..b4f55e4ebe 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
 
+- HIP: the OCR families share one host protocol, with the same replies (#2040). Thanks to @laudney.
+
 - Lane host: the paged KV pool's promises now count in the host's one ledger, with no change in behaviour (#2127). Thanks to @jschmied.
 
 - Qwen3.5 on HIP: the prompt conv writes q, k and v to their own rows, which drops 144 of the 27B's 160 column copies, with the same results (#2117). Thanks to @AbrahamCete.