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 <