qwen27: pin the pair against the single Mac in the test suite

build.zig

@@ -160,7 +160,7 @@ pub fn build(b: *std.Build) void {
         const gemma4_host_mods = metalEngineModules(b, b.graph.host, .Debug, draft_ids);
         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);
+        test_step.dependOn(@import("zig/build/qwen27_host.zig").compileHost(b, gemma4_host_mods.metal, gemma4_host_mods.engine, gemma4_host_mods.lanes));
     }
     @import("zig/build/qwen27_host.zig").tests(b);
     const xgrammar = grammar_build.option(b);

zig/build/qwen27.zig

@@ -107,6 +107,9 @@ pub fn targets(b: *std.Build, target: std.Build.ResolvedTarget, optimize: std.bu
 
     const family: []const std.Build.Module.Import = &.{ .{ .name = "qwen27", .module = qwen27 }, .{ .name = "core", .module = qwen_core }, .{ .name = "metal", .module = metal } };
     _ = program(b, "tf-qwen27-run", "Build the native greedy, chunk and prompt-speed check on real weights", "zig/tests/qwen27/run.zig", target, true, &.{ .{ .name = "qwen27", .module = qwen27 }, .{ .name = "core", .module = qwen_core }, .{ .name = "metal", .module = metal }, .{ .name = "core_lane", .module = core_lane }, .{ .name = "qwen_runtime_sources", .module = s.runtime_sources }, .{ .name = "qwen_gdn_source", .module = s.gdn_source } });
+    // the TP gate before any wire: ranks = 1 parity, then two rank-local shards through the loopback link
+    const tp_check = program(b, "tf-qwen27-tp-check", "The 27B TP gate: ranks = 1 parity, then the two-rank loopback stream (SKIP without a checkpoint)", "zig/tests/qwen27/tp_check.zig", target, true, family);
+    test_step.dependOn(&b.addRunArtifact(tp_check).step);
     _ = program(b, "tf-qwen27-attention-fused", "Build the fused-chain versus tail-kernel attention byte check on this Mac's tensor units, no models", "zig/tests/qwen27/attention_fused.zig", target, false, &.{ .{ .name = "qwen27", .module = qwen27 }, .{ .name = "metal", .module = metal }, .{ .name = "qwen_runtime_sources", .module = s.runtime_sources } });
     _ = program(b, "tf-qwen27-attention-precision", "Build decode attention's error against a float64 reference on this Mac, no models", "zig/tests/qwen27/attention_precision.zig", target, false, &.{ .{ .name = "qwen27", .module = qwen27 }, .{ .name = "metal", .module = metal }, .{ .name = "qwen_runtime_sources", .module = s.runtime_sources } });
     _ = program(b, "tf-qwen27-projection-cases", "Build real-input production projection precision replay", "zig/tests/qwen27/projection_cases.zig", target, true, &.{ .{ .name = "qwen27", .module = qwen27 }, .{ .name = "metal", .module = metal } });

zig/build/qwen27_host.zig

@@ -61,9 +61,14 @@ pub fn tests(b: *std.Build) void {
     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);
 }
 
-pub fn compileHost(b: *std.Build, metal: *std.Build.Module, engine: *std.Build.Module, lanes: *std.Build.Module) void {
+pub fn compileHost(b: *std.Build, metal: *std.Build.Module, engine: *std.Build.Module, lanes: *std.Build.Module) *std.Build.Step {
     const 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 = lanes }} });
     const host = b.createModule(.{ .root_source_file = b.path("zig/src/native/qwen27_served.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{ .{ .name = "metal", .module = metal }, .{ .name = "tensorfold", .module = engine }, .{ .name = "engine_api", .module = api } } });
+    // pair refusals are decided before a model loads: the host's own tests run on CPU alone
+    const served = b.step("test-qwen27-served-host", "Rank 1's named refusal and the pair's --vision refusal, no shaders or GPU");
+    const served_run = b.addRunArtifact(b.addTest(.{ .root_module = host }));
+    served_run.has_side_effects = true; // a rerun must re-execute the tests, not replay a cached pass
+    served.dependOn(&served_run.step);
     const boundary = b.addTest(.{ .root_module = b.createModule(.{ .root_source_file = b.path("zig/qwen27_bank_boundary_host.zig"), .target = b.graph.host, .optimize = .debug, .link_libc = true, .imports = &.{ .{ .name = "tensorfold", .module = engine }, .{ .name = "metal", .module = metal } } }), .filters = &.{"qwen27 bank boundary:"} });
     b.step("test-qwen27-bank-boundary-host", "Runner pending-keep guards with no device or shader execution").dependOn(&b.addRunArtifact(boundary).step);
     const facts = b.createModule(.{ .root_source_file = b.path("zig/src/native/draft_sources_metal_test.zig"), .target = b.graph.host, .link_libc = true, .imports = &.{ .{ .name = "metal", .module = metal }, .{ .name = "tensorfold", .module = engine }, .{ .name = "engine_api", .module = api }, .{ .name = "lanes", .module = lanes } } });
@@ -81,6 +86,7 @@ pub fn compileHost(b: *std.Build, metal: *std.Build.Module, engine: *std.Build.M
     }) });
     _ = object.getEmittedBin();
     b.step("check-qwen27-speed-host", "Compile the actual Qwen27 startup and encoding without shaders").dependOn(&object.step);
+    return served;
 }
 
 fn tryNames(b: *std.Build, names: *std.ArrayList([]const u8), text: []const u8) void {

zig/kernels/metal/qwen27/attention_mpp.metal

@@ -263,8 +263,71 @@ kernel void q27_attn_tail(device const bfloat* QB [[buffer(0)]], device const Q2
   }
 }
 
-// A row's chunk partials combined in key order, a block of loads ahead; `gated` applies the query gate (q27_gate).
+// A row's chunk partials in key order, loads ahead: `gated` the query gate, `heads` the rank's (12 = one Mac's 24 / 2).
 kernel void q27_attn_merge(device const float* POA [[buffer(0)]], device const float* PMA [[buffer(1)]], device const float* PLA [[buffer(2)]],
+    device const float* POB [[buffer(3)]], device const float* PMB [[buffer(4)]], device const float* PLB [[buffer(5)]],
+    device const int* meta [[buffer(6)]], device const int* nodes [[buffer(7)]], device bfloat* OUT [[buffer(8)]],
+    device const bfloat* QG [[buffer(9)]], constant uint& gated [[buffer(10)]], constant uint& heads [[buffer(11)]],
+    uint3 threadgroup_position_in_grid [[threadgroup_position_in_grid]], uint thread_index_in_simdgroup [[thread_index_in_simdgroup]]) {
+  const uint lane = thread_index_in_simdgroup;
+  const uint hk = threadgroup_position_in_grid.x;
+  const uint r = threadgroup_position_in_grid.y;
+  const int NCB = meta[3], W = meta[4], RPA = meta[5], CA = meta[6];
+  const int st = nodes[NW * (int(r) / G) + 1];
+  const int NCBS = meta[MG + st * MS + 5];
+  const int ra = meta[MG + st * MS + 3] * 16 + (int(r) / G - meta[MG + st * MS + 6]) * G + int(r) % G;
+  const int CT = nodes[NW * (int(r) / G) + 3]; // the row's own prefix chunks
+  constexpr int DP = D / 32, MB = 8;
+  const int node = r / G, g = r % G;
+  float m = -INFINITY, l = 0.0f, o[DP];
+  for (int i = 0; i < DP; i++) o[i] = 0.0f;
+  for (int c0 = 0; c0 < CT; c0 += MB) {
+    float mcs[MB], lcs[MB];
+    float4 lo[MB], hi[MB];
+    for (int j = 0; j < MB; j++) {
+      const int64_t row = ((int64_t)hk * CA + min(c0 + j, CT - 1)) * RPA + ra;
+      mcs[j] = PMA[row];
+      lcs[j] = PLA[row];
+      lo[j] = *(device const float4*)(POA + row * D + lane * DP);
+      hi[j] = *(device const float4*)(POA + row * D + lane * DP + 4);
+    }
+    for (int j = 0; j < MB; j++) {
+      if (c0 + j >= CT) break;
+      const float mc = mcs[j];
+      if (mc == -INFINITY) continue;
+      const float lc = lcs[j];
+      const float nm = max(m, mc);
+      const float f1 = fast::exp(m - nm), f2 = fast::exp(mc - nm);
+      l = l * f1 + lc * f2;
+      for (int i = 0; i < 4; i++) o[i] = o[i] * f1 + lo[j][i] * f2;
+      for (int i = 0; i < 4; i++) o[4 + i] = o[4 + i] * f1 + hi[j][i] * f2;
+      m = nm;
+    }
+  }
+  for (int c = 0; c < NCBS; c++) {
+    const int64_t row = (((int64_t)hk * NCB + c) * W + node) * 16 + g;
+    const float mc = PMB[row];
+    if (mc == -INFINITY) continue;
+    const float lc = PLB[row];
+    const float nm = max(m, mc);
+    const float f1 = fast::exp(m - nm), f2 = fast::exp(mc - nm);
+    l = l * f1 + lc * f2;
+    for (int i = 0; i < DP; i++) o[i] = o[i] * f1 + POB[row * D + lane * DP + i] * f2;
+    m = nm;
+  }
+  const int h = hk * G + g;
+  for (int i = 0; i < DP; i++) {
+    const int64_t index = ((int64_t)node * (int64_t)heads + h) * D + lane * DP + i;
+    const bfloat v = static_cast<bfloat>(o[i] / l);
+    if (gated) {
+      const float gate = float(QG[((int64_t)node * (int64_t)heads + h) * 2 * D + D + lane * DP + i]);
+      OUT[index] = static_cast<bfloat>(float(v) * q27_round(1.0f / (1.0f + exp(-gate)), 0));
+    } else OUT[index] = v;
+  }
+}
+
+// A single Mac's 24 heads are known at compile time: main's kernel, so one Mac runs unpaired code.
+kernel void q27_attn_merge_single(device const float* POA [[buffer(0)]], device const float* PMA [[buffer(1)]], device const float* PLA [[buffer(2)]],
     device const float* POB [[buffer(3)]], device const float* PMB [[buffer(4)]], device const float* PLB [[buffer(5)]],
     device const int* meta [[buffer(6)]], device const int* nodes [[buffer(7)]], device bfloat* OUT [[buffer(8)]],
     device const bfloat* QG [[buffer(9)]], constant uint& gated [[buffer(10)]],

zig/kernels/metal/qwen27/attention_prompt.metal

@@ -4,7 +4,8 @@
 using namespace metal;
 using namespace mpp::tensor_ops;
 constant constexpr int PG = 6, PD = 256, PH = 128, PCK = 512, PTK = 32;
-struct Q27PromptArgs { int start, rows, chunks, rp, key_stride, value_stride; float scale2; };
+// kv_heads is the rank's, not the model's: a split attention writes local node rows.
+struct Q27PromptArgs { int start, rows, chunks, rp, key_stride, value_stride; float scale2; int kv_heads; };
 typedef vec<float, 8> q27f8;
 typedef vec<bfloat, 8> q27b8;
 typedef vec<half, 8> q27h8;
@@ -166,5 +167,29 @@ kernel void q27_attn_prompt_merge(device const float* PO [[buffer(0)]], device c
     m = nm;
   }
   const int node = r / PG, head = hk * PG + r % PG;
+  const int64_t at = ((int64_t)node * (PG * args.kv_heads) + head) * PD + lane * DP;
+  for (int i = 0; i < DP; i++) OUT[at + i] = static_cast<bfloat>(o[i] / l);
+}
+
+// A single Mac's 24 query heads are known at compile time: main's kernel, so one Mac runs unpaired code.
+kernel void q27_attn_prompt_merge_single(device const float* PO [[buffer(0)]], device const float* PM [[buffer(1)]], device const float* PL [[buffer(2)]],
+    constant Q27PromptArgs& args [[buffer(3)]], device bfloat* OUT [[buffer(4)]],
+    uint3 tg [[threadgroup_position_in_grid]], uint lane [[thread_index_in_simdgroup]]) {
+  const int hk = int(tg.x), r = int(tg.y);
+  constexpr int DP = PD / 32;
+  float m = -INFINITY, l = 0.0f, o[DP];
+  for (int i = 0; i < DP; i++) o[i] = 0.0f;
+  for (int c = 0; c < args.chunks; c++) {
+    const int64_t row = ((int64_t)hk * args.chunks + c) * args.rp + r;
+    const float mc = PM[row];
+    if (mc == -INFINITY) continue;
+    const float lc = PL[row];
+    const float nm = max(m, mc);
+    const float f1 = fast::exp2(m - nm), f2 = fast::exp2(mc - nm);
+    l = l * f1 + lc * f2;
+    for (int i = 0; i < DP; i++) o[i] = o[i] * f1 + PO[row * PD + lane * DP + i] * f2;
+    m = nm;
+  }
+  const int node = r / PG, head = hk * PG + r % PG;
   for (int i = 0; i < DP; i++) OUT[((int64_t)node * (PG * 4) + head) * PD + lane * DP + i] = static_cast<bfloat>(o[i] / l);
 }

zig/qwen27_bank_boundary_host.zig

@@ -38,6 +38,34 @@ test "qwen27 bank boundary: a verify tail cannot run on prompt heads or accepted
     try std.testing.expect(!called);
 }
 
+fn unusedSeamPart(_: *anyopaque) q.forward.Seam {
+    unreachable;
+}
+fn unusedSeamStage(_: *anyopaque, _: *@import("metal").ComputeEncoder, _: u32, _: q.projection.Ref, _: u32, _: u32, _: u32) void {
+    unreachable;
+}
+fn unusedSeamGather(_: *anyopaque, _: *@import("metal").ComputeEncoder, _: u32, _: q.projection.Ref, _: q.projection.Ref, _: u32, _: u32, _: u32) void {
+    unreachable;
+}
+fn unusedSeamFailed(_: *anyopaque) bool {
+    unreachable;
+}
+test "qwen27 bank boundary: a rank's verify never runs a tail" {
+    var model: q.model.Model = undefined;
+    model.frame.capacity = 16;
+    var runner: q.decode_round.Runner = undefined;
+    runner.failed = false;
+    runner.active = null;
+    runner.model = &model;
+    var called = false;
+    runner.tail = .{ .ptr = &called, .encode = unusedTail };
+    runner.reduce = .{ .ptr = undefined, .part = unusedSeamPart, .stage = unusedSeamStage, .gather = unusedSeamGather, .failed = unusedSeamFailed };
+    var round: q.round_plan.Round = undefined;
+    round.ids = @constCast(&[_]u32{0});
+    try std.testing.expectError(error.TailOnRank, runner.verifyHead(&round, .all));
+    try std.testing.expect(!called);
+}
+
 test "qwen27 bank boundary: snapshot drain preserves the picked bonus while reset forgets it" {
     var g: q.dflash.generation.Generation = undefined;
     g.keep = null;

zig/qwen27_core.zig

@@ -33,3 +33,6 @@ pub const bf16_topk_gpu = @import("bf16_topk_gpu");
 pub const tree_commit_gpu = @import("tree_commit_gpu");
 
 pub const tree_slots = @import("src/core/tree_slots.zig");
+
+/// The peer window and its seam kernel: the 27B's tensor-parallel exchange, and its one-Mac loopback.
+pub const pair = @import("src/core/pair_rows.zig");

zig/src/core/simd_qmm.zig

@@ -184,6 +184,12 @@ test "chunks follow output width only and tiles fit the reduction buffer" {
     try std.testing.expectEqual(@as(u32, 8), splits(34816));
     try std.testing.expectEqual(splits(6240), chunks(3168, 6240));
     try std.testing.expectEqual(splits(6240), chunks(6240, 0));
+    // a pair's q and q | k | v key their checkpoint width: 12,288 and 10,240 over 6,144 and 5,120 a rank's.
+    try std.testing.expectEqual(@as(u32, 8), chunks(6144, 12288));
+    try std.testing.expectEqual(@as(u32, 8), chunks(5120, 10240));
+    // the rank's window alone would take 16 chunks, the split this key keeps a pair from forking.
+    try std.testing.expectEqual(@as(u32, 16), chunks(6144, 0));
+    try std.testing.expectEqual(@as(u32, 16), chunks(5120, 0));
     for (s_values) |s| for ([_]u32{ 1, 2 }) |rt| {
         const nt = tiles(34816, rt, s);
         try std.testing.expect(s * rt * nt * 64 * 4 <= 16384);

zig/src/core/simd_qmm_bits.zig

@@ -61,11 +61,12 @@ pub const Pipelines = struct {
 pub const Call = struct { pipeline: mtl.Pipeline, args: Args, groups: mtl.Size, threads: usize };
 
 /// y[m * ld + j] = x W^T, g64 codes of `bits`: the twin at 1-2 rows when `scalar` (checked equal), else tiles.
-pub fn call(p: *Pipelines, bits: u8, n: usize, k: usize, rows: usize, ld: usize, scalar: bool) !Call {
+pub fn call(p: *Pipelines, bits: u8, n: usize, k: usize, rows: usize, ld: usize, scalar: bool, whole: usize) !Call {
+    // `whole` is the split rank's full width, whose own tree its K chunks keep (0: `n`).
     if (n == 0 or n % 8 != 0 or n > std.math.maxInt(i32) or k == 0 or k % 64 != 0 or k > std.math.maxInt(i32)) return error.UnsupportedSimdShape;
     if (ld < n or ld > std.math.maxInt(i32)) return error.UnsupportedSimdStride;
     if (rows == 0 or rows > std.math.maxInt(i32)) return error.NoRows;
-    const s = simd.splits(n);
+    const s = simd.chunks(n, whole);
     const args = Args{ .k = @intCast(k), .n = @intCast(n), .r = @intCast(rows), .ld = @intCast(ld) };
     if (scalar and simd.scalarBlock(rows, s) != 0) {
         const nr: usize = if (n > 2048) 2 else 1;
@@ -99,8 +100,8 @@ pub fn encode(e: mtl.ComputeEncoder, c: Call, m: Matrix, x: mtl.Buffer, x_off: u
     e.dispatchGroups(c.groups, .{ .width = c.threads });
 }
 
-/// Whether the scalar twin's 1-2 row results equal the tiles' for an [n, k] matrix of `bits` codes.
-pub fn twinExact(p: *Pipelines, device: mtl.Device, bits: u8, m: Matrix, n: usize, k: usize) !bool {
+/// Whether the scalar twin's 1-2 row results equal the tiles' for an [n, k] `bits` matrix, chunked as `call` takes it.
+pub fn twinExact(p: *Pipelines, device: mtl.Device, bits: u8, m: Matrix, n: usize, k: usize, whole: usize) !bool {
     const rows = 8;
     const pool = mtl.objc.Pool.push();
     defer pool.pop();
@@ -118,9 +119,9 @@ pub fn twinExact(p: *Pipelines, device: mtl.Device, bits: u8, m: Matrix, n: usiz
     defer queue.deinit();
     const cb = queue.commandBuffer();
     const e = cb.compute(.serial);
-    encode(e, try call(p, bits, n, k, rows, n, false), m, x, 0, tiled, 0);
-    for (0..rows) |r| encode(e, try call(p, bits, n, k, 1, n, true), m, x, r * k * 2, twin, r * n * 2);
-    for ([_]usize{ 0, rows - 2 }, 0..) |r, i| encode(e, try call(p, bits, n, k, 2, n, true), m, x, r * k * 2, pairs, i * 2 * n * 2);
+    encode(e, try call(p, bits, n, k, rows, n, false, whole), m, x, 0, tiled, 0);
+    for (0..rows) |r| encode(e, try call(p, bits, n, k, 1, n, true, whole), m, x, r * k * 2, twin, r * n * 2);
+    for ([_]usize{ 0, rows - 2 }, 0..) |r, i| encode(e, try call(p, bits, n, k, 2, n, true, whole), m, x, r * k * 2, pairs, i * 2 * n * 2);
     e.end();
     cb.commit();
     cb.wait();

zig/src/families/gemma4/weights.zig

@@ -304,7 +304,7 @@ pub fn load(gpa: std.mem.Allocator, device: mtl.Device, c: *const ckpt.Checkpoin
 pub fn projectRow(ps: *Projections, device: mtl.Device, e: mtl.ComputeEncoder, raw: rowp.Weights, l: Linear, x: lane.Ref, y: mtl.Buffer, rows: usize) !void {
     const p = try ps.simd(device);
     const mtx = rowp.simd_bits.Matrix{ .w = raw.w, .scales = raw.scales, .biases = raw.biases };
-    const call = try rowp.simd_bits.call(p, raw.bits, raw.n, raw.k, rows, l.y_stride orelse l.n, false);
+    const call = try rowp.simd_bits.call(p, raw.bits, raw.n, raw.k, rows, l.y_stride orelse l.n, false, 0);
     rowp.simd_bits.encode(e, call, mtx, x.buf, x.off, y, l.col * 2);
 }
 

zig/src/families/nemotron/forward.zig

@@ -135,7 +135,7 @@ pub const Forward = struct {
         var calls: [3]row.simd_bits.Call = undefined;
         for (parts, 0..) |m, i| {
             if (!tiled(m.rows)) return false;
-            calls[i] = row.simd_bits.call(p, m.rows.bits, m.rows.n, m.rows.k, rows, lin.n, false) catch return false;
+            calls[i] = row.simd_bits.call(p, m.rows.bits, m.rows.n, m.rows.k, rows, lin.n, false, 0) catch return false;
         }
         for (parts, calls[0..parts.len]) |m, c| {
             e.pipe(c.pipeline);

zig/src/families/nemotron/model.zig

@@ -124,10 +124,10 @@ pub const Model = struct {
         var it = m.weights.linears.valueIterator();
         while (it.next()) |lin| if (lin.tiles == null) {
             for (lin.stack[0..lin.stacked]) |s| if (fwd.tiled(s.rows)) {
-                _ = try row.simd_bits.call(p, s.rows.bits, s.rows.n, s.rows.k, prefill.lane_rows + 1, lin.n, false);
+                _ = try row.simd_bits.call(p, s.rows.bits, s.rows.n, s.rows.k, prefill.lane_rows + 1, lin.n, false, 0);
             };
             if (lin.rows) |r| if (fwd.tiled(r)) {
-                _ = try row.simd_bits.call(p, r.bits, r.n, r.k, prefill.lane_rows + 1, r.n, false);
+                _ = try row.simd_bits.call(p, r.bits, r.n, r.k, prefill.lane_rows + 1, r.n, false, 0);
             };
         };
         var gathers: std.ArrayList(prefill_w.Gathers) = .empty;