Repository navigation
t27c PTX backend: ternary x int8 GEMM kernel emitted by a .t27 spec (int8 mma.sync), ptx-run glue #5963
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| # Three checks that existed and that nothing ran. | |
| # | |
| # Each was written after a defect it would have caught, and each has since | |
| # caught a live one -- but only because someone remembered to type it. That is | |
| # not a guard, it is a habit, and the next session does not inherit habits. | |
| # | |
| # check_seal_currency.py 116 seals described output the compiler had | |
| # stopped producing, left by five merged backend | |
| # repairs. spec_hash was stale for ZERO of them, so | |
| # every existing coverage and staleness check was | |
| # green (#3415). | |
| # check_ring_spec_drift.py ten of seventeen rings/* crates name a spec that | |
| # exists; one had drifted and one was two designs | |
| # sharing a vocabulary (#3422). | |
| # ring_spec_differential.py a matching signature is not matching behaviour: | |
| # ring-090 read 16 of 16 identical and disagreed on | |
| # 126 of 1190 cases (#3420, #3424). | |
| # diagnostic_capacity.py clang's default -ferror-limit=20 had been | |
| # censoring every C error total published for this | |
| # repository: 15133 errors, reported as 3849. A | |
| # ratchet on the MEASURING APPARATUS -- it fails if | |
| # an instrument that reported completely starts | |
| # truncating (#3450). | |
| # check_duplicate_declarations.py | |
| # three specs declare 25 top-level names twice and | |
| # `t27c check` exited 0 on all of them. Every | |
| # backend rejects a redeclaration -- measured, not | |
| # assumed: cc `redefinition`, rustc E0428, zig | |
| # `duplicate`, iverilog `already been declared` | |
| # (#3438). | |
| # | |
| # NOT a required context, and it must not become one without the owner: a | |
| # `paths:`-filtered workflow that is required hangs every pull request it does | |
| # not match, forever, on a check that never posts. The `push:` trigger is here | |
| # so master HAS a baseline -- a paths-filtered workflow with no push trigger | |
| # runs zero times on the default branch and its history is empty (#2925 family). | |
| # | |
| # Cost, measured rather than estimated: 12s for the seal scan, 0s for the drift | |
| # scan, ~1s per differential pair, 7s for the duplicate-declaration ratchet, | |
| # under 1s for the diagnostic-capacity ratchet, on top of the compiler build. | |
| name: Spec Guards | |
| on: | |
| pull_request: | |
| paths: | |
| - 'bootstrap/src/**' | |
| - 'specs/**' | |
| - 'rings/**' | |
| - '.trinity/seals/**' | |
| - 'tools/check_seal_currency.py' | |
| - 'tools/check_seal_coverage.py' | |
| - 'tools/seal_baseline.txt' | |
| - 'tools/check_ring_spec_drift.py' | |
| - 'tools/ring_spec_differential.py' | |
| - 'tools/check_duplicate_declarations.py' | |
| - 'tools/duplicate_declarations_baseline.txt' | |
| - 'tools/lowered_collisions_baseline.txt' | |
| - 'tools/published_figures.py' | |
| - 'tools/diagnostic_capacity.py' | |
| - '.github/workflows/spec-guards.yml' | |
| push: | |
| branches: [master] | |
| paths: | |
| - 'bootstrap/src/**' | |
| - 'specs/**' | |
| - 'rings/**' | |
| - '.trinity/seals/**' | |
| - 'tools/check_seal_currency.py' | |
| - 'tools/check_seal_coverage.py' | |
| - 'tools/seal_baseline.txt' | |
| - 'tools/check_ring_spec_drift.py' | |
| - 'tools/ring_spec_differential.py' | |
| - 'tools/check_duplicate_declarations.py' | |
| - 'tools/duplicate_declarations_baseline.txt' | |
| - 'tools/lowered_collisions_baseline.txt' | |
| - 'tools/published_figures.py' | |
| - 'tools/diagnostic_capacity.py' | |
| - '.github/workflows/spec-guards.yml' | |
| workflow_dispatch: | |
| permissions: | |
| contents: read | |
| concurrency: | |
| group: spec-guards-${{ github.ref }} | |
| cancel-in-progress: true | |
| jobs: | |
| spec-guards: | |
| runs-on: ubuntu-latest | |
| steps: | |
| # Depth 2, not 1: the published-figures step compares this commit with the one before it (on a pull | |
| # request the merge ref's parents are the base tip and the PR head). Still a shallow clone, so | |
| # nothing that asks "is this repository shallow" changes its answer. | |
| - uses: actions/checkout@v4 | |
| with: | |
| fetch-depth: 2 | |
| - uses: actions/setup-python@v5 | |
| with: | |
| python-version: '3.12' | |
| # LANG-EN for the specs this change touches: specs/policy/lang_en.t27, generated to C (#7548). | |
| # bootstrap/build.rs no longer reruns on a spec edit, so this is where a changed spec is checked. | |
| - name: No Cyrillic in a changed spec (specs/policy/lang_en.t27) | |
| run: | | |
| d=$(mktemp -d) | |
| printf '#include "gen/c/policy/lang_en.c"\n#include <stdio.h>\nstatic uint8_t b[1<<24];int main(void){size_t n=fread(b,1,sizeof b,stdin);if(n==sizeof b)return 2;unsigned l=first_cyrillic_line(b,n);if(l)printf("%%u\\n",l);return l!=0;}\n' | cc -w -I . -x c -o "$d/gate" - | |
| if printf 'a\n\320\277\n' | "$d/gate" >/dev/null; then echo "negative control passed: the gate is broken"; exit 1; fi | |
| bad=0 | |
| for f in $(git diff --name-only --diff-filter=AMR HEAD^1 HEAD -- 'specs/*.t27' 'specs/*.tri'); do | |
| l=$("$d/gate" < "$f") || { echo "::error file=$f,line=$l::Cyrillic in a spec (LANG-EN)"; bad=1; } | |
| done | |
| exit $bad | |
| - name: Build the compiler every check is measured against | |
| run: cargo build --release -p t27c | |
| # gen/c/claw/gates.c is committed (force-added; gen/ is ignored but | |
| # tracked) and the claw-drift step in gate-topology.yml feeds those | |
| # COMMITTED bytes -- nothing else holds them to their spec. An edit to | |
| # specs/claw/gates.t27 without a regen would drift silently. So would the IGLA board's | |
| # own_model event, which a Railway cron reposts from master's committed JSON (#10321). | |
| # Same shape as t27o-golden's check mode (#8996): regenerate, then diff --exit-code. | |
| - name: committed gen/c/claw/*.c and gen/json/igla/own_model_board.json are what their specs lower to today | |
| run: | | |
| ./target/release/t27c gen-c specs/claw/gates.t27 > gen/c/claw/gates.c | |
| ./target/release/t27c gen-c specs/claw/t27f_gate.t27 > gen/c/claw/t27f_gate.c | |
| ./target/release/t27c gen-c specs/claw/sarif_feed.t27 > gen/c/claw/sarif_feed.c | |
| ./target/release/t27c gen-c specs/claw/suggest_feed.t27 > gen/c/claw/suggest_feed.c | |
| ./target/release/t27c gen-c specs/claw/refresh.t27 > gen/c/claw/refresh.c | |
| ./target/release/t27c gen-c specs/claw/ledger_union.t27 > gen/c/claw/ledger_union.c | |
| ./target/release/t27c gen-rust specs/tri/fmt_source.t27 > bootstrap/gen/rust/tri/fmt_source.rs | |
| ./target/release/t27c gen-json specs/igla/own_model_board.t27 > gen/json/igla/own_model_board.json | |
| git diff --exit-code -- gen/c/claw/gates.c gen/c/claw/t27f_gate.c gen/c/claw/sarif_feed.c gen/c/claw/suggest_feed.c gen/c/claw/refresh.c gen/c/claw/ledger_union.c bootstrap/gen/rust/tri/fmt_source.rs gen/json/igla/own_model_board.json | |
| # #7565: test only the specs this change can affect. specs/ci/affected.t27 picks them (the | |
| # changed specs and every spec that imports one through `use`, transitively); the graph is the | |
| # base tree's edges plus the head tree's, as use_resolve drops an import of a missing file. Each | |
| # pick is measured in both trees by the same t27c: red if it fails more tests than on the base, | |
| # or stops building. RUN_ALL (a change outside specs/) leaves the corpus to corpus-ratchet. | |
| - name: Affected specs -- select with specs/ci/affected.t27, test only those | |
| shell: bash | |
| run: | | |
| set -euo pipefail | |
| t=$PWD/target/release/t27c d=$RUNNER_TEMP/aff; mkdir -p "$d/base" "$d/zig" | |
| curl -fsSL https://ziglang.org/download/0.16.0/zig-x86_64-linux-0.16.0.tar.xz | tar -xJ -C "$d/zig" --strip-components=1 | |
| export PATH="$d/zig:$PATH" | |
| # #8095: what a build may take as already verified (t27c frontier, specs/verified/seal_identity.t27), | |
| # always against the BASE revision's seals (a pull request cannot vouch for itself) and this head's | |
| # test runner. A base spec frontier marks REUSE scores 0 failing tests without being run: 0 is the | |
| # strictest baseline, so a wrong reuse can only turn this step red, never hide a regression. | |
| git archive HEAD^1 specs .trinity/seals | tar -x -C "$d/base" | |
| mkdir -p "$d/base/bootstrap/src" && git show HEAD:bootstrap/src/test_report.rs > "$d/base/bootstrap/src/test_report.rs" | |
| (cd "$d/base" && "$t" frontier --list) > "$d/rebuild.txt"; tail -1 "$d/rebuild.txt" | tee -a "$GITHUB_STEP_SUMMARY" | |
| printf '#include "gen/c/ci/affected.c"\n#include <stdio.h>\n#include <stdlib.h>\nint main(void){size_t c=1<<26;uint8_t*b=malloc(c),*s=malloc(c),*o=calloc(c+1,1);size_t n=fread(b,1,c,stdin);if(n==c)return 2;if(select_affected(b,n,s,c,o,c)==RUN_ALL){puts("RUN_ALL");return 0;}fputs((char*)o,stdout);return 0;}\n' | cc -w -I . -x c -o "$d/sel" - | |
| { git diff --name-status HEAD^1 HEAD; echo --graph; "$t" use-edges; (cd "$d/base" && "$t" use-edges) | |
| echo --corpus; git ls-files 'specs/*.t27' 'specs/*.tri'; } | "$d/sel" > "$d/sel.txt" | |
| if grep -qx RUN_ALL "$d/sel.txt"; then echo "RUN_ALL: a change outside specs/; corpus-ratchet runs the corpus" | tee -a "$GITHUB_STEP_SUMMARY"; exit 0; fi | |
| echo "affected: $(wc -l < "$d/sel.txt") of $(git ls-files 'specs/*.t27' 'specs/*.tri' | wc -l) specs" | tee -a "$GITHUB_STEP_SUMMARY" | |
| score() { T27C_TEST_REPORT_EXIT_ZERO=1 "$t" test-report "$1" --specs-dir specs | awk '/^ BLOCKED/{b=1} /^ FAIL +[0-9]+$/{f=$2} END{print (b ? 1000000 : f+0)}'; } | |
| bad=0 | |
| while IFS= read -r f; do | |
| case $f in *.t27) ;; *) continue ;; esac | |
| h=$(score "$f"); b=0 | |
| if [ -f "$d/base/$f" ] && ! grep -q " $f\$" "$d/rebuild.txt"; then echo "$f: base reused (seal v2, t27c frontier)" | |
| elif [ -f "$d/base/$f" ]; then b=$(cd "$d/base" && score "$f"); fi | |
| echo "$f: head $h, base $b (failing tests; 1000000 = does not build)" | |
| if [ "$h" -gt "$b" ]; then echo "::error file=$f::$h failing vs $b on the base"; bad=1; fi | |
| done < "$d/sel.txt" | |
| exit "$bad" | |
| # The self-checks run FIRST and separately. A checker whose own negative | |
| # control fails cannot be believed when it then reports zero, and running | |
| # them in the same step as the real scan would let a green scan hide a | |
| # broken instrument. | |
| - name: Self-check -- seal currency | |
| run: python3 tools/check_seal_currency.py --self-check | |
| - name: Self-check -- ring/spec drift | |
| run: python3 tools/check_ring_spec_drift.py --self-check | |
| - name: Self-check -- differential harness | |
| run: python3 tools/ring_spec_differential.py --self-check | |
| - name: Self-check -- duplicate declarations | |
| run: python3 tools/check_duplicate_declarations.py --self-check | |
| # The fixtures, before the counts. Two of the first five produced the | |
| # wrong KIND of diagnostic and both read as findings about the compiler. | |
| - name: Self-check -- diagnostic capacity fixtures | |
| run: python3 tools/diagnostic_capacity.py --self-check | |
| - name: Does every seal still describe what the compiler emits | |
| run: python3 tools/check_seal_currency.py | |
| - name: Does every ring still agree with the spec it names | |
| run: python3 tools/check_ring_spec_drift.py | |
| # A RATCHET against a per-spec baseline, not a total. A ceiling on the | |
| # sum is satisfied by zero and says nothing about a spec that grew while | |
| # another shrank. | |
| - name: Does any spec declare one name twice | |
| run: python3 tools/check_duplicate_declarations.py | |
| - name: Published figures -- self-check | |
| run: python3 tools/published_figures.py --self-check | |
| - name: Published figures -- every pinned spec population re-derived | |
| # A figure without its matcher is not reproducible and a figure whose | |
| # unit is implicit is not comparable. This re-derives each from the | |
| # specs, in CODE ONLY, and asks specs/ci/derived_data.t27 for the | |
| # verdict (#7174): red for a matcher that changed meaning without its | |
| # pin, or a pin the tool would not write; a pin that only lags merges is | |
| # a warning and a job summary, never red -- a re-pin cannot win a race | |
| # against every lane that merges specs. | |
| env: | |
| T27C: ./target/release/t27c | |
| run: python3 tools/published_figures.py --check | |
| # A ratchet on the instruments, not on the code: if a compiler that | |
| # reported completely starts truncating, every total measured with it | |
| # silently becomes a floor. | |
| - name: Does every compiler still report what it finds | |
| run: python3 tools/diagnostic_capacity.py | |
| - name: Run each CONVERGED pair on the same inputs | |
| shell: bash | |
| run: | | |
| set -euo pipefail | |
| # The pair list comes from the drift tool, not from this file: a | |
| # hard-coded list goes stale the first time a pair converges. | |
| # | |
| # `while read`, not `mapfile`: mapfile is bash 4+ and macOS ships bash | |
| # 3.2, so this step could not be run locally before pushing it. A | |
| # construct that only the runner can execute is a construct nobody | |
| # tested. The `|| [ -n "$line" ]` is not decoration -- a plain | |
| # `while read` DROPS a final line with no trailing newline. | |
| python3 tools/check_ring_spec_drift.py --converged > /tmp/converged.txt | |
| n=0 | |
| while IFS= read -r line || [ -n "$line" ]; do | |
| [ -z "$line" ] && continue | |
| n=$((n + 1)) | |
| # shellcheck disable=SC2086 | |
| python3 tools/ring_spec_differential.py $line | |
| done < /tmp/converged.txt | |
| if [ "$n" -eq 0 ]; then | |
| echo "::error::no CONVERGED pair -- the differential step measured nothing" | |
| exit 1 | |
| fi | |
| echo "CONVERGED pairs run: $n" |