Ship the PCIe probe instruments with a guard that refuses to contend with a service
The RDNA4 offload campaign left three HIP probes and a driver sitting loose in $HOME, where the next agent cannot find them and a stray `rm -rf` ends them. They answer a recurring question on the R9700 box — is decode limited by the link or by the engine's read path — so they move into scripts/probes/ with a driver, tests, and a page. Two things changed in the code, not just the address: - The driver measures only after a guard that runs BEFORE anything is built: gpu_busy_percent over PCIE_PROBE_MAX_BUSY, or any process holding /dev/kfd, is a refusal (exit 3) rather than a measurement. A benchmark taken beside a live endpoint is both disruptive and worthless, and silently unloading a production service is not a tool's decision — --force is the human saying otherwise. The refusal-before-build ordering is pinned by a test, because the sysfs and kfd paths are overridable and a refusal is otherwise untestable off the card box. - hip_smoke.hip and pcie_probe.hip called hipGetDeviceProperties, which SEGFAULTS on this ROCm build even though hipGetDeviceCount works; only pcie_bw.hip avoided it, and that is the one that produced the campaign's numbers. Device identity now comes from the sysfs topology block the driver prints, so all three probes are runnable on the build we ship. Verified here: refusal on a real busy APU (exit 3), missing toolchain (exit 4), missing sources (exit 2). Compile-and-run still needs the box with hipcc. Refs #53
This commit is contained in:
Frank committed
1 parent
35ceae6d7b
commit
56fc899763
8 files changed
+573
No files matched your search
Executable
+126
@@ -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
|
||||
@@ -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;
|
||||
}
|
||||
@@ -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;
|
||||
}
|
||||
@@ -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;
|
||||
}
|
||||
Reference in new issue
Block a user