From 29f91f695e3dcb13c88878987b1cacf10d45e24d Mon Sep 17 00:00:00 2001
From: Elevatormusic <22101396+Elevatormusic@users.noreply.github.com>
Date: Fri, 23 Oct 2026 15:34:51 -0700
Subject: [PATCH 1/5] dsv41: dense bf16 linears of a mixed pack (EAQ v1.1's
 indexer path)

A mixed EXL3 pack stores some whole linears as a plain bf16 <prefix>.weight
[n, k] instead of trellis parts (Mia-AiLab EAQ v1.1: wq_a, the indexer wq_b
and wk, and the compressor wkv and wgate of the indexer layers).

- plan: wantsDense plans a dense linear as one bf16 .weight entry; a split
  (TP column or row part) dense linear is refused. The rank cache builder
  takes the dense set from the shards, the checker from the cache.
- weights: Linear.dense; linear() loads a dense entry when the source has no
  trellis for the prefix.
- prompt / round: dense linears run through cuBLAS GemmEx (bf16 operands,
  COMPUTE_32F, bf16 or fp32 out) in grouped(), mm and the decode q_proj (the
  indexer wq_b from q's normed rows); no counters, rotated rows or partials.
---
 zig/src/families/dsv41/check_main.zig      |  3 +-
 zig/src/families/dsv41/checkpoint.zig      | 14 +++++
 zig/src/families/dsv41/cublas.zig          | 11 ++++
 zig/src/families/dsv41/exl3_linear.zig     |  1 +
 zig/src/families/dsv41/plan.zig            | 71 +++++++++++++++++++++-
 zig/src/families/dsv41/prompt.zig          | 33 ++++++++--
 zig/src/families/dsv41/rank_cache_main.zig |  3 +-
 zig/src/families/dsv41/round.zig           |  6 ++
 zig/src/families/dsv41/weights.zig         | 56 ++++++++++++++++-
 9 files changed, 187 insertions(+), 11 deletions(-)

diff --git a/zig/src/families/dsv41/check_main.zig b/zig/src/families/dsv41/check_main.zig
index 6ac51fafcf..e6b2ba1091 100644
--- a/zig/src/families/dsv41/check_main.zig
+++ b/zig/src/families/dsv41/check_main.zig
@@ -37,7 +37,8 @@ pub fn main(init: std.process.Init) !u8 {
     defer ix.deinit();
     // four ranks: the exact 2D split (node g = TP2 rank g % 2 of pair g / 2); the cache file is still rank g of 4
     const split: dsv41.plan.Split = if (world == 4) .{ .rank = rank % 2, .world = 2, .pair = rank / 2, .parity = parity } else .{ .rank = rank, .world = world };
-    const want = try dsv41.plan.wants(a, cfg, split, dspark);
+    // a mixed pack's dense linears: planned as the cache holds them
+    const want = try dsv41.plan.wantsDense(a, cfg, split, dspark, dsv41.plan.Dense.inCache(&ix));
     var out_buf: [4096]u8 = undefined;
     var out = std.Io.File.stdout().writer(io, &out_buf);
     const w = &out.interface;
diff --git a/zig/src/families/dsv41/checkpoint.zig b/zig/src/families/dsv41/checkpoint.zig
index 383f701592..883ee48290 100644
--- a/zig/src/families/dsv41/checkpoint.zig
+++ b/zig/src/families/dsv41/checkpoint.zig
@@ -79,6 +79,20 @@ pub const Shards = struct {
         const i = sh.names.get(name) orelse return null;
         return sh.files[i].get(name);
     }
+
+    /// The linears this checkpoint holds dense (plan.Dense): a `<prefix>.weight` and no `<prefix>.trellis`.
+    pub fn dense(sh: *const Shards) plan.Dense {
+        return .{ .ctx = sh, .has = hasDense };
+    }
+
+    fn hasDense(ctx: *const anyopaque, prefix: []const u8) bool {
+        const sh: *const Shards = @ptrCast(@alignCast(ctx));
+        var buf: [256]u8 = undefined;
+        const tr = std.fmt.bufPrint(&buf, "{s}.trellis", .{prefix}) catch return false;
+        if (sh.names.contains(tr)) return false;
+        const w = std.fmt.bufPrint(&buf, "{s}.weight", .{prefix}) catch return false;
+        return sh.names.contains(w);
+    }
 };
 
 /// A loader key split into its fields: "name|dtype" (plain) or "prefix|cols|rows|part" (an EXL3 part).
diff --git a/zig/src/families/dsv41/cublas.zig b/zig/src/families/dsv41/cublas.zig
index 0411372578..5019eb35a0 100644
--- a/zig/src/families/dsv41/cublas.zig
+++ b/zig/src/families/dsv41/cublas.zig
@@ -73,6 +73,17 @@ pub const Blas = struct {
         try b.check(b.api.cublasGemmEx(b.handle, .t, .n, @intCast(n), @intCast(rows), @intCast(k), &one, w, bf16_type, @intCast(k), x, bf16_type, @intCast(k), &zero, out, bf16_type, @intCast(n), 68, 99));
     }
 
+    /// A dense linear of a mixed pack (weights.Linear.dense): x bf16 [rows, k] (row stride ldx) @ w bf16 [n, k]^T into
+    /// out [rows, n] (row stride ldo), bf16 or fp32. linearBf16's call (COMPUTE_32F, fp32 alpha 1 and beta 0,
+    /// CUBLAS_GEMM_DEFAULT_TENSOR_OP) with the strides and the output dtype as the caller's buffers have them.
+    pub fn linearDense(b: *const Blas, x: u64, ldx: usize, w: u64, out: u64, ldo: usize, out_f32: bool, rows: usize, k: usize, n: usize) Error!void {
+        const one: f32 = 1;
+        const zero: f32 = 0;
+        const bf16_type: c_int = 14; // CUDA_R_16BF
+        const out_type: c_int = if (out_f32) 0 else bf16_type; // CUDA_R_32F
+        try b.check(b.api.cublasGemmEx(b.handle, .t, .n, @intCast(n), @intCast(rows), @intCast(k), &one, w, bf16_type, @intCast(k), x, bf16_type, @intCast(ldx), &zero, out, out_type, @intCast(ldo), 68, 99));
+    }
+
     /// torch.matmul of x [rows, k] fp32 (row-major) and w [n, k] fp32 transposed: out [rows, n] fp32, as torch issues
     /// it (column-major C^T = W x^T: transa T, transb N, m = n, n = rows, lda = ldb = k, ldc = n, alpha 1, beta 0).
     pub fn xwT(b: *const Blas, x: u64, w: u64, out: u64, rows: usize, k: usize, n: usize) Error!void {
diff --git a/zig/src/families/dsv41/exl3_linear.zig b/zig/src/families/dsv41/exl3_linear.zig
index 25a9241925..21d83cdfff 100644
--- a/zig/src/families/dsv41/exl3_linear.zig
+++ b/zig/src/families/dsv41/exl3_linear.zig
@@ -224,6 +224,7 @@ pub fn glinear(k: *const Kernels, stream: cuda.Stream, calls: []const Call, m: u
 pub fn zFloats(layers: []const weights.Linear, m: usize) usize {
     var n: usize = 0;
     for (layers) |l| {
+        if (l.isDense()) continue; // runs through cuBLAS (prompt.denseRows)
         const p = plan(l.k, planN(l));
         if (p.sk > 1) n += p.sk * m * l.n;
     }
diff --git a/zig/src/families/dsv41/plan.zig b/zig/src/families/dsv41/plan.zig
index 0b045244cf..a3a186a9bc 100644
--- a/zig/src/families/dsv41/plan.zig
+++ b/zig/src/families/dsv41/plan.zig
@@ -7,7 +7,8 @@
 //! and keeps pair p's whole 128-column blocks of each of its rank's split outputs; DSpark stays TP2 inside the pair.
 const std = @import("std");
 const Config = @import("config.zig").Config;
-const DType = @import("rank_cache.zig").DType;
+const rank_cache = @import("rank_cache.zig");
+const DType = rank_cache.DType;
 
 /// What a split means for one rank.
 pub const Split = struct {
@@ -138,9 +139,34 @@ pub fn partKey(buf: []u8, prefix: []const u8, cols: ?[2]usize, rows: ?[2]usize,
     return std.fmt.bufPrint(buf, "{s}|{s}|{s}|{s}", .{ prefix, cs, rs, part });
 }
 
+/// A dense linear's key: its plain `<prefix>.weight` in the checkpoint's own dtype (bf16 [n, k]).
+pub fn denseKey(buf: []u8, prefix: []const u8) ![]const u8 {
+    return std.fmt.bufPrint(buf, "{s}.weight|None", .{prefix});
+}
+
+/// Which linears a mixed pack holds dense (a plain `<prefix>.weight` and no trellis): `has(ctx, prefix)`. A whole
+/// linear the predicate names is planned as its dense entry; a split one (2D / TP column or row parts) is refused.
+pub const Dense = struct {
+    ctx: *const anyopaque,
+    has: *const fn (ctx: *const anyopaque, prefix: []const u8) bool,
+
+    /// The linears a rank cache holds dense: a dense entry and no trellis part.
+    pub fn inCache(ix: *const rank_cache.Index) Dense {
+        return .{ .ctx = ix, .has = cacheHas };
+    }
+
+    fn cacheHas(ctx: *const anyopaque, prefix: []const u8) bool {
+        const ix: *const rank_cache.Index = @ptrCast(@alignCast(ctx));
+        var buf: [256]u8 = undefined;
+        if (ix.get(partKey(&buf, prefix, null, null, "tr") catch return false) != null) return false;
+        return ix.get(denseKey(&buf, prefix) catch return false) != null;
+    }
+};
+
 const Builder = struct {
     a: std.mem.Allocator,
     out: std.ArrayList(Want) = .empty,
+    dense: ?Dense = null,
 
     fn plain(b: *Builder, tensor: []const u8, dtype: ?DType) !void {
         var buf: [256]u8 = undefined;
@@ -148,6 +174,11 @@ const Builder = struct {
     }
 
     fn exl3(b: *Builder, prefix: []const u8, cols: ?[2]usize, rows: ?[2]usize, k: ?usize, n: ?usize) !void {
+        if (b.dense) |d| if (d.has(d.ctx, prefix)) {
+            if (cols != null or rows != null) return error.DenseLinearSplit;
+            var buf: [256]u8 = undefined;
+            return b.out.append(b.a, .{ .key = try b.a.dupe(u8, try denseKey(&buf, prefix)), .dtype = .bf16 });
+        };
         const k16: ?usize = if (rows) |r| (r[1] - r[0]) / 16 else if (k) |x| x / 16 else null;
         const n16: ?usize = if (cols) |c| (c[1] - c[0]) / 16 else if (n) |x| x / 16 else null;
         for ([_][]const u8{ "tr", "suh", "svh" }) |part| {
@@ -220,7 +251,12 @@ fn block(b: *Builder, c: Config, s: Split, p: []const u8, i: usize, experts: usi
 
 /// Every tensor rank `s` loads, in load order (`dspark`: the DSpark stages and heads too). Keys live in `a`.
 pub fn wants(a: std.mem.Allocator, c: Config, s: Split, dspark: bool) ![]Want {
-    var b: Builder = .{ .a = a };
+    return wantsDense(a, c, s, dspark, null);
+}
+
+/// wants, the linears `dense` names planned as their dense entries (a mixed pack).
+pub fn wantsDense(a: std.mem.Allocator, c: Config, s: Split, dspark: bool, dense: ?Dense) ![]Want {
+    var b: Builder = .{ .a = a, .dense = dense };
     try b.plain("embed.weight", null);
     try b.plain("norm.weight", null);
     try b.exl3("head", s.headCols(c), null, c.hidden, null);
@@ -266,6 +302,37 @@ test "a rank's tensors in load order: the split's ranges, keys as Python prints
     try std.testing.expectEqualStrings("mtp.2.confidence_head.proj.weight|None", w[w.len - 1].key);
 }
 
+test "a mixed pack's dense linears: one bf16 .weight entry for a whole linear, a split one refused" {
+    var arena: std.heap.ArenaAllocator = .init(std.testing.allocator);
+    defer arena.deinit();
+    const a = arena.allocator();
+    var why: @import("config.zig").Why = .{};
+    const c = try @import("config.zig").parse(a, @import("config.zig").test_config, &why);
+    const names = struct {
+        fn has(_: *const anyopaque, prefix: []const u8) bool {
+            return std.mem.eql(u8, prefix, "layers.2.attn.wq_a") or std.mem.eql(u8, prefix, "layers.2.attn.indexer.wq_b");
+        }
+        fn split(_: *const anyopaque, prefix: []const u8) bool {
+            return std.mem.eql(u8, prefix, "layers.0.attn.wq_b");
+        }
+    };
+    const w = try wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &names, .has = names.has });
+    const plain = try wants(a, c, .{ .rank = 0, .world = 2 }, true);
+    // each dense linear: its three parts give way to one entry
+    try std.testing.expectEqual(plain.len - 2 * 2, w.len);
+    var seen: usize = 0;
+    for (w) |x| {
+        if (std.mem.startsWith(u8, x.key, "layers.2.attn.wq_a|") or std.mem.startsWith(u8, x.key, "layers.2.attn.indexer.wq_b|")) return error.PartsLeft;
+        if (std.mem.eql(u8, x.key, "layers.2.attn.wq_a.weight|None") or std.mem.eql(u8, x.key, "layers.2.attn.indexer.wq_b.weight|None")) {
+            try std.testing.expectEqual(@as(?DType, .bf16), x.dtype);
+            seen += 1;
+        }
+    }
+    try std.testing.expectEqual(@as(usize, 2), seen);
+    // wq_b is a column slice of the rank's heads: no dense form of it
+    try std.testing.expectError(error.DenseLinearSplit, wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &names, .has = names.split }));
+}
+
 test "each split tensor's ranges: TP2's slices, a 2D node's pair parts of them" {
     var arena: std.heap.ArenaAllocator = .init(std.testing.allocator);
     defer arena.deinit();
diff --git a/zig/src/families/dsv41/prompt.zig b/zig/src/families/dsv41/prompt.zig
index 2fb961c169..38987bc0db 100644
--- a/zig/src/families/dsv41/prompt.zig
+++ b/zig/src/families/dsv41/prompt.zig
@@ -498,6 +498,7 @@ fn layerLinears(e: *const Engine, ch: *Chunk, a: *Arena, lay: weights.Layer, max
 /// One linear alone at 128 rows in the chunk's grouped-linear sizes, and its split-K counters (int32 [8 N / 128], zeros,
 /// left zero by each launch).
 fn addLinear(e: *const Engine, ch: *Chunk, a: *Arena, l: weights.Linear, max_gk: *usize, max_gz: *usize) !void {
+    if (l.isDense()) return; // cuBLAS: no rotated rows, partials or counters
     max_gk.* = @max(max_gk.*, l.k);
     max_gz.* = @max(max_gz.*, exl3_linear.zFloats(&[_]weights.Linear{l}, 128));
     if (ch.n_counters == ch.counters.len) return error.TooManyLinears;
@@ -555,19 +556,40 @@ pub fn counterOf(ch: *const Chunk, l: weights.Linear) !u64 {
 pub fn grouped(e: *const Engine, ch: *const Chunk, m: usize, ls: []const weights.Linear, xs: []const u64, ldxs: []const usize, outs: []const u64, ldos: []const usize, types: []const exl3_linear.DType) !void {
     if (m == 0 or m > 128 or ls.len > 8) return error.NotADecodeWindow;
     var calls: [8]exl3_linear.Call = undefined;
+    var n: usize = 0;
     var off: usize = 0;
     for (ls, 0..) |l, i| {
-        calls[i] = .{ .layer = l, .x = xs[i], .ldx = @intCast(ldxs[i]), .x_dtype = .bf16, .xh = ch.gxh + off * 2, .y = outs[i], .ldy = @intCast(ldos[i]), .y_dtype = types[i], .counters = try counterOf(ch, l) };
+        // a mixed pack's dense linears through cuBLAS, the group's EXL3 ones as before
+        if (l.isDense()) {
+            try denseRows(e, m, l, xs[i], ldxs[i], outs[i], ldos[i], types[i]);
+            continue;
+        }
+        calls[n] = .{ .layer = l, .x = xs[i], .ldx = @intCast(ldxs[i]), .x_dtype = .bf16, .xh = ch.gxh + off * 2, .y = outs[i], .ldy = @intCast(ldos[i]), .y_dtype = types[i], .counters = try counterOf(ch, l) };
         off += m * l.k;
+        n += 1;
     }
-    try exl3_linear.rotMany(e.lin, e.s, calls[0..ls.len], m, true);
-    try exl3_linear.glinear(e.lin, e.s, calls[0..ls.len], m, ch.gz, true, true);
+    if (n == 0) return;
+    try exl3_linear.rotMany(e.lin, e.s, calls[0..n], m, true);
+    try exl3_linear.glinear(e.lin, e.s, calls[0..n], m, ch.gz, true, true);
+}
+
+/// A dense linear (weights.Linear.dense) on m rows: x bf16 (row stride ldx) into out (row stride ldo, bf16 or fp32).
+pub fn denseRows(e: *const Engine, m: usize, l: weights.Linear, x: u64, ldx: usize, out: u64, ldo: usize, dt: exl3_linear.DType) !void {
+    if (!l.isDense()) return error.NotDense;
+    const f32_out = switch (dt) {
+        .bf16 => false,
+        .f32 => true,
+        .f16 => return error.UnsupportedDType,
+    };
+    try e.blas.linearDense(x, ldx, l.dense, out, ldo, f32_out, m, l.k, l.n);
 }
 
 /// Exl3Group.rotated on m rows already rotated into each call's xh (a kernel that folds the rotation in wrote them):
-/// the glinear launches alone, each call's counters and folds as given.
+/// the glinear launches alone, each call's counters and folds as given. A dense linear has no rotated rows: its
+/// caller runs it from the normed rows (denseRows).
 pub fn groupedRotated(e: *const Engine, ch: *const Chunk, m: usize, calls: []exl3_linear.Call) !void {
     if (m == 0 or m > 128 or calls.len > exl3_linear.gmax) return error.NotADecodeWindow;
+    for (calls) |c| if (c.layer.isDense()) return error.DenseNotRotated;
     for (calls) |*c| c.counters = try counterOf(ch, c.layer);
     try exl3_linear.glinear(e.lin, e.s, calls, m, ch.gz, true, true);
 }
@@ -711,11 +733,12 @@ pub fn attnMixes(e: *const Engine, ch: *Chunk, li: usize) !void {
 
 /// mm(layer, x, out_dtype) of m rows: the prompt GEMM beyond 128 rows, the layer's own group (GROUPED) up to 128.
 fn mmRows(e: *const Engine, ch: *const Chunk, m: usize, l: weights.Linear, x: u64, ldx: usize, out: u64, out_type: tri_markov.OutType, os: usize) !void {
-    if (m > 128) return exl3_prefill.matmul(e.pf, e.t, ch.ws, l, x, .bf16, ldx, out, out_type, os, m);
     const dt: exl3_linear.DType = switch (out_type) {
         .bf16 => .bf16,
         .fp32 => .f32,
     };
+    if (l.isDense()) return denseRows(e, m, l, x, ldx, out, os, dt);
+    if (m > 128) return exl3_prefill.matmul(e.pf, e.t, ch.ws, l, x, .bf16, ldx, out, out_type, os, m);
     try grouped(e, ch, m, &.{l}, &.{x}, &.{ldx}, &.{out}, &.{os}, &.{dt});
 }
 
diff --git a/zig/src/families/dsv41/rank_cache_main.zig b/zig/src/families/dsv41/rank_cache_main.zig
index 879be3ccdc..2a1a3e4be2 100644
--- a/zig/src/families/dsv41/rank_cache_main.zig
+++ b/zig/src/families/dsv41/rank_cache_main.zig
@@ -49,9 +49,10 @@ pub fn main(init: std.process.Init) !u8 {
     // four ranks: the exact 2D split (node g = TP2 rank g % 2 of pair g / 2); the file is still rank g of 4
     const split: dsv41.plan.Split = if (world == 4) .{ .rank = rank % 2, .world = 2, .pair = rank / 2, .parity = parity } else .{ .rank = rank, .world = world };
     const cfg = try dsv41.Config.read(a, io, args[1]);
-    const want = try dsv41.plan.wants(a, cfg, split, dspark);
     var sh = try dsv41.checkpoint.Shards.open(gpa, io, args[1]);
     defer sh.close(io);
+    // a mixed pack's dense linears (a plain <prefix>.weight, no trellis): one bf16 entry each
+    const want = try dsv41.plan.wantsDense(a, cfg, split, dspark, sh.dense());
 
     // every entry's layout first, and from the layouts the index: a missing tensor or a layout the split does not
     // expect stops before a byte is written
diff --git a/zig/src/families/dsv41/round.zig b/zig/src/families/dsv41/round.zig
index c674aa0c3f..b7de6ffa71 100644
--- a/zig/src/families/dsv41/round.zig
+++ b/zig/src/families/dsv41/round.zig
@@ -590,6 +590,7 @@ fn attention(e: *const Engine, rd: *Round, ch: *const Chunk, cs: *const Caches,
         const outs = [_]u64{ rd.q, rd.iq };
         for (qls) |ol| {
             const l = ol orelse continue;
+            if (l.isDense()) continue; // a mixed pack's dense linear: from q's normed rows below
             const xh = rd.xhq + n * R * c.q_lora * 2;
             rot[n] = .{ .suh = l.suh, .h = xh };
             calls[n] = .{ .layer = l, .x = 0, .ldx = 0, .x_dtype = .bf16, .xh = xh, .y = outs[n], .ldy = @intCast(l.n), .y_dtype = .bf16, .counters = 0 };
@@ -598,6 +599,10 @@ fn attention(e: *const Engine, rd: *Round, ch: *const Chunk, cs: *const Caches,
         }
         try tri_norm.qKvNorm(t, rd.qa, lay.wq_a.n, lay.q_norm, c.eps, rot[0..n], rd.ykv, lay.kv_norm, rope.cos, rope.sin, pos, ring, ring_rows, rd.g(.wslot), true, rd_dim, rd.qr, rd.kout, R, lay.wq_a.n, hd);
         try prompt.groupedRotated(e, ch, R, calls[0..n]);
+        for (qls, outs) |ol, out| {
+            const l = ol orelse continue;
+            if (l.isDense()) try prompt.denseRows(e, R, l, rd.qr, lay.wq_a.n, out, l.n, .bf16);
+        }
     }
     // the prefetch: this layer's wo_a and wo_b behind q's linears
     {
@@ -798,6 +803,7 @@ fn moe(e: *const Engine, rd: *Round, ch: *const Chunk, li: usize) !void {
 
 /// An EXL3 linear's words: their first byte and length (k * n values of k2 / 2 bits).
 fn linearBytes(l: weights.Linear) [2]u64 {
+    if (l.isDense()) return .{ l.dense, @as(u64, l.k) * l.n * 2 };
     return .{ l.words, @as(u64, l.k) * l.n * l.k2 / 16 };
 }
 
diff --git a/zig/src/families/dsv41/weights.zig b/zig/src/families/dsv41/weights.zig
index 3997ca531d..ebef6d7203 100644
--- a/zig/src/families/dsv41/weights.zig
+++ b/zig/src/families/dsv41/weights.zig
@@ -17,7 +17,22 @@ const span_piece = 64 << 20;
 /// An EXL3 linear in the strips layout: words int32 [N/128, K/16, 8, 8 * bits], fp16 suh [K] and svh [N].
 /// plan_n: the N its decode plan is made for, when not n: a 2D node's column part of a TP2 slice keeps that slice's
 /// plan (linear.py pinned plans), so its columns sum in the slice's K order.
-pub const Linear = struct { words: u64, suh: u64, svh: u64, k: u32, n: u32, k2: u32, plan_n: u32 = 0 };
+/// An EXL3 linear (words, suh, svh), or a dense one: `dense` bf16 [n, k] row-major (a mixed pack stores some linears
+/// as plain `<prefix>.weight` tensors at source precision, EAQ v1.1's indexer path), words / suh / svh / k2 then 0.
+pub const Linear = struct {
+    words: u64,
+    suh: u64,
+    svh: u64,
+    k: u32,
+    n: u32,
+    k2: u32,
+    plan_n: u32 = 0,
+    dense: u64 = 0,
+
+    pub fn isDense(l: Linear) bool {
+        return l.dense != 0;
+    }
+};
 
 /// A layer's routed experts and its shared expert (the last entry), as exl3_experts.prepare lays them out. On a 2D node
 /// (exl3/experts2d.py prepare) gate / up hold its intermediate blocks (width of the half's down_k) and down its
@@ -305,10 +320,12 @@ const Loader = struct {
         return L.upload(tensor, L.aux.items);
     }
 
-    /// linear(prefix, cols, rows): Exl3Linear.from_tensors in strips.
+    /// linear(prefix, cols, rows): Exl3Linear.from_tensors in strips, or a whole linear the pack holds dense (see
+    /// denseLinear).
     fn linear(L: *Loader, prefix: []const u8, cols: ?[2]usize, rows: ?[2]usize) !Linear {
         var kb: [256]u8 = undefined;
         var nb: [256]u8 = undefined;
+        if (cols == null and rows == null) if (try L.denseLinear(prefix)) |d| return d;
         const tr = try L.entry(try plan.partKey(&kb, prefix, cols, rows, "tr"));
         const t = try L.shape(try plan.partKey(&kb, prefix, cols, rows, "tr"));
         const kt = t[0];
@@ -323,6 +340,41 @@ const Loader = struct {
         return .{ .words = words, .suh = suh, .svh = svh, .k = @intCast(kt * 16), .n = @intCast(nt * 16), .k2 = @intCast(w16 / 8) };
     }
 
+    /// A whole linear the source holds as a plain bf16 `<prefix>.weight` [n, k] (plan.zig's dense entry) instead of
+    /// trellis parts: uploaded as it is. Null when the source has its trellis (every linear of an all-EXL3 pack).
+    fn denseLinear(L: *Loader, prefix: []const u8) !?Linear {
+        var kb: [256]u8 = undefined;
+        var nb: [256]u8 = undefined;
+        if (L.has(try plan.partKey(&kb, prefix, null, null, "tr"))) return null;
+        const key = try plan.denseKey(&kb, prefix);
+        if (!L.has(key)) return null;
+        const s = try L.dims(key);
+        if (s.dtype != .bf16 or s.rank != 2 or s.n % 8 != 0 or s.k % 8 != 0) return error.UnexpectedTensor;
+        const w = try L.upload(try std.fmt.bufPrint(&nb, "{s}.weight", .{prefix}), try L.entry(key));
+        return .{ .words = 0, .suh = 0, .svh = 0, .k = @intCast(s.k), .n = @intCast(s.n), .k2 = 0, .dense = w };
+    }
+
+    fn has(L: *Loader, key: []const u8) bool {
+        return switch (L.src) {
+            .cache => |c| c.index.get(key) != null,
+            .shards => |sh| if (checkpoint.describe(sh, key)) |_| true else |_| false,
+        };
+    }
+
+    /// An entry's dtype, rank and first two dimensions ([n, k] for a dense linear).
+    fn dims(L: *Loader, key: []const u8) !struct { dtype: rank_cache.DType, rank: u8, n: usize, k: usize } {
+        return switch (L.src) {
+            .cache => |c| blk: {
+                const e = c.index.get(key) orelse return error.MissingEntry;
+                break :blk .{ .dtype = e.dtype, .rank = e.rank, .n = e.shape[0], .k = e.shape[1] };
+            },
+            .shards => |sh| blk: {
+                const s = try checkpoint.describe(sh, key);
+                break :blk .{ .dtype = s.dtype, .rank = s.rank, .n = s.shape[0], .k = s.shape[1] };
+            },
+        };
+    }
+
     /// A trellis entry's [kt, nt, w16].
     fn shape(L: *Loader, key: []const u8) ![3]usize {
         return switch (L.src) {

From 64bb1e48a08ef53d99cce59788b1843c444bcf15 Mon Sep 17 00:00:00 2001
From: Elevatormusic <22101396+Elevatormusic@users.noreply.github.com>
Date: Fri, 23 Oct 2026 15:35:24 -0700
Subject: [PATCH 2/5] dsv41: the dense plan test's context is a byte, not a
 type

---
 zig/src/families/dsv41/plan.zig | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/zig/src/families/dsv41/plan.zig b/zig/src/families/dsv41/plan.zig
index a3a186a9bc..937af75f08 100644
--- a/zig/src/families/dsv41/plan.zig
+++ b/zig/src/families/dsv41/plan.zig
@@ -316,7 +316,8 @@ test "a mixed pack's dense linears: one bf16 .weight entry for a whole linear, a
             return std.mem.eql(u8, prefix, "layers.0.attn.wq_b");
         }
     };
-    const w = try wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &names, .has = names.has });
+    const ctx: u8 = 0;
+    const w = try wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &ctx, .has = names.has });
     const plain = try wants(a, c, .{ .rank = 0, .world = 2 }, true);
     // each dense linear: its three parts give way to one entry
     try std.testing.expectEqual(plain.len - 2 * 2, w.len);
@@ -330,7 +331,7 @@ test "a mixed pack's dense linears: one bf16 .weight entry for a whole linear, a
     }
     try std.testing.expectEqual(@as(usize, 2), seen);
     // wq_b is a column slice of the rank's heads: no dense form of it
-    try std.testing.expectError(error.DenseLinearSplit, wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &names, .has = names.split }));
+    try std.testing.expectError(error.DenseLinearSplit, wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &ctx, .has = names.split }));
 }
 
 test "each split tensor's ranges: TP2's slices, a 2D node's pair parts of them" {

From 5ea721f071c2d1a1191ac21234c412f8a4b4ba82 Mon Sep 17 00:00:00 2001
From: Elevatormusic <22101396+Elevatormusic@users.noreply.github.com>
Date: Sat, 24 Oct 2026 11:17:09 -0700
Subject: [PATCH 3/5] zig: zig fmt on five dsv41-zig files (the Zig format
 check fails on every pull request)

zig fmt 0.17.0 changes only layout and @intFromEnum to @backingInt.
The Zig format check fails on these five base-branch files in every
open dsv41-zig pull request, also where the pull request does not touch them.
---
 zig/src/families/dsv41/exl3_linear.zig      | 16 ++++++++--------
 zig/src/families/dsv41/round.zig            |  5 ++---
 zig/src/families/dsv41/split_bench_main.zig | 19 ++++++++++++++++---
 zig/src/families/dsv41/tri_basic.zig        |  6 +++---
 zig/src/server/effort_values_test.zig       |  2 +-
 5 files changed, 30 insertions(+), 18 deletions(-)

diff --git a/zig/src/families/dsv41/exl3_linear.zig b/zig/src/families/dsv41/exl3_linear.zig
index 25a9241925..dcce543ad1 100644
--- a/zig/src/families/dsv41/exl3_linear.zig
+++ b/zig/src/families/dsv41/exl3_linear.zig
@@ -164,7 +164,7 @@ pub fn rotMany(k: *const Kernels, stream: cuda.Stream, calls: []const Call, m: u
     var a: RArgs = .{ .n = @intCast(calls.len), .M = @intCast(m) };
     var warps: usize = 0;
     for (calls, 0..) |c, i| {
-        a.l[i] = .{ .x = c.x, .suh = c.layer.suh, .xh = c.xh, .ldx = c.ldx, .x_dtype = @intFromEnum(c.x_dtype), .K = @intCast(c.layer.k), .first = @intCast(warps) };
+        a.l[i] = .{ .x = c.x, .suh = c.layer.suh, .xh = c.xh, .ldx = c.ldx, .x_dtype = @backingInt(c.x_dtype), .K = @intCast(c.layer.k), .first = @intCast(warps) };
         warps += m * (c.layer.k / 128);
     }
     var args: cuda.Args = .{};
@@ -189,7 +189,7 @@ pub fn glinear(k: *const Kernels, stream: cuda.Stream, calls: []const Call, m: u
             done[j] = true;
             const s = strides(c.layer);
             const n: usize = @intCast(a.n);
-            a.l[n] = .{ .xh = c.xh, .ldx = @intCast(c.layer.k), .T = c.layer.words, .stride_k = s[0], .stride_nb = s[1], .svh = c.layer.svh, .y = c.y, .y_dtype = @intFromEnum(c.y_dtype), .ldy = c.ldy, .K = @intCast(c.layer.k), .N = @intCast(c.layer.n), .SK = @intCast(pj.sk), .counters = c.counters, .first = @intCast(blocks) };
+            a.l[n] = .{ .xh = c.xh, .ldx = @intCast(c.layer.k), .T = c.layer.words, .stride_k = s[0], .stride_nb = s[1], .svh = c.layer.svh, .y = c.y, .y_dtype = @backingInt(c.y_dtype), .ldy = c.ldy, .K = @intCast(c.layer.k), .N = @intCast(c.layer.n), .SK = @intCast(pj.sk), .counters = c.counters, .first = @intCast(blocks) };
             if (c.rot) |r| {
                 a.l[n].rsuh = r.suh;
                 a.l[n].rxh = r.xh;
@@ -243,12 +243,12 @@ test "plans as linear.py makes them" {
 test "plans and strides equal every recorded glinear layer of the served build" {
     // (K, N, SK, WK, K2, stride_k, stride_nb) of each distinct layer the recording's glinear calls passed
     const seen = [_][7]u32{
-        .{ 512, 128, 1, 4, 16, 512, 16384 },       .{ 1280, 4096, 2, 4, 10, 320, 25600 },     .{ 1280, 16384, 1, 8, 8, 256, 20480 },
-        .{ 1280, 16384, 1, 8, 10, 320, 25600 },    .{ 3072, 25600, 1, 8, 8, 256, 49152 },     .{ 3072, 25600, 1, 8, 10, 320, 61440 },
-        .{ 4096, 1024, 8, 4, 8, 256, 65536 },      .{ 4096, 1024, 8, 4, 10, 320, 81920 },     .{ 4096, 5120, 8, 4, 8, 256, 65536 },
-        .{ 4096, 5120, 8, 4, 10, 320, 81920 },     .{ 5120, 512, 8, 4, 8, 256, 81920 },       .{ 5120, 512, 8, 4, 10, 320, 102400 },
-        .{ 5120, 512, 8, 4, 12, 384, 122880 },     .{ 5120, 1280, 8, 4, 8, 256, 81920 },      .{ 5120, 1280, 8, 4, 10, 320, 102400 },
-        .{ 5120, 1280, 8, 4, 12, 384, 122880 },    .{ 5120, 64640, 1, 8, 12, 384, 122880 },   .{ 15360, 5120, 8, 4, 8, 256, 245760 },
+        .{ 512, 128, 1, 4, 16, 512, 16384 },    .{ 1280, 4096, 2, 4, 10, 320, 25600 },   .{ 1280, 16384, 1, 8, 8, 256, 20480 },
+        .{ 1280, 16384, 1, 8, 10, 320, 25600 }, .{ 3072, 25600, 1, 8, 8, 256, 49152 },   .{ 3072, 25600, 1, 8, 10, 320, 61440 },
+        .{ 4096, 1024, 8, 4, 8, 256, 65536 },   .{ 4096, 1024, 8, 4, 10, 320, 81920 },   .{ 4096, 5120, 8, 4, 8, 256, 65536 },
+        .{ 4096, 5120, 8, 4, 10, 320, 81920 },  .{ 5120, 512, 8, 4, 8, 256, 81920 },     .{ 5120, 512, 8, 4, 10, 320, 102400 },
+        .{ 5120, 512, 8, 4, 12, 384, 122880 },  .{ 5120, 1280, 8, 4, 8, 256, 81920 },    .{ 5120, 1280, 8, 4, 10, 320, 102400 },
+        .{ 5120, 1280, 8, 4, 12, 384, 122880 }, .{ 5120, 64640, 1, 8, 12, 384, 122880 }, .{ 15360, 5120, 8, 4, 8, 256, 245760 },
     };
     for (seen) |r| {
         const p = plan(r[0], r[1]);
diff --git a/zig/src/families/dsv41/round.zig b/zig/src/families/dsv41/round.zig
index c674aa0c3f..59dac0c8de 100644
--- a/zig/src/families/dsv41/round.zig
+++ b/zig/src/families/dsv41/round.zig
@@ -446,9 +446,9 @@ pub const Round = struct {
     fn mark(rd: *Round, e: *const Engine, p: PhaseTimer.Phase) !void {
         if (rd.clk != 0 and rd.clk_marks and (rd.clk_ratio == null or rd.clk_ratio.? == rd.cur_ratio)) {
             const now = cuda.launch.count;
-            rd.launches[@intFromEnum(p)] += now -% rd.launch_mark;
+            rd.launches[@backingInt(p)] += now -% rd.launch_mark;
             rd.launch_mark = now;
-            try rd.clock(e, @intFromEnum(p));
+            try rd.clock(e, @backingInt(p));
         }
         if (rd.timer) |t| try t.mark(e.s, p);
     }
@@ -801,7 +801,6 @@ fn linearBytes(l: weights.Linear) [2]u64 {
     return .{ l.words, @as(u64, l.k) * l.n * l.k2 / 16 };
 }
 
-
 /// A measurement's L2 prefetches left out (TF_DS_L2_SKIP, read at Model.open: bit 1 the "attn" fork after q, 2 the "moe"
 /// fork after wo_b, 4 the "next" fork after the experts); 0 in service.
 pub var l2_skip: u8 = 0;
diff --git a/zig/src/families/dsv41/split_bench_main.zig b/zig/src/families/dsv41/split_bench_main.zig
index fc8952dd2a..7980a7e5ef 100644
--- a/zig/src/families/dsv41/split_bench_main.zig
+++ b/zig/src/families/dsv41/split_bench_main.zig
@@ -320,10 +320,23 @@ pub fn main(init: std.process.Init) !u8 {
             try scales.upload(0, hs);
         }
         const run: Run = .{
-            .t = t, .o = &o, .s = stream, .top = top, .n = n, .rb = rb,
+            .t = t,
+            .o = &o,
+            .s = stream,
+            .top = top,
+            .n = n,
+            .rb = rb,
             .kx = .{ .fp4 = .{ .codes = codes.ptr, .scales = scales.ptr } },
-            .q = q.ptr, .w = w.ptr, .vis = vis.ptr, .nvis = nvis.ptr, .keys = keys.ptr, .tmax = tmax.ptr, .bufs = bufs,
-            .parts = parts.ptr, .merged = merged.ptr, .vblk = vblk.ptr,
+            .q = q.ptr,
+            .w = w.ptr,
+            .vis = vis.ptr,
+            .nvis = nvis.ptr,
+            .keys = keys.ptr,
+            .tmax = tmax.ptr,
+            .bufs = bufs,
+            .parts = parts.ptr,
+            .merged = merged.ptr,
+            .vblk = vblk.ptr,
         };
 
         // every output byte, both splits, both patterns; the served keys past vis
diff --git a/zig/src/families/dsv41/tri_basic.zig b/zig/src/families/dsv41/tri_basic.zig
index 66cb3470a5..34acadee40 100644
--- a/zig/src/families/dsv41/tri_basic.zig
+++ b/zig/src/families/dsv41/tri_basic.zig
@@ -19,9 +19,9 @@ fn ci(name: []const u8, v: usize) aot.Const {
 pub fn hcPre(t: Tri, h: u64, fn_w: u64, scale: u64, base: u64, pre_in: u64, norm_w: u64, eps: f32, hc_eps: f32, iters: usize, out: u64, pre_out: u64, post: u64, comb: u64, part: u64, rows: usize, d: usize) !void {
     try t.run("_hc_partial", .{ u(rows), hc_blocks, 1 }, &.{ p("X", "*bf16", h), p("FN", "*fp32", fn_w), p("PART", "*fp32", part) }, &.{ ci("WIDE", 4 * d), ci("NB", hc_blocks), ci("SUB", 128) });
     try t.run("_hc_finish", .{ u(rows), 1, 1 }, &.{
-        p("X", "*bf16", h),          p("PART", "*fp32", part),       p("BASE", "*fp32", base), p("SCALE", "*fp32", scale),
-        p("PRE_IN", "*fp32", pre_in), p("NW", "*bf16", norm_w),      p("OUT", "*bf16", out),   p("PRE_OUT", "*fp32", pre_out),
-        p("POST", "*fp32", post),    p("COMB", "*fp32", comb),       aot.float("eps", eps),    aot.float("hc_eps", hc_eps),
+        p("X", "*bf16", h),           p("PART", "*fp32", part), p("BASE", "*fp32", base), p("SCALE", "*fp32", scale),
+        p("PRE_IN", "*fp32", pre_in), p("NW", "*bf16", norm_w), p("OUT", "*bf16", out),   p("PRE_OUT", "*fp32", pre_out),
+        p("POST", "*fp32", post),     p("COMB", "*fp32", comb), aot.float("eps", eps),    aot.float("hc_eps", hc_eps),
     }, &.{ ci("D", d), ci("NB", hc_blocks), ci("ITERS", iters), ci("BLOCK", 1024) });
 }
 
diff --git a/zig/src/server/effort_values_test.zig b/zig/src/server/effort_values_test.zig
index f8b01f8ae9..e3ec47b9ef 100644
--- a/zig/src/server/effort_values_test.zig
+++ b/zig/src/server/effort_values_test.zig
@@ -10,7 +10,7 @@ const Conn = @import("http_conn.zig").Conn;
 
 /// DeepSeek-V4.1's numbers (families/dsv41/text.zig; the served lane's EFFORT_VALUES).
 const numbers = [_]api.Effort{
-    .{ .name = "minimal", .value = 25 }, .{ .name = "low", .value = 50 },  .{ .name = "medium", .value = 62 },
+    .{ .name = "minimal", .value = 25 }, .{ .name = "low", .value = 50 },   .{ .name = "medium", .value = 62 },
     .{ .name = "high", .value = 75 },    .{ .name = "xhigh", .value = 88 }, .{ .name = "max", .value = 100 },
 };
 

From 1104203a46f1088a5bd4614a9bd618c1fbb9e9d1 Mon Sep 17 00:00:00 2001
From: Elevatormusic <22101396+Elevatormusic@users.noreply.github.com>
Date: Sat, 24 Oct 2026 11:53:09 -0700
Subject: [PATCH 4/5] dsv41: dense linear comments are one line (lean check)

---
 zig/src/families/dsv41/cublas.zig  |  4 +---
 zig/src/families/dsv41/plan.zig    |  3 +--
 zig/src/families/dsv41/prompt.zig  |  4 +---
 zig/src/families/dsv41/weights.zig | 14 +++++---------
 4 files changed, 8 insertions(+), 17 deletions(-)

diff --git a/zig/src/families/dsv41/cublas.zig b/zig/src/families/dsv41/cublas.zig
index 5019eb35a0..f5e47e31be 100644
--- a/zig/src/families/dsv41/cublas.zig
+++ b/zig/src/families/dsv41/cublas.zig
@@ -73,9 +73,7 @@ pub const Blas = struct {
         try b.check(b.api.cublasGemmEx(b.handle, .t, .n, @intCast(n), @intCast(rows), @intCast(k), &one, w, bf16_type, @intCast(k), x, bf16_type, @intCast(k), &zero, out, bf16_type, @intCast(n), 68, 99));
     }
 
-    /// A dense linear of a mixed pack (weights.Linear.dense): x bf16 [rows, k] (row stride ldx) @ w bf16 [n, k]^T into
-    /// out [rows, n] (row stride ldo), bf16 or fp32. linearBf16's call (COMPUTE_32F, fp32 alpha 1 and beta 0,
-    /// CUBLAS_GEMM_DEFAULT_TENSOR_OP) with the strides and the output dtype as the caller's buffers have them.
+    /// A dense linear (weights.Linear.dense): x bf16 [rows, k] @ w bf16 [n, k]^T into bf16 or fp32 out, as linearBf16.
     pub fn linearDense(b: *const Blas, x: u64, ldx: usize, w: u64, out: u64, ldo: usize, out_f32: bool, rows: usize, k: usize, n: usize) Error!void {
         const one: f32 = 1;
         const zero: f32 = 0;
diff --git a/zig/src/families/dsv41/plan.zig b/zig/src/families/dsv41/plan.zig
index 937af75f08..190393fe00 100644
--- a/zig/src/families/dsv41/plan.zig
+++ b/zig/src/families/dsv41/plan.zig
@@ -144,8 +144,7 @@ pub fn denseKey(buf: []u8, prefix: []const u8) ![]const u8 {
     return std.fmt.bufPrint(buf, "{s}.weight|None", .{prefix});
 }
 
-/// Which linears a mixed pack holds dense (a plain `<prefix>.weight` and no trellis): `has(ctx, prefix)`. A whole
-/// linear the predicate names is planned as its dense entry; a split one (2D / TP column or row parts) is refused.
+/// Which whole linears a mixed pack holds dense (`has(ctx, prefix)`); a split (2D / TP part) dense linear is refused.
 pub const Dense = struct {
     ctx: *const anyopaque,
     has: *const fn (ctx: *const anyopaque, prefix: []const u8) bool,
diff --git a/zig/src/families/dsv41/prompt.zig b/zig/src/families/dsv41/prompt.zig
index 38987bc0db..084edd0e8d 100644
--- a/zig/src/families/dsv41/prompt.zig
+++ b/zig/src/families/dsv41/prompt.zig
@@ -584,9 +584,7 @@ pub fn denseRows(e: *const Engine, m: usize, l: weights.Linear, x: u64, ldx: usi
     try e.blas.linearDense(x, ldx, l.dense, out, ldo, f32_out, m, l.k, l.n);
 }
 
-/// Exl3Group.rotated on m rows already rotated into each call's xh (a kernel that folds the rotation in wrote them):
-/// the glinear launches alone, each call's counters and folds as given. A dense linear has no rotated rows: its
-/// caller runs it from the normed rows (denseRows).
+/// Exl3Group.rotated on m rows already rotated into each call's xh: the glinear launches alone (dense: denseRows).
 pub fn groupedRotated(e: *const Engine, ch: *const Chunk, m: usize, calls: []exl3_linear.Call) !void {
     if (m == 0 or m > 128 or calls.len > exl3_linear.gmax) return error.NotADecodeWindow;
     for (calls) |c| if (c.layer.isDense()) return error.DenseNotRotated;
diff --git a/zig/src/families/dsv41/weights.zig b/zig/src/families/dsv41/weights.zig
index ebef6d7203..313492e92f 100644
--- a/zig/src/families/dsv41/weights.zig
+++ b/zig/src/families/dsv41/weights.zig
@@ -14,11 +14,7 @@ const max_span = 4 << 30;
 /// The size of one direct read within such a stretch.
 const span_piece = 64 << 20;
 
-/// An EXL3 linear in the strips layout: words int32 [N/128, K/16, 8, 8 * bits], fp16 suh [K] and svh [N].
-/// plan_n: the N its decode plan is made for, when not n: a 2D node's column part of a TP2 slice keeps that slice's
-/// plan (linear.py pinned plans), so its columns sum in the slice's K order.
-/// An EXL3 linear (words, suh, svh), or a dense one: `dense` bf16 [n, k] row-major (a mixed pack stores some linears
-/// as plain `<prefix>.weight` tensors at source precision, EAQ v1.1's indexer path), words / suh / svh / k2 then 0.
+/// An EXL3 linear in the strips layout: words int32 [N/128, K/16, 8, 8 * bits], fp16 suh [K] and svh [N]; or dense.
 pub const Linear = struct {
     words: u64,
     suh: u64,
@@ -26,7 +22,9 @@ pub const Linear = struct {
     k: u32,
     n: u32,
     k2: u32,
+    /// The N of its decode plan when not n: a 2D column part keeps its TP2 slice's plan (linear.py pinned plans).
     plan_n: u32 = 0,
+    /// A mixed pack's plain bf16 `<prefix>.weight` [n, k] row-major (EAQ v1.1); words, suh, svh and k2 are then 0.
     dense: u64 = 0,
 
     pub fn isDense(l: Linear) bool {
@@ -320,8 +318,7 @@ const Loader = struct {
         return L.upload(tensor, L.aux.items);
     }
 
-    /// linear(prefix, cols, rows): Exl3Linear.from_tensors in strips, or a whole linear the pack holds dense (see
-    /// denseLinear).
+    /// linear(prefix, cols, rows): Exl3Linear.from_tensors in strips, or a whole linear the pack holds dense.
     fn linear(L: *Loader, prefix: []const u8, cols: ?[2]usize, rows: ?[2]usize) !Linear {
         var kb: [256]u8 = undefined;
         var nb: [256]u8 = undefined;
@@ -340,8 +337,7 @@ const Loader = struct {
         return .{ .words = words, .suh = suh, .svh = svh, .k = @intCast(kt * 16), .n = @intCast(nt * 16), .k2 = @intCast(w16 / 8) };
     }
 
-    /// A whole linear the source holds as a plain bf16 `<prefix>.weight` [n, k] (plan.zig's dense entry) instead of
-    /// trellis parts: uploaded as it is. Null when the source has its trellis (every linear of an all-EXL3 pack).
+    /// A whole linear the source holds as a plain bf16 `<prefix>.weight` [n, k], uploaded as it is; null: a trellis.
     fn denseLinear(L: *Loader, prefix: []const u8) !?Linear {
         var kb: [256]u8 = undefined;
         var nb: [256]u8 = undefined;

From 52173eadd2900cdc2cf6d156a1ffa4ea82b21c8d Mon Sep 17 00:00:00 2001
From: Elevatormusic <22101396+Elevatormusic@users.noreply.github.com>
Date: Sat, 24 Oct 2026 12:17:49 -0700
Subject: [PATCH 5/5] dsv41: a dense wo_a slice is refused at plan and load
 time (no dense run path for it yet)

Round.init copies each wo_a slice's suh, and the wo_a decode runs
groupedRotated, which refuses dense layers. A mixed pack with a dense
wo_a slice thus planned and loaded, then failed at init. plan.denseRunnable
now refuses it in the planner (error.DenseLinearUnsupported) and the
loader checks it too. New test case in the dense plan test.
---
 zig/src/families/dsv41/plan.zig    | 11 +++++++++++
 zig/src/families/dsv41/weights.zig |  5 ++++-
 2 files changed, 15 insertions(+), 1 deletion(-)

diff --git a/zig/src/families/dsv41/plan.zig b/zig/src/families/dsv41/plan.zig
index 190393fe00..69741e570c 100644
--- a/zig/src/families/dsv41/plan.zig
+++ b/zig/src/families/dsv41/plan.zig
@@ -144,6 +144,11 @@ pub fn denseKey(buf: []u8, prefix: []const u8) ![]const u8 {
     return std.fmt.bufPrint(buf, "{s}.weight|None", .{prefix});
 }
 
+/// Whether a dense linear has a run path: not wo_a's slices (Round.init reads their suh; decode runs groupedRotated).
+pub fn denseRunnable(prefix: []const u8) bool {
+    return std.mem.indexOf(u8, prefix, ".attn.wo_a.") == null;
+}
+
 /// Which whole linears a mixed pack holds dense (`has(ctx, prefix)`); a split (2D / TP part) dense linear is refused.
 pub const Dense = struct {
     ctx: *const anyopaque,
@@ -175,6 +180,7 @@ const Builder = struct {
     fn exl3(b: *Builder, prefix: []const u8, cols: ?[2]usize, rows: ?[2]usize, k: ?usize, n: ?usize) !void {
         if (b.dense) |d| if (d.has(d.ctx, prefix)) {
             if (cols != null or rows != null) return error.DenseLinearSplit;
+            if (!denseRunnable(prefix)) return error.DenseLinearUnsupported;
             var buf: [256]u8 = undefined;
             return b.out.append(b.a, .{ .key = try b.a.dupe(u8, try denseKey(&buf, prefix)), .dtype = .bf16 });
         };
@@ -314,6 +320,9 @@ test "a mixed pack's dense linears: one bf16 .weight entry for a whole linear, a
         fn split(_: *const anyopaque, prefix: []const u8) bool {
             return std.mem.eql(u8, prefix, "layers.0.attn.wq_b");
         }
+        fn woA(_: *const anyopaque, prefix: []const u8) bool {
+            return std.mem.eql(u8, prefix, "layers.0.attn.wo_a.slice.0");
+        }
     };
     const ctx: u8 = 0;
     const w = try wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &ctx, .has = names.has });
@@ -331,6 +340,8 @@ test "a mixed pack's dense linears: one bf16 .weight entry for a whole linear, a
     try std.testing.expectEqual(@as(usize, 2), seen);
     // wq_b is a column slice of the rank's heads: no dense form of it
     try std.testing.expectError(error.DenseLinearSplit, wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &ctx, .has = names.split }));
+    // a whole wo_a slice has no dense run path yet: refused at plan time, not at Round.init
+    try std.testing.expectError(error.DenseLinearUnsupported, wantsDense(a, c, .{ .rank = 0, .world = 2 }, true, .{ .ctx = &ctx, .has = names.woA }));
 }
 
 test "each split tensor's ranges: TP2's slices, a 2D node's pair parts of them" {
diff --git a/zig/src/families/dsv41/weights.zig b/zig/src/families/dsv41/weights.zig
index 313492e92f..c25d23c84a 100644
--- a/zig/src/families/dsv41/weights.zig
+++ b/zig/src/families/dsv41/weights.zig
@@ -322,7 +322,10 @@ const Loader = struct {
     fn linear(L: *Loader, prefix: []const u8, cols: ?[2]usize, rows: ?[2]usize) !Linear {
         var kb: [256]u8 = undefined;
         var nb: [256]u8 = undefined;
-        if (cols == null and rows == null) if (try L.denseLinear(prefix)) |d| return d;
+        if (cols == null and rows == null) if (try L.denseLinear(prefix)) |d| {
+            if (!plan.denseRunnable(prefix)) return error.DenseLinearUnsupported;
+            return d;
+        };
         const tr = try L.entry(try plan.partKey(&kb, prefix, cols, rows, "tr"));
         const t = try L.shape(try plan.partKey(&kb, prefix, cols, rows, "tr"));
         const kt = t[0];