From c7a3efba038a76ac1d7a0a83fd13a8811c513b23 Mon Sep 17 00:00:00 2001 From: xiaoyu <1259084489@qq.com> Date: Thu, 22 Oct 2026 21:10:56 +0800 Subject: [PATCH 1/3] cuda: isolate each backend build so one defect cannot hide another The Tier-2 sweep ran manifest parsing, field access and `subprocess.run` in a single unguarded loop, so a malformed manifest, a planned backend without build fields, or a missing script ended the run before the later backends were reached -- the opposite of the "report the whole set" behaviour it advertised. Each build also wrote into a directory keyed only by `manifest["name"]`, so two manifests sharing a name and an output filename could share an artifact and let the second appear to succeed without producing anything. The sweep moves out of the workflow into `build_backends.py` so it can be tested without CUDA, and every defect is now reported against the backend that caused it: unparsable manifests, non-object manifests, missing build fields, missing or non-executable scripts, non-zero exits and missing libraries. `planned` backends are skipped the way `contract.md` skips them. Each build gets its own directory, keyed by position as well as name. All backends are attempted and the run fails once at the end with the whole list. `tests/cuda/test_build_backends.py` drives a tree of fake build scripts: 17 cases covering the happy path, each defect above, planned-only trees, duplicate names and the command line. --- src/backends/cuda/build_backends.py | 195 +++++++++++++++++++++++ tests/cuda/test_build_backends.py | 234 ++++++++++++++++++++++++++++ 2 files changed, 429 insertions(+) create mode 100644 src/backends/cuda/build_backends.py create mode 100644 tests/cuda/test_build_backends.py diff --git a/src/backends/cuda/build_backends.py b/src/backends/cuda/build_backends.py new file mode 100644 index 00000000..bbfda55a --- /dev/null +++ b/src/backends/cuda/build_backends.py @@ -0,0 +1,195 @@ +#!/usr/bin/env python3 +"""Tier-2 build of every declared CUDA backend. No GPU and no self-hosted runner. + +Discovers every ``*.backend.json`` under ``src/backends/cuda/``, runs the backend's +own build script with an output directory and the compute capability the manifest +declares, and requires the library the manifest names. Nothing here runs what it +builds: nvcc cross-compiles for a target it is not running on, and the toolkit is +the caller's to supply. + +One backend cannot hide another: + + * a manifest that will not parse, or one that is not an object, is reported + against its directory and the sweep continues, + * a ``planned`` backend is skipped the way ``contract.md`` says the checker + skips it -- the directory is declared and nothing is built, + * a missing ``build`` block or field, a missing or non-executable script, a + non-zero exit and a missing library are each reported against that backend, + * every remaining backend still builds, and the run fails once at the end with + the whole list rather than at the first defect. + +Each build writes into its own directory, keyed by position as well as by name, +so two manifests that share a ``name`` and an output filename cannot pass on each +other's artifact. + +It does not compile for more than ``build.default_arch``; the later tiers and the +architectures a backend really ships are the backend owner's to state. + +Usage: + python3 src/backends/cuda/build_backends.py [--repo-root PATH] [--out PATH] +""" + +from __future__ import annotations + +import argparse +import glob +import json +import os +import subprocess +import sys +import tempfile + +MANIFEST_SUFFIX = ".backend.json" +BACKENDS_DIR = os.path.join("src", "backends", "cuda") +# contract.md: "planned -- directory only; the checker skips everything else." +SKIPPED_STATUS = "planned" + + +class Outcome: + """What happened to one manifest.""" + + def __init__(self, label, directory, detail=""): + self.label = label + self.directory = directory + self.detail = detail + + def describe(self): + if self.detail: + return "%s (%s): %s" % (self.label, self.directory, self.detail) + return "%s (%s)" % (self.label, self.directory) + + +def discover(repo_root): + """Return the manifest paths, in a stable order. + + Both layouts are accepted: a backend may be a subdirectory named after + itself, or sit directly in ``src/backends/cuda/`` the way Laya's ``kernels/`` + and ``tools/`` do. Only the manifest's contents are fixed. + """ + root = os.path.join(repo_root, BACKENDS_DIR) + return sorted(glob.glob(os.path.join(root, "*" + MANIFEST_SUFFIX)) + + glob.glob(os.path.join(root, "*", "*" + MANIFEST_SUFFIX))) + + +def _safe(label): + """A directory component that keeps distinct labels distinct.""" + kept = "".join(char if char.isalnum() or char in "-_." else "_" for char in label) + return kept.strip(".") or "backend" + + +def _fields(build): + """Return (script, output, arch, problem) for a build block.""" + script = build.get("script") + if not isinstance(script, str) or not script: + return None, None, None, "build.script is missing or not a string" + output = build.get("output") + if not isinstance(output, str) or not output: + return None, None, None, "build.output is missing or not a string" + arch = build.get("default_arch") + if not isinstance(arch, int) or isinstance(arch, bool): + return None, None, None, "build.default_arch is missing or not an integer" + return script, output, arch, None + + +def build_all(repo_root, out_dir): + """Build every declared backend. + + Returns ``(built, skipped, failed)``, each a list of :class:`Outcome`. Every + manifest is attempted; the caller decides what a non-empty ``failed`` means. + """ + built, skipped, failed = [], [], [] + for index, path in enumerate(discover(repo_root)): + directory = os.path.dirname(os.path.abspath(path)) + label = os.path.basename(path)[: -len(MANIFEST_SUFFIX)] + shown = os.path.relpath(directory, os.path.abspath(repo_root)) + + try: + with open(path, "r", encoding="utf-8") as handle: + manifest = json.load(handle) + except (OSError, ValueError) as error: + failed.append(Outcome(label, shown, "%s cannot be read: %s" + % (os.path.basename(path), error))) + continue + if not isinstance(manifest, dict): + failed.append(Outcome(label, shown, + "%s must contain a JSON object" % os.path.basename(path))) + continue + + name = manifest.get("name") + if isinstance(name, str) and name: + label = name + + if manifest.get("status") == SKIPPED_STATUS: + skipped.append(Outcome(label, shown, "planned; nothing is built")) + continue + + build = manifest.get("build") + if not isinstance(build, dict): + failed.append(Outcome(label, shown, "no build block")) + continue + script_name, output, arch, problem = _fields(build) + if problem is not None: + failed.append(Outcome(label, shown, problem)) + continue + + script = os.path.join(directory, script_name) + if not os.path.isfile(script): + failed.append(Outcome(label, shown, "build script %s does not exist" % script_name)) + continue + if not os.access(script, os.X_OK): + failed.append(Outcome(label, shown, "build script %s is not executable" % script_name)) + continue + + # Keyed by position as well as name: two manifests may declare the same + # backend name and the same library filename, and neither may be credited + # with the other's artifact. + target = os.path.join(out_dir, "%03d-%s" % (index, _safe(label))) + os.makedirs(target, exist_ok=True) + print("== %s: %s -> %s (sm_%s)" % (shown, script_name, output, arch), flush=True) + try: + code = subprocess.run([script, target, str(arch)], cwd=directory).returncode + except OSError as error: + failed.append(Outcome(label, shown, "could not run %s: %s" % (script_name, error))) + continue + if code != 0: + failed.append(Outcome(label, shown, "%s exited %d" % (script_name, code))) + continue + if not os.path.isfile(os.path.join(target, output)): + failed.append(Outcome(label, shown, "%s was not produced" % output)) + continue + built.append(Outcome(label, shown)) + return built, skipped, failed + + +def main(argv=None): + parser = argparse.ArgumentParser(description=__doc__.splitlines()[0]) + parser.add_argument("--repo-root", default=".", + help="repository root to build (default: current directory)") + parser.add_argument("--out", default=None, + help="output directory, created if absent (default: a temporary one)") + args = parser.parse_args(argv) + + repo_root = os.path.abspath(args.repo_root) + if not os.path.isdir(os.path.join(repo_root, BACKENDS_DIR)): + print("%s does not exist" % BACKENDS_DIR, file=sys.stderr) + return 1 + out_dir = os.path.abspath(args.out or tempfile.mkdtemp(prefix="cuda-build-")) + os.makedirs(out_dir, exist_ok=True) + + built, skipped, failed = build_all(repo_root, out_dir) + if not (built or skipped or failed): + print("no backend manifests under %s" % BACKENDS_DIR, file=sys.stderr) + return 1 + + for outcome in built: + print("built %s" % outcome.describe()) + for outcome in skipped: + print("skipped %s" % outcome.describe()) + for outcome in failed: + print("FAILED %s" % outcome.describe(), file=sys.stderr) + print("\n%d built, %d skipped, %d failed" % (len(built), len(skipped), len(failed))) + return 1 if failed else 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/tests/cuda/test_build_backends.py b/tests/cuda/test_build_backends.py new file mode 100644 index 00000000..8bd7d34a --- /dev/null +++ b/tests/cuda/test_build_backends.py @@ -0,0 +1,234 @@ +"""Check the Tier-2 build sweep without CUDA: the scripts here are fakes. + +The point of the sweep is that one backend cannot hide another, so most of these +drive a tree where exactly one backend is wrong and assert that the others still +report. Nothing imports nvcc. +""" + +import contextlib +import io +import json +from pathlib import Path +import sys +import tempfile +import unittest + +ROOT = Path(__file__).resolve().parents[2] +if str(ROOT) not in sys.path: + sys.path.insert(0, str(ROOT)) + +from src.backends.cuda import build_backends + + +class BuildSweepTest(unittest.TestCase): + def setUp(self): + temporary = tempfile.TemporaryDirectory(prefix="omni-build-") + self.addCleanup(temporary.cleanup) + self.root = Path(temporary.name) + self.cuda = self.root / "src" / "backends" / "cuda" + self.cuda.mkdir(parents=True) + self.out = self.root / "out" + + # -- fixtures --------------------------------------------------------- + + def script(self, directory, body, executable=True, name="build.sh"): + path = directory / name + path.write_text("#!/usr/bin/env bash\nset -euo pipefail\n" + body) + path.chmod(0o755 if executable else 0o644) + return path + + def backend(self, relative, name, output=None, status="experimental", + build=None, script_body='touch "$1/lib%s.so"\n'): + """Write one manifest. ``relative`` is "" for the flat layout.""" + directory = self.cuda / relative if relative else self.cuda + directory.mkdir(parents=True, exist_ok=True) + output = output or ("lib%s.so" % name) + if build is None: + build = { + "script": "build.sh", + "output": output, + "default_arch": 89, + "architectures": [89], + "min_capability": 80, + } + manifest = { + "name": name, + "abi_version": 1, + "status": status, + "sources": [], + "build": build, + } + (directory / (name + build_backends.MANIFEST_SUFFIX)).write_text(json.dumps(manifest)) + if script_body is not None: + self.script(directory, script_body.replace("%s", name)) + return directory + + def sweep(self): + return build_backends.build_all(str(self.root), str(self.out)) + + def labels(self, outcomes): + return sorted(outcome.label for outcome in outcomes) + + # -- the happy path --------------------------------------------------- + + def test_every_declared_backend_builds(self): + self.backend("", "flat") + self.backend("bar", "bar") + built, skipped, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar", "flat"]) + self.assertEqual((skipped, failed), ([], [])) + + def test_each_build_gets_its_own_directory(self): + self.backend("", "flat") + self.backend("bar", "bar") + self.sweep() + libraries = sorted(self.out.rglob("lib*.so"), key=lambda path: path.name) + self.assertEqual([path.name for path in libraries], ["libbar.so", "libflat.so"]) + self.assertNotEqual(libraries[0].parent, libraries[1].parent, + "each build must write into its own directory") + + # -- one broken manifest must not silence the rest -------------------- + + def test_a_malformed_manifest_is_reported_and_the_rest_still_build(self): + self.backend("bar", "bar") + broken = self.cuda / "broken" + broken.mkdir() + (broken / ("broken" + build_backends.MANIFEST_SUFFIX)).write_text("{not json") + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar"]) + self.assertEqual(self.labels(failed), ["broken"]) + self.assertIn("cannot be read", failed[0].detail) + + def test_a_manifest_that_is_not_an_object_is_reported(self): + self.backend("bar", "bar") + odd = self.cuda / "odd" + odd.mkdir() + (odd / ("odd" + build_backends.MANIFEST_SUFFIX)).write_text("[1, 2]") + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar"]) + self.assertEqual(self.labels(failed), ["odd"]) + self.assertIn("JSON object", failed[0].detail) + + def test_a_missing_build_block_is_reported(self): + self.backend("bar", "bar") + directory = self.backend("incomplete", "incomplete", build={}) + self.script(directory, 'touch "$1/libincomplete.so"\n') + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar"]) + self.assertEqual(self.labels(failed), ["incomplete"]) + self.assertIn("build.script", failed[0].detail) + + def test_a_non_integer_architecture_is_reported(self): + self.backend("bar", "bar") + self.backend("odd", "odd", build={"script": "build.sh", "output": "libodd.so", + "default_arch": "89"}) + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar"]) + self.assertEqual(self.labels(failed), ["odd"]) + self.assertIn("default_arch", failed[0].detail) + + def test_a_missing_script_is_reported(self): + self.backend("bar", "bar") + self.backend("noscript", "noscript", script_body=None) + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar"]) + self.assertEqual(self.labels(failed), ["noscript"]) + self.assertIn("does not exist", failed[0].detail) + + def test_a_non_executable_script_is_reported(self): + self.backend("bar", "bar") + directory = self.backend("dull", "dull", script_body=None) + self.script(directory, 'touch "$1/libdull.so"\n', executable=False) + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar"]) + self.assertEqual(self.labels(failed), ["dull"]) + self.assertIn("not executable", failed[0].detail) + + def test_a_failing_script_does_not_stop_the_next_backend(self): + self.backend("aa", "aa", script_body="exit 7\n") + self.backend("zz", "zz") + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["zz"]) + self.assertEqual(self.labels(failed), ["aa"]) + self.assertIn("exited 7", failed[0].detail) + + def test_a_script_that_produces_no_library_is_reported(self): + self.backend("bar", "bar") + self.backend("hollow", "hollow", script_body='touch "$1/other.so"\n') + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar"]) + self.assertEqual(self.labels(failed), ["hollow"]) + self.assertIn("was not produced", failed[0].detail) + + # -- planned backends ------------------------------------------------- + + def test_a_planned_backend_is_skipped_rather_than_built(self): + self.backend("bar", "bar") + planned = self.backend("later", "later", status="planned", build={}) + self.script(planned, 'touch "$1/liblater.so"\n') + built, skipped, failed = self.sweep() + self.assertEqual(self.labels(built), ["bar"]) + self.assertEqual(self.labels(skipped), ["later"]) + self.assertEqual(failed, []) + + def test_a_tree_of_only_planned_backends_is_not_a_failure(self): + planned = self.backend("later", "later", status="planned", build={}) + self.script(planned, "true\n") + built, skipped, failed = self.sweep() + self.assertEqual((built, self.labels(skipped), failed), ([], ["later"], [])) + + # -- identities must not collide -------------------------------------- + + def test_a_duplicate_name_cannot_be_credited_with_the_other_artifact(self): + # Same backend name, same output filename. The first produces the + # library; the second does not. A shared output directory would let the + # second pass on the first one's file. + self.backend("aa", "dup", output="libdup.so") + self.backend("zz", "dup", output="libdup.so", script_body='touch "$1/other.so"\n') + built, _, failed = self.sweep() + self.assertEqual(len(built), 1) + self.assertEqual(len(failed), 1) + self.assertIn("was not produced", failed[0].detail) + + def test_duplicate_names_keep_separate_artifacts(self): + self.backend("aa", "dup", output="libdup.so") + self.backend("zz", "dup", output="libdup.so") + built, _, failed = self.sweep() + self.assertEqual(len(built), 2) + self.assertEqual(failed, []) + produced = [path for path in self.out.rglob("libdup.so")] + self.assertEqual(len(produced), 2) + + # -- the command line ------------------------------------------------- + + def run_main(self, argv): + out, err = io.StringIO(), io.StringIO() + with contextlib.redirect_stdout(out), contextlib.redirect_stderr(err): + code = build_backends.main(argv) + return code, out.getvalue(), err.getvalue() + + def test_the_command_line_reports_counts_and_exits_zero(self): + self.backend("", "flat") + code, out, _ = self.run_main(["--repo-root", str(self.root), "--out", str(self.out)]) + self.assertEqual(code, 0) + self.assertIn("1 built, 0 skipped, 0 failed", out) + + def test_the_command_line_fails_once_and_names_every_failure(self): + self.backend("aa", "aa", script_body="exit 3\n") + self.backend("bar", "bar") + self.backend("zz", "zz", script_body='touch "$1/other.so"\n') + code, out, err = self.run_main(["--repo-root", str(self.root), "--out", str(self.out)]) + self.assertEqual(code, 1) + self.assertIn("1 built, 0 skipped, 2 failed", out) + self.assertIn("aa", err) + self.assertIn("zz", err) + self.assertIn("bar", out) + + def test_a_tree_with_no_manifest_fails_rather_than_passing_quietly(self): + code, _, err = self.run_main(["--repo-root", str(self.root), "--out", str(self.out)]) + self.assertEqual(code, 1) + self.assertIn("no backend manifests", err) + + +if __name__ == "__main__": + unittest.main(verbosity=2) From 9f96df9d62a6cb7ef879fee8a869a8546fc8b730 Mon Sep 17 00:00:00 2001 From: xiaoyu <1259084489@qq.com> Date: Thu, 22 Oct 2026 21:11:31 +0800 Subject: [PATCH 2/3] ci: compile every declared backend The contract's Tier 2 -- build each backend, so a kernel that stops compiling cannot land -- was the only tier with no wiring at all, so nothing in CI has ever compiled a `.cu` file. The job runs the sweep from `build_backends.py`, which is tested without CUDA in `tests/cuda/`, and then reports the ELF images inside each library. A file of the declared name proves that something was linked, not that it was built for a target: an empty metadata image and a correctly targeted kernel look the same until the architectures are listed. No GPU and no self-hosted runner: nvcc cross-compiles for a target it is not running on, and nothing here runs what it builds. The container carries nvcc and cuBLASLt, and the TileLang the Laya recipe records is installed because that backend's export needs it. Depends on the TileLang target pin in the preceding commit: without it Laya's export falls back to sm_50 and this job cannot be green. --- .github/workflows/ci.yml | 43 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 43 insertions(+) diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 01a74893..7058006a 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -83,3 +83,46 @@ jobs: - name: Contract check if: steps.present.outputs.ready == 'true' run: python src/backends/cuda/check_contract.py --repo-root . + + cuda-compile: + # Tier 2 of the contract: build every declared backend, so a kernel that + # stops compiling cannot land. No GPU and no self-hosted runner -- nvcc + # cross-compiles for a target it is not running on, and nothing here runs + # what it builds. Tier 3, which does need a card, stays separate. + # + # The container carries nvcc and cuBLASLt. Laya exports its CUDA from + # TileLang before compiling it, so the TileLang the recipe records is + # installed too. The CUDA line is 13.0 because that is what + # `recipe/laya/native/README.md` records for the native build. + runs-on: ubuntu-latest + timeout-minutes: 30 + container: nvidia/cuda:13.0.0-devel-ubuntu24.04 + steps: + - uses: actions/checkout@v5 + - name: Build toolchain + run: | + apt-get update + apt-get install -y --no-install-recommends python3 python3-pip + python3 -m pip install --no-cache-dir --break-system-packages tilelang==0.1.14 + - name: Compile every declared backend + # The manifests decide what is built: a backend added anywhere under + # src/backends/cuda/ is covered without an edit here. One backend's + # defect is reported against that backend and does not stop the others. + run: python3 src/backends/cuda/build_backends.py --repo-root . --out "$RUNNER_TEMP/cuda-build" + - name: Report the architecture inside every library + # A file of the declared name can exist and still carry the wrong targets, + # or no device code at all: its presence proves neither. Listing the ELF + # images is what distinguishes a library built for a target from an empty + # metadata image. + run: | + missing=0 + for library in "$RUNNER_TEMP"/cuda-build/*/*.so; do + [ -f "$library" ] || continue + echo "== $library" + cuobjdump --list-elf "$library" | tee "$library.elf.txt" + if ! grep -q "sm_" "$library.elf.txt"; then + echo " no sm_ target inside this library" >&2 + missing=1 + fi + done + exit $missing From d87226039c8257928ac7306cff5fb570103309e3 Mon Sep 17 00:00:00 2001 From: xiaoyu <1259084489@qq.com> Date: Fri, 23 Oct 2026 23:48:22 +0800 Subject: [PATCH 3/3] cuda: do not reuse a library from an earlier run The sweep cleared nothing before building: it created the per-backend directory with `exist_ok=True` and afterwards only asked whether the declared library was there. Position and name keep two backends in one invocation apart, but pointing two invocations at the same `--out` left the first run's library in place, so a build script that produced nothing was reported as built. The directory belongs to this tool, so it is cleared first: whatever the check finds afterwards was produced by this invocation. A regression drives the same `--out` twice, the second time with a script that writes nothing, and requires the second to fail; it fails without this change. --- src/backends/cuda/build_backends.py | 9 ++++++++- tests/cuda/test_build_backends.py | 13 +++++++++++++ 2 files changed, 21 insertions(+), 1 deletion(-) diff --git a/src/backends/cuda/build_backends.py b/src/backends/cuda/build_backends.py index bbfda55a..598bd44d 100644 --- a/src/backends/cuda/build_backends.py +++ b/src/backends/cuda/build_backends.py @@ -35,6 +35,7 @@ import glob import json import os +import shutil import subprocess import sys import tempfile @@ -143,8 +144,14 @@ def build_all(repo_root, out_dir): # Keyed by position as well as name: two manifests may declare the same # backend name and the same library filename, and neither may be credited # with the other's artifact. + # + # Cleared first, because the directory belongs to this tool. Reusing an + # existing `--out` would otherwise leave the previous invocation's library + # in place, and a build that produces nothing would be reported as having + # produced it. target = os.path.join(out_dir, "%03d-%s" % (index, _safe(label))) - os.makedirs(target, exist_ok=True) + shutil.rmtree(target, ignore_errors=True) + os.makedirs(target) print("== %s: %s -> %s (sm_%s)" % (shown, script_name, output, arch), flush=True) try: code = subprocess.run([script, target, str(arch)], cwd=directory).returncode diff --git a/tests/cuda/test_build_backends.py b/tests/cuda/test_build_backends.py index 8bd7d34a..831fdeee 100644 --- a/tests/cuda/test_build_backends.py +++ b/tests/cuda/test_build_backends.py @@ -179,6 +179,19 @@ def test_a_tree_of_only_planned_backends_is_not_a_failure(self): # -- identities must not collide -------------------------------------- + def test_reusing_the_output_directory_does_not_reuse_a_library(self): + # The same --out twice: the first build writes the library, the second + # writes nothing. Keeping the directory would credit the second with the + # first one's file and report a build that never happened. + directory = self.backend("", "flat") + built, _, failed = self.sweep() + self.assertEqual(self.labels(built), ["flat"]) + self.script(directory, "true\n") + built, _, failed = self.sweep() + self.assertEqual(built, [], "a build that produced nothing must not be reported as built") + self.assertEqual(self.labels(failed), ["flat"]) + self.assertIn("was not produced", failed[0].detail) + def test_a_duplicate_name_cannot_be_credited_with_the_other_artifact(self): # Same backend name, same output filename. The first produces the # library; the second does not. A shared output directory would let the