diff --git a/.github/qualify-rocm/README.md b/.github/qualify-rocm/README.md new file mode 100644 index 000000000000..4131a47fe489 --- /dev/null +++ b/.github/qualify-rocm/README.md @@ -0,0 +1,162 @@ +# ROCm build qualification + +Verification of the ROCm archives this repository publishes. Driven from +`.github/workflows/release.yml`, but every script here can be run by hand +against a downloaded release asset, which is the easiest way to understand what +any of them does. + +## Why this exists + +A broken ROCm build does not look broken. + +llama.cpp loads its GPU support as a separate plug-in at run time, and treats a +plug-in that fails to load as a non-event: it writes a line to the log, skips +it, and carries on using the CPU. So a build whose GPU support is completely +broken still starts, still answers every request, and still returns correct +answers. It is simply ten times slower, with no error anywhere. + +The same quietness applies further up. If the compiler produces GPU code for +only some of the gfx targets it was asked to cover, nothing reports an error — the +output file is still created, with the expected name and a plausible size. + +Both of those ship green unless something deliberately looks inside the +archive. That is what this suite is for. + +## Tiers + +| Tier | Runs on | Blocks a release? | +|---|---|---| +| Tier 0 — static checks | standard GitHub runners, including on PRs | yes, affected ROCm version only | +| Tier 1 — GPU inference | self-hosted machines with a real AMD GPU | yes, affected ROCm version only | + +Each ROCm version is treated as an independent product. A consumer downloads +exactly one of them, so a failure in one version removes only that version's +Windows and Linux pair; the other versions, and every non-ROCm backend, publish +as normal. The release job then exits non-zero so the gap is noticed. + +### Tier 0 — `tier0_static.py`, `check_post_merge.py` + +Runs against the archive as it will be published, not against the build +directory, because most of what it looks for only exists once the files have +been packaged. + +| ID | Check | Blocks release | +|---|---|---| +| T0.1 | the binaries a user needs are all present | yes | +| T0.2 | the GPU backend exists on both operating systems and is big enough to hold GPU code | yes | +| T0.3 | the GPU backend really contains code for every gfx target the build asked for | yes | +| T0.4 | the CPU backend variants are present | yes | +| T0.5 | every library the archive needs is either inside it, part of the OS, or a ROCm library | yes | +| T0.6 | the archive still runs on the oldest Linux release we support | yes | +| T0.7 | no build-machine paths left inside the binaries | yes | +| T0.8 | `llama-server` still accepts the flags it is launched with | yes | +| T0.9 | the folder inside the archive matches the archive's own name | no | +| T0.10 | the Windows CPU backend merge landed | yes | +| T0.11 | the finished archive still opens and is complete | yes | +| T0.12 | the ROCm libraries the archive needs exist in the ROCm a user installs | yes | + +T0.10 and T0.11 come from `check_post_merge.py`, which runs later than the +rest: the release job merges the Windows CPU backend into the Windows ROCm +archives, and those two questions can only be answered after that has happened. + +### Tier 1 — `tier1_gpu.py` + +Starts `llama-server` with the flags the downstream consumer uses, loads a +small model, and talks to it over HTTP. + +| ID | Check | Blocks release | +|---|---|---| +| T1.1 | the server starts and reports itself healthy | yes | +| T1.2 | the GPU actually did the work | yes | +| T1.3 | the reply is non-empty, valid text of a sensible length, with a proper stop reason | yes | +| T1.4 | two identical requests with randomness disabled give identical replies | yes | +| T1.5 | the model closed its reasoning block | no | + +**T1.2 is the one that matters.** Every other check in this tier passes +perfectly well when the GPU was never used at all. + +T1.4 is worth a word too: with randomness switched off, the same question must +produce the same answer every time. When it does not, something underneath is +reading uninitialised memory or producing invalid numbers — a problem that +otherwise shows up as occasional nonsense rather than a clean failure. + +## How some of this works + +### Counting the gfx targets inside the backend + +`binfmt.py` searches the GPU backend file for the short text labels the +compiler writes next to each gfx target's code, rather than asking an AMD tool to +list them. + +The reason is that the GPU code lives in a named block inside the file, and +Windows shortens those names to eight characters once a program is linked — so +a block Linux calls `.hip_fatbin` is called `.hip_fat` there. Any tool asked to +find the block by name comes back empty on Windows. Searching the file's bytes +does not care what the block is called, so one implementation covers both +operating systems and needs no AMD tooling installed, which in turn means the +check can run on an ordinary runner against the archive we actually ship. + +Two deliberate behaviours: + +- Finding **no** gfx targets at all is a failure, not a pass. A check that can + silently come up empty is worse than no check. +- The number of entries per target is reported, not just which targets are + present. The compiler writes one entry per source file, so in a healthy build + every target has the same count; a target with a lower count was built into + only part of the project, which a present/absent check would miss. + +### Which ROCm the checks verify against + +The archives do not contain ROCm. A user installs it separately, and llama.cpp +loads it at run time. So "are this archive's libraries available?" depends +entirely on what the user installed. + +Checking against the ROCm used to *build* the archive would always pass — those +are the exact files it was compiled against. `install_rocm.py` therefore +installs the components an end user gets (the kernels for their GPU) rather +than the components a build needs (compilers), and T0.12 and Tier 1 both verify +against that. Everything goes into a throwaway virtual environment under the +runner's temporary directory; nothing is installed system-wide. + +### Which Linux versions the archive runs on + +Every Linux program records the newest version of the system C library it +relied on, and refuses to start on a machine with an older one. Newer Linux +releases ship newer versions of that library, so building on a newer Linux +silently narrows who can run the result — nothing warns about it at build time. + +T0.6 reads that number out of the archive and compares it against +`SUPPORTED_GLIBC` in `tier0_static.py`, which is set to what the oldest Linux +release we support provides. The check passes today and would fail if the build +job's runner image moved to a newer Linux without that decision being made +deliberately. + +## Running by hand + +```bash +# Check a downloaded release asset +python3 .github/qualify-rocm/tier0_static.py \ + --archive llama-bNNNNN-bin-ubuntu-rocm-10.1-x64.tar.gz \ + --os ubuntu --rocm-version 10.1 \ + --gpu-targets "gfx900;gfx906;...;gfx1201" \ + --output tier0.json + +# Combine reports into one record per archive and decide what may ship +python3 .github/qualify-rocm/aggregate.py \ + --fragments-dir ./fragments --tag bNNNNN \ + --records-dir ./qualification --fail-on-no-promote +``` + +Each run writes a `qualification-.json` record that is attached to the +release. It lists which gfx targets the archive contains code for, which of them +were actually tested on hardware, and the result of every check — so questions +like "was this gfx target ever verified in that build?" stay answerable long after +the CI logs have expired. + +## Out of scope + +- **Performance.** Tracked downstream against these same artifacts; a second + set of numbers here would split the history. +- **Answer quality.** T1.3 and T1.4 cover "the output is well-formed and + reproducible", which is what catches broken GPU kernels. Judging whether an + answer is *good* is a separate problem. diff --git a/.github/qualify-rocm/aggregate.py b/.github/qualify-rocm/aggregate.py new file mode 100644 index 000000000000..8fbc60866413 --- /dev/null +++ b/.github/qualify-rocm/aggregate.py @@ -0,0 +1,246 @@ +#!/usr/bin/env python3 +""" +Merge tier reports into one qualification record per published archive, decide +which ROCm versions may be published, and write the job summary. + +Reports are grouped by (OS, ROCm version), because that pair identifies exactly +one release asset -- the unit a user downloads and the unit a bug report names. + +A ROCm version is publishable when BOTH its archives (Windows and Linux) pass. +Versions are judged independently of each other: they are separate products and +a consumer installs exactly one of them, so a problem in one is no reason to +withhold another. + +Within one archive: + - the static checks must be present and must pass. + - the GPU checks must pass IF they ran. An archive with no GPU report is + recorded as not hardware-validated rather than being failed -- each build + covers 20 or more gfx targets and only one of them has test hardware, so + the record states which is which instead of implying they are all equal. + +With --fail-on-no-promote the exit code is non-zero if any version failed. The +release job uses that to drop only the affected ROCm versions, leaving every +other backend untouched. + +Usage: + python3 aggregate.py --fragments-dir ./fragments --tag bNNNNN \ + --records-dir ./qualification --output-summary summary.json \ + --fail-on-no-promote +""" + +import argparse +import glob +import json +import os +import sys + +sys.path.insert(0, os.path.dirname(os.path.abspath(__file__))) + +import report # noqa: E402 + + +def load_fragments(fragments_dir): + fragments = [] + for path in sorted(glob.glob(os.path.join(fragments_dir, "*.json"))): + try: + with open(path, encoding="utf-8") as handle: + data = json.load(handle) + except (OSError, ValueError) as exc: + print(f"::warning::skipping unreadable fragment {path}: {exc}") + continue + if data.get("kind") == "tier_fragment": + fragments.append(data) + return fragments + + +def group_key(fragment): + build = fragment.get("build", {}) + return (build.get("os") or "?", build.get("rocm_version") or "?") + + +def published_name(tag, os_name, rocm_version): + """The release asset filename for this (os, rocm) pair. + """ + if not tag: + return None + if os_name == "win": + return f"llama-{tag}-bin-win-rocm-{rocm_version}-x64.zip" + if os_name == "ubuntu": + return f"llama-{tag}-bin-ubuntu-rocm-{rocm_version}-x64.tar.gz" + return None + + +def decide(record): + """Return (publishable, reasons) for one archive's merged record.""" + reasons = [] + tiers = record.get("tiers", {}) + + if "tier0" not in tiers: + reasons.append("tier0 missing") + elif tiers["tier0"]["status"] == report.STATUS_FAIL: + reasons.append("tier0=fail") + + if "tier1" in tiers and tiers["tier1"]["status"] == report.STATUS_FAIL: + reasons.append("tier1=fail") + + return not reasons, reasons + + +def gfx_summary(record): + """Pull the gfx audit result out of T0.3 for the record's top level. + + This is the field that answers "was my GPU ever verified?" -- present in + the binary is not the same claim as validated on hardware, and the record + should not blur them. + """ + tests = record.get("tiers", {}).get("tier0", {}).get("tests", []) + audit = next((t for t in tests if t.get("id") == "T0.3"), None) + details = (audit or {}).get("details", {}) + validated = [] + if record.get("build", {}).get("hardware_validated"): + target = record.get("build", {}).get("gfx_target") + if target: + validated = [target] + return { + "requested": details.get("expected", []), + "present_in_binary": details.get("present", []), + "missing_from_binary": details.get("missing", []), + "hardware_validated": validated, + } + + +def write_step_summary(records, promote_all): + path = os.environ.get("GITHUB_STEP_SUMMARY") + if not path: + return + lines = ["## ROCm build qualification", "", + f"**ROCm family publishable: {'YES' if promote_all else 'NO'}**", "", + "| Artifact | Tier 0 | Tier 1 | gfx in binary | HW validated | Promote |", + "|---|---|---|---|---|---|"] + for record in records: + build = record["build"] + tiers = record.get("tiers", {}) + gfx = record["gfx"] + t0 = tiers.get("tier0", {}).get("status", "missing").upper() + t1 = tiers.get("tier1", {}).get("status", "not run").upper() + hw = ", ".join(gfx["hardware_validated"]) or "none" + lines.append( + f"| `{build.get('artifact')}` | {t0} | {t1} | " + f"{len(gfx['present_in_binary'])}/{len(gfx['requested'])} | {hw} | " + f"{'yes' if record['promoted'] else 'NO'} |") + + problems = [] + for record in records: + for tier in sorted(record.get("tiers", {})): + for test in record["tiers"][tier].get("tests", []): + if test["status"] in (report.STATUS_FAIL, report.STATUS_WARN): + gate = "" if test.get("gating") else " _(non-gating)_" + problems.append( + f"- `{record['build'].get('artifact')}` `{test['id']}` " + f"**{test['status']}**{gate} — {test['name']}: " + f"{test.get('error') or ''}") + if problems: + lines += ["", "### Findings", *problems] + + with open(path, "a", encoding="utf-8") as handle: + handle.write("\n".join(lines) + "\n") + + +def main(): + parser = argparse.ArgumentParser(description="Aggregate ROCm tier fragments") + parser.add_argument("--fragments-dir", required=True) + parser.add_argument("--tag", default=None, help="llama.cpp release tag") + parser.add_argument("--records-dir", default=None, + help="write one qualification-.json per archive here") + parser.add_argument("--output-summary", default=None, + help="write the combined decision here") + parser.add_argument("--fail-on-no-promote", action="store_true") + args = parser.parse_args() + + fragments = load_fragments(args.fragments_dir) + if not fragments: + print(f"::error::No tier fragments found in {args.fragments_dir}. The " + "qualification jobs produced nothing, so nothing can be vouched for.") + if args.fail_on_no_promote: + sys.exit(1) + return + + grouped = {} + for fragment in fragments: + grouped.setdefault(group_key(fragment), []).append(fragment) + + records = [] + for (os_name, rocm_version), group in sorted(grouped.items()): + record = report.merge_fragments(group, extra_meta={ + "llamacpp_tag": args.tag, + "artifact": published_name(args.tag, os_name, rocm_version), + }) + promote, reasons = decide(record) + record["promoted"] = promote + record["promotion"] = {"blocked_by": reasons} + record["gfx"] = gfx_summary(record) + records.append(record) + + artifact = record["build"].get("artifact") or f"{os_name}-rocm-{rocm_version}" + if args.records_dir: + os.makedirs(args.records_dir, exist_ok=True) + stem = artifact + for suffix in (".tar.gz", ".zip", ".tgz"): + if stem.endswith(suffix): + stem = stem[: -len(suffix)] + break + out = os.path.join(args.records_dir, f"qualification-{stem}.json") + with open(out, "w", encoding="utf-8") as handle: + json.dump(record, handle, indent=2) + handle.write("\n") + print(f"Wrote {out}") + + print(f"{artifact}: overall={record['overall']} promote={promote} " + f"blocked_by={reasons} " + f"gfx={len(record['gfx']['present_in_binary'])}/" + f"{len(record['gfx']['requested'])} " + f"hw_validated={record['gfx']['hardware_validated'] or 'none'}") + + # A ROCm version is publishable only when every archive carrying it passed. + # Versions are independent, so one bad version never withholds the others. + by_version = {} + for record in records: + version = record["build"].get("rocm_version") or "?" + by_version.setdefault(version, []).append(record) + blocked = sorted(v for v, group in by_version.items() + if not all(r["promoted"] for r in group)) + promote_all = not blocked + + write_step_summary(records, promote_all) + + if args.output_summary: + with open(args.output_summary, "w", encoding="utf-8") as handle: + json.dump({ + "tag": args.tag, + "all_versions_publishable": promote_all, + "blocked_rocm_versions": blocked, + "artifacts": [{"artifact": r["build"].get("artifact"), + "rocm_version": r["build"].get("rocm_version"), + "publishable": r["promoted"], + "overall": r["overall"], + "blocked_by": r["promotion"]["blocked_by"]} + for r in records], + }, handle, indent=2) + handle.write("\n") + + github_output = os.environ.get("GITHUB_OUTPUT") + if github_output: + with open(github_output, "a", encoding="utf-8") as handle: + handle.write(f"promote_rocm={'true' if promote_all else 'false'}\n") + handle.write(f"blocked_rocm_versions={' '.join(blocked)}\n") + + if blocked: + print(f"::error::ROCm {', '.join(blocked)} did not qualify and will be " + "dropped from this release. Other ROCm versions and every other " + "backend are unaffected.") + if args.fail_on_no_promote and blocked: + sys.exit(1) + + +if __name__ == "__main__": + main() diff --git a/.github/qualify-rocm/binfmt.py b/.github/qualify-rocm/binfmt.py new file mode 100644 index 000000000000..cf5d5111f320 --- /dev/null +++ b/.github/qualify-rocm/binfmt.py @@ -0,0 +1,413 @@ +#!/usr/bin/env python3 +""" +Read-only inspection of compiled binaries: which gfx targets they support, +which libraries they need, and which version of the C library they require. + +Background, because none of this is obvious from the outside. + +A compiled program is stored in a structured file format. Linux uses ELF (.so +shared libraries, and executables); Windows uses PE (.dll and .exe). They hold +the same kinds of information -- machine code, data, and a list of the other +libraries the program needs -- just arranged differently. Both divide the file +into named blocks called "sections". + +AMD GPUs come in different architectures, called gfx targets -- gfx1151, +gfx1030, and so on. Each one needs machine code built specifically for it. When +llama.cpp's HIP backend is compiled it is asked to produce code for a whole list +of gfx targets at once, and the result is a single "fat binary": one file with a +stack of per-target builds inside it. That is why the HIP backend is around a +gigabyte while every other backend is a few megabytes. + +The problem this module exists to solve: if the compiler produces code for only +some of the requested gfx targets, NOTHING reports an error. The file is still +created, the build still succeeds, and the file still has the expected name. A +user whose gfx target was dropped then runs llama.cpp, which looks for code for +their GPU, does not find it, and quietly falls back to the CPU -- much slower, +with no error message anywhere. So the only way to know is to open the fat +binary and count what is actually inside it. + +Everything here uses the Python standard library only, so it can run on any +runner without installing a GPU toolchain. +""" + +import os +import re +import struct + +# --------------------------------------------------------------------------- +# Counting the gfx targets inside a fat binary +# --------------------------------------------------------------------------- + +# Each per-target build inside the fat binary is labelled with a plain readable +# tag, stored as ordinary text right in the file. The tags look like this: +# +# hip-amdgcn-amd-amdhsa--gfx1151 +# | | | +# | | the gfx target this build is for +# | the AMD GPU platform the code was built for +# the kind of GPU offloading +# +# So finding every gfx target is simply a matter of searching the file's bytes +# for that pattern and collecting the names. The double dash is not a typo -- +# it is an intentionally empty field in the middle of the platform name. +# +# Two details the pattern handles deliberately: +# * The leading word is matched loosely ("hip", "hipv4", ...) because +# different compiler versions spell it differently. +# * Requiring "amdgcn-amd-amdhsa" means only GPU builds match. The same file +# also contains a regular CPU build labelled with an Intel/AMD CPU platform +# name, and we do not want to count that as a gfx target. +BUNDLE_ENTRY_RE = re.compile(rb"[a-z0-9_]+-amdgcn-amd-amdhsa-[-0-9a-zA-Z_.:+]*") + +# Marker text that begins each group of per-target builds. Counted purely as a +# sanity signal, so the report can show how many groups were found. +BUNDLE_MAGIC = b"__CLANG_OFFLOAD_BUNDLE__" + +# gfx target names look like gfx1151, gfx90a, gfx1030. +_GFX_RE = re.compile(r"(gfx[0-9a-fA-F]+)") + +# The fat binary is around a gigabyte, so read it in pieces instead of loading +# the whole thing into memory at once. +_CHUNK = 64 << 20 +# When reading in pieces, a tag could land half in one piece and half in the +# next. Carrying this many bytes over to the next piece guarantees every tag is +# seen whole. It comfortably exceeds the length of any tag. +_OVERLAP = 256 + + +def strip_target_id(entry_id): + """Pull the gfx target name out of a tag. + + Some builds append extra hardware options after the name, separated by a + colon, e.g. 'gfx90a:xnack+'. Those describe optional features, not a + different target, so everything from the colon onward is discarded. + + 'hip-amdgcn-amd-amdhsa--gfx90a:xnack+' -> 'gfx90a' + """ + match = _GFX_RE.search(entry_id.split(":", 1)[0]) + return match.group(1) if match else None + + +def scan_offload_bundles(path): + """Search a binary for per-target build tags. + + Returns (how many times each tag appears, how many build groups were found, + how many bytes were read). + + This reads the file's raw bytes rather than asking a compiler tool to list + the contents, for two reasons. + + First, it works the same on Linux and Windows. The GPU code lives in a + section whose name is '.hip_fatbin' on Linux -- but Windows caps section + names at eight characters once a file is linked, so the very same section + is called '.hip_fat' there. Tools that are told to look up the section by + name therefore find nothing on Windows. Searching the bytes does not care + what the section is called, so one piece of code covers both. + + Second, it needs no AMD toolchain installed, which means this check can run + on an ordinary runner against the archive we actually publish, instead of + only on the machine that compiled it. + """ + counts = {} + magic = 0 + total = 0 + + def tally(window, limit): + """Count tags that begin before `limit`. + + The bytes from `limit` onward get carried into the next piece of the + file, where they will be examined then. That way each position in the + file is considered exactly once: nothing is counted twice, and no tag + is missed because it straddled a boundary. + """ + nonlocal magic + for match in BUNDLE_ENTRY_RE.finditer(window): + if match.start() >= limit: + break + entry = match.group(0).decode("ascii", "replace") + counts[entry] = counts.get(entry, 0) + 1 + start = 0 + while True: + found = window.find(BUNDLE_MAGIC, start) + if found < 0 or found >= limit: + break + magic += 1 + start = found + 1 + + tail = b"" + with open(path, "rb") as handle: + while True: + buf = handle.read(_CHUNK) + if not buf: + tally(tail, len(tail)) + break + total += len(buf) + window = tail + buf + safe = max(len(window) - _OVERLAP, 0) + tally(window, safe) + tail = window[safe:] + return counts, magic, total + + +def audit_gfx_targets(path, expected): + """Check which gfx targets a fat binary actually contains code for. + + `expected` is the list of gfx targets the build was asked to produce, taken + from the build job itself so Linux and Windows are each compared against + their own list rather than a shared guess. + + The returned dict reports which targets are present, which are missing, and + how many times each one appears. That last number is useful on its own: the + compiler emits one entry per source file it compiles, so in a healthy build + every target has the SAME count. A target with a noticeably lower count was + built into only part of the project, which a simple present/absent check + would miss entirely. + """ + counts, magic, total = scan_offload_bundles(path) + per_target = {} + for entry, n in counts.items(): + target = strip_target_id(entry) + if target: + per_target[target] = per_target.get(target, 0) + n + + expected_set = {strip_target_id(t) or t for t in expected} + present = set(per_target) + uniform = len(set(per_target.values())) <= 1 + + return { + "binary": os.path.basename(path), + "bytes_scanned": total, + "bundle_magic_count": magic, + "expected_count": len(expected_set), + "present_count": len(present), + "expected": sorted(expected_set), + "present": sorted(present), + "missing": sorted(expected_set - present), + "extra": sorted(present - expected_set), + "entries_per_target": dict(sorted(per_target.items())), + "entry_count_uniform": uniform, + } + + +# --------------------------------------------------------------------------- +# ELF -- the Linux program file format +# --------------------------------------------------------------------------- + +def is_elf(path): + try: + with open(path, "rb") as handle: + return handle.read(4) == b"\x7fELF" + except OSError: + return False + + +def elf_sections(path): + """List the named blocks ("sections") an ELF file is divided into. + + Each entry gives the section's name, where it starts in the file, and how + big it is. The rest of this module uses it to find specific sections such + as the list of required libraries. + """ + with open(path, "rb") as handle: + header = handle.read(64) + if header[:4] != b"\x7fELF" or header[4] != 2: + raise ValueError(f"{path}: not ELF64") + e_shoff, = struct.unpack_from(" "rocm[libraries,device-]==" + +The part in brackets matters. A BUILD needs ROCm's compilers, so the build jobs +ask for the "devel" component. A USER needs the kernels for their own GPU, so +they get "device-" instead. Verifying against the build's own components +would be circular -- those are the exact files the binary was compiled against, +so of course they are present. Installing the user's set instead means the +checks answer the question that actually matters: will this archive find what +it needs on a real machine? + +Everything lands under --prefix and nothing is installed system-wide, so +repeated runs on a persistent machine stay isolated from each other. + +Usage: + python3 install_rocm.py --version 10.1.0 --arch gfx1151 --prefix /tmp/rocm +""" + +import argparse +import glob +import json +import os +import subprocess +import sys + +# AMD publishes ROCm wheels on different package indexes depending on the +# major version. Keep this in step with the index each build job installs from. +INDEX_BY_MAJOR = { + "7": "https://repo.amd.com/rocm/whl-multi-arch/", + "10": "https://stable.repo.amd.com/rocm/whl-next/", +} + + +def venv_python(venv_dir): + if os.name == "nt": + return os.path.join(venv_dir, "Scripts", "python.exe") + return os.path.join(venv_dir, "bin", "python") + + +def venv_script(venv_dir, name): + if os.name == "nt": + return os.path.join(venv_dir, "Scripts", f"{name}.exe") + return os.path.join(venv_dir, "bin", name) + + +def run(cmd, **kwargs): + print("+ " + " ".join(str(c) for c in cmd), flush=True) + return subprocess.run(cmd, check=False, **kwargs) + + +def site_packages(venv_dir): + if os.name == "nt": + matches = [os.path.join(venv_dir, "Lib", "site-packages")] + else: + matches = sorted(glob.glob(os.path.join(venv_dir, "lib", "python3*", "site-packages"))) + for path in matches: + if os.path.isdir(path): + return path + raise SystemExit(f"no site-packages under {venv_dir}") + + +def runtime_dirs(venv_dir): + """Every directory that has to be on the library search path. + + This returns a list, not a single root directory, because ROCm spreads its + runtime across several packages: the core HIP libraries live in one, and + the maths libraries (rocBLAS, hipBLAS and friends) live in another. Adding + only the first means the core loads and then the maths libraries cannot be + found, which surfaces as a confusing "DLL not found" on Windows or a + missing-library error on Linux. + """ + site = site_packages(venv_dir) + dirs = [] + for pkg in sorted(glob.glob(os.path.join(site, "_rocm_sdk_*"))): + for sub in ("bin", "lib", os.path.join("lib", "llvm", "bin"), + os.path.join("lib", "llvm", "lib")): + candidate = os.path.join(pkg, sub) + if os.path.isdir(candidate) and candidate not in dirs: + dirs.append(candidate) + return dirs + + +def main(): + parser = argparse.ArgumentParser( + description="Install a user-style ROCm runtime into a throwaway venv") + parser.add_argument("--version", required=True, help="e.g. 10.1.0") + parser.add_argument("--arch", required=True, help="e.g. gfx1151") + parser.add_argument("--prefix", required=True, help="scratch dir for the venv") + parser.add_argument("--index-url", default=None, + help="override the index chosen from the ROCm major version") + parser.add_argument("--output-env", default=None, + help="write KEY=VALUE loader-path lines here (e.g. $GITHUB_ENV)") + parser.add_argument("--output-json", default=None, + help="write the resolved runtime dirs here") + args = parser.parse_args() + + major = args.version.split(".")[0] + index = args.index_url or INDEX_BY_MAJOR.get(major) + if not index: + raise SystemExit( + f"No package index known for ROCm {args.version} (major {major}). " + f"Known: {sorted(INDEX_BY_MAJOR)}. Add it to INDEX_BY_MAJOR.") + + venv_dir = os.path.join(args.prefix, f"rocm-{args.version}-{args.arch}") + os.makedirs(args.prefix, exist_ok=True) + + if not os.path.exists(venv_python(venv_dir)): + if run([sys.executable, "-m", "venv", venv_dir]).returncode != 0: + raise SystemExit(f"could not create venv at {venv_dir}") + + python = venv_python(venv_dir) + run([python, "-m", "pip", "install", "--upgrade", "pip", "--quiet"]) + + spec = f"rocm[libraries,device-{args.arch}]=={args.version}" + proc = run([python, "-m", "pip", "install", "--only-binary", ":all:", + "--index-url", index, spec]) + if proc.returncode != 0: + raise SystemExit( + f"\n::error::Could not install '{spec}' from {index}.\n" + f"::error::ROCm {args.version} device wheels for {args.arch} were not " + f"resolvable at that index. If this is the 7.x leg, the multi-arch " + f"index may not publish per-device wheels -- check what " + f"`pip index versions rocm-sdk-device-{args.arch}` offers there, and " + f"either point --index-url elsewhere or drop this leg from the matrix.") + + dirs = runtime_dirs(venv_dir) + if not dirs: + raise SystemExit( + "::error::ROCm installed but no _rocm_sdk_*/{bin,lib} directories were " + f"found under {site_packages(venv_dir)} -- the wheel layout changed.") + + # Informational only -- the search path comes from runtime_dirs() above. + sdk = venv_script(venv_dir, "rocm-sdk") + if os.path.exists(sdk): + run([sdk, "path", "--root"]) + + print(f"\nROCm {args.version} ({args.arch}) runtime directories:") + for directory in dirs: + print(f" {directory}") + + if args.output_json: + with open(args.output_json, "w", encoding="utf-8") as handle: + json.dump({"version": args.version, "arch": args.arch, + "index": index, "venv": venv_dir, + "runtime_dirs": dirs}, handle, indent=2) + + if args.output_env: + joined = os.pathsep.join(dirs) + with open(args.output_env, "a", encoding="utf-8") as handle: + handle.write(f"ROCM_RUNTIME_DIRS={joined}\n") + handle.write(f"ROCM_VENV={venv_dir}\n") + print("\nROCm runtime ready.") + + +if __name__ == "__main__": + main() diff --git a/.github/qualify-rocm/report.py b/.github/qualify-rocm/report.py new file mode 100644 index 000000000000..566d3c3738bb --- /dev/null +++ b/.github/qualify-rocm/report.py @@ -0,0 +1,249 @@ +#!/usr/bin/env python3 +""" +Shared reporting for the ROCm build-qualification suite. + +Every tier emits a JSON *fragment* with an identical schema; `aggregate.py` +merges the fragments for one artifact into a single qualification record that +is attached to the GitHub release. + +The schema is deliberately flat and stable so a dashboard could consume it +later without reshaping: + + - Every test has a stable `id` ("T0.3", "T1.2") that never changes meaning. + - `status` is always one of the STATUS_* values below. + - Numbers live under `details` as numbers, never strings. + - `gating` says whether a FAIL blocks the ROCm family from the release. + +This module is stdlib-only: Tier 1 runs on self-hosted machines where nothing +is installed but ROCm itself. +""" + +import datetime +import json +import time +import traceback + +SCHEMA_VERSION = 1 + +STATUS_PASS = "pass" +STATUS_FAIL = "fail" +STATUS_WARN = "warn" +STATUS_SKIP = "skip" + +_SEVERITY = {STATUS_SKIP: 0, STATUS_PASS: 1, STATUS_WARN: 2, STATUS_FAIL: 3} + +TIER_NAMES = { + "tier0": "Static artifact verification", + "tier1": "GPU functional inference", +} + + +def now_iso(): + """Current UTC time as an ISO-8601 string with a trailing 'Z'.""" + return ( + datetime.datetime.now(datetime.timezone.utc) + .replace(microsecond=0) + .isoformat() + .replace("+00:00", "Z") + ) + + +def build_meta( + artifact, + os_name=None, + rocm_version=None, + llamacpp_tag=None, + upstream_sha=None, + gfx_target=None, + hardware_validated=False, + run_id=None, + run_attempt=None, +): + """Per-artifact metadata block shared by every fragment. + + `artifact` is the release asset filename, which is the only identifier a + user filing a bug report will have in front of them. + """ + build_id = artifact or "local" + if run_id: + build_id = f"{build_id}@{run_id}" + return { + "build_id": build_id, + "artifact": artifact, + "os": os_name, + "rocm_version": rocm_version, + "llamacpp_tag": llamacpp_tag, + # The upstream llama.cpp commit this archive was built from. Pinned once + # per run, so every artifact in a release reports the same one. + "upstream_sha": upstream_sha, + "gfx_target": gfx_target, + "hardware_validated": bool(hardware_validated), + "run_id": run_id, + "run_attempt": run_attempt, + } + + +class TierReport: + """Collects test results for one tier and serializes a JSON fragment.""" + + def __init__(self, tier, meta): + if tier not in TIER_NAMES: + raise ValueError(f"unknown tier {tier!r}") + self.tier = tier + self.meta = meta + self.tests = [] + self._started = time.monotonic() + self._generated_at = now_iso() + + def add(self, test_id, name, status, gating=True, duration_ms=None, + error=None, details=None): + """Record a single test result.""" + if status not in _SEVERITY: + raise ValueError(f"invalid status {status!r}") + self.tests.append({ + "id": test_id, + "name": name, + "tier": self.tier, + "status": status, + "gating": bool(gating), + "duration_ms": duration_ms, + "error": error, + "details": details or {}, + }) + return status + + def run(self, test_id, name, fn, gating=True): + """Run a check function, timing it and trapping exceptions. + + `fn` returns a status string, or (status, error), or + (status, error, details). A raised exception is recorded as a FAIL with + the traceback in details, so one broken check never aborts the tier -- + we always want the full picture of a bad build in one run. + """ + start = time.monotonic() + try: + result = fn() + if isinstance(result, str): + status, error, details = result, None, None + elif len(result) == 2: + (status, error), details = result, None + else: + status, error, details = result + except Exception as exc: # noqa: BLE001 - capture everything + status = STATUS_FAIL + error = f"{type(exc).__name__}: {exc}" + details = {"traceback": traceback.format_exc()} + duration_ms = int((time.monotonic() - start) * 1000) + return self.add(test_id, name, status, gating=gating, + duration_ms=duration_ms, error=error, details=details) + + def counts(self): + out = {STATUS_PASS: 0, STATUS_FAIL: 0, STATUS_WARN: 0, STATUS_SKIP: 0, + "total": len(self.tests)} + for test in self.tests: + out[test["status"]] += 1 + return out + + def rollup_status(self): + """Tier status = worst severity among *gating* tests.""" + worst = STATUS_PASS + for test in self.tests: + if not test["gating"]: + continue + if _SEVERITY[test["status"]] > _SEVERITY[worst]: + worst = test["status"] + return worst + + def to_dict(self): + return { + "schema_version": SCHEMA_VERSION, + "kind": "tier_fragment", + "tier": self.tier, + "tier_name": TIER_NAMES[self.tier], + "status": self.rollup_status(), + "generated_at": self._generated_at, + "duration_ms": int((time.monotonic() - self._started) * 1000), + "build": self.meta, + "counts": self.counts(), + "tests": self.tests, + } + + def write(self, path): + with open(path, "w", encoding="utf-8") as handle: + json.dump(self.to_dict(), handle, indent=2) + handle.write("\n") + + def print_summary(self): + """Human-readable summary to stdout (the CI log).""" + rollup = self.rollup_status() + print(f"\n=== {self.tier} ({TIER_NAMES[self.tier]}): {rollup.upper()} ===") + for test in self.tests: + gate = "" if test["gating"] else " (non-gating)" + line = f" [{test['status'].upper()}] {test['id']} {test['name']}{gate}" + if test["error"]: + line += f" -- {test['error']}" + print(line) + counts = self.counts() + print(f" totals: {counts[STATUS_PASS]} pass, {counts[STATUS_FAIL]} fail, " + f"{counts[STATUS_WARN]} warn, {counts[STATUS_SKIP]} skip") + + +def merge_fragments(fragments, extra_meta=None): + """Merge per-tier fragments for one artifact into one record.""" + fragments = sorted(fragments, key=lambda frag: frag.get("tier", "")) + meta = {} + for frag in fragments: + for key, value in frag.get("build", {}).items(): + # Never let a tier that did not touch hardware clear the flag a + # tier that did has already set. + if key == "hardware_validated": + meta[key] = bool(meta.get(key)) or bool(value) + elif value is not None or key not in meta: + meta[key] = value + if extra_meta: + meta.update({k: v for k, v in extra_meta.items() if v is not None}) + + tiers = {} + for frag in fragments: + tier = frag["tier"] + bucket = tiers.setdefault(tier, { + "status": STATUS_SKIP, + "tier_name": frag.get("tier_name"), + "counts": {}, + "tests": [], + }) + seen = {t.get("id") for t in bucket["tests"]} + for test in frag.get("tests", []): + # Identical ids from a re-run would double-count; keep the first. + if test.get("id") in seen: + continue + bucket["tests"].append(test) + seen.add(test.get("id")) + + overall = STATUS_PASS + totals = {STATUS_PASS: 0, STATUS_FAIL: 0, STATUS_WARN: 0, STATUS_SKIP: 0, "total": 0} + for tier, bucket in tiers.items(): + counts = {STATUS_PASS: 0, STATUS_FAIL: 0, STATUS_WARN: 0, STATUS_SKIP: 0, + "total": len(bucket["tests"])} + status = STATUS_PASS + for test in bucket["tests"]: + counts[test["status"]] += 1 + if test.get("gating") and _SEVERITY[test["status"]] > _SEVERITY[status]: + status = test["status"] + bucket["counts"] = counts + bucket["status"] = status + if _SEVERITY[status] > _SEVERITY[overall]: + overall = status + for key, value in counts.items(): + totals[key] = totals.get(key, 0) + value + + return { + "schema_version": SCHEMA_VERSION, + "kind": "qualification_record", + "generated_at": now_iso(), + "build": meta, + "overall": overall, + "promoted": False, + "counts": totals, + "tiers": tiers, + } diff --git a/.github/qualify-rocm/tier0_static.py b/.github/qualify-rocm/tier0_static.py new file mode 100644 index 000000000000..80782ca6cec5 --- /dev/null +++ b/.github/qualify-rocm/tier0_static.py @@ -0,0 +1,675 @@ +#!/usr/bin/env python3 +""" +Tier 0 -- static verification of a published ROCm archive. No GPU required. + +Runs on a free hosted runner against the ARCHIVE WE ARE ABOUT TO PUBLISH (or +have just published), not against the build tree. That distinction matters because +every bug this tier exists to catch -- a dropped gfx target, a missing CPU +backend after the Windows merge, a leaked build path -- is invisible in the +build tree and only appears in the packaged bytes. + +Checks: + T0.1 archive contents -- every binary a user needs is present. + T0.2 HIP backend present -- ggml-hip exists on BOTH OSes and is + large enough to contain GPU code. + T0.3 gfx target audit -- the GPU backend really contains code for + every gfx target the build asked for. + See binfmt.py for how and why. + T0.4 CPU backend variants -- ggml-cpu-* present. On Windows after the + release job's CPU merge this is also the + proof that the merge landed. + T0.5 runtime dependency survey -- every non-ROCm dependency resolves from + the archive or the base system; the ROCm + ones are recorded as expected-external. + T0.6 Linux version support -- Linux only. The archive must not require + a newer C library than the oldest + distribution we support. + T0.7 library search path -- Linux only. No absolute build-machine + paths baked into the binaries. + T0.8 llama-server flag surface -- the flags Lemonade launches with still + parse (needs --run-binaries). + T0.9 archive tag consistency -- the directory inside the tarball matches + the tag in the asset filename. + T0.12 ROCm dependency resolution -- the ROCm libraries this archive needs + exist in the ROCm a user will install + (needs --rocm-dirs). + +(T0.10 and T0.11 are emitted separately by check_post_merge.py, which runs on +the final archives after the release job assembles them.) + +Usage: + python3 tier0_static.py --archive llama-bin-ubuntu-rocm-10.1-x64.tar.gz \ + --os ubuntu --rocm-version 10.1 --gpu-targets "gfx900;gfx906;..." \ + --output tier0-ubuntu-10.1.json +""" + +import argparse +import fnmatch +import json +import os +import re +import shutil +import struct +import subprocess +import sys +import tarfile +import tempfile +import zipfile + +sys.path.insert(0, os.path.dirname(os.path.abspath(__file__))) + +import binfmt # noqa: E402 +import report # noqa: E402 + +# Tools a consumer actually invokes. +REQUIRED_TOOLS = ["llama-server", "llama-cli", "llama-bench", "llama-perplexity"] + +REQUIRED_LIBS = { + "ubuntu": ["libggml-base.so", "libggml.so", "libllama.so", "libggml-hip.so"], + "win": ["ggml-base.dll", "ggml.dll", "llama.dll", "ggml-hip.dll"], +} + +CPU_VARIANT_GLOB = {"ubuntu": "libggml-cpu-*.so", "win": "ggml-cpu-*.dll"} +HIP_BACKEND = {"ubuntu": "libggml-hip.so", "win": "ggml-hip.dll"} + +MIN_HIP_BACKEND_BYTES = 64 << 20 + +# Libraries that come with any Linux installation, so the archive does not and +# should not carry them. +BASE_LIBS_LINUX = { + "libc.so.6", "libm.so.6", "libdl.so.2", "librt.so.1", "libpthread.so.0", + "libstdc++.so.6", "libgcc_s.so.1", "ld-linux-x86-64.so.2", "libutil.so.1", + "libresolv.so.2", "libgomp.so.1", + # llama-server uses the distribution's OpenSSL 3 rather than a copy built + # into the binary. Every supported distribution ships it. + "libssl.so.3", "libcrypto.so.3", +} + +# ROCm libraries. These are deliberately NOT in the archive: a user installs +# ROCm separately, and llama.cpp loads it from there at run time. T0.5 records +# which ones an archive needs; T0.12 confirms they exist in the ROCm that will +# actually be installed. +ROCM_LIB_PREFIXES = ( + "libamdhip64", "libhipblas", "librocblas", "libhipblaslt", "librocsolver", + "libamd_comgr", "libhsa-runtime64", "librocm", "libroctx", "librocprofiler", +) +ROCM_DLL_PREFIXES = ( + "amdhip64", "hipblas", "rocblas", "hipblaslt", "rocsolver", "amd_comgr", + "amdhsa", "rocm", +) + +# Windows DLLs supplied by the OS or the MSVC redistributable. +BASE_DLLS_WIN = ("kernel32", "user32", "advapi32", "ws2_32", "shell32", + "ole32", "oleaut32", "bcrypt", "crypt32", "dbghelp", + "msvcp140", "vcruntime140", "vcruntime140_1", "msvcrt", + "api-ms-win-", "ucrtbase", "concrt140", "powrprof", + "psapi", "setupapi", "cfgmgr32", "version", "winmm") + +# --- oldest Linux we support ------------------------------------------------ +# The Linux ROCm artifact is built on Ubuntu 24.04, which is also the oldest +# release the downstream consumer supports. Ubuntu 24.04 ships glibc 2.39, so +# an archive that needs anything newer than that would fail to start even on a +# supported machine. +# +# Raise this only together with the runner image in the build job, and only +# after deciding that dropping the older distributions is acceptable. See +# binfmt.elf_glibc_floor() for what this number means. +SUPPORTED_GLIBC = (2, 39) +SUPPORTED_DISTRO = "Ubuntu 24.04" + +# Flags Lemonade passes to llama-server. +# Upstream renames flags routinely and this repo rebuilds upstream every 12 hours, +# so a rename should fail here within hours rather than surfacing as a Lemonade bug. +LEMONADE_SERVER_FLAGS = [ + "--model", "--ctx-size", "--device", "--port", "--jinja", "--metrics", + "--mmproj", "--no-mmproj-offload", "--model-draft", "--embeddings", + "--reranking", +] + + +# --------------------------------------------------------------------------- +# Archive handling +# --------------------------------------------------------------------------- + +def extract(archive, dest): + """Extract `archive` into `dest`; return (content_root, top_level_names). + + Linux tarballs nest everything under llama-/; Windows zips put the + files at the archive root. Both layouts are normal -- return the directory + that actually holds the binaries. + """ + if archive.endswith((".tar.gz", ".tgz")): + with tarfile.open(archive) as tar: + names = tar.getnames() + # data filter: never let a published archive write outside dest. + try: + tar.extractall(dest, filter="data") + except TypeError: # Python < 3.12 + tar.extractall(dest) + elif archive.endswith(".zip"): + with zipfile.ZipFile(archive) as zf: + names = zf.namelist() + zf.extractall(dest) + else: + raise ValueError(f"unsupported archive type: {archive}") + + tops = sorted({n.split("/")[0] for n in names if n.split("/")[0] not in (".", "")}) + root = dest + if len(tops) == 1 and os.path.isdir(os.path.join(dest, tops[0])): + root = os.path.join(dest, tops[0]) + return root, tops + + +def listing(root): + return sorted(os.listdir(root)) + + +def exe(name, os_name): + return f"{name}.exe" if os_name == "win" else name + + +# --------------------------------------------------------------------------- +# Checks +# --------------------------------------------------------------------------- + +def check_contents(root, os_name): + files = set(listing(root)) + missing = [t for t in (exe(n, os_name) for n in REQUIRED_TOOLS) if t not in files] + missing += [lib for lib in REQUIRED_LIBS[os_name] if lib not in files] + if "LICENSE" not in files: + missing.append("LICENSE") + + # The windows-rocm build job is the only one that does not pass the shared + # CMAKE_ARGS (which carry -DLLAMA_BUILD_TESTS=OFF -DLLAMA_BUILD_EXAMPLES=OFF), + # so its zip ships the upstream unit-test binaries to end users. + tests = sorted(f for f in files if f.startswith("test-")) + details = {"entries": len(files), "missing": missing, + "test_binaries_shipped": len(tests), "test_binaries": tests[:10]} + + if missing: + return (report.STATUS_FAIL, f"missing from archive: {', '.join(missing)}", details) + if tests: + return (report.STATUS_WARN, + f"{len(tests)} upstream test binary/binaries shipped to users " + "(this job does not pass -DLLAMA_BUILD_TESTS=OFF)", details) + return (report.STATUS_PASS, None, details) + + +def check_hip_backend(root, os_name): + name = HIP_BACKEND[os_name] + path = os.path.join(root, name) + details = {"backend": name, "exists": os.path.exists(path)} + if not details["exists"]: + return (report.STATUS_FAIL, + f"{name} is missing -- the HIP backend silently failed to build " + "or was dropped at packaging", details) + size = os.path.getsize(path) + details["bytes"] = size + if size < MIN_HIP_BACKEND_BYTES: + return (report.STATUS_FAIL, + f"{name} is only {size / 1e6:.1f} MB; a multi-target HIP fat " + "binary is ~1 GB, so this contains little or no device code", + details) + return (report.STATUS_PASS, None, details) + + +def check_gfx_targets(root, os_name, gpu_targets): + name = HIP_BACKEND[os_name] + path = os.path.join(root, name) + if not os.path.exists(path): + return (report.STATUS_FAIL, f"{name} missing (see T0.2)", {}) + + expected = [t.strip() for t in re.split(r"[;,\s]+", gpu_targets) if t.strip()] + if not expected: + return (report.STATUS_FAIL, "no expected gfx target list was supplied", {}) + + result = binfmt.audit_gfx_targets(path, expected) + + # An audit that can silently find nothing is worse than no audit: if the + # scan returns zero targets the METHOD is broken (toolchain changed the + # bundle-entry format), and that must fail loudly rather than pass. + if result["present_count"] == 0: + return (report.STATUS_FAIL, + "found NO gfx code objects at all -- the audit method is broken, " + "not necessarily the build. Check the offload bundle entry format " + "against binfmt.BUNDLE_ENTRY_RE.", result) + + if result["missing"]: + return (report.STATUS_FAIL, + f"{len(result['missing'])} gfx target(s) requested but NOT in the " + f"fat binary: {', '.join(result['missing'])}. Users with those GPUs " + "fall back to CPU with no error message.", result) + + if not result["entry_count_uniform"]: + counts = result["entries_per_target"] + low = min(counts, key=counts.get) + return (report.STATUS_WARN, + f"all targets present but entry counts differ (e.g. {low} has " + f"{counts[low]}, max is {max(counts.values())}) -- a target may be " + "compiled into only some translation units", result) + + return (report.STATUS_PASS, None, result) + + +def check_cpu_variants(root, os_name, post_merge): + files = listing(root) + variants = fnmatch.filter(files, CPU_VARIANT_GLOB[os_name]) + details = {"count": len(variants), "variants": sorted(variants), + "post_merge": post_merge} + if not variants: + extra = "" + if os_name == "win" and post_merge: + extra = (" -- the release job's Windows CPU-backend merge did not land " + "in this zip") + return (report.STATUS_FAIL, + f"no CPU backend matching {CPU_VARIANT_GLOB[os_name]}{extra}", details) + # GGML_CPU_ALL_VARIANTS=ON produces well over a dozen; a single one means + # the variant build collapsed or only part of the merge arrived. + if len(variants) < 5: + return (report.STATUS_WARN, + f"only {len(variants)} CPU backend variant(s); " + "GGML_CPU_ALL_VARIANTS=ON normally yields 14+", details) + return (report.STATUS_PASS, None, details) + + +def _classify_linux(needed, present): + rocm, missing = [], [] + for soname in needed: + if soname.startswith(ROCM_LIB_PREFIXES): + rocm.append(soname) + elif soname in present or soname in BASE_LIBS_LINUX: + continue + else: + missing.append(soname) + return rocm, missing + + +def _classify_win(imports, present_lower): + rocm, missing = [], [] + for dll in imports: + low = dll.lower() + if low.startswith(ROCM_DLL_PREFIXES): + rocm.append(dll) + elif low in present_lower or low.startswith(BASE_DLLS_WIN): + continue + else: + missing.append(dll) + return rocm, missing + + +def check_dependencies(root, os_name): + """Survey every binary's dynamic dependencies. + + ROCm libraries are EXPECTED to be absent -- this repo ships no ROCm. + They are recorded so the set a consumer must provide is written down per release. + Anything else that does not resolve is a packaging defect and fails. + """ + files = listing(root) + present = set(files) + present_lower = {f.lower() for f in files} + rocm_needed, unresolved, scanned = set(), {}, 0 + + for name in files: + path = os.path.join(root, name) + if not os.path.isfile(path) or os.path.islink(path): + continue + try: + if os_name == "ubuntu": + if not binfmt.is_elf(path): + continue + needed, _ = binfmt.elf_dynamic(path) + rocm, missing = _classify_linux(needed, present) + else: + if not binfmt.is_pe(path): + continue + imports = binfmt.pe_imports(path) + if not imports: + continue + rocm, missing = _classify_win(imports, present_lower) + except (ValueError, OSError, struct.error) as exc: + unresolved[name] = [f""] + continue + scanned += 1 + rocm_needed.update(rocm) + if missing: + unresolved[name] = sorted(missing) + + details = { + "binaries_scanned": scanned, + "rocm_runtime_required": sorted(rocm_needed), + "unresolved": unresolved, + } + if unresolved: + flat = sorted({d for deps in unresolved.values() for d in deps}) + return (report.STATUS_FAIL, + f"{len(unresolved)} binary/binaries have dependencies that are " + f"neither in the archive nor part of the base system: " + f"{', '.join(flat[:6])}", details) + if not rocm_needed: + return (report.STATUS_WARN, + "no ROCm runtime dependency found at all -- expected the HIP " + "backend to link one", details) + return (report.STATUS_PASS, None, details) + + +def check_linux_version_support(root, os_name): + """Confirm the archive still runs on the oldest Linux we support. + + Each binary records the newest version of the system C library it needs. + Whichever binary needs the newest one decides the oldest Linux release the + whole archive can run on. If that goes past what SUPPORTED_GLIBC allows, + users on a supported distribution get a startup error instead of a program. + + This normally only changes when the build job's runner image changes, which + is exactly the moment it should be noticed. + """ + if os_name != "ubuntu": + return (report.STATUS_SKIP, "Linux-only check", {}) + + needs, needs_file = (0, 0), None + per_file = {} + for name in listing(root): + path = os.path.join(root, name) + if not os.path.isfile(path) or os.path.islink(path) or not binfmt.is_elf(path): + continue + floor = binfmt.elf_glibc_floor(path) + if not floor: + continue + per_file[name] = ".".join(map(str, floor)) + if floor > needs: + needs, needs_file = floor, name + + allowed = ".".join(map(str, SUPPORTED_GLIBC)) + actual = ".".join(map(str, needs)) if needs_file else None + details: dict = { + "glibc_required": actual, + "glibc_allowed": allowed, + "oldest_supported_distro": SUPPORTED_DISTRO, + "required_by": needs_file, + "per_file": per_file, + } + if not needs_file: + return (report.STATUS_WARN, + "could not determine which C library version this archive needs", + details) + if needs > SUPPORTED_GLIBC: + return (report.STATUS_FAIL, + f"this archive needs glibc {actual}, but the oldest Linux we " + f"support ({SUPPORTED_DISTRO}) only provides {allowed}. " + f"{needs_file} is what requires it. Users on a supported " + "distribution will see \"GLIBC_... not found\" and the program " + "will not start. This usually means the build job's runner image " + "moved to a newer Linux.", details) + return (report.STATUS_PASS, None, details) + + +def check_rpath(root, os_name): + """No absolute build-machine paths in RUNPATH. + + The ROCm Linux job never runs the `patchelf --set-rpath '$ORIGIN'` normalisation. + """ + if os_name != "ubuntu": + return (report.STATUS_SKIP, "Linux-only check", {}) + + leaks = {} + for name in listing(root): + path = os.path.join(root, name) + if not os.path.isfile(path) or os.path.islink(path) or not binfmt.is_elf(path): + continue + _, runpath = binfmt.elf_dynamic(path) + bad = [part for entry in runpath for part in entry.split(":") + if part and not part.startswith("$ORIGIN")] + if bad: + leaks[name] = bad + + details = {"leaked": leaks} + if leaks: + flat = sorted({p for parts in leaks.values() for p in parts}) + return (report.STATUS_FAIL, + f"{len(leaks)} binary/binaries carry a build-machine path in " + f"their library search path: {flat[0]}. The build job should " + "reset it to '$ORIGIN' before packaging.", details) + return (report.STATUS_PASS, None, details) + + +def check_rocm_dependencies(root, os_name, rocm_dirs): + """Confirm the ROCm libraries this archive needs will actually be there. + + This archive does not contain ROCm. The user installs ROCm separately and + llama.cpp loads it at run time. So the archive says "I need a library called + libamdhip64.so.7", and whether that is true depends entirely on what the + user installed -- something the build machine can tell us nothing about. + + Checking against the ROCm the archive was COMPILED with would be circular: + of course the libraries it linked against are present there. The useful + question is whether they are present in the ROCm a user ends up with, which + is installed from the same release but with a different set of components: + the build needs compilers, the user needs per-GPU kernels. So `rocm_dirs` + must point at a user-style ROCm install, not at the build toolchain. + + When nothing is supplied, this is skipped rather than failed -- the rest of + Tier 0 does not depend on having a ROCm install available. + """ + if not rocm_dirs: + return (report.STATUS_SKIP, + "no ROCm install supplied via --rocm-dirs; nothing to resolve " + "the archive's ROCm dependencies against", {}) + + available = {} + for directory in rocm_dirs: + if not os.path.isdir(directory): + continue + for entry in os.listdir(directory): + available.setdefault(entry.lower(), os.path.join(directory, entry)) + + required, unresolved = set(), {} + for name in listing(root): + path = os.path.join(root, name) + if not os.path.isfile(path) or os.path.islink(path): + continue + try: + if os_name == "ubuntu": + if not binfmt.is_elf(path): + continue + deps = [d for d in binfmt.elf_dynamic(path)[0] + if d.startswith(ROCM_LIB_PREFIXES)] + else: + if not binfmt.is_pe(path): + continue + deps = [d for d in binfmt.pe_imports(path) + if d.lower().startswith(ROCM_DLL_PREFIXES)] + except (ValueError, OSError, struct.error): + continue + required.update(deps) + missing = [d for d in deps if d.lower() not in available] + if missing: + unresolved[name] = sorted(missing) + + details: dict = { + "rocm_dirs_searched": len(rocm_dirs), + "files_available": len(available), + "rocm_libraries_required": sorted(required), + "unresolved": unresolved, + } + if not required: + return (report.STATUS_WARN, + "this archive does not reference any ROCm library, which is " + "unexpected for a ROCm build", details) + if unresolved: + flat = sorted({d for deps in unresolved.values() for d in deps}) + return (report.STATUS_FAIL, + f"the installed ROCm does not provide {', '.join(flat)}, which " + f"{len(unresolved)} binary/binaries in this archive need. On " + "Linux llama-server will refuse to start; on Windows the GPU " + "backend fails to load and llama.cpp quietly runs on the CPU " + "instead.", details) + return (report.STATUS_PASS, None, details) + + +def check_server_flags(root, os_name, run_binaries): + """The llama-server flags Lemonade launches with must still parse.""" + if not run_binaries: + return (report.STATUS_SKIP, + "not running binaries on this runner (OS mismatch)", {}) + + server = os.path.join(root, exe("llama-server", os_name)) + if not os.path.exists(server): + return (report.STATUS_FAIL, "llama-server missing (see T0.1)", {}) + if os_name == "ubuntu": + os.chmod(server, 0o755) + + proc = subprocess.run([server, "--help"], capture_output=True, text=True, + check=False, timeout=180, cwd=root) + text = (proc.stdout or "") + (proc.stderr or "") + details: dict = {"returncode": proc.returncode, "help_bytes": len(text)} + if proc.returncode != 0 or not text.strip(): + details["output_tail"] = text[-600:] + return (report.STATUS_FAIL, + f"llama-server --help exited {proc.returncode} with " + f"{len(text)} bytes of output", details) + + missing = [f for f in LEMONADE_SERVER_FLAGS + if not re.search(rf"(? {root} ({len(listing(root))} entries)") + + tier.run("T0.1", "archive contents", + lambda: check_contents(root, args.os_name)) + tier.run("T0.2", "HIP backend present", + lambda: check_hip_backend(root, args.os_name)) + tier.run("T0.3", "gfx target audit", + lambda: check_gfx_targets(root, args.os_name, args.gpu_targets)) + tier.run("T0.4", "CPU backend variants", + lambda: check_cpu_variants(root, args.os_name, args.post_merge)) + tier.run("T0.5", "runtime dependency survey", + lambda: check_dependencies(root, args.os_name)) + tier.run("T0.6", "Linux version support", + lambda: check_linux_version_support(root, args.os_name)) + tier.run("T0.7", "library search path", + lambda: check_rpath(root, args.os_name)) + tier.run("T0.8", "llama-server flag surface", + lambda: check_server_flags(root, args.os_name, args.run_binaries)) + tier.run("T0.9", "archive tag consistency", + lambda: check_tag_consistency(args.archive, tops), + gating=False) + tier.run("T0.12", "ROCm dependency resolution", + lambda: check_rocm_dependencies(root, args.os_name, rocm_dirs)) + finally: + shutil.rmtree(workdir, ignore_errors=True) + + tier.write(args.output) + tier.print_summary() + print(f"\nWrote {args.output}") + + if args.fail_on_error and tier.rollup_status() == report.STATUS_FAIL: + sys.exit(1) + + +if __name__ == "__main__": + main() diff --git a/.github/qualify-rocm/tier1_gpu.py b/.github/qualify-rocm/tier1_gpu.py new file mode 100644 index 000000000000..b38a7b379a52 --- /dev/null +++ b/.github/qualify-rocm/tier1_gpu.py @@ -0,0 +1,412 @@ +#!/usr/bin/env python3 +""" +Tier 1 -- GPU functional inference. Requires a real AMD GPU (gfx1151 runners). + +Drives llama-server over HTTP rather than llama-cli, because llama-server is +the binary the downstream consumer actually launches, with the flag set +reproduced in SERVER_ARGS below. + +T1.2 is the check that carries the weight here. A GPU backend that fails to +load is not treated as an error by llama.cpp: it logs the failure, skips the +backend and continues on the CPU. So a test that only asks "did the server +answer?" passes happily on a build whose GPU support is entirely broken -- the +answers are just much slower. T1.2 confirms the GPU really did the work. + +Checks: + T1.1 server boot -- llama-server starts and serves /health. + T1.2 GPU offload confirmed -- the device/offload lines prove work landed on + the GPU (gating; the anti-CPU-fallback gate). + T1.3 response well-formed -- non-empty, valid UTF-8, no replacement chars, + sane length, proper finish_reason. + T1.4 greedy determinism -- two identical temperature=0 requests return + byte-identical text (gating; model-agnostic + evidence the sampling loop saw sane logits). + T1.5 closed tag -- NON-gating. A Qwen3-thinking artifact kept as + a data point; T1.3 and T1.4 are the real + correctness signals. + +Usage: + python3 tier1_gpu.py --archive-root ./extracted --model ./Qwen3-0.6B-Q4_0.gguf \ + --rocm-version 10.1 --gfx-target gfx1151 --output tier1.json +""" + +import argparse +import glob +import json +import os +import re +import shutil +import signal +import subprocess +import sys +import time +import urllib.error +import urllib.request + +sys.path.insert(0, os.path.dirname(os.path.abspath(__file__))) + +import report # noqa: E402 + +READY_TIMEOUT = 600 +REQUEST_TIMEOUT = 180 + +# The flags the downstream consumer launches llama-server with. Using the same +# ones here means this tier exercises the real integration surface rather than +# a simplified version of it. +SERVER_ARGS = ["--ctx-size", "4096", "--jinja", "--metrics"] + +# Device/offload lines that prove llama-server placed work on the GPU. +GPU_OFFLOAD_PATTERNS = [ + re.compile(r"using device ROCm"), + re.compile(r"offloaded\s+[1-9]\d*/\d+\s+layers to GPU"), +] + +CPU_FALLBACK_HINTS = [ + re.compile(r"failed to load .*ggml-hip", re.I), + re.compile(r"no usable GPU", re.I), + re.compile(r"offloaded\s+0/\d+\s+layers to GPU"), +] + + +# --------------------------------------------------------------------------- +# HTTP +# --------------------------------------------------------------------------- + +def http_get(url, timeout=10): + try: + with urllib.request.urlopen(url, timeout=timeout) as resp: + return resp.status, resp.read() + except (urllib.error.URLError, OSError): + return None, b"" + + +def http_post_json(url, payload, timeout=REQUEST_TIMEOUT): + data = json.dumps(payload).encode("utf-8") + req = urllib.request.Request(url, data=data, + headers={"Content-Type": "application/json"}) + with urllib.request.urlopen(req, timeout=timeout) as resp: + raw = resp.read() + # Decode strictly: a malformed byte sequence here is itself a finding, and + # T1.3 wants to know about it rather than have it silently replaced. + return json.loads(raw.decode("utf-8")) + + +# --------------------------------------------------------------------------- +# Server lifecycle +# --------------------------------------------------------------------------- + +def stage_hip_runtime(rocm_dirs, exe_dir): + """Windows only: copy the HIP runtime next to llama-server.exe. + + Windows resolves DLLs from System32 BEFORE PATH, so a stale System32 + amdhip64_*.dll shadows the ROCm we just installed and the binary silently + binds to the wrong runtime. + """ + if os.name != "nt": + return [] + staged = [] + for directory in rocm_dirs: + for pattern in ("amdhip64*.dll", "amd_comgr*.dll", "rocm_kpack*.dll"): + for src in glob.glob(os.path.join(directory, pattern)): + dst = os.path.join(exe_dir, os.path.basename(src)) + if not os.path.exists(dst): + shutil.copy2(src, dst) + staged.append(os.path.basename(src)) + return staged + + +def build_env(rocm_dirs, exe_dir, gfx_target): + env = dict(os.environ) + joined = os.pathsep.join(rocm_dirs + [exe_dir]) + if os.name == "nt": + env["PATH"] = joined + os.pathsep + env.get("PATH", "") + # Raises the shared-memory limit on these targets so larger models fit. + if gfx_target in ("gfx1151", "gfx1152"): + env["OCL_SET_SVM_SIZE"] = "262144" + else: + env["LD_LIBRARY_PATH"] = joined + os.pathsep + env.get("LD_LIBRARY_PATH", "") + return env + + +class Server: + def __init__(self, exe, model, port, log_path, env, cwd): + self.exe = exe + self.model = model + self.port = port + self.log_path = log_path + self.env = env + self.cwd = cwd + self.proc = None + self._log = None + + def start(self): + cmd = [self.exe, "--model", self.model, "--host", "127.0.0.1", + "--port", str(self.port), "-ngl", "99", *SERVER_ARGS] + print("+ " + " ".join(cmd), flush=True) + self._log = open(self.log_path, "w", encoding="utf-8", errors="replace") + kwargs = {} + if os.name != "nt": + kwargs["start_new_session"] = True + self.proc = subprocess.Popen(cmd, stdout=self._log, + stderr=subprocess.STDOUT, env=self.env, + cwd=self.cwd, **kwargs) + + def wait_ready(self, timeout): + proc = self.proc + if proc is None: + return False, "llama-server was never started" + deadline = time.monotonic() + timeout + while time.monotonic() < deadline: + if proc.poll() is not None: + return False, f"llama-server exited early (rc={proc.returncode})" + status, _ = http_get(f"http://127.0.0.1:{self.port}/health") + if status == 200: + return True, None + time.sleep(3) + return False, f"llama-server not ready within {timeout}s" + + def stop(self): + if self.proc and self.proc.poll() is None: + try: + if os.name == "nt": + self.proc.terminate() + else: + os.killpg(os.getpgid(self.proc.pid), signal.SIGTERM) + self.proc.wait(timeout=30) + except (ProcessLookupError, OSError, subprocess.TimeoutExpired): + try: + self.proc.kill() + except OSError: + pass + if self._log: + self._log.close() + + def log_text(self): + try: + with open(self.log_path, encoding="utf-8", errors="replace") as handle: + return handle.read() + except OSError: + return "" + + +# --------------------------------------------------------------------------- +# Checks +# --------------------------------------------------------------------------- + +def check_gpu_offload(server, base): + log = server.log_text() + matched = [p.pattern for p in GPU_OFFLOAD_PATTERNS if p.search(log)] + + props_status, props_raw = http_get(f"{base}/props") + props = {} + if props_status == 200: + try: + props = json.loads(props_raw.decode("utf-8")) + except (ValueError, UnicodeDecodeError): + props = {} + + details = { + "patterns_matched": matched, + "offload_line": next( + (m.group(0) for m in + [re.search(r"offloaded\s+\d+/\d+\s+layers to GPU", log)] if m), None), + "device_line": next( + (m.group(0) for m in [re.search(r"using device ROCm\S*.*", log)] if m), None), + # /props is supplementary: its schema is not stable across upstream + # releases, so it is recorded but never gates. + "props_status": props_status, + "props_n_gpu_layers": (props.get("default_generation_settings") or {}) + .get("n_gpu_layers"), + } + + if matched: + return (report.STATUS_PASS, None, details) + + hints = [p.pattern for p in CPU_FALLBACK_HINTS if p.search(log)] + details["cpu_fallback_hints"] = hints + details["log_tail"] = log[-2000:] + reason = ("the GPU backend did not load and llama.cpp fell back to the CPU" + if hints else "no device/offload line was logged at all") + return (report.STATUS_FAIL, + f"GPU offload NOT confirmed -- {reason}. With GGML_BACKEND_DL=ON this " + "is silent: the server still answers requests, just on the CPU.", + details) + + +def _content(body): + choice = (body.get("choices") or [{}])[0] + message = choice.get("message") or {} + text = (message.get("content") or "") + (message.get("reasoning_content") or "") + return text, choice.get("finish_reason") + + +def chat(base, prompt, max_tokens=192, seed=0): + return http_post_json(f"{base}/v1/chat/completions", { + "messages": [{"role": "user", "content": prompt}], + "max_tokens": max_tokens, + "temperature": 0.0, + "top_k": 1, + "seed": seed, + }) + + +def check_well_formed(base, state): + body = chat(base, "In one short paragraph, what is the capital of France?") + text, finish = _content(body) + state["sample"] = text + + bad_chars = text.count("�") + details = { + "finish_reason": finish, + "chars": len(text), + "sample": text[:200], + "replacement_chars": bad_chars, + } + problems = [] + if not text.strip(): + problems.append("response is empty") + if bad_chars: + problems.append(f"{bad_chars} U+FFFD replacement character(s) " + "-- the decoder produced invalid UTF-8") + if 0 < len(text.strip()) < 8: + problems.append(f"response is implausibly short ({len(text.strip())} chars)") + if finish not in ("stop", "length"): + problems.append(f"finish_reason is {finish!r}, expected 'stop' or 'length'") + + if problems: + return (report.STATUS_FAIL, "; ".join(problems), details) + return (report.STATUS_PASS, None, details) + + +def check_determinism(base): + prompt = "Count from one to five, separated by commas." + outputs = [] + for _ in range(2): + body = chat(base, prompt, max_tokens=64) + outputs.append(_content(body)[0]) + + details = {"first": outputs[0][:200], "second": outputs[1][:200], + "identical": outputs[0] == outputs[1]} + if not outputs[0].strip(): + return (report.STATUS_FAIL, "greedy request returned nothing", details) + if outputs[0] != outputs[1]: + return (report.STATUS_FAIL, + "two identical temperature=0 requests produced different text -- " + "the sampling path is non-deterministic, which usually means " + "uninitialised memory or NaN logits in the GPU kernels", details) + return (report.STATUS_PASS, None, details) + + +def check_think_tag(state): + text = state.get("sample") or "" + closed = "" in text + details = {"found": closed, "sample": text[:200]} + if closed: + return (report.STATUS_PASS, None, details) + return (report.STATUS_WARN, + "no closed tag. This is a Qwen3-thinking chat-template " + "artifact, not a correctness signal -- T1.3/T1.4 are authoritative.", + details) + + +def main(): + parser = argparse.ArgumentParser(description="Tier 1 GPU functional inference") + parser.add_argument("--archive-root", required=True, + help="directory holding the extracted llama.cpp binaries") + parser.add_argument("--model", required=True, help="path to the .gguf model") + parser.add_argument("--rocm-version", required=True) + parser.add_argument("--gfx-target", default="gfx1151") + parser.add_argument("--artifact", default=None, help="release asset filename") + parser.add_argument("--tag", default=None) + parser.add_argument("--rocm-dirs", default=os.environ.get("ROCM_RUNTIME_DIRS", ""), + help="os.pathsep-separated ROCm runtime directories") + parser.add_argument("--port", type=int, default=8231) + parser.add_argument("--logs-dir", default=".") + parser.add_argument("--ready-timeout", type=int, default=READY_TIMEOUT) + parser.add_argument("--run-id", default=os.environ.get("GITHUB_RUN_ID")) + parser.add_argument("--run-attempt", default=os.environ.get("GITHUB_RUN_ATTEMPT")) + parser.add_argument("--output", default="tier1-report.json") + parser.add_argument("--fail-on-error", action="store_true") + args = parser.parse_args() + + root = os.path.abspath(args.archive_root) + exe = os.path.join(root, "llama-server.exe" if os.name == "nt" else "llama-server") + if not os.path.exists(exe): + raise SystemExit(f"::error::llama-server not found at {exe}") + if os.name != "nt": + for entry in os.listdir(root): + path = os.path.join(root, entry) + if os.path.isfile(path): + os.chmod(path, 0o755) + + rocm_dirs = [d for d in args.rocm_dirs.split(os.pathsep) if d] + if not rocm_dirs: + raise SystemExit( + "::error::No ROCm runtime directories supplied. This repo's archives " + "do not bundle ROCm -- run install_rocm.py first and pass its " + "ROCM_RUNTIME_DIRS through --rocm-dirs.") + + staged = stage_hip_runtime(rocm_dirs, root) + if staged: + print(f"Staged HIP runtime next to llama-server: {', '.join(staged)}") + + meta = report.build_meta( + artifact=args.artifact, + os_name="win" if os.name == "nt" else "ubuntu", + rocm_version=args.rocm_version, + llamacpp_tag=args.tag, + gfx_target=args.gfx_target, + hardware_validated=True, + run_id=args.run_id, + run_attempt=args.run_attempt, + ) + tier = report.TierReport("tier1", meta) + + os.makedirs(args.logs_dir, exist_ok=True) + log_path = os.path.join( + args.logs_dir, f"llama-server-{args.gfx_target}-rocm{args.rocm_version}.log") + base = f"http://127.0.0.1:{args.port}" + env = build_env(rocm_dirs, root, args.gfx_target) + server = Server(exe, os.path.abspath(args.model), args.port, log_path, env, root) + state = {} + + try: + server.start() + ready, err = server.wait_ready(args.ready_timeout) + if ready: + tier.add("T1.1", "server boot", report.STATUS_PASS, + details={"port": args.port, "args": SERVER_ARGS}) + else: + tier.add("T1.1", "server boot", report.STATUS_FAIL, error=err, + details={"log_tail": server.log_text()[-3000:]}) + + if ready: + tier.run("T1.2", "GPU offload confirmed", + lambda: check_gpu_offload(server, base)) + tier.run("T1.3", "response well-formed", + lambda: check_well_formed(base, state)) + tier.run("T1.4", "greedy determinism", lambda: check_determinism(base)) + tier.run("T1.5", "closed tag", + lambda: check_think_tag(state), gating=False) + else: + for test_id, name, gating in [ + ("T1.2", "GPU offload confirmed", True), + ("T1.3", "response well-formed", True), + ("T1.4", "greedy determinism", True), + ("T1.5", "closed tag", False), + ]: + tier.add(test_id, name, report.STATUS_SKIP, + gating=gating, error="server not ready") + finally: + server.stop() + + tier.write(args.output) + tier.print_summary() + print(f"\nWrote {args.output}") + + if args.fail_on_error and tier.rollup_status() == report.STATUS_FAIL: + sys.exit(1) + + +if __name__ == "__main__": + main() diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index 47d70eeb2b29..fef8f8d2ec8d 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -15,6 +15,9 @@ on: # action is consumed via a pinned @lemonade ref, so only changes to this # workflow itself can affect the outcome of a PR run. Scope PR runs to it. - '.github/workflows/release.yml' + # The qualification suite decides whether the ROCm family ships, so a + # change to it must be validated the same way. + - '.github/qualify-rocm/**' concurrency: group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }} @@ -25,7 +28,34 @@ env: CMAKE_ARGS: "-DLLAMA_BUILD_EXAMPLES=OFF -DLLAMA_BUILD_TESTS=OFF -DLLAMA_BUILD_TOOLS=ON -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON" jobs: + # Every build job below clones llama.cpp from upstream. Left to themselves + # they each clone at their own start time, so a push upstream mid-run means + # different jobs compile DIFFERENT source and the release then labels them + # all with one tag. Resolve the commit once here and have every job check out + # that exact commit, so one release is one upstream revision. + resolve-upstream: + runs-on: ubuntu-24.04 + outputs: + sha: ${{ steps.resolve.outputs.sha }} + tag: ${{ steps.resolve.outputs.tag }} + steps: + - name: Clone + uses: actions/checkout@v6 + with: + fetch-depth: 0 + repository: 'ggml-org/llama.cpp' + + - name: Resolve commit and tag + id: resolve + run: | + sha=$(git rev-parse HEAD) + tag="b$(git rev-list --count HEAD)" + echo "sha=$sha" >> "$GITHUB_OUTPUT" + echo "tag=$tag" >> "$GITHUB_OUTPUT" + echo "Pinning this run to upstream $sha ($tag)" + ubuntu-24-rocm: + needs: resolve-upstream runs-on: ubuntu-24.04 strategy: @@ -51,6 +81,7 @@ jobs: with: fetch-depth: 0 repository: 'ggml-org/llama.cpp' + ref: ${{ needs.resolve-upstream.outputs.sha }} - name: ccache uses: ggml-org/ccache-action@v1.2.21 @@ -75,7 +106,7 @@ jobs: - name: Dependencies id: depends run: | - sudo apt install -y build-essential git cmake wget + sudo apt install -y build-essential git cmake wget patchelf - name: Setup TheRock with Wheels id: therock_env @@ -127,9 +158,22 @@ jobs: ${{ env.CMAKE_ARGS }} cmake --build build --config Release -j $(nproc) + # The HIP link adds the pip-installed ROCm devel tree's lib dir to RUNPATH, + # so without this the shipped libggml-hip.so carries an absolute build-machine path + # (.../_rocm_sdk_devel/lib) that means nothing on a user's box. + - name: Normalize RPATH for portable distribution + run: | + for f in build/bin/*; do + [ -f "$f" ] && ! [ -L "$f" ] || continue + if file "$f" | grep -q 'ELF'; then + patchelf --set-rpath '$ORIGIN' "$f" + fi + done + - name: Determine tag name id: tag - uses: lemonade-sdk/llama.cpp/.github/actions/get-tag-name@lemonade + shell: bash + run: echo "name=${{ needs.resolve-upstream.outputs.tag }}" >> "$GITHUB_OUTPUT" - name: Get ROCm short version run: echo "ROCM_VERSION_SHORT=$(echo '${{ matrix.ROCM_VERSION }}' | cut -d '.' -f 1,2)" >> $GITHUB_ENV @@ -140,13 +184,38 @@ jobs: cp LICENSE ./build/bin/ tar -czvf llama-bin-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.tar.gz --transform "s,./,llama-${{ steps.tag.outputs.name }}/," -C ./build/bin . + # The qualification jobs must audit the fat binary against the SAME target + # list this job compiled, so publish the matrix entry alongside the + # artifact instead of duplicating it in another job. + - name: Record qualification metadata + run: | + cat > qualify-meta-ubuntu-rocm-${{ env.ROCM_VERSION_SHORT }}.json <> "$GITHUB_OUTPUT" - name: Pack artifacts id: pack_artifacts @@ -341,6 +412,7 @@ jobs: name: llama-ubuntu-cuda-${{ matrix.sm }}-x64.tar.xz ubuntu-22-cuda-arm64: + needs: resolve-upstream runs-on: ubuntu-22.04-arm strategy: @@ -357,6 +429,7 @@ jobs: with: fetch-depth: 0 repository: 'ggml-org/llama.cpp' + ref: ${{ needs.resolve-upstream.outputs.sha }} - name: ccache uses: ggml-org/ccache-action@v1.2.21 @@ -500,7 +573,8 @@ jobs: - name: Determine tag name id: tag - uses: lemonade-sdk/llama.cpp/.github/actions/get-tag-name@lemonade + shell: bash + run: echo "name=${{ needs.resolve-upstream.outputs.tag }}" >> "$GITHUB_OUTPUT" - name: Pack artifacts id: pack_artifacts @@ -523,6 +597,7 @@ jobs: name: llama-ubuntu-cuda-${{ matrix.sm }}-arm64.tar.xz ubuntu-22-openvino: + needs: resolve-upstream runs-on: ubuntu-22.04 env: @@ -536,6 +611,7 @@ jobs: with: fetch-depth: 0 repository: 'ggml-org/llama.cpp' + ref: ${{ needs.resolve-upstream.outputs.sha }} - name: ccache uses: ggml-org/ccache-action@v1.2.21 @@ -647,7 +723,8 @@ jobs: - name: Determine tag name id: tag - uses: lemonade-sdk/llama.cpp/.github/actions/get-tag-name@lemonade + shell: bash + run: echo "name=${{ needs.resolve-upstream.outputs.tag }}" >> "$GITHUB_OUTPUT" - name: Pack artifacts run: | @@ -663,6 +740,7 @@ jobs: name: llama-bin-ubuntu-openvino-${{ env.OPENVINO_VERSION_MAJOR }}-x64.tar.gz windows-cpu: + needs: resolve-upstream runs-on: windows-2025 strategy: @@ -677,6 +755,7 @@ jobs: with: fetch-depth: 0 repository: 'ggml-org/llama.cpp' + ref: ${{ needs.resolve-upstream.outputs.sha }} - name: ccache uses: ggml-org/ccache-action@v1.2.21 @@ -734,6 +813,7 @@ jobs: name: llama-bin-win-cpu-${{ matrix.arch }}.zip windows-rocm: + needs: resolve-upstream runs-on: windows-2022 strategy: @@ -759,6 +839,7 @@ jobs: with: fetch-depth: 0 repository: 'ggml-org/llama.cpp' + ref: ${{ needs.resolve-upstream.outputs.sha }} - name: Cache ROCm Installation id: cache-rocm @@ -867,7 +948,8 @@ jobs: - name: Determine tag name id: tag - uses: lemonade-sdk/llama.cpp/.github/actions/get-tag-name@lemonade + shell: bash + run: echo "name=${{ needs.resolve-upstream.outputs.tag }}" >> "$GITHUB_OUTPUT" - name: Get ROCm short version run: | @@ -879,15 +961,37 @@ jobs: cp "LICENSE" "build\bin\" 7z a -snl llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip .\build\bin\* + - name: Record qualification metadata + run: | + $meta = [ordered]@{ + os = 'win' + rocm_version = '${{ env.ROCM_VERSION_SHORT }}' + rocm_version_full= '${{ matrix.ROCM_VERSION }}' + gpu_targets = '${{ matrix.gpu_targets }}' + tag = '${{ steps.tag.outputs.name }}' + upstream_sha = '${{ needs.resolve-upstream.outputs.sha }}' + artifact = 'llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip' + } + $path = "qualify-meta-win-rocm-${{ env.ROCM_VERSION_SHORT }}.json" + $meta | ConvertTo-Json | Set-Content -Path $path -Encoding utf8 + Get-Content $path + - name: Upload artifacts uses: actions/upload-artifact@v6 with: path: llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip name: llama-bin-win-rocm-${{ env.ROCM_VERSION_SHORT }}-${{ matrix.build }}.zip + - name: Upload qualification metadata + uses: actions/upload-artifact@v6 + with: + path: qualify-meta-win-rocm-${{ env.ROCM_VERSION_SHORT }}.json + name: qualify-meta-win-rocm-${{ env.ROCM_VERSION_SHORT }} + windows-cuda: runs-on: windows-2022 needs: + - resolve-upstream - windows-cpu strategy: @@ -904,6 +1008,7 @@ jobs: with: fetch-depth: 0 repository: 'ggml-org/llama.cpp' + ref: ${{ needs.resolve-upstream.outputs.sha }} - name: Install CUDA Toolkit uses: Jimver/cuda-toolkit@v0.2.35 @@ -1025,6 +1130,235 @@ jobs: path: llama-windows-cuda-${{ matrix.sm }}-x64.7z name: llama-windows-cuda-${{ matrix.sm }}-x64.7z + # --------------------------------------------------------------------------- + # ROCm build qualification. See .github/qualify-rocm/README.md. + # + # Tier 0 is static and runs on free hosted runners, including on PRs. + # Tier 1 runs real inference on the gfx1151 self-hosted boxes. + # Neither can block the CUDA/OpenVINO/CPU families: the release job drops + # only the ROCm family when qualification fails (see "Qualify ROCm release + # artifacts" there). + # --------------------------------------------------------------------------- + + qualify-tier0: + # !cancelled() rather than always(): if a ROCm build leg failed there is no + # archive to audit, and the release job's existing per-family completeness + # check is what reports that gap. This job vouches for whatever was built. + if: ${{ !cancelled() }} + needs: + - ubuntu-24-rocm + - windows-rocm + + strategy: + fail-fast: false + matrix: + # One leg per published archive. Each leg must run on the archive's own + # OS, because two of the checks launch llama-server out of the archive + # and install the ROCm a user of that archive would install. + os: [ubuntu, win] + rocm: ["7.14", "10.0", "10.1"] + include: + - os: ubuntu + runner: ubuntu-24.04 + - os: win + runner: windows-2025 + + runs-on: ${{ matrix.runner }} + + steps: + - name: Checkout qualification suite + uses: actions/checkout@v6 + + - name: Set up Python + uses: actions/setup-python@v6 + with: + python-version: '3.12' + + - name: Download ROCm archive + uses: actions/download-artifact@v7 + with: + name: llama-bin-${{ matrix.os }}-rocm-${{ matrix.rocm }}-x64.${{ matrix.os == 'win' && 'zip' || 'tar.gz' }} + path: ./artifact + + - name: Download qualification metadata + uses: actions/download-artifact@v7 + with: + name: qualify-meta-${{ matrix.os }}-rocm-${{ matrix.rocm }} + path: ./meta + + # Installs the ROCm an end user would install, so T0.12 can confirm the + # archive's GPU libraries will actually be found at run time. Allowed to + # fail: if the install does not succeed, T0.12 reports itself as skipped + # and every other Tier 0 check still runs. + - name: Install a user-style ROCm runtime + id: rocm + continue-on-error: true + shell: bash + run: | + python .github/qualify-rocm/install_rocm.py \ + --version "${{ matrix.rocm }}.0" \ + --arch gfx1151 \ + --prefix "${RUNNER_TEMP}/rocm-check" \ + --output-env "$GITHUB_ENV" + + - name: Run Tier 0 static verification + shell: bash + run: | + mkdir -p fragments + # --fail-on-error makes THIS leg go red as soon as a check fails, which is the only signal a pull request gets. + # The release job and the aggregate that normally decides what ships does not run on PR's. + python .github/qualify-rocm/tier0_static.py \ + --meta "meta/qualify-meta-${{ matrix.os }}-rocm-${{ matrix.rocm }}.json" \ + --archive-dir ./artifact \ + --run-binaries \ + --fail-on-error \ + --output "fragments/tier0-${{ matrix.os }}-rocm-${{ matrix.rocm }}.json" + + - name: Upload Tier 0 report + if: always() + uses: actions/upload-artifact@v6 + with: + name: qualify-fragments-tier0-${{ matrix.os }}-${{ matrix.rocm }} + path: fragments/*.json + if-no-files-found: warn + + qualify-tier1: + # Hardware validation on the Strix Halo (gfx1151) runners. + # Only one of the 20+ gfx targets in a build has test hardware. + # The qualification record states that explicitly rather than implying the rest were tested too. + # Runs on every trigger -- schedule, manual dispatch and pull request -- EXCEPT a pull request opened from a fork. + if: >- + ${{ !cancelled() && + (github.event_name != 'pull_request' || + github.event.pull_request.head.repo.full_name == github.repository) }} + needs: + - ubuntu-24-rocm + - windows-rocm + + strategy: + fail-fast: false + # Six legs: three ROCm versions (7.14, 10.0, 10.1) on each of Windows and Linux. + matrix: + include: + - os: ubuntu + rocm: "7.14" + runner: [stx-halo, Linux] + - os: ubuntu + rocm: "10.0" + runner: [stx-halo, Linux] + - os: ubuntu + rocm: "10.1" + runner: [stx-halo, Linux] + - os: win + rocm: "7.14" + runner: [xsj-dispatch, strix-halo, Windows] + - os: win + rocm: "10.0" + runner: [xsj-dispatch, strix-halo, Windows] + - os: win + rocm: "10.1" + runner: [xsj-dispatch, strix-halo, Windows] + + runs-on: ${{ matrix.runner }} + + steps: + - name: Checkout qualification suite + uses: actions/checkout@v6 + + - name: Download ROCm archive + uses: actions/download-artifact@v7 + with: + name: llama-bin-${{ matrix.os }}-rocm-${{ matrix.rocm }}-x64.${{ matrix.os == 'win' && 'zip' || 'tar.gz' }} + path: ./artifact + + - name: Download qualification metadata + uses: actions/download-artifact@v7 + with: + name: qualify-meta-${{ matrix.os }}-rocm-${{ matrix.rocm }} + path: ./meta + + - name: Resolve Python + shell: bash + run: | + PY=$(command -v python3 || command -v python) || { + echo "::error::No Python interpreter on this runner."; exit 1; } + echo "PY=$PY" >> "$GITHUB_ENV" + "$PY" --version + + - name: Extract archive + shell: bash + run: | + rm -rf extracted && mkdir -p extracted + "$PY" -I -c " + import os, sys, tarfile, zipfile + src = sys.argv[1] + if src.endswith('.zip'): + with zipfile.ZipFile(src) as z: z.extractall('extracted') + else: + with tarfile.open(src) as t: + try: t.extractall('extracted', filter='data') + except TypeError: t.extractall('extracted') + " "artifact/llama-bin-${{ matrix.os }}-rocm-${{ matrix.rocm }}-x64.${{ matrix.os == 'win' && 'zip' || 'tar.gz' }}" + # Linux tarballs nest under llama-/; Windows zips are flat. + root=$(dirname "$(find extracted -name 'llama-server*' -print -quit)") + [ -n "$root" ] || { echo "::error::llama-server not found in archive"; exit 1; } + echo "ARCHIVE_ROOT=$root" >> "$GITHUB_ENV" + echo "Archive root: $root" + + # The archives do not contain ROCm -- a user installs it separately. Without + # this step the Linux binary never starts, and the Windows one is worse: a + # GPU backend that fails to load is not treated as an error, so llama.cpp + # quietly runs on the CPU and the test passes while proving nothing. + - name: Install a user-style ROCm runtime + shell: bash + run: | + "$PY" .github/qualify-rocm/install_rocm.py \ + --version "${{ matrix.rocm }}.0" \ + --arch gfx1151 \ + --prefix "${RUNNER_TEMP}/rocm-qual" \ + --output-env "$GITHUB_ENV" \ + --output-json rocm-install.json + + - name: Download test model + shell: bash + run: | + url="https://huggingface.co/unsloth/Qwen3-0.6B-GGUF/resolve/main/Qwen3-0.6B-Q4_0.gguf?download=true" + curl -fL --retry 3 --retry-delay 5 -o Qwen3-0.6B-Q4_0.gguf "$url" + ls -l Qwen3-0.6B-Q4_0.gguf + + - name: Run Tier 1 GPU inference + shell: bash + run: | + mkdir -p fragments + "$PY" .github/qualify-rocm/tier1_gpu.py \ + --archive-root "$ARCHIVE_ROOT" \ + --model Qwen3-0.6B-Q4_0.gguf \ + --rocm-version "${{ matrix.rocm }}" \ + --gfx-target gfx1151 \ + --artifact "llama-bin-${{ matrix.os }}-rocm-${{ matrix.rocm }}-x64" \ + --logs-dir fragments \ + --fail-on-error \ + --output "fragments/tier1-${{ matrix.os }}-rocm-${{ matrix.rocm }}.json" + + - name: Upload Tier 1 fragments + if: always() + uses: actions/upload-artifact@v6 + with: + name: qualify-fragments-tier1-${{ matrix.os }}-${{ matrix.rocm }} + path: | + fragments/*.json + fragments/*.log + if-no-files-found: warn + + - name: Tear down + if: always() + shell: bash + run: | + # The stx-halo boxes are persistent; never leave a server holding VRAM. + pkill -f llama-server 2>/dev/null || true + taskkill //F //IM llama-server.exe 2>/dev/null || true + rm -rf extracted artifact Qwen3-0.6B-Q4_0.gguf || true + release: # always() overrides the default behavior of skipping this job when a # needed build job fails or is skipped (e.g. one CUDA sm_* matrix leg @@ -1041,6 +1375,7 @@ jobs: runs-on: ubuntu-24.04 needs: + - resolve-upstream - windows-cpu - windows-rocm - windows-cuda @@ -1048,6 +1383,8 @@ jobs: - ubuntu-22-cuda - ubuntu-22-cuda-arm64 - ubuntu-22-openvino + - qualify-tier0 + - qualify-tier1 steps: - name: Clone @@ -1056,10 +1393,17 @@ jobs: with: fetch-depth: 0 repository: 'ggml-org/llama.cpp' + ref: ${{ needs.resolve-upstream.outputs.sha }} + + - name: Checkout qualification suite + uses: actions/checkout@v6 + with: + path: lemonade-ci - name: Determine tag name id: tag - uses: lemonade-sdk/llama.cpp/.github/actions/get-tag-name@lemonade + shell: bash + run: echo "name=${{ needs.resolve-upstream.outputs.tag }}" >> "$GITHUB_OUTPUT" # Downloaded in batches rather than one call: actions/download-artifact # starts every matched artifact concurrently, and 28 parallel extractions @@ -1195,7 +1539,13 @@ jobs: # overlap with any other family's files. declare -A family_glob=( [cuda]="release/llama-${TAG}-*cuda-sm_*" - [rocm]="release/llama-${TAG}-bin-*rocm-*" + # Each ROCm version is its own family. They are independent + # products: a consumer downloads exactly one of them, so a gap in + # 7.14 is no reason to withhold a perfectly good 10.0 build. What + # must stay consistent is the Windows+Linux pair WITHIN a version. + [rocm-7.14]="release/llama-${TAG}-bin-*rocm-7.14-*" + [rocm-10.0]="release/llama-${TAG}-bin-*rocm-10.0-*" + [rocm-10.1]="release/llama-${TAG}-bin-*rocm-10.1-*" [openvino]="release/llama-${TAG}-bin-*openvino-*" [cpu]="release/llama-${TAG}-bin-win-cpu-*" ) @@ -1220,8 +1570,17 @@ jobs: check_glob cuda "release/llama-${TAG}-windows-cuda-${sm}-x64.7z" done - check_glob rocm "release/llama-${TAG}-bin-ubuntu-rocm-*-x64.tar.gz" - check_glob rocm "release/llama-${TAG}-bin-win-rocm-*-x64.zip" + # Enumerate every ROCm version explicitly, the way the CUDA sm_* loop + # above does. A bare `-rocm-*-` glob is satisfied by ANY single + # version, so a release missing one version entirely would still pass. + # That matters because a consumer asks for one exact ROCm version by + # name -- a missing version is a failed download on every machine + # that wanted it, not a graceful fallback to another one. + # Keep this list in step with the build matrices. + for rocm in 7.14 10.0 10.1; do + check_glob "rocm-${rocm}" "release/llama-${TAG}-bin-ubuntu-rocm-${rocm}-x64.tar.gz" + check_glob "rocm-${rocm}" "release/llama-${TAG}-bin-win-rocm-${rocm}-x64.zip" + done check_glob openvino "release/llama-${TAG}-bin-ubuntu-openvino-*-x64.tar.gz" check_glob cpu "release/llama-${TAG}-bin-win-cpu-x64.zip" @@ -1251,6 +1610,110 @@ jobs: echo "All expected artifacts present. Found $files artifact(s) ready to upload." fi + - name: Download qualification fragments + uses: actions/download-artifact@v7 + with: + pattern: qualify-fragments-* + path: ./fragments + merge-multiple: true + + - name: Download qualification metadata + uses: actions/download-artifact@v7 + with: + pattern: qualify-meta-* + path: ./meta + merge-multiple: true + + - name: Post-merge checks on final artifacts + id: post_merge + env: + TAG: ${{ steps.tag.outputs.name }} + run: | + shopt -s nullglob + mkdir -p fragments + for archive in release/llama-${TAG}-bin-ubuntu-rocm-*.tar.gz \ + release/llama-${TAG}-bin-win-rocm-*.zip; do + base=$(basename "$archive") + if [[ "$base" == *-win-rocm-* ]]; then os=win; else os=ubuntu; fi + ver=$(sed -E 's/.*-rocm-([0-9]+\.[0-9]+)-x64\..*/\1/' <<<"$base") + echo "::group::Post-merge — $base" + python3 lemonade-ci/.github/qualify-rocm/check_post_merge.py \ + --archive "$archive" --os "$os" --rocm-version "$ver" --tag "$TAG" \ + --output "fragments/tier0-postmerge-$os-rocm-$ver.json" || true + echo "::endgroup::" + done + + - name: Qualify ROCm release artifacts + id: qualify + continue-on-error: true + env: + TAG: ${{ steps.tag.outputs.name }} + run: | + mkdir -p qualification + python3 lemonade-ci/.github/qualify-rocm/aggregate.py \ + --fragments-dir ./fragments \ + --tag "$TAG" \ + --records-dir ./qualification \ + --output-summary ./qualification/qualification-summary.json \ + --fail-on-no-promote + + - name: Drop unqualified ROCm versions + id: rocm_gate + env: + TAG: ${{ steps.tag.outputs.name }} + PROMOTE: ${{ steps.qualify.outputs.promote_rocm }} + BLOCKED: ${{ steps.qualify.outputs.blocked_rocm_versions }} + run: | + if [ "$PROMOTE" = "true" ]; then + echo "Every ROCm version qualified — all ROCm artifacts will be published." + echo "dropped=0" >> "$GITHUB_OUTPUT" + exit 0 + fi + + # The aggregate step names the versions that failed. If it could not + # run at all we have no evidence for any of them, so withhold the lot + # rather than publish something unverified. + blocked="$BLOCKED" + if [ -z "$blocked" ]; then + blocked="7.14 10.0 10.1" + echo "::warning::Qualification produced no verdict — withholding every \ + ROCm version rather than publishing unverified artifacts." + fi + + dropped=0 + for version in $blocked; do + count=$(ls release/llama-${TAG}-bin-*rocm-${version}-* 2>/dev/null | wc -l) + [ "$count" -gt 0 ] || continue + echo "::warning::ROCm ${version} did not qualify — removing its \ + ${count} artifact(s) from this release. Other ROCm versions and every \ + other backend are unaffected. The attached qualification record lists \ + the failing checks." + rm -fv release/llama-${TAG}-bin-*rocm-${version}-* || true + dropped=$((dropped + count)) + done + echo "dropped=$dropped" >> "$GITHUB_OUTPUT" + echo "blocked=$blocked" >> "$GITHUB_OUTPUT" + + # Lets a user prove the bytes they downloaded are the bytes we uploaded. + - name: Publish SHA-256 checksums + run: | + shopt -s nullglob + files=(release/*.zip release/*.tar.gz release/*.tar.xz release/*.7z) + if [ ${#files[@]} -eq 0 ]; then + echo "No artifacts to checksum." + exit 0 + fi + ( cd release && sha256sum $(for f in "${files[@]}"; do basename "$f"; done) \ + > checksums.txt ) + cat release/checksums.txt + + - name: Stage qualification records as release assets + run: | + shopt -s nullglob + for record in qualification/*.json; do + cp -v "$record" release/ + done + # Get the release for this tag, creating it if it does not exist yet. # This is idempotent: if a previous run created the release but failed # partway through uploading (e.g. a transient GitHub outage), re-running @@ -1335,8 +1798,13 @@ jobs: } } + // checksums.txt and the qualification-*.json records ship alongside + // the archives: the first lets a user verify a download, the second + // records which gfx targets were present in the binary vs actually + // validated on hardware. const files = fs.readdirSync('./release').filter((f) => - f.endsWith('.zip') || f.endsWith('.tar.gz') || f.endsWith('.tar.xz') || f.endsWith('.7z')); + f.endsWith('.zip') || f.endsWith('.tar.gz') || f.endsWith('.tar.xz') || + f.endsWith('.7z') || f === 'checksums.txt' || f.endsWith('.json')); const wanted = new Set(files); // Reconcile existing release assets against what this run decided to @@ -1395,3 +1863,9 @@ jobs: run: | echo "::error::${{ steps.check_artifacts.outputs.missing }} artifact(s) were missing (${{ steps.check_artifacts.outputs.summary }}) — the incomplete families were dropped from this release; other families were still published. See the 'Check release artifacts' step for the full per-file list." exit 1 + + - name: Fail if ROCm qualification did not pass + if: steps.rocm_gate.outputs.dropped != '0' + run: | + echo "::error::ROCm ${{ steps.rocm_gate.outputs.blocked }} did not qualify — ${{ steps.rocm_gate.outputs.dropped }} artifact(s) were dropped from this release. Every other ROCm version and every other backend was published normally. The qualification-*.json records attached to the release list each failing check." + exit 1