feat(hip): one host protocol for the HIP OCR families

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);
@@ -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);
@@ -380,11 +381,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 +417,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 +437,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)}));

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));

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");
 }
 

zig/src/core/contract/hip_records.zig

@@ -497,8 +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_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{

zig/src/core/contract/hip_records_families.zig

@@ -0,0 +1,20 @@
+//! 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 = .{
+    @import("hip_records_glm_ocr.zig").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 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, 182), every.len); // main's records, GLM-OCR's 12 among them
+    try std.testing.expectEqual(@as(usize, 170), every.len - all.len);
+}

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,30 +84,27 @@ 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,
         .markers = &.{ 151665, 151666, 151667 },
         .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 {

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", .{});

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");
@@ -23,12 +24,15 @@ 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;
 
-/// 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,
@@ -37,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};
@@ -221,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,
+        },
     };
 }