From 422d5ce6699ef9624ca80bde9decb66a19980c55 Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Sun, 18 Oct 2026 18:16:31 +0700 Subject: [PATCH 01/11] igla ptx: spec that writes the PTX of the ternary x int8 mma.sync GEMM kernel, with its host reference (Refs #10404) Co-Authored-By: Claude Opus 5.5 --- specs/igla/coder/ternary_gemm_ptx.t27 | 758 ++++++++++++++++++++++++++ 1 file changed, 758 insertions(+) create mode 100644 specs/igla/coder/ternary_gemm_ptx.t27 diff --git a/specs/igla/coder/ternary_gemm_ptx.t27 b/specs/igla/coder/ternary_gemm_ptx.t27 new file mode 100644 index 0000000000..4689974176 --- /dev/null +++ b/specs/igla/coder/ternary_gemm_ptx.t27 @@ -0,0 +1,758 @@ +// SPDX-License-Identifier: Apache-2.0 +; specs/igla/coder/ternary_gemm_ptx.t27 -- the PTX text of the ternary x int8 GEMM kernel, written by t27. +; Issue #10404, plan #10360. Owner override 2026-10-18: build the backend, and build it in t27. +; -- +; What this is. specs/igla/coder/ternary_gemm_gpu.t27 pins the math of a tiled ternary x int8 GEMM and +; checks it bit-exactly on two compilers. This spec is its GPU lowering: emit_ptx() writes the PTX of +; one kernel, `ternary_gemm_s8`, for int8 tensor cores (mma.sync m16n8k32 s8 x s8 -> s32, sm_80 and +; later; the RTX 4090 is sm_89). `t27c compile specs/igla/coder/ternary_gemm_gpu.t27 --backend ptx` +; writes this text; `t27c gen-rust` lowers this spec into bootstrap/gen/rust/igla/coder/ and Rust +; only copies the bytes to a file or to the CUDA driver. +; The honest limit: this is a kernel-specific emitter, not a general t27 -> PTX lowering. It writes +; exactly one kernel, and t27c refuses `--backend ptx` for every other spec. +; -- +; Why int8 tensor cores and not dp4a. On the RTX 4090 bench of #10360 (CUTLASS 3.5.1, one pod), +; int8 tensor cores ran 2.9-3.5x bf16 and dp4a ran 0.91-0.94x bf16; only the mma path can win. +; -- +; The kernel, v1. +; C[m][n] = sum_k A[m][k] * w(k, n); A is i8 row-major (m x k); W is the k x n code matrix of +; ternary_gemm_gpu.t27, row-major, four 2-bit codes per byte, code i at byte i/4, lane i%4, lane 0 +; in the low bits; C is i32 row-major (m x n). +; Grid (ceil(n/128), ceil(m/128)), 256 threads. Each block computes a 128 x 128 tile of C in k-steps +; of 32. Per k-step: +; - A tile: 128 rows x 32 bytes; thread t loads 16 bytes, row t/2, bytes (t%2)*16, zero past m +; and past k, and stores them in sA at row * 48 + chunk; +; - W tile: 32 k-rows x 128 columns of codes = 256 words; thread t loads one u32 (16 codes) at +; k-row t/8, columns (t%8)*16, decodes each code once as w = (c & 1) - ((c >> 1) & 1), and +; stores it TRANSPOSED as an s8 in sB at column * 48 + k-row ("col" operand of the mma); +; - warp v = t/32 owns the 64 x 32 sub-tile at rows (v/4)*64, columns (v%4)*32: 4 x 4 mma tiles +; of 16 x 8, so 16 mma.sync per k-step and 64 i32 accumulators per thread. +; The 48-byte row stride (32 data bytes + 16 pad) puts the 32 lanes of every fragment load on 32 +; distinct banks: lane (g, t) reads word 12g + t, and 12g mod 32 for g = 0..7 is 0,12,24,4,16,28,8,20. +; Fragment layouts are the PTX ISA's for m16n8k32 .s8 (g = lane/4, t = lane%4): +; A: a0 = row g, k 4t..4t+3; a1 = row g+8; a2 = row g, k 16+4t..; a3 = row g+8, k 16+4t.. +; B: b0 = column g, k 4t..4t+3; b1 = column g, k 16+4t.. +; C: c0, c1 = row g, columns 2t, 2t+1; c2, c3 = row g+8, same columns. +; v1 has no cp.async and no double buffering: it is the first kernel that runs, not the fastest. +; Preconditions, checked by shape_ok() before a launch: n and k are multiples of 16 (aligned 16-byte +; A loads and 4-byte W loads; a k chunk or a 16-column W word is then either wholly inside or wholly +; outside), 1 <= m, and k <= K_MAX_I32 so that no i32 partial sum can overflow. +; -- +; The host side lives here too, so that the bit-exact check is t27 as well: gen_a, gen_wp (the data), +; ref_row (the CPU reference, one output row at a time) and mismatches (the count the run reports). +; Bit-exact is valid for the reason ternary_gemm_gpu.t27 states: every product is in [-128, 128], the +; sum is i32 and cannot overflow while k <= K_MAX_I32, so integer addition is associative and any +; order of the k-sum gives the same bits. +; ASCII only (L3). phi^2 + 1/phi^2 = 3 | TRINITY + +module igla-coder-ternary-gemm-ptx; + +/// The kernel's entry name, the one cuModuleGetFunction looks up. +pub const KERNEL_NAME : str = "ternary_gemm_s8"; +/// Block tile and k-step of the kernel, and its thread count. +pub const BM : u32 = 128; +pub const BN : u32 = 128; +pub const BK : u32 = 32; +pub const THREADS : u32 = 256; +/// Bytes per row of a shared tile: 32 data bytes and 16 bytes of padding. +pub const SMEM_STRIDE : u32 = 48; +/// mma.sync instructions per k-step: 4 x 4 tiles of 16 x 8 per warp. +pub const MMA_PER_STEP : u32 = 16; +/// Largest k for which no i32 partial sum can overflow: 128 * 16777215 = 2147483520. +pub const K_MAX_I32 : usize = 16777215; +/// Largest grid y the driver accepts. +pub const GRID_Y_MAX : usize = 65535; +/// Buffer the emitted text needs, with room to spare; emit_ptx returns the exact length. +pub const PTX_CAP : usize = 32768; +/// tern-tc-500m: d_model and d_ff (read from the trainer config, see #10360). +pub const D_MODEL_500M : usize = 1280; +pub const D_FF_500M : usize = 4096; + +// ============================================================================ +// 1. A bounded text writer +// ============================================================================ + +/// Write s at out[w..]; bytes past cap are counted but not written. Returns the new end. +fn put_s(out : []u8, cap : usize, w : usize, s : str) -> usize { + var i : usize = 0; + while (i < s.len) { + if (w + i < cap) { out[w + i] = s[i]; } + i += 1; + } + return w + s.len; +} + +/// Write v in decimal at out[w..], bounded by cap like put_s. +fn put_u(out : []u8, cap : usize, w : usize, v : u32) -> usize { + var d : u32 = 1; + while (v / d >= 10) { d = d * 10; } + var x : usize = w; + while (d > 0) { + const c : u8 = '0' + (((v / d) % 10) as u8); + if (x < cap) { out[x] = c; } + x += 1; + d = d / 10; + } + return x; +} + +/// One indented instruction line that has no operand to fill in. +fn ins(out : []u8, cap : usize, w : usize, s : str) -> usize { + const x : usize = put_s(out, cap, w, " "); + return put_s(out, cap, put_s(out, cap, x, s), "\n"); +} + +/// "{%name, %name, %name, %name}" +fn put_quad(out : []u8, cap : usize, w : usize, name : str, first : u32) -> usize { + var x : usize = put_s(out, cap, w, "{"); + var j : u32 = 0; + while (j < 4) { + if (j > 0) { x = put_s(out, cap, x, ", "); } + x = put_u(out, cap, put_s(out, cap, x, name), first + j); + j += 1; + } + return put_s(out, cap, x, "}"); +} + +// ============================================================================ +// 2. The kernel, section by section +// ============================================================================ + +fn emit_header(out : []u8, cap : usize, w : usize) -> usize { + var x : usize = put_s(out, cap, w, "// ternary_gemm_s8 -- emitted by t27c from specs/igla/coder/ternary_gemm_ptx.t27 (#10404). Do not edit.\n"); + x = put_s(out, cap, x, ".version 7.0\n.target sm_80\n.address_size 64\n\n"); + x = put_s(out, cap, x, ".visible .entry ternary_gemm_s8(\n"); + x = put_s(out, cap, x, " .param .u64 p_a,\n .param .u64 p_w,\n .param .u64 p_c,\n"); + x = put_s(out, cap, x, " .param .u32 p_m,\n .param .u32 p_n,\n .param .u32 p_k\n)\n"); + x = put_s(out, cap, x, ".maxntid 256, 1, 1\n{\n"); + x = ins(out, cap, x, ".reg .pred %p<8>;"); + x = ins(out, cap, x, ".reg .b32 %r<48>;"); + x = ins(out, cap, x, ".reg .b32 %acc<64>;"); + x = ins(out, cap, x, ".reg .b32 %fa<16>;"); + x = ins(out, cap, x, ".reg .b32 %fb<8>;"); + x = ins(out, cap, x, ".reg .b64 %rd<12>;"); + x = ins(out, cap, x, ".shared .align 16 .b8 sA[6144];"); + x = ins(out, cap, x, ".shared .align 16 .b8 sB[6144];"); + return x; +} + +/// Parameters, thread coordinates, and the addresses each thread reuses every k-step. +fn emit_prologue(out : []u8, cap : usize, w : usize) -> usize { + var x : usize = ins(out, cap, w, "ld.param.u64 %rd1, [p_a];"); + x = ins(out, cap, x, "ld.param.u64 %rd2, [p_w];"); + x = ins(out, cap, x, "ld.param.u64 %rd3, [p_c];"); + x = ins(out, cap, x, "ld.param.u32 %r1, [p_m];"); + x = ins(out, cap, x, "ld.param.u32 %r2, [p_n];"); + x = ins(out, cap, x, "ld.param.u32 %r3, [p_k];"); + x = ins(out, cap, x, "cvta.to.global.u64 %rd1, %rd1;"); + x = ins(out, cap, x, "cvta.to.global.u64 %rd2, %rd2;"); + x = ins(out, cap, x, "cvta.to.global.u64 %rd3, %rd3;"); + x = ins(out, cap, x, "mov.u32 %r4, %tid.x;"); + x = ins(out, cap, x, "mov.u32 %r5, %ctaid.x;"); + x = ins(out, cap, x, "mov.u32 %r6, %ctaid.y;"); + ; A tile: row r7 = tid/2, chunk r8 = (tid%2)*16, global row r9, p1 = row inside m. + x = ins(out, cap, x, "shr.u32 %r7, %r4, 1;"); + x = ins(out, cap, x, "and.b32 %r8, %r4, 1;"); + x = ins(out, cap, x, "shl.b32 %r8, %r8, 4;"); + x = ins(out, cap, x, "shl.b32 %r9, %r6, 7;"); + x = ins(out, cap, x, "add.u32 %r9, %r9, %r7;"); + x = ins(out, cap, x, "setp.lt.u32 %p1, %r9, %r1;"); + x = ins(out, cap, x, "mov.u32 %r10, sA;"); + x = ins(out, cap, x, "mad.lo.u32 %r11, %r7, 48, %r8;"); + x = ins(out, cap, x, "add.u32 %r11, %r11, %r10;"); + x = ins(out, cap, x, "mul.wide.u32 %rd4, %r9, %r3;"); + x = ins(out, cap, x, "add.u64 %rd4, %rd4, %rd1;"); + x = ins(out, cap, x, "cvt.u64.u32 %rd5, %r8;"); + x = ins(out, cap, x, "add.u64 %rd4, %rd4, %rd5;"); + ; W tile: k-row r12 = tid/8, column group r13 = (tid%8)*16, global column r14, p2 = inside n. + x = ins(out, cap, x, "shr.u32 %r12, %r4, 3;"); + x = ins(out, cap, x, "and.b32 %r13, %r4, 7;"); + x = ins(out, cap, x, "shl.b32 %r13, %r13, 4;"); + x = ins(out, cap, x, "shl.b32 %r14, %r5, 7;"); + x = ins(out, cap, x, "add.u32 %r14, %r14, %r13;"); + x = ins(out, cap, x, "setp.lt.u32 %p2, %r14, %r2;"); + x = ins(out, cap, x, "mov.u32 %r15, sB;"); + x = ins(out, cap, x, "mad.lo.u32 %r16, %r13, 48, %r12;"); + x = ins(out, cap, x, "add.u32 %r16, %r16, %r15;"); + x = ins(out, cap, x, "shr.u32 %r17, %r2, 2;"); + x = ins(out, cap, x, "shr.u32 %r18, %r14, 2;"); + ; Warp r19, lane r20, g r21, t r22, warp row r23, warp column r24. + x = ins(out, cap, x, "shr.u32 %r19, %r4, 5;"); + x = ins(out, cap, x, "and.b32 %r20, %r4, 31;"); + x = ins(out, cap, x, "shr.u32 %r21, %r20, 2;"); + x = ins(out, cap, x, "and.b32 %r22, %r20, 3;"); + x = ins(out, cap, x, "shr.u32 %r23, %r19, 2;"); + x = ins(out, cap, x, "and.b32 %r24, %r19, 3;"); + ; Fragment bases: r25 = sA + (64 wm + g) * 48 + 4t, r27 = sB + (32 wn + g) * 48 + 4t. + x = ins(out, cap, x, "shl.b32 %r26, %r22, 2;"); + x = ins(out, cap, x, "shl.b32 %r25, %r23, 6;"); + x = ins(out, cap, x, "add.u32 %r25, %r25, %r21;"); + x = ins(out, cap, x, "mul.lo.u32 %r25, %r25, 48;"); + x = ins(out, cap, x, "add.u32 %r25, %r25, %r26;"); + x = ins(out, cap, x, "add.u32 %r25, %r25, %r10;"); + x = ins(out, cap, x, "shl.b32 %r27, %r24, 5;"); + x = ins(out, cap, x, "add.u32 %r27, %r27, %r21;"); + x = ins(out, cap, x, "mul.lo.u32 %r27, %r27, 48;"); + x = ins(out, cap, x, "add.u32 %r27, %r27, %r26;"); + x = ins(out, cap, x, "add.u32 %r27, %r27, %r15;"); + x = ins(out, cap, x, "mov.u32 %r28, 0;"); + var i : u32 = 0; + while (i < 64) { + x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " mov.b32 %acc"), i), ", 0;\n"); + i += 1; + } + return x; +} + +/// The cooperative tile loads of one k-step (k0 in r28), A first, then W decoded into sB. +fn emit_tile_loads(out : []u8, cap : usize, w : usize) -> usize { + var x : usize = ins(out, cap, w, "add.u32 %r29, %r28, %r8;"); + x = ins(out, cap, x, "setp.lt.u32 %p4, %r29, %r3;"); + x = ins(out, cap, x, "and.pred %p4, %p4, %p1;"); + x = ins(out, cap, x, "mov.b32 %r30, 0;"); + x = ins(out, cap, x, "mov.b32 %r31, 0;"); + x = ins(out, cap, x, "mov.b32 %r32, 0;"); + x = ins(out, cap, x, "mov.b32 %r33, 0;"); + x = ins(out, cap, x, "cvt.u64.u32 %rd5, %r28;"); + x = ins(out, cap, x, "add.u64 %rd6, %rd4, %rd5;"); + x = ins(out, cap, x, "@%p4 ld.global.v4.u32 {%r30, %r31, %r32, %r33}, [%rd6];"); + x = ins(out, cap, x, "st.shared.v4.u32 [%r11], {%r30, %r31, %r32, %r33};"); + x = ins(out, cap, x, "add.u32 %r34, %r28, %r12;"); + x = ins(out, cap, x, "setp.lt.u32 %p5, %r34, %r3;"); + x = ins(out, cap, x, "and.pred %p5, %p5, %p2;"); + x = ins(out, cap, x, "mov.b32 %r35, 0;"); + x = ins(out, cap, x, "mul.wide.u32 %rd7, %r34, %r17;"); + x = ins(out, cap, x, "cvt.u64.u32 %rd8, %r18;"); + x = ins(out, cap, x, "add.u64 %rd7, %rd7, %rd8;"); + x = ins(out, cap, x, "add.u64 %rd9, %rd7, %rd2;"); + x = ins(out, cap, x, "@%p5 ld.global.u32 %r35, [%rd9];"); + var j : u32 = 0; + while (j < 16) { + x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " shr.u32 %r36, %r35, "), 2 * j), ";\n"); + x = ins(out, cap, x, "and.b32 %r37, %r36, 1;"); + x = ins(out, cap, x, "shr.u32 %r38, %r36, 1;"); + x = ins(out, cap, x, "and.b32 %r38, %r38, 1;"); + x = ins(out, cap, x, "sub.s32 %r37, %r37, %r38;"); + x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " st.shared.u8 [%r16+"), j * SMEM_STRIDE), "], %r37;\n"); + j += 1; + } + return x; +} + +/// " ld.shared.u32 %, [%+];" +fn emit_lds(out : []u8, cap : usize, w : usize, reg : str, i : u32, base : str, off : u32) -> usize { + var x : usize = put_u(out, cap, put_s(out, cap, put_s(out, cap, w, " ld.shared.u32 "), reg), i); + x = put_s(out, cap, put_s(out, cap, put_s(out, cap, x, ", ["), base), "+"); + return put_s(out, cap, put_u(out, cap, x, off), "];\n"); +} + +/// The fragment loads and the 16 mma.sync of one k-step. +fn emit_mma_step(out : []u8, cap : usize, w : usize) -> usize { + var x : usize = w; + var mi : u32 = 0; + while (mi < 4) { + const off : u32 = mi * 16 * SMEM_STRIDE; + x = emit_lds(out, cap, x, "%fa", mi * 4, "%r25", off); + x = emit_lds(out, cap, x, "%fa", mi * 4 + 1, "%r25", off + 8 * SMEM_STRIDE); + x = emit_lds(out, cap, x, "%fa", mi * 4 + 2, "%r25", off + 16); + x = emit_lds(out, cap, x, "%fa", mi * 4 + 3, "%r25", off + 8 * SMEM_STRIDE + 16); + mi += 1; + } + var ni : u32 = 0; + while (ni < 4) { + x = emit_lds(out, cap, x, "%fb", ni * 2, "%r27", ni * 8 * SMEM_STRIDE); + x = emit_lds(out, cap, x, "%fb", ni * 2 + 1, "%r27", ni * 8 * SMEM_STRIDE + 16); + ni += 1; + } + mi = 0; + while (mi < 4) { + ni = 0; + while (ni < 4) { + const b : u32 = (mi * 4 + ni) * 4; + x = put_s(out, cap, x, " mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 "); + x = put_s(out, cap, put_quad(out, cap, x, "%acc", b), ", "); + x = put_s(out, cap, put_quad(out, cap, x, "%fa", mi * 4), ", {%fb"); + x = put_u(out, cap, put_s(out, cap, put_u(out, cap, x, ni * 2), ", %fb"), ni * 2 + 1); + x = put_s(out, cap, put_quad(out, cap, put_s(out, cap, x, "}, "), "%acc", b), ";\n"); + ni += 1; + } + mi += 1; + } + return x; +} + +/// One guarded pair store: accumulators b, b+1 to C[row0 + dr][col0 + dc .. +1]. +fn emit_store(out : []u8, cap : usize, w : usize, dr : u32, dc : u32, b : u32) -> usize { + var x : usize = put_s(out, cap, put_u(out, cap, put_s(out, cap, w, " add.u32 %r41, %r39, "), dr), ";\n"); + x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " add.u32 %r42, %r40, "), dc), ";\n"); + x = ins(out, cap, x, "setp.lt.u32 %p6, %r41, %r1;"); + x = ins(out, cap, x, "setp.lt.u32 %p7, %r42, %r2;"); + x = ins(out, cap, x, "and.pred %p6, %p6, %p7;"); + x = ins(out, cap, x, "cvt.u64.u32 %rd10, %r42;"); + x = ins(out, cap, x, "mad.wide.u32 %rd10, %r41, %r2, %rd10;"); + x = ins(out, cap, x, "shl.b64 %rd10, %rd10, 2;"); + x = ins(out, cap, x, "add.u64 %rd10, %rd10, %rd3;"); + x = put_u(out, cap, put_s(out, cap, x, " @%p6 st.global.v2.u32 [%rd10], {%acc"), b); + return put_s(out, cap, put_u(out, cap, put_s(out, cap, x, ", %acc"), b + 1), "};\n"); +} + +/// Row base r39 = 128 by + 64 wm + g and column base r40 = 128 bx + 32 wn + 2t, then 32 stores. +fn emit_epilogue(out : []u8, cap : usize, w : usize) -> usize { + var x : usize = ins(out, cap, w, "shl.b32 %r39, %r6, 7;"); + x = ins(out, cap, x, "shl.b32 %r41, %r23, 6;"); + x = ins(out, cap, x, "add.u32 %r39, %r39, %r41;"); + x = ins(out, cap, x, "add.u32 %r39, %r39, %r21;"); + x = ins(out, cap, x, "shl.b32 %r40, %r5, 7;"); + x = ins(out, cap, x, "shl.b32 %r41, %r24, 5;"); + x = ins(out, cap, x, "add.u32 %r40, %r40, %r41;"); + x = ins(out, cap, x, "shl.b32 %r41, %r22, 1;"); + x = ins(out, cap, x, "add.u32 %r40, %r40, %r41;"); + var mi : u32 = 0; + while (mi < 4) { + var ni : u32 = 0; + while (ni < 4) { + const b : u32 = (mi * 4 + ni) * 4; + x = emit_store(out, cap, x, mi * 16, ni * 8, b); + x = emit_store(out, cap, x, mi * 16 + 8, ni * 8, b + 2); + ni += 1; + } + mi += 1; + } + return x; +} + +/// The whole kernel into out[0..cap). Returns the bytes it needs; when that exceeds cap the buffer +/// holds a clipped text the caller must not use. +pub fn emit_ptx(out : []u8, cap : usize) -> usize { + var x : usize = emit_header(out, cap, 0); + x = emit_prologue(out, cap, x); + x = put_s(out, cap, x, "LOOP_K:\n"); + x = ins(out, cap, x, "setp.ge.u32 %p3, %r28, %r3;"); + x = ins(out, cap, x, "@%p3 bra DONE_K;"); + x = emit_tile_loads(out, cap, x); + x = ins(out, cap, x, "bar.sync 0;"); + x = emit_mma_step(out, cap, x); + x = ins(out, cap, x, "bar.sync 0;"); + x = ins(out, cap, x, "add.u32 %r28, %r28, 32;"); + x = ins(out, cap, x, "bra LOOP_K;"); + x = put_s(out, cap, x, "DONE_K:\n"); + x = emit_epilogue(out, cap, x); + x = ins(out, cap, x, "ret;"); + return put_s(out, cap, x, "}\n"); +} + +// ============================================================================ +// 3. Launch geometry and the shapes the kernel accepts +// ============================================================================ + +/// Blocks along a dimension of x elements: ceil(x / 128). +pub fn grid_dim(x : usize) -> u32 { + return ((x + 127) / 128) as u32; +} + +/// Whether the kernel can run m x k times k x n. n and k multiples of 16; k small enough for i32. +pub fn shape_ok(m : usize, n : usize, k : usize) -> bool { + if (m < 1 or n < 16 or k < 16) { return false; } + if (n % 16 != 0 or k % 16 != 0) { return false; } + if (k > K_MAX_I32) { return false; } + return (m + 127) / 128 <= GRID_Y_MAX; +} + +// ============================================================================ +// 4. Host data, the CPU reference and the mismatch count +// ============================================================================ + +/// 31-bit LCG (glibc constants); its period is 2^31, so no row of a 500m shape repeats another. +pub fn lcg(s : u64) -> u64 { + return (s * 1103515245 + 12345) % 2147483648; +} + +/// count activations in [-128, 127] from bits 16..23 of the LCG; returns the last state. +pub fn gen_a(a : []i8, count : usize, seed : u64) -> u64 { + var s : u64 = seed; + var i : usize = 0; + while (i < count) { + s = lcg(s); + a[i] = ((((s >> 16) % 256) as i32) - 128) as i8; + i += 1; + } + return s; +} + +/// count 2-bit codes (all four values, so code 3 = 0 is exercised too), packed four per byte with +/// lane 0 in the low bits. count is a multiple of 4. Returns the last state. +pub fn gen_wp(wp : []u8, count : usize, seed : u64) -> u64 { + var s : u64 = seed; + var i : usize = 0; + while (i < count) { + var byte : u32 = 0; + var lane : u32 = 0; + while (lane < 4) { + s = lcg(s); + const code : u32 = ((s >> 16) % 4) as u32; + byte = byte | (code << (lane * 2)); + lane += 1; + } + wp[i / 4] = byte as u8; + i += 4; + } + return s; +} + +/// Code i of the packed matrix. +pub fn code_at(wp : []const u8, i : usize) -> u32 { + const b : u32 = wp[i / 4] as u32; + const sh : u32 = ((i % 4) * 2) as u32; + return (b >> sh) & 3; +} + +/// The decode the kernel runs: (c & 1) - ((c >> 1) & 1); 0 -> 0, 1 -> +1, 2 -> -1, 3 -> 0. +pub fn weight(code : u32) -> i32 { + return ((code & 1) as i32) - (((code >> 1) & 1) as i32); +} + +/// acc[0..n) = row `row` of A x W, k-outer so the codes are read in memory order. +pub fn ref_row(a : []const i8, wp : []const u8, n : usize, k : usize, row : usize, acc : []i32) -> void { + var j : usize = 0; + while (j < n) { + acc[j] = 0; + j += 1; + } + var kk : usize = 0; + while (kk < k) { + const v : i32 = a[row * k + kk] as i32; + if (v != 0) { + const base : usize = kk * n; + var col : usize = 0; + while (col < n) { + acc[col] += v * weight(code_at(wp, base + col)); + col += 1; + } + } + kk += 1; + } +} + +/// Entries of c (m x n) that differ from the reference, over rows 0, stride, 2 stride, ... and always +/// the last row; stride 1 checks every entry. acc is n entries of scratch. +pub fn mismatches(a : []const i8, wp : []const u8, c : []const i32, m : usize, n : usize, k : usize, stride : usize, acc : []i32) -> usize { + var st : usize = stride; + if (st == 0) { st = 1; } + var bad : usize = 0; + var row : usize = 0; + while (row < m) { + ref_row(a, wp, n, k, row, acc); + var col : usize = 0; + while (col < n) { + if (c[row * n + col] != acc[col]) { bad += 1; } + col += 1; + } + if (row + 1 < m and row + st >= m) { + row = m - 1; + } else { + row += st; + } + } + return bad; +} + +// ============================================================================ +// 5. Test helpers: search the emitted text +// ============================================================================ + +/// Whether t[at..] starts with s. +fn at_text(t : []const u8, n : usize, at : usize, s : str) -> bool { + if (at + s.len > n) { return false; } + var i : usize = 0; + while (i < s.len) { + if (t[at + i] != s[i]) { return false; } + i += 1; + } + return true; +} + +/// How many times s occurs in t[0..n). +fn count_text(t : []const u8, n : usize, s : str) -> usize { + var c : usize = 0; + var i : usize = 0; + while (i < n) { + if (at_text(t, n, i, s)) { c += 1; } + i += 1; + } + return c; +} + +// ============================================================================ +// 6. Tests +// ============================================================================ + +test the_text_opens_with_the_ptx_header { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert at_text(&out, m, 0, "// ternary_gemm_s8 -- emitted by t27c from specs/igla/coder/ternary_gemm_ptx.t27"); + assert count_text(&out, m, "\n.version 7.0\n.target sm_80\n.address_size 64\n\n") == 1; +} + +test the_entry_takes_three_pointers_and_three_sizes { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert count_text(&out, m, ".visible .entry ternary_gemm_s8(\n .param .u64 p_a,\n .param .u64 p_w,\n .param .u64 p_c,\n .param .u32 p_m,\n .param .u32 p_n,\n .param .u32 p_k\n)\n.maxntid 256, 1, 1\n{\n") == 1; + assert count_text(&out, m, " .shared .align 16 .b8 sA[6144];\n") == 1; + assert count_text(&out, m, " .shared .align 16 .b8 sB[6144];\n") == 1; +} + +test sixteen_int8_mma_per_k_step { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert count_text(&out, m, "mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 ") == 16; + assert count_text(&out, m, " mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 {%acc0, %acc1, %acc2, %acc3}, {%fa0, %fa1, %fa2, %fa3}, {%fb0, %fb1}, {%acc0, %acc1, %acc2, %acc3};\n") == 1; + assert count_text(&out, m, " mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 {%acc20, %acc21, %acc22, %acc23}, {%fa4, %fa5, %fa6, %fa7}, {%fb2, %fb3}, {%acc20, %acc21, %acc22, %acc23};\n") == 1; + assert count_text(&out, m, " mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 {%acc60, %acc61, %acc62, %acc63}, {%fa12, %fa13, %fa14, %fa15}, {%fb6, %fb7}, {%acc60, %acc61, %acc62, %acc63};\n") == 1; + assert count_text(&out, m, "dp4a") == 0; +} + +test every_accumulator_starts_at_zero_once { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert count_text(&out, m, " mov.b32 %acc") == 64; + assert count_text(&out, m, " mov.b32 %acc0, 0;\n") == 1; + assert count_text(&out, m, " mov.b32 %acc63, 0;\n") == 1; +} + +test fragment_offsets_follow_the_48_byte_stride { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert count_text(&out, m, " ld.shared.u32 ") == 24; + ; A: tile mi at 16 rows x 48 = 768 bytes, the g+8 rows 384 further, the upper k half 16 further. + assert count_text(&out, m, " ld.shared.u32 %fa0, [%r25+0];\n") == 1; + assert count_text(&out, m, " ld.shared.u32 %fa1, [%r25+384];\n") == 1; + assert count_text(&out, m, " ld.shared.u32 %fa2, [%r25+16];\n") == 1; + assert count_text(&out, m, " ld.shared.u32 %fa3, [%r25+400];\n") == 1; + assert count_text(&out, m, " ld.shared.u32 %fa13, [%r25+2688];\n") == 1; + assert count_text(&out, m, " ld.shared.u32 %fa15, [%r25+2704];\n") == 1; + ; B: tile ni at 8 columns x 48 = 384 bytes. + assert count_text(&out, m, " ld.shared.u32 %fb6, [%r27+1152];\n") == 1; + assert count_text(&out, m, " ld.shared.u32 %fb7, [%r27+1168];\n") == 1; +} + +test sixteen_codes_decoded_per_thread_per_k_step { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert count_text(&out, m, " st.shared.u8 ") == 16; + assert count_text(&out, m, " shr.u32 %r36, %r35, 0;\n") == 1; + assert count_text(&out, m, " shr.u32 %r36, %r35, 30;\n") == 1; + assert count_text(&out, m, " st.shared.u8 [%r16+0], %r37;\n") == 1; + assert count_text(&out, m, " st.shared.u8 [%r16+720], %r37;\n") == 1; + assert count_text(&out, m, " sub.s32 %r37, %r37, %r38;\n") == 16; +} + +test thirty_two_guarded_pair_stores { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert count_text(&out, m, "st.global.v2.u32 ") == 32; + assert count_text(&out, m, " @%p6 st.global.v2.u32 [%rd10], {%acc0, %acc1};\n") == 1; + assert count_text(&out, m, " @%p6 st.global.v2.u32 [%rd10], {%acc62, %acc63};\n") == 1; + assert count_text(&out, m, " add.u32 %r41, %r39, 56;\n") == 1; + assert count_text(&out, m, " add.u32 %r42, %r40, 24;\n") == 8; +} + +test one_loop_two_barriers_one_return { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert count_text(&out, m, "\nLOOP_K:\n") == 1; + assert count_text(&out, m, "\nDONE_K:\n") == 1; + assert count_text(&out, m, " bar.sync 0;\n") == 2; + assert count_text(&out, m, " add.u32 %r28, %r28, 32;\n bra LOOP_K;\n") == 1; + assert at_text(&out, m, m - 11, " ret;\n}\n"); +} + +test the_text_is_ascii_and_fits_its_buffer { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + assert m > 16000; + assert m < PTX_CAP; + var i : usize = 0; + while (i < m) { + assert out[i] == '\n' or (out[i] >= 32 and out[i] < 127); + i += 1; + } +} + +test a_short_buffer_reports_the_full_length_and_writes_nothing_past_cap { + var out : [32768]u8 = undefined; + const m : usize = emit_ptx(&out, PTX_CAP); + var small : [64]u8 = [7; 64]; + const s : usize = emit_ptx(&small, 32); + assert s == m; + assert small[31] == out[31]; + assert small[32] == 7; + assert small[63] == 7; +} + +test emitting_twice_gives_the_same_bytes { + var a : [32768]u8 = undefined; + var b : [32768]u8 = undefined; + const m : usize = emit_ptx(&a, PTX_CAP); + const k : usize = emit_ptx(&b, PTX_CAP); + assert k == m; + var i : usize = 0; + while (i < m) { + assert a[i] == b[i]; + i += 1; + } +} + +test the_decode_is_the_gpu_spec_decode { + ; ternary_gemm_gpu.t27 decode: 0 -> 0, 1 -> +1, 2 -> -1, 3 -> 0. + assert weight(0) == 0; + assert weight(1) == 1; + assert weight(2) == -1; + assert weight(3) == 0; +} + +test packing_puts_lane_zero_in_the_low_bits { + var wp : [4]u8 = [0; 4]; + wp[0] = 0x1B; + wp[3] = 0x80; + ; 0x1B = 00 01 10 11: lanes 0..3 hold 3, 2, 1, 0. + assert code_at(&wp, 0) == 3; + assert code_at(&wp, 1) == 2; + assert code_at(&wp, 2) == 1; + assert code_at(&wp, 3) == 0; + assert code_at(&wp, 15) == 2; + assert code_at(&wp, 12) == 0; +} + +test generated_codes_round_trip_through_the_packing { + var wp : [8]u8 = [0; 8]; + const last : u64 = gen_wp(&wp, 32, 7); + var s : u64 = 7; + var i : usize = 0; + var seen : [4]usize = [0; 4]; + while (i < 32) { + s = lcg(s); + const code : u32 = ((s >> 16) % 4) as u32; + assert code_at(&wp, i) == code; + seen[code] += 1; + i += 1; + } + assert s == last; + assert seen[0] > 0 and seen[1] > 0 and seen[2] > 0 and seen[3] > 0; +} + +test generated_activations_span_the_i8_range { + var a : [4096]i8 = [0; 4096]; + const s : u64 = gen_a(&a, 4096, 1); + assert s != 1; + var lo : i32 = 0; + var hi : i32 = 0; + var i : usize = 0; + while (i < 4096) { + const v : i32 = a[i] as i32; + if (v < lo) { lo = v; } + if (v > hi) { hi = v; } + i += 1; + } + assert lo == -128; + assert hi == 127; +} + +test the_reference_row_on_a_hand_case { + ; k = 16, n = 16: a[kk] = kk + 1; w(kk, kk) = +1 for even kk, -1 for odd kk, w(kk, 15 - kk) = 3 + ; (a zero), everything else 0. Then acc[j] = (j + 1) for even j and -(j + 1) for odd j. + var a : [16]i8 = [0; 16]; + var wp : [64]u8 = [0; 64]; + var kk : usize = 0; + while (kk < 16) { + a[kk] = (kk + 1) as i8; + const i : usize = kk * 16 + kk; + var code : u32 = 1; + if (kk % 2 == 1) { code = 2; } + wp[i / 4] = ((wp[i / 4] as u32) | (code << (((i % 4) * 2) as u32))) as u8; + const z : usize = kk * 16 + 15 - kk; + if (z != i) { wp[z / 4] = ((wp[z / 4] as u32) | (3 << (((z % 4) * 2) as u32))) as u8; } + kk += 1; + } + var acc : [16]i32 = [5; 16]; + ref_row(&a, &wp, 16, 16, 0, &acc); + assert acc[0] == 1; + assert acc[1] == -2; + assert acc[6] == 7; + assert acc[15] == -16; +} + +test the_extremes_reach_128_k_without_overflow { + ; a = -128 everywhere, w = -1 everywhere: every entry is 128 * k. + var a : [64]i8 = [-128; 64]; + var wp : [256]u8 = [0xAA; 256]; + var acc : [16]i32 = [0; 16]; + ref_row(&a, &wp, 16, 64, 0, &acc); + assert acc[0] == 8192; + assert acc[15] == 8192; +} + +test mismatches_count_exactly_the_wrong_entries { + ; m = 3, n = 16, k = 32: c filled from the reference has 0 mismatches; one wrong entry has 1. + var a : [96]i8 = [0; 96]; + var wp : [128]u8 = [0; 128]; + const s : u64 = gen_a(&a, 96, 3); + const t : u64 = gen_wp(&wp, 512, s); + assert t != s; + var c : [48]i32 = [0; 48]; + var acc : [16]i32 = [0; 16]; + var row : usize = 0; + while (row < 3) { + ref_row(&a, &wp, 16, 32, row, &acc); + var j : usize = 0; + while (j < 16) { + c[row * 16 + j] = acc[j]; + j += 1; + } + row += 1; + } + assert mismatches(&a, &wp, &c, 3, 16, 32, 1, &acc) == 0; + c[16 + 5] += 1; + assert mismatches(&a, &wp, &c, 3, 16, 32, 1, &acc) == 1; + ; stride 2 checks rows 0 and 2 only, so the wrong entry in row 1 is not seen; + assert mismatches(&a, &wp, &c, 3, 16, 32, 2, &acc) == 0; + ; a wrong entry in the last row is seen at any stride. + c[47] -= 1; + assert mismatches(&a, &wp, &c, 3, 16, 32, 100, &acc) == 1; + assert mismatches(&a, &wp, &c, 3, 16, 32, 0, &acc) == 2; +} + +test the_500m_shapes_are_accepted { + assert shape_ok(4096, D_MODEL_500M, D_MODEL_500M) == true; + assert shape_ok(4096, D_FF_500M, D_MODEL_500M) == true; + assert shape_ok(4096, D_MODEL_500M, D_FF_500M) == true; + assert shape_ok(16384, D_FF_500M, D_MODEL_500M) == true; + assert shape_ok(17, 16, 16) == true; +} + +test unsupported_shapes_are_refused { + assert shape_ok(0, 1280, 1280) == false; + assert shape_ok(4096, 1288, 1280) == false; + assert shape_ok(4096, 1280, 1288) == false; + assert shape_ok(4096, 8, 1280) == false; + assert shape_ok(4096, 1280, 16777216) == false; + assert shape_ok(4096, 1280, 16777200) == true; + assert shape_ok(65535 * 128 + 1, 1280, 1280) == false; +} + +test the_grid_covers_every_row_and_column { + assert grid_dim(1280) == 10; + assert grid_dim(4096) == 32; + assert grid_dim(16384) == 128; + assert grid_dim(129) == 2; + assert grid_dim(1) == 1; +} + +invariant accumulator_cannot_overflow { 128 * K_MAX_I32 <= 2147483647 } + +invariant shared_tiles_fit_static_shared_memory { 2 * BM * SMEM_STRIDE <= 49152 } + +invariant the_stride_holds_a_k_step_and_keeps_16_byte_rows { SMEM_STRIDE >= BK and SMEM_STRIDE % 16 == 0 } + +invariant eight_warps_cover_the_block { (THREADS / 32) * 64 * 32 == BM * BN } + +invariant the_500m_shapes_are_whole_blocks { D_MODEL_500M % 128 == 0 and D_FF_500M % 128 == 0 and D_MODEL_500M % BK == 0 } From cce2fe5c25293a3c32b5accda455e4aaf5db02ed Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Sun, 18 Oct 2026 18:17:59 +0700 Subject: [PATCH 02/11] igla ptx: comments inside fn bodies are // (Refs #10404) Co-Authored-By: Claude Opus 5.5 --- specs/igla/coder/ternary_gemm_ptx.t27 | 28 +++++++++++++-------------- 1 file changed, 14 insertions(+), 14 deletions(-) diff --git a/specs/igla/coder/ternary_gemm_ptx.t27 b/specs/igla/coder/ternary_gemm_ptx.t27 index 4689974176..22382803d9 100644 --- a/specs/igla/coder/ternary_gemm_ptx.t27 +++ b/specs/igla/coder/ternary_gemm_ptx.t27 @@ -150,7 +150,7 @@ fn emit_prologue(out : []u8, cap : usize, w : usize) -> usize { x = ins(out, cap, x, "mov.u32 %r4, %tid.x;"); x = ins(out, cap, x, "mov.u32 %r5, %ctaid.x;"); x = ins(out, cap, x, "mov.u32 %r6, %ctaid.y;"); - ; A tile: row r7 = tid/2, chunk r8 = (tid%2)*16, global row r9, p1 = row inside m. + // A tile: row r7 = tid/2, chunk r8 = (tid%2)*16, global row r9, p1 = row inside m. x = ins(out, cap, x, "shr.u32 %r7, %r4, 1;"); x = ins(out, cap, x, "and.b32 %r8, %r4, 1;"); x = ins(out, cap, x, "shl.b32 %r8, %r8, 4;"); @@ -164,7 +164,7 @@ fn emit_prologue(out : []u8, cap : usize, w : usize) -> usize { x = ins(out, cap, x, "add.u64 %rd4, %rd4, %rd1;"); x = ins(out, cap, x, "cvt.u64.u32 %rd5, %r8;"); x = ins(out, cap, x, "add.u64 %rd4, %rd4, %rd5;"); - ; W tile: k-row r12 = tid/8, column group r13 = (tid%8)*16, global column r14, p2 = inside n. + // W tile: k-row r12 = tid/8, column group r13 = (tid%8)*16, global column r14, p2 = inside n. x = ins(out, cap, x, "shr.u32 %r12, %r4, 3;"); x = ins(out, cap, x, "and.b32 %r13, %r4, 7;"); x = ins(out, cap, x, "shl.b32 %r13, %r13, 4;"); @@ -176,14 +176,14 @@ fn emit_prologue(out : []u8, cap : usize, w : usize) -> usize { x = ins(out, cap, x, "add.u32 %r16, %r16, %r15;"); x = ins(out, cap, x, "shr.u32 %r17, %r2, 2;"); x = ins(out, cap, x, "shr.u32 %r18, %r14, 2;"); - ; Warp r19, lane r20, g r21, t r22, warp row r23, warp column r24. + // Warp r19, lane r20, g r21, t r22, warp row r23, warp column r24. x = ins(out, cap, x, "shr.u32 %r19, %r4, 5;"); x = ins(out, cap, x, "and.b32 %r20, %r4, 31;"); x = ins(out, cap, x, "shr.u32 %r21, %r20, 2;"); x = ins(out, cap, x, "and.b32 %r22, %r20, 3;"); x = ins(out, cap, x, "shr.u32 %r23, %r19, 2;"); x = ins(out, cap, x, "and.b32 %r24, %r19, 3;"); - ; Fragment bases: r25 = sA + (64 wm + g) * 48 + 4t, r27 = sB + (32 wn + g) * 48 + 4t. + // Fragment bases: r25 = sA + (64 wm + g) * 48 + 4t, r27 = sB + (32 wn + g) * 48 + 4t. x = ins(out, cap, x, "shl.b32 %r26, %r22, 2;"); x = ins(out, cap, x, "shl.b32 %r25, %r23, 6;"); x = ins(out, cap, x, "add.u32 %r25, %r25, %r21;"); @@ -523,14 +523,14 @@ test fragment_offsets_follow_the_48_byte_stride { var out : [32768]u8 = undefined; const m : usize = emit_ptx(&out, PTX_CAP); assert count_text(&out, m, " ld.shared.u32 ") == 24; - ; A: tile mi at 16 rows x 48 = 768 bytes, the g+8 rows 384 further, the upper k half 16 further. + // A: tile mi at 16 rows x 48 = 768 bytes, the g+8 rows 384 further, the upper k half 16 further. assert count_text(&out, m, " ld.shared.u32 %fa0, [%r25+0];\n") == 1; assert count_text(&out, m, " ld.shared.u32 %fa1, [%r25+384];\n") == 1; assert count_text(&out, m, " ld.shared.u32 %fa2, [%r25+16];\n") == 1; assert count_text(&out, m, " ld.shared.u32 %fa3, [%r25+400];\n") == 1; assert count_text(&out, m, " ld.shared.u32 %fa13, [%r25+2688];\n") == 1; assert count_text(&out, m, " ld.shared.u32 %fa15, [%r25+2704];\n") == 1; - ; B: tile ni at 8 columns x 48 = 384 bytes. + // B: tile ni at 8 columns x 48 = 384 bytes. assert count_text(&out, m, " ld.shared.u32 %fb6, [%r27+1152];\n") == 1; assert count_text(&out, m, " ld.shared.u32 %fb7, [%r27+1168];\n") == 1; } @@ -603,7 +603,7 @@ test emitting_twice_gives_the_same_bytes { } test the_decode_is_the_gpu_spec_decode { - ; ternary_gemm_gpu.t27 decode: 0 -> 0, 1 -> +1, 2 -> -1, 3 -> 0. + // ternary_gemm_gpu.t27 decode: 0 -> 0, 1 -> +1, 2 -> -1, 3 -> 0. assert weight(0) == 0; assert weight(1) == 1; assert weight(2) == -1; @@ -614,7 +614,7 @@ test packing_puts_lane_zero_in_the_low_bits { var wp : [4]u8 = [0; 4]; wp[0] = 0x1B; wp[3] = 0x80; - ; 0x1B = 00 01 10 11: lanes 0..3 hold 3, 2, 1, 0. + // 0x1B = 00 01 10 11: lanes 0..3 hold 3, 2, 1, 0. assert code_at(&wp, 0) == 3; assert code_at(&wp, 1) == 2; assert code_at(&wp, 2) == 1; @@ -658,8 +658,8 @@ test generated_activations_span_the_i8_range { } test the_reference_row_on_a_hand_case { - ; k = 16, n = 16: a[kk] = kk + 1; w(kk, kk) = +1 for even kk, -1 for odd kk, w(kk, 15 - kk) = 3 - ; (a zero), everything else 0. Then acc[j] = (j + 1) for even j and -(j + 1) for odd j. + // k = 16, n = 16: a[kk] = kk + 1; w(kk, kk) = +1 for even kk, -1 for odd kk, w(kk, 15 - kk) = 3 + // (a zero), everything else 0. Then acc[j] = (j + 1) for even j and -(j + 1) for odd j. var a : [16]i8 = [0; 16]; var wp : [64]u8 = [0; 64]; var kk : usize = 0; @@ -682,7 +682,7 @@ test the_reference_row_on_a_hand_case { } test the_extremes_reach_128_k_without_overflow { - ; a = -128 everywhere, w = -1 everywhere: every entry is 128 * k. + // a = -128 everywhere, w = -1 everywhere: every entry is 128 * k. var a : [64]i8 = [-128; 64]; var wp : [256]u8 = [0xAA; 256]; var acc : [16]i32 = [0; 16]; @@ -692,7 +692,7 @@ test the_extremes_reach_128_k_without_overflow { } test mismatches_count_exactly_the_wrong_entries { - ; m = 3, n = 16, k = 32: c filled from the reference has 0 mismatches; one wrong entry has 1. + // m = 3, n = 16, k = 32: c filled from the reference has 0 mismatches; one wrong entry has 1. var a : [96]i8 = [0; 96]; var wp : [128]u8 = [0; 128]; const s : u64 = gen_a(&a, 96, 3); @@ -713,9 +713,9 @@ test mismatches_count_exactly_the_wrong_entries { assert mismatches(&a, &wp, &c, 3, 16, 32, 1, &acc) == 0; c[16 + 5] += 1; assert mismatches(&a, &wp, &c, 3, 16, 32, 1, &acc) == 1; - ; stride 2 checks rows 0 and 2 only, so the wrong entry in row 1 is not seen; + // stride 2 checks rows 0 and 2 only, so the wrong entry in row 1 is not seen; assert mismatches(&a, &wp, &c, 3, 16, 32, 2, &acc) == 0; - ; a wrong entry in the last row is seen at any stride. + // a wrong entry in the last row is seen at any stride. c[47] -= 1; assert mismatches(&a, &wp, &c, 3, 16, 32, 100, &acc) == 1; assert mismatches(&a, &wp, &c, 3, 16, 32, 0, &acc) == 2; From 772be75248f5364dfed66ea6c1d45783112d119f Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Sun, 18 Oct 2026 18:19:22 +0700 Subject: [PATCH 03/11] igla ptx: pin the emitted length, fix the store-offset count (Refs #10404) Co-Authored-By: Claude Opus 5.5 --- specs/igla/coder/ternary_gemm_ptx.t27 | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/specs/igla/coder/ternary_gemm_ptx.t27 b/specs/igla/coder/ternary_gemm_ptx.t27 index 22382803d9..2dc9ea15bd 100644 --- a/specs/igla/coder/ternary_gemm_ptx.t27 +++ b/specs/igla/coder/ternary_gemm_ptx.t27 @@ -552,7 +552,7 @@ test thirty_two_guarded_pair_stores { assert count_text(&out, m, "st.global.v2.u32 ") == 32; assert count_text(&out, m, " @%p6 st.global.v2.u32 [%rd10], {%acc0, %acc1};\n") == 1; assert count_text(&out, m, " @%p6 st.global.v2.u32 [%rd10], {%acc62, %acc63};\n") == 1; - assert count_text(&out, m, " add.u32 %r41, %r39, 56;\n") == 1; + assert count_text(&out, m, " add.u32 %r41, %r39, 56;\n") == 4; assert count_text(&out, m, " add.u32 %r42, %r40, 24;\n") == 8; } @@ -569,7 +569,7 @@ test one_loop_two_barriers_one_return { test the_text_is_ascii_and_fits_its_buffer { var out : [32768]u8 = undefined; const m : usize = emit_ptx(&out, PTX_CAP); - assert m > 16000; + assert m == 21403; assert m < PTX_CAP; var i : usize = 0; while (i < m) { From 36b674bb3354942129eea37425a85374ac03fcd6 Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Sun, 18 Oct 2026 18:25:15 +0700 Subject: [PATCH 04/11] t27c: --backend ptx and ptx-run (cuda feature) as thin glue over the t27 emitter; delete orphan math_pslq.rs to pay the budget (Refs #10404) Co-Authored-By: Claude Opus 5.5 --- bootstrap/Cargo.toml | 3 + .../gen/rust/igla/coder/ternary_gemm_ptx.rs | 422 +++++++++++++ bootstrap/src/host/driver.rs | 39 ++ bootstrap/src/main.rs | 29 +- bootstrap/src/math_pslq.rs | 589 ------------------ docs/reports/orphan_modules.json | 4 +- 6 files changed, 494 insertions(+), 592 deletions(-) create mode 100644 bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs delete mode 100644 bootstrap/src/math_pslq.rs diff --git a/bootstrap/Cargo.toml b/bootstrap/Cargo.toml index c0d9e52da0..451804c899 100644 --- a/bootstrap/Cargo.toml +++ b/bootstrap/Cargo.toml @@ -14,6 +14,9 @@ license = "MIT" # `audio` (Google Discovery Engine). Built without it, they still parse and # exit with a message that names this feature. net = ["dep:reqwest"] +# `cuda`: `t27c ptx-run` (#10404) links libcuda and runs the kernel that +# specs/igla/coder/ternary_gemm_ptx.t27 emits. Built without it, the subcommand does not exist. +cuda = [] # `server`: `t27c serve` (axum), with its sandbox proxy, Railway client and # JWT sandbox tokens. Implies `net`. server = [ diff --git a/bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs b/bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs new file mode 100644 index 0000000000..edddd86a72 --- /dev/null +++ b/bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs @@ -0,0 +1,422 @@ +// Generated from .t27 spec +// DO NOT EDIT — generated by t27c +// NOT LOWERED BY THIS BACKEND: 21 test(s), 5 invariant(s), 0 function(s) with no body. +// This backend emits declarations only. The spec's checks live in +// the Zig and Verilog outputs; do not read this file as verified. + +pub const KERNEL_NAME: &'static str = "ternary_gemm_s8"; + +pub const BM: u32 = 128; + +pub const BN: u32 = 128; + +pub const BK: u32 = 32; + +pub const THREADS: u32 = 256; + +pub const SMEM_STRIDE: u32 = 48; + +pub const MMA_PER_STEP: u32 = 16; + +pub const K_MAX_I32: usize = 16777215; + +pub const GRID_Y_MAX: usize = 65535; + +pub const PTX_CAP: usize = 32768; + +pub const D_MODEL_500M: usize = 1280; + +pub const D_FF_500M: usize = 4096; + +pub fn put_s(out: &mut [u8], cap: usize, w: usize, s: &'static str) -> usize { + let mut i: usize = 0; + while (i < s.len()) { + if ((w + i) < cap) { + out[((w + i)) as usize] = s.as_bytes()[(i) as usize]; + } + i += 1; + } + return (w + s.len()); +} + +pub fn put_u(out: &mut [u8], cap: usize, w: usize, v: u32) -> usize { + let mut d: u32 = 1; + while ((v / d) >= 10) { + d = (d * 10); + } + let mut x: usize = w; + while (d > 0) { + let c: u8 = (b'0' + (((v / d) % 10) as u8)); + if (x < cap) { + out[(x) as usize] = c; + } + x += 1; + d = (d / 10); + } + return x; +} + +pub fn ins(out: &mut [u8], cap: usize, w: usize, s: &'static str) -> usize { + let x: usize = put_s(out, cap, w, " "); + return put_s(out, cap, put_s(out, cap, x, s), "\n"); +} + +pub fn put_quad(out: &mut [u8], cap: usize, w: usize, name: &'static str, first: u32) -> usize { + let mut x: usize = put_s(out, cap, w, "{"); + let mut j: u32 = 0; + while (j < 4) { + if (j > 0) { + x = put_s(out, cap, x, ", "); + } + x = put_u(out, cap, put_s(out, cap, x, name), (first + j)); + j += 1; + } + return put_s(out, cap, x, "}"); +} + +pub fn emit_header(out: &mut [u8], cap: usize, w: usize) -> usize { + let mut x: usize = put_s(out, cap, w, "// ternary_gemm_s8 -- emitted by t27c from specs/igla/coder/ternary_gemm_ptx.t27 (#10404). Do not edit.\n"); + x = put_s(out, cap, x, ".version 7.0\n.target sm_80\n.address_size 64\n\n"); + x = put_s(out, cap, x, ".visible .entry ternary_gemm_s8(\n"); + x = put_s(out, cap, x, " .param .u64 p_a,\n .param .u64 p_w,\n .param .u64 p_c,\n"); + x = put_s(out, cap, x, " .param .u32 p_m,\n .param .u32 p_n,\n .param .u32 p_k\n)\n"); + x = put_s(out, cap, x, ".maxntid 256, 1, 1\n{\n"); + x = ins(out, cap, x, ".reg .pred %p<8>;"); + x = ins(out, cap, x, ".reg .b32 %r<48>;"); + x = ins(out, cap, x, ".reg .b32 %acc<64>;"); + x = ins(out, cap, x, ".reg .b32 %fa<16>;"); + x = ins(out, cap, x, ".reg .b32 %fb<8>;"); + x = ins(out, cap, x, ".reg .b64 %rd<12>;"); + x = ins(out, cap, x, ".shared .align 16 .b8 sA[6144];"); + x = ins(out, cap, x, ".shared .align 16 .b8 sB[6144];"); + return x; +} + +pub fn emit_prologue(out: &mut [u8], cap: usize, w: usize) -> usize { + let mut x: usize = ins(out, cap, w, "ld.param.u64 %rd1, [p_a];"); + x = ins(out, cap, x, "ld.param.u64 %rd2, [p_w];"); + x = ins(out, cap, x, "ld.param.u64 %rd3, [p_c];"); + x = ins(out, cap, x, "ld.param.u32 %r1, [p_m];"); + x = ins(out, cap, x, "ld.param.u32 %r2, [p_n];"); + x = ins(out, cap, x, "ld.param.u32 %r3, [p_k];"); + x = ins(out, cap, x, "cvta.to.global.u64 %rd1, %rd1;"); + x = ins(out, cap, x, "cvta.to.global.u64 %rd2, %rd2;"); + x = ins(out, cap, x, "cvta.to.global.u64 %rd3, %rd3;"); + x = ins(out, cap, x, "mov.u32 %r4, %tid.x;"); + x = ins(out, cap, x, "mov.u32 %r5, %ctaid.x;"); + x = ins(out, cap, x, "mov.u32 %r6, %ctaid.y;"); + x = ins(out, cap, x, "shr.u32 %r7, %r4, 1;"); + x = ins(out, cap, x, "and.b32 %r8, %r4, 1;"); + x = ins(out, cap, x, "shl.b32 %r8, %r8, 4;"); + x = ins(out, cap, x, "shl.b32 %r9, %r6, 7;"); + x = ins(out, cap, x, "add.u32 %r9, %r9, %r7;"); + x = ins(out, cap, x, "setp.lt.u32 %p1, %r9, %r1;"); + x = ins(out, cap, x, "mov.u32 %r10, sA;"); + x = ins(out, cap, x, "mad.lo.u32 %r11, %r7, 48, %r8;"); + x = ins(out, cap, x, "add.u32 %r11, %r11, %r10;"); + x = ins(out, cap, x, "mul.wide.u32 %rd4, %r9, %r3;"); + x = ins(out, cap, x, "add.u64 %rd4, %rd4, %rd1;"); + x = ins(out, cap, x, "cvt.u64.u32 %rd5, %r8;"); + x = ins(out, cap, x, "add.u64 %rd4, %rd4, %rd5;"); + x = ins(out, cap, x, "shr.u32 %r12, %r4, 3;"); + x = ins(out, cap, x, "and.b32 %r13, %r4, 7;"); + x = ins(out, cap, x, "shl.b32 %r13, %r13, 4;"); + x = ins(out, cap, x, "shl.b32 %r14, %r5, 7;"); + x = ins(out, cap, x, "add.u32 %r14, %r14, %r13;"); + x = ins(out, cap, x, "setp.lt.u32 %p2, %r14, %r2;"); + x = ins(out, cap, x, "mov.u32 %r15, sB;"); + x = ins(out, cap, x, "mad.lo.u32 %r16, %r13, 48, %r12;"); + x = ins(out, cap, x, "add.u32 %r16, %r16, %r15;"); + x = ins(out, cap, x, "shr.u32 %r17, %r2, 2;"); + x = ins(out, cap, x, "shr.u32 %r18, %r14, 2;"); + x = ins(out, cap, x, "shr.u32 %r19, %r4, 5;"); + x = ins(out, cap, x, "and.b32 %r20, %r4, 31;"); + x = ins(out, cap, x, "shr.u32 %r21, %r20, 2;"); + x = ins(out, cap, x, "and.b32 %r22, %r20, 3;"); + x = ins(out, cap, x, "shr.u32 %r23, %r19, 2;"); + x = ins(out, cap, x, "and.b32 %r24, %r19, 3;"); + x = ins(out, cap, x, "shl.b32 %r26, %r22, 2;"); + x = ins(out, cap, x, "shl.b32 %r25, %r23, 6;"); + x = ins(out, cap, x, "add.u32 %r25, %r25, %r21;"); + x = ins(out, cap, x, "mul.lo.u32 %r25, %r25, 48;"); + x = ins(out, cap, x, "add.u32 %r25, %r25, %r26;"); + x = ins(out, cap, x, "add.u32 %r25, %r25, %r10;"); + x = ins(out, cap, x, "shl.b32 %r27, %r24, 5;"); + x = ins(out, cap, x, "add.u32 %r27, %r27, %r21;"); + x = ins(out, cap, x, "mul.lo.u32 %r27, %r27, 48;"); + x = ins(out, cap, x, "add.u32 %r27, %r27, %r26;"); + x = ins(out, cap, x, "add.u32 %r27, %r27, %r15;"); + x = ins(out, cap, x, "mov.u32 %r28, 0;"); + let mut i: u32 = 0; + while (i < 64) { + x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " mov.b32 %acc"), i), ", 0;\n"); + i += 1; + } + return x; +} + +pub fn emit_tile_loads(out: &mut [u8], cap: usize, w: usize) -> usize { + let mut x: usize = ins(out, cap, w, "add.u32 %r29, %r28, %r8;"); + x = ins(out, cap, x, "setp.lt.u32 %p4, %r29, %r3;"); + x = ins(out, cap, x, "and.pred %p4, %p4, %p1;"); + x = ins(out, cap, x, "mov.b32 %r30, 0;"); + x = ins(out, cap, x, "mov.b32 %r31, 0;"); + x = ins(out, cap, x, "mov.b32 %r32, 0;"); + x = ins(out, cap, x, "mov.b32 %r33, 0;"); + x = ins(out, cap, x, "cvt.u64.u32 %rd5, %r28;"); + x = ins(out, cap, x, "add.u64 %rd6, %rd4, %rd5;"); + x = ins(out, cap, x, "@%p4 ld.global.v4.u32 {%r30, %r31, %r32, %r33}, [%rd6];"); + x = ins(out, cap, x, "st.shared.v4.u32 [%r11], {%r30, %r31, %r32, %r33};"); + x = ins(out, cap, x, "add.u32 %r34, %r28, %r12;"); + x = ins(out, cap, x, "setp.lt.u32 %p5, %r34, %r3;"); + x = ins(out, cap, x, "and.pred %p5, %p5, %p2;"); + x = ins(out, cap, x, "mov.b32 %r35, 0;"); + x = ins(out, cap, x, "mul.wide.u32 %rd7, %r34, %r17;"); + x = ins(out, cap, x, "cvt.u64.u32 %rd8, %r18;"); + x = ins(out, cap, x, "add.u64 %rd7, %rd7, %rd8;"); + x = ins(out, cap, x, "add.u64 %rd9, %rd7, %rd2;"); + x = ins(out, cap, x, "@%p5 ld.global.u32 %r35, [%rd9];"); + let mut j: u32 = 0; + while (j < 16) { + x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " shr.u32 %r36, %r35, "), (2 * j)), ";\n"); + x = ins(out, cap, x, "and.b32 %r37, %r36, 1;"); + x = ins(out, cap, x, "shr.u32 %r38, %r36, 1;"); + x = ins(out, cap, x, "and.b32 %r38, %r38, 1;"); + x = ins(out, cap, x, "sub.s32 %r37, %r37, %r38;"); + x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " st.shared.u8 [%r16+"), (j * SMEM_STRIDE)), "], %r37;\n"); + j += 1; + } + return x; +} + +pub fn emit_lds(out: &mut [u8], cap: usize, w: usize, reg: &'static str, i: u32, base: &'static str, off: u32) -> usize { + let mut x: usize = put_u(out, cap, put_s(out, cap, put_s(out, cap, w, " ld.shared.u32 "), reg), i); + x = put_s(out, cap, put_s(out, cap, put_s(out, cap, x, ", ["), base), "+"); + return put_s(out, cap, put_u(out, cap, x, off), "];\n"); +} + +pub fn emit_mma_step(out: &mut [u8], cap: usize, w: usize) -> usize { + let mut x: usize = w; + let mut mi: u32 = 0; + while (mi < 4) { + let off: u32 = ((mi * 16) * SMEM_STRIDE); + x = emit_lds(out, cap, x, "%fa", (mi * 4), "%r25", off); + x = emit_lds(out, cap, x, "%fa", ((mi * 4) + 1), "%r25", (off + (8 * SMEM_STRIDE))); + x = emit_lds(out, cap, x, "%fa", ((mi * 4) + 2), "%r25", (off + 16)); + x = emit_lds(out, cap, x, "%fa", ((mi * 4) + 3), "%r25", ((off + (8 * SMEM_STRIDE)) + 16)); + mi += 1; + } + let mut ni: u32 = 0; + while (ni < 4) { + x = emit_lds(out, cap, x, "%fb", (ni * 2), "%r27", ((ni * 8) * SMEM_STRIDE)); + x = emit_lds(out, cap, x, "%fb", ((ni * 2) + 1), "%r27", (((ni * 8) * SMEM_STRIDE) + 16)); + ni += 1; + } + mi = 0; + while (mi < 4) { + ni = 0; + while (ni < 4) { + let b: u32 = (((mi * 4) + ni) * 4); + x = put_s(out, cap, x, " mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 "); + x = put_s(out, cap, put_quad(out, cap, x, "%acc", b), ", "); + x = put_s(out, cap, put_quad(out, cap, x, "%fa", (mi * 4)), ", {%fb"); + x = put_u(out, cap, put_s(out, cap, put_u(out, cap, x, (ni * 2)), ", %fb"), ((ni * 2) + 1)); + x = put_s(out, cap, put_quad(out, cap, put_s(out, cap, x, "}, "), "%acc", b), ";\n"); + ni += 1; + } + mi += 1; + } + return x; +} + +pub fn emit_store(out: &mut [u8], cap: usize, w: usize, dr: u32, dc: u32, b: u32) -> usize { + let mut x: usize = put_s(out, cap, put_u(out, cap, put_s(out, cap, w, " add.u32 %r41, %r39, "), dr), ";\n"); + x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " add.u32 %r42, %r40, "), dc), ";\n"); + x = ins(out, cap, x, "setp.lt.u32 %p6, %r41, %r1;"); + x = ins(out, cap, x, "setp.lt.u32 %p7, %r42, %r2;"); + x = ins(out, cap, x, "and.pred %p6, %p6, %p7;"); + x = ins(out, cap, x, "cvt.u64.u32 %rd10, %r42;"); + x = ins(out, cap, x, "mad.wide.u32 %rd10, %r41, %r2, %rd10;"); + x = ins(out, cap, x, "shl.b64 %rd10, %rd10, 2;"); + x = ins(out, cap, x, "add.u64 %rd10, %rd10, %rd3;"); + x = put_u(out, cap, put_s(out, cap, x, " @%p6 st.global.v2.u32 [%rd10], {%acc"), b); + return put_s(out, cap, put_u(out, cap, put_s(out, cap, x, ", %acc"), (b + 1)), "};\n"); +} + +pub fn emit_epilogue(out: &mut [u8], cap: usize, w: usize) -> usize { + let mut x: usize = ins(out, cap, w, "shl.b32 %r39, %r6, 7;"); + x = ins(out, cap, x, "shl.b32 %r41, %r23, 6;"); + x = ins(out, cap, x, "add.u32 %r39, %r39, %r41;"); + x = ins(out, cap, x, "add.u32 %r39, %r39, %r21;"); + x = ins(out, cap, x, "shl.b32 %r40, %r5, 7;"); + x = ins(out, cap, x, "shl.b32 %r41, %r24, 5;"); + x = ins(out, cap, x, "add.u32 %r40, %r40, %r41;"); + x = ins(out, cap, x, "shl.b32 %r41, %r22, 1;"); + x = ins(out, cap, x, "add.u32 %r40, %r40, %r41;"); + let mut mi: u32 = 0; + while (mi < 4) { + let mut ni: u32 = 0; + while (ni < 4) { + let b: u32 = (((mi * 4) + ni) * 4); + x = emit_store(out, cap, x, (mi * 16), (ni * 8), b); + x = emit_store(out, cap, x, ((mi * 16) + 8), (ni * 8), (b + 2)); + ni += 1; + } + mi += 1; + } + return x; +} + +pub fn emit_ptx(out: &mut [u8], cap: usize) -> usize { + let mut x: usize = emit_header(out, cap, 0); + x = emit_prologue(out, cap, x); + x = put_s(out, cap, x, "LOOP_K:\n"); + x = ins(out, cap, x, "setp.ge.u32 %p3, %r28, %r3;"); + x = ins(out, cap, x, "@%p3 bra DONE_K;"); + x = emit_tile_loads(out, cap, x); + x = ins(out, cap, x, "bar.sync 0;"); + x = emit_mma_step(out, cap, x); + x = ins(out, cap, x, "bar.sync 0;"); + x = ins(out, cap, x, "add.u32 %r28, %r28, 32;"); + x = ins(out, cap, x, "bra LOOP_K;"); + x = put_s(out, cap, x, "DONE_K:\n"); + x = emit_epilogue(out, cap, x); + x = ins(out, cap, x, "ret;"); + return put_s(out, cap, x, "}\n"); +} + +pub fn grid_dim(x: usize) -> u32 { + return (((x + 127) / 128) as u32); +} + +pub fn shape_ok(m: usize, n: usize, k: usize) -> bool { + if (((m < 1) || (n < 16)) || (k < 16)) { + return false; + } + if (((n % 16) != 0) || ((k % 16) != 0)) { + return false; + } + if (k > K_MAX_I32) { + return false; + } + return (((m + 127) / 128) <= GRID_Y_MAX); +} + +pub fn lcg(s: u64) -> u64 { + return (((s * 1103515245) + 12345) % 2147483648); +} + +pub fn gen_a(a: &mut [i8], count: usize, seed: u64) -> u64 { + let mut s: u64 = seed; + let mut i: usize = 0; + while (i < count) { + s = lcg(s); + a[(i) as usize] = (((((s >> 16) % 256) as i32) - 128) as i8); + i += 1; + } + return s; +} + +pub fn gen_wp(wp: &mut [u8], count: usize, seed: u64) -> u64 { + let mut s: u64 = seed; + let mut i: usize = 0; + while (i < count) { + let mut byte: u32 = 0; + let mut lane: u32 = 0; + while (lane < 4) { + s = lcg(s); + let code: u32 = (((s >> 16) % 4) as u32); + byte = (byte | (code << (lane * 2))); + lane += 1; + } + wp[((i / 4)) as usize] = (byte as u8); + i += 4; + } + return s; +} + +pub fn code_at(wp: &[u8], i: usize) -> u32 { + let b: u32 = (wp[((i / 4)) as usize] as u32); + let sh: u32 = (((i % 4) * 2) as u32); + return ((b >> sh) & 3); +} + +pub fn weight(code: u32) -> i32 { + return (((code & 1) as i32) - (((code >> 1) & 1) as i32)); +} + +pub fn ref_row(a: impl AsRef<[i8]>, wp: &[u8], n: usize, k: usize, row: usize, acc: &mut [i32]) -> () { + let a: &[i8] = a.as_ref(); + let mut j: usize = 0; + while (j < n) { + acc[(j) as usize] = 0; + j += 1; + } + let mut kk: usize = 0; + while (kk < k) { + let v: i32 = (a[(((row * k) + kk)) as usize] as i32); + if (v != 0) { + let base: usize = (kk * n); + let mut col: usize = 0; + while (col < n) { + acc[(col) as usize] += (v * weight(code_at(wp, (base + col)))); + col += 1; + } + } + kk += 1; + } +} + +pub fn mismatches(a: impl AsRef<[i8]>, wp: &[u8], c: impl AsRef<[i32]>, m: usize, n: usize, k: usize, stride: usize, acc: &mut [i32]) -> usize { + let a: &[i8] = a.as_ref(); + let c: &[i32] = c.as_ref(); + let mut st: usize = stride; + if (st == 0) { + st = 1; + } + let mut bad: usize = 0; + let mut row: usize = 0; + while (row < m) { + ref_row(a, wp, n, k, row, acc); + let mut col: usize = 0; + while (col < n) { + if (c[(((row * n) + col)) as usize] != acc[(col) as usize]) { + bad += 1; + } + col += 1; + } + if (((row + 1) < m) && ((row + st) >= m)) { + row = (m - 1); + } else { + row += st; + } + } + return bad; +} + +pub fn at_text(t: &[u8], n: usize, at: usize, s: &'static str) -> bool { + if ((at + s.len()) > n) { + return false; + } + let mut i: usize = 0; + while (i < s.len()) { + if (t[((at + i)) as usize] != s.as_bytes()[(i) as usize]) { + return false; + } + i += 1; + } + return true; +} + +pub fn count_text(t: &[u8], n: usize, s: &'static str) -> usize { + let mut c: usize = 0; + let mut i: usize = 0; + while (i < n) { + if at_text(t, n, i, s) { + c += 1; + } + i += 1; + } + return c; +} + diff --git a/bootstrap/src/host/driver.rs b/bootstrap/src/host/driver.rs index 575cde4ee7..6b8d07f9b1 100644 --- a/bootstrap/src/host/driver.rs +++ b/bootstrap/src/host/driver.rs @@ -298,3 +298,42 @@ mod tests { assert_eq!(snap.weight_addr_64(), 0x2000_0000_1000_0000); } } + +// `t27c ptx-run` (#10404): load the PTX that specs/igla/coder/ternary_gemm_ptx.t27 wrote, run it, count mismatches with its reference. +#[cfg(feature = "cuda")] #[link(name = "cuda")] +extern "C" { + fn cuInit(flags: u32) -> i32; fn cuDeviceGet(dev: *mut i32, ordinal: i32) -> i32; + fn cuCtxCreate_v2(ctx: *mut usize, flags: u32, dev: i32) -> i32; fn cuCtxSynchronize() -> i32; + fn cuModuleLoadData(module: *mut usize, image: *const u8) -> i32; fn cuModuleGetFunction(func: *mut usize, module: usize, name: *const u8) -> i32; + fn cuMemAlloc_v2(ptr: *mut u64, bytes: usize) -> i32; fn cuMemcpyHtoD_v2(dst: u64, src: *const u8, bytes: usize) -> i32; + fn cuMemcpyDtoH_v2(dst: *mut u8, src: u64, bytes: usize) -> i32; + fn cuLaunchKernel(f: usize, gx: u32, gy: u32, gz: u32, bx: u32, by: u32, bz: u32, smem: u32, stream: usize, params: *mut *mut u8, extra: usize) -> i32; +} +#[cfg(feature = "cuda")] +pub fn ptx_run(m: usize, n: usize, k: usize, iters: u32, stride: usize) -> anyhow::Result<()> { + use crate::ternary_gemm_ptx as t; + anyhow::ensure!(t::shape_ok(m, n, k), "m={} n={} k={}: refused by shape_ok in ternary_gemm_ptx.t27", m, n, k); + let ck = |r: i32| if r == 0 { Ok(()) } else { Err(anyhow::anyhow!("CUDA driver error {}", r)) }; + let mut ptx = vec![0u8; t::PTX_CAP + 1]; + ptx.truncate(t::emit_ptx(&mut ptx, t::PTX_CAP)); ptx.push(0); + let (mut a, mut wp, mut c, mut acc) = (vec![0i8; m * k], vec![0u8; k * n / 4], vec![0i32; m * n], vec![0i32; n]); + t::gen_wp(&mut wp, k * n, t::gen_a(&mut a, m * k, 1)); + let (mut dev, mut ctx, mut module, mut f, mut da, mut dw, mut dc, mut pm, mut pn, mut pk) = (0i32, 0usize, 0usize, 0usize, 0u64, 0u64, 0u64, m as u32, n as u32, k as u32); + let ms = unsafe { + ck(cuInit(0))?; ck(cuDeviceGet(&mut dev, 0))?; ck(cuCtxCreate_v2(&mut ctx, 0, dev))?; + ck(cuModuleLoadData(&mut module, ptx.as_ptr()))?; ck(cuModuleGetFunction(&mut f, module, b"ternary_gemm_s8\0".as_ptr()))?; + ck(cuMemAlloc_v2(&mut da, a.len()))?; ck(cuMemAlloc_v2(&mut dw, wp.len()))?; ck(cuMemAlloc_v2(&mut dc, c.len() * 4))?; + ck(cuMemcpyHtoD_v2(da, a.as_ptr() as *const u8, a.len()))?; ck(cuMemcpyHtoD_v2(dw, wp.as_ptr(), wp.len()))?; + let mut args: [*mut u8; 6] = [&mut da as *mut u64 as *mut u8, &mut dw as *mut u64 as *mut u8, &mut dc as *mut u64 as *mut u8, + &mut pm as *mut u32 as *mut u8, &mut pn as *mut u32 as *mut u8, &mut pk as *mut u32 as *mut u8]; + let mut launch = || ck(cuLaunchKernel(f, t::grid_dim(n), t::grid_dim(m), 1, t::THREADS, 1, 1, 0, 0, args.as_mut_ptr(), 0)); + launch()?; ck(cuCtxSynchronize())?; ck(cuMemcpyDtoH_v2(c.as_mut_ptr() as *mut u8, dc, c.len() * 4))?; + let t0 = std::time::Instant::now(); + for _ in 0..iters { launch()?; } ck(cuCtxSynchronize())?; + t0.elapsed().as_secs_f64() * 1e3 / iters.max(1) as f64 + }; + let bad = t::mismatches(&a, &wp, &c, m, n, k, stride, &mut acc); + println!("ptx-run m={} n={} k={} rows_checked_every={} mismatches={} iters={} ms_per_call={:.4} tops={:.2}", m, n, k, stride.max(1), bad, iters, ms, 2.0 * (m * n * k) as f64 / ms / 1e9); + anyhow::ensure!(bad == 0, "{} entries differ from the t27 reference", bad); + Ok(()) +} diff --git a/bootstrap/src/main.rs b/bootstrap/src/main.rs index 9c5481aa0b..07f21bf390 100644 --- a/bootstrap/src/main.rs +++ b/bootstrap/src/main.rs @@ -95,6 +95,10 @@ mod ed25519; #[path = "../gen/rust/tri/fmt_source.rs"] #[allow(dead_code, unused_parens, unused_mut, unused_assignments, unused_variables, non_snake_case, non_upper_case_globals, clippy::all)] mod fmt_source; +// specs/igla/coder/ternary_gemm_ptx.t27 (#10404), lowered by `t27c gen-rust`: the PTX of the ternary GEMM kernel. Never hand-edit; regenerate. +#[path = "../gen/rust/igla/coder/ternary_gemm_ptx.rs"] +#[allow(dead_code, unused_parens, unused_mut, unused_assignments, unused_variables, non_snake_case, non_upper_case_globals, clippy::all)] +mod ternary_gemm_ptx; mod memory; mod trit_stdlib; mod behavior_sva; @@ -1342,7 +1346,7 @@ enum Commands { Compile { /// Input file path input: String, - /// Backend: zig, verilog, c, or ml + /// Backend: zig, verilog, c, ml, or ptx (one kernel only: specs/igla/coder/ternary_gemm_gpu.t27) #[arg(long, default_value = "zig")] backend: String, /// Output file path (default: input with backend extension) @@ -1553,6 +1557,18 @@ enum Commands { input: String, }, + /// Run the PTX kernel t27 emitted on a CUDA GPU and check it against the t27 reference (#10404) + #[cfg(feature = "cuda")] + #[command(name = "ptx-run")] + PtxRun { + #[arg(long, default_value_t = 4096)] m: usize, + #[arg(long, default_value_t = 1280)] n: usize, + #[arg(long, default_value_t = 1280)] k: usize, + #[arg(long, default_value_t = 20)] iters: u32, + /// Check every stride-th row, and always the last one, against the reference + #[arg(long, default_value_t = 1)] stride: usize, + }, + /// Dependency graph of .t27 modules Graph { #[arg(long, default_value = ".")] @@ -6224,6 +6240,7 @@ fn backend_extension(backend: &str) -> &str { "verilog" => ".v", "c" => ".c", "rust" => ".rs", + "ptx" => ".ptx", _ => ".zig", } } @@ -6264,6 +6281,12 @@ fn run_compile(input_path: &str, backend: &str, output: Option<&str>) -> anyhow: } let code = match backend { + "ptx" if path.ends_with("specs/igla/coder/ternary_gemm_gpu.t27") => { + let mut buf = vec![0u8; ternary_gemm_ptx::PTX_CAP]; + let n = ternary_gemm_ptx::emit_ptx(&mut buf, ternary_gemm_ptx::PTX_CAP); + String::from_utf8_lossy(&buf[..n]).into_owned() + } + "ptx" => anyhow::bail!("--backend ptx lowers one kernel only: specs/igla/coder/ternary_gemm_gpu.t27 (#10404)"), "verilog" => { let mut cg = compiler::VerilogCodegen::new(); cg.gen_verilog(&ast); @@ -11797,6 +11820,8 @@ async fn main() -> anyhow::Result<()> { Commands::Bench { input } => run_bench(&input)?, Commands::Explain { input } => run_explain(&input)?, Commands::Fmt { input } => run_fmt(&input)?, + #[cfg(feature = "cuda")] + Commands::PtxRun { m, n, k, iters, stride } => host::driver::ptx_run(m, n, k, iters, stride)?, Commands::Graph { repo_root, format } => run_graph(&repo_root, &format)?, Commands::Doc { input, output_dir } => run_doc(&input, &output_dir)?, Commands::DocAll { repo_root, output_dir } => run_doc_all(&repo_root, &output_dir)?, @@ -12239,6 +12264,8 @@ fn main() -> anyhow::Result<()> { Commands::Bench { input } => run_bench(&input)?, Commands::Explain { input } => run_explain(&input)?, Commands::Fmt { input } => run_fmt(&input)?, + #[cfg(feature = "cuda")] + Commands::PtxRun { m, n, k, iters, stride } => host::driver::ptx_run(m, n, k, iters, stride)?, Commands::Graph { repo_root, format } => run_graph(&repo_root, &format)?, Commands::Doc { input, output_dir } => run_doc(&input, &output_dir)?, Commands::DocAll { repo_root, output_dir } => run_doc_all(&repo_root, &output_dir)?, diff --git a/bootstrap/src/math_pslq.rs b/bootstrap/src/math_pslq.rs deleted file mode 100644 index b8fc5bf27c..0000000000 --- a/bootstrap/src/math_pslq.rs +++ /dev/null @@ -1,589 +0,0 @@ -//! `tri math pslq` — PSLQ integer relation finder with high-precision arithmetic. -//! -//! PSLQ (Bailey-Borwein-Fein) algorithm finds integer relations among -//! high-precision floating-point numbers. This implementation uses rug for arbitrary -//! precision rational arithmetic. -//! -//! # Mathematical Background -//! -//! Given a vector `x = (x_1, ..., x_n)` of floating-point numbers, PSLQ finds -//! integers `a = (a_1, ..., a_n)` such that: -//! -//! ```text -//! a_1*x_1 + a_2*x_2 + ... + a_n*x_n = 0 -//! ``` -//! -//! within a specified tolerance. The algorithm builds a relation matrix incrementally -//! using the LLL lattice basis reduction. -//! -//! # Trinity Basis -//! -//! The canonical Trinity basis for logarithmic constant space is: -//! `{ln 3, ln π, ln φ, 1}` -//! -//! This corresponds to the Sacred Formula parameter space: -//! `V = n * 3^m * π^k * φ^p * e^q * 1^r` -//! -//! # Usage -//! -//! ```bash -//! tri math pslq --vector "ln(3),ln(pi),ln(phi),1" --max-coeff 12 --precision 150 -//! tri math pslq --basis trinity --check-targets -//! ``` - -use anyhow::Context; -use clap::Subcommand; -use rug::{ops::Pow, Integer, Rational}; -use std::f64::consts::PI; - -/// Golden ratio φ = (1 + √5) / 2 -const PHI: f64 = 1.618033988749895; - -/// Default precision for high-precision calculations (bits) -const DEFAULT_PRECISION: u32 = 150; - -/// Default maximum coefficient bound -const DEFAULT_MAX_COEFF: i32 = 12; - -/// Default tolerance for relation detection (norm of vector) -const DEFAULT_TOLERANCE: f64 = 1e-50; - -#[derive(Subcommand, Debug, Clone)] -pub enum PslqCommands { - /// Find integer relations for a given vector - Vector { - /// Comma-separated values (e.g., "ln(phi),ln(pi),1,ln(2)") - #[arg(long)] - vector: String, - - /// Target value to find relation for (e.g., "137.036") - #[arg(long)] - target: Option, - - /// Maximum coefficient magnitude (default: 12) - #[arg(long, default_value = "12")] - max_coeff: i32, - - /// Precision in bits (default: 150) - #[arg(long, default_value = "150")] - precision: u32, - - /// Tolerance for relation detection (default: 1e-50) - #[arg(long, default_value = "1e-50")] - tolerance: Option, - }, - - /// Use predefined Trinity basis with canonical vectors - Basis { - /// Basis name: trinity, extended, reduced - #[arg(long, default_value = "trinity")] - basis: String, - - /// Check against PDG targets - #[arg(long)] - check_targets: bool, - - /// Check ZIP-derived formulas - #[arg(long)] - check_zip: bool, - }, - - /// Run verification on Sacred Formula catalog - Verify { - /// Path to catalog JSON (default: research/sacred_formula_catalog.json) - #[arg(long)] - catalog: Option, - - /// Minimum error % to include (default: 0.1) - #[arg(long, default_value = "0.1")] - max_error: f64, - }, -} - -/// PSLQ relation result -#[derive(Debug, Clone)] -pub struct PslqRelation { - /// Coefficients (a_1, ..., a_n) - pub coefficients: Vec, - /// Norm of the coefficient vector - pub norm: f64, - /// L2 norm of residual (how close to zero) - pub residual_norm: f64, - /// Quality score (lower is better) - pub quality: f64, -} - -/// High-precision float context -pub struct HighPrecisionContext { - precision: u32, -} - -impl HighPrecisionContext { - pub fn new(precision: u32) -> Self { - Self { precision } - } -} - -/// Compute ln(x) with high precision -fn hp_ln(context: &HighPrecisionContext, x: f64) -> rug::Float { - let f = rug::Float::with_val(context.precision, x); - f.ln() -} - -/// Compute φ^p with high precision -fn hp_phi_pow(context: &HighPrecisionContext, p: i32) -> rug::Float { - let phi = rug::Float::with_val(context.precision, PHI); - phi.pow(p) -} - -/// Compute π^k with high precision -fn hp_pi_pow(context: &HighPrecisionContext, k: i32) -> rug::Float { - let pi = rug::Float::with_val(context.precision, PI); - pi.pow(k) -} - -/// Compute ln(3) with high precision -fn hp_ln_3(context: &HighPrecisionContext) -> rug::Float { - let three = rug::Float::with_val(context.precision, 3.0); - three.ln() -} - -/// Parse a vector of values to high-precision Floats -fn parse_vector(context: &HighPrecisionContext, vec_str: &str) -> Vec { - vec_str - .split(',') - .map(|s| s.trim()) - .filter(|s| !s.is_empty()) - .map(|s| parse_math_value(context, s)) - .filter_map(|r| r.ok()) - .collect() -} - -/// Parse a single mathematical expression to high-precision Float -fn parse_math_value(context: &HighPrecisionContext, s: &str) -> Result { - let s_lower = s.to_lowercase(); - - // Handle ln() function calls - if s_lower.starts_with("ln(") && s_lower.ends_with(')') { - let inner = &s[3..s.len()-1]; - let val = parse_simple_float(context, inner)?; - return Ok(val.ln()); - } - - // Handle phi^n notation - if s_lower.starts_with("phi") { - if s_lower.len() > 4 && &s[4..5] == '^' { - let exp: i32 = s[5..].parse().map_err(|e| format!("Invalid phi exponent: {}", e))?; - return Ok(hp_phi_pow(context, exp)); - } - return Ok(hp_phi_pow(context, 0)); - } - - // Handle pi^n notation - if s_lower.starts_with("pi") { - if s_lower.len() > 2 && &s[2..3] == '^' { - let exp: i32 = s[3..].parse().map_err(|e| format!("Invalid pi exponent: {}", e))?; - return Ok(hp_pi_pow(context, exp)); - } - return Ok(hp_pi_pow(context, 0)); - } - - // Simple number - parse_simple_float(context, s) -} - -/// Parse a simple floating-point number -fn parse_simple_float(context: &HighPrecisionContext, s: &str) -> Result { - rug::Float::parse(context.precision, s) - .map_err(|e| format!("Failed to parse '{}': {}", s, e)) -} - -/// Build relation matrix for PSLQ algorithm -/// -/// Matrix M has columns for each basis vector b_i: -/// M = [b_1 | b_2 | ... | b_k] -fn build_relation_matrix(vectors: &[Vec], max_coeff: i32) -> Vec> { - let k = vectors.len(); - let m = max_coeff as usize; - - // For each coefficient position, build rational approximation - let mut matrix = Vec::with_capacity(m); - - for i in 0..m { - let mut row = Vec::with_capacity(k); - for j in 0..k { - // Round vector[j] to rational with denominator 2^(m-i-1) - let denominator = Integer::from(2).pow(i32::try_from(m - i - 1).unwrap_or(0)); - let rounded = Rational::from(&vectors[j][i], denominator); - row.push(rounded); - } - matrix.push(row); - } - - matrix -} - -/// LLL lattice basis reduction (simplified) -/// -/// This is a simplified version using Gram-Schmidt orthogonalization -fn lll_reduction(matrix: &mut Vec>) -> Option> { - let m = matrix.len(); - if m == 0 { - return None; - } - let k = matrix[0].len(); - - // Simplified reduction: find smallest norm vector - // Full LLL is more complex; this gives a good heuristic - let mut best_idx = 0; - let mut best_norm = Rational::from(1_000_000_i64); - - for i in 0..m { - let mut sum_sq = Rational::from(0); - for j in 0..k { - let val = &matrix[i][j]; - sum_sq += val * val; - } - // Simplified norm (actual LLL uses Gram matrix) - if sum_sq < best_norm { - best_norm = sum_sq; - best_idx = i; - } - } - - Some(matrix[best_idx].clone()) -} - -/// Run PSLQ algorithm to find integer relations -fn run_pslq( - vectors: Vec, - max_coeff: i32, - tolerance: f64, -) -> Option { - let k = vectors.len(); - if k == 0 { - return None; - } - - // Build relation matrix - let mut matrix = build_relation_matrix(&vectors, max_coeff); - - // Apply LLL reduction - if let Some(coefficients) = lll_reduction(&mut matrix) { - // Verify relation - let residual_norm = verify_relation(&vectors, &coefficients); - - Some(PslqRelation { - coefficients: coefficients.iter().map(|r| r.to_integer().unwrap_or(0)).collect(), - norm: norm_f64(&coefficients), - residual_norm: residual_norm, - quality: compute_quality(&coefficients, residual_norm), - }) - } else { - None - } -} - -/// Verify that the coefficients produce zero (within tolerance) -fn verify_relation(vectors: &[Vec], coeffs: &[Rational]) -> f64 { - let k = vectors.len(); - let mut result = rug::Float::with_val(DEFAULT_PRECISION, 0.0); - - for i in 0..k { - let term = rug::Float::with_val(DEFAULT_PRECISION, 0.0); - // term = coefficient * vector[i] - let coeff_val = rug::Float::with_val(DEFAULT_PRECISION, coeffs[i].to_f64()); - let vector_val = &vectors[i]; - term.assign(&coeff_val * vector_val); - result += term; - } - - result.to_f64().abs() -} - -/// Compute L2 norm of a rational vector -fn norm_f64(coeffs: &[Rational]) -> f64 { - let mut sum = 0.0_f64; - for c in coeffs { - let val = c.to_f64(); - sum += val * val; - } - sum.sqrt() -} - -/// Compute quality score (lower is better) -fn compute_quality(coeffs: &[Rational], residual: f64) -> f64 { - let norm = norm_f64(coeffs); - let max_coeff = coeffs.iter().map(|c| c.abs().to_f64()).fold(0.0_f64, f64::max); - // Quality: small residual + small coefficients - residual + 0.01 * norm + 0.001 * max_coeff -} - -/// Get Trinity basis vectors -fn get_trinity_basis() -> Vec { - vec![ - "ln(3)".to_string(), - "ln(pi)".to_string(), - "ln(phi)".to_string(), - "1".to_string(), - ] -} - -/// Check ZIP-derived formulas (Phase K) -fn get_zip_basis() -> Vec { - vec![ - "sin2_theta12_zip".to_string(), // ZIP formula: n=14,k=0,m=0,p=0,q=0,r=0 - "sin2_theta23_zip".to_string(), // ZIP formula: n=14,k=0,m=0,p=1,q=0,r=0 - "sin2_theta13_zip".to_string(), // ZIP formula: n=14,k=0,m=0,p=0,q=0,r=0 - ] -} - -/// Get PDG target constants for verification -fn get_pdg_targets() -> Vec<(&'static str, f64)> { - vec![ - ("alpha_inv", 137.036), // α⁻¹ from Pellis expansion - ("sin2_theta12", 0.307), // JUNO 2025 measured - ("sin2_theta23", 0.546), // PDG 2024 - ("sin2_theta13", 0.0222), // PDG 2024 - ("delta_cp", 197.0), // Trinity prediction - ("alpha_s", 0.118034), // Trinity Sacred Formula - ] -} - -/// Run verification check on Sacred Formula catalog -fn run_verification(catalog_path: Option, max_error: f64) -> anyhow::Result<()> { - use std::collections::HashMap; - - // Load catalog - let path = catalog_path.unwrap_or_else(|| "research/sacred_formula_catalog.json".to_string()); - - let catalog_content = std::fs::read_to_string(&path) - .with_context(|| format!("Failed to read catalog: {}", path))?; - - let catalog: serde_json::from_str(&catalog_content) - .with_context(|| format!("Failed to parse catalog JSON"))?; - - let context = HighPrecisionContext::new(DEFAULT_PRECISION); - - // Collect values to test - let mut test_values: HashMap = HashMap::new(); - let trinity_basis = get_trinity_basis(); - - // Add basis vectors - for basis_name in &trinity_basis { - if let Ok(val) = parse_math_value(&context, basis_name) { - test_values.insert(basis_name.clone(), val); - } - } - - // Add catalog formulas - let mut vectors = Vec::new(); - vectors.push(trinity_basis.clone()); - - // Check each entry in catalog - let mut verified = 0; - let mut candidates = 0; - - for entry in catalog.as_array().unwrap_or(&serde_json::json!([])).iter() { - let name = entry["name"].as_str().unwrap_or(""); - let n = entry["n"].as_i64().unwrap_or(1) as i32; - let k = entry["k"].as_i64().unwrap_or(0) as i32; - let m = entry["m"].as_i64().unwrap_or(0) as i32; - let p = entry["p"].as_i64().unwrap_or(0) as i32; - let q = entry["q"].as_i64().unwrap_or(0) as i32; - let r = entry["r"].as_i64().unwrap_or(0) as i32; - - // Compute Sacred Formula value: V = n * 3^m * π^k * φ^p * e^q * 1^r - let ln_3 = hp_ln_3(&context); - let ln_pi = hp_pi_pow(&context, k).ln(); - let ln_phi = hp_phi_pow(&context, p).ln(); - let ln_e = rug::Float::with_val(context.precision, std::f64::consts::E).ln(); - - let value = rug::Float::with_val(context.precision, n as f64) - * rug::Float::with_val(context.precision, 3).pow(m as i32) - * hp_pi_pow(&context, k) - * hp_phi_pow(&context, p) - * ln_e.pow(q as i32); - - let value_f64 = value.to_f64(); - - // Check for relation with integer coefficients - let test_vector = vec![ - ln_3.clone(), - ln_pi.clone(), - ln_phi.clone(), - rug::Float::with_val(context.precision, 1.0), - ]; - - if let Some(relation) = run_pslq(test_vector, 20, DEFAULT_TOLERANCE) { - let error_pct = if value_f64.abs() > 1e-10 { - relation.residual_norm / value_f64.abs() * 100.0 - } else { - relation.residual_norm * 100.0 - }; - - if error_pct < max_error { - let coeff_str = relation.coefficients.iter() - .enumerate() - .map(|(i, c)| format!("{}*{}", if i == 3 { format!("{}", c) } else { format!("{}x_{}", i, c) })) - .collect::>() - .join(" + "); - println!("✓ {} : V = {:.12} : PSLQ = {{{}}} : Δ = {:.6}%", - name, value_f64, - coeff_str, - error_pct - ); - verified += 1; - } else { - candidates += 1; - } - } - } - - println!("\n=== PSLQ Verification Summary ==="); - println!("Total entries: {}", verified + candidates); - println!("Verified (Δ < {}%): {}", max_error, verified); - println!("Candidates (Δ >= {}%): {}", max_error, candidates); - - Ok(()) -} - -/// Run vector analysis -fn run_vector(vector_str: String, target_str: Option, max_coeff: i32, precision: u32, tolerance: Option) -> anyhow::Result<()> { - let context = HighPrecisionContext::new(precision); - let vectors = parse_vector(&context, &vector_str); - - if vectors.is_empty() { - anyhow::bail!("No valid values provided"); - } - - println!("=== PSLQ: Integer Relation Finder ==="); - println!("Vector: [{}]", vector_str); - println!("Precision: {} bits, Max coeff: |a_i| ≤ {}", precision, max_coeff); - - if let Some(relation) = run_pslq(vectors, max_coeff, tolerance.unwrap_or_else(|| format!("{}", DEFAULT_TOLERANCE)).parse().unwrap_or(DEFAULT_TOLERANCE)) { - println!("\nFound relation:"); - for (i, coeff) in relation.coefficients.iter().enumerate() { - println!(" a_{} = {}", i + 1, coeff); - } - println!("Norm: {:.6}", relation.norm); - println!("Residual |L₂|: {:.2e}", relation.residual_norm); - println!("Quality score: {:.4}", relation.quality); - - // Check against target if provided - if let Some(target_val_str) = target_str { - if let Ok(target_val) = rug::Float::parse(precision, target_val_str) { - println!("\nTarget value: {}", target_val_str); - - // Compute target value from relation - let mut computed = rug::Float::with_val(precision, 0.0); - for (coeff, vector) in relation.coefficients.iter().zip(vectors.iter()) { - let c = rug::Float::with_val(precision, *coeff as f64); - computed += c * vector; - } - println!("Computed from relation: {}", computed); - println!("Difference: {}", (computed - target_val).abs()); - } - } - } else { - println!("\nNo integer relation found within tolerance {}", tolerance.unwrap_or_else(|| format!("{}", DEFAULT_TOLERANCE))); - } - - Ok(()) -} - -/// Run basis check (trinity or extended) -fn run_basis(basis_name: String, check_targets: bool, check_zip: bool) -> anyhow::Result<()> { - let context = HighPrecisionContext::new(DEFAULT_PRECISION); - - let basis = match basis_name.as_str() { - "trinity" => get_trinity_basis(), - "extended" => { - let mut b = get_trinity_basis(); - b.push("ln(2)".to_string()); // Extended basis - b - } - "reduced" => vec![ - "ln(phi)".to_string(), - "ln(pi)".to_string(), - "1".to_string(), - ] - _ => anyhow::bail!("Unknown basis: {}", basis_name), - }; - - let vectors: Vec = basis - .iter() - .filter_map(|s| parse_math_value(&context, s).ok()) - .collect(); - - if vectors.is_empty() { - anyhow::bail!("No valid basis vectors"); - } - - println!("=== PSLQ: Basis Check ==="); - println!("Basis: {} ({} vectors)", basis_name, basis.len()); - println!("Vectors: {:?}", basis); - - if let Some(relation) = run_pslq(vectors, DEFAULT_MAX_COEFF, DEFAULT_TOLERANCE) { - println!("\nFound relation:"); - for (i, coeff) in relation.coefficients.iter().enumerate() { - println!(" a_{} = {}", i + 1, coeff); - } - println!("Norm: {:.6}", relation.norm); - println!("Residual |L₂|: {:.2e}", relation.residual_norm); - println!("Quality score: {:.4}", relation.quality); - - // This proves linear dependence - println!("\n>>> Linear dependence detected!"); - println!(" Basis vectors are NOT linearly independent."); - println!(" Relation: sum_i(a_i * b_i) = 0"); - } else { - println!("\nNo relation found: basis vectors appear linearly independent"); - } - - // Check against targets if requested - if check_targets { - println!("\n=== Checking against PDG targets ==="); - for (target_name, target_val) in get_pdg_targets() { - println!("Target: {} = {}", target_name, target_val); - } - } - - // Check ZIP formulas if requested - if check_zip { - println!("\n=== Checking ZIP-derived formulas ==="); - let zip_basis = get_zip_basis(); - let zip_vectors: Vec = zip_basis - .iter() - .filter_map(|s| parse_math_value(&context, s).ok()) - .collect(); - - if !zip_vectors.is_empty() { - if let Some(relation) = run_pslq(zip_vectors, DEFAULT_MAX_COEFF, DEFAULT_TOLERANCE) { - println!("\nZIP formulas relation:"); - println!("{:?}", relation.coefficients); - } else { - println!("ZIP formulas appear independent"); - } - } - } - - Ok(()) -} - -pub fn run_pslq_command( - cmd: PslqCommands, - _repo_root: &std::path::Path, -) -> anyhow::Result<()> { - match cmd { - PslqCommands::Vector { vector, target, max_coeff, precision, tolerance } => { - run_vector(vector, target, max_coeff, precision, tolerance)?; - } - PslqCommands::Basis { basis, check_targets, check_zip } => { - run_basis(basis, check_targets, check_zip)?; - } - PslqCommands::Verify { catalog, max_error } => { - run_verification(catalog, max_error)?; - } - } -} diff --git a/docs/reports/orphan_modules.json b/docs/reports/orphan_modules.json index bf4cde3c5b..0156bf30ac 100644 --- a/docs/reports/orphan_modules.json +++ b/docs/reports/orphan_modules.json @@ -2,9 +2,9 @@ "_what": "Per-crate ceiling on .rs files under src/ that no `mod` declaration reaches from the crate root. Measured by `tri mods orphan --gate`.", "_why": "#2427, a Zig-lexer PR, removed `mod elab;` from cli/tri/src/main.rs in the same hunk as two other lines and left the 319-line file. cargo build cannot error on a file it does not compile, the suite fell 358 -> 354, and no gate reads the test count. Nothing printed a word for six days (#2900).", "_rule": "Exact match, not an upper bound. A count BELOW its ceiling fails too, so the next orphan cannot hide in the slack -- lower the number in the commit that connects a file.", - "_bootstrap_6": "Pre-existing, not repaired here: _enrich/mod.rs, math_pslq.rs, neural/attention.rs, neural/hslm.rs, runtime_minimal.rs, runtime_minimal_test.rs. tooling.rs (7 unconnected gate bodies, #2905) was deleted by #10165 to pay the foreign-line budget; the rest are unexamined. Connecting any of them is a decision about what goes red, not a repair.", + "_bootstrap_6": "Pre-existing, not repaired here: _enrich/mod.rs, neural/attention.rs, neural/hslm.rs, runtime_minimal.rs, runtime_minimal_test.rs. tooling.rs (7 unconnected gate bodies, #2905) was deleted by #10165 and math_pslq.rs by #10404, each to pay the foreign-line budget; the rest are unexamined. Connecting any of them is a decision about what goes red, not a repair.", "ceilings": { - "bootstrap": 6, + "bootstrap": 5, "cli/tri": 0, "cli/trios-bridge": 0, "cli/flash-spi": 0, From 87ea11f5e7e8413575135bdd1041acf53c894df0 Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Sun, 18 Oct 2026 18:27:18 +0700 Subject: [PATCH 05/11] seal: igla-coder-ternary-gemm-ptx, 21/21 tests, 0 vacuous (Refs #10404) Co-Authored-By: Claude Opus 5.5 --- .../coder_igla-coder-ternary-gemm-ptx.json | 28 +++++++++++++++++++ 1 file changed, 28 insertions(+) create mode 100644 .trinity/seals/coder_igla-coder-ternary-gemm-ptx.json diff --git a/.trinity/seals/coder_igla-coder-ternary-gemm-ptx.json b/.trinity/seals/coder_igla-coder-ternary-gemm-ptx.json new file mode 100644 index 0000000000..5291598a4b --- /dev/null +++ b/.trinity/seals/coder_igla-coder-ternary-gemm-ptx.json @@ -0,0 +1,28 @@ +{ + "built_by": "t27c-bootstrap@0.5.2+2406a4804", + "closure": [], + "config": "gen=c,rust,verilog,zig;tests=zig", + "gen_hash_c": "sha256:c6e3f013bea9fc3a9cb29847dc6afca703ca48be4120bc0f0d37513e7142eeb4", + "gen_hash_rust": "sha256:5a0b5af78cfd53229c272a35cb6fbbcc161345f01ed527f7701969be2e74354b", + "gen_hash_verilog": "sha256:c62893a10c45ed25c7145fb54d1a3e3ed5fea21041e5dd5fd70a4b1db301d850", + "gen_hash_zig": "sha256:5b05f3a643b9adc744f7595c518e306489936858478f86a98abbf9000720302a", + "module": "igla-coder-ternary-gemm-ptx", + "ring": 12, + "seal_schema": 2, + "sealed_at": "2026-10-18T11:26:31Z", + "sealed_by": "t27c-bootstrap@0.5.2", + "spec_hash": "sha256:c032dd4bfe08352725bc8eefa687021033d037a975931673e189dcc4d04c19f6", + "spec_path": "specs/igla/coder/ternary_gemm_ptx.t27", + "tests": { + "failed": 0, + "failing": [], + "forced": false, + "passed": 21, + "total": 21, + "vacuous": 0 + }, + "toolchain": { + "test_runner": "sha256:1982ae948e24d00c5d221ccec12dfce1eb01204b53e37b0f857f93b16d64629d", + "zig": "0.16.0" + } +} \ No newline at end of file From bd2f63231e5bf31384817238088d721799b1066e Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Sun, 18 Oct 2026 18:35:08 +0700 Subject: [PATCH 06/11] igla ptx: one writer call per statement, so gen-rust output borrows out once; driver: split the truncate borrow (Refs #10404) Co-Authored-By: Claude Opus 5.5 --- bootstrap/src/host/driver.rs | 2 +- specs/igla/coder/ternary_gemm_ptx.t27 | 61 +++++++++++++++++++-------- 2 files changed, 45 insertions(+), 18 deletions(-) diff --git a/bootstrap/src/host/driver.rs b/bootstrap/src/host/driver.rs index 6b8d07f9b1..6632a5ac15 100644 --- a/bootstrap/src/host/driver.rs +++ b/bootstrap/src/host/driver.rs @@ -315,7 +315,7 @@ pub fn ptx_run(m: usize, n: usize, k: usize, iters: u32, stride: usize) -> anyho anyhow::ensure!(t::shape_ok(m, n, k), "m={} n={} k={}: refused by shape_ok in ternary_gemm_ptx.t27", m, n, k); let ck = |r: i32| if r == 0 { Ok(()) } else { Err(anyhow::anyhow!("CUDA driver error {}", r)) }; let mut ptx = vec![0u8; t::PTX_CAP + 1]; - ptx.truncate(t::emit_ptx(&mut ptx, t::PTX_CAP)); ptx.push(0); + let len = t::emit_ptx(&mut ptx, t::PTX_CAP); ptx.truncate(len); ptx.push(0); let (mut a, mut wp, mut c, mut acc) = (vec![0i8; m * k], vec![0u8; k * n / 4], vec![0i32; m * n], vec![0i32; n]); t::gen_wp(&mut wp, k * n, t::gen_a(&mut a, m * k, 1)); let (mut dev, mut ctx, mut module, mut f, mut da, mut dw, mut dc, mut pm, mut pn, mut pk) = (0i32, 0usize, 0usize, 0usize, 0u64, 0u64, 0u64, m as u32, n as u32, k as u32); diff --git a/specs/igla/coder/ternary_gemm_ptx.t27 b/specs/igla/coder/ternary_gemm_ptx.t27 index 2dc9ea15bd..9df0bb900a 100644 --- a/specs/igla/coder/ternary_gemm_ptx.t27 +++ b/specs/igla/coder/ternary_gemm_ptx.t27 @@ -98,8 +98,10 @@ fn put_u(out : []u8, cap : usize, w : usize, v : u32) -> usize { /// One indented instruction line that has no operand to fill in. fn ins(out : []u8, cap : usize, w : usize, s : str) -> usize { - const x : usize = put_s(out, cap, w, " "); - return put_s(out, cap, put_s(out, cap, x, s), "\n"); + // One writer call per statement: a nested call would borrow `out` twice in the Rust that gen-rust writes. + var x : usize = put_s(out, cap, w, " "); + x = put_s(out, cap, x, s); + return put_s(out, cap, x, "\n"); } /// "{%name, %name, %name, %name}" @@ -108,7 +110,8 @@ fn put_quad(out : []u8, cap : usize, w : usize, name : str, first : u32) -> usiz var j : u32 = 0; while (j < 4) { if (j > 0) { x = put_s(out, cap, x, ", "); } - x = put_u(out, cap, put_s(out, cap, x, name), first + j); + x = put_s(out, cap, x, name); + x = put_u(out, cap, x, first + j); j += 1; } return put_s(out, cap, x, "}"); @@ -198,7 +201,9 @@ fn emit_prologue(out : []u8, cap : usize, w : usize) -> usize { x = ins(out, cap, x, "mov.u32 %r28, 0;"); var i : u32 = 0; while (i < 64) { - x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " mov.b32 %acc"), i), ", 0;\n"); + x = put_s(out, cap, x, " mov.b32 %acc"); + x = put_u(out, cap, x, i); + x = put_s(out, cap, x, ", 0;\n"); i += 1; } return x; @@ -228,12 +233,16 @@ fn emit_tile_loads(out : []u8, cap : usize, w : usize) -> usize { x = ins(out, cap, x, "@%p5 ld.global.u32 %r35, [%rd9];"); var j : u32 = 0; while (j < 16) { - x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " shr.u32 %r36, %r35, "), 2 * j), ";\n"); + x = put_s(out, cap, x, " shr.u32 %r36, %r35, "); + x = put_u(out, cap, x, 2 * j); + x = put_s(out, cap, x, ";\n"); x = ins(out, cap, x, "and.b32 %r37, %r36, 1;"); x = ins(out, cap, x, "shr.u32 %r38, %r36, 1;"); x = ins(out, cap, x, "and.b32 %r38, %r38, 1;"); x = ins(out, cap, x, "sub.s32 %r37, %r37, %r38;"); - x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " st.shared.u8 [%r16+"), j * SMEM_STRIDE), "], %r37;\n"); + x = put_s(out, cap, x, " st.shared.u8 [%r16+"); + x = put_u(out, cap, x, j * SMEM_STRIDE); + x = put_s(out, cap, x, "], %r37;\n"); j += 1; } return x; @@ -241,9 +250,14 @@ fn emit_tile_loads(out : []u8, cap : usize, w : usize) -> usize { /// " ld.shared.u32 %, [%+];" fn emit_lds(out : []u8, cap : usize, w : usize, reg : str, i : u32, base : str, off : u32) -> usize { - var x : usize = put_u(out, cap, put_s(out, cap, put_s(out, cap, w, " ld.shared.u32 "), reg), i); - x = put_s(out, cap, put_s(out, cap, put_s(out, cap, x, ", ["), base), "+"); - return put_s(out, cap, put_u(out, cap, x, off), "];\n"); + var x : usize = put_s(out, cap, w, " ld.shared.u32 "); + x = put_s(out, cap, x, reg); + x = put_u(out, cap, x, i); + x = put_s(out, cap, x, ", ["); + x = put_s(out, cap, x, base); + x = put_s(out, cap, x, "+"); + x = put_u(out, cap, x, off); + return put_s(out, cap, x, "];\n"); } /// The fragment loads and the 16 mma.sync of one k-step. @@ -270,10 +284,16 @@ fn emit_mma_step(out : []u8, cap : usize, w : usize) -> usize { while (ni < 4) { const b : u32 = (mi * 4 + ni) * 4; x = put_s(out, cap, x, " mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 "); - x = put_s(out, cap, put_quad(out, cap, x, "%acc", b), ", "); - x = put_s(out, cap, put_quad(out, cap, x, "%fa", mi * 4), ", {%fb"); - x = put_u(out, cap, put_s(out, cap, put_u(out, cap, x, ni * 2), ", %fb"), ni * 2 + 1); - x = put_s(out, cap, put_quad(out, cap, put_s(out, cap, x, "}, "), "%acc", b), ";\n"); + x = put_quad(out, cap, x, "%acc", b); + x = put_s(out, cap, x, ", "); + x = put_quad(out, cap, x, "%fa", mi * 4); + x = put_s(out, cap, x, ", {%fb"); + x = put_u(out, cap, x, ni * 2); + x = put_s(out, cap, x, ", %fb"); + x = put_u(out, cap, x, ni * 2 + 1); + x = put_s(out, cap, x, "}, "); + x = put_quad(out, cap, x, "%acc", b); + x = put_s(out, cap, x, ";\n"); ni += 1; } mi += 1; @@ -283,8 +303,12 @@ fn emit_mma_step(out : []u8, cap : usize, w : usize) -> usize { /// One guarded pair store: accumulators b, b+1 to C[row0 + dr][col0 + dc .. +1]. fn emit_store(out : []u8, cap : usize, w : usize, dr : u32, dc : u32, b : u32) -> usize { - var x : usize = put_s(out, cap, put_u(out, cap, put_s(out, cap, w, " add.u32 %r41, %r39, "), dr), ";\n"); - x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " add.u32 %r42, %r40, "), dc), ";\n"); + var x : usize = put_s(out, cap, w, " add.u32 %r41, %r39, "); + x = put_u(out, cap, x, dr); + x = put_s(out, cap, x, ";\n"); + x = put_s(out, cap, x, " add.u32 %r42, %r40, "); + x = put_u(out, cap, x, dc); + x = put_s(out, cap, x, ";\n"); x = ins(out, cap, x, "setp.lt.u32 %p6, %r41, %r1;"); x = ins(out, cap, x, "setp.lt.u32 %p7, %r42, %r2;"); x = ins(out, cap, x, "and.pred %p6, %p6, %p7;"); @@ -292,8 +316,11 @@ fn emit_store(out : []u8, cap : usize, w : usize, dr : u32, dc : u32, b : u32) - x = ins(out, cap, x, "mad.wide.u32 %rd10, %r41, %r2, %rd10;"); x = ins(out, cap, x, "shl.b64 %rd10, %rd10, 2;"); x = ins(out, cap, x, "add.u64 %rd10, %rd10, %rd3;"); - x = put_u(out, cap, put_s(out, cap, x, " @%p6 st.global.v2.u32 [%rd10], {%acc"), b); - return put_s(out, cap, put_u(out, cap, put_s(out, cap, x, ", %acc"), b + 1), "};\n"); + x = put_s(out, cap, x, " @%p6 st.global.v2.u32 [%rd10], {%acc"); + x = put_u(out, cap, x, b); + x = put_s(out, cap, x, ", %acc"); + x = put_u(out, cap, x, b + 1); + return put_s(out, cap, x, "};\n"); } /// Row base r39 = 128 by + 64 wm + g and column base r40 = 128 bx + 32 wn + 2t, then 32 stores. From c9a20a08c0a86022e47dad49d17623521ad71d75 Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Sun, 18 Oct 2026 18:39:29 +0700 Subject: [PATCH 07/11] gen + seal: igla-coder-ternary-gemm-ptx regenerated after the writer split, 21/21 (Refs #10404) Co-Authored-By: Claude Opus 5.5 --- .../coder_igla-coder-ternary-gemm-ptx.json | 12 ++-- .../gen/rust/igla/coder/ternary_gemm_ptx.rs | 60 +++++++++++++------ 2 files changed, 49 insertions(+), 23 deletions(-) diff --git a/.trinity/seals/coder_igla-coder-ternary-gemm-ptx.json b/.trinity/seals/coder_igla-coder-ternary-gemm-ptx.json index 5291598a4b..d6c0212a91 100644 --- a/.trinity/seals/coder_igla-coder-ternary-gemm-ptx.json +++ b/.trinity/seals/coder_igla-coder-ternary-gemm-ptx.json @@ -2,16 +2,16 @@ "built_by": "t27c-bootstrap@0.5.2+2406a4804", "closure": [], "config": "gen=c,rust,verilog,zig;tests=zig", - "gen_hash_c": "sha256:c6e3f013bea9fc3a9cb29847dc6afca703ca48be4120bc0f0d37513e7142eeb4", - "gen_hash_rust": "sha256:5a0b5af78cfd53229c272a35cb6fbbcc161345f01ed527f7701969be2e74354b", - "gen_hash_verilog": "sha256:c62893a10c45ed25c7145fb54d1a3e3ed5fea21041e5dd5fd70a4b1db301d850", - "gen_hash_zig": "sha256:5b05f3a643b9adc744f7595c518e306489936858478f86a98abbf9000720302a", + "gen_hash_c": "sha256:7ac537a19563bff12ac28e1a1669de534e956d202b1adf4dc30f2e94d9f96b47", + "gen_hash_rust": "sha256:df047da747040c5a33657f13557fee1d3469efc588c9771ab969b612ba14531e", + "gen_hash_verilog": "sha256:8f0b6dd58e6c0094639dbb97e1d308dd21c19936c66c6db3c9043ed17c03c3aa", + "gen_hash_zig": "sha256:fd96535350bba53d44225aa26dba60a9dcf420d3eb3eaf63565a07864cf688fb", "module": "igla-coder-ternary-gemm-ptx", "ring": 12, "seal_schema": 2, - "sealed_at": "2026-10-18T11:26:31Z", + "sealed_at": "2026-10-18T11:38:54Z", "sealed_by": "t27c-bootstrap@0.5.2", - "spec_hash": "sha256:c032dd4bfe08352725bc8eefa687021033d037a975931673e189dcc4d04c19f6", + "spec_hash": "sha256:2ead1690d3f7fdbff7c9ac62cc6f98a89bb210bf359dccc225d63f236ba03c98", "spec_path": "specs/igla/coder/ternary_gemm_ptx.t27", "tests": { "failed": 0, diff --git a/bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs b/bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs index edddd86a72..b57d529427 100644 --- a/bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs +++ b/bootstrap/gen/rust/igla/coder/ternary_gemm_ptx.rs @@ -57,8 +57,9 @@ pub fn put_u(out: &mut [u8], cap: usize, w: usize, v: u32) -> usize { } pub fn ins(out: &mut [u8], cap: usize, w: usize, s: &'static str) -> usize { - let x: usize = put_s(out, cap, w, " "); - return put_s(out, cap, put_s(out, cap, x, s), "\n"); + let mut x: usize = put_s(out, cap, w, " "); + x = put_s(out, cap, x, s); + return put_s(out, cap, x, "\n"); } pub fn put_quad(out: &mut [u8], cap: usize, w: usize, name: &'static str, first: u32) -> usize { @@ -68,7 +69,8 @@ pub fn put_quad(out: &mut [u8], cap: usize, w: usize, name: &'static str, first: if (j > 0) { x = put_s(out, cap, x, ", "); } - x = put_u(out, cap, put_s(out, cap, x, name), (first + j)); + x = put_s(out, cap, x, name); + x = put_u(out, cap, x, (first + j)); j += 1; } return put_s(out, cap, x, "}"); @@ -149,7 +151,9 @@ pub fn emit_prologue(out: &mut [u8], cap: usize, w: usize) -> usize { x = ins(out, cap, x, "mov.u32 %r28, 0;"); let mut i: u32 = 0; while (i < 64) { - x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " mov.b32 %acc"), i), ", 0;\n"); + x = put_s(out, cap, x, " mov.b32 %acc"); + x = put_u(out, cap, x, i); + x = put_s(out, cap, x, ", 0;\n"); i += 1; } return x; @@ -178,21 +182,30 @@ pub fn emit_tile_loads(out: &mut [u8], cap: usize, w: usize) -> usize { x = ins(out, cap, x, "@%p5 ld.global.u32 %r35, [%rd9];"); let mut j: u32 = 0; while (j < 16) { - x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " shr.u32 %r36, %r35, "), (2 * j)), ";\n"); + x = put_s(out, cap, x, " shr.u32 %r36, %r35, "); + x = put_u(out, cap, x, (2 * j)); + x = put_s(out, cap, x, ";\n"); x = ins(out, cap, x, "and.b32 %r37, %r36, 1;"); x = ins(out, cap, x, "shr.u32 %r38, %r36, 1;"); x = ins(out, cap, x, "and.b32 %r38, %r38, 1;"); x = ins(out, cap, x, "sub.s32 %r37, %r37, %r38;"); - x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " st.shared.u8 [%r16+"), (j * SMEM_STRIDE)), "], %r37;\n"); + x = put_s(out, cap, x, " st.shared.u8 [%r16+"); + x = put_u(out, cap, x, (j * SMEM_STRIDE)); + x = put_s(out, cap, x, "], %r37;\n"); j += 1; } return x; } pub fn emit_lds(out: &mut [u8], cap: usize, w: usize, reg: &'static str, i: u32, base: &'static str, off: u32) -> usize { - let mut x: usize = put_u(out, cap, put_s(out, cap, put_s(out, cap, w, " ld.shared.u32 "), reg), i); - x = put_s(out, cap, put_s(out, cap, put_s(out, cap, x, ", ["), base), "+"); - return put_s(out, cap, put_u(out, cap, x, off), "];\n"); + let mut x: usize = put_s(out, cap, w, " ld.shared.u32 "); + x = put_s(out, cap, x, reg); + x = put_u(out, cap, x, i); + x = put_s(out, cap, x, ", ["); + x = put_s(out, cap, x, base); + x = put_s(out, cap, x, "+"); + x = put_u(out, cap, x, off); + return put_s(out, cap, x, "];\n"); } pub fn emit_mma_step(out: &mut [u8], cap: usize, w: usize) -> usize { @@ -218,10 +231,16 @@ pub fn emit_mma_step(out: &mut [u8], cap: usize, w: usize) -> usize { while (ni < 4) { let b: u32 = (((mi * 4) + ni) * 4); x = put_s(out, cap, x, " mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 "); - x = put_s(out, cap, put_quad(out, cap, x, "%acc", b), ", "); - x = put_s(out, cap, put_quad(out, cap, x, "%fa", (mi * 4)), ", {%fb"); - x = put_u(out, cap, put_s(out, cap, put_u(out, cap, x, (ni * 2)), ", %fb"), ((ni * 2) + 1)); - x = put_s(out, cap, put_quad(out, cap, put_s(out, cap, x, "}, "), "%acc", b), ";\n"); + x = put_quad(out, cap, x, "%acc", b); + x = put_s(out, cap, x, ", "); + x = put_quad(out, cap, x, "%fa", (mi * 4)); + x = put_s(out, cap, x, ", {%fb"); + x = put_u(out, cap, x, (ni * 2)); + x = put_s(out, cap, x, ", %fb"); + x = put_u(out, cap, x, ((ni * 2) + 1)); + x = put_s(out, cap, x, "}, "); + x = put_quad(out, cap, x, "%acc", b); + x = put_s(out, cap, x, ";\n"); ni += 1; } mi += 1; @@ -230,8 +249,12 @@ pub fn emit_mma_step(out: &mut [u8], cap: usize, w: usize) -> usize { } pub fn emit_store(out: &mut [u8], cap: usize, w: usize, dr: u32, dc: u32, b: u32) -> usize { - let mut x: usize = put_s(out, cap, put_u(out, cap, put_s(out, cap, w, " add.u32 %r41, %r39, "), dr), ";\n"); - x = put_s(out, cap, put_u(out, cap, put_s(out, cap, x, " add.u32 %r42, %r40, "), dc), ";\n"); + let mut x: usize = put_s(out, cap, w, " add.u32 %r41, %r39, "); + x = put_u(out, cap, x, dr); + x = put_s(out, cap, x, ";\n"); + x = put_s(out, cap, x, " add.u32 %r42, %r40, "); + x = put_u(out, cap, x, dc); + x = put_s(out, cap, x, ";\n"); x = ins(out, cap, x, "setp.lt.u32 %p6, %r41, %r1;"); x = ins(out, cap, x, "setp.lt.u32 %p7, %r42, %r2;"); x = ins(out, cap, x, "and.pred %p6, %p6, %p7;"); @@ -239,8 +262,11 @@ pub fn emit_store(out: &mut [u8], cap: usize, w: usize, dr: u32, dc: u32, b: u32 x = ins(out, cap, x, "mad.wide.u32 %rd10, %r41, %r2, %rd10;"); x = ins(out, cap, x, "shl.b64 %rd10, %rd10, 2;"); x = ins(out, cap, x, "add.u64 %rd10, %rd10, %rd3;"); - x = put_u(out, cap, put_s(out, cap, x, " @%p6 st.global.v2.u32 [%rd10], {%acc"), b); - return put_s(out, cap, put_u(out, cap, put_s(out, cap, x, ", %acc"), (b + 1)), "};\n"); + x = put_s(out, cap, x, " @%p6 st.global.v2.u32 [%rd10], {%acc"); + x = put_u(out, cap, x, b); + x = put_s(out, cap, x, ", %acc"); + x = put_u(out, cap, x, (b + 1)); + return put_s(out, cap, x, "};\n"); } pub fn emit_epilogue(out: &mut [u8], cap: usize, w: usize) -> usize { From 39095f448a6faaa50287858b34580b5b91a0da7b Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Sun, 18 Oct 2026 21:51:19 +0700 Subject: [PATCH 08/11] t27b ledger: list specs/igla/coder/ternary_gemm_ptx.t27 as pass, the one spec this PR adds (Refs #10404) Co-Authored-By: Claude Opus 5.5 --- docs/reports/t27b_expectations.json | 1 + 1 file changed, 1 insertion(+) diff --git a/docs/reports/t27b_expectations.json b/docs/reports/t27b_expectations.json index 8d982e2ade..e5e1913512 100644 --- a/docs/reports/t27b_expectations.json +++ b/docs/reports/t27b_expectations.json @@ -438,6 +438,7 @@ {"path": "specs/igla/coder/pipeline.t27", "t27b": "blocked", "blocker": "ConstDecl(slice)", "reason": "unimplemented"}, {"path": "specs/igla/coder/prm.t27", "t27b": "pass"}, {"path": "specs/igla/coder/ternary_gemm_gpu.t27", "t27b": "pass"}, + {"path": "specs/igla/coder/ternary_gemm_ptx.t27", "t27b": "pass"}, {"path": "specs/igla/coder/tokenizer.t27", "t27b": "pass"}, {"path": "specs/igla/coder/training.t27", "t27b": "blocked", "blocker": "ExprCall(undeclared fn)", "reason": "unimplemented"}, {"path": "specs/igla/coder/weights.t27", "t27b": "blocked", "blocker": "ConstDecl(slice)", "reason": "unimplemented"}, From b082cc9bf40998efe03b40e9900f3ff48722346f Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Mon, 19 Oct 2026 00:31:34 +0700 Subject: [PATCH 09/11] t27b ledger: list the four specs/analog specs master added unlisted, with the verdicts CI measured (Refs #10404) bench, linear and mos_table pass; linear_zoh is blocked on ExprCall(@sin). Co-Authored-By: Claude Opus 5.5 --- docs/reports/t27b_expectations.json | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/docs/reports/t27b_expectations.json b/docs/reports/t27b_expectations.json index fb1dfbbef3..232c7139bd 100644 --- a/docs/reports/t27b_expectations.json +++ b/docs/reports/t27b_expectations.json @@ -43,6 +43,10 @@ {"path": "specs/agents/x.t27", "t27b": "pass"}, {"path": "specs/agents/y.t27", "t27b": "pass"}, {"path": "specs/agents/z.t27", "t27b": "pass"}, + {"path": "specs/analog/bench.t27", "t27b": "pass"}, + {"path": "specs/analog/linear.t27", "t27b": "pass"}, + {"path": "specs/analog/linear_zoh.t27", "t27b": "blocked", "blocker": "ExprCall(@sin)", "reason": "unimplemented"}, + {"path": "specs/analog/mos_table.t27", "t27b": "pass"}, {"path": "specs/api/c_abi.t27", "t27b": "pass"}, {"path": "specs/api/tri_api_context.t27", "t27b": "pass"}, {"path": "specs/api/tri_api_loop.t27", "t27b": "pass"}, From 68fe141a24a05d084b9c263b53751b7307d44ad8 Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Mon, 19 Oct 2026 09:16:27 +0700 Subject: [PATCH 10/11] t27b ledger: take master 3a00c389c's ledger, plus this PR's one row for ternary_gemm_ptx.t27 (Refs #10404) Master's ledger now lists the analog, sar, igla coder and t27mcp specs that were unlisted here; the only difference left is this PR's own spec. Co-Authored-By: Claude Opus 5.5 --- docs/reports/t27b_expectations.json | 67 +++++++++++++++++++++++++++-- 1 file changed, 63 insertions(+), 4 deletions(-) diff --git a/docs/reports/t27b_expectations.json b/docs/reports/t27b_expectations.json index 232c7139bd..ca46a8ffdd 100644 --- a/docs/reports/t27b_expectations.json +++ b/docs/reports/t27b_expectations.json @@ -1,10 +1,10 @@ { "schema_version": 1, - "generated_by": "tri t27b ratchet --bless (scripts/tri_loop/t27b.py), #6115", + "generated_by": "t27b ledger steward (specs/jobs/t27b_ledger_steward.t27), #8654", "source": { - "commit": "1ef60ae3486bd83404e398d59fe5b286c1636a19", + "commit": "2520ee1a5985688b0ec2fc5c0d8867d6add1a6a9", "ref": "master", - "finished": null, + "finished": "2026-10-19T01:07:28Z", "asserts_counted": true }, "reasons": { @@ -12,7 +12,7 @@ "reference-bug": "the reference passes it for the wrong reason; t27b is right not to", "n/a": "outside what t27b is for (say why in the PR that sets it)" }, - "max_not_pass": 134, + "max_not_pass": 138, "entries": [ {"path": "specs/a/b_c.t27", "t27b": "pass"}, {"path": "specs/a_b/c.t27", "t27b": "pass"}, @@ -47,6 +47,8 @@ {"path": "specs/analog/linear.t27", "t27b": "pass"}, {"path": "specs/analog/linear_zoh.t27", "t27b": "blocked", "blocker": "ExprCall(@sin)", "reason": "unimplemented"}, {"path": "specs/analog/mos_table.t27", "t27b": "pass"}, + {"path": "specs/analog/sar_cdac.t27", "t27b": "pass"}, + {"path": "specs/analog/sar_comparator.t27", "t27b": "pass"}, {"path": "specs/api/c_abi.t27", "t27b": "pass"}, {"path": "specs/api/tri_api_context.t27", "t27b": "pass"}, {"path": "specs/api/tri_api_loop.t27", "t27b": "pass"}, @@ -65,6 +67,7 @@ {"path": "specs/automation/agent-neuro-kie.t27", "t27b": "pass"}, {"path": "specs/automation/agent-provider-chain.t27", "t27b": "pass"}, {"path": "specs/automation/agent-spec-bank.t27", "t27b": "pass"}, + {"path": "specs/automation/app-tabs.t27", "t27b": "pass"}, {"path": "specs/automation/ball-board.t27", "t27b": "pass"}, {"path": "specs/automation/ball-grants.t27", "t27b": "pass"}, {"path": "specs/automation/blog-share.t27", "t27b": "pass"}, @@ -184,6 +187,8 @@ {"path": "specs/ci/gates/verilog_widths.t27", "t27b": "pass"}, {"path": "specs/ci/schema.t27", "t27b": "pass"}, {"path": "specs/ci/suite_ledger_steward.t27", "t27b": "pass"}, + {"path": "specs/claw/check_latest.t27", "t27b": "pass"}, + {"path": "specs/claw/check_state.t27", "t27b": "pass"}, {"path": "specs/claw/control.t27", "t27b": "pass"}, {"path": "specs/claw/control_pins.t27", "t27b": "pass"}, {"path": "specs/claw/engine.t27", "t27b": "pass"}, @@ -193,19 +198,26 @@ {"path": "specs/claw/gates_pins.t27", "t27b": "pass"}, {"path": "specs/claw/github_surface.t27", "t27b": "pass"}, {"path": "specs/claw/ledger_union.t27", "t27b": "pass"}, + {"path": "specs/claw/merge_lane.t27", "t27b": "pass"}, {"path": "specs/claw/model_route.t27", "t27b": "pass"}, {"path": "specs/claw/preview.t27", "t27b": "pass"}, + {"path": "specs/claw/push_gate.t27", "t27b": "pass"}, + {"path": "specs/claw/queue_cost.t27", "t27b": "pass"}, {"path": "specs/claw/recipes.t27", "t27b": "pass"}, {"path": "specs/claw/refresh.t27", "t27b": "pass"}, + {"path": "specs/claw/required_red.t27", "t27b": "pass"}, {"path": "specs/claw/sarif.t27", "t27b": "pass"}, {"path": "specs/claw/sarif_feed.t27", "t27b": "pass"}, {"path": "specs/claw/sarif_pins.t27", "t27b": "pass"}, {"path": "specs/claw/service.t27", "t27b": "pass"}, + {"path": "specs/claw/shape_race.t27", "t27b": "pass"}, + {"path": "specs/claw/social_visuals.t27", "t27b": "pass"}, {"path": "specs/claw/suggest.t27", "t27b": "pass"}, {"path": "specs/claw/suggest_feed.t27", "t27b": "pass"}, {"path": "specs/claw/t27f_gate.t27", "t27b": "pass"}, {"path": "specs/claw/t27f_gate_pins.t27", "t27b": "pass"}, {"path": "specs/claw/text.t27", "t27b": "pass"}, + {"path": "specs/claw/token_shape.t27", "t27b": "pass"}, {"path": "specs/compiler/c_array_repeat_wrapper.t27", "t27b": "pass"}, {"path": "specs/compiler/c_fixed_array_return.t27", "t27b": "pass"}, {"path": "specs/compiler/codegen_conformance.t27", "t27b": "pass"}, @@ -327,6 +339,7 @@ {"path": "specs/fpga/apb_bridge.t27", "t27b": "pass"}, {"path": "specs/fpga/assembler.t27", "t27b": "pass"}, {"path": "specs/fpga/axi4.t27", "t27b": "pass"}, + {"path": "specs/fpga/board_receipt.t27", "t27b": "pass"}, {"path": "specs/fpga/boards/arty_a7_integration.t27", "t27b": "pass"}, {"path": "specs/fpga/boards/qmtech_a100t_integration.t27", "t27b": "pass"}, {"path": "specs/fpga/bootrom.t27", "t27b": "pass"}, @@ -336,6 +349,7 @@ {"path": "specs/fpga/crossopt.t27", "t27b": "pass"}, {"path": "specs/fpga/cts.t27", "t27b": "pass"}, {"path": "specs/fpga/ddr3_head_bandwidth.t27", "t27b": "pass"}, + {"path": "specs/fpga/ddr3_ternary_matvec.t27", "t27b": "pass"}, {"path": "specs/fpga/dft.t27", "t27b": "pass"}, {"path": "specs/fpga/die_probe.t27", "t27b": "pass"}, {"path": "specs/fpga/dsp/adders.t27", "t27b": "pass"}, @@ -368,6 +382,8 @@ {"path": "specs/fpga/secded.t27", "t27b": "pass"}, {"path": "specs/fpga/simulator.t27", "t27b": "pass"}, {"path": "specs/fpga/stdlib.t27", "t27b": "pass"}, + {"path": "specs/fpga/systolic.t27", "t27b": "pass"}, + {"path": "specs/fpga/tern9m_l0_k_weights.t27", "t27b": "pass"}, {"path": "specs/fpga/ternary_isa.t27", "t27b": "pass"}, {"path": "specs/fpga/ternary_link.t27", "t27b": "pass"}, {"path": "specs/fpga/testbench/assembler_tb.t27", "t27b": "pass"}, @@ -441,7 +457,9 @@ {"path": "specs/igla/coder/benchmark.t27", "t27b": "blocked", "blocker": "ExprIdentifier(undeclared)", "reason": "unimplemented"}, {"path": "specs/igla/coder/dataset.t27", "t27b": "blocked", "blocker": "ExprIdentifier(undeclared)", "reason": "unimplemented"}, {"path": "specs/igla/coder/eval.t27", "t27b": "blocked", "blocker": "ExprIdentifier(undeclared)", "reason": "unimplemented"}, + {"path": "specs/igla/coder/kill_rule.t27", "t27b": "blocked", "blocker": "ExprCall(undeclared fn)", "reason": "unimplemented"}, {"path": "specs/igla/coder/pipeline.t27", "t27b": "blocked", "blocker": "ConstDecl(slice)", "reason": "unimplemented"}, + {"path": "specs/igla/coder/preflight.t27", "t27b": "pass"}, {"path": "specs/igla/coder/prm.t27", "t27b": "pass"}, {"path": "specs/igla/coder/ternary_gemm_gpu.t27", "t27b": "pass"}, {"path": "specs/igla/coder/ternary_gemm_ptx.t27", "t27b": "pass"}, @@ -652,6 +670,7 @@ {"path": "specs/policy/json_int53.t27", "t27b": "pass"}, {"path": "specs/policy/l2_generation.t27", "t27b": "pass"}, {"path": "specs/policy/lang_en.t27", "t27b": "pass"}, + {"path": "specs/policy/linked_issue_prose.t27", "t27b": "pass"}, {"path": "specs/policy/merge_ours.t27", "t27b": "pass"}, {"path": "specs/policy/own_language.t27", "t27b": "pass"}, {"path": "specs/policy/pr_budget.t27", "t27b": "pass"}, @@ -737,6 +756,7 @@ {"path": "specs/port/fpga/verilog/gft_train1_jtag.t27", "t27b": "pass"}, {"path": "specs/port/fpga/verilog/link_node.t27", "t27b": "pass"}, {"path": "specs/port/fpga/verilog/mvp_ternary_classifier_check.t27", "t27b": "pass"}, + {"path": "specs/port/fpga/verilog/mvp_ternary_classifier_jtag_noport.t27", "t27b": "pass"}, {"path": "specs/port/fpga/verilog/mvp_ternary_classifier_top.t27", "t27b": "pass"}, {"path": "specs/port/fpga/verilog/phi_weights_jtag.t27", "t27b": "pass"}, {"path": "specs/port/fpga/verilog/ternary_mac_demo_top.t27", "t27b": "pass"}, @@ -745,6 +765,7 @@ {"path": "specs/port/fpga/vivado/blinky.t27", "t27b": "pass"}, {"path": "specs/port/fpga/vivado/gf16_matmul4x4.t27", "t27b": "pass"}, {"path": "specs/port/fpga/vivado/gf16_matmul4x4_top.t27", "t27b": "pass"}, + {"path": "specs/port/fpga/vivado/gf16_matmul_top.t27", "t27b": "pass"}, {"path": "specs/port/fpga/vivado/gf16_top.t27", "t27b": "pass"}, {"path": "specs/port/fpga/vivado/gf16_uart_sim_bench.t27", "t27b": "pass"}, {"path": "specs/port/fpga/vivado/uart_detect_top.t27", "t27b": "pass"}, @@ -1107,6 +1128,7 @@ {"path": "specs/port/trinity/src/wasm_stubs/auto_shard_stub.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/bin/seed-agent/src/main.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/asha.t27", "t27b": "pass"}, + {"path": "specs/port/trios-railway/crates/trios-igla-race/src/bin/reset_queue.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/ema.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/invariants.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/race.t27", "t27b": "pass"}, @@ -1141,6 +1163,10 @@ {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-02/src/double_ratchet_sending_chain_rotation_guard.t27", "t27b": "blocked", "blocker": "type mismatch", "reason": "unimplemented"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-02/src/epoch_rollover_wraparound_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-02/src/skipped_message_key_gap_bound_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/commit_path_secret_aead_keying_mismatch.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/commit_secret_export_collision.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/commit_signature.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/concurrent_add_remove.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/confirmation_tag_chain.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/external_commit.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/external_commit_resumption_psk_misbinding.t27", "t27b": "pass"}, @@ -1158,6 +1184,8 @@ {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/pcs_healing.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/proposal_ref_collision.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/proposal_validation.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/psk_external_injection.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/ratchet_tree_blank_node_depth_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/ratchet_tree_resolution_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/ratchet_tree_update_integrity_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-03/src/unmerged_leaves_bound_guard.t27", "t27b": "pass"}, @@ -1167,9 +1195,12 @@ {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/mac_truncation.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/mls_plaintext_framing_integrity.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_alignment_boundary_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_block_alignment_uniformity_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_block_cipher_mode_uniformity_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_block_size_alignment_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_byte_entropy_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_byte_pattern_uniformity_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_ciphertext_length_class_balance_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_ciphertext_length_consistency_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_class_collision_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_class_downgrade_guard.t27", "t27b": "pass"}, @@ -1181,11 +1212,20 @@ {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_crypto_binding_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_fill_byte_randomness_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_key_derivation_domain_separation_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_key_derivation_uniqueness_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_key_rotation_uniformity_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_length_class_transition_uniformity_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_length_distribution_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_metadata_leak_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_nonce_reuse_detection_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_oracle_chosen_ct.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_oracle_timing_sidechannel.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_payload_entropy_minimum_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_response_timing_fingerprint_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_timing_side_channel_guard.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/padding_timing_uniformity_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/safety_number.t27", "t27b": "pass"}, + {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/welcome_path_secret_unmasking.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-04/src/wire_frame_integrity_guard.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-05/src/envelope_size_distribution_uniformity.t27", "t27b": "pass"}, {"path": "specs/port/trios/crates/trios-chat/rings/CR-CHAT-05/src/store_compaction_integrity_guard.t27", "t27b": "blocked", "blocker": "ExprArrayLiteral", "reason": "unimplemented"}, @@ -1378,6 +1418,7 @@ {"path": "specs/queen/replay/golden_chain_history.t27", "t27b": "pass"}, {"path": "specs/queen/replay/reviewer_decisions.t27", "t27b": "pass"}, {"path": "specs/queen/replay/telemetry_decisions.t27", "t27b": "pass"}, + {"path": "specs/queen/review_first.t27", "t27b": "pass"}, {"path": "specs/queen/review_lanes.t27", "t27b": "pass"}, {"path": "specs/queen/review_log.t27", "t27b": "pass"}, {"path": "specs/queen/review_valve.t27", "t27b": "pass"}, @@ -1407,15 +1448,22 @@ {"path": "specs/server/router.t27", "t27b": "blocked", "blocker": "type (alias)", "reason": "unimplemented"}, {"path": "specs/server/session.t27", "t27b": "pass"}, {"path": "specs/server/sse.t27", "t27b": "blocked", "blocker": "type (alias)", "reason": "unimplemented"}, + {"path": "specs/skills/t27-deps-card.t27", "t27b": "pass"}, {"path": "specs/skills/t27-issue-shape.t27", "t27b": "pass"}, {"path": "specs/skills/t27-measure-corpus.t27", "t27b": "pass"}, {"path": "specs/skills/t27-phi-loop.t27", "t27b": "pass"}, {"path": "specs/skills/t27-providers-steward.t27", "t27b": "pass"}, + {"path": "specs/skills/t27-reel-beat.t27", "t27b": "pass"}, {"path": "specs/skills/t27-self-host.t27", "t27b": "pass"}, + {"path": "specs/skills/t27-social-visual-consent.t27", "t27b": "pass"}, + {"path": "specs/skills/t27-social-visual-format.t27", "t27b": "pass"}, + {"path": "specs/skills/t27-social-visual-gate.t27", "t27b": "pass"}, + {"path": "specs/skills/t27-social-visual-scene.t27", "t27b": "pass"}, {"path": "specs/skills/t27-t27f.t27", "t27b": "pass"}, {"path": "specs/skills/t27-tri-pipeline.t27", "t27b": "pass"}, {"path": "specs/skills/t27-tri.t27", "t27b": "pass"}, {"path": "specs/skills/t27-wave-audit.t27", "t27b": "pass"}, + {"path": "specs/skills/t27-why-card.t27", "t27b": "pass"}, {"path": "specs/skills/t27-wrap-up.t27", "t27b": "pass"}, {"path": "specs/skills/trinity-blog-post.t27", "t27b": "pass"}, {"path": "specs/skills/trinity-board-sync.t27", "t27b": "pass"}, @@ -1442,8 +1490,10 @@ {"path": "specs/system/unified-agent-operations.t27", "t27b": "pass"}, {"path": "specs/t27f/build.t27", "t27b": "pass"}, {"path": "specs/t27f/deck.t27", "t27b": "pass"}, + {"path": "specs/t27f/deps_card.t27", "t27b": "blocked", "blocker": "ExprCall(std.*)", "reason": "unimplemented"}, {"path": "specs/t27f/frame_cache.t27", "t27b": "pass"}, {"path": "specs/t27f/frame_shard.t27", "t27b": "blocked", "blocker": "type [N]T", "reason": "unimplemented"}, + {"path": "specs/t27f/frontend.t27", "t27b": "pass"}, {"path": "specs/t27f/hot.t27", "t27b": "pass"}, {"path": "specs/t27f/lesson_reel.t27", "t27b": "blocked", "blocker": "type [N]T", "reason": "unimplemented"}, {"path": "specs/t27f/outfit_atlas.t27", "t27b": "pass"}, @@ -1452,12 +1502,16 @@ {"path": "specs/t27f/pbr_sphere.t27", "t27b": "pass"}, {"path": "specs/t27f/quest.t27", "t27b": "pass"}, {"path": "specs/t27f/reel.t27", "t27b": "pass"}, + {"path": "specs/t27f/reel_beats.t27", "t27b": "pass"}, + {"path": "specs/t27f/reel_factory.t27", "t27b": "pass"}, + {"path": "specs/t27f/reel_factory_plan.t27", "t27b": "pass"}, {"path": "specs/t27f/reel_gf16_nmse.t27", "t27b": "pass"}, {"path": "specs/t27f/reel_marketing.t27", "t27b": "pass"}, {"path": "specs/t27f/reel_particles.t27", "t27b": "blocked", "blocker": "ExprCall(@sin)", "reason": "unimplemented"}, {"path": "specs/t27f/reel_particles_mux.t27", "t27b": "pass"}, {"path": "specs/t27f/reel_t27_basics.t27", "t27b": "pass"}, {"path": "specs/t27f/reel_topics.t27", "t27b": "pass"}, + {"path": "specs/t27f/spec_deps.t27", "t27b": "pass"}, {"path": "specs/t27f/spec_graph.t27", "t27b": "pass"}, {"path": "specs/t27f/spec_graph_derive.t27", "t27b": "pass"}, {"path": "specs/t27f/spec_graph_export.t27", "t27b": "pass"}, @@ -1470,6 +1524,7 @@ {"path": "specs/t27f/topic_engine.t27", "t27b": "pass"}, {"path": "specs/t27f/topic_reel.t27", "t27b": "pass"}, {"path": "specs/t27f/ui_plan.t27", "t27b": "pass"}, + {"path": "specs/t27f/why_card.t27", "t27b": "blocked", "blocker": "type [N]T", "reason": "unimplemented"}, {"path": "specs/t27mcp/check_hints.t27", "t27b": "pass"}, {"path": "specs/t27mcp/dashboard.t27", "t27b": "pass"}, {"path": "specs/t27mcp/examples.t27", "t27b": "pass"}, @@ -1479,6 +1534,7 @@ {"path": "specs/t27mcp/spec_brief.t27", "t27b": "pass"}, {"path": "specs/t27mcp/telegram.t27", "t27b": "pass"}, {"path": "specs/t27mcp/tool_schema.t27", "t27b": "codegen", "blocker": "t27b: unsupported construct FnDecl(frame or branch range) at line 0 (fn the round trip fixtures: a tuple, arrays, a boolean, nested objects: conditional branch out of range (function too large))", "reason": "unimplemented"}, + {"path": "specs/t27mcp/visual_output.t27", "t27b": "pass"}, {"path": "specs/ternary/activation_quantizer.t27", "t27b": "pass"}, {"path": "specs/ternary/bitnet_layer.t27", "t27b": "pass"}, {"path": "specs/ternary/bitnet_majority.t27", "t27b": "pass"}, @@ -1563,6 +1619,7 @@ {"path": "specs/tools/t27o/rust_convert.t27", "t27b": "pass"}, {"path": "specs/tools/t27o/rust_std.t27", "t27b": "pass"}, {"path": "specs/tools/t27o/service.t27", "t27b": "pass"}, + {"path": "specs/tools/t27o/spice_convert.t27", "t27b": "pass"}, {"path": "specs/tools/tri/abandoned.t27", "t27b": "pass"}, {"path": "specs/tools/tri/cell.t27", "t27b": "pass"}, {"path": "specs/tools/tri/census.t27", "t27b": "pass"}, @@ -1897,6 +1954,7 @@ {"path": "specs/tri/utils/time.t27", "t27b": "blocked", "blocker": "ExprCall(@divFloor)", "reason": "unimplemented"}, {"path": "specs/tri/utils/utf8.t27", "t27b": "blocked", "blocker": "ExprCall(@truncate)", "reason": "unimplemented"}, {"path": "specs/tri/utils/version.t27", "t27b": "pass"}, + {"path": "specs/trinet/eth_top_ax7203.t27", "t27b": "pass"}, {"path": "specs/trinet/etx.t27", "t27b": "pass"}, {"path": "specs/trinet/ledger.t27", "t27b": "pass"}, {"path": "specs/trinet/mint_authority.t27", "t27b": "pass"}, @@ -1966,6 +2024,7 @@ {"path": "specs/tutorial/07_tests_invariants_benches.t27", "t27b": "pass"}, {"path": "specs/tutorial/08_modules_and_arrays.t27", "t27b": "pass"}, {"path": "specs/ui/queen_evidence.t27", "t27b": "pass"}, + {"path": "specs/ui/viewport.t27", "t27b": "pass"}, {"path": "specs/vault/audit.t27", "t27b": "pass"}, {"path": "specs/vault/merge.t27", "t27b": "pass"}, {"path": "specs/vault/policy.t27", "t27b": "pass"}, From f3628bd5f618a3b7ab8f629579d244d958cd3013 Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Mon, 19 Oct 2026 09:30:36 +0700 Subject: [PATCH 11/11] t27b ledger: list three specs master added unlisted, all t27b pass as CI measured (Refs #10404) specs/fpga/analog_cell_ax7203.t27 and trios-igla-race bin/debug_queue.t27, bin/smoke_agent.t27. Co-Authored-By: Claude Opus 5.5 --- docs/reports/t27b_expectations.json | 3 +++ 1 file changed, 3 insertions(+) diff --git a/docs/reports/t27b_expectations.json b/docs/reports/t27b_expectations.json index ca46a8ffdd..c1d836ecb9 100644 --- a/docs/reports/t27b_expectations.json +++ b/docs/reports/t27b_expectations.json @@ -336,6 +336,7 @@ {"path": "specs/effects/replay_oracle.t27", "t27b": "pass"}, {"path": "specs/file/schema.t27", "t27b": "pass"}, {"path": "specs/fpga/adapter.t27", "t27b": "pass"}, + {"path": "specs/fpga/analog_cell_ax7203.t27", "t27b": "pass"}, {"path": "specs/fpga/apb_bridge.t27", "t27b": "pass"}, {"path": "specs/fpga/assembler.t27", "t27b": "pass"}, {"path": "specs/fpga/axi4.t27", "t27b": "pass"}, @@ -1128,7 +1129,9 @@ {"path": "specs/port/trinity/src/wasm_stubs/auto_shard_stub.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/bin/seed-agent/src/main.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/asha.t27", "t27b": "pass"}, + {"path": "specs/port/trios-railway/crates/trios-igla-race/src/bin/debug_queue.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/bin/reset_queue.t27", "t27b": "pass"}, + {"path": "specs/port/trios-railway/crates/trios-igla-race/src/bin/smoke_agent.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/ema.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/invariants.t27", "t27b": "pass"}, {"path": "specs/port/trios-railway/crates/trios-igla-race/src/race.t27", "t27b": "pass"},