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 f990fd2..91a202e 100644 --- a/omp-docs/tools/index.md +++ b/omp-docs/tools/index.md @@ -20,6 +20,8 @@ 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. ## Adding a New Tool 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/omp-docs/tools/remote-job.md b/omp-docs/tools/remote-job.md new file mode 100644 index 0000000..94b526d --- /dev/null +++ b/omp-docs/tools/remote-job.md @@ -0,0 +1,83 @@ +--- +type: reference +title: remote-job — detached jobs that outlive the session +description: Run a command on a remote box so its result survives the connection, the restart, and the agent that started it — the replacement for ssh piped through tail and sed, which reported success on jobs that never finished. +tags: + - ssh + - fleet + - automation + - verification +verified: 2026-10-20 +timestamp: 2026-10-20 +--- + +# remote-job + +`scripts/remote-job.sh` starts a command **detached on the host**, and reports +what actually happened from files on disk. It exists because +`ssh host 'cmd | tail -f log | sed -n /ready/q'` lied repeatedly during the +worktree campaign. + +## Why the pipe lies + +Three independent failure modes, all of them structural: + +- **The broker kills long-lived ssh.** A job longer than the connection's + patience dies mid-run; the caller sees a closed pipe, not a failure. +- **A readiness pattern can match a log header.** `sed -n /building/q` fires on + a banner printed before the work starts — the job is at line 1 and the caller + already moved on. +- **`sed -n /x/q` exits the instant it matches.** Even when the pattern is + genuinely ready-shaped, the pipeline tears down the remote command it was + watching. Result: green report, dead job. + +The only reliable signal is a **completion sentinel written by the job itself**, +read afterwards by a short-lived connection. That is what this is. + +## Contract + +```bash +remote-job.sh start NAME [--] CMD... # launch detached, return at once +remote-job.sh status [NAME] # one line per job; exit 0 only if all DONE +remote-job.sh wait NAME [SEC] [N] # block until DONE or timeout, then tail N lines +remote-job.sh out NAME [N] # tail a job's output +remote-job.sh tail [NAME] # follow one job, or every job at once +remote-job.sh watch # refreshing status, for a spare pane +``` + +Job dir `$REMOTE_JOBS_DIR` (default `$HOME/jobs`). Per job: `NAME.out` (stdout+ +stderr), `NAME.rc` (exit code), `NAME.done` (sentinel, **written last**). + +Exit status of the *report*: `0` all reported jobs finished · `1` unfinished or +timed out · `2` usage. The `rc=` column is the job's own code — a job that +failed is still `DONE rc=7` and `status` exits `0`, because "is it finished" and +"did it work" are different questions a supervisor must be able to ask +separately. + +Written for the case where the job **restarts the endpoint the agent itself is +running on**: the ssh connection dies, the session may die, and the answer is +still sitting in `$JOBS` when a session comes back. + +## Use over ssh + +```bash +ssh dankserver ~/.omp/scripts/remote-job.sh start swap -- \ + bash -c 'systemctl --user restart frank-archon && ' +ssh dankserver ~/.omp/scripts/remote-job.sh wait swap 300 20 # short-lived reconnect +``` + +Each ssh call is short, so there is nothing there to abort, and the second call +can be made from a fresh session on a machine that rebooted underneath it. + +## Failure modes + +| Event | Result | +|-------|--------| +| Session/connection dies mid-job | job keeps running; `status` from anywhere reports it | +| Command name is wrong | `PENDING (no output yet)`, exit 1 — never a false DONE | +| Job exits non-zero | `DONE rc=N`; the code is data, not the report's status | +| Two jobs, same NAME | the second replaces the first's state files | +| Job dir missing | created on first use | + +Pinned by `tests/remote-job.bats` (black-box: output, exit codes, sentinel +ordering — never internals). 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/scripts/remote-job.sh b/scripts/remote-job.sh new file mode 100755 index 0000000..20082b3 --- /dev/null +++ b/scripts/remote-job.sh @@ -0,0 +1,115 @@ +#!/usr/bin/env bash +# scripts/remote-job.sh — a detached job whose result outlives the session that started it +# +# Why this exists rather than `ssh host 'cmd | tail -f | sed -n /ready/q'`, which lied about success +# repeatedly, for three structural reasons — all of them the pipe's fault: the broker aborts +# long-lived ssh; a readiness pattern can match a log *header* before the work has begun; and +# `sed -n /x/q` exits the instant it matches, so the supervisor reports green on a job that never +# finished. +# +# Here the job runs detached on the box it is invoked on (over ssh, or locally). It writes NAME.out, +# then NAME.rc, then touches NAME.done LAST. Status is then a cheap read of a directory: no +# long-lived connection to lose, and the answer survives the session, the terminal, and the agent — +# including the case where the job restarts the very service the agent is running on. +# +# remote-job.sh start NAME [--] CMD... launch detached, return at once +# remote-job.sh status [NAME] one line per job +# remote-job.sh wait NAME [SECONDS] [N] block until DONE or timeout, then tail N lines +# remote-job.sh out NAME [LINES] print the tail of a job's output +# remote-job.sh tail [NAME] follow one job, or every job at once +# remote-job.sh watch refreshing status, for a spare pane +# +# Job directory: $REMOTE_JOBS_DIR, default $HOME/jobs. Ship this file to the host and run it there; +# nothing about it is box-specific, which is the point. +# +# Exit status: 0 when every job reported is finished · 1 when something is unfinished or a wait +# timed out · 2 usage. A job that FAILED still counts as finished for `status`: "is it done" and +# "did it work" are different questions, and the second one is the rc= column. + +set -uo pipefail + +JOBS="${REMOTE_JOBS_DIR:-$HOME/jobs}" +mkdir -p "$JOBS" + +case "${1:-}" in + start) + [[ $# -ge 2 ]] || { echo "usage: remote-job.sh start NAME [--] CMD..." >&2; exit 2; } + name="$2"; shift 2 + [[ "${1:-}" == "--" ]] && shift + [[ $# -ge 1 ]] || { echo "remote-job: start needs a command for '$name'" >&2; exit 2; } + # Clear the previous run entirely: a stale .out would make the new run's line count a lie. + rm -f "$JOBS/$name.out" "$JOBS/$name.rc" "$JOBS/$name.done" + # setsid leaves the invoking process group, so a dropped ssh or a killed agent takes nothing + # down with it. rc lands BEFORE the sentinel: a waiter must never see DONE with no exit code. + JOBS="$JOBS" setsid nohup bash -c ' + name="$1"; shift + "$@" > "$JOBS/$name.out" 2>&1 + echo $? > "$JOBS/$name.rc" + touch "$JOBS/$name.done" + ' _ "$name" "$@" < /dev/null > /dev/null 2>&1 & + printf 'started %s\n' "$name" + ;; + + status) + if [[ -n "${2:-}" ]]; then + names="$2" + else + names=$(cd "$JOBS" && ls *.out 2>/dev/null | sed 's/\.out$//' | sort) + fi + pending=0 + for n in $names; do + if [[ -f "$JOBS/$n.done" ]]; then + printf ' %-20s DONE rc=%s %s lines\n' "$n" \ + "$(cat "$JOBS/$n.rc" 2>/dev/null)" "$(wc -l < "$JOBS/$n.out" 2>/dev/null | tr -d ' ')" + elif [[ -f "$JOBS/$n.out" ]]; then + printf ' %-20s RUN %s lines so far\n' "$n" "$(wc -l < "$JOBS/$n.out" 2>/dev/null | tr -d ' ')" + pending=1 + else + printf ' %-20s PENDING (no output yet)\n' "$n" + pending=1 + fi + done + if [[ -z "$names" ]]; then echo " (no jobs)"; exit 0; fi + exit "$pending" + ;; + + wait) + [[ $# -ge 2 ]] || { echo "usage: remote-job.sh wait NAME [SECONDS] [TAIL]" >&2; exit 2; } + name="$2"; limit="${3:-3600}"; lines="${4:-25}"; waited=0 + while [[ ! -e "$JOBS/$name.done" ]] && (( waited < limit )); do + sleep 1; waited=$((waited + 1)) + done + if [[ -e "$JOBS/$name.done" ]]; then + printf 'DONE %s rc=%s\n' "$name" "$(cat "$JOBS/$name.rc" 2>/dev/null)"; rc=0 + else + printf 'TIMEOUT %s after %ss\n' "$name" "$waited"; rc=1 + fi + tail -n "$lines" "$JOBS/$name.out" 2>/dev/null + exit "$rc" + ;; + + out) + [[ $# -ge 2 ]] || { echo "usage: remote-job.sh out NAME [LINES]" >&2; exit 2; } + tail -n "${3:-40}" "$JOBS/$2.out" 2>/dev/null + ;; + + tail) + # Multiplexed follow: no NAME means every job at once, so one pane shows the whole harness. + if [[ -n "${2:-}" ]]; then exec tail -f "$JOBS/$2.out"; fi + exec tail -f "$JOBS"/*.out + ;; + + watch) + while true; do + clear 2>/dev/null || true + date '+%H:%M:%S' + "$0" status + sleep 5 + done + ;; + + *) + echo "usage: remote-job.sh {start|status|wait|out|tail|watch} ..." >&2 + exit 2 + ;; +esac 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 ] +} diff --git a/tests/remote-job.bats b/tests/remote-job.bats new file mode 100755 index 0000000..c980ffc --- /dev/null +++ b/tests/remote-job.bats @@ -0,0 +1,135 @@ +#!/usr/bin/env bats +# tests/remote-job.bats — black-box tests for scripts/remote-job.sh (Refs #53) +# +# Contract: a job runs detached, writes its output, then its exit code, then touches .done LAST. +# Status is then a cheap independent read that survives the session that started the job — the +# whole point is replacing `ssh host 'tail -f | sed'`, which reported success on jobs that had not +# finished. Exit status answers "is everything finished"; the rc= field answers "did it work". +# +# Jobs really run here, against a temp job dir. No fake filesystem state is pre-seeded, because a +# fixture that pre-creates .done proves nothing about the code that is supposed to create it. + +setup() { + TEST_TMP="$(mktemp -d)" + SCRIPT="$BATS_TEST_DIRNAME/../scripts/remote-job.sh" + export REMOTE_JOBS_DIR="$TEST_TMP/jobs" +} + +teardown() { + # Jobs are detached on purpose; leave nothing behind that could outlive the test run. + pkill -P 1 -f "$TEST_TMP" 2>/dev/null || true + rm -rf "$TEST_TMP" +} + +# until_done NAME SECONDS — poll instead of sleeping a fixed amount, so the suite is neither slow +# nor flaky on a loaded machine. +until_done() { + local name="$1" limit="${2:-10}" waited=0 + while [[ ! -e "$REMOTE_JOBS_DIR/$name.done" && $waited -lt $((limit * 10)) ]]; do + sleep 0.1; waited=$((waited + 1)) + done + [[ -e "$REMOTE_JOBS_DIR/$name.done" ]] +} + +@test "start returns immediately while the job is still running" { + run timeout 10 "$SCRIPT" start slow -- sleep 2 + [ "$status" -eq 0 ] + [[ "$output" == "started slow" ]] + # A 2-second job must not have blocked a start that is supposed to return at once, and while it + # is still in flight the report must say unfinished — never DONE. + run "$SCRIPT" status slow + [[ "$output" == *RUN* || "$output" == *PENDING* ]] + [ "$status" -eq 1 ] +} + +@test "a successful job is reported DONE with its output" { + "$SCRIPT" start ok -- bash -c 'echo marker-4711' + until_done ok + run "$SCRIPT" status ok + [[ "$output" == *"ok"*DONE*rc=0* ]] + [ "$status" -eq 0 ] + run "$SCRIPT" out ok + [[ "$output" == *marker-4711* ]] +} + +@test "a failing job's exit code survives into the report" { + # This is the reason the .rc file exists: a supervisor that cannot see the failure code is the + # bug the runner was written to fix. + "$SCRIPT" start boom -- bash -c 'echo it failed here >&2; exit 7' + until_done boom + run "$SCRIPT" status boom + [[ "$output" == *"boom"*DONE*rc=7* ]] + run "$SCRIPT" out boom + [[ "$output" == *"it failed here"* ]] +} + +@test "the .done sentinel is never written before the exit code" { + # If .done landed first, a waiter could read an empty rc and call the job successful. + "$SCRIPT" start order -- bash -c 'echo x' + until_done order + [ -f "$REMOTE_JOBS_DIR/order.rc" ] + [ -s "$REMOTE_JOBS_DIR/order.rc" ] + run cat "$REMOTE_JOBS_DIR/order.rc" + [[ "$output" =~ ^[0-9]+$ ]] +} + +@test "wait prints DONE with the exit code and a tail of the output" { + "$SCRIPT" start fin -- bash -c 'echo one; echo two' + run "$SCRIPT" wait fin 15 + [ "$status" -eq 0 ] + [[ "$output" == *"DONE fin rc=0"* ]] + [[ "$output" == *two* ]] +} + +@test "wait gives up on its own and says so" { + "$SCRIPT" start never -- sleep 30 + run "$SCRIPT" wait never 1 + [ "$status" -eq 1 ] + [[ "$output" == *TIMEOUT*never* ]] + "$SCRIPT" status never | grep -q . # the job is still there, unfinished +} + +@test "status distinguishes running from pending from done" { + run "$SCRIPT" status + [[ "$output" == *"(no jobs)"* ]] + [ "$status" -eq 0 ] + + "$SCRIPT" start later -- sleep 3 + run "$SCRIPT" status later + [[ "$output" == *later* ]] + [[ "$output" == *RUN* || "$output" == *PENDING* ]] + [ "$status" -eq 1 ] # unfinished: a supervisor must be able to branch on this +} + +@test "two jobs in flight keep separate state" { + "$SCRIPT" start a -- bash -c 'echo A; exit 0' + "$SCRIPT" start b -- bash -c 'echo B; exit 2' + until_done a && until_done b + run "$SCRIPT" status + [[ "$output" == *"a"*DONE*rc=0* ]] + [[ "$output" == *"b"*DONE*rc=2* ]] + run "$SCRIPT" out a + [[ "$output" == A ]] + run "$SCRIPT" out b + [[ "$output" == B ]] +} + +@test "out honours a line limit" { + "$SCRIPT" start many -- bash -c 'seq 1 40' + until_done many + run "$SCRIPT" out many 3 + [ "$(printf '%s\n' "$output" | grep -c .)" -eq 3 ] + [[ "$output" == *38* ]] # the tail, not the head +} + +@test "an unknown verb is a usage error" { + run "$SCRIPT" frobnicate + [ "$status" -eq 2 ] + [[ "$output" == *usage* ]] +} + +@test "the job directory is created on first use" { + [ ! -d "$REMOTE_JOBS_DIR" ] + "$SCRIPT" status + [ -d "$REMOTE_JOBS_DIR" ] +}