From 80e0287d0fc7c2273501868abb22147273504f84 Mon Sep 17 00:00:00 2001 From: Vyncint Ng <115854244+vyncint@users.noreply.github.com> Date: Fri, 21 Aug 2026 07:07:36 +0700 Subject: [PATCH] ci(book): gate the dialect page's catalog stamp against catalog.json Closes #1090. `mlir-dialects.md` states its own freshness rule: At catalog SHA-256 `26a9cc13` (the stamp in every `ops/generated/` file header), the dialect holds 570 operations across 42 modules ... If the header stamp no longer starts with `26a9cc13`, the counts on this page predate the catalog you are reading. Nothing enforced that rule, so the rule that would expose stale counts was itself unchecked. The page failed its own test for weeks before #1052 (560 printed against 570 in the tree), and **it is failing again right now**: the catalog hashes to `00372dfe`, five hand-written ops and five catalog entries after that refresh landed. Two hand audits in two batches is the pattern this closes. ## The gate One file hashed, one page grepped, no cargo and no network. The hash is the one `cuda-intrinsics-gen` stamps into every generated file -- `sha256_bytes(catalog_json.as_bytes())` over the file as committed -- so a plain `sha256sum` reproduces it, and the page prints its first eight digits. Three checks rather than one, so each failure says which thing moved: - The generated sources must stamp the hash `catalog.json` actually has. Failing here means the generated outputs are stale, and the message says to run `cuda-intrinsics-gen -- generate` -- not to edit the book. - The page must print exactly one distinct stamp. It names it twice, once stating it and once in the rule that tells a reader what a mismatch means, so updating one and not the other leaves the rule contradicting the statement above it. - The stamps must agree, and the failure prints both sides plus the two commands that re-derive the counts the stamp guards. It refuses to report success from a check that found nothing: a page reworded so no stamp is left fails with that stated, rather than passing vacuously. That path needed the `|| true` on the capture rather than on a later pipeline -- under `pipefail` a grep matching nothing aborted the script before the diagnostic, turning a reworded page into a silent non-zero exit. ## The refresh that makes it pass Five numbers and the stamp, all re-derived from the tree: | claim | was | now | |---|--:|--:| | catalog stamp (x2) | `26a9cc13` | `00372dfe` | | total operations | 570 | 575 | | hand-written operations | 21 | 26 | | hand-written `wgmma` row | 9 | 14 | | catalog entries | 1010 | 1015 | Generated stays 549 across 35 modules and every area row is unchanged, because #1085's five new catalog entries are aliases onto existing structural ops -- exactly what the page's "several intrinsics can share one structural op" sentence predicts. The +5 hand-written are all `wgmma`, from #1075, #1077 and #1079. The `wgmma` row's description moves with its count, as in #1052: it said "the m64n64k16 bf16/f16 shapes", and the module now holds four shape/type pairs -- `m64n64k16` bf16 and f16, `m64n128k16` bf16 (#1079), `m64n64k8` tf32 (#1077). ## Verification - Four mutations, each failing with its own message: a stale stamp, only one of the two stamps updated, the page reworded so none is found, and an edited `catalog.json`. The healthy tree passes. - Every count re-checked against the tree: 26 + 549 = 575, both tables sum to their stated totals (7 rows to 26, 11 rows to 549), and each still lists every module on disk. - All eleven guard scripts pass, `check-spdx-headers.sh` accepts the new file, `cuda-intrinsics-gen -- check` reports outputs current, the workflow parses with all nine jobs on `ubuntu-latest`, and the book builds clean under `sphinx-build -W`. `just check-guards` should list this script too, one line beside the other ten. Its recipe is being restructured in #1102, so that line is left for a follow-up rather than conflicting; no guard enforces that list, so CI is unaffected either way. Signed-off-by: Vyncint Ng <115854244+vyncint@users.noreply.github.com> --- .github/workflows/status-guard.yml | 20 ++++ cuda-oxide-book/compiler/mlir-dialects.md | 12 +-- scripts/check-book-catalog-stamp.sh | 112 ++++++++++++++++++++++ 3 files changed, 138 insertions(+), 6 deletions(-) create mode 100755 scripts/check-book-catalog-stamp.sh diff --git a/.github/workflows/status-guard.yml b/.github/workflows/status-guard.yml index 5d323de528..810e166217 100644 --- a/.github/workflows/status-guard.yml +++ b/.github/workflows/status-guard.yml @@ -112,6 +112,26 @@ jobs: - name: Verify every device call the book shows resolves run: bash scripts/check-book-api-names.sh + # The dialect-nvvm book page prints the catalog SHA-256 and states its own + # freshness rule: if the stamp no longer matches, every count on the page + # predates the catalog being read. Nothing enforced that rule, so the page + # failed its own test for weeks -- 560 operations printed against 570 in the + # tree -- and only a hand audit caught it, twice. Machine-checking the stamp + # turns "the counts are stale" into a red check on the pull request that + # rotated the catalog. The counts stay hand-written; only the stamp is gated, + # because the page derives staleness from it. + book-catalog-stamp: + name: book catalog stamp + runs-on: ubuntu-latest + timeout-minutes: 5 + steps: + - name: Clone repository + uses: actions/checkout@v4 + + # Hashes one file and greps one page: no cargo, no build, no network. + - name: Verify the dialect page's catalog stamp is current + run: bash scripts/check-book-catalog-stamp.sh + # `unit-tests.yml`'s matrix is the only thing deciding which crates get # `cargo test`, and a member missing from it fails nothing: the workflow is # green, `just check` is green, and the tests never ran. #971 found four diff --git a/cuda-oxide-book/compiler/mlir-dialects.md b/cuda-oxide-book/compiler/mlir-dialects.md index 6b8d23669c..02cdc46832 100644 --- a/cuda-oxide-book/compiler/mlir-dialects.md +++ b/cuda-oxide-book/compiler/mlir-dialects.md @@ -296,16 +296,16 @@ they become `call` instructions to `@llvm.nvvm.*` intrinsics. ### Architecture Coverage -At catalog SHA-256 `26a9cc13` (the stamp in every `ops/generated/` file -header), the dialect holds 570 operations across 42 modules, and they come +At catalog SHA-256 `00372dfe` (the stamp in every `ops/generated/` file +header), the dialect holds 575 operations across 42 modules, and they come from two different places. The split is the first thing to know about it, because it decides where -- and whether -- you would add one. If the header -stamp no longer starts with `26a9cc13`, the counts on this page predate the +stamp no longer starts with `00372dfe`, the counts on this page predate the catalog you are reading. **Hand-written**, directly under `crates/dialect-nvvm/src/ops/`. These are the ops with bespoke verification or lowering that the intrinsic catalog does not -describe. There are seven modules and 21 operations: +describe. There are seven modules and 26 operations: | Module | Description | Ops | | :-------- | :---------------------------------------------------------- | --: | @@ -315,12 +315,12 @@ describe. There are seven modules and 21 operations: | `debug` | `assertfail`, `vprintf` | 2 | | `grid` | Cooperative `grid_sync` | 1 | | `memory` | Generic-to-shared address conversion with a byte offset | 1 | -| `wgmma` | Warpgroup MMA descriptors and the m64n64k16 bf16/f16 shapes | 9 | +| `wgmma` | Warpgroup MMA descriptors; bf16/f16 at m64n64k16, bf16 at m64n128k16, tf32 at m64n64k8 | 14 | **Generated**, under `ops/generated/`, from `intrinsics/catalog.json` by `cuda-intrinsics-gen`. Every file there opens with `// @generated ... DO NOT EDIT.`, and editing one by hand is undone by the next generator run. This is -the large majority -- 35 modules and 549 operations, resolved from 1010 catalog +the large majority -- 35 modules and 549 operations, resolved from 1015 catalog entries, since several intrinsics can share one structural op: | Area | Modules | Ops | diff --git a/scripts/check-book-catalog-stamp.sh b/scripts/check-book-catalog-stamp.sh new file mode 100755 index 0000000000..1722fe2426 --- /dev/null +++ b/scripts/check-book-catalog-stamp.sh @@ -0,0 +1,112 @@ +#!/usr/bin/env bash +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# Verify the dialect-nvvm book page's catalog stamp still matches +# intrinsics/catalog.json. +# +# The page states its own freshness rule: it prints the catalog SHA-256 and +# says that if the stamp no longer matches, every count on the page predates +# the catalog being read. Nothing checked it. The page failed that test for +# weeks -- 560 operations printed against 570 in the tree -- and only a hand +# audit caught it, twice (#1052 and the batch before it). +# +# So the counts rot silently while the rule that would expose them is itself +# unenforced. This makes the rule executable: the stamp is machine-checked, +# which turns "the counts are stale" into a red check on the op-adding pull +# request rather than a discovery months later. The counts stay hand-written; +# only the stamp needs a gate, because the page derives staleness from it. +# +# The hash is the one `cuda-intrinsics-gen` stamps into every generated file: +# `sha256_bytes(catalog_json.as_bytes())` over the file as committed, so a +# plain sha256sum of catalog.json reproduces it. The page prints the first +# eight hex digits, which is what this compares. +# +# Run this after any change to intrinsics/catalog.json. +set -euo pipefail + +export LC_ALL=C + +cd "$(dirname "$0")/.." + +CATALOG=intrinsics/catalog.json +PAGE=cuda-oxide-book/compiler/mlir-dialects.md +GENERATED_DIR=crates/dialect-nvvm/src/ops/generated + +test -f "${CATALOG}" +test -f "${PAGE}" + +actual="$(sha256sum "${CATALOG}" | cut -c1-8)" + +# Self-test: a truncated or reworded hash pipeline must not read as agreement. +case "${actual}" in + [0-9a-f][0-9a-f][0-9a-f][0-9a-f][0-9a-f][0-9a-f][0-9a-f][0-9a-f]) ;; + *) + echo "error: could not hash ${CATALOG} (got '${actual}')" >&2 + exit 1 + ;; +esac + +# Cross-check against the generated sources, which carry the full hash in +# their headers. If the generator's notion of the catalog hash ever stops +# being a plain sha256sum of the file, this fails here rather than silently +# comparing the page against the wrong number. +stamped="$(grep -ohE 'catalog SHA-256: [0-9a-f]{64}' "${GENERATED_DIR}"/*.rs | + sed 's/.*: //' | sort -u)" +if [ "$(printf '%s\n' "${stamped}" | grep -c .)" -ne 1 ]; then + echo "error: ${GENERATED_DIR} carries more than one catalog SHA-256:" >&2 + printf ' %s\n' ${stamped} >&2 + echo " regenerate with 'cargo run -p cuda-intrinsics-gen -- generate'" >&2 + exit 1 +fi +if [ "${stamped#"${actual}"}" = "${stamped}" ]; then + echo "error: ${CATALOG} hashes to ${actual}, but the generated sources stamp" >&2 + echo " ${stamped}. The generated outputs are stale; run" >&2 + echo " 'cargo run -p cuda-intrinsics-gen -- generate'." >&2 + exit 1 +fi + +# The page names the stamp twice: once stating it, once in the rule that tells +# a reader what a mismatch means. Both have to move together, or the rule +# contradicts the statement above it. +# `|| true` on the capture, not on a later pipeline: with pipefail a grep that +# matches nothing would abort the script before the diagnostics below, turning +# "the page was reworded" into a silent non-zero exit. +stamps="$(grep -oE '`[0-9a-f]{8}`' "${PAGE}" | tr -d '`' || true)" +count="$(printf '%s' "${stamps}" | grep -c . || true)" +found="$(printf '%s\n' "${stamps}" | sort -u | grep -c . >/dev/null 2>&1 && + printf '%s\n' "${stamps}" | sort -u || true)" + +if [ "${count}" -lt 2 ]; then + echo "error: found ${count} catalog stamps in ${PAGE}, expected at least 2" >&2 + echo " (the statement and the staleness rule). The page was reworded;" >&2 + echo " refusing to report success from a check that found nothing." >&2 + exit 1 +fi + +if [ "$(printf '%s\n' "${found}" | grep -c .)" -ne 1 ]; then + echo "error: ${PAGE} prints more than one catalog stamp:" >&2 + printf ' %s\n' ${found} >&2 + echo " the statement and the staleness rule must name the same hash." >&2 + exit 1 +fi + +if [ "${found}" != "${actual}" ]; then + cat >&2 <