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 diff --git a/src/backends/cuda/build_backends.py b/src/backends/cuda/build_backends.py new file mode 100644 index 00000000..598bd44d --- /dev/null +++ b/src/backends/cuda/build_backends.py @@ -0,0 +1,202 @@ +#!/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 shutil +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. + # + # 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))) + 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 + 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..831fdeee --- /dev/null +++ b/tests/cuda/test_build_backends.py @@ -0,0 +1,247 @@ +"""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_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 + # 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)