From 52f8ba964ec329cacd8c0666c2a9a6a64b4b28bb Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 16 Aug 2026 22:07:26 -0400 Subject: [PATCH 01/12] rdna-port: scoping doc (issue #1196) CCA (Claude Code Augmented) --- docs/rdna-port/SCOPE.md | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) create mode 100644 docs/rdna-port/SCOPE.md diff --git a/docs/rdna-port/SCOPE.md b/docs/rdna-port/SCOPE.md new file mode 100644 index 0000000000..01f43a314c --- /dev/null +++ b/docs/rdna-port/SCOPE.md @@ -0,0 +1,29 @@ +# RDNA support port — private scoping (pre-PR) + +Issue: https://github.com/AMD-AGI/Hyperloom/issues/1196 + +## Integration points found (all Hyperloom-side; Magpie untouched) + +1. `src/hyperloom/inference_optimizer/gpu_types.py` + - `_AMD_GPU_TYPES` += "r9700" (RDNA4 wave-1 target) + - `_GFX_TO_RUNNER` += {"gfx1201": "r9700", "gfx1200": "r9700"} + - `_AMD_GPU_DISPATCH_IDENTITIES` += {"r9700": ("gfx1201", 64)} + - `_autodetect_gpu_type`: add R9700/RADEON product tags + gfx fallthrough +2. Runner: `custom_{runner_type}.sh` convention in + `orchestrator/actions/executors/_workload_envs.py` — deliver + `custom_r9700.sh` wrapping llama-bench/llama-server (tg/pp/PPL gates, + interleaved A/B, anchor+config emission). No vLLM/SGLang dependency. +3. Quant schemes: `orchestrator/phases/quantization_schemes.py` — add RDNA4 + table (dp4a-universal set; WMMA int8/fp8; NO mxfp4 [gfx950-only], + NO MFMA; sparse24 available). +4. Profiling guards (TraceLens/Magpie consumers): document + guard + gfx1201 realities: host_trap-only PC sampling (min interval 512), + FETCH_SIZE/PMC-derived counters read zero, GL2C ratios-only under + profile_standard, setperflevel high = -10% throughput trap. +5. Tests: flip `test_profile_and_kernel_handlers.py` gfx1100 expectation + pattern into positive gfx1201 coverage. + +## Evidence base (ours, measured on gfx1201) +- instrument map + measurement protocol: roc10 ledger 2026-08-11..16 +- scheme/datapath grid: tools/bench-card isa-grid (17 arches x 32 instrs) +- llama.cpp harness: bench-card provenance + PPL gates + interleaved A/B From ac5ef2fc6e76b02817110072c315cde3cd81db55 Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 18:07:25 -0400 Subject: [PATCH 02/12] RDNA4 (gfx1201 / Radeon AI PRO R9700) support MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Upstream maps consumer arches to None, so a Radeon box resolves to no runner and no dispatch identity. Makes gfx1201 a first-class target. gpu_identity.py r9700 -> (gfx1201, 64). CU count read from the board (amd-smi NUM_COMPUTE_UNITS), not a spec sheet. Because _AMD_GPU_TYPES and the rocm-smi product tags derive from this table, one row also makes the board accepted by the CLI and detectable by autodetect. gpu_types.py gfx1201 -> r9700 runner. gfx1200 (Navi 44) deliberately left unmapped: same ISA family, never measured here, and a runner label is a claim of support. examples/rdna/ custom_r9700.sh — llama.cpp scriptable runner. No core change was needed: the bypass path already resolves an operator script by the custom_{runner_type}.sh convention. quantization Documents why mxfp4 is unavailable on RDNA4 — the MX scaled-converts are gfx950/gfx1250-gated silicon, so a capability probe grades them REJECTED. fp8 stays available. VERIFIED ON HARDWARE (2x Radeon AI PRO R9700, gfx1201), not asserted: - rocm-smi autodetect resolves r9700 -> ('gfx1201', 64) - the runner benchmarks end to end: prefill 4488.33 t/s, decode 165.81 t/s (Bonsai-8B Q2_0, pp512/tg64/r2) - its inferencex_result.json is accepted by Hyperloom's OWN validator: is_valid_measurement() -> True, scriptable detected. The harness judges the output, not the author of the script. - 756 upstream tests pass (gpu/quant/preflight/parser/provenance), 0 failures Two negatives are deliberate and tested: gfx1200 unmapped, mxfp4 refused. docs/rdna-port/SCOPE.md is corrected rather than deleted — it was written against a 402-commit-stale clone and scoped a much larger change than upstream now needs. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- docs/rdna-port/README.md | 95 ++++++++++ docs/rdna-port/SCOPE.md | 39 ++-- examples/rdna/custom_r9700.sh | 167 ++++++++++++++++++ src/hyperloom/common/gpu_identity.py | 7 + .../inference_optimizer/gpu_types.py | 5 + .../tests/test_rdna4_r9700_support.py | 73 ++++++++ .../phases/quantization_schemes.py | 9 + 7 files changed, 371 insertions(+), 24 deletions(-) create mode 100644 docs/rdna-port/README.md create mode 100755 examples/rdna/custom_r9700.sh create mode 100644 src/hyperloom/inference_optimizer/tests/test_rdna4_r9700_support.py diff --git a/docs/rdna-port/README.md b/docs/rdna-port/README.md new file mode 100644 index 0000000000..e97bbe062f --- /dev/null +++ b/docs/rdna-port/README.md @@ -0,0 +1,95 @@ +# 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.** Board autodetects; runner benchmarked end-to-end on 2× R9700 | +| 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 | +| gfx11 (RDNA3) | gfx1100/1101 | 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. + +## 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 index 01f43a314c..6b388028b7 100644 --- a/docs/rdna-port/SCOPE.md +++ b/docs/rdna-port/SCOPE.md @@ -1,29 +1,20 @@ -# RDNA support port — private scoping (pre-PR) +# RDNA support port — scoping (superseded by the implementation) Issue: https://github.com/AMD-AGI/Hyperloom/issues/1196 -## Integration points found (all Hyperloom-side; Magpie untouched) +**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: -1. `src/hyperloom/inference_optimizer/gpu_types.py` - - `_AMD_GPU_TYPES` += "r9700" (RDNA4 wave-1 target) - - `_GFX_TO_RUNNER` += {"gfx1201": "r9700", "gfx1200": "r9700"} - - `_AMD_GPU_DISPATCH_IDENTITIES` += {"r9700": ("gfx1201", 64)} - - `_autodetect_gpu_type`: add R9700/RADEON product tags + gfx fallthrough -2. Runner: `custom_{runner_type}.sh` convention in - `orchestrator/actions/executors/_workload_envs.py` — deliver - `custom_r9700.sh` wrapping llama-bench/llama-server (tg/pp/PPL gates, - interleaved A/B, anchor+config emission). No vLLM/SGLang dependency. -3. Quant schemes: `orchestrator/phases/quantization_schemes.py` — add RDNA4 - table (dp4a-universal set; WMMA int8/fp8; NO mxfp4 [gfx950-only], - NO MFMA; sparse24 available). -4. Profiling guards (TraceLens/Magpie consumers): document + guard - gfx1201 realities: host_trap-only PC sampling (min interval 512), - FETCH_SIZE/PMC-derived counters read zero, GL2C ratios-only under - profile_standard, setperflevel high = -10% throughput trap. -5. Tests: flip `test_profile_and_kernel_handlers.py` gfx1100 expectation - pattern into positive gfx1201 coverage. +| 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 | -## Evidence base (ours, measured on gfx1201) -- instrument map + measurement protocol: roc10 ledger 2026-08-11..16 -- scheme/datapath grid: tools/bench-card isa-grid (17 arches x 32 instrs) -- llama.cpp harness: bench-card provenance + PPL gates + interleaved A/B +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/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/src/hyperloom/common/gpu_identity.py b/src/hyperloom/common/gpu_identity.py index 669f89cb43..014c0d8d41 100644 --- a/src/hyperloom/common/gpu_identity.py +++ b/src/hyperloom/common/gpu_identity.py @@ -20,6 +20,13 @@ "mi308x": ("gfx942", 304), "mi325x": ("gfx942", 304), "mi355x": ("gfx950", 256), + # RDNA4 consumer/workstation. Unlike the Instinct rows this is a wave-32 + # arch with no MFMA, so anything that assumes CDNA matrix instructions must + # branch on the arch rather than on "is AMD". CU count read from the board + # (amd-smi NUM_COMPUTE_UNITS on an AMD Radeon AI PRO R9700), not a spec + # sheet. gfx1200 (Navi 44) is deliberately absent: same ISA family, but no + # one has measured it here, and a runner label is a claim of support. + "r9700": ("gfx1201", 64), } diff --git a/src/hyperloom/inference_optimizer/gpu_types.py b/src/hyperloom/inference_optimizer/gpu_types.py index fcbf98fec8..483c621ddb 100644 --- a/src/hyperloom/inference_optimizer/gpu_types.py +++ b/src/hyperloom/inference_optimizer/gpu_types.py @@ -26,6 +26,11 @@ # the arch a runner is reached by is a deliberate choice, not an inverse. "gfx942": "mi300x", "gfx950": "mi355x", + # RDNA4. Reached by the llama.cpp scriptable runner (custom_r9700.sh), not + # vLLM/SGLang: below ~32GB the serving stacks' VRAM and concurrency + # assumptions do not hold, and llama.cpp is what this class of card is + # actually served with. + "gfx1201": "r9700", } #: Re-exported from ``hyperloom.common`` so provenance and this module cannot diff --git a/src/hyperloom/inference_optimizer/tests/test_rdna4_r9700_support.py b/src/hyperloom/inference_optimizer/tests/test_rdna4_r9700_support.py new file mode 100644 index 0000000000..31f30e6255 --- /dev/null +++ b/src/hyperloom/inference_optimizer/tests/test_rdna4_r9700_support.py @@ -0,0 +1,73 @@ +# SPDX-FileCopyrightText: 2026 The-Monk +# SPDX-License-Identifier: MIT + +"""RDNA4 (gfx1201 / Radeon AI PRO R9700) support. + +Upstream maps consumer arches to ``None``, so a gfx1201 box resolves to no +runner and no dispatch identity. These lock the port's behaviour, including +the parts that must stay NEGATIVE: gfx1200 is deliberately unmapped (same ISA +family, never measured here) and mxfp4 must stay unavailable, because the MX +scaled-converts are gfx950/gfx1250-gated silicon rather than a missing kernel. +""" + +from __future__ import annotations + +import pytest + +from hyperloom.common.gpu_identity import AMD_GPU_DISPATCH_IDENTITIES, gfx_arch_for_gpu_type +from hyperloom.inference_optimizer.gpu_types import ( + _AMD_GPU_TYPES, + _GFX_TO_RUNNER, + _PRODUCT_TAGS, + _gpu_runner_type, + amd_gpu_dispatch_identity, +) +from hyperloom.orchestrator.phases.quantization_schemes import ( + SchemeNotSupportedError, + supported_schemes, + validate_scheme, +) + + +def test_r9700_has_a_dispatch_identity(): + # 64 CUs read from the board (amd-smi NUM_COMPUTE_UNITS), not a spec sheet. + assert AMD_GPU_DISPATCH_IDENTITIES["r9700"] == ("gfx1201", 64) + assert gfx_arch_for_gpu_type("r9700") == "gfx1201" + + +def test_r9700_is_an_accepted_board_and_resolves_explicitly(): + assert "r9700" in _AMD_GPU_TYPES + assert amd_gpu_dispatch_identity("r9700") == ("gfx1201", 64) + + +def test_gfx1201_selects_the_r9700_runner(): + assert _GFX_TO_RUNNER["gfx1201"] == "r9700" + # Not aliased onto an Instinct runner the way mi325x/mi308x fold to mi300x. + assert _gpu_runner_type("r9700") == "r9700" + + +def test_r9700_product_tag_is_derived_for_autodetect(): + # rocm-smi reports "AMD Radeon AI PRO R9700"; autodetect substring-matches + # the derived tag, so the board is detectable without a bespoke branch. + assert "R9700" in _PRODUCT_TAGS + + +def test_gfx1200_is_deliberately_unmapped(): + # Navi 44 is the same ISA family but has not been measured on this port. + # A runner label is a claim of support; absence here is the honest state. + assert "gfx1200" not in _GFX_TO_RUNNER + + +def test_mxfp4_is_unavailable_on_rdna4(): + schemes = supported_schemes("r9700") + assert "fp8" in schemes and "ptpc_fp8" in schemes + assert "mxfp4" not in schemes and "mxfp4_fp8" not in schemes + + +def test_mxfp4_request_on_rdna4_fails_loudly(): + with pytest.raises(SchemeNotSupportedError): + validate_scheme("mxfp4", "r9700") + + +def test_fp8_is_accepted_on_rdna4(): + validate_scheme("fp8", "r9700") # must not raise diff --git a/src/hyperloom/orchestrator/phases/quantization_schemes.py b/src/hyperloom/orchestrator/phases/quantization_schemes.py index ffe85040c6..81230ae9dd 100644 --- a/src/hyperloom/orchestrator/phases/quantization_schemes.py +++ b/src/hyperloom/orchestrator/phases/quantization_schemes.py @@ -14,6 +14,15 @@ ``mxfp4`` / ``mxfp4_fp8`` are MI355X-only. :func:`validate_scheme` enforces this so an mxfp4 request on an mi300x target fails loudly instead of producing an unservable artifact. + + On RDNA4 (``r9700``/gfx1201) the same exclusion holds, for a different + reason worth stating: the MX scaled-convert instructions are gfx950/gfx1250 + gated, so mxfp4 has no hardware datapath here at all -- a capability probe on + gfx1201 grades them REJECTED, not merely slow. ``fp8`` remains available + (RDNA4 has WMMA fp8 and a native E4M3 path); there is no MFMA on this arch, + so anything reaching for CDNA matrix cores must branch on the gfx arch rather + than on "is AMD". See ``docs/rdna-port/README.md`` for the measured + scheme-to-datapath table. """ from __future__ import annotations From 3bc9510deac93fa13ec229bb17c9229fa1588d5e Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 18:12:23 -0400 Subject: [PATCH 03/12] README: say what this fork is on the landing page The issue #1196 comment sends people here, and the landing page was still upstream's README with nothing identifying it as the RDNA port or disclaiming AMD endorsement. Banner only; upstream's text is untouched below it, and PR branches are cut from main, so this adds no diff noise upstream. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- README.md | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/README.md b/README.md index 616f79fe6b..5e51d505e2 100755 --- a/README.md +++ b/README.md @@ -1,3 +1,22 @@ +> ### 🔴 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). +> - Tracking issue upstream: [AMD-AGI/Hyperloom#1196](https://github.com/AMD-AGI/Hyperloom/issues/1196). +> - 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) From 5035f6512f17cb41cc8833ac08ee10e3eeacffbc Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 18:21:33 -0400 Subject: [PATCH 04/12] docs: record the RDNA3 evidence, scoped to what it actually shows MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Upstream #1041 (@kyuubyN) is the RDNA3 counterpart to #1196: Hyperloom's methodology independently re-implemented against their OWN HIP engine on gfx1102 (RX 7600) and gfx1100 (W7900), with measured results including a run whose honest verdict was "revert everything". Stated precisely, because the easy overstatement is wrong: that validates the APPROACH on gfx11. It does NOT establish Hyperloom running on RDNA3 — the runner map still excludes gfx11 and upstream's suite asserts the exclusion (_GFX_TO_RUNNER.get("gfx1100") is None). A result measured on A does not close B until B is measured; that rule is what the rest of this port is built on. Also carries over their engagement check: compare the actual rendered source each instance compiled rather than inferring engagement from log markers. It caught a knob that looked plausible but never reached the GPU — the same failure mode as optimizing a path that is not the live dispatched one. Invites RDNA3 support here and says what it needs. #1041 was auto-close-bot pressured and closed by its own author earlier today; the work behind it is real and should not be stranded. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- docs/rdna-port/README.md | 41 +++++++++++++++++++++++++++++++++++++++- 1 file changed, 40 insertions(+), 1 deletion(-) diff --git a/docs/rdna-port/README.md b/docs/rdna-port/README.md index e97bbe062f..7e4fd68983 100644 --- a/docs/rdna-port/README.md +++ b/docs/rdna-port/README.md @@ -16,7 +16,8 @@ branch makes gfx1201 a first-class target. | Radeon AI PRO R9700 | gfx1201 | **MEASURED.** Board autodetects; runner benchmarked end-to-end on 2× R9700 | | 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 | -| gfx11 (RDNA3) | gfx1100/1101 | Not in this branch yet | +| 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 @@ -88,6 +89,44 @@ 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. +## 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- From a25a9e05d970dc5675a19615b8f4d820d260d9d3 Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 18:21:49 -0400 Subject: [PATCH 05/12] README: cite the RDNA3 issue on the landing page too The banner named only #1196. Anyone arriving from the RDNA3 side should see their thread acknowledged and the invitation, without digging into docs/. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- README.md | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/README.md b/README.md index 5e51d505e2..1add4de91d 100755 --- a/README.md +++ b/README.md @@ -9,7 +9,9 @@ > 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). -> - Tracking issue upstream: [AMD-AGI/Hyperloom#1196](https://github.com/AMD-AGI/Hyperloom/issues/1196). +> - 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). > From 8db710f2639d8a48c55a5d42e2cc9683bf3ad7a9 Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 18:29:06 -0400 Subject: [PATCH 06/12] AGENTS.md + support probe + hardware-report template MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Agents arriving on someone else's RDNA box need three things: what to do, a way to produce evidence, and somewhere to put it. AGENTS.md THE ONE RULE: a claim is MEASURED or ASSUMED, and an ASSUMED claim never becomes a support-matrix entry. A support entry is a promise to a stranger whose hardware you cannot see. §1 confirm support on hardware you have; §2 add a target that is not mapped (one identities row + _GFX_TO_RUNNER + runner + tests incl. the negatives + a matrix row); §3 the playbook for when it does not work. §3 is the part that makes a remote agent useful rather than just loud: exhaust the mundane before the exotic (perf-level gating -> env traps -> vendor tracker -> versions -> build flags -> only then kernel internals); isolate bottom-up by layer with that layer's own ground-truth tool, because cross-layer misattribution is the biggest time-waster there is; a symptom table for the failures this port actually produces; and the rules for claiming a fix (it ran, a correctness gate passed, the path you changed is the one that executed, speedup = baseline / candidate, suspect the baseline when a win looks huge, and a result above the roofline means you measured cache). examples/rdna/rdna-support-probe.sh Turns all of that into one command. Prints a paste-ready report and exits CONFIRMED (0) / PARTIAL (1) / UNSUPPORTED (2) — and on UNSUPPORTED it emits the exact rows needed to map that arch. The CONFIRMED verdict requires Hyperloom's own is_valid_measurement() to accept the run: the harness grades the output, never the agent that produced it. Dogfooded here, and three bugs it exposed are fixed: VRAM lives in a nested block (not a flat key), ROCm version comes from amd-smi, and the GPU count now counts amdgpu cards ONLY -- the first version counted this board's ASPEED BMC and reported a 2-GPU host as 3-GPU, which is exactly the sort of wrong detail that makes a whole report untrustworthy. Verified end to end on 2x R9700: CONFIRMED, exit 0, prefill 4448.74 t/s / decode 168.11 t/s -- within ~1% of the independent earlier run. .github/ISSUE_TEMPLATE/rdna-hardware-report.yml Requires the verbatim probe block AND a "what you did NOT test" field, since that is what keeps the matrix honest. Says plainly that negative results are welcome: nobody without your hardware can produce them. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- .../ISSUE_TEMPLATE/rdna-hardware-report.yml | 71 ++++++ AGENTS.md | 232 ++++++++++++++++++ README.md | 4 +- docs/rdna-port/README.md | 13 + examples/rdna/rdna-support-probe.sh | 164 +++++++++++++ 5 files changed, 483 insertions(+), 1 deletion(-) create mode 100644 .github/ISSUE_TEMPLATE/rdna-hardware-report.yml create mode 100644 AGENTS.md create mode 100755 examples/rdna/rdna-support-probe.sh diff --git a/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml b/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml new file mode 100644 index 0000000000..92a6532168 --- /dev/null +++ b/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml @@ -0,0 +1,71 @@ +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. + + 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: 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..7c40c5537a --- /dev/null +++ b/AGENTS.md @@ -0,0 +1,232 @@ +# 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 a benchmark you can only confirm that the board is +*recognised*, which is the weaker half of the claim. + +The probe prints a paste-ready markdown block and one of three verdicts: + +| Verdict | Means | +|---|---| +| `CONFIRMED` | Board resolves **and** the runner produced a measurement Hyperloom's own validator accepted | +| `PARTIAL` | Board resolves, but no accepted measurement (no model, no llama.cpp, or the run failed) | +| `UNSUPPORTED` | The arch is not in `_GFX_TO_RUNNER` — go to §2 | + +### 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" +./examples/rdna/rdna-support-probe.sh --bench +``` + +Both must pass, and the probe must say `CONFIRMED`, before the matrix changes. + +--- + +## §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 1add4de91d..110d255084 100755 --- a/README.md +++ b/README.md @@ -8,7 +8,9 @@ > - **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). +> - 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.** diff --git a/docs/rdna-port/README.md b/docs/rdna-port/README.md index 7e4fd68983..615a996020 100644 --- a/docs/rdna-port/README.md +++ b/docs/rdna-port/README.md @@ -89,6 +89,19 @@ 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). 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 diff --git a/examples/rdna/rdna-support-probe.sh b/examples/rdna/rdna-support-probe.sh new file mode 100755 index 0000000000..76e3cc3fa1 --- /dev/null +++ b/examples/rdna/rdna-support-probe.sh @@ -0,0 +1,164 @@ +#!/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 board resolves AND a measurement passed Hyperloom's validator +# 1 = PARTIAL board resolves, no harness-accepted measurement +# 2 = UNSUPPORTED arch is not mapped (see AGENTS.md §2) +# +# ./rdna-support-probe.sh identity + resolution only +# ./rdna-support-probe.sh --bench also run the runner and have Hyperloom judge it +# +# Env for --bench: MODEL=/path/*.gguf plus LLAMA_BENCH= or LLAMA_CPP_DIR= +# Optional: PY= GPU= +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 + +# --- optional: run the real thing and let Hyperloom judge it +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" + 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* ]]; then + VERDICT="CONFIRMED"; RC=0 +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 < Date: Sun, 23 Aug 2026 18:55:06 -0400 Subject: [PATCH 07/12] =?UTF-8?q?Correctness=20gates=20the=20verdict=20?= =?UTF-8?q?=E2=80=94=20throughput=20alone=20could=20earn=20CONFIRMED?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Reported by JM: correctness was never mentioned. It was worse than an omission. AGENTS.md preached "a wrong-but-fast kernel is worth zero" in §3.4 while the probe stamped CONFIRMED on throughput alone, and the run I reported as our own confirmation had quality_gate: None — no correctness evidence of any kind. Hyperloom's is_valid_measurement() treats a MISSING quality_gate as non-blocking. That is right for the harness (a scriptable run need not carry one) and the wrong bar for a support claim, so the probe now adds its own gate instead of leaning on the harness for something the harness never promised. probe --bench now runs llama.cpp test-backend-ops FIRST, comparing ops against the CPU reference on this arch, and CONFIRMED requires it to pass AND a harness-accepted measurement. A fast run with no correctness evidence now prints, in the verdict line itself, "PARTIAL (throughput only — CORRECTNESS UNVERIFIED)" rather than passing quietly. Default CORRECTNESS_OPS=MUL_MAT,MUL_MAT_ID (the matmul paths a quant or arch port actually breaks); CORRECTNESS_OPS=ALL widens it. The model-level PPL gate is surfaced as its own row. AGENTS New §1.1a: correctness is half the verdict, not a footnote. Says explicitly what is_valid_measurement() does not do, so nobody infers a correctness claim from an ACCEPTED measurement. §2.3: if correctness FAILS on a new arch, report it and STOP — do not tune a kernel computing the wrong answer, and do not add a matrix row. docs Matrix cites the correctness gate for gfx1201, not just the benchmark. template Correctness is now a REQUIRED field, and a correctness failure is declared a first-class report. Op-level and model-level gates fail differently and both are worth having: ops prove the kernels compute the right thing, perplexity proves the assembled model still produces sane output. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- .../ISSUE_TEMPLATE/rdna-hardware-report.yml | 18 +++++ AGENTS.md | 44 +++++++++-- docs/rdna-port/README.md | 8 +- examples/rdna/rdna-support-probe.sh | 76 ++++++++++++++++--- 4 files changed, 129 insertions(+), 17 deletions(-) diff --git a/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml b/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml index 92a6532168..94d45396a9 100644 --- a/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml +++ b/.github/ISSUE_TEMPLATE/rdna-hardware-report.yml @@ -7,6 +7,10 @@ body: 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 @@ -19,6 +23,20 @@ body: **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: diff --git a/AGENTS.md b/AGENTS.md index 7c40c5537a..166a05d633 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -57,17 +57,42 @@ Corollaries that have already bitten this project: ``` Use `--bench` whenever a model is available (`MODEL=/path/to/*.gguf` and a -llama.cpp build). Without a benchmark you can only confirm that the board is -*recognised*, which is the weaker half of the claim. +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` | Board resolves **and** the runner produced a measurement Hyperloom's own validator accepted | -| `PARTIAL` | Board resolves, but no accepted measurement (no model, no llama.cpp, or the run failed) | +| `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. Widen with `CORRECTNESS_OPS=ALL` (slower, more thorough) or +narrow it if 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 @@ -130,10 +155,17 @@ generosity — it is a support claim written by someone who cannot honour it. PY= PYTHONPATH=src $PY -m pytest src/hyperloom/inference_optimizer/tests/ -q \ -k "gpu or quant or preflight or parser or provenance" -./examples/rdna/rdna-support-probe.sh --bench +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`, before the matrix changes. +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. --- diff --git a/docs/rdna-port/README.md b/docs/rdna-port/README.md index 615a996020..5c97ec0bea 100644 --- a/docs/rdna-port/README.md +++ b/docs/rdna-port/README.md @@ -13,7 +13,7 @@ branch makes gfx1201 a first-class target. | Target | Arch | Status | |---|---|---| -| Radeon AI PRO R9700 | gfx1201 | **MEASURED.** Board autodetects; runner benchmarked end-to-end on 2× R9700 | +| Radeon AI PRO R9700 | gfx1201 | **MEASURED.** Board autodetects; `test-backend-ops` passes vs the CPU reference; runner benchmarked end-to-end on 2× R9700 | | 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 | @@ -98,7 +98,11 @@ Run the probe and open an **RDNA hardware report** issue with its output: ``` It prints a paste-ready block and a verdict — `CONFIRMED`, `PARTIAL`, or -`UNSUPPORTED` (with the exact change needed to map your arch). Agents working +`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. diff --git a/examples/rdna/rdna-support-probe.sh b/examples/rdna/rdna-support-probe.sh index 76e3cc3fa1..2712f92030 100755 --- a/examples/rdna/rdna-support-probe.sh +++ b/examples/rdna/rdna-support-probe.sh @@ -6,15 +6,22 @@ # machine, or to show exactly what is missing to ADD it. # # Prints a paste-ready markdown report and exits with the verdict: -# 0 = CONFIRMED board resolves AND a measurement passed Hyperloom's validator -# 1 = PARTIAL board resolves, no harness-accepted measurement +# 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 also run the runner and have Hyperloom judge it +# ./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= +# 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 @@ -75,7 +82,41 @@ else echo " Set PY=/path/to/python (the one with the repo deps installed)." >&2 fi -# --- optional: run the real thing and let Hyperloom judge it +# --- 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 + OPS="${CORRECTNESS_OPS:-MUL_MAT,MUL_MAT_ID}" + CLOG="$(mktemp)"; CORRECT_OK=1; PASSED_TOTAL=0 + 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" @@ -100,6 +141,9 @@ print("ACCEPTED by is_valid_measurement()" if is_valid_measurement(m) else "REJE 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)" @@ -110,8 +154,13 @@ fi # --- verdict if [ "$RUNNER" = "-" ] || [ -z "$RESOLVED" ] || [ "$RESOLVED" = "-" ]; then VERDICT="UNSUPPORTED"; RC=2 -elif [[ "$VALID" == ACCEPTED* ]]; then +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 @@ -134,8 +183,10 @@ cat < Date: Sun, 23 Aug 2026 18:59:24 -0400 Subject: [PATCH 08/12] probe: fix stray accumulator in the correctness row; record measured cost MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The row read "(MUL_MAT,MUL_MAT_ID):0 MUL_MAT:1196/1196 ..." because the accumulator was initialised to 0 rather than empty. Cosmetic, but this string is the artifact strangers paste into issues, so it should not carry a phantom count next to real ones. Default gate measured at ~90s on gfx1201 (MUL_MAT 1196/1196, MUL_MAT_ID 865/865 vs the CPU reference) — cheap enough to be non-negotiable. Documented in AGENTS.md and carried into the support matrix alongside the throughput, so the matrix states both halves of the claim rather than speed alone. Our own CONFIRMED re-earned under the corrected gate: correctness PASSED + prefill 4476.34 t/s / decode 169.14 t/s ACCEPTED by is_valid_measurement(). Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- AGENTS.md | 6 ++++-- docs/rdna-port/README.md | 2 +- examples/rdna/rdna-support-probe.sh | 4 +++- 3 files changed, 8 insertions(+), 4 deletions(-) diff --git a/AGENTS.md b/AGENTS.md index 166a05d633..0c89d599aa 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -80,8 +80,10 @@ 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. Widen with `CORRECTNESS_OPS=ALL` (slower, more thorough) or -narrow it if you are iterating. +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 diff --git a/docs/rdna-port/README.md b/docs/rdna-port/README.md index 5c97ec0bea..9918b7c1f3 100644 --- a/docs/rdna-port/README.md +++ b/docs/rdna-port/README.md @@ -13,7 +13,7 @@ branch makes gfx1201 a first-class target. | Target | Arch | Status | |---|---|---| -| Radeon AI PRO R9700 | gfx1201 | **MEASURED.** Board autodetects; `test-backend-ops` passes vs the CPU reference; runner benchmarked end-to-end on 2× R9700 | +| Radeon AI PRO R9700 | gfx1201 | **MEASURED.** 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 | diff --git a/examples/rdna/rdna-support-probe.sh b/examples/rdna/rdna-support-probe.sh index 2712f92030..e99083059c 100755 --- a/examples/rdna/rdna-support-probe.sh +++ b/examples/rdna/rdna-support-probe.sh @@ -97,8 +97,10 @@ if [ "$BENCH" = 1 ]; then 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=0 + 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 From c92d454def06028234210e6b92c15b5c92e84c8f Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 19:22:04 -0400 Subject: [PATCH 09/12] =?UTF-8?q?container:=20portability=20test=20against?= =?UTF-8?q?=20STOCK=20upstream=20llama.cpp=20=E2=80=94=20CONFIRMED?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The port was validated only against our own llama.cpp fork and a pinned ROCm. That proves it runs here and says nothing about anyone else's machine, which is the only question that matters for someone with their own card. examples/rdna/container/ builds a clean rocm/dev-ubuntu-24.04 image, compiles STOCK upstream llama.cpp from source for gfx1201, clones this fork from GitHub, and runs the probe. Nothing of ours is in the stack: if the runner only worked against our kernels, that would be a fork feature, not RDNA support. RESULT — CONFIRMED, exit 0, in-container, ROCm 7.14.60850, upstream llama.cpp: correctness MUL_MAT 1193/1193, MUL_MAT_ID 869/869 vs the CPU reference benchmark prefill 3883.66 t/s / decode 93.97 t/s (Qwen3-8B-Q4_K_M) harness ACCEPTED by is_valid_measurement() Not comparable to the host figures: different model AND different quant (Bonsai-8B Q2_0 there, Qwen3-8B Q4_K_M here, since Q2_0 is a fork-specific type upstream cannot load). The claim is portability, not a speed comparison. TWO REAL BUGS, both found by running it, both of which would have read as "RDNA support is broken" to a newcomer: libhipblas.so.3 not found — the ROCm dev image ships 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. Compilation succeeds (cmake passes -L), then every ROCm-linked binary dies at LOAD time. Fixed in the image, and the fix verifies itself in the same layer so a regression fails the build instead of the run. unable to find group render — --group-add resolves the group name INSIDE the image, which has none. The host's numeric GID is what grants /dev/kfd and /dev/dri. Corrected in both the Containerfile header and the README. Worth recording: during the broken run the probe reported PARTIAL with the linker error in the correctness row rather than a false CONFIRMED. The gate refusing to pass is the gate working. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- docs/rdna-port/README.md | 2 +- examples/rdna/container/Containerfile | 88 +++++++++++++++++++++++++ examples/rdna/container/README.md | 94 +++++++++++++++++++++++++++ 3 files changed, 183 insertions(+), 1 deletion(-) create mode 100644 examples/rdna/container/Containerfile create mode 100644 examples/rdna/container/README.md diff --git a/docs/rdna-port/README.md b/docs/rdna-port/README.md index 9918b7c1f3..66096c8abd 100644 --- a/docs/rdna-port/README.md +++ b/docs/rdna-port/README.md @@ -13,7 +13,7 @@ branch makes gfx1201 a first-class target. | Target | Arch | Status | |---|---|---| -| Radeon AI PRO R9700 | gfx1201 | **MEASURED.** 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) | +| 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 | diff --git a/examples/rdna/container/Containerfile b/examples/rdna/container/Containerfile new file mode 100644 index 0000000000..b454cbeed7 --- /dev/null +++ b/examples/rdna/container/Containerfile @@ -0,0 +1,88 @@ +# 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 . +# +# docker run --rm --device /dev/kfd --device /dev/dri \ +# --group-add "$(getent group render | cut -d: -f3)" \ +# --group-add "$(getent group video | cut -d: -f3)" \ +# --security-opt seccomp=unconfined \ +# -e MODEL=/models/your.gguf -v /path/to/models:/models:ro \ +# hyperloom-rdna:test +# +# NOTE the numeric GIDs. `--group-add render` resolves the name INSIDE the +# image, which has no render group, and fails with "unable to find group +# render". The host's GID is what grants access to /dev/kfd and /dev/dri. +# +ARG ROCM_IMAGE=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..d9dc01d0ed --- /dev/null +++ b/examples/rdna/container/README.md @@ -0,0 +1,94 @@ +# 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 +docker run --rm \ + --device /dev/kfd --device /dev/dri \ + --group-add "$(getent group render | cut -d: -f3)" \ + --group-add "$(getent group video | cut -d: -f3)" \ + --security-opt seccomp=unconfined \ + -v /path/to/models:/models:ro \ + -e MODEL=/models/your-model.gguf \ + hyperloom-rdna:test +``` + +Use the **numeric** GIDs. `--group-add render` resolves the group name inside +the image, which has no `render` group, and fails with `unable to find group +render`; the host's GID is what actually grants access to `/dev/kfd` and +`/dev/dri`. + +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. +- **`unable to find group render`** — the `--group-add` name-vs-GID trap above. + +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. From 0d54ae60e756d14d9dbffa77ac6b5ea620cf5d57 Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 19:33:29 -0400 Subject: [PATCH 10/12] container: actually fix both bugs, rather than documenting one of them MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Bug 1 (libhipblas.so.3 not found) was already fixed in the image. Verified it is a real fix and not a masked one: ldconfig -p now reports 4 hipblas entries, and llama-bench and test-backend-ops each have 0 unresolved libraries. Bug 2 (unable to find group render) had only been documented — I wrote the numeric-GID workaround into the README and moved on without asking whether the flag was needed at all. It is not. The container runs as root, which reaches /dev/kfd and /dev/dri through CAP_DAC_OVERRIDE; verified with no group flags at all: rocminfo sees both gfx1201 GPUs and test-backend-ops passes. --group-add is now gone from the documented command and is added only when dropping to a non-root user, where the host's numeric GID is genuinely required. Adds run.sh so the incantation does not have to be rediscovered: resolves the model path and mounts its directory read-only, passes the devices, forwards the knobs, supports DOCKER=podman and RUN_AS_USER, and fails immediately with a clear message when /dev/kfd is absent instead of 90 seconds into a run. Verified end to end through the wrapper: CONFIRMED, exit 0, correctness MUL_MAT_ID 869/869 vs the CPU reference, on stock upstream llama.cpp. A workaround in a README is not a fix; it is a note asking the next person to pay the same cost. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- examples/rdna/container/Containerfile | 16 ++++--- examples/rdna/container/README.md | 36 ++++++++++++---- examples/rdna/container/run.sh | 62 +++++++++++++++++++++++++++ 3 files changed, 101 insertions(+), 13 deletions(-) create mode 100755 examples/rdna/container/run.sh diff --git a/examples/rdna/container/Containerfile b/examples/rdna/container/Containerfile index b454cbeed7..fc0984bd55 100644 --- a/examples/rdna/container/Containerfile +++ b/examples/rdna/container/Containerfile @@ -12,16 +12,22 @@ # --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 \ -# --group-add "$(getent group render | cut -d: -f3)" \ -# --group-add "$(getent group video | cut -d: -f3)" \ # --security-opt seccomp=unconfined \ # -e MODEL=/models/your.gguf -v /path/to/models:/models:ro \ # hyperloom-rdna:test # -# NOTE the numeric GIDs. `--group-add render` resolves the name INSIDE the -# image, which has no render group, and fails with "unable to find group -# render". The host's GID is what grants access to /dev/kfd and /dev/dri. +# 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. # ARG ROCM_IMAGE=rocm/dev-ubuntu-24.04:latest FROM ${ROCM_IMAGE} diff --git a/examples/rdna/container/README.md b/examples/rdna/container/README.md index d9dc01d0ed..18233c6527 100644 --- a/examples/rdna/container/README.md +++ b/examples/rdna/container/README.md @@ -28,21 +28,33 @@ Pin the inputs when you want a reproducible run: ## 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 \ - --group-add "$(getent group render | cut -d: -f3)" \ - --group-add "$(getent group video | cut -d: -f3)" \ --security-opt seccomp=unconfined \ -v /path/to/models:/models:ro \ -e MODEL=/models/your-model.gguf \ hyperloom-rdna:test ``` -Use the **numeric** GIDs. `--group-add render` resolves the group name inside -the image, which has no `render` group, and fails with `unable to find group -render`; the host's GID is what actually grants access to `/dev/kfd` and -`/dev/dri`. +**No `--group-add` is needed.** The container runs as root, which reaches +`/dev/kfd` and `/dev/dri` through `CAP_DAC_OVERRIDE` — verified: `rocminfo` +sees both GPUs and `test-backend-ops` passes without it. It is required *only* +when dropping to a non-root user, and then it must be the host's **numeric** +GID, because `--group-add render` resolves the group name inside the image, +which has no `render` group. `run.sh` handles that via `RUN_AS_USER=1000:1000`. 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 @@ -74,8 +86,16 @@ Both of these were found by *running* it, and both would have looked like 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. -- **`unable to find group render`** — the `--group-add` name-vs-GID trap above. + 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 diff --git a/examples/rdna/container/run.sh b/examples/rdna/container/run.sh new file mode 100755 index 0000000000..9487ee56af --- /dev/null +++ b/examples/rdna/container/run.sh @@ -0,0 +1,62 @@ +#!/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")") + +# The container runs as root by default, and root reaches /dev/kfd and /dev/dri +# via CAP_DAC_OVERRIDE -- no --group-add needed. It is needed ONLY when dropping +# to a non-root user, 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". +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" "$@" From d652e2e7bdd81e0c2aae9c1ae0e6a0421fc12666 Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 20:24:53 -0400 Subject: [PATCH 11/12] container: fully-qualify the base image so podman works MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The README claimed DOCKER=podman works. It did not: podman has no unqualified-search registry by default, so the short name rocm/dev-ubuntu-24.04 fails with "did not resolve to an alias and no unqualified-search registries are defined". Docker silently assumes docker.io; podman does not. Found while porting this recipe into the Rocky toolkit, where podman is the auto-detected engine — i.e. the claim was falsified by the first person to use the other engine, which is the whole reason to state such claims testably. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- examples/rdna/container/Containerfile | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/examples/rdna/container/Containerfile b/examples/rdna/container/Containerfile index fc0984bd55..688df23835 100644 --- a/examples/rdna/container/Containerfile +++ b/examples/rdna/container/Containerfile @@ -29,7 +29,10 @@ # 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. # -ARG ROCM_IMAGE=rocm/dev-ubuntu-24.04:latest +# 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 From 7340d1de732504d4e77d6ebf9eaca53c98f8792f Mon Sep 17 00:00:00 2001 From: jmonk Date: Sun, 23 Aug 2026 20:33:35 -0400 Subject: [PATCH 12/12] =?UTF-8?q?container:=20rootless=20podman=20needs=20?= =?UTF-8?q?keep-groups=20=E2=80=94=20my=20"no=20--group-add"=20claim=20was?= =?UTF-8?q?=20engine-specific?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit I verified that root in the container reaches /dev/kfd via CAP_DAC_OVERRIDE and wrote it up as a general fact. It is true for DOCKER (rootful) and false for ROOTLESS PODMAN, which is what podman defaults to here. Under rootless podman, in-container uid 0 maps to the invoking host uid, so /dev/kfd shows as nobody:nogroup and CAP_DAC_OVERRIDE does not apply. Host supplementary groups are dropped without --group-add keep-groups. The failure mode is the dangerous kind, and I hit it: rocminfo reported ZERO GPUs, test-backend-ops silently fell back to CPU-only and printed "1/1 backends passed" (vs 2/2 with a GPU), and the run exited 0. A green verify with no GPU attached is worse than a red one. run.sh now detects rootless podman and adds keep-groups; the README documents both engines with the measured evidence for each. Co-Authored-By: Claude Opus 5 Claude-Session: https://claude.ai/code/session_013oPUZALTu4ZPwMdHDdwiQt --- examples/rdna/container/README.md | 24 ++++++++++++++++++------ examples/rdna/container/run.sh | 22 +++++++++++++++++----- 2 files changed, 35 insertions(+), 11 deletions(-) diff --git a/examples/rdna/container/README.md b/examples/rdna/container/README.md index 18233c6527..52f83ce3d6 100644 --- a/examples/rdna/container/README.md +++ b/examples/rdna/container/README.md @@ -49,12 +49,24 @@ docker run --rm \ hyperloom-rdna:test ``` -**No `--group-add` is needed.** The container runs as root, which reaches -`/dev/kfd` and `/dev/dri` through `CAP_DAC_OVERRIDE` — verified: `rocminfo` -sees both GPUs and `test-backend-ops` passes without it. It is required *only* -when dropping to a non-root user, and then it must be the host's **numeric** -GID, because `--group-add render` resolves the group name inside the image, -which has no `render` group. `run.sh` handles that via `RUN_AS_USER=1000:1000`. +**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 diff --git a/examples/rdna/container/run.sh b/examples/rdna/container/run.sh index 9487ee56af..b61ec2edf9 100755 --- a/examples/rdna/container/run.sh +++ b/examples/rdna/container/run.sh @@ -32,11 +32,23 @@ ARGS=(run --rm -v "$MODEL_DIR:/models:ro" -e "MODEL=/models/$(basename "$MODEL_ABS")") -# The container runs as root by default, and root reaches /dev/kfd and /dev/dri -# via CAP_DAC_OVERRIDE -- no --group-add needed. It is needed ONLY when dropping -# to a non-root user, 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". +# 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