Merge pull request 'Rescue the detached job runner and the PCIe probe instruments from $HOME sediment' (#54) from feat/53-fleet-probes into main

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

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

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.

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 && <the health probe that endpoint answers>'
+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).

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    (<repo>/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

scripts/probes/hip_smoke.hip

@@ -0,0 +1,36 @@
+// Minimal HIP smoke test: is the host ROCm usable from a plain process?
+#include <cstdio>
+#include <cstring>
+#include <hip/hip_runtime.h>
+
+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;
+}

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 <cstdio>
+#include <cstring>
+#include <cstdlib>
+#include <hip/hip_runtime.h>
+
+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;
+}

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 <cstdio>
+#include <cstring>
+#include <cstdlib>
+#include <hip/hip_runtime.h>
+
+#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;
+}

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

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 ]
+}

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" ]
+}