diff --git a/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml b/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml new file mode 100644 index 0000000000..94d45396a9 --- /dev/null +++ b/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml @@ -0,0 +1,89 @@ +name: "RDNA hardware report" +description: "Confirm support on RDNA hardware you have, or report what is missing to add it" +labels: ["rdna", "hardware-report"] +body: + - type: markdown + attributes: + value: | + Use this to report **measured** results from RDNA hardware — yours or a machine you have access to. + + **A correctness failure is a first-class report.** If `test-backend-ops` fails on your arch, + post that and stop — do not tune a kernel that computes the wrong answer. That report is + worth more than a fast number. + + Run the probe first and paste its output below: + + ```bash + ./examples/rdna/rdna-support-probe.sh # identity only + ./examples/rdna/rdna-support-probe.sh --bench # with a real benchmark (preferred) + ``` + + Agents working this repo: read `AGENTS.md` first — §1 confirm, §2 add, §3 playbook when it breaks. + + **Negative results are welcome.** "This failed on gfx1102, here is the exact error" is a real + contribution and nobody without your hardware can produce it. + + - type: dropdown + id: correctness + attributes: + label: Correctness gate + description: | + Did `test-backend-ops` pass against the CPU reference on this arch? Throughput without + this proves nothing — a kernel that is fast and wrong benchmarks beautifully. + options: + - "PASSED — test-backend-ops passed on this arch" + - "FAILED — ops failed vs the CPU reference (please paste which ones below)" + - "NOT RUN — no test-backend-ops build available" + validations: + required: true + + - type: dropdown + id: verdict + attributes: + label: Probe verdict + description: What did the probe print? + options: + - "CONFIRMED — board resolves and a measurement was accepted by the harness" + - "PARTIAL — board resolves, no accepted measurement yet" + - "UNSUPPORTED — arch is not mapped" + - "The probe itself failed to run" + validations: + required: true + + - type: textarea + id: probe_output + attributes: + label: Probe output + description: Paste the whole block verbatim. Do not tidy the numbers or omit failures. + render: markdown + validations: + required: true + + - type: textarea + id: not_tested + attributes: + label: What you did NOT test + description: | + Be specific — this is what keeps the support matrix honest. GPU count, other models, + multi-GPU, other quantisations, anything you skipped or could not run. + placeholder: "Single GPU only; one model (Q4_K_M); did not test multi-GPU or fp8." + validations: + required: true + + - type: textarea + id: model_and_stack + attributes: + label: Model and llama.cpp build + description: Which GGUF, and which llama.cpp (upstream commit, or a fork + branch)? + validations: + required: false + + - type: textarea + id: notes + attributes: + label: Anything surprising + description: | + Hangs, wrong output, throughput that looks too good (often a cache artefact), counters + reading zero, anything you had to work around to get it running. + validations: + required: false diff --git a/AGENTS.md b/AGENTS.md new file mode 100644 index 0000000000..0c89d599aa --- /dev/null +++ b/AGENTS.md @@ -0,0 +1,266 @@ +# AGENTS.md — instructions for coding agents working in this repo + +This is the RDNA fork of [AMD-AGI/Hyperloom](https://github.com/AMD-AGI/Hyperloom). +Its job is to make Radeon-class AMD GPUs usable with Hyperloom **and to be +honest about which ones are actually proven**. + +If you are an agent operating here, your task is almost always one of two +things. Both end with **posting a report**, because a result nobody can see is +the same as no result. + +1. **CONFIRM support** on hardware the user physically has → §1 +2. **ADD support** for a gfx target that is not mapped yet → §2 + +and when either one goes wrong, **§3 is the playbook** for working out why. + +--- + +## THE ONE RULE + +**A claim is MEASURED or it is ASSUMED. Never publish an ASSUMED claim as support.** + +- MEASURED = the command ran **on this machine, this session**, and you are + pasting its real output. +- ASSUMED = theory, a spec sheet, a related arch, a previous session, "it should + work because gfx1100 and gfx1102 are both RDNA3." + +Assumptions are useful for deciding what to *test*. They never *close* a +question. A support matrix entry is a promise to a stranger whose hardware you +cannot see — if you did not run it, say "untested" and move on. That is a +perfectly good answer and costs nobody anything. A wrong "supported" costs +someone a day. + +Corollaries that have already bitten this project: + +- **Do not grade your own output.** Where a harness can judge a result, let it + (§1.3 uses Hyperloom's own `is_valid_measurement()`, not your reading of the + JSON). "The JSON looks right" is not a verdict. +- **A missing tool must fail loudly.** A silent empty result is + indistinguishable from "the feature is absent" — which is the exact wrong + answer from a probe whose whole job is establishing truth. +- **Verify the path you are testing is the live one.** From upstream + [#1041](https://github.com/AMD-AGI/Hyperloom/issues/1041): compare the actual + rendered source that was compiled, rather than inferring engagement from log + markers. A knob that looks plausible but never reaches the GPU will otherwise + "work". +- **Measured on A does not close B.** Same ISA family is not the same board. + +--- + +## §1 — CONFIRM support on hardware you have + +### 1.1 Run the probe + +```bash +./examples/rdna/rdna-support-probe.sh # identity + resolution only, no GPU work +./examples/rdna/rdna-support-probe.sh --bench # adds a real benchmark + harness verdict +``` + +Use `--bench` whenever a model is available (`MODEL=/path/to/*.gguf` and a +llama.cpp build). Without it you can only confirm that the board is +*recognised* — which proves the identity plumbing works and nothing about +whether this GPU computes correct answers. + +The probe prints a paste-ready markdown block and one of three verdicts: + +| Verdict | Means | +|---|---| +| `CONFIRMED` | Correctness **passed** against the CPU reference **and** the runner produced a measurement Hyperloom's own validator accepted. Both halves, or it is not confirmed | +| `PARTIAL` | One half is missing. `PARTIAL (throughput only — CORRECTNESS UNVERIFIED)` is called out separately, because it is the dangerous one | +| `UNSUPPORTED` | The arch is not in `_GFX_TO_RUNNER` — go to §2 | + +### 1.1a Correctness is half the verdict, not a footnote + +**A throughput number with no correctness gate is not evidence of support.** A +kernel that is fast and *wrong* produces a beautiful benchmark — that is the +normal failure mode of a bad port, not an exotic one. This is also the first +thing Hyperloom's own methodology insists on: correctness gate *before* any +benchmark is trusted. + +`--bench` therefore runs llama.cpp's `test-backend-ops`, which compares backend +ops against the **CPU reference on your arch**, and refuses `CONFIRMED` without +it. Defaults to `MUL_MAT,MUL_MAT_ID` — the matmul paths a quant or arch port +actually breaks — which takes **~90 s on gfx1201** (measured: 1196/1196 and +865/865 against the CPU reference). That is a one-time cost for a support +claim. Widen with `CORRECTNESS_OPS=ALL` (slower, more thorough), or narrow it +while you are iterating. + +Add a model-level gate on top when you can — `PPL_FILE= PPL_MAX=` +runs perplexity through the runner. Op-level correctness proves the kernels +compute the right thing; a perplexity bound proves the assembled model still +produces sane output. They fail differently, so both are worth having. + +Note what `is_valid_measurement()` does **not** do: a *missing* `quality_gate` +is non-blocking there, so the harness will happily accept a fast run with no +correctness evidence at all. That is correct behaviour for the harness and the +wrong bar for a support claim, which is why this probe adds its own. + +### 1.2 Do not edit the matrix from a PARTIAL + +`docs/rdna-port/README.md`'s support matrix moves to "MEASURED" only on +`CONFIRMED`. A `PARTIAL` is still worth posting — it tells us the identity +plumbing works on your board — but it is not support. + +### 1.3 Post it + +Open an issue using the **RDNA hardware report** template and paste the probe +block verbatim. Do not summarise it, do not tidy the numbers, and do not omit +the failures: a report whose honest verdict is "this regressed" or "this did +not build" is worth more than a clean one, because it is the one nobody else +can produce. + +State plainly what you did **not** test. Card count matters (a 4-GPU host is +not a validated 2-GPU host), and so does anything you skipped. + +--- + +## §2 — ADD support for a new gfx target + +The port for gfx1201 is the template — read that commit first. Adding a board +is small, because upstream derives the accepted-board set and the `rocm-smi` +product tags from **one table**. + +### 2.1 The change + +1. **`src/hyperloom/common/gpu_identity.py`** — one row: + `"": ("", )`. + Read the CU count **off the board** (`amd-smi static -g 0`), never a spec + sheet or a web search. This single row also makes the board CLI-accepted and + autodetectable, because `_AMD_GPU_TYPES` and `_PRODUCT_TAGS` derive from it. + The board key must be a substring of what `rocm-smi --showproductname` + prints, or autodetect will not match it. +2. **`src/hyperloom/inference_optimizer/gpu_types.py`** — a `_GFX_TO_RUNNER` + entry mapping the gfx arch to the runner label. +3. **A runner script** in `examples/rdna/`. `custom_r9700.sh` is llama.cpp-based + and should port to any RDNA target with a rename and *its own* measured + caveats. Keep the contract: read `$MODEL`/`$RESULT_DIR`/`$RESULT_FILENAME`, + write an InferenceX-shaped `inferencex_result.json`, refuse to emit a + zero-throughput "success". +4. **Tests** in `src/hyperloom/inference_optimizer/tests/`. Mirror + `test_rdna4_r9700_support.py`, **including its negative tests**. If upstream + asserts your arch is unmapped (it asserts this for gfx1100 in + `test_profile_and_kernel_handlers.py`), flip that assertion into positive + coverage in the same commit. +5. **The support matrix** in `docs/rdna-port/README.md`, plus any profiling + caveats that differ from gfx1201's. + +### 2.2 Map only what you ran + +Do not add a `_GFX_TO_RUNNER` entry for a sibling chip you have not booted. +gfx1200 (Navi 44) is deliberately absent here for exactly this reason, and +there is a test locking that absence. Adding an untested arch is not +generosity — it is a support claim written by someone who cannot honour it. + +### 2.3 Verify before you post + +```bash +PY= +PYTHONPATH=src $PY -m pytest src/hyperloom/inference_optimizer/tests/ -q \ + -k "gpu or quant or preflight or parser or provenance" +MODEL=/path/model.gguf LLAMA_CPP_DIR=/path/llama.cpp \ + ./examples/rdna/rdna-support-probe.sh --bench +``` + +Both must pass, and the probe must say `CONFIRMED` — which means the +correctness gate passed on your arch, not merely that it went fast — before the +matrix changes. + +If correctness FAILS on a new arch, that is the most valuable thing you can +report. Post it with the failing op list and stop; do not tune a kernel that is +computing the wrong answer, and do not add a matrix row. + +--- + +## §3 — Playbook: when it does not work + +Most "unsupported hardware" is not unsupported hardware. Work it in this order. + +### 3.1 Exhaust the mundane before the exotic + +When a counter reads zero, a feature reports unsupported, or a kernel +underperforms, check **in this order**, and stop at the first thing that +explains it: + +1. **perf-level / power gating** (a zero counter is usually a gate, not silicon) +2. **env-var traps** (a stray `ROCR_VISIBLE_DEVICES`, `HSA_*`, `HIP_*` left set) +3. **the vendor tracker** — search ROCm/llama.cpp issues before theorising; the + bug is often known, with a workaround +4. **driver / firmware / runtime versions** (and whether the one you *think* is + loaded is the one actually loaded — check `ldd`, not your memory) +5. **build flags / target arch** (was it even compiled for this gfx?) +6. **only then** tile sizes, occupancy, ILP, kernel internals + +Most "impossible" walls are a mundane switch. Reaching for the exotic +explanation first is the single most expensive habit in this work. + +### 3.2 Isolate by layer, bottom-up + +Name the layer the fault lives in, verify with **that layer's own ground-truth +tool**, and do not blame one layer for another's problem. Cross-layer +misattribution — blaming silicon for a compiler artifact, or the model for an +environment trap — sends you optimising the wrong thing entirely. + +| Layer | Ground truth | Classic trap | +|---|---|---| +| L0 board / link / topology | `lspci -vv`, transfer benchmarks | Checking only the GPU's own link. Navi boards sit behind an on-board switch whose **upstream** port is where the motherboard downgrades you — the leaf can read Gen5 while the real path is Gen3 | +| L1 silicon / ISA | the assembler (`llvm-mc`) — it cannot lie | Assuming a marketing feature is present | +| L2 firmware / driver | `amd-smi`, `dmesg`, sysfs perf level | A counter reading 0 is a **gate**, not broken hardware | +| L3 ROCm runtime / env | `ldd`, `HIP_VISIBLE_DEVICES`, `LD_LIBRARY_PATH` | A stale env var causing silent CPU fallback; the loader picking a different ROCm than you meant | +| L4 compiler / toolchain | `--save-temps` + disassembly, target features | A builtin that **arity-checks before target-feature**, so an unavailable instruction looks available | +| L5 framework kernels | disassembly; correctness vs a CPU reference | "Instruction unused" concluded from a *sample*; the kernel may exist but be dormant or undispatched | +| L6 model / quant format | bit round-trip vs reference | Trusting a pack layout without a bit-identical proof | +| L7 serving | served tokens, VRAM, throughput | Blaming the model for an L2/L3 fault (garbled output, CPU fallback, a reaped process) | + +### 3.3 Specific failures, and what they usually are + +| Symptom | Look here first | +|---|---| +| Probe says `UNSUPPORTED` but the GPU works | The arch simply is not mapped — §2. Not a hardware problem | +| Board not detected, arch reads `unknown` | `rocm-smi`/`rocminfo` absent (a minimal runtime install omits them; the probe falls back to `amd-smi`), or the board key is not a substring of the product name | +| Runner exits 3, "llama-bench not found" | Working as designed — it refuses to report a zero rather than let a missing binary read as a slow GPU. Set `LLAMA_BENCH` / `LLAMA_CPP_DIR` | +| Benchmark runs, validator REJECTS | `output_throughput` is 0/absent, or a `quality_gate` failed. A failed gate makes a run unselectable **no matter how fast it was** — that is the gate working | +| Throughput looks impossibly good | You are probably measuring cache, not the real path. If a result exceeds the memory roofline, the working set fits in cache — enlarge it until it does not | +| A/B shows a win | Check the direction: `speedup = baseline ÷ candidate`. Inverted ratios have reported regressions as wins here before | +| A/B shows a huge win | Suspect the **baseline**, not the candidate. Compare against a competent baseline, not a strawman; a broken baseline (NaN/Inf, a pathological path) manufactures spectacular speedups | +| Numbers move run to run | Interleave A/B, discard warmup, take medians of ≥3, and re-measure cold. A number you did not measure *this run* does not exist | + +### 3.4 Before you claim a fix + +- It **compiled and ran on the target**, and you have the output. +- A **correctness gate passed** — versus a CPU reference or a quality metric. + A wrong-but-fast kernel is worth zero. +- The path you changed is the one that actually **executed**. Verify by + comparing the rendered source or the disassembly, not by trusting a log line. +- Say which claims are MEASURED and which are ASSUMED. Both belong in the + report; only one belongs in the support matrix. + +Negative results are first-class here. "This does not work on gfx1102, here is +the exact failure" is a real contribution, and often a more useful one than a +win, because nobody else can produce it without your hardware. + +## Repo conventions + +- **Branches.** `rdna` is the working branch. **`main` tracks upstream + byte-for-byte — never commit to it.** PR branches for upstream are cut from + `main` and cherry-pick only port commits, so they stay reviewable. +- **Never open a PR against AMD-AGI/Hyperloom without explicit human approval.** + Same for posting on upstream issues. Draft it and ask. +- **Licence.** Upstream files stay MIT © AMD with their SPDX headers intact. + New files carry their own SPDX header. This fork is not an AMD product and + must not present itself as one. +- **Commits** explain *why*, and say which claims were measured. + +## RDNA profiling traps + +Do not report profiling numbers from an RDNA4 box without accounting for these +(all measured on gfx1201; re-verify on your own arch rather than assuming they +carry): + +| Trap | Why it misleads | +|---|---| +| `--setperflevel high` | **Reduces** inference throughput ~10% vs `auto` while tightening variance — pinning clocks for "stable" absolutes makes the number worse | +| `profile_standard` | Un-gates the memory counters but throttles to ~1593 MHz vs ~2330 — it suppresses any benchmark running beside it. Prefer `profile_peak` | +| `GL2C` reading zero | Perf-level gating, not broken firmware | +| `FETCH_SIZE` and some derived PMC counters | Collect zeros — reads as "no traffic" rather than "unsupported" | +| PC sampling | `host_trap` only, min interval 512; `stochastic` is unsupported | +| No MFMA on RDNA | Anything reaching for CDNA matrix cores must branch on gfx arch, not on "is AMD" | diff --git a/README.md b/README.md index 616f79fe6b..110d255084 100755 --- a/README.md +++ b/README.md @@ -1,3 +1,26 @@ +> ### 🔴 RDNA fork — community port, not an AMD product +> +> This is a fork of [AMD-AGI/Hyperloom](https://github.com/AMD-AGI/Hyperloom) +> (MIT, © Advanced Micro Devices, Inc.) adding **RDNA4 / gfx1201** support — +> Radeon AI PRO R9700 and friends. It is **not endorsed by or affiliated with +> AMD**, and upstream is the place for anything not RDNA-specific. +> +> - **Start here → [`docs/rdna-port/README.md`](docs/rdna-port/README.md)** — +> support matrix (what is measured vs. untested), quickstart, and the RDNA4 +> profiling caveats that silently break Instinct-shaped tooling. +> - The llama.cpp runner is [`examples/rdna/custom_r9700.sh`](examples/rdna/custom_r9700.sh); +> report your own hardware with [`examples/rdna/rdna-support-probe.sh`](examples/rdna/rdna-support-probe.sh). +> - **Agents:** read [`AGENTS.md`](AGENTS.md) — confirm support, add a target, or work a failure. +> - Tracking issues upstream: [#1196](https://github.com/AMD-AGI/Hyperloom/issues/1196) (RDNA4, ours) and +> [#1041](https://github.com/AMD-AGI/Hyperloom/issues/1041) (RDNA3, [@kyuubyN](https://github.com/kyuubyN) — +> methodology independently validated on gfx1100/gfx1102). **RDNA3 runner support is welcome here.** +> - Branches: `rdna` (this work, default) · `main` (tracks upstream byte-for-byte +> so PR branches cherry-pick cleanly). +> +> Everything below is upstream's README, unmodified. + +--- + # ROCm Hyperloom [![Tests](https://github.com/AMD-AGI/Hyperloom/actions/workflows/tests-coverage.yml/badge.svg)](https://github.com/AMD-AGI/Hyperloom/actions/workflows/tests-coverage.yml) diff --git a/docs/rdna-port/README.md b/docs/rdna-port/README.md new file mode 100644 index 0000000000..66096c8abd --- /dev/null +++ b/docs/rdna-port/README.md @@ -0,0 +1,151 @@ +# RDNA support for Hyperloom (gfx12 / RDNA4) + +Community port. **Not an AMD product and not endorsed by AMD** — this is a fork +of [AMD-AGI/Hyperloom](https://github.com/AMD-AGI/Hyperloom) (MIT, © AMD) +adding consumer/workstation Radeon support. Upstream tracks this as +[issue #1196](https://github.com/AMD-AGI/Hyperloom/issues/1196). + +Upstream validates against Instinct only; `_GFX_TO_RUNNER` maps consumer arches +to `None`, so a Radeon box resolves to no runner and no dispatch identity. This +branch makes gfx1201 a first-class target. + +## Support matrix — what is measured vs. what is not + +| Target | Arch | Status | +|---|---|---| +| Radeon AI PRO R9700 | gfx1201 | **MEASURED, and reproduced in a clean container against stock upstream llama.cpp** (see `examples/rdna/container/`). Board autodetects; `test-backend-ops` passes vs the CPU reference (MUL_MAT 1196/1196, MUL_MAT_ID 865/865); runner benchmarked end-to-end on 2× R9700 (prefill ~4476 t/s, decode ~169 t/s, 8B Q2_0) | +| RX 9070 / 9070 XT | gfx1201 | Same arch, board never tested here. Autodetect keys on arch, so it should resolve — unverified | +| RX 9060 XT | gfx1200 | **Deliberately unmapped.** Navi 44 is the same ISA family but nobody has measured it; a runner label is a claim of support | +| W7900 / RX 7600 | gfx1100 / gfx1102 | Methodology **independently validated** by [@kyuubyN](https://github.com/kyuubyN) — see below. Runner map not covered here yet | +| gfx11 (other RDNA3) | gfx1101 etc. | Not in this branch yet | +| Instinct | gfx942/gfx950 | Unchanged from upstream | + +Multi-GPU: measured on a **2-card** host only. On this box the two GPUs report +`peers: false` and host↔GPU is capped at PCIe Gen3 (~12.9 GB/s measured, an +on-board switch upstream port downgrades the link). A 4-card rig will hit +different walls; nothing here is validated at that width. + +## Quickstart + +```bash +export HYPERLOOM_BYPASS_SCRIPTS_DIR=$PWD/examples/rdna +export MODEL=/path/to/model.gguf +export LLAMA_CPP_DIR=/path/to/llama.cpp # expects build/bin/llama-bench +``` + +Hyperloom resolves `custom_r9700.sh` from that directory by the +`custom_{runner_type}.sh` convention once the board is detected as `r9700`. + +Verify detection before running anything long: + +```python +from hyperloom.inference_optimizer.gpu_types import amd_gpu_dispatch_identity +assert amd_gpu_dispatch_identity() == ("gfx1201", 64) +``` + +Run the runner standalone (it is a normal script — no Hyperloom required): + +```bash +MODEL=/path/model.gguf LLAMA_BENCH=/path/llama-bench RESULT_DIR=/tmp/out \ + PP=512 TG=64 REPS=3 ./examples/rdna/custom_r9700.sh +``` + +### Why llama.cpp and not vLLM/SGLang + +Below ~32 GB the serving stacks' VRAM and concurrency assumptions do not +transfer, and llama.cpp is what this class of card is actually served with. The +runner is a **scriptable** (server-less) workload: no OpenAI endpoint, no HTTP +client. It writes the InferenceX-shaped `inferencex_result.json` that +Hyperloom's collectors already consume, so nothing downstream needed changing. + +### Correctness gate + +`PPL_FILE` + `PPL_MAX` run `llama-perplexity` and emit a `quality_gate`. A run +that fails the gate is **unselectable regardless of throughput** — which is the +point of shipping one, since a wrong-but-fast kernel otherwise wins the +benchmark. With `PPL_FILE` set but no `PPL_MAX`, the perplexity is recorded and +explicitly marked `skipped` rather than being silently treated as a pass. + +## RDNA4 profiling realities + +These will silently break tooling that assumes Instinct. All measured on +gfx1201: + +| Reality | Consequence | +|---|---| +| PC sampling is **host_trap only**, min interval 512 | `stochastic` is unsupported; a harness requesting it gets nothing | +| Several PMC derived counters (e.g. `FETCH_SIZE`) collect **zeros** | Reads as "no memory traffic" rather than as an unsupported counter | +| `GL2C` counters need a **fixed DPM profile** to leave zero | The counters are perf-level gated, not broken firmware | +| `profile_peak` vs `profile_standard` | Both un-gate the counters, but `profile_standard` **throttles to ~1593 MHz** vs ~2330 MHz — it suppresses any benchmark running alongside it. Prefer `profile_peak` | +| `--setperflevel high` **reduces** inference throughput ~10% vs `auto` | A trap for any harness that pins clocks for "stable" absolutes | +| No MFMA on this arch | Anything reaching for CDNA matrix cores must branch on gfx arch, not on "is AMD" | + +## Quantization schemes on gfx1201 + +`fp8` / `ptpc_fp8` are available; RDNA4 has WMMA fp8 and a native E4M3 path. + +`mxfp4` / `mxfp4_fp8` are **unavailable**, and for a stronger reason than the +upstream MI355X gate: the MX scaled-convert instructions are gfx950/gfx1250 +gated, so an assembler capability probe on gfx1201 grades them REJECTED. There +is no hardware datapath to fall back to. + +## Reporting from your own hardware + +Run the probe and open an **RDNA hardware report** issue with its output: + +```bash +./examples/rdna/rdna-support-probe.sh --bench # MODEL=... LLAMA_CPP_DIR=... +``` + +It prints a paste-ready block and a verdict — `CONFIRMED`, `PARTIAL`, or +`UNSUPPORTED` (with the exact change needed to map your arch). `CONFIRMED` +requires **both** a passing correctness gate (`test-backend-ops` vs the CPU +reference on your arch) and a measurement Hyperloom's validator accepts: +throughput alone cannot confirm support, because a kernel that is fast and +wrong benchmarks beautifully. Agents working +in this repo should read [`AGENTS.md`](../../AGENTS.md): §1 confirm, §2 add a +target, §3 the playbook for when it does not work. + +## RDNA3 (gfx11) — what is actually established + +Upstream [issue #1041](https://github.com/AMD-AGI/Hyperloom/issues/1041) is the +RDNA3 counterpart to #1196, opened by [@kyuubyN](https://github.com/kyuubyN). +Be precise about what it demonstrates, because it is easy to overstate: + +**What was validated:** Hyperloom's *methodology* — bounded discrete catalog, +correctness gate before any benchmark is trusted, interleaved A/B with warmup +discarded, non-overlapping-IQR significance test, graded +KEEP/NEEDS_REVIEW/REVERT verdict, and an engagement proof that the candidate +actually compiled and ran — independently re-implemented against **their own +HIP inference engine**, on gfx1102 (RX 7600, 8 GB) and gfx1100 (W7900, 48 GB), +with real measured results including a run whose honest verdict was "revert +everything". Write-up: +[Radeon-hackathon-2026-07-Aetheris](https://github.com/kyuubyN/Radeon-hackathon-2026-07-Aetheris/blob/main/PROJECT_SPECIFICATION.md#61-self-optimizing-kernels-built-on-amds-own-methodology). + +**What was NOT validated:** Hyperloom itself running on RDNA3. The runner map +still excludes gfx11, and upstream's suite *asserts* the exclusion +(`test_profile_and_kernel_handlers.py`: `_GFX_TO_RUNNER.get("gfx1100") is None`). +So "verified on RDNA3" means the approach transfers to gfx11 — not that this +port runs there. Nobody should read the matrix above as more than that. + +That distinction is the same discipline the port itself is built on: a result +measured on A does not close B until B is measured. + +One finding from #1041 worth carrying over regardless of arch: their engagement +check compares the **actual rendered source each instance compiled**, rather +than inferring engagement from log markers — and it caught a tuning knob that +looked plausible but never reached the GPU. That is the same failure mode as +optimizing a code path that is not the live dispatched one, and it is worth +having on any arch. + +**RDNA3 contributions are welcome here.** What it needs: a `gfx1100`/`gfx1102` +entry in the identities table plus `_GFX_TO_RUNNER`, a runner script (the +llama.cpp one in `examples/rdna/` should port with little more than a name and +different measured caveats), and flipping upstream's negative gfx1100 +assertion into positive coverage. The gfx1201 work is the template. + +## Upstream posture + +`main` on this fork tracks `upstream/main` byte-for-byte so PR branches cherry- +pick cleanly. AMD has said RDNA is triaged after 2026-08-30; this branch exists +so the hardware is usable now, not to route around that. diff --git a/docs/rdna-port/SCOPE.md b/docs/rdna-port/SCOPE.md new file mode 100644 index 0000000000..6b388028b7 --- /dev/null +++ b/docs/rdna-port/SCOPE.md @@ -0,0 +1,20 @@ +# RDNA support port — scoping (superseded by the implementation) + +Issue: https://github.com/AMD-AGI/Hyperloom/issues/1196 + +**This document is kept for history and is partly WRONG.** It was written +against a clone that was 402 commits stale, and upstream has since refactored +the integration points it names. Recorded here because the correction is the +useful part: + +| Scoped (stale) | Actual | +|---|---| +| Edit `_AMD_GPU_TYPES`, `_GFX_TO_RUNNER` and `_AMD_GPU_DISPATCH_IDENTITIES` separately in `gpu_types.py` | Identities moved to `hyperloom/common/gpu_identity.py`; `_AMD_GPU_TYPES` and the rocm-smi product tags are **derived** from it. Adding a board is one row + one `_GFX_TO_RUNNER` entry | +| Add R9700/RADEON product tags to `_autodetect_gpu_type` | Not needed — tags derive from the identities table, so autodetect works the moment the row exists | +| Deliver `custom_r9700.sh` as a core change | The runner is operator-supplied via `$HYPERLOOM_BYPASS_SCRIPTS_DIR`; no core change. Shipped as `examples/rdna/custom_r9700.sh` | +| Add an RDNA4 quant table to `quantization_schemes.py` | The existing MI355X gate already excludes mxfp4 everywhere else. What was missing was the *reason* (gfx950-gated silicon, not a missing kernel) and a test locking it | +| Flip `test_profile_and_kernel_handlers.py`'s gfx1100 expectation | Left alone; coverage added as `test_rdna4_r9700_support.py` instead | + +Net: the port is far smaller than scoped, because upstream had already done the +single-source-of-truth refactor this needed. See `README.md` in this directory +for the delivered state. diff --git a/examples/rdna/container/Containerfile b/examples/rdna/container/Containerfile new file mode 100644 index 0000000000..688df23835 --- /dev/null +++ b/examples/rdna/container/Containerfile @@ -0,0 +1,97 @@ +# SPDX-FileCopyrightText: 2026 The-Monk +# SPDX-License-Identifier: MIT +# +# Portability test for the RDNA port: a CLEAN ROCm image, STOCK upstream +# llama.cpp built from source, and this fork installed from git. +# +# The point is that none of it is ours. The port was first validated against a +# private llama.cpp fork and a pinned ROCm, which proves nothing about anybody +# else's machine. If it works here it works for a stranger. +# +# docker build -t hyperloom-rdna:test \ +# --build-arg AMDGPU_TARGETS=gfx1201 \ +# -f examples/rdna/container/Containerfile . +# +# Then just use the wrapper, which handles device passthrough and mounts: +# +# MODEL=/path/to/model.gguf examples/rdna/container/run.sh +# +# Or by hand. No --group-add: the container runs as root, which reaches +# /dev/kfd and /dev/dri via CAP_DAC_OVERRIDE. +# +# docker run --rm --device /dev/kfd --device /dev/dri \ +# --security-opt seccomp=unconfined \ +# -e MODEL=/models/your.gguf -v /path/to/models:/models:ro \ +# hyperloom-rdna:test +# +# Only when dropping to a non-root user (--user) do you need --group-add, and +# then it must be the host's NUMERIC gid: `--group-add render` resolves the name +# INSIDE the image, which has no render group, and fails with "unable to find +# group render". run.sh does this for you via RUN_AS_USER. +# +# Fully qualified: podman has no unqualified-search registry by default and a +# short name fails with "did not resolve to an alias". Docker assumes docker.io; +# podman does not. +ARG ROCM_IMAGE=docker.io/rocm/dev-ubuntu-24.04:latest +FROM ${ROCM_IMAGE} + +# gfx1201 = RDNA4 (R9700 / RX 9070). Override for your arch; the build is +# arch-specific, so a wrong value here produces a binary that cannot run. +ARG AMDGPU_TARGETS=gfx1201 +ARG LLAMA_CPP_REF=master +ARG FORK_REF=rdna + +ENV DEBIAN_FRONTEND=noninteractive +RUN apt-get update && apt-get install -y --no-install-recommends \ + git cmake ninja-build build-essential ccache \ + libcurl4-openssl-dev python3 python3-pip python3-venv \ + && rm -rf /var/lib/apt/lists/* + +# The ROCm dev image ships the libraries in /opt/rocm/lib but registers that +# path with neither ld.so.conf.d nor LD_LIBRARY_PATH, so `ldconfig -p` knows no +# hipblas at all. Compilation still succeeds (cmake passes -L explicitly), and +# then every ROCm-linked binary dies at LOAD time with +# libhipblas.so.3: cannot open shared object file +# which reads as "the GPU support is broken" rather than "the linker was never +# told where the libraries live". Found by running this container; fixed here so +# nobody else loses an afternoon to it. +RUN echo "/opt/rocm/lib" > /etc/ld.so.conf.d/rocm.conf \ + && echo "/opt/rocm/lib64" >> /etc/ld.so.conf.d/rocm.conf \ + && ldconfig \ + && ldconfig -p | grep -q libhipblas || { echo "ldconfig still cannot see hipblas" >&2; exit 1; } + +# --- STOCK upstream llama.cpp. Not our fork: if the runner only works against +# --- our kernels, that is a fork feature, not RDNA support. +WORKDIR /opt +RUN git clone --depth 1 --branch ${LLAMA_CPP_REF} https://github.com/ggml-org/llama.cpp.git +WORKDIR /opt/llama.cpp +RUN cmake -B build -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DGGML_HIP=ON \ + -DAMDGPU_TARGETS=${AMDGPU_TARGETS} \ + -DLLAMA_CURL=OFF \ + -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang \ + -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ \ + && cmake --build build --target llama-bench llama-perplexity test-backend-ops -j"$(nproc)" + +# --- the fork, from git. Installed the way a stranger would install it. +WORKDIR /opt +RUN git clone --depth 1 --branch ${FORK_REF} https://github.com/The-Monk/hyperloom-rdna.git +WORKDIR /opt/hyperloom-rdna +# The probe needs the package importable and PyYAML for the validator import +# chain. A venv keeps us off the image's system python (PEP 668). +RUN python3 -m venv /opt/venv \ + && /opt/venv/bin/pip install --no-cache-dir --upgrade pip \ + && /opt/venv/bin/pip install --no-cache-dir pyyaml + +ENV PY=/opt/venv/bin/python \ + LLAMA_BENCH=/opt/llama.cpp/build/bin/llama-bench \ + TEST_BACKEND_OPS=/opt/llama.cpp/build/bin/test-backend-ops \ + LLAMA_PERPLEXITY=/opt/llama.cpp/build/bin/llama-perplexity \ + HIP_VISIBLE_DEVICES=0 \ + PP=512 TG=64 REPS=2 + +# Default: the full support probe. Correctness gate included — it is the half +# that a fresh toolchain is most likely to break. +ENTRYPOINT ["/opt/hyperloom-rdna/examples/rdna/rdna-support-probe.sh"] +CMD ["--bench"] diff --git a/examples/rdna/container/README.md b/examples/rdna/container/README.md new file mode 100644 index 0000000000..52f83ce3d6 --- /dev/null +++ b/examples/rdna/container/README.md @@ -0,0 +1,126 @@ +# Container portability test + +Proves the port works on a stack that is **not ours**: a clean ROCm image, +**stock upstream llama.cpp** built from source inside the container, and this +fork installed from git. + +This matters because the port was first validated against a private llama.cpp +fork and a pinned ROCm userspace. That proves the code runs *here*. It says +nothing about anyone else's machine, which is the only question that counts for +someone trying to use their own card. + +## Build + +```bash +docker build -t hyperloom-rdna:test \ + --build-arg AMDGPU_TARGETS=gfx1201 \ + -f examples/rdna/container/Containerfile /tmp/empty-ctx +``` + +`AMDGPU_TARGETS` must match your silicon — the HIP build is arch-specific, and +a wrong value produces binaries that cannot run on your GPU. `mkdir /tmp/empty-ctx` +first: nothing is copied from the working tree, everything is cloned from git, +so the build context is deliberately empty. + +Pin the inputs when you want a reproducible run: +`--build-arg LLAMA_CPP_REF=`, `--build-arg FORK_REF=`, +`--build-arg ROCM_IMAGE=rocm/dev-ubuntu-24.04:`. + +## Run + +```bash +MODEL=/path/to/model.gguf examples/rdna/container/run.sh +``` + +`run.sh` resolves the model path, mounts its directory read-only, passes the +devices, forwards the knobs (`PP`, `TG`, `REPS`, `CORRECTNESS_OPS`, `PPL_FILE`, +`PPL_MAX`, `HIP_VISIBLE_DEVICES`), and fails immediately with a clear message if +`/dev/kfd` is missing rather than 90 seconds into a run. Arguments pass through +to the probe. It works with `DOCKER=podman` too. + +By hand, if you prefer: + +```bash +docker run --rm \ + --device /dev/kfd --device /dev/dri \ + --security-opt seccomp=unconfined \ + -v /path/to/models:/models:ro \ + -e MODEL=/models/your-model.gguf \ + hyperloom-rdna:test +``` + +**Device access differs by engine** — and getting it wrong looks like a broken +GPU rather than a permissions problem: + +- **Docker (rootful):** no `--group-add` needed. The container runs as root and + reaches `/dev/kfd` and `/dev/dri` through `CAP_DAC_OVERRIDE` — verified, + `rocminfo` sees both GPUs and `test-backend-ops` passes without it. +- **Podman rootless:** in-container uid 0 maps to *your* host uid, so `/dev/kfd` + appears as `nobody:nogroup` and `CAP_DAC_OVERRIDE` does **not** apply. Host + supplementary groups are dropped unless you pass **`--group-add keep-groups`**. + Measured: without it `rocminfo` reports **zero** GPUs and `test-backend-ops` + runs CPU-only and reports `1/1 backends passed` — a green result that proves + nothing. With it: `gfx1201`, `2/2 backends passed`. + +`run.sh` detects rootless podman and adds `keep-groups` for you. + +`--group-add render` is wrong on both engines: the *name* resolves inside the +image, which has no `render` group. A non-root user needs the host's **numeric** +GID (`RUN_AS_USER=1000:1000` in `run.sh`). + +Defaults to the full probe (`--bench`), so it runs the correctness gate and the +benchmark and prints a paste-ready report. Add `-e CORRECTNESS_OPS=ALL` for a +thorough gate, or `-e PPL_FILE=/models/corpus.txt -e PPL_MAX=` for a +model-level gate too. + +**On a multi-GPU host, pick your compute GPU deliberately** (`-e HIP_VISIBLE_DEVICES=0`, +the default here). If one card is driving a display, do not benchmark on it — +a GPU hang there takes the desktop with it. + +## Use a STOCK quant format + +The container builds **upstream** llama.cpp, which only knows upstream types. +A GGUF in a fork-specific quant (`Q2_0`, `IU4`, `F8E4M3`, and friends) will not +load, and the failure looks like a broken port rather than an unknown type. + +Use `Q4_K_M`, `Q8_0`, `Q4_0` or another upstream format here. That is the point +of the exercise: if the runner only works against a private fork's kernels, +that is a fork feature, not RDNA support. + +## What this test already caught + +Both of these were found by *running* it, and both would have looked like +"RDNA support is broken" to someone trying the port for the first time: + +- **`libhipblas.so.3: cannot open shared object file`.** The ROCm dev image + ships its libraries in `/opt/rocm/lib` but registers that path with neither + `ld.so.conf.d` nor `LD_LIBRARY_PATH` — `ldconfig -p` knows no hipblas at all. + Compilation still succeeds, because cmake passes `-L` explicitly; every + ROCm-linked binary then dies at **load** time. The Containerfile now writes + the `ld.so.conf.d` entry, runs `ldconfig`, and **verifies** the result in the + same layer, so a regression fails the build rather than the run. Verified + after the fix: 4 hipblas entries in `ldconfig -p`, and **0** unresolved + libraries in both `llama-bench` and `test-backend-ops`. +- **`unable to find group render`.** `--group-add` resolves the group name + inside the image, which has no `render` group. The first fix was to document + the numeric-GID workaround; the actual fix was to check whether the flag was + needed at all. It is not — the container runs as root and reaches the devices + via `CAP_DAC_OVERRIDE`, confirmed by `rocminfo` and a passing + `test-backend-ops` with no group flags. It is now absent from the documented + command, and `run.sh` adds it only under `RUN_AS_USER`. + +Note also that the probe behaved correctly during the broken run: it reported +`PARTIAL` with the linker error in the correctness row, rather than a false +`CONFIRMED`. The gate refusing to pass is the gate working. + +## Reading the result + +The probe's exit code is the verdict: `0` CONFIRMED, `1` PARTIAL, `2` UNSUPPORTED. +`CONFIRMED` requires **both** the correctness gate (`test-backend-ops` vs the +CPU reference) and a measurement Hyperloom's own validator accepts — see +[`AGENTS.md`](../../../AGENTS.md). + +A container result is worth more than a host result, because it rules out +everything your host happens to have lying around: your ROCm, your compiler, +your fork, your environment variables. If it passes here, the port is the +reason. diff --git a/examples/rdna/container/run.sh b/examples/rdna/container/run.sh new file mode 100755 index 0000000000..b61ec2edf9 --- /dev/null +++ b/examples/rdna/container/run.sh @@ -0,0 +1,74 @@ +#!/usr/bin/env bash +# SPDX-FileCopyrightText: 2026 The-Monk +# SPDX-License-Identifier: MIT +# +# Run the RDNA portability test container correctly, so nobody has to rediscover +# the device-passthrough incantation. +# +# ./run.sh # uses $MODEL from the environment +# MODEL=/path/model.gguf ./run.sh +# MODEL=... ./run.sh --bench # args pass through to the probe +# +# Knobs (all optional): IMAGE, DOCKER (docker|podman), HIP_VISIBLE_DEVICES, +# PP, TG, REPS, CORRECTNESS_OPS, PPL_FILE, PPL_MAX, RUN_AS_USER +set -euo pipefail + +IMAGE="${IMAGE:-hyperloom-rdna:test}" +DOCKER="${DOCKER:-docker}" +command -v "$DOCKER" >/dev/null 2>&1 || { echo "ERROR: $DOCKER not found. Set DOCKER=podman?" >&2; exit 3; } + +# Fail here, not 90 seconds into a run, if the GPU was never exposed to us. +[ -e /dev/kfd ] || { echo "ERROR: /dev/kfd missing — no ROCm-capable GPU visible on this host." >&2; exit 3; } +[ -d /dev/dri ] || { echo "ERROR: /dev/dri missing." >&2; exit 3; } + +: "${MODEL:?set MODEL=/path/to/model.gguf (a STOCK quant: Q4_K_M, Q8_0, Q4_0 ...)}" +[ -f "$MODEL" ] || { echo "ERROR: MODEL not found: $MODEL" >&2; exit 3; } +MODEL_ABS="$(readlink -f "$MODEL")" +MODEL_DIR="$(dirname "$MODEL_ABS")" + +ARGS=(run --rm + --device /dev/kfd --device /dev/dri + --security-opt seccomp=unconfined + -v "$MODEL_DIR:/models:ro" + -e "MODEL=/models/$(basename "$MODEL_ABS")") + +# Device access differs by engine, and getting it wrong looks like a broken GPU +# rather than a permissions problem: +# docker (rootful) root in-container reaches the devices via CAP_DAC_OVERRIDE; +# no --group-add needed. +# podman ROOTLESS in-container uid 0 maps to your host uid, so /dev/kfd is +# nobody:nogroup and CAP_DAC_OVERRIDE does NOT apply. Host +# supplementary groups are dropped unless --group-add +# keep-groups is passed. Without it rocminfo reports ZERO +# GPUs and test-backend-ops quietly runs CPU-only, which +# "passes" while proving nothing. +if [ "$(basename "$DOCKER")" = "podman" ] \ + && [ "$($DOCKER info --format '{{.Host.Security.Rootless}}' 2>/dev/null)" = "true" ]; then + ARGS+=(--group-add keep-groups) +fi + +# --group-add render is wrong on both engines: the NAME resolves inside the +# image, which has no render group. A non-root user needs the host NUMERIC gid. +if [ -n "${RUN_AS_USER:-}" ]; then + ARGS+=(--user "$RUN_AS_USER") + for g in render video; do + gid="$(getent group "$g" | cut -d: -f3)" + [ -n "$gid" ] && ARGS+=(--group-add "$gid") + done +fi + +# A perplexity corpus lives on the host too; mount it or the gate silently +# cannot run. +if [ -n "${PPL_FILE:-}" ]; then + [ -f "$PPL_FILE" ] || { echo "ERROR: PPL_FILE not found: $PPL_FILE" >&2; exit 3; } + PPL_ABS="$(readlink -f "$PPL_FILE")" + ARGS+=(-v "$(dirname "$PPL_ABS"):/ppl:ro" -e "PPL_FILE=/ppl/$(basename "$PPL_ABS")") + [ -n "${PPL_MAX:-}" ] && ARGS+=(-e "PPL_MAX=$PPL_MAX") +fi + +for v in HIP_VISIBLE_DEVICES PP TG REPS CORRECTNESS_OPS; do + [ -n "${!v:-}" ] && ARGS+=(-e "$v=${!v}") +done + +echo ">> $DOCKER ${ARGS[*]} $IMAGE $*" +exec "$DOCKER" "${ARGS[@]}" "$IMAGE" "$@" diff --git a/examples/rdna/custom_r9700.sh b/examples/rdna/custom_r9700.sh new file mode 100755 index 0000000000..53fda309d9 --- /dev/null +++ b/examples/rdna/custom_r9700.sh @@ -0,0 +1,167 @@ +#!/usr/bin/env bash +# SPDX-FileCopyrightText: 2026 The-Monk +# SPDX-License-Identifier: MIT +# +# Hyperloom scriptable benchmark runner for AMD RDNA4 (gfx1201) via llama.cpp. +# +# WHY llama.cpp AND NOT vLLM/SGLANG: below ~32GB the serving stacks' VRAM and +# concurrency assumptions do not transfer, and llama.cpp is what this class of +# card is actually served with. This is a server-less "scriptable" workload: +# no OpenAI endpoint, no HTTP client -- run a benchmark, write one JSON. +# +# CONTRACT (from bypass_scriptable.py / benchmark_result.py): +# in : $MODEL $RESULT_DIR $RESULT_FILENAME $RUNNER_TYPE (Hyperloom sets these) +# out: $RESULT_DIR/$RESULT_FILENAME.json, InferenceX-shaped. To be selectable, +# output_throughput MUST be > 0 and, if a quality_gate is present, it +# must pass. A gate that is absent is non-blocking; a gate that FAILS +# makes the run unselectable no matter how fast it was -- which is the +# point of shipping one. +# +# INSTALL: point Hyperloom at this file's directory -- +# export HYPERLOOM_BYPASS_SCRIPTS_DIR=/path/to/examples/rdna +# +# KNOBS (env): LLAMA_CPP_DIR|LLAMA_BENCH, PP, TG, REPS, NGL, EXTRA_BENCH_ARGS, +# PPL_FILE, PPL_MAX, LLAMA_PERPLEXITY +set -euo pipefail + +MODEL="${MODEL:?MODEL is required (path to a .gguf)}" +RESULT_DIR="${RESULT_DIR:-$PWD}" +RESULT_FILENAME="${RESULT_FILENAME:-inferencex_result}" +RUNNER_TYPE="${RUNNER_TYPE:-r9700}" +PP="${PP:-1024}" # prefill tokens -> ttft / prompt throughput +TG="${TG:-128}" # decode tokens -> output_throughput (the headline) +REPS="${REPS:-3}" # llama-bench repetitions; it reports avg + stddev +NGL="${NGL:-999}" # offload everything to the GPU by default + +# --- resolve llama-bench. FAIL LOUDLY: a missing binary must not read as a +# --- zero-throughput result, which is indistinguishable from "the GPU is slow". +LLAMA_BENCH="${LLAMA_BENCH:-}" +if [ -z "$LLAMA_BENCH" ]; then + for c in "${LLAMA_CPP_DIR:-}/build/bin/llama-bench" \ + "${LLAMA_CPP_DIR:-}/llama-bench" \ + "$(command -v llama-bench 2>/dev/null || true)"; do + [ -n "$c" ] && [ -x "$c" ] && { LLAMA_BENCH="$c"; break; } + done +fi +[ -n "$LLAMA_BENCH" ] && [ -x "$LLAMA_BENCH" ] || { + echo "ERROR: llama-bench not found. Set LLAMA_BENCH=/path/to/llama-bench" >&2 + echo " or LLAMA_CPP_DIR=/path/to/llama.cpp (expects build/bin/llama-bench)." >&2 + exit 3; } +[ -f "$MODEL" ] || { echo "ERROR: MODEL not found: $MODEL" >&2; exit 3; } + +mkdir -p "$RESULT_DIR" +OUT="$RESULT_DIR/$RESULT_FILENAME.json" +BENCH_JSON="$RESULT_DIR/llama_bench.json" + +# --- provenance. A number without the stack that produced it is not comparable +# --- across sessions, which is the whole reason Hyperloom keeps a KB. +# rocminfo is not always present (or not on PATH) even on a working ROCm box -- +# it ships in rocminfo/rocm-core, which a minimal runtime install omits. Fall +# back to amd-smi, then to llama.cpp's own device line, so provenance degrades +# to "unknown" only when nothing on the system can name the arch. +GFX="$(rocminfo 2>/dev/null | grep -oE 'gfx[0-9a-f]+' | head -1 || true)" +[ -z "$GFX" ] && GFX="$(amd-smi static -g 0 2>/dev/null | grep -oE 'gfx[0-9a-f]+' | head -1 || true)" +DRIVER="$(cat /sys/module/amdgpu/version 2>/dev/null || true)" +HIP_VER="$(hipconfig --version 2>/dev/null || true)" + +echo ">> llama-bench: $LLAMA_BENCH" +echo ">> model : $MODEL" +echo ">> config : pp=$PP tg=$TG reps=$REPS ngl=$NGL arch=${GFX:-unknown}" + +START=$(date +%s) +# shellcheck disable=SC2086 # EXTRA_BENCH_ARGS is a deliberate word-split knob +"$LLAMA_BENCH" -m "$MODEL" -p "$PP" -n "$TG" -ngl "$NGL" -r "$REPS" \ + ${EXTRA_BENCH_ARGS:-} -o json > "$BENCH_JSON" +END=$(date +%s) + +# --- optional correctness gate. Speed without a correctness gate is how a +# --- wrong-but-fast kernel wins a benchmark; PPL_MAX makes that unselectable. +PPL_JSON=null +if [ -n "${PPL_FILE:-}" ]; then + LLAMA_PPL="${LLAMA_PERPLEXITY:-$(dirname "$LLAMA_BENCH")/llama-perplexity}" + if [ -x "$LLAMA_PPL" ]; then + echo ">> perplexity gate: $PPL_FILE (max ${PPL_MAX:-unset})" + PPL_LOG="$RESULT_DIR/perplexity.log" + "$LLAMA_PPL" -m "$MODEL" -f "$PPL_FILE" -ngl "$NGL" > "$PPL_LOG" 2>&1 || true + PPL_VAL="$(grep -oE 'Final estimate: PPL = [0-9.]+' "$PPL_LOG" | grep -oE '[0-9.]+$' | tail -1 || true)" + PPL_JSON="{\"metric\":\"perplexity\",\"value\":${PPL_VAL:-null},\"threshold\":${PPL_MAX:-null}}" + else + echo "WARN: PPL_FILE set but llama-perplexity not found; no gate emitted." >&2 + fi +fi + +RESULT_DIR="$RESULT_DIR" OUT="$OUT" BENCH_JSON="$BENCH_JSON" MODEL="$MODEL" \ +RUNNER_TYPE="$RUNNER_TYPE" PP="$PP" TG="$TG" REPS="$REPS" NGL="$NGL" \ +GFX="$GFX" DRIVER="$DRIVER" HIP_VER="$HIP_VER" PPL_JSON="$PPL_JSON" \ +DURATION="$((END - START))" python3 - <<'PY' +import json, os + +rows = json.load(open(os.environ["BENCH_JSON"])) +def pick(kind): + """avg_ts for the pp (n_prompt>0) or tg (n_gen>0) row.""" + for r in rows: + if kind == "pp" and int(r.get("n_prompt") or 0) > 0 and int(r.get("n_gen") or 0) == 0: + return r + if kind == "tg" and int(r.get("n_gen") or 0) > 0 and int(r.get("n_prompt") or 0) == 0: + return r + return None + +pp, tg = pick("pp"), pick("tg") +pp_ts = float(pp["avg_ts"]) if pp else None # prompt tokens/s +tg_ts = float(tg["avg_ts"]) if tg else None # decode tokens/s == the headline + +if not tg_ts or tg_ts <= 0: + # Refuse to emit a "valid" result with no decode throughput: Hyperloom would + # treat 0 as a legitimate slow variant rather than as a broken run. + raise SystemExit("FATAL: no decode (tg) throughput parsed from llama-bench output") + +n_pp, n_tg = int(os.environ["PP"]), int(os.environ["TG"]) +ppl = os.environ.get("PPL_JSON", "null") +gate = json.loads(ppl) if ppl and ppl != "null" else None +if gate is not None: + v, thr = gate.get("value"), gate.get("threshold") + gate["passed"] = bool(v is not None and thr is not None and float(v) <= float(thr)) + if thr is None: + # Measured but no threshold to judge against: record it, do not claim a verdict. + gate.pop("passed", None) + gate["skipped"] = True + gate["reason"] = "no PPL_MAX threshold supplied" + +report = { + "success": True, + "framework": "llamacpp", + "model": os.environ["MODEL"], + "workload_kind": "scriptable", + "throughput_unit": "tokens/s", + "throughput": { + "output_throughput": tg_ts, + "total_token_throughput": tg_ts, + "completed_requests": int(os.environ["REPS"]), + "duration_seconds": float(os.environ["DURATION"]), + }, + "latency": { + # llama-bench reports rates, not per-request latency. TPOT is the exact + # reciprocal of decode t/s; TTFT is DERIVED from prefill rate (prompt + # tokens / prompt t/s) and is therefore a model of first-token time, not + # a measurement of one. Labelled so nobody reads it as measured. + "tpot": {"mean_ms": 1000.0 / tg_ts}, + "ttft": {"mean_ms": (n_pp / pp_ts * 1000.0) if pp_ts else None}, + }, + "quality_gate": gate, + "rdna": { + "runner_type": os.environ["RUNNER_TYPE"], + "gfx_arch": os.environ.get("GFX") or None, + "amdgpu_driver": os.environ.get("DRIVER") or None, + "hip_version": os.environ.get("HIP_VER") or None, + "prompt_tokens": n_pp, + "gen_tokens": n_tg, + "n_gpu_layers": int(os.environ["NGL"]), + "prompt_throughput_tps": pp_ts, + "ttft_ms_is_derived": True, + "llama_bench_rows": rows, + }, +} +json.dump(report, open(os.environ["OUT"], "w"), indent=2) +print(f">> wrote {os.environ['OUT']}") +print(f">> prefill {pp_ts:.2f} t/s | decode {tg_ts:.2f} t/s" if pp_ts else f">> decode {tg_ts:.2f} t/s") +PY diff --git a/examples/rdna/rdna-support-probe.sh b/examples/rdna/rdna-support-probe.sh new file mode 100755 index 0000000000..e99083059c --- /dev/null +++ b/examples/rdna/rdna-support-probe.sh @@ -0,0 +1,224 @@ +#!/usr/bin/env bash +# SPDX-FileCopyrightText: 2026 The-Monk +# SPDX-License-Identifier: MIT +# +# RDNA support probe — produces the evidence needed to CONFIRM support on this +# machine, or to show exactly what is missing to ADD it. +# +# Prints a paste-ready markdown report and exits with the verdict: +# 0 = CONFIRMED correctness PASSED *and* a measurement passed Hyperloom's validator +# 1 = PARTIAL board resolves, but one of those two halves is missing +# 2 = UNSUPPORTED arch is not mapped (see AGENTS.md §2) +# +# ./rdna-support-probe.sh identity + resolution only +# ./rdna-support-probe.sh --bench correctness gate + benchmark + harness verdict +# +# CORRECTNESS IS NOT OPTIONAL. Throughput alone cannot confirm support: a kernel +# that is fast and WRONG will produce a beautiful number. --bench runs +# llama.cpp's test-backend-ops (every op compared against the CPU reference on +# this arch) and refuses CONFIRMED without it. +# +# Env for --bench: MODEL=/path/*.gguf plus LLAMA_BENCH= or LLAMA_CPP_DIR= +# Optional: PY= GPU= TEST_BACKEND_OPS= +# CORRECTNESS_OPS=MUL_MAT,MUL_MAT_ID (comma-separated, or ALL) +# PPL_FILE= PPL_MAX= (adds a model-level gate too) +set -uo pipefail + +BENCH=0; [ "${1:-}" = "--bench" ] && BENCH=1 +GPU="${GPU:-0}" +REPO="$(cd "$(dirname "${BASH_SOURCE[0]}")/../.." && pwd)" +u(){ echo "${1:-unknown}"; } + +# --- identity. amd-smi first (present on modern ROCm), then rocm-smi, then +# --- rocminfo; each can be absent on a perfectly working box. +SMI_STATIC="$(amd-smi static -g "$GPU" 2>/dev/null || true)" +PRODUCT="$(sed -n 's/.*MARKET_NAME:[[:space:]]*//p' <<<"$SMI_STATIC" | head -1)" +[ -z "$PRODUCT" ] && PRODUCT="$(rocm-smi --showproductname 2>/dev/null | sed -n 's/.*Card Series:[[:space:]]*//p' | head -1)" +CUS="$(sed -n 's/.*NUM_COMPUTE_UNITS:[[:space:]]*//p' <<<"$SMI_STATIC" | head -1)" +GFX="$(grep -oE 'gfx[0-9a-f]+' <<<"$SMI_STATIC" | head -1)" +[ -z "$GFX" ] && GFX="$(rocminfo 2>/dev/null | grep -oE 'gfx[0-9a-f]+' | head -1)" +# VRAM size is nested under a "VRAM:" block, not a flat VRAM_SIZE key. +VRAM="$(awk '/^[[:space:]]*VRAM:/{f=1;next} f&&/SIZE:/{print $2" "$3;exit}' <<<"$SMI_STATIC")" +# Count amdgpu cards ONLY. A server board's BMC (ASPEED etc.) and any +# simple-framebuffer appear under /sys/class/drm and would inflate this -- +# reporting a 2-GPU host as 3-GPU is exactly the kind of wrong detail that +# makes a support report untrustworthy. +NGPU="$(for d in /sys/class/drm/card*/device/driver; do + [ -e "$d" ] && basename "$(readlink -f "$d")"; done 2>/dev/null | grep -c '^amdgpu$')" +[ "${NGPU:-0}" = "0" ] && NGPU="$(amd-smi list 2>/dev/null | grep -c '^GPU:')" +[ "${NGPU:-0}" = "0" ] && NGPU="unknown" + +# --- stack +KERNEL="$(uname -r)" +DRIVER="$(cat /sys/module/amdgpu/version 2>/dev/null || true)" +ROCM_V="$(amd-smi version 2>/dev/null | sed -n 's/.*ROCm version:[[:space:]]*\([^ |]*\).*/\1/p' | head -1)" +[ -z "$ROCM_V" ] && ROCM_V="$(cat /opt/rocm/.info/version 2>/dev/null || hipconfig --version 2>/dev/null || true)" + +# --- pick a python that can import the package +PY="${PY:-}" +if [ -z "$PY" ]; then + for c in python3 "$REPO/.venv/bin/python" /usr/bin/python3; do + command -v "$c" >/dev/null 2>&1 || [ -x "$c" ] || continue + PYTHONPATH="$REPO/src" "$c" -c "import hyperloom.common.gpu_identity" >/dev/null 2>&1 \ + && { PY="$c"; break; } + done +fi + +RESOLVED="-"; IDENTITY="-"; RUNNER="-" +if [ -n "$PY" ]; then + read -r RESOLVED IDENTITY RUNNER <<<"$(GFX="$GFX" PYTHONPATH="$REPO/src" "$PY" - <<'PYRESOLVE' 2>/dev/null +import os +from hyperloom.inference_optimizer.gpu_types import ( + _GFX_TO_RUNNER, _resolve_amd_gpu_type, amd_gpu_dispatch_identity) +ident = amd_gpu_dispatch_identity() +print(_resolve_amd_gpu_type() or "-", + f"{ident[0]}/{ident[1]}CU" if ident else "-", + _GFX_TO_RUNNER.get(os.environ.get("GFX", ""), "-")) +PYRESOLVE +)" + : "${RESOLVED:=-}" "${IDENTITY:=-}" "${RUNNER:=-}" +else + echo "WARN: no python could import hyperloom from $REPO/src -- resolution unknown." >&2 + echo " Set PY=/path/to/python (the one with the repo deps installed)." >&2 +fi + +# --- CORRECTNESS FIRST. A wrong-but-fast kernel is worth zero, so this gates +# --- the verdict; the benchmark below is only meaningful once this passes. +CORRECT="not run"; CORRECT_OK=0 +if [ "$BENCH" = 1 ]; then + TBO="${TEST_BACKEND_OPS:-}" + if [ -z "$TBO" ]; then + for c in "$(dirname "${LLAMA_BENCH:-/nonexistent}")/test-backend-ops" \ + "${LLAMA_CPP_DIR:-}/build/bin/test-backend-ops" \ + "$(command -v test-backend-ops 2>/dev/null || true)"; do + [ -n "$c" ] && [ -x "$c" ] && { TBO="$c"; break; } + done + fi + if [ -z "$TBO" ] || [ ! -x "$TBO" ]; then + CORRECT="NOT RUN — test-backend-ops not found (build it, or set TEST_BACKEND_OPS)" + else + # ~90s on gfx1201 for the default pair. A one-time cost for a support + # claim; narrow it with CORRECTNESS_OPS while iterating. + OPS="${CORRECTNESS_OPS:-MUL_MAT,MUL_MAT_ID}" + CLOG="$(mktemp)"; CORRECT_OK=1; PASSED_TOTAL="" + if [ "$OPS" = "ALL" ]; then OPLIST=""; else OPLIST="$(tr ',' ' ' <<<"$OPS")"; fi + for op in ${OPLIST:-__all__}; do + if [ "$op" = "__all__" ]; then ARGS=(); else ARGS=(-o "$op"); fi + if HIP_VISIBLE_DEVICES="${HIP_VISIBLE_DEVICES:-$GPU}" "$TBO" "${ARGS[@]}" >"$CLOG" 2>&1; then + n="$(grep -oE '[0-9]+/[0-9]+ tests passed' "$CLOG" | tail -1)" + PASSED_TOTAL="$PASSED_TOTAL ${op}:${n:-ok}" + else + CORRECT_OK=0 + CORRECT="FAILED on ${op} — $(grep -iE 'FAIL|error' "$CLOG" | head -2 | tr '\n' ' ')" + break + fi + done + [ "$CORRECT_OK" = 1 ] && CORRECT="PASSED vs CPU reference (${OPS}):${PASSED_TOTAL}" + rm -f "$CLOG" + fi +fi + +# --- then measure, and let Hyperloom judge the measurement +BENCH_MD=""; VALID="not run"; PP_TS=""; TG_TS="" +if [ "$BENCH" = 1 ]; then + SCRIPT_PATH="$REPO/examples/rdna/custom_${RUNNER}.sh" + if [ "$RUNNER" = "-" ] || [ ! -x "$SCRIPT_PATH" ]; then + VALID="no runner script for this arch ($SCRIPT_PATH)" + elif [ -z "${MODEL:-}" ]; then + VALID="skipped: set MODEL=/path/to/model.gguf" + else + OUTDIR="$(mktemp -d)" + if RESULT_DIR="$OUTDIR" HIP_VISIBLE_DEVICES="${HIP_VISIBLE_DEVICES:-$GPU}" \ + "$SCRIPT_PATH" >"$OUTDIR/run.log" 2>&1; then + RES="$OUTDIR/inferencex_result.json" + TG_TS="$($PY -c "import json;print(json.load(open('$RES'))['throughput']['output_throughput'])" 2>/dev/null)" + PP_TS="$($PY -c "import json;print(json.load(open('$RES'))['rdna']['prompt_throughput_tps'])" 2>/dev/null)" + # THE VERDICT IS THE HARNESS'S, NOT OURS. + VALID="$(PYTHONPATH="$REPO/src" "$PY" - "$RES" <<'PY' 2>/dev/null || echo "validator unavailable" +import json,sys +from hyperloom.orchestrator.actions.executors.benchmark_result import ( + extract_benchmark_measurement, is_valid_measurement) +m = extract_benchmark_measurement(json.load(open(sys.argv[1]))) +print("ACCEPTED by is_valid_measurement()" if is_valid_measurement(m) else "REJECTED by is_valid_measurement()") +PY +)" + BENCH_MD="prefill ${PP_TS:-?} t/s / decode ${TG_TS:-?} t/s" + QGATE="$($PY -c " +import json;g=json.load(open('$RES')).get('quality_gate') +print('none supplied' if not g else ('PASSED' if g.get('passed') else ('SKIPPED: '+str(g.get('reason'))) if g.get('skipped') else 'FAILED')+' '+str({k:v for k,v in g.items() if k in ('metric','value','threshold')}))" 2>/dev/null)" + else + VALID="runner FAILED — see log tail below" + BENCH_MD="$(tail -5 "$OUTDIR/run.log" 2>/dev/null)" + fi + fi +fi + +# --- verdict +if [ "$RUNNER" = "-" ] || [ -z "$RESOLVED" ] || [ "$RESOLVED" = "-" ]; then + VERDICT="UNSUPPORTED"; RC=2 +elif [[ "$VALID" == ACCEPTED* ]] && [ "$CORRECT_OK" = 1 ]; then + VERDICT="CONFIRMED"; RC=0 +elif [[ "$VALID" == ACCEPTED* ]]; then + # Fast is not the same as right. Say so in the verdict, not in a footnote. + VERDICT="PARTIAL (throughput only — CORRECTNESS UNVERIFIED)"; RC=1 +elif [ "$CORRECT_OK" = 1 ]; then + VERDICT="PARTIAL (correctness passed, no accepted measurement)"; RC=1 +else + VERDICT="PARTIAL"; RC=1 +fi + +cat <": ("${GFX:-}", ${CUS:-})\` + — board key must be a substring of \`$(u "$PRODUCT")\` for autodetect to match +- \`src/hyperloom/inference_optimizer/gpu_types.py\`: \`_GFX_TO_RUNNER["${GFX:-}"] = ""\` +- a runner \`examples/rdna/custom_.sh\` (copy \`custom_r9700.sh\`) +- tests mirroring \`test_rdna4_r9700_support.py\`, negatives included +- a row in the \`docs/rdna-port/README.md\` matrix + +Map only the arch you actually ran. Do not add sibling chips you have not booted. +MD +elif [ "$RC" = 1 ]; then cat <