Repository navigation
t27c PTX backend: ternary x int8 GEMM kernel emitted by a .t27 spec (int8 mma.sync), ptx-run glue #5863
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
| # This repository has 35 workflows and, until this one, not a single one built | |
| # the `tri` crate. The consequence was exactly what you would expect: `cargo | |
| # build -p tri` on main failed because dlc10 embeds a bitstream with | |
| # `include_bytes!` that was committed to a feature branch and never to main, so | |
| # nobody could build the CLI from a clean checkout. An autonomous loop went on | |
| # committing into cli/tri/src/ the whole time. | |
| # | |
| # A crate nothing builds is a crate nobody can use, and the breakage is silent | |
| # because there is no signal to be red. | |
| name: cli-tri | |
| on: | |
| push: | |
| branches: [main, master] | |
| paths: | |
| - 'cli/**' | |
| - '.claude/skills/**' | |
| - 'Cargo.toml' | |
| - 'Cargo.lock' | |
| # `tri census pin --gate` pins `gates quiet` and `gates shell`, whose | |
| # subject is this directory. A trigger narrower than the subject cannot | |
| # fire on the change it exists for -- and worse, the ledger would go | |
| # stale on a workflow-only commit and the NEXT cli/** commit would fail, | |
| # blaming an author who changed nothing. Priced over the last 200 commits | |
| # on master: 24 touch .github/workflows/, 2 of those already touch cli/, | |
| # so this adds this job on 22 of 200 commits (11%). | |
| - '.github/workflows/**' | |
| - 'scripts/ci/test_a_moved_census_names_its_file.py' | |
| # NO `paths:` here, deliberately. A required context that does not POST is not | |
| # "passed" -- it is pending forever, so a paths-filtered workflow cannot be | |
| # required without hanging every PR that skips it. Measured over the last 40 | |
| # merged PRs: 15 of them (38%) touched none of the paths this filter listed. | |
| # The `push:` filter above is kept: a master baseline only needs taking when | |
| # the subject moves. | |
| # | |
| # Cost of running on every PR, measured rather than estimated: +38% of runs on | |
| # a job whose median is 2.9 minutes, green 58 of its last 60 (43 of 45 on | |
| # pull_request). A gate that reddens rarely is one that can be required. | |
| pull_request: | |
| workflow_dispatch: | |
| # A newer push or PR edit makes this run's verdict stale: cancel it (queue jam 2026-10-09). | |
| concurrency: | |
| group: cli-tri-${{ github.ref }} | |
| cancel-in-progress: true | |
| jobs: | |
| build: | |
| # A UNIQUE context name. GitHub matches a required check by CONTEXT, and | |
| # `build-vivado-image.yml` and `coq-kernel.yml` both emit `build` too -- so | |
| # "require build" names three workflows and can be satisfied by the wrong | |
| # one, or hang on a job that does not run. First of two things that had to | |
| # be true before this gate could be required at all. | |
| name: cli-tri | |
| runs-on: ubuntu-latest | |
| timeout-minutes: 25 | |
| steps: | |
| # chips/phi is a gitlink (160000) to gHashTag/tt-trinity-phi, and | |
| # checkout's default is submodules: false -- so the directory the last | |
| # step points `tri rtl check` at was EMPTY on the runner, and the CLI died | |
| # with a bare "No such file or directory (os error 2)". The binary was | |
| # fine; it had nothing to read. That step had never executed once in this | |
| # workflow's history, because the two causes before it (#2303 yosys | |
| # ordering, #2305 dry_run_sweep_ok) aborted the job first. | |
| # | |
| # All three chips/* submodules are public, so the default GITHUB_TOKEN -- | |
| # which is scoped to this repo alone -- can clone them; ~10.2 MB total. | |
| # None of the three has a nested .gitmodules, so `true` is the minimal | |
| # setting that works and `recursive` would buy nothing. | |
| - uses: actions/checkout@v4 | |
| with: | |
| submodules: true | |
| - uses: dtolnay/rust-toolchain@stable | |
| # From a clean checkout, with nothing copied in by hand. That is the | |
| # condition that was broken and the only one worth asserting. | |
| - name: cargo build -p tri --all-targets | |
| # Tests are targets too. The plain build stayed green twice while the | |
| # test build was broken on master (#2227, #2236) -- every PR inherited | |
| # a break no gate reported. --all-targets closes that gap. | |
| run: cargo build -p tri --all-targets | |
| # #5482. `tri hooks pre-push` runs the NOW gate's script in a | |
| # contributor's own clone (on request only since #5935; `tri gates | |
| # preview` did too until then), and until #5482 every pass there | |
| # rewrote .git/config and appended to .gitattributes. This job has the | |
| # binary, so the `tri` caller is asked here, in a scratch repository; | |
| # untrusted-input-gate.yml runs the script's own arms without one. | |
| # Placed before the tests so a red test does not skip it. | |
| - name: The NOW gate's tri caller writes nothing to the clone | |
| run: python3 scripts/ci/test_now_gate_writes_nothing.py --tri ./target/debug/tri | |
| # yosys has to be on PATH BEFORE the tests, not after them. This install | |
| # used to sit inside the last step's run block, which is ordered after | |
| # `cargo test -p tri`, and ubuntu-latest ships no yosys. So | |
| # fpga::tests::test_smoke_gate_json_synthetic_verify_lean -- which calls | |
| # the real smoke_gate(), whose `passed` flag ANDs in `yosys_ok` -- hit | |
| # "[smoke-gate] SKIP: yosys not on PATH" and failed on every single run. | |
| # 11 runs on master, 11 failures, no green in this workflow's history: | |
| # the job could not pass by construction, and the job never even reached | |
| # the install because a failed step aborts the rest. | |
| # | |
| # The test is right to demand the binary. A smoke gate that "passes" | |
| # while skipping its only synthesis phase asserts nothing -- the same | |
| # absence-is-not-a-value defect closed in #2285 and #2287. Fix the | |
| # ordering, not the assertion. | |
| - name: install yosys | |
| run: sudo apt-get update -qq && sudo apt-get install -y -qq yosys | |
| # zig before the tests, for the same reason as yosys above: `tri mutate | |
| # spec` compiles a mutant with zig and runs its tests, and its own tests | |
| # (#7148) do both on small lowered specs. Without zig they would fail, and | |
| # rightly: a hang or a compile error told apart from a kill without | |
| # running zig would assert nothing. Same pin as bootstrap-tests.yml. | |
| - name: Install zig | |
| env: | |
| ZIG_VERSION: "0.16.0" | |
| run: | | |
| set -euo pipefail | |
| curl -fsSL "https://ziglang.org/download/${ZIG_VERSION}/zig-x86_64-linux-${ZIG_VERSION}.tar.xz" -o /tmp/zig.tar.xz | |
| mkdir -p "$RUNNER_TEMP/zig" | |
| tar -xJf /tmp/zig.tar.xz -C "$RUNNER_TEMP/zig" --strip-components=1 | |
| echo "$RUNNER_TEMP/zig" >> "$GITHUB_PATH" | |
| "$RUNNER_TEMP/zig/zig" version | |
| # --nocapture --test-threads=1, and the log tee'd out, so the ratchet | |
| # below can measure something `cargo test` otherwise throws away. | |
| # | |
| # Three smoke-gate tests need a Lean toolchain that no workflow installs | |
| # (#2370). They do not use #[ignore] -- they println! a reason and return. | |
| # So libtest counts them as PASSED and the summary line reads | |
| # "0 ignored", and `cargo test` CAPTURES the stdout of a passing test, so | |
| # the reason never reaches the log at all. Without --nocapture there is | |
| # nothing to measure. Without --test-threads=1 the reason cannot be | |
| # attributed to a test, because libtest lands it on the "test NAME ... " | |
| # line it has already opened and parallel workers interleave those. | |
| # | |
| # set -o pipefail keeps tee from swallowing a real test failure. | |
| - name: cargo test -p tri | |
| run: | | |
| set -o pipefail | |
| cargo test -p tri -- --nocapture --test-threads=1 2>&1 | tee "$RUNNER_TEMP/tri-test.log" | |
| # A green run of this job is consistent with those three never having | |
| # executed. This step puts the count and the reasons in the job summary so | |
| # a reader of a green run cannot mistake it for coverage, and ratchets the | |
| # set: it FAILS if the set grows. A fourth test starting to skip silently | |
| # is the regression nobody would otherwise see. | |
| # | |
| # It deliberately does not fail on today's three. Turning existing skips | |
| # red is the owner's policy decision and a skip is not a failure; #2370 | |
| # stays open for the real fix, which is installing the toolchain. | |
| # | |
| # The baseline is emitted by this same script (--emit-baseline) from a real | |
| # test log. It is never hand-written: a hand-typed "3" would restate what | |
| # we already know and would not survive a rename. | |
| - name: runtime skip ratchet | |
| if: always() | |
| run: python3 cli/tri/skipwatch.py "$RUNNER_TEMP/tri-test.log" cli/tri/skip-baseline.txt | |
| # `tri rtl check` reports numbers, and a binary that runs but reports | |
| # nothing would pass the two steps above. yosys is already on PATH from | |
| # the install step above, so the command is exercised against a real | |
| # design rather than assumed. | |
| # `tri skill check` finds two sections sharing a number, prints PROBLEMS | |
| # and exits 1 -- and until now NOTHING called it. `grep -rn "skill check"` | |
| # across .github/workflows/, scripts/ and tools/ returned zero. | |
| # | |
| # It was not hypothetical: two sections were both numbered 504 on master, | |
| # landed 26 minutes apart by two sessions, each appending above the highest | |
| # number it had read. The checker would have caught it on the second PR. | |
| # | |
| # This job already builds the binary, so the marginal cost is one command. | |
| # Its `paths:` grows to .claude/skills/** in the same commit, because a | |
| # gate whose trigger is narrower than its subject cannot fire on the change | |
| # it exists for. | |
| - name: No two skill sections share a number | |
| run: ./target/debug/tri skill check | |
| # A census that moves silently is a number nobody re-read. Measured over | |
| # the 39 most recent transitions on master with one fixed instrument: | |
| # 8 commits moved a census and only 4 said so -- and one of the silent | |
| # four had made 45 of 50 red workflows in a sibling repository invisible | |
| # (#3176). The A/B that finds it is two commands, and nothing ran them. | |
| # | |
| # This is cheap for a reason that is structural rather than hopeful: each | |
| # pinned census's POPULATION IS A DIRECTORY, so it cannot move unless a | |
| # file in that directory changes. All 8 of the 8 observed moves had edited | |
| # that census's own subject; not one was a side effect. So the re-bless | |
| # falls only on commits already working in that area. | |
| # | |
| # Deliberately excluded: `dead` and `unmeasured` read the GitHub API, so | |
| # their answer moves when the world moves and pinning them would redden | |
| # this gate on somebody else's push. A test refuses their addition. | |
| - name: No census moved without saying so | |
| run: ./target/debug/tri census pin --gate | |
| # The gate above prints the FIRST line that differs and stops, which says a | |
| # census moved and not what moved it. On #5787 that first line was `files | |
| # read 22 -> 23` and the cause sat further down (`0 do not -> 1 do not`: a | |
| # helper spelling "lower bound" in lowercase), and the PR read "open" for six | |
| # ticks while this job was red on it. `tri census explain` lists every moved | |
| # line and names the changed file behind each by putting that file back | |
| # alone in a scratch copy. This step proves it names the mover and not the | |
| # bystander, writes nothing, and leaves no copy behind (#5825). | |
| - name: A moved census names the file that moved it | |
| run: python3 scripts/ci/test_a_moved_census_names_its_file.py --tri ./target/debug/tri | |
| # `cargo test` cannot report this: a test that lost its `#[test]` and one | |
| # that gained a second CANCEL in the total, so the suite count is | |
| # unchanged and green. Commit f7c1ff5 shipped exactly that pair -- one | |
| # test stopped running, another was registered twice -- and the pass that | |
| # wrote it verified "11 attributes against 11 functions" and read a match. | |
| - name: Every test runs, and runs once | |
| run: ./target/debug/tri gates tests --gate | |
| - name: the CLI actually produces a report | |
| run: | | |
| set -uo pipefail | |
| ./target/debug/tri rtl check chips/phi --json > /tmp/r.json 2>/tmp/r.err || true | |
| cat /tmp/r.err | head -5 | |
| N=$(python3 -c "import json;print(len(json.load(open('/tmp/r.json'))['checks']))" 2>/dev/null || echo 0) | |
| echo "verdict lines: $N" | |
| if [ "${N:-0}" -lt 5 ]; then | |
| echo "::error::tri rtl check emitted $N verdicts; five checks should each emit one" | |
| exit 1 | |
| fi |