qwen27: canonical prompt chains follow the admitted frame on pre-M5 chips, past 128 rows

build.zig

@@ -174,8 +174,10 @@ pub fn build(b: *std.Build) void {
         gemma4_host_mods.sources.addImport("nemotron_packed_metallib", embedded(b, b.addWriteFiles().add("gemma4-compile-empty.metallib", ""), "nemotron_packed.metallib"));
         @import("zig/build/gemma4_host.zig").compileHost(b, gemma4_host_mods.metal, gemma4_host_mods.engine, gemma4_host_mods.lanes, gemma4_host_mods.sources);
         @import("zig/build/qwen27_host.zig").compileHost(b, gemma4_host_mods.metal, gemma4_host_mods.engine, gemma4_host_mods.lanes);
+        @import("zig/build/prefill_widths_metal.zig").compileHost(b, gemma4_host_mods.engine, gemma4_host_mods.lanes);
     }
     @import("zig/build/qwen27_host.zig").tests(b, test_step);
+    @import("zig/build/prefill_widths_metal.zig").tests(b, test_step);
     const xgrammar = grammar_build.option(b);
     dist_build.targets(b, draft_ids, build_options, release_version, xgrammar);
     @import("zig/build/release_paths_step.zig").tests(b, test_step);

docs/recipes/qwen3.8-27b.md

@@ -64,6 +64,12 @@ Larger or unbounded cuts keep the general draw. Sampling time has its own pass-p
 On Metal, M5 prompt passes request 2,048 rows when no other request is decoding; earlier chips keep the 512-row
 default. The larger frame must fit the configured context and all streams beside loaded weights and the 8 GiB
 margin, otherwise M5 keeps 512 rows. Startup reserves the actual prompt scratch sizes as well as stream caches.
+Qwen3.8-27B's pre-M5 GQA6 prompt chains follow the admitted frame, up to 2,048 rows, using the
+one-row tail's key chunks, half probabilities and merge. Branching and mixed-stream rounds retain 128 rows.
+The older simdgroup prompt kernel uses different probability and tile arithmetic, so this path uses the tail at
+short widths too. The M5 path and other attention geometries keep their existing selection.
+Host planning and address checks pass; Metal byte replacement, checkpoint equality and timing remain to qualify.
+No speedup over an earlier release is claimed. A larger frame must still fit the existing admission budget.
 Wide prompt frames keep separate 512-row activation scratch for small rounds, so decoding binds the small buffers.
 Admission counts both banks.
 `TF_QWEN27_WORKING_SET_BYTES` can lower the device working-set budget for admission checks. `TF_PROMPT_ROWS` selects another multiple of 16 from 16 to 2,048. Each chunk takes at most 512 rows beside

zig/build/gemma4_host.zig

@@ -40,6 +40,11 @@ pub fn compileHost(b: *std.Build, metal: *std.Build.Module, engine: *std.Build.M
     runner.addImport("h3_kernels", b.createModule(.{ .root_source_file = b.path("zig/kernels/metal/h3/kernels.zig") }));
     runner.addAnonymousImport("ling_metal", .{ .root_source_file = b.path("zig/kernels/metal/ling/sources.zig") });
     runner.addAnonymousImport("ling_route", .{ .root_source_file = b.path("zig/src/families/ling/route.zig") });
+    const chain_root = b.createModule(.{ .root_source_file = b.path("zig/q27_contract_test.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{ .{ .name = "metal", .module = metal }, .{ .name = "tensorfold", .module = engine }, .{ .name = "contract", .module = @import("kernel_contract.zig").module(b, b.graph.host, .debug) } } });
+    const chain_filter = b.option([]const u8, "q27-contract-host-filter", "One chain contract host regression");
+    const chain_run = b.addRunArtifact(b.addTest(.{ .root_module = chain_root, .filters = if (chain_filter) |f| &.{f} else &.{"q27 contract:"} }));
+    chain_run.has_side_effects = true;
+    b.step("test-q27-contract-host", "Actual chain source, stable hooks and named pipeline refusals without GPU calls").dependOn(&chain_run.step);
     const contract_root = b.createModule(.{ .root_source_file = b.path("zig/gemma4_contract_compile.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{.{ .name = "runner", .module = runner }} });
     const contract_object = b.addObject(.{ .name = "gemma4-contract-compile", .root_module = contract_root });
     _ = contract_object.getEmittedBin();

zig/build/glm_expiry.zig

@@ -52,7 +52,9 @@ pub fn steps(b: *std.Build, test_step: *std.Build.Step, lanes: *std.Build.Module
         .imports = &.{ .{ .name = "metal", .module = metal }, .{ .name = "lanes", .module = lanes }, .{ .name = "fabric", .module = fabric }, .{ .name = "kernel_sources", .module = sources } },
     }) });
     if (b.graph.host.result.os.tag == .macos) schedule.root_module.linkSystemLibrary("objc", .{});
-    b.step("test-glm-scheduling-host", "Actual GLM prefill chunks and lane-host decode fairness using a fake engine").dependOn(&b.addRunArtifact(schedule).step);
+    const schedule_run = b.addRunArtifact(schedule);
+    b.step("test-glm-scheduling-host", "Actual GLM prefill chunks and lane-host decode fairness using a fake engine").dependOn(&schedule_run.step);
+    b.step("test-glm-prefill-widths-host", "Actual GLM pass widths, restore cuts and shared lane stats without devices").dependOn(&schedule_run.step);
     if (b.graph.host.result.os.tag == .macos) {
         const compile = b.addObject(.{ .name = "glm-restore-host", .root_module = b.createModule(.{
             .root_source_file = b.path("zig/src/glm_restore_compile.zig"),

zig/build/prefill_widths_metal.zig

@@ -0,0 +1,18 @@
+//! Metal prompt observations tested at their real completion boundaries; no shaders or model execution.
+const std = @import("std");
+
+pub fn tests(b: *std.Build, parent: *std.Build.Step) void {
+    const lanes = b.createModule(.{ .root_source_file = b.path("zig/src/core/lanes/lanes.zig"), .target = b.graph.host, .link_libc = true });
+    const root = b.createModule(.{ .root_source_file = b.path("zig/src/families/nemotron/prompt_flight_test.zig"), .target = b.graph.host, .optimize = .debug, .link_libc = true, .imports = &.{.{ .name = "lanes", .module = lanes }} });
+    const run = b.addRunArtifact(b.addTest(.{ .root_module = root, .filters = &.{"Nemotron prefill widths:"} }));
+    run.has_side_effects = true;
+    b.step("test-nemotron-prefill-widths-host", "Actual Nemotron cut and flight completion with recording buffers").dependOn(&run.step);
+    parent.dependOn(&run.step);
+}
+
+pub fn compileHost(b: *std.Build, engine: *std.Build.Module, lanes: *std.Build.Module) void {
+    const root = b.createModule(.{ .root_source_file = b.path("zig/src/metal_prefill_widths_compile.zig"), .target = b.graph.host, .optimize = .debug, .link_libc = true, .imports = &.{ .{ .name = "tensorfold", .module = engine }, .{ .name = "lanes", .module = lanes } } });
+    const object = b.addObject(.{ .name = "metal-prefill-widths-compile", .root_module = root });
+    _ = object.getEmittedBin();
+    b.step("check-nemotron-prefill-widths-host", "Compile actual Metal prefill, read, land and release without shaders").dependOn(&object.step);
+}

zig/build/q27_profile_host.zig

@@ -11,4 +11,5 @@ pub fn steps(b: *std.Build, server_tests: *std.Build.Module) void {
     step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = profile })).step);
     step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = api, .filters = &.{"gpu_seconds:"} })).step);
     step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = server_tests, .filters = &.{"gpu_seconds:"} })).step);
+    b.step("test-qwen27-prefill-widths-http-host", "Completed prompt widths on raw and chat HTTP routes without devices").dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = server_tests, .filters = &.{"prefill widths:"} })).step);
 }

zig/build/qwen27_host.zig

@@ -19,6 +19,10 @@ pub fn tests(b: *std.Build, test_step: *std.Build.Step) void {
     const allocated = b.createModule(.{ .root_source_file = b.path("zig/qwen27_allocation_host.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{.{ .name = "lanes", .module = pricing_lanes }} });
     const window_step = b.step("test-qwen27-window-host", "Device width, copied trees and counted zero-grant drafting without shaders");
     for ([_]*std.Build.Module{ window, taps, allocated }) |m| window_step.dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = m })).step);
+    const chain_ref = b.createModule(.{ .root_source_file = b.path("zig/tests/contract/q27_chain_ref.zig"), .target = b.graph.host, .imports = &.{.{ .name = "contract", .module = @import("kernel_contract.zig").module(b, b.graph.host, .debug) }} });
+    const chain_run = b.addRunArtifact(b.addTest(.{ .root_module = chain_ref }));
+    test_step.dependOn(&chain_run.step);
+    b.step("test-qwen27-chain-reference-host", "FP64 chain attention and reduced-sum witness without a GPU").dependOn(&chain_run.step);
     const rows = b.createModule(.{ .root_source_file = b.path("zig/src/qwen27_rows_test.zig"), .target = b.graph.host });
     b.step("test-qwen27-prompt-rows-host", "Device-aware prompt rows beside owned decoding streams without a GPU").dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = rows, .filters = &.{"qwen27 rows:"} })).step);
     const ane_policy = b.createModule(.{ .root_source_file = b.path("zig/src/families/qwen27/ane_policy.zig"), .target = b.graph.host });
@@ -51,6 +55,7 @@ pub fn tests(b: *std.Build, test_step: *std.Build.Step) void {
     var names: std.ArrayList([]const u8) = .empty;
     tryNames(b, &names, @embedFile("../src/families/qwen27/gdn_test.zig"));
     tryNames(b, &names, @embedFile("../src/families/qwen27/attention_plan.zig"));
+    tryNames(b, &names, @embedFile("../src/families/qwen27/attention_chain_test.zig"));
     const plans = b.addRunArtifact(b.addTest(.{ .root_module = keep_module, .filters = names.items }));
     plans.has_side_effects = true;
     b.step("test-qwen27-plans-host", "Shared GDN and attention plans with no shaders or GPU").dependOn(&plans.step);
@@ -66,6 +71,11 @@ pub fn tests(b: *std.Build, test_step: *std.Build.Step) void {
     b.step("test-qwen27-equality-host", "Shared lanes drafted, resumed and concurrent equality on the fake target").dependOn(&equal.step);
     const resume_lanes = b.createModule(.{ .root_source_file = b.path("zig/src/core/lanes/lanes.zig"), .target = b.graph.host, .link_libc = true });
     const resume_api = b.createModule(.{ .root_source_file = b.path("zig/src/core/engine_api.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{.{ .name = "lanes", .module = resume_lanes }} });
+    const prompt_stats = b.createModule(.{ .root_source_file = b.path("zig/src/qwen27_prefill_stats_test.zig"), .target = b.graph.host, .optimize = .debug, .link_libc = true, .imports = &.{ .{ .name = "lanes", .module = resume_lanes }, .{ .name = "engine_api", .module = resume_api } } });
+    const prompt_stats_run = b.addRunArtifact(b.addTest(.{ .root_module = prompt_stats, .filters = &.{"qwen27 prefill stats:"} }));
+    prompt_stats_run.has_side_effects = true;
+    test_step.dependOn(&prompt_stats_run.step);
+    b.step("test-qwen27-prefill-stats-host", "Completed prompt passes through the lane host and server mailbox without a GPU").dependOn(&prompt_stats_run.step);
     const resume_rules = b.createModule(.{ .root_source_file = b.path("zig/src/native/qwen27_resume.zig"), .target = b.graph.host, .optimize = .debug, .imports = &.{.{ .name = "engine_api", .module = resume_api }} });
     b.step("test-qwen27-resume-host", "Slot 0's in-place resume and rewind point: the cache's rule, with no GPU").dependOn(&b.addRunArtifact(b.addTest(.{ .root_module = resume_rules, .filters = &.{"qwen27 resume:"} })).step);
 }

zig/kernels/metal/qwen27/attention_mpp.metal

@@ -239,7 +239,7 @@ kernel void q27_attn_tail(device const bfloat* QB [[buffer(0)]], device const Q2
         const int q = kt + h * 32 + row;
         int phys = -1;
         if (q < P) phys = q;
-        else if (q < nmax) phys = P + paths[node * MAXD + (q - P)];
+        else if (q < nmax) phys = meta[MG + st * MS + 10] ? q : P + paths[node * MAXD + (q - P)];
         ((threadgroup vec<bfloat, 8>*)KV)[e] = phys >= 0 ? *(const device vec<bfloat, 8>*)(kbase + phys * kstep + col) : vec<bfloat, 8>(0);
       }
       threadgroup_barrier(mem_flags::mem_threadgroup);
@@ -287,7 +287,7 @@ kernel void q27_attn_tail(device const bfloat* QB [[buffer(0)]], device const Q2
         const int q = kt + row;
         int phys = -1;
         if (q < P) phys = q;
-        else if (q < nmax) phys = P + paths[node * MAXD + (q - P)];
+        else if (q < nmax) phys = meta[MG + st * MS + 10] ? q : P + paths[node * MAXD + (q - P)];
         ((threadgroup vec<bfloat, 8>*)KV)[e] = phys >= 0 ? *(const device vec<bfloat, 8>*)(vbase + phys * vstep + col) : vec<bfloat, 8>(0);
       }
       threadgroup_barrier(mem_flags::mem_threadgroup);

zig/q27_contract_test.zig

@@ -0,0 +1,49 @@
+//! Host-only source and lifecycle checks for the real chain contract; no Metal operation is called.
+const std = @import("std");
+const tf = @import("tensorfold");
+const g = @import("contract").gates;
+const chain = @import("tests/contract/metal_q27_chain.zig");
+const t = std.testing;
+
+test "q27 contract: both chain libraries contain the served pre-M5 prompt shader" {
+    for ([_]bool{ false, true }) |baseline| for ([_]bool{ false, true }) |units| {
+        const text = try chain.source(t.allocator, units, baseline, false, .q27b);
+        defer t.allocator.free(text);
+        const symbol = if (units) "Q27_PROMPT(8, false," else "Q27_PROMPT_SG(8, false,";
+        try t.expect(std.mem.indexOf(u8, text, symbol) != null);
+        try t.expect(std.mem.indexOf(u8, text, "q27_attn_prompt_merge_single") != null);
+    };
+    const text = try tf.qwen27.simd_attention.library(t.allocator, false, true);
+    defer t.allocator.free(text);
+    try t.expect(std.mem.indexOf(u8, text, "Q27_PROMPT_SG(8, false,") != null);
+}
+
+noinline fn fromAnotherFrame(x: *chain.Case) g.Kernel {
+    return x.kernel();
+}
+
+test "q27 contract: returned hooks survive the registering stack frame" {
+    var x: chain.Case = undefined;
+    x.reg = .{ .heads = .q27b };
+    x.shape = .{ .rows = 3, .n = 24, .k = 0 };
+    const k = fromAnotherFrame(&x);
+    const expected = comptime g.hooks(chain.Case);
+    try t.expectEqualDeep(expected, k.vt.*);
+    try t.expectEqual(@as(usize, 3 * 24 * 256), k.vt.count(k.ctx));
+}
+
+test "q27 contract: missing prompt pipelines become named gate failures without Metal calls" {
+    var x: chain.Case = undefined;
+    x.rig = .{ .a = t.allocator, .device = undefined, .queue = undefined };
+    defer x.rig.deinit();
+    x.reg = .{ .heads = .q27b };
+    x.op.prompt = null;
+    x.old.prompt = null;
+    const shape = @import("contract").spec.Shape{ .rows = 16, .n = 24, .k = 0 };
+    try t.expectError(error.Qwen27ChainPromptUnavailable, x.prepare(shape, .random, 1));
+    var tally: g.Tally = .{};
+    try g.runAll(t.allocator, &.{x.kernel()}, 1, &tally);
+    try t.expectEqual(@as(usize, 1), tally.kernels);
+    try t.expectEqual(@as(usize, 1), tally.failed);
+    try t.expectEqual(@as(usize, 0), tally.passed);
+}

zig/src/core/contract/q27_chain_records.zig

@@ -0,0 +1,10 @@
+//! Canonical chain attention keeps the tail's 128 threads and the ordered merge's 32 threads.
+const limits = @import("limits.zig");
+
+// Metal 4 tensor descriptors compile at runtime; M1/M2 ahead-of-time limits are not claimed.
+pub const records = [_]limits.MetalLaunch{
+    .{ .name = "qwen27-chain one q27_attn_tail", .threads = 128, .tg_memory = 18432, .offline = false },
+    .{ .name = "qwen27-chain one q27_attn_merge_single", .threads = 32, .tg_memory = 0, .offline = false },
+    .{ .name = "qwen27-chain rank q27_attn_tail", .threads = 128, .tg_memory = 18432, .offline = false },
+    .{ .name = "qwen27-chain rank q27_attn_merge_single", .threads = 32, .tg_memory = 0, .offline = false },
+};

zig/src/core/contract/registry.zig

@@ -568,7 +568,7 @@ test {
 }
 const paging = @import("paging_pending.zig");
 /// The Metal collector keeps measured slots and the restored vision declarations.
-pub const all_metal = @import("attention_pv_records.zig").records ++ metal ++ @import("delta_slots_records.zig").records ++ @import("flashnext_pending.zig").records ++ @import("simd_pairs_records.zig").records ++ @import("metal_pair_records.zig").records ++ @import("glm_rows_records.zig").records ++ @import("weight_only_records.zig").metal ++ @import("hc_activation_records.zig").metal ++ @import("qwen4_records.zig").metal ++ @import("lfm2_records.zig").records ++ @import("nemotron_experts_records.zig").records ++ @import("nemotron_pages_records.zig").records ++ @import("nemotron_attention_records.zig").records ++
+pub const all_metal = @import("q27_chain_records.zig").records ++ @import("attention_pv_records.zig").records ++ metal ++ @import("delta_slots_records.zig").records ++ @import("flashnext_pending.zig").records ++ @import("simd_pairs_records.zig").records ++ @import("metal_pair_records.zig").records ++ @import("glm_rows_records.zig").records ++ @import("weight_only_records.zig").metal ++ @import("hc_activation_records.zig").metal ++ @import("qwen4_records.zig").metal ++ @import("lfm2_records.zig").records ++ @import("nemotron_experts_records.zig").records ++ @import("nemotron_pages_records.zig").records ++ @import("nemotron_attention_records.zig").records ++
     @import("nemotron_ssd_records.zig").records ++ @import("flashnext_attention_records.zig").records ++ @import("spark25_records.zig").records ++ @import("gemma_attention_records.zig").records ++ @import("moss_image_records.zig").records ++ @import("dsv41_metal_records.zig").records ++ @import("minimax_records.zig").records ++ @import("ling_route_records.zig").records ++ @import("qwen27_moe_records.zig").records ++ @import("affine_gather_records.zig").records;
 const learning = @import("registry_learning.zig");
 pub const gradient_pending = learning.gradient_pending;

zig/src/core/lane_host.zig

@@ -300,7 +300,7 @@ pub const LaneHost = struct {
         h.gpa.free(job.marks);
         h.gpa.free(job.key);
         const s = &job.stream;
-        const stats: Stats = if (job.started) .{ .rounds = s.rounds, .grammar = if (s.grammar_time.windows > 0) s.grammar_time else null, .drafted = s.drafted, .accepted = s.accepted, .min_rows = s.min_rows, .loop_period = s.loop_period, .gpu_seconds = s.gpu_seconds, .prefill_seconds = if (job.prefilled) |done| @as(f64, @floatFromInt(@as(i64, @intCast(@max(0, done - job.began))))) / 1e9 else null } else .{};
+        const stats: Stats = if (job.started) .{ .rounds = s.rounds, .grammar = if (s.grammar_time.windows > 0) s.grammar_time else null, .drafted = s.drafted, .accepted = s.accepted, .min_rows = s.min_rows, .prefill_widths = s.prefill_widths.items, .loop_period = s.loop_period, .gpu_seconds = s.gpu_seconds, .prefill_seconds = if (job.prefilled) |done| @as(f64, @floatFromInt(@as(i64, @intCast(@max(0, done - job.began))))) / 1e9 else null } else .{};
         var told = if (job.started) recalls.finished(h, job, reason, stats) else stats;
         told.queued_seconds = if (job.admitted_ns) |at| @as(f64, @floatFromInt(@max(0, at - job.submitted))) / 1e9 else 0;
         emit(job, .{ .finished = .{ .reason = reason, .stats = told, .message = message } });

zig/src/core/lanes/lanes.zig

@@ -11,6 +11,7 @@ pub const recall = @import("recall.zig");
 pub const recall_store = @import("recall_store.zig");
 pub const sampling = @import("sampling.zig");
 pub const fill = @import("fill.zig");
+pub const prefill_widths = @import("prefill_widths.zig");
 pub const shape = @import("shape.zig");
 pub const plan_lanes = @import("plan_lanes.zig");
 pub const cost_cache = @import("cost_cache.zig");

zig/src/core/lanes/prefill_widths.zig

@@ -0,0 +1,11 @@
+//! Optional request-owned observations of completed target prompt passes, independent of backend arithmetic.
+const std = @import("std");
+
+/// Observation failure never fails inference or reports a partial list.
+pub fn completed(gpa: std.mem.Allocator, stream: anytype, rows: usize) void {
+    if (stream.prefill_widths_failed) return;
+    stream.prefill_widths.append(gpa, @intCast(rows)) catch {
+        stream.prefill_widths.clearRetainingCapacity();
+        stream.prefill_widths_failed = true;
+    };
+}