diff --git a/skills/hrr-replay-findings/SKILL.md b/skills/hrr-replay-findings/SKILL.md new file mode 100644 index 0000000..b40a71e --- /dev/null +++ b/skills/hrr-replay-findings/SKILL.md @@ -0,0 +1,103 @@ +--- +name: hrr-replay-findings +description: >- + Replay and analyze HIP Runtime Replay (HRR) capture archives. Classifies GPU + faults and replay divergence (read-only/OOB write, illegal access, NaN/Inf, + hang, OOM) and extracts fault address, event index, and kernel. Use when the + user provides an HRR archive, asks to replay a recording, triage a GPU crash, + or mentions hrr-playback, capture.hrr, or Memory access fault during replay. +--- + +# HRR Replay Findings + +Replay an HRR archive on the host GPU, then produce a structured finding. + +## What the user should say + +The user only needs to point at the recording. Examples: + +- *"Replay and analyze this HRR archive: `/data/capture.hrr/pid-1842`"* +- *"What's in this capture? `capture.hrr/pid-1842`"* +- *"Analyze this replay log from an HRR run"* (log-only path) + +The user should **not** need to name scripts, set env vars, pick a GPU, or know where ROCm is installed. + +## What to ask the user (only if missing) + +| Missing | Ask once | +|---------|----------| +| Archive path | *"Which `capture.hrr/pid-*` directory should I use?"* | +| `hrr-playback` not found after discovery | *"Where is `hrr-playback` installed on this machine?"* | + +Do **not** ask for: GPU index, Docker, source trees, HIP library paths, ROCm install path (assume `/opt/rocm`). + +## Agent workflow + +``` +1. Resolve archive — user path, or largest events.bin under capture.hrr/pid-* +2. Discover hrr-playback (see below); ask user only if not found +3. Run skill scripts/run_hrr_replay.sh --archive --analyze +4. Read .finding.md and explain in plain language +``` + +**Execute in the same turn** — do not narrate planning steps. + +### Discover `hrr-playback` (in order) + +1. `command -v hrr-playback` +2. `$ROCM_PATH/bin/hrr-playback` (default `ROCM_PATH=/opt/rocm`) +3. `/opt/rocm/bin/hrr-playback` +4. User-provided path → set `HRR_PLAYBACK` for that run only + +`run_hrr_replay.sh` adds `/opt/rocm/lib` and a sibling `lib/` next to the playback binary to `LD_LIBRARY_PATH` automatically. + +### Discover archive + +If the user gives `capture.hrr/` without a `pid-*` child, pick the `pid-*` directory whose `events.bin` is largest. + +### Commands (agent runs these — not the user) + +```bash +SKILL=/hrr-replay-findings # from .cursor/skills or installed skills dir + +# Metadata only (~seconds, no full replay): +"$SKILL/scripts/run_hrr_replay.sh" --archive --info + +# Full replay + structured finding: +"$SKILL/scripts/run_hrr_replay.sh" --archive --analyze +``` + +Log-only (no replay): + +```bash +python3 "$SKILL/scripts/analyze_replay_finding.py" \ + --log --archive --format markdown -o finding.md +``` + +## System assumptions + +| Assumption | Default | +|------------|---------| +| ROCm install | `/opt/rocm` (`$ROCM_PATH` override) | +| GPU | Auto-pick device with most free VRAM | +| Replay mode | Native host (`/dev/kfd`); no Docker | + +## Fault taxonomy + +| `fault_class` | Meaning | +|---------------|---------| +| `replay_pass` | Clean replay | +| `read_only_page_fault` | Write to read-only page | +| `illegal_memory_access` | Other GPU memory fault | +| `nan_inf_divergence` | D2H numerical mismatch | +| `hang` | Device/queue hang | +| `replay_oom` | Out of VRAM — report environment issue | +| `replay_fatal_api` | HIP API error stopped replay | + +## Capture explainer (short) + +- **events.bin** — recorded HIP API sequence +- **blobs/** — code objects and payloads +- **Trailer** — absent when the original run crashed; reader recovers complete events + +See [reference.md](reference.md) and [examples.md](examples.md). diff --git a/skills/hrr-replay-findings/examples.md b/skills/hrr-replay-findings/examples.md new file mode 100644 index 0000000..8af9a52 --- /dev/null +++ b/skills/hrr-replay-findings/examples.md @@ -0,0 +1,32 @@ +# Examples + +## What the user says + +**Full replay + analysis** + +> Replay and analyze my HRR archive at `/data/crash/capture.hrr/pid-1842` + +**Archive summary only** + +> Summarize this HRR capture: `capture.hrr/pid-1842` + +**Existing log** + +> Analyze this HRR replay log: `replay.log` (archive is `capture.hrr/pid-1842`) + +The user does not mention scripts, `HRR_PLAYBACK`, or GPU numbers. + +## What the agent does + +1. Finds `hrr-playback` on `PATH` or `/opt/rocm/bin/hrr-playback` +2. If missing, asks: *"Where is hrr-playback installed?"* +3. Runs `run_hrr_replay.sh --archive ... --analyze` +4. Presents the finding + +## If hrr-playback is not in a standard location + +User answers: *"It's in `/opt/amd-hrr/bin/hrr-playback`"* + +Agent sets `HRR_PLAYBACK=/opt/amd-hrr/bin/hrr-playback` for that session only and re-runs. + +If a `lib/` directory sits beside `bin/`, the runner picks it up automatically. diff --git a/skills/hrr-replay-findings/reference.md b/skills/hrr-replay-findings/reference.md new file mode 100644 index 0000000..5cbc4e0 --- /dev/null +++ b/skills/hrr-replay-findings/reference.md @@ -0,0 +1,123 @@ +# HRR capture reference + +## Directory layout + +``` +capture.hrr/ + pid-/ + events.bin # primary event stream (may be GB-scale) + blobs/ # code objects, graphs, memcpy payloads + (optional metadata files) +``` + +Pick the `pid-*` directory with the **largest `events.bin`** for the faulting process. + +## events.bin record model (conceptual) + +Each event has: + +| Field | Role | +|-------|------| +| Thread id | Capturing host thread | +| Sequence / event index | Monotonic call index in replay | +| API id | HIP API (malloc, launch, memcpy, sync, …) | +| Payload | API-specific bytes (variable-length for kernel launches) | + +**Kernel launch payload** includes: stream, kernel name, code-object hash, grid, block, shared memory, **kernarg blob** (pointer table + struct args), optional D2H snapshot descriptors. + +## Completeness markers + +| Signal | Meaning | +|--------|---------| +| `Complete: YES` (`--info`) | Clean shutdown trailer present | +| `recovered N events` | Crash capture; trailer missing; reader kept all complete records | +| `Torn trailing record` | Last record partial; preceding events valid | + +On recent ROCm builds with crash-resilient HRR capture, crash captures are **expected** to lack a trailer and still be replayable (reader reports `recovered N events`). + +## Replay log lines + +### Progress + +``` +[HRR progress] elapsed_s=... seq=13118764 kernels=797227 d2h_pass=4303 d2h_fail=0 ... +``` + +- `seq` — last replayed event sequence number (use as **failing_call_index** proxy when fault follows) +- `kernels` — kernel launch count so far +- `d2h_*` — device-to-host validation counters + +### GPU memory fault (ROCr) + +``` +Memory access fault by GPU node-N (Agent handle: 0x...) on address 0xADDR. Reason: ... +:0:rocdevice.cpp:NNNN: Memory Fault Error [..., faulting addr: 0xADDR, kernel: Cijk_...] +``` + +Extract: **fault_address**, **kernel_name**, **gpu_node**, **fault_reason**. + +### Hang analysis block + +``` +Dispatch Header = 0x..., grid=[...], workgroup=[...], kernarg_address=0x..., kernel_obj=0x... +``` + +Extract: **kernarg_address**, **grid**, **workgroup** — ties fault to launch packet. + +### Fatal API abort + +``` +[HRR] Fatal: T146 Event 9268 (hipMalloc) returned 2 (out of memory) — aborting replay +``` + +Extract: **failing_thread**, **failing_call_index**, **failing_api**. + +### Suballoc fidelity (optional playback feature) + +``` +[HRR] SUBALLOC OOB: kernel arg[10] rec 0x... resolves inside a captured segment but in no active tensor block +``` + +High count on `arg[10]` with `d2h_fail=0` and later MAF → likely **stale/OOB device pointer** in kernarg, not host numerics. + +## Tensile / hipBLASLt kernel name cheat sheet + +Example: + +``` +Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x192x128_..._SK3_..._WS64_WG16_16_1 +``` + +| Token | Meaning | +|-------|---------| +| `Cijk_*` | Contraction GEMM family | +| `MT128x192x128` | Macro-tile dimensions | +| `SK3` | StreamK variant | +| `WS64` | Workspace-related sizing hint | +| `Bias_HA` | Bias + HPA layout flags | + +**read_only_page_fault** on StreamK GEMM → investigate edge tile / workspace (`AddressWS`) / output (`AddressD`) stores. + +## Playback build fidelity + +The same archive can **PASS** on one `hrr-playback` build and fault on another. Record which playback binary produced each result. + +## ROCm layout (assumed) + +| Path | Role | +|------|------| +| `/opt/rocm/bin/hrr-playback` | Default playback tool location | +| `/opt/rocm/lib` | HIP/HSA and ROCm runtime libraries | +| `$ROCM_PATH` | Override prefix if ROCm is not under `/opt/rocm` | + +`run_hrr_replay.sh` prepends these to `LD_LIBRARY_PATH`. If `hrr-playback` lives in `/bin/`, `/lib` is added automatically. + +## Parser script + +```bash +python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py --help +``` + +Outputs JSON or Markdown `Finding` with fields: + +`outcome`, `fault_class`, `fault_address`, `failing_event_seq`, `failing_call_index`, `kernel_name`, `kernarg_address`, `d2h_fail`, `archive_events`, … diff --git a/skills/hrr-replay-findings/scripts/analyze_replay_finding.py b/skills/hrr-replay-findings/scripts/analyze_replay_finding.py new file mode 100755 index 0000000..cafb9d5 --- /dev/null +++ b/skills/hrr-replay-findings/scripts/analyze_replay_finding.py @@ -0,0 +1,346 @@ +#!/usr/bin/env python3 +"""Parse HRR replay/capture logs into a structured finding (read-only).""" +from __future__ import annotations + +import argparse +import json +import re +import subprocess +import sys +from dataclasses import asdict, dataclass, field +from pathlib import Path +from typing import Any + +# --- regex library (diverse workloads) --- + +RE_PROGRESS = re.compile( + r"\[HRR progress\].*seq=(\d+).*kernels=(\d+).*d2h_pass=(\d+).*" + r"d2h_fail=(\d+).*d2h_attempted=(\d+).*last=\"([^\"]+)\"" +) +RE_FATAL_EVENT = re.compile( + r"\[HRR\] Fatal: T(\d+) Event (\d+) \(([^)]+)\) returned (\d+) \(([^)]+)\)" +) +RE_FATAL_GPU = re.compile( + r"\[HRR\] Fatal: GPU error after T(\d+) Event (\d+) \(([^)]+)\): (\d+) \(([^)]+)\)" +) +RE_FATAL_GENERIC = re.compile(r"\[HRR\] Fatal: ([^\n]+)") +RE_MAF = re.compile( + r"Memory access fault by GPU node-(\d+).*on address (0x[0-9a-fA-F]+)\.\s*" + r"Reason:\s*([^.\n]+)" +) +RE_MEM_FAULT_ERR = re.compile( + r"Memory Fault Error \[host: [^,]+, GPU index: \d+, faulting addr: (0x[0-9a-fA-F]+), " + r"kernel: ([^\]]+)\]" +) +RE_HANG = re.compile(r"HSA_STATUS_ERROR_(MEMORY_FAULT|ABORTED|EXCEPTION)") +RE_PASS = re.compile(r"\[HRR\] PASS\b") +RE_FAIL = re.compile(r"\[HRR\] FAIL\b") +RE_ARCHIVE_RECOVERED = re.compile( + r"recovered (\d+) events|Archive : (\d+) events, (\d+) kernels, (\d+) blobs, (\d+) code objects" +) +RE_ARCHIVE_COMPLETE = re.compile(r"Complete:\s+(YES|NO)") +RE_CAPTURE_MAF = RE_MAF +RE_SUBALLOC_OOB = re.compile( + r"\[HRR\] SUBALLOC OOB: kernel arg\[(\d+)\] rec (0x[0-9a-fA-F]+)" +) +RE_D2H_SUMMARY = re.compile( + r"D2H checks\s+: (\d+) pass.*?, (\d+) fail, (\d+) skipped" +) +RE_KERNARG = re.compile(r"kernarg_address=(0x[0-9a-fA-F]+)") +RE_GRID = re.compile(r"grid=\[([^\]]+)\], workgroup=\[([^\]]+)\]") +RE_CIJK = re.compile(r"(Cijk_[A-Za-z0-9_]+)") +RE_CAPTURE_HIP = re.compile(r"\[capture\] HIP_SO=(\S+)") + + +@dataclass +class Finding: + outcome: str + fault_class: str + fault_address: str | None = None + fault_reason: str | None = None + failing_event_seq: int | None = None + failing_call_index: int | None = None + failing_thread: int | None = None + failing_api: str | None = None + kernel_name: str | None = None + kernel_family: str | None = None + kernarg_address: str | None = None + grid: str | None = None + workgroup: str | None = None + gpu_node: str | None = None + last_progress_kernel: str | None = None + kernels_launched: int | None = None + d2h_pass: int | None = None + d2h_fail: int | None = None + d2h_attempted: int | None = None + suballoc_oob_count: int = 0 + suballoc_oob_args: list[int] = field(default_factory=list) + archive_events: int | None = None + archive_kernels: int | None = None + archive_complete: str | None = None + capture_hip_so: str | None = None + sources: list[str] = field(default_factory=list) + notes: list[str] = field(default_factory=list) + + def to_dict(self) -> dict[str, Any]: + return asdict(self) + + +def _classify(text: str, finding: Finding) -> str: + if RE_PASS.search(text): + if finding.d2h_fail and finding.d2h_fail > 0: + return "nan_inf_divergence" + return "replay_pass" + if "out of memory" in text.lower() or "hipErrorOutOfMemory" in text: + return "replay_oom" + if RE_FATAL_EVENT.search(text) or RE_FATAL_GPU.search(text) or RE_FATAL_GENERIC.search(text): + if "out of memory" in text.lower(): + return "replay_oom" + return "replay_fatal_api" + if RE_MAF.search(text) or RE_MEM_FAULT_ERR.search(text): + reason = (finding.fault_reason or "").lower() + if "read-only" in reason: + return "read_only_page_fault" + return "illegal_memory_access" + if RE_HANG.search(text) and not RE_PASS.search(text): + return "hang" + if RE_FAIL.search(text) or (finding.d2h_fail and finding.d2h_fail > 0): + return "nan_inf_divergence" + if "Replay aborted" in text or "aborting replay" in text: + return "replay_aborted" + return "unknown" + + +def _kernel_family(name: str | None) -> str | None: + if not name: + return None + if name.startswith("Cijk_"): + m = re.search(r"_MT(\d+x\d+x\d+)", name) + sk = "_SK3_" if "_SK3_" in name else ("_SK2_" if "_SK2_" in name else None) + parts = ["hipblaslt_gemm"] + if m: + parts.append(f"MT{m.group(1)}") + if sk: + parts.append("streamk" if "SK3" in sk else "streamk_variant") + return "/".join(parts) + if name.startswith("_ZN"): + return "pytorch_kernel" + return "other" + + +def parse_text(text: str, source: str, finding: Finding) -> Finding: + finding.sources.append(source) + + for m in RE_CAPTURE_HIP.finditer(text): + finding.capture_hip_so = m.group(1) + + for m in RE_ARCHIVE_RECOVERED.finditer(text): + g = m.groups() + if g[0]: + finding.archive_events = int(g[0]) + if len(g) >= 5 and g[1]: + finding.archive_events = int(g[1]) + finding.archive_kernels = int(g[2]) + + m = RE_ARCHIVE_COMPLETE.search(text) + if m: + finding.archive_complete = m.group(1) + + oob_args: set[int] = set() + for m in RE_SUBALLOC_OOB.finditer(text): + finding.suballoc_oob_count += 1 + oob_args.add(int(m.group(1))) + finding.suballoc_oob_args = sorted(oob_args) + + last_prog = None + for m in RE_PROGRESS.finditer(text): + finding.failing_event_seq = int(m.group(1)) + finding.kernels_launched = int(m.group(2)) + finding.d2h_pass = int(m.group(3)) + finding.d2h_fail = int(m.group(4)) + finding.d2h_attempted = int(m.group(5)) + last_prog = m.group(6) + finding.last_progress_kernel = last_prog + + for m in (RE_FATAL_EVENT, RE_FATAL_GPU): + hit = m.search(text) + if hit: + finding.failing_thread = int(hit.group(1)) + finding.failing_call_index = int(hit.group(2)) + finding.failing_api = hit.group(3) + break + + m = RE_MAF.search(text) + if m: + finding.gpu_node = m.group(1) + finding.fault_address = m.group(2) + finding.fault_reason = m.group(3).strip() + + m = RE_MEM_FAULT_ERR.search(text) + if m: + finding.fault_address = finding.fault_address or m.group(1) + finding.kernel_name = m.group(2).strip() + + if not finding.kernel_name: + cijk = RE_CIJK.search(text) + if cijk: + finding.kernel_name = cijk.group(1) + + m = RE_KERNARG.search(text) + if m: + finding.kernarg_address = m.group(1) + + m = RE_GRID.search(text) + if m: + finding.grid = m.group(1) + finding.workgroup = m.group(2) + + m = RE_D2H_SUMMARY.search(text) + if m: + finding.d2h_pass = int(m.group(1)) + finding.d2h_fail = int(m.group(2)) + + if RE_PASS.search(text): + finding.outcome = "PASS" + elif RE_MAF.search(text) or RE_MEM_FAULT_ERR.search(text): + finding.outcome = "MAF" + elif RE_FAIL.search(text): + finding.outcome = "FAIL" + elif "aborting replay" in text or RE_FATAL_EVENT.search(text): + finding.outcome = "ABORT" + else: + finding.outcome = "UNKNOWN" + + finding.fault_class = _classify(text, finding) + finding.kernel_family = _kernel_family(finding.kernel_name) + return finding + + +def run_archive_info(archive: Path, hrr_playback: str | None) -> str: + play = hrr_playback or "hrr-playback" + try: + proc = subprocess.run( + [play, str(archive), "--info"], + capture_output=True, + text=True, + timeout=120, + check=False, + ) + return proc.stdout + proc.stderr + except FileNotFoundError: + return "" + except subprocess.TimeoutExpired: + return "[timeout running hrr-playback --info]" + + +def parse_sweep_tsv(path: Path) -> list[dict[str, Any]]: + rows: list[dict[str, Any]] = [] + lines = path.read_text(encoding="utf-8", errors="replace").splitlines() + if not lines: + return rows + header = lines[0].split("\t") + for line in lines[1:]: + if not line.strip(): + continue + cols = line.split("\t") + rows.append(dict(zip(header, cols))) + return rows + + +def render_markdown(f: Finding, sweep: list[dict[str, Any]] | None = None) -> str: + lines = [ + "# HRR replay finding", + "", + "## Summary", + f"- **Outcome**: {f.outcome}", + f"- **Fault class**: `{f.fault_class}`", + f"- **Kernel**: `{f.kernel_name or 'unknown'}`", + f"- **Kernel family**: `{f.kernel_family or 'unknown'}`", + "", + "## Fault details", + f"- **Fault address**: `{f.fault_address or 'n/a'}`", + f"- **Fault reason**: {f.fault_reason or 'n/a'}", + f"- **Failing event seq**: {f.failing_event_seq or 'n/a'}", + f"- **Failing call index**: {f.failing_call_index or 'n/a'}", + f"- **Failing API**: {f.failing_api or 'n/a'}", + f"- **Kernarg address**: `{f.kernarg_address or 'n/a'}`", + f"- **GPU node**: {f.gpu_node or 'n/a'}", + f"- **Grid / workgroup**: {f.grid or 'n/a'} / {f.workgroup or 'n/a'}", + "", + "## Replay progress at fault", + f"- **Kernels launched**: {f.kernels_launched or 'n/a'}", + f"- **D2H**: pass={f.d2h_pass or 0} fail={f.d2h_fail or 0} attempted={f.d2h_attempted or 0}", + f"- **Last progress kernel**: `{f.last_progress_kernel or 'n/a'}`", + "", + "## Archive / capture", + f"- **Events**: {f.archive_events or 'n/a'}", + f"- **Kernels (archive)**: {f.archive_kernels or 'n/a'}", + f"- **Complete**: {f.archive_complete or 'n/a'}", + f"- **Capture HIP**: `{f.capture_hip_so or 'n/a'}`", + f"- **Suballoc OOB reports**: {f.suballoc_oob_count} (args: {f.suballoc_oob_args or []})", + "", + "## Sources", + ] + for s in f.sources: + lines.append(f"- `{s}`") + if f.notes: + lines.extend(["", "## Notes"]) + lines.extend(f"- {n}" for n in f.notes) + if sweep: + lines.extend(["", "## Multi-run sweep"]) + lines.append("| run | gpu | outcome | fault_addr |") + lines.append("|-----|-----|---------|------------|") + for r in sweep: + lines.append( + f"| {r.get('run','')} | {r.get('gpu','')} | {r.get('outcome','')} | {r.get('fault_addr','')} |" + ) + return "\n".join(lines) + "\n" + + +def main() -> int: + ap = argparse.ArgumentParser(description=__doc__) + ap.add_argument("--log", action="append", default=[], help="Replay or capture log (repeatable)") + ap.add_argument("--archive", help="HRR archive pid-* directory for --info") + ap.add_argument("--sweep-tsv", help="multi-replay sweep summary TSV") + ap.add_argument("--hrr-playback", help="Path to hrr-playback binary") + ap.add_argument("--format", choices=("json", "markdown"), default="markdown") + ap.add_argument("-o", "--output", help="Write report to file") + args = ap.parse_args() + + if not args.log and not args.archive and not args.sweep_tsv: + ap.error("provide --log, --archive, and/or --sweep-tsv") + + finding = Finding(outcome="UNKNOWN", fault_class="unknown") + for log_path in args.log: + p = Path(log_path) + if not p.is_file(): + finding.notes.append(f"log not found: {p}") + continue + parse_text(p.read_text(encoding="utf-8", errors="replace"), str(p), finding) + + if args.archive: + arch = Path(args.archive) + info = run_archive_info(arch, args.hrr_playback) + if info: + parse_text(info, f"{arch} (--info)", finding) + else: + finding.notes.append("hrr-playback --info unavailable; archive path recorded only") + finding.sources.append(str(arch)) + + sweep = parse_sweep_tsv(Path(args.sweep_tsv)) if args.sweep_tsv else None + if sweep: + finding.notes.append(f"multi-replay sweep: {len(sweep)} runs") + + out = ( + json.dumps(finding.to_dict(), indent=2) + if args.format == "json" + else render_markdown(finding, sweep) + ) + if args.output: + Path(args.output).write_text(out, encoding="utf-8") + print(out) + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/skills/hrr-replay-findings/scripts/run_hrr_replay.sh b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh new file mode 100755 index 0000000..8987128 --- /dev/null +++ b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh @@ -0,0 +1,137 @@ +#!/usr/bin/env bash +# Run HRR replay on the host GPU and optionally analyze the log. +set -euo pipefail + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +ANALYZER="$SCRIPT_DIR/analyze_replay_finding.py" +ROCM_PATH="${ROCM_PATH:-/opt/rocm}" + +ARCHIVE="" +LOG="" +DO_ANALYZE=0 +DO_INFO=0 +EXTRA_ARGS=() + +while [[ $# -gt 0 ]]; do + case "$1" in + --archive) ARCHIVE="$2"; shift 2 ;; + --log) LOG="$2"; shift 2 ;; + --analyze) DO_ANALYZE=1; shift ;; + --info) DO_INFO=1; shift ;; + --) shift; EXTRA_ARGS+=("$@"); break ;; + *) EXTRA_ARGS+=("$1"); shift ;; + esac +done + +[[ -n "$ARCHIVE" ]] || { echo "error: --archive required" >&2; exit 1; } +ARCHIVE="$(readlink -f "$ARCHIVE" 2>/dev/null || realpath "$ARCHIVE" 2>/dev/null || echo "$ARCHIVE")" +[[ -d "$ARCHIVE" ]] || { echo "error: archive not found: $ARCHIVE" >&2; exit 1; } + +resolve_playback() { + local c candidates=() + [[ -n "${HRR_PLAYBACK:-}" ]] && candidates+=("$HRR_PLAYBACK") + if command -v hrr-playback >/dev/null 2>&1; then + candidates+=("$(command -v hrr-playback)") + fi + candidates+=( + "$ROCM_PATH/bin/hrr-playback" + "/opt/rocm/bin/hrr-playback" + ) + local p + for p in "${candidates[@]}"; do + [[ -n "$p" && -x "$p" ]] || continue + echo "$p" + return + done + echo "" +} + +setup_library_path() { + local play="$1" + local bin_dir lib_dir paths=() + bin_dir="$(cd "$(dirname "$play")" && pwd)" + lib_dir="$(cd "$bin_dir/.." && pwd)/lib" + paths+=("$ROCM_PATH/lib" "/opt/rocm/lib") + [[ -d "$lib_dir" ]] && paths+=("$lib_dir") + [[ -d "$bin_dir/../lib" ]] && paths+=("$(cd "$bin_dir/../lib" && pwd)") + local seen="" p + for p in "${paths[@]}"; do + [[ -d "$p" ]] || continue + [[ ":$seen:" == *":$p:"* ]] && continue + seen="${seen:+$seen:}$p" + LD_LIBRARY_PATH="${p}${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" + done + export LD_LIBRARY_PATH +} + +pick_gpu() { + if [[ -n "${GPU:-}" ]]; then + echo "$GPU" + return + fi + if command -v rocm-smi >/dev/null 2>&1; then + local best="" best_free=-1 idx free + while read -r idx free; do + [[ -n "$idx" ]] || continue + if (( free > best_free )); then + best_free=$free + best=$idx + fi + done < <(rocm-smi --showmeminfo vram 2>/dev/null | awk ' + /GPU\[/ { gsub(/[^0-9]/,"",$1); idx=$1 } + /Used Memory/ { used=$NF } + /Total Memory/ { total=$NF; if (idx!="") { print idx, total-used; idx="" } } + ') + if [[ -n "$best" ]]; then + echo "[run_hrr_replay] auto-selected GPU $best (most free VRAM)" >&2 + echo "$best" + return + fi + fi + echo "[run_hrr_replay] default GPU 0" >&2 + echo "0" +} + +HRR_PLAY="$(resolve_playback)" +GPU="$(pick_gpu)" + +if [[ -z "$HRR_PLAY" ]]; then + echo "error: hrr-playback not found. Checked PATH, \$ROCM_PATH/bin ($ROCM_PATH), /opt/rocm/bin." >&2 + echo "error: Ask the user where hrr-playback is installed, then set HRR_PLAYBACK for this run." >&2 + exit 1 +fi + +setup_library_path "$HRR_PLAY" +echo "[run_hrr_replay] playback=$HRR_PLAY" >&2 + +if [[ "$DO_INFO" == "1" ]]; then + exec "$HRR_PLAY" "$ARCHIVE" --info "${EXTRA_ARGS[@]}" +fi + +[[ -r /dev/kfd ]] || { + echo "error: /dev/kfd not accessible — AMD GPU driver required" >&2 + exit 1 +} + +if [[ -z "$LOG" ]]; then + LOG="hrr-replay-$(basename "$ARCHIVE")-$(date -u +%Y%m%dT%H%M%SZ).log" +fi +mkdir -p "$(dirname "$LOG")" 2>/dev/null || true + +echo "[run_hrr_replay] GPU=$GPU archive=$ARCHIVE" >&2 +set +e +ROCR_VISIBLE_DEVICES="$GPU" "$HRR_PLAY" "$ARCHIVE" "${EXTRA_ARGS[@]}" 2>&1 | tee "$LOG" +RC=${PIPESTATUS[0]} +set -e + +echo "[run_hrr_replay] log=$LOG exit=$RC" + +if [[ "$DO_ANALYZE" == "1" ]]; then + FINDING="${LOG%.log}.finding.md" + python3 "$ANALYZER" --log "$LOG" --archive "$ARCHIVE" \ + --hrr-playback "$HRR_PLAY" \ + --format markdown -o "$FINDING" + echo "[run_hrr_replay] finding=$FINDING" +fi + +exit "$RC" diff --git a/skills/hrr-replay-findings/skill-card.md b/skills/hrr-replay-findings/skill-card.md new file mode 100644 index 0000000..d9e5d4d --- /dev/null +++ b/skills/hrr-replay-findings/skill-card.md @@ -0,0 +1,13 @@ +# Skill Card + +## Description + +Read-only triage of HIP Runtime Replay (HRR) captures and replay logs. Produces a structured finding: fault class (OOB/read-only write, illegal access, NaN/Inf divergence, hang, OOM), fault address, failing event/call index, implicated kernel and kernarg, plus a capture layout explainer. Works across vLLM, PyTorch, and hipBLASLt workloads. + +## Owner + +AMD + +## License + +MIT