diff --git a/omp-docs/log.md b/omp-docs/log.md index a7255a5..207ad8b 100644 --- a/omp-docs/log.md +++ b/omp-docs/log.md @@ -16,6 +16,29 @@ lesson, a few lines. Full incident narratives are recoverable from git (`git log --follow omp-docs/log.md`); knowledge that outlives an entry lives on its concept page, linked here. Never re-expand an entry in place. +## 2026-10-20 — two instruments rescued from `$HOME` script sediment + +`$HOME` had accumulated 123 loose scripts at its root (archived to +`.AGENTTRASH/home-scripts-2026-10-20`, ~1.8 GiB of stray audio consolidated into +`~/audiosamples`). Two were durable instruments and now live in the repo with +black-box tests and pages: [tools/remote-job](tools/remote-job.md) and +[tools/pcie-probe](tools/pcie-probe.md). Refs #53. Laws learned: + +- **A readiness pattern in a log is not evidence of completion.** `ssh host 'cmd | + tail -f log | sed -n /ready/q'` fails three independent ways — the broker aborts + the long connection, the pattern can match a header printed before the work + starts, and `sed -n /x/q` exits the moment it matches, tearing down the job it was + watching. The only honest signal is a sentinel the job writes last, read over a + short-lived connection. +- **Never benchmark beside a live service, and never let a tool decide to stop one.** + The guard refuses before anything is built; unloading a production endpoint is a + human's call. A number taken as noise beside serving traffic is both disruptive and + wrong. +- **An instrument that segfaults is not an instrument.** Two of the three shipped + `.hip` probes called `hipGetDeviceProperties`, which crashes on this ROCm build + while `hipGetDeviceCount` works fine; device identity now comes from sysfs, and the + trap is recorded in the sources so it is not reintroduced. + ## 2026-10-20 — the secret gate was reading the wrong files, and one credential had already escaped Found while deciding what else belongs in `dotfiles/`. (1) The gate whitelisted diff --git a/omp-docs/tools/index.md b/omp-docs/tools/index.md index e8e0abe..91a202e 100644 --- a/omp-docs/tools/index.md +++ b/omp-docs/tools/index.md @@ -20,6 +20,7 @@ Configuration guides and reference material for installed software. Each tool ha * [niri](niri/index.md) — Scrollable-tiling Wayland compositor — configuration, keybindings, and architecture for the workstation. * [noctalia-shell](noctalia-shell.md) — Native desktop environment for Niri — launcher, control center, settings, lock screen, session menu, hotkey overlay, media controls, and system management. * [packages](packages.md) — `packages/*.pkgs` desired state and the two tools that keep a machine and its manifest agreeing: `install.sh` installs the delta, `pkg-inventory.sh` reports what the machine has outgrown. +* [pcie-probe](pcie-probe.md) — HIP probes that separate PCIe link bandwidth from the engine's read path at the expert-slab size, and the refuse-by-default guard against benchmarking beside a live service. * [remote-job](remote-job.md) — Detached job runner for `ssh` work: completion sentinels on the host replace the `tail -f | sed -n /ready/q` pattern that reported success on jobs that never finished. * [tmux](tmux.md) — Terminal multiplexer configuration, extended-keys setup, keyboard input best practices, and reload/restart procedures. diff --git a/omp-docs/tools/pcie-probe.md b/omp-docs/tools/pcie-probe.md new file mode 100644 index 0000000..5a17f05 --- /dev/null +++ b/omp-docs/tools/pcie-probe.md @@ -0,0 +1,92 @@ +--- +type: reference +title: pcie-probe — link or read path +description: HIP instruments that separate PCIe link bandwidth from the engine's read path at the real expert-slab size, plus the refuse-by-default guard that stops a benchmark from running beside a live service. +tags: + - gpu + - rocm + - rdna4 + - benchmarking + - dankserver +verified: 2026-10-20 +timestamp: 2026-10-20 +--- + +# pcie-probe + +`scripts/pcie-probe.sh` answers one question, which comes up every time MoE +expert offload is discussed on the R9700 box +([dankserver](../system/dankserver.md)): **is the stall the PCIe link, or is it +the engine's read path?** The two look identical from the outside — decode is +slow while weights stream in — and they have opposite fixes. + +## Instruments + +Sources in `scripts/probes/`, built with the host `hipcc`: + +| Probe | Question it answers | +|-------|---------------------| +| `hip_smoke.hip` | is host ROCm usable from a plain process at all (count, alloc, copy) | +| `pcie_bw.hip` | host→device bandwidth at the engine's exact slab size, pinned vs pageable | +| `pcie_probe.hip` | the size sweep — where throughput collapses as transfers shrink | + +```bash +scripts/pcie-probe.sh # guard, build, topology, measure +scripts/pcie-probe.sh --device 0 --device 1 +scripts/pcie-probe.sh --build-only # compile, never touch the cards +``` + +## The law: never benchmark beside a live service + +The guard runs **before anything is built or run** and refuses when the GPU is +occupied — `gpu_busy_percent` above `PCIE_PROBE_MAX_BUSY` (default 5 %) or any +process holding `/dev/kfd`. Refusal is exit `3`, and it stops nothing. + +Auto-stopping a production endpoint is not this tool's decision to make, and a +measurement taken as noise beside serving traffic is worthless anyway: it is +both disruptive and wrong. Take the service down deliberately, or pass `--force` +and mean it. `tests/pcie-probe.bats` pins that the refusal happens before the +build step, so a contended box cannot be measured by accident. + +## Reading the numbers + +The reference is the **link**, printed as a topology block from sysfs +(`current_link_width` / `current_link_speed` against `max_*`): a card trained to +fewer lanes or a lower generation than it is wired for caps everything +downstream, and no engine change will recover it. + +Then the sweep table, whose first row is the engine's actual expert-slab size — +throughput at *that* size is what matters, not the big-transfer figure: + +- Small transfers reach the trained-link ceiling → the link is healthy, the + engine's read path (fetch granularity, pipelining, pinning) is the limit. +- Small transfers stall far below a large-transfer figure → the wall is the + link or the IOMMU path, not the engine. +- Pinned versus pageable columns apart means pinning is load-bearing; together + means the copies are already pinned by someone else. + +Measure one device at a time (`--device`). A run that lets the runtime pick +several devices cannot attribute bandwidth to a card at all. + +## The trap carried in the sources + +`hipGetDeviceProperties` **segfaults** on the ROCm build these probes were +written against, while `hipGetDeviceCount` works fine and reports every device. +None of the three sources calls it; device identity comes from sysfs instead. +The setvbuf call in `hip_smoke.hip` is a fossil of discovering that — an +unflushed stdout hid the crash point for an afternoon. Keep both facts intact +when touching these files. + +## Exit status and seams + +`0` measured · `2` usage or no sources · `3` refused, GPU busy · `4` no HIP +toolchain or build failed. A client box without the SDK exits `4` loudly rather +than reporting an empty success; the probes belong on the box with the cards. + +Test seams (all documented in the script header): `PCIE_PROBE_SRC`, +`PCIE_PROBE_BUILD`, `PCIE_PROBE_SYSFS`, `PCIE_PROBE_KFD`, `PCIE_PROBE_HIPCC`, +`PCIE_PROBE_MAX_BUSY` — the sysfs/kfd paths are what make a refusal testable on +a machine with no cards, which is the only honest way to test a refusal. + +Measured campaign numbers are not kept here; they belong with the endpoint they +were taken on, in the serverAppConfigs LLM runbook vault. diff --git a/scripts/pcie-probe.sh b/scripts/pcie-probe.sh new file mode 100755 index 0000000..dd1b03d --- /dev/null +++ b/scripts/pcie-probe.sh @@ -0,0 +1,126 @@ +#!/usr/bin/env bash +# scripts/pcie-probe.sh — is the decode bottleneck the LINK, or the engine's read path? +# +# Instruments live in scripts/probes/: hip_smoke.hip answers "is the host ROCm usable from a plain +# process", pcie_bw.hip measures host→device bandwidth at the engine's real expert-slab size, and +# pcie_probe.hip sweeps sizes so a stall can be attributed to the link rather than to the engine. +# Written during the RDNA4 expert-offload campaign, kept because the same cards are in Dankserver. +# +# LAW — never benchmark beside a live service. The guard below runs BEFORE anything is built or run, +# and it stops nothing: putting a production endpoint down is not this tool's decision to make. When +# you have deliberately decided that contention is acceptable, say so with --force. +# +# Trap carried in the sources: hipGetDeviceProperties SEGFAULTS on this ROCm build even though +# hipGetDeviceCount works fine and reports every device. Nothing here calls it. Keep it that way. +# +# usage: pcie-probe.sh [--device N]... [--build-only] [--force] +# +# --device N measure this device; repeat for each card. Without it the probes run once with the +# environment's own device selection — one shared run cannot attribute bandwidth. +# --build-only compile and stop; never touches the cards. +# --force measure even though the guard says the GPU is busy. +# +# Environment (defaults in parentheses) — the sysfs and kfd paths are overridable so the guard can +# be tested on a machine with no cards, which is the only honest way to test a refusal: +# PCIE_PROBE_SRC probe sources (/scripts/probes) +# PCIE_PROBE_BUILD build output ($HOME/.cache/pcie-probe) +# PCIE_PROBE_SYSFS drm sysfs root (/sys/class/drm) +# PCIE_PROBE_KFD KFD device node (/dev/kfd) +# PCIE_PROBE_HIPCC compiler (hipcc) +# PCIE_PROBE_MAX_BUSY refusal threshold, percent (5) +# +# Exit status: 0 measured · 2 usage or missing sources · 3 REFUSED, GPU busy · 4 no toolchain or +# build failed. + +set -uo pipefail + +SRC="${PCIE_PROBE_SRC:-$(dirname "$(realpath "$0")")/probes}" +BUILD="${PCIE_PROBE_BUILD:-$HOME/.cache/pcie-probe}" +SYSFS="${PCIE_PROBE_SYSFS:-/sys/class/drm}" +KFD="${PCIE_PROBE_KFD:-/dev/kfd}" +HIPCC="${PCIE_PROBE_HIPCC:-hipcc}" +MAX_BUSY="${PCIE_PROBE_MAX_BUSY:-5}" + +force=0 +build_only=0 +devices=() +while [[ $# -gt 0 ]]; do + case "$1" in + --device) [[ $# -ge 2 ]] || { echo "pcie-probe: --device needs an argument" >&2; exit 2; } + devices+=("$2"); shift 2 ;; + --build-only) build_only=1; shift ;; + --force) force=1; shift ;; + *) echo "usage: pcie-probe.sh [--device N]... [--build-only] [--force]" >&2; exit 2 ;; + esac +done + +mapfile -t sources < <(find "$SRC" -maxdepth 1 -name '*.hip' 2>/dev/null | sort) +if (( ${#sources[@]} == 0 )); then + echo "pcie-probe: no .hip probe sources in $SRC" >&2 + exit 2 +fi + +# ---- the guard, before anything touches the card -------------------------------- +if (( ! force )); then + busy="" + if [[ -e "$KFD" ]] && command -v fuser > /dev/null 2>&1; then + holders="$(fuser "$KFD" 2>/dev/null | tr -s '[:space:]' ' ')" + [[ -n "${holders// /}" ]] && busy="KFD in use by:${holders}" + fi + for card in "$SYSFS"/card*; do + [[ -f "$card/device/gpu_busy_percent" ]] || continue + pct="$(cat "$card/device/gpu_busy_percent" 2>/dev/null)" + [[ "$pct" =~ ^[0-9]+$ ]] || continue + (( pct > MAX_BUSY )) && busy="${busy:+$busy; }$(basename "$card") at ${pct}%" + done + if [[ -n "$busy" ]]; then + printf 'pcie-probe: REFUSED — GPU busy (%s)\n' "$busy" >&2 + printf ' never benchmark beside a live service. Take the service down deliberately,\n' >&2 + printf ' or pass --force if you have decided the contention is acceptable.\n' >&2 + exit 3 + fi +fi + +# ---- build ---------------------------------------------------------------------- +if ! command -v "$HIPCC" > /dev/null 2>&1; then + printf 'pcie-probe: no HIP toolchain (%s) — run this on the box with the cards, not here\n' "$HIPCC" >&2 + exit 4 +fi +mkdir -p "$BUILD" +for src in "${sources[@]}"; do + bin="$BUILD/$(basename "${src%.hip}")" + if ! "$HIPCC" -O2 -o "$bin" "$src" 2>&1 | tail -3; then + printf 'pcie-probe: build failed for %s\n' "$(basename "$src")" >&2 + exit 4 + fi +done +if (( build_only )); then + printf 'pcie-probe: built %s probe(s) into %s\n' "${#sources[@]}" "$BUILD" + exit 0 +fi + +# ---- topology: the link the numbers must be read against ------------------------ +printf 'pcie-probe host=%s hipcc=%s probes=%s\n' "$(hostname)" "$HIPCC" "${#sources[@]}" +for card in "$SYSFS"/card*; do + [[ -f "$card/device/current_link_width" ]] || continue + printf ' link %-9s current x%s @ %s max x%s @ %s\n' "$(basename "$card")" \ + "$(cat "$card/device/current_link_width")" \ + "$(cat "$card/device/current_link_speed" 2>/dev/null || echo '?')" \ + "$(cat "$card/device/max_link_width" 2>/dev/null || echo '?')" \ + "$(cat "$card/device/max_link_speed" 2>/dev/null || echo '?')" +done + +# ---- measure, per device -------------------------------------------------------- +(( ${#devices[@]} > 0 )) || devices=(default) +for d in "${devices[@]}"; do + if [[ $d == default ]]; then + printf '\n########## engine device selection ##########\n' + else + printf '\n########## device %s ##########\n' "$d" + fi + for name in hip_smoke pcie_bw pcie_probe; do + bin="$BUILD/$name" + [[ -x "$bin" ]] || continue + if [[ $d == default ]]; then "$bin"; else HIP_VISIBLE_DEVICES="$d" "$bin" "$d"; fi + done +done diff --git a/scripts/probes/hip_smoke.hip b/scripts/probes/hip_smoke.hip new file mode 100644 index 0000000..450a668 --- /dev/null +++ b/scripts/probes/hip_smoke.hip @@ -0,0 +1,36 @@ +// Minimal HIP smoke test: is the host ROCm usable from a plain process? +#include +#include +#include + +int main(void) { + int n = -1; + hipError_t e = hipGetDeviceCount(&n); + printf("hipGetDeviceCount -> %d (%s), count=%d\n", + (int)e, hipGetErrorString(e), n); + if (e != hipSuccess || n <= 0) return 1; + setvbuf(stdout, NULL, _IONBF, 0); /* an earlier crash lost the buffer, hiding the crash point */ + // Device properties are deliberately NOT queried here: hipGetDeviceProperties segfaults on the + // ROCm build this probe was written for, even though hipGetDeviceCount above works fine. Card + // identity comes from sysfs anyway — scripts/pcie-probe.sh prints the link topology. + + const size_t bytes = 2459136; /* 2.344 MiB: the engine's expert slab */ + void *dev = NULL, *host = NULL; + if (hipMalloc(&dev, bytes) != hipSuccess) { printf("hipMalloc failed\n"); return 1; } + if (hipHostMalloc(&host, bytes, hipHostMallocDefault) != hipSuccess) { + printf("hipHostMalloc failed\n"); return 1; + } + memset(host, 0x5a, bytes); + + hipEvent_t a, b; hipEventCreate(&a); hipEventCreate(&b); + hipDeviceSynchronize(); + hipEventRecord(a, 0); + const int iters = 200; + for (int i = 0; i < iters; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice); + hipEventRecord(b, 0); hipEventSynchronize(b); + float ms = 0; hipEventElapsedTime(&ms, a, b); + printf("\n %.3f MiB x %d iters in %.2f ms -> %.2f GB/s H2D (pinned)\n", + bytes/1048576.0, iters, ms, (double)bytes*iters/(ms/1000.0)/1e9); + printf(" Gen5 x8 ceiling = 31.5 GB/s\n"); + return 0; +} diff --git a/scripts/probes/pcie_bw.hip b/scripts/probes/pcie_bw.hip new file mode 100644 index 0000000..85702e2 --- /dev/null +++ b/scripts/probes/pcie_bw.hip @@ -0,0 +1,64 @@ +// H2D bandwidth at the engine's expert-slab size. +// Deliberately avoids hipGetDeviceProperties: it segfaults on this ROCm build +// (hipGetDeviceCount works fine and reports 3 devices). +#include +#include +#include +#include + +static double one(size_t bytes, bool pinned, int iters) { + void *dev = NULL, *host = NULL; + if (hipMalloc(&dev, bytes) != hipSuccess) return -1; + if (pinned) { if (hipHostMalloc(&host, bytes, hipHostMallocDefault) != hipSuccess) return -1; } + else { host = malloc(bytes); } + memset(host, 0x5a, bytes); + + hipEvent_t a, b; hipEventCreate(&a); hipEventCreate(&b); + for (int i = 0; i < 3; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice); + hipDeviceSynchronize(); + + hipEventRecord(a, 0); + for (int i = 0; i < iters; i++) hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice); + hipEventRecord(b, 0); hipEventSynchronize(b); + float ms = 0; hipEventElapsedTime(&ms, a, b); + double gb = (double)bytes * iters / (ms / 1000.0) / 1e9; + + if (!pinned) free(host); else hipHostFree(host); + hipFree(dev); hipEventDestroy(a); hipEventDestroy(b); + return gb; +} + +int main(int argc, char **argv) { + setvbuf(stdout, NULL, _IONBF, 0); + int dev = (argc > 1) ? atoi(argv[1]) : 0; + if (hipSetDevice(dev) != hipSuccess) { printf("setDevice(%d) failed\n", dev); return 1; } + printf("device %d selected\n", dev); + + size_t sz[] = {2459136UL, 3UL<<20, 16UL<<20, 64UL<<20, 256UL<<20, 1UL<<30}; + printf(" %10s | %7s | %7s | %7s\n", "MiB", "pinned", "pagin.", "xfer us"); + printf(" %s\n", "-------------------------------------------"); + for (int i = 0; i < 6; i++) { + size_t b = sz[i]; + int iters = (b >= (1UL<<30)) ? 10 : (b >= (256UL<<20)) ? 20 : 100; + double p = one(b, true, iters); + double g = one(b, false, iters); + // single-transfer latency + void *dv, *hv; + if (hipMalloc(&dv, b) == hipSuccess && hipHostMalloc(&hv, b, hipHostMallocDefault) == hipSuccess) { + memset(hv, 0x5a, b); + hipEvent_t a, c; hipEventCreate(&a); hipEventCreate(&c); + hipMemcpy(dv, hv, b, hipMemcpyHostToDevice); hipDeviceSynchronize(); + hipEventRecord(a, 0); + hipMemcpy(dv, hv, b, hipMemcpyHostToDevice); + hipEventRecord(c, 0); hipEventSynchronize(c); + float ms = 0; hipEventElapsedTime(&ms, a, c); + printf(" %10.3f | %7.2f | %7.2f | %7.1f\n", b/1048576.0, p, g, ms*1000.0); + hipHostFree(hv); hipFree(dv); + } else { + printf(" %10.3f | %7.2f | %7.2f | (alloc failed)\n", b/1048576.0, p, g); + } + } + printf("\n 2.344 MiB is the engine's exact expert slab size.\n"); + printf(" Gen5 x8 ceiling = 31.5 GB/s per direction.\n"); + return 0; +} diff --git a/scripts/probes/pcie_probe.hip b/scripts/probes/pcie_probe.hip new file mode 100644 index 0000000..b35fd4a --- /dev/null +++ b/scripts/probes/pcie_probe.hip @@ -0,0 +1,87 @@ +// HIP adaptation of the pcie-link-capability-probe. +// Question: is the decode bottleneck the LINK, or the engine's read PATH? +// Measure H2D at the engine's actual expert slab size (2.344 MiB) and across +// a size sweep. If 2.344 MiB copies reach ~31 GB/s, the link is healthy and +// the engine is leaving bandwidth on the table. If they stall near 10 GB/s, +// the link is the wall. +// +// build: hipcc -O2 -o pcie_probe pcie_probe.hip (scripts/pcie-probe.sh does this and guards the run) +#include +#include +#include +#include + +#define CK(x) do { hipError_t e = (x); if (e != hipSuccess) { \ + fprintf(stderr, "HIP %s (%s:%d)\n", hipGetErrorString(e), __FILE__, __LINE__); \ + return 1; } } while (0) + +static double ms_between(hipEvent_t a, hipEvent_t b) { + float ms = 0; hipEventElapsedTime(&ms, a, b); return (double)ms; +} + +static int bench_size(size_t bytes) { + const int iters = bytes >= (1UL<<30) ? 20 + : bytes >= (256UL<<20) ? 30 + : bytes >= (16UL<<20) ? 60 : 150; + void *dev = NULL, *host = NULL, *pageable = NULL; + hipStream_t s = NULL; + CK(hipMalloc(&dev, bytes)); + CK(hipHostMalloc(&host, bytes, hipHostMallocDefault)); // PINNED, like the engine's slabs + memset(host, 0x5a, bytes); + CK(hipStreamCreate(&s)); + hipEvent_t a, b; CK(hipEventCreate(&a)); CK(hipEventCreate(&b)); + + for (int i = 0; i < 3; i++) CK(hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice)); + CK(hipDeviceSynchronize()); + + // sync copy + CK(hipEventRecord(a, 0)); + for (int i = 0; i < iters; i++) CK(hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice)); + CK(hipEventRecord(b, 0)); CK(hipEventSynchronize(b)); + double h2d = (double)bytes * iters / (ms_between(a,b)/1000.0) / 1e9; + + // async copy on a stream (the engine's fill pipeline shape) + CK(hipEventRecord(a, 0)); + for (int i = 0; i < iters; i++) CK(hipMemcpyAsync(dev, host, bytes, hipMemcpyHostToDevice, s)); + CK(hipEventRecord(b, 0)); CK(hipEventSynchronize(b)); + double h2da = (double)bytes * iters / (ms_between(a,b)/1000.0) / 1e9; + + // single-transfer latency at this size + CK(hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice)); + CK(hipDeviceSynchronize()); + CK(hipEventRecord(a, 0)); + CK(hipMemcpy(dev, host, bytes, hipMemcpyHostToDevice)); + CK(hipEventRecord(b, 0)); CK(hipEventSynchronize(b)); + double one_us = ms_between(a,b) * 1000.0; + + // pageable, to see whether pinning is load-bearing here + pageable = malloc(bytes); memset(pageable, 0x5a, bytes); + CK(hipMemcpy(dev, pageable, bytes, hipMemcpyHostToDevice)); + CK(hipEventRecord(a, 0)); + for (int i = 0; i < iters; i++) CK(hipMemcpy(dev, pageable, bytes, hipMemcpyHostToDevice)); + CK(hipEventRecord(b, 0)); CK(hipEventSynchronize(b)); + double pag = (double)bytes * iters / (ms_between(a,b)/1000.0) / 1e9; + + printf(" %8.3f MiB | %6.2f | %6.2f | %8.1f us | %6.2f\n", + bytes/1048576.0, h2d, h2da, one_us, pag); + + free(pageable); hipHostFree(host); hipFree(dev); + hipEventDestroy(a); hipEventDestroy(b); hipStreamDestroy(s); + return 0; +} + +int main(int argc, char **argv) { + int dev_id = 0; + if (argc > 1) dev_id = atoi(argv[1]); + // No hipGetDeviceProperties here: it segfaults on this ROCm build while hipGetDeviceCount is + // fine, and the link topology these numbers must be read against comes from sysfs, which + // scripts/pcie-probe.sh prints. Only the device index is needed to pick what to measure. + printf("device %d (PCI bus from sysfs, see the topology block above)\n", dev_id); + printf(" size | H2D | async | 1x xfer | pageable [GB/s]\n"); + printf(" %s\n", "-----------------------------------------------------------"); + size_t sz[] = {2459136UL, /* 2.344 MiB: the engine's exact expert slab */ + 3UL<<20, 16UL<<20, 64UL<<20, 256UL<<20, 1UL<<30, 2UL<<30}; + for (int i = 0; i < 7; i++) if (bench_size(sz[i])) return 1; + printf("\nGen5 x8 ceiling = 31.5 GB/s per direction\n"); + return 0; +} diff --git a/tests/pcie-probe.bats b/tests/pcie-probe.bats new file mode 100755 index 0000000..9c74afc --- /dev/null +++ b/tests/pcie-probe.bats @@ -0,0 +1,144 @@ +#!/usr/bin/env bats +# tests/pcie-probe.bats — black-box tests for scripts/pcie-probe.sh (Refs #53) +# +# Contract: measure host→device bandwidth on an RDNA4 box, and refuse to measure while the GPU is +# in use. Fleet law: never benchmark beside a live service — so the guard runs BEFORE anything is +# built or run, and the tool stops nothing and starts nothing (auto-stopping a production service is +# the thing it must not do unasked). +# +# Exit status: 0 measured · 2 usage · 3 refused (GPU busy) · 4 no toolchain / build failed. +# +# hipcc and the compiled probes are stubbed: this box has no HIP SDK (it is a client laptop) and a +# real measurement belongs on a machine with the cards, beside no running model. + +setup() { + TEST_TMP="$(mktemp -d)" + SCRIPT="$BATS_TEST_DIRNAME/../scripts/pcie-probe.sh" + SRC="$TEST_TMP/src" + BUILD="$TEST_TMP/build" + SYSFS="$TEST_TMP/drm" + BIN="$TEST_TMP/bin" + mkdir -p "$SRC" "$BUILD" "$BIN" "$SYSFS/card0/device" "$SYSFS/card1/device" + + # Fixtures stand in for the real sysfs files the driver reads. + printf '8\n' > "$SYSFS/card0/device/current_link_width" + printf '32.0 GT/s PCIe\n' > "$SYSFS/card0/device/current_link_speed" + printf '16\n' > "$SYSFS/card0/device/max_link_width" + printf '32.0 GT/s PCIe\n' > "$SYSFS/card0/device/max_link_speed" + printf '0\n' > "$SYSFS/card0/device/gpu_busy_percent" + printf '8\n' > "$SYSFS/card1/device/current_link_width" + printf '32.0 GT/s PCIe\n' > "$SYSFS/card1/device/current_link_speed" + printf '16\n' > "$SYSFS/card1/device/max_link_width" + printf '32.0 GT/s PCIe\n' > "$SYSFS/card1/device/max_link_speed" + printf '0\n' > "$SYSFS/card1/device/gpu_busy_percent" + + # Probe sources: the real .hip files are shipped by the script's own directory, so the fixtures + # only need to exist for the build step to be invoked. + : > "$SRC/pcie_probe.hip"; : > "$SRC/pcie_bw.hip"; : > "$SRC/hip_smoke.hip" + + # Stub toolchain: records that it was called, and emits a probe binary that records that IT ran. + cat > "$BIN/hipcc" <<'STUB' +#!/usr/bin/env bash +[ "${STUB_HIPCC_FAILS:-0}" = 1 ] && { echo "stub: no matching GPU target" >&2; exit 1; } +out="" +while [ $# -gt 0 ]; do [ "$1" = "-o" ] && out="$2"; shift; done +[ -n "$out" ] || exit 2 +: > "$PCIE_PROBE_BUILD/.built" +{ echo '#!/usr/bin/env bash' + echo ': > "$PCIE_PROBE_BUILD/.ran"' + echo 'echo "PROBE RAN dev=${HIP_VISIBLE_DEVICES:-none} bin=$(basename "$0")"' + echo 'echo " slab 2.344 MiB: 31.4 GB/s"' +} > "$out" +chmod +x "$out" +STUB + cat > "$BIN/fuser" <<'STUB' +#!/usr/bin/env bash +[ -n "${STUB_KFD_PIDS:-}" ] && printf '%s\n' $STUB_KFD_PIDS +exit 0 +STUB + chmod +x "$BIN/hipcc" "$BIN/fuser" + + export PATH="$BIN:$PATH" PCIE_PROBE_SRC="$SRC" PCIE_PROBE_BUILD="$BUILD" \ + PCIE_PROBE_SYSFS="$SYSFS" PCIE_PROBE_KFD="$TEST_TMP/kfd" PCIE_PROBE_HIPCC="$BIN/hipcc" + : > "$TEST_TMP/kfd" +} + +teardown() { rm -rf "$TEST_TMP"; } + +@test "a missing HIP toolchain is a loud failure, not a silent skip" { + run env PCIE_PROBE_HIPCC="$TEST_TMP/nope-hipcc" "$SCRIPT" + [ "$status" -eq 4 ] + [[ "$output" == *"hip"* ]] + [ ! -e "$BUILD/.built" ] +} + +@test "a failed build is reported as a failure" { + run env STUB_HIPCC_FAILS=1 "$SCRIPT" + [ "$status" -eq 4 ] + [[ "$output" == *"build"* ]] +} + +@test "a GPU held by another process refuses the measurement before anything is built" { + # Order is the contract: build-and-then-ask would already have loaded the card beside a live + # service. The .built marker proves the guard ran first. + run env STUB_KFD_PIDS="4242" "$SCRIPT" + [ "$status" -eq 3 ] + [[ "$output" == *"busy"* ]] + [ ! -e "$BUILD/.built" ] +} + +@test "a busy GPU percent refuses too" { + printf '91\n' > "$SYSFS/card0/device/gpu_busy_percent" + run "$SCRIPT" + [ "$status" -eq 3 ] + [[ "$output" == *"card0"* ]] + [ ! -e "$BUILD/.built" ] +} + +@test "--force overrides the guard for a deliberately contended measurement" { + run env STUB_KFD_PIDS="4242" "$SCRIPT" --force + [ "$status" -eq 0 ] + [ -e "$BUILD/.ran" ] +} + +@test "an idle GPU builds and runs the probe" { + run "$SCRIPT" + [ "$status" -eq 0 ] + [ -e "$BUILD/.built" ] + [ -e "$BUILD/.ran" ] + [[ "$output" == *"PROBE RAN"* ]] + [[ "$output" == *"31.4 GB/s"* ]] # probe output passes through untouched +} + +@test "--build-only compiles without touching the card" { + run "$SCRIPT" --build-only + [ "$status" -eq 0 ] + [ -e "$BUILD/.built" ] + [ ! -e "$BUILD/.ran" ] +} + +@test "the link topology is reported from sysfs, current against max" { + run "$SCRIPT" + [[ "$output" == *"card0"*"x8"*"32.0 GT/s"* ]] + [[ "$output" == *"max x16"* ]] +} + +@test "each requested device is measured separately" { + # One shared run cannot attribute bandwidth to a card; the whole point is per-card numbers. + run "$SCRIPT" --device 0 --device 1 + [ "$status" -eq 0 ] + [[ "$output" == *"dev=0"* ]] + [[ "$output" == *"dev=1"* ]] +} + +@test "no probe sources is a usage-level error, not an empty success" { + rm -f "$SRC"/*.hip + run "$SCRIPT" + [ "$status" -ne 0 ] + [[ "$output" == *"probe"* ]] +} + +@test "an unknown flag is a usage error" { + run "$SCRIPT" --frobnicate + [ "$status" -eq 2 ] +}