gHashTag · Oct 19, 2026 · Oct 19, 2026
@@ -67,7 +67,7 @@ Then: close the sub-issue with `Closes #N`, tick it in the epic, open the next m
 
 **The board flow is local; the Railway lab is not needed.** `O="$(git rev-parse --git-common-dir)/../build/fpga/openxc7"` (the main checkout's `build/`, from any worktree) holds `np-main/build/nextpnr-xilinx`, the AX7203 chipdb `xc7a200tfbg484-2.bin`, `prjxray/utils/fasm2frames.py` (run with `$O/venv/bin/python3`) and `px-main/build/tools/xc7frames2bit`; yosys and openFPGALoader come from Homebrew. nextpnr on this Mac takes about 22 s and 2.4 GB: check `memory_pressure` (>= 25 % free) and `df -h /` (>= 8 GB) first and never run two at once. Read the memory notes before declaring a blocker: the recipe was there.
 
-Sandboxed sessions: `export ZIG_GLOBAL_CACHE_DIR=$TMPDIR/zig-cache ZIG_LOCAL_CACHE_DIR=$TMPDIR/zig-local`; commit with `git -c core.hooksPath=.githooks commit`; `gh` fails TLS, use curl with `$(gh auth token)`. There is no `t27t` in this repository (searched 2026-10-18); do not invent one.
+Sandboxed sessions: `export ZIG_GLOBAL_CACHE_DIR=$TMPDIR/zig-cache ZIG_LOCAL_CACHE_DIR=$TMPDIR/zig-local`; commit with `git -c core.hooksPath=.githooks commit`; `gh` fails TLS, use curl with `$(gh auth token)`. The t27 trainer is t27t: `specs/t27t/t27t.t27` (`tri t27t`); its runner is gHashTag/igla-coder-gpu.
 
 ## The overnight loop tick
 

@@ -178,6 +178,33 @@ derived, suggested or unknown. The surfaces and their status are
   node and edge instead of the whole file. Use it before reading a big spec in
   full.
 
+### t27t: the t27 trainer
+
+t27t is the t27 training product. Its law is
+[`specs/t27t/t27t.t27`](specs/t27t/t27t.t27), run with `tri t27t`. Its runner
+today is the PyTorch trainer gHashTag/igla-coder-gpu (`train.py`, tern-tc
+models). The spec says what the runner must do; it does not rewrite the runner.
+
+- **What makes training cheaper** (measured on an RTX 4090 with CUTLASS 3.5.1,
+  2026-10-18, #10360): the forward GEMMs run on int8 tensor cores, at 2.9x to
+  3.5x bf16, and the gradients stay in bf16. PyTorch reaches those cores through
+  `torch._int_mm`, so no t27c CUDA backend is needed for the saving.
+- **What makes it more expensive:** int8 dp4a on CUDA cores runs at 0.91x to
+  0.94x of bf16. `pick_path` never chooses it.
+- **The saving law:** f x (1 - 1/s), which is 16% to 24%, or $16 to $24 on a
+  $100 run. f is the share of a step spent in forward GEMMs, and it is **an
+  estimate, not a measurement** (`F_MEASURED = false`). Do not quote the saving
+  without saying so.
+- **When a generated kernel takes over:** only on a shape where it beats tuned
+  CUTLASS s8 by 10% (`GEN_CUDA_BAR_E2`).
+- **Open:** the runner is not wired to `torch._int_mm` yet
+  (`RUNNER_USES_INT_MM = false`). The spec's `NOT_MEASURED` lists what else was
+  not measured.
+- **Numbers about our formats** go through the papers check,
+  [`specs/numeric/paper_claims.t27`](specs/numeric/paper_claims.t27) (`tri papers`).
+  Of 45 numbers printed in three papers, 34 recompute, 8 do not, 3 cannot be
+  checked.
+
 ### Hot reload and incremental work: only what changed is redone
 
 The rule: key every result on the content hash of its inputs and reuse it while

@@ -1553,6 +1553,7 @@
   {"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/t27t/t27t.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"},

@@ -41,7 +41,7 @@ tri -- the t27 front door.
 
 Compiler commands (forwarded to t27c):
   tri test [--json] [...]     run the conformance suite (t27c suite --repo-root .)
-  tri papers | factory        the numbers three format papers print (specs/numeric/paper_claims.t27) | the reel factory's uniqueness gate (specs/t27f/reel_factory.t27)
+  tri papers | factory | t27t the numbers three format papers print (specs/numeric/paper_claims.t27) | the reel factory's uniqueness gate (specs/t27f/reel_factory.t27) | the t27 trainer t27t (specs/t27t/t27t.t27)
   tri <anything else>         forwarded to t27c unchanged (gen, seal, fpga-build, ...)
 
 Local commands (no compiler needed):
@@ -532,7 +532,7 @@ if [[ ! -x "$T27C" ]]; then
 fi
 
 case "$cmd" in papers) exec "$T27C" test-report "$REPO_ROOT/specs/numeric/paper_claims.t27" "${@:2}";;
-  factory) exec "$T27C" test-report "$REPO_ROOT/specs/t27f/reel_factory.t27" "${@:2}";; esac
+  factory) exec "$T27C" test-report "$REPO_ROOT/specs/t27f/reel_factory.t27" "${@:2}";; t27t) exec "$T27C" test-report "$REPO_ROOT/specs/t27t/t27t.t27" "${@:2}";; esac
 # Most subcommands are implemented in the Rust `tri` binary (cli/tri),
 # which is a DIFFERENT binary from t27c. Everything not handled above was
 # forwarded to t27c, so every one of them existed and was unreachable through the

@@ -0,0 +1,11 @@
+// SPDX-License-Identifier: Apache-2.0
+# OWNERS -- specs/t27t/
+
+## Primary
+
+**T-Trainer** -- t27t, the t27 trainer: measured GEMM speed-ups, the saving law, the GEMM path rule.
+
+## Dependencies
+
+- The bench of gHashTag/t27#10360 (numbers are copied in with their source, and recomputed by the tests).
+- The runner gHashTag/igla-coder-gpu, which t27t constrains and does not rewrite.
@@ -0,0 +1,276 @@
+// SPDX-License-Identifier: Apache-2.0
+; specs/t27t/t27t.t27 -- t27t, the t27 trainer: how a training run gets cheaper, measured, not promised.
+; --
+; What t27t is.
+; - t27t is the name of the t27 training product. Its law is this file; its runner today is the
+;   PyTorch trainer gHashTag/igla-coder-gpu (train.py, the tern-tc 9m / 100m / 500m models).
+;   The runner does not change its language: t27t says what the runner must do and checks it.
+; - Its first rule comes from one bench (gHashTag/t27#10360, RTX 4090, CUTLASS 3.5.1, 2026-10-18):
+;   a training step gets cheaper when the FORWARD GEMMs run on int8 tensor cores (mma.sync s8,
+;   reached from PyTorch by torch._int_mm) and the gradients stay in bf16. No t27c CUDA backend is
+;   needed for that.
+; - Its second rule is the papers check (specs/numeric/paper_claims.t27, `tri papers`): a number
+;   t27t prints about its formats is one a spec can recompute. 45 printed numbers in three papers
+;   were recomputed: 34 fit, 8 do not, 3 cannot be checked. Those counts live in that spec, not here.
+; --
+; The rules of this file.
+; - Times are in units of 10^-4 ms (0.2254 ms is 2254). Speed ratios are in hundredths (2.93x is 293).
+;   Savings are in tenths of a percent (21.8% is 218), which is also dimes saved on a $100 run.
+; - Every ratio the bench comment printed is kept beside the times it came from, and a test
+;   recomputes it from those times. The times were printed to four digits, so a recomputed ratio
+;   may differ from the printed one by one hundredth (3.5456 from 0.8506 / 0.2399 against a printed
+;   3.54); the tests allow exactly that and nothing more.
+; - What was not measured is listed as not measured (NOT_MEASURED). Above all f, the share of a
+;   real tern-tc-500m step spent in forward GEMMs, is an estimate: F_MEASURED is false.
+; phi^2 + 1/phi^2 = 3 | TRINITY
+
+module T27t;
+
+pub const NAME : str = "t27t";
+pub const RUNNER : str = "gHashTag/igla-coder-gpu train.py (PyTorch)";
+pub const BENCH : str = "gHashTag/t27#10360: RTX 4090, CUTLASS 3.5.1, sm_89, CUDA 12.4, 2026-10-18";
+
+; The runner's forward GEMM is not yet wired to torch._int_mm, and f has not been measured.
+pub const RUNNER_USES_INT_MM : bool = false;
+pub const F_MEASURED : bool = false;
+
+// ---------------------------------------------------------------- the bench
+
+pub const SHAPES : u32 = 5;
+
+pub const SHAPE_NAME : [5]str = [
+    "4096 x 3456 x 1280 (FFN up)",
+    "4096 x 1280 x 3456 (FFN down)",
+    "16384 x 3456 x 1280",
+    "16384 x 1280 x 3456",
+    "16384 x 1280 x 1280 (attn proj)"
+];
+
+pub const SHAPE_M : [5]u64 = [4096, 4096, 16384, 16384, 16384];
+pub const SHAPE_N : [5]u64 = [3456, 1280, 3456, 1280, 1280];
+pub const SHAPE_K : [5]u64 = [1280, 3456, 1280, 3456, 1280];
+
+; Fastest CUTLASS kernel per shape and data type, 50 timed iterations after 10 warm-up.
+pub const T_BF16 : [5]u64 = [2254, 2239, 8540, 8506, 3215];
+pub const T_INT8_TC : [5]u64 = [770, 773, 2701, 2399, 1007];
+pub const T_DP4A : [5]u64 = [2406, 2470, 9234, 9278, 3460];
+
+; Ratios as the bench comment printed them.
+pub const X_INT8_TC_PRINTED : [5]u64 = [293, 290, 316, 354, 319];
+pub const X_DP4A_PRINTED : [5]u64 = [94, 91, 92, 92, 93];
+
+; The datasheet promise for int8 over bf16 tensor cores on Ada, in hundredths.
+pub const X_DATASHEET : u64 = 400;
+
+; Speed of a path against bf16, rounded half up to hundredths.
+pub fn ratio_e2(t_bf16: u64, t_path: u64) -> u64 {
+    return (t_bf16 * 200 + t_path) / (2 * t_path);
+}
+
+pub fn within_one(a: u64, b: u64) -> bool {
+    if (a > b) { return a - b <= 1; }
+    return b - a <= 1;
+}
+
+pub fn x_int8_tc(i: u32) -> u64 {
+    return ratio_e2(T_BF16[i], T_INT8_TC[i]);
+}
+
+pub fn x_dp4a(i: u32) -> u64 {
+    return ratio_e2(T_BF16[i], T_DP4A[i]);
+}
+
+pub const COL_INT8_TC : u32 = 0;
+pub const COL_DP4A : u32 = 1;
+
+pub fn printed(col: u32, i: u32) -> u64 {
+    if (col == COL_DP4A) { return X_DP4A_PRINTED[i]; }
+    return X_INT8_TC_PRINTED[i];
+}
+
+pub fn min_printed(col: u32) -> u64 {
+    var i : u32 = 1;
+    var m : u64 = printed(col, 0);
+    while (i < SHAPES) {
+        if (printed(col, i) < m) { m = printed(col, i); }
+        i = i + 1;
+    }
+    return m;
+}
+
+pub fn max_printed(col: u32) -> u64 {
+    var i : u32 = 1;
+    var m : u64 = printed(col, 0);
+    while (i < SHAPES) {
+        if (printed(col, i) > m) { m = printed(col, i); }
+        i = i + 1;
+    }
+    return m;
+}
+
+/// How many shapes have a recomputed ratio within one hundredth of the printed one.
+pub fn printed_ratios_that_hold() -> u32 {
+    var i : u32 = 0;
+    var c : u32 = 0;
+    while (i < SHAPES) {
+        if (within_one(x_int8_tc(i), X_INT8_TC_PRINTED[i]) and within_one(x_dp4a(i), X_DP4A_PRINTED[i])) {
+            c = c + 1;
+        }
+        i = i + 1;
+    }
+    return c;
+}
+
+/// How many shapes have the CUDA-core int8 path (dp4a) slower than bf16.
+pub fn dp4a_slower_count() -> u32 {
+    var i : u32 = 0;
+    var c : u32 = 0;
+    while (i < SHAPES) {
+        if (T_DP4A[i] > T_BF16[i]) { c = c + 1; }
+        i = i + 1;
+    }
+    return c;
+}
+
+// ---------------------------------------------------------------- the saving law
+
+; A run saves f x (1 - 1/s): f is the share of step time that moves to int8, s its speed-up.
+; f = f_num / f_den, s = s_e2 / 100; the result is in tenths of a percent, rounded half up.
+pub fn saving_e3(f_num: u64, f_den: u64, s_e2: u64) -> u64 {
+    if (s_e2 <= 100) { return 0; }
+    return (2 * f_num * 1000 * (s_e2 - 100) + f_den * s_e2) / (2 * f_den * s_e2);
+}
+
+; The two estimates of f the bench used: forward GEMMs are one of three matmuls per layer
+; (forward, grad-input, grad-weight); with matmul at 75% of the step that becomes one quarter.
+pub const F_ALL_MATMUL_NUM : u64 = 1;
+pub const F_ALL_MATMUL_DEN : u64 = 3;
+pub const F_MATMUL_75_NUM : u64 = 1;
+pub const F_MATMUL_75_DEN : u64 = 4;
+
+; The cost table of the bench comment, in tenths of a percent: rows f = 1/3 and 1/4, columns
+; s = 2.9, 3.2, 3.5.
+pub const S_COLUMNS : [3]u64 = [290, 320, 350];
+pub const SAVING_THIRD_PRINTED : [3]u64 = [218, 229, 238];
+pub const SAVING_QUARTER_PRINTED : [3]u64 = [164, 172, 179];
+
+/// Dollars saved on a $100 run, rounded half up to whole dollars.
+pub fn dollars_per_100(saving: u64) -> u64 {
+    return (saving + 5) / 10;
+}
+
+// ---------------------------------------------------------------- which path a GEMM takes
+
+pub const PASS_FORWARD : u32 = 0;
+pub const PASS_BACKWARD : u32 = 1;
+
+pub const PATH_BF16 : u32 = 0;
+pub const PATH_INT8_TC : u32 = 1;
+pub const PATH_DP4A : u32 = 2;
+pub const PATH_GEN_CUDA : u32 = 3;
+
+pub const PATHS : [4]str = [
+    "bf16 tensor cores",
+    "int8 tensor cores via torch._int_mm",
+    "int8 dp4a on CUDA cores",
+    "a t27c-generated CUDA kernel"
+];
+
+; A generated kernel replaces torch._int_mm only when it beats tuned CUTLASS s8 by 10% (RFC #10360).
+pub const GEN_CUDA_BAR_E2 : u64 = 110;
+
+/// The path t27t gives one GEMM.
+/// pass: forward or backward. has_int8_tc: the device has int8 tensor cores.
+/// gen_vs_cutlass_e2: a generated kernel's measured speed against CUTLASS s8 on this shape, in
+/// hundredths, or 0 when none has been measured.
+pub fn pick_path(pass: u32, has_int8_tc: bool, gen_vs_cutlass_e2: u64) -> u32 {
+    // Gradients stay bf16: the saving is counted with them in bf16.
+    if (pass == PASS_BACKWARD) { return PATH_BF16; }
+    // Without int8 tensor cores the only int8 left is dp4a, measured slower than bf16 on every shape.
+    if (!has_int8_tc) { return PATH_BF16; }
+    if (gen_vs_cutlass_e2 >= GEN_CUDA_BAR_E2) { return PATH_GEN_CUDA; }
+    return PATH_INT8_TC;
+}
+
+// ---------------------------------------------------------------- what is known and what is not
+
+pub const NOT_MEASURED : [4]str = [
+    "f, the share of a tern-tc-500m step spent in forward GEMMs",
+    "cuBLAS itself; torch._int_mm is represented by the best CUTLASS s8 kernel",
+    "activation quantization from bf16 to int8 and its kernel launch",
+    "the runner wired to torch._int_mm, end to end"
+];
+
+// ---------------------------------------------------------------- tests
+
+test "every printed ratio recomputes from its times" {
+    assert x_int8_tc(0) == 293;
+    assert x_int8_tc(1) == 290;
+    assert x_int8_tc(2) == 316;
+    assert x_int8_tc(3) == 355;
+    assert x_int8_tc(4) == 319;
+    assert x_dp4a(0) == 94;
+    assert x_dp4a(1) == 91;
+    assert x_dp4a(2) == 92;
+    assert x_dp4a(3) == 92;
+    assert x_dp4a(4) == 93;
+    assert printed_ratios_that_hold() == 5;
+}
+
+test "int8 tensor cores: 2.9x to 3.5x, under the datasheet 4x" {
+    assert min_printed(COL_INT8_TC) == 290;
+    assert max_printed(COL_INT8_TC) == 354;
+    assert max_printed(COL_INT8_TC) < X_DATASHEET;
+}
+
+test "dp4a on CUDA cores is slower than bf16 on every shape" {
+    assert dp4a_slower_count() == SHAPES;
+    assert max_printed(COL_DP4A) == 94;
+    assert min_printed(COL_DP4A) == 91;
+}
+
+test "the cost table: 16 to 24 dollars on a 100 dollar run" {
+    assert saving_e3(F_ALL_MATMUL_NUM, F_ALL_MATMUL_DEN, S_COLUMNS[0]) == SAVING_THIRD_PRINTED[0];
+    assert saving_e3(F_ALL_MATMUL_NUM, F_ALL_MATMUL_DEN, S_COLUMNS[1]) == SAVING_THIRD_PRINTED[1];
+    assert saving_e3(F_ALL_MATMUL_NUM, F_ALL_MATMUL_DEN, S_COLUMNS[2]) == SAVING_THIRD_PRINTED[2];
+    assert saving_e3(F_MATMUL_75_NUM, F_MATMUL_75_DEN, S_COLUMNS[0]) == SAVING_QUARTER_PRINTED[0];
+    assert saving_e3(F_MATMUL_75_NUM, F_MATMUL_75_DEN, S_COLUMNS[1]) == SAVING_QUARTER_PRINTED[1];
+    assert saving_e3(F_MATMUL_75_NUM, F_MATMUL_75_DEN, S_COLUMNS[2]) == SAVING_QUARTER_PRINTED[2];
+    assert dollars_per_100(SAVING_QUARTER_PRINTED[0]) == 16;
+    assert dollars_per_100(SAVING_THIRD_PRINTED[2]) == 24;
+}
+
+test "no speed-up, no saving; a slower path is never a saving" {
+    assert saving_e3(1, 3, 100) == 0;
+    assert saving_e3(1, 3, 94) == 0;
+    assert saving_e3(1, 1, 200) == 500;
+}
+
+test "the path rule" {
+    assert pick_path(PASS_FORWARD, true, 0) == PATH_INT8_TC;
+    assert pick_path(PASS_BACKWARD, true, 0) == PATH_BF16;
+    assert pick_path(PASS_FORWARD, false, 0) == PATH_BF16;
+    assert pick_path(PASS_FORWARD, true, 109) == PATH_INT8_TC;
+    assert pick_path(PASS_FORWARD, true, 110) == PATH_GEN_CUDA;
+    assert pick_path(PASS_BACKWARD, true, 150) == PATH_BF16;
+    assert PATHS[PATH_INT8_TC] == "int8 tensor cores via torch._int_mm";
+}
+
+test "negative control: a ratio off by two is caught" {
+    assert !within_one(x_int8_tc(3), 353);
+    assert within_one(x_int8_tc(3), 354);
+    assert saving_e3(1, 3, 290) != 219;
+}
+
+test "what is not known stays written down" {
+    assert !F_MEASURED;
+    assert !RUNNER_USES_INT_MM;
+    assert NOT_MEASURED[0] == "f, the share of a tern-tc-500m step spent in forward GEMMs";
+    assert NAME == "t27t";
+}
+
+invariant dp4a_is_never_picked { pick_path(PASS_FORWARD, false, 0) != PATH_DP4A and pick_path(PASS_FORWARD, true, 0) != PATH_DP4A }
+
+invariant gradients_stay_bf16 { pick_path(PASS_BACKWARD, true, 200) == PATH_BF16 }
+
+invariant the_bench_has_five_shapes { SHAPE_M[0] == 4096 and SHAPE_N[1] == 1280 and SHAPE_K[3] == 3456 }