From 2a755bf2f6e0d4a204f40b8a70c3901d3681e5fd Mon Sep 17 00:00:00 2001 From: Dmitrii Vasilev Date: Mon, 19 Oct 2026 09:42:19 +0700 Subject: [PATCH] t27t: the t27 trainer, with the int8 tensor-core saving law in t27 (Closes #10789) specs/t27t/t27t.t27 is t27t's law: the RTX 4090 bench of #10360 with every ratio recomputed from its times, the saving f x (1 - 1/s) recomputed cell by cell (16% to 24%, f still an estimate), and pick_path, which sends forward GEMMs to int8 tensor cores via torch._int_mm, keeps gradients in bf16 and never picks dp4a. Runner: gHashTag/igla-coder-gpu. `tri t27t` runs it. AGENTS.md section 8 points at it. Co-Authored-By: Claude Opus 5.5 --- .claude/skills/analog-emulation/SKILL.md | 2 +- AGENTS.md | 27 +++ docs/reports/t27b_expectations.json | 1 + scripts/tri | 4 +- specs/t27t/OWNERS.md | 11 + specs/t27t/t27t.t27 | 276 +++++++++++++++++++++++ 6 files changed, 318 insertions(+), 3 deletions(-) create mode 100644 specs/t27t/OWNERS.md create mode 100644 specs/t27t/t27t.t27 diff --git a/.claude/skills/analog-emulation/SKILL.md b/.claude/skills/analog-emulation/SKILL.md index 14e7424966..4506361afe 100644 --- a/.claude/skills/analog-emulation/SKILL.md +++ b/.claude/skills/analog-emulation/SKILL.md @@ -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 diff --git a/AGENTS.md b/AGENTS.md index 314394b4e4..0910430a24 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -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 diff --git a/docs/reports/t27b_expectations.json b/docs/reports/t27b_expectations.json index 83b8b9a91d..3ebbbd2de9 100644 --- a/docs/reports/t27b_expectations.json +++ b/docs/reports/t27b_expectations.json @@ -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"}, diff --git a/scripts/tri b/scripts/tri index 0f8b3635aa..2e835df73d 100755 --- a/scripts/tri +++ b/scripts/tri @@ -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 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 diff --git a/specs/t27t/OWNERS.md b/specs/t27t/OWNERS.md new file mode 100644 index 0000000000..35c125709b --- /dev/null +++ b/specs/t27t/OWNERS.md @@ -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. diff --git a/specs/t27t/t27t.t27 b/specs/t27t/t27t.t27 new file mode 100644 index 0000000000..35bd47b023 --- /dev/null +++ b/specs/t27t/t27t.t27 @@ -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 }