nihalpasham · Aug 21, 2026 · Aug 21, 2026
@@ -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

@@ -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 |

@@ -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 <<EOF
+error: ${PAGE} is stale.
+
+  page stamp:      ${found}
+  ${CATALOG}: ${actual}
+
+By the page's own rule, every count on it now predates the catalog. Update the
+stamp in both places and re-derive the counts it guards:
+
+  total operations   grep -c '^pub struct .*Op;' ${GENERATED_DIR}/*.rs \\
+                       crates/dialect-nvvm/src/ops/*.rs
+  catalog entries    jq '.intrinsics | length' ${CATALOG}
+
+then check that both tables still sum to their stated totals.
+EOF
+    exit 1
+fi
+
+echo "OK: ${PAGE} stamps ${actual}, matching ${CATALOG} and ${GENERATED_DIR}."