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. diff --git a/zig/build/hip.zig b/zig/build/hip.zig index f1fa189e86..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); @@ -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)})); 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 ` 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..4dd38d02d4 100644 --- a/zig/src/core/contract/hip_records.zig +++ b/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{ 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..43fffd7901 --- /dev/null +++ b/zig/src/core/contract/hip_records_families.zig @@ -0,0 +1,20 @@ +//! The HIP families' gate-5 records, one array a family (`hip_records_.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); +} 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/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..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"); @@ -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, + }, }; } 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/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..06c9f54fe6 --- /dev/null +++ b/zig/src/hip_serve/offer.zig @@ -0,0 +1,41 @@ +//! 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). + 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.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/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); +} diff --git a/zig/src/native/hip.zig b/zig/src/native/hip.zig index 472c36a42b..9457fde272 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; @@ -297,12 +281,15 @@ 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, }; 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); errdefer gpa.destroy(h); h.gpa = gpa; @@ -315,48 +302,37 @@ 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), .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(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 = 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, + .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; 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(); @@ -366,7 +342,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 +370,25 @@ 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, 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")); + } + 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")); + 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 { + _ = 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 +420,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); @@ -560,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 new file mode 100644 index 0000000000..39ac1f322b --- /dev/null +++ b/zig/src/native/hip_families.zig @@ -0,0 +1,110 @@ +//! 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 .{}; +} + +/// 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; +} + +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 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 }; + 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)); +} 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..b881aa950a --- /dev/null +++ b/zig/tests/hip/contract_families.zig @@ -0,0 +1,13 @@ +//! The HIP families' contract entries, one line a family: `tf-hip-kernel-contract image ` runs one. +pub const Profile = @import("contract_entry.zig").Profile; +pub const Entry = @import("contract_entry.zig").Entry; + +pub const entries = [_]Entry{ + .{ .group = .glm_ocr, .run = @import("contract_glm_ocr.zig").run }, +}; + +comptime { // one entry a group, so `image ` 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_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_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..a2252abf96 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"); @@ -24,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"); @@ -39,9 +38,19 @@ 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")) { - 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 +63,57 @@ 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| { + 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}); + 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 @@ -221,23 +276,16 @@ 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(); + 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 new file mode 100644 index 0000000000..9076511d37 --- /dev/null +++ b/zig/tests/hip/contract_select.zig @@ -0,0 +1,211 @@ +//! `tf-hip-kernel-contract image `: 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'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); +} + +/// 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 { + 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 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 .{}; +} + +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 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); + 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.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()); +} + +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); +} 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);