From 11435ff6d7cc5f932e9a6d12905de2876edaad6d Mon Sep 17 00:00:00 2001 From: AMD HRR Team Date: Wed, 15 Jul 2026 15:28:21 +0000 Subject: [PATCH 1/6] Add hrr-replay-findings skill for read-only HRR replay triage. Structured parser classifies replay faults (read-only/OOB, illegal access, NaN/Inf divergence, hang, OOM) and extracts fault address, event index, kernel, and kernarg from logs and archives across diverse GPU workloads. --- skills/hrr-replay-findings/SKILL.md | 120 ++++++ skills/hrr-replay-findings/examples.md | 70 ++++ skills/hrr-replay-findings/reference.md | 113 ++++++ .../scripts/analyze_replay_finding.py | 346 ++++++++++++++++++ skills/hrr-replay-findings/skill-card.md | 13 + 5 files changed, 662 insertions(+) create mode 100644 skills/hrr-replay-findings/SKILL.md create mode 100644 skills/hrr-replay-findings/examples.md create mode 100644 skills/hrr-replay-findings/reference.md create mode 100755 skills/hrr-replay-findings/scripts/analyze_replay_finding.py create mode 100644 skills/hrr-replay-findings/skill-card.md diff --git a/skills/hrr-replay-findings/SKILL.md b/skills/hrr-replay-findings/SKILL.md new file mode 100644 index 0000000..f84826f --- /dev/null +++ b/skills/hrr-replay-findings/SKILL.md @@ -0,0 +1,120 @@ +--- +name: hrr-replay-findings +description: >- + Read-only analysis of HIP Runtime Replay (HRR) recordings and replay logs. + Classifies replay divergence or GPU faults (OOB/read-only write, illegal access, + NaN/Inf D2H divergence, hang, OOM), extracts fault address, failing event/call + index, and implicated kernel/allocation. Summarizes what an HRR capture holds. + Use when given an HRR archive (capture.hrr/pid-*), hrr-playback log, capture log, + multi-replay sweep TSV, Memory access fault output, or when triaging replay PASS + vs MAF across workloads (vLLM, PyTorch finetune, hipBLASLt StreamK). +--- + +# HRR Replay Findings (read-only) + +Turn **replay output + optional archive metadata** into a **structured finding**. +This skill is **read-only** — it does not patch, rebuild, or re-run capture unless +the user explicitly asks in a follow-up. + +## When to use + +- Replay or capture log shows `[HRR] PASS`, `Memory access fault`, `Fatal:`, or `d2h_fail` +- Comparing replay outcomes across GPUs, runs, or playback builds +- First triage step **before** end-to-end crash analysis (patch/build/validate workflow) + +## Quick workflow + +``` +1. Identify inputs (archive, replay log, capture log, sweep TSV) +2. Run analyze_replay_finding.py (below) +3. Present the structured finding + capture explainer (reference.md) +4. Only if user asks: reproduce, bisect, or patch +``` + +## Step 1 — Gather inputs + +| Input | How to find | +|-------|-------------| +| **Archive** | Largest `events.bin` under `capture.hrr/pid-*` | +| **Replay log** | `hrr-playback` stdout/stderr (saved to a file) | +| **Capture log** | Application log from the captured run; optional `[capture] HIP_SO=` line | +| **Sweep TSV** | Tab-separated summary from multiple replay runs (`*.summary.tsv`) | + +## Step 2 — Run the parser (default) + +```bash +SCRIPT=skills/hrr-replay-findings/scripts/analyze_replay_finding.py + +python3 "$SCRIPT" \ + --log /path/to/replay.log \ + --log /path/to/capture.log \ + --archive /path/to/capture.hrr/pid-NNN \ + --hrr-playback /path/to/hrr-playback \ + --format markdown \ + -o finding.md +``` + +Optional multi-run table: + +```bash +python3 "$SCRIPT" \ + --log replay-gpu0.log --log replay-gpu1.log \ + --sweep-tsv replay-sweep.summary.tsv \ + --format json +``` + +## Step 3 — Report template + +Use the parser output and add **one paragraph** of interpretation: + +```markdown +## Interpretation + +- Fault class `read_only_page_fault` on `Cijk_*_SK3_*` → likely hipBLASLt StreamK edge/OOB write. +- `d2h_fail=0` at MAF → replay diverged on-GPU before any host numerical check failed. +- Archive incomplete (no EOF trailer) → common for crash captures with event recovery enabled. +``` + +## Fault taxonomy + +| `fault_class` | Meaning | Typical log signals | +|---------------|---------|---------------------| +| `replay_pass` | Replay finished cleanly | `[HRR] PASS`, `d2h_fail=0` | +| `read_only_page_fault` | Write to non-writable page | `Reason: Write access to a read-only page` | +| `illegal_memory_access` | Other GPU memory fault | `Memory access fault`, `MEMORY_FAULT` | +| `nan_inf_divergence` | Host saw bad numerics | `d2h_fail>0`, `[HRR] FAIL` | +| `hang` | Queue/device hang without clean summary | `HSA_STATUS_ERROR_*`, no PASS/MAF line | +| `replay_oom` | Replay ran out of VRAM | `out of memory`, `hipMalloc` + `Fatal:` | +| `replay_fatal_api` | HIP API error aborted replay | `[HRR] Fatal: T* Event *` | +| `replay_aborted` | Aborted without classified fault | `aborting replay` | +| `unknown` | Insufficient signals | — | + +## Capture explainer (short) + +An HRR **capture** is a time-ordered **event stream** plus **sidecar blobs**: + +- **events.bin** — HIP API calls (alloc, memcpy, kernel launch, sync), kernel names, kernarg payloads, optional D2H snapshots +- **blobs/** — code objects, graph capture data, memcpy payloads +- **Trailer** — clean shutdown marker when capture exits normally; **absent** when the app crashed (reader recovers complete events) + +Replay **re-executes** API sequence on a live GPU, remapping recorded device pointers to new allocations. Divergence means replay state ≠ capture state at some event index — not necessarily that capture missed data. + +Full layout: [reference.md](reference.md). + +## Cross-workload notes + +- **hipBLASLt / Tensile `Cijk_*`**: decode `MTxx`, `_SK3_` = StreamK; see reference.md +- **PyTorch `_*` kernels**: last `[HRR progress] last=` demangled name is the active kernel before fault +- **Suballoc OOB lines**: `SUBALLOC OOB: kernel arg[N]` — pointer in segment but outside active tensor block +- **D2H checks**: `d2h_pass` / `d2h_fail` — host numerical validation; independent of on-GPU MAF + +## Do not + +- Guess ROCm commit, hipBLASLt tag, or root cause without evidence +- Treat `PASS` on one archive as proof all archives pass on the same build +- Modify archives (use `hrr-playback --repair` only when user requests) + +## Additional resources + +- Capture layout and log patterns: [reference.md](reference.md) +- Worked examples: [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..c6bed09 --- /dev/null +++ b/skills/hrr-replay-findings/examples.md @@ -0,0 +1,70 @@ +# Examples + +## Example 1 — GPU memory fault during replay (StreamK GEMM) + +**Input:** `replay.log` + archive `capture.hrr/pid-12345` + +**Command:** + +```bash +python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py \ + --log replay.log \ + --archive capture.hrr/pid-12345 \ + --format markdown +``` + +**Expected finding (abridged):** + +| Field | Value | +|-------|-------| +| outcome | MAF | +| fault_class | read_only_page_fault | +| kernel_name | Cijk_..._MT128x192x128_..._SK3_... | +| fault_address | 0x7b9224c09000 | +| failing_event_seq | (from last `[HRR progress]` line) | +| d2h_fail | 0 | + +**Interpretation:** On-GPU fault at a hipBLASLt StreamK GEMM; `d2h_fail=0` means host numerical checks had not failed yet — not a NaN/Inf divergence. + +--- + +## Example 2 — clean replay pass + +**Input:** `replay-pass.log` + +| Field | Value | +|-------|-------| +| outcome | PASS | +| fault_class | replay_pass | +| kernels_launched | (from replay summary) | +| d2h_fail | 0 | + +**Interpretation:** Replay completed without GPU fault or D2H mismatch. + +--- + +## Example 3 — replay OOM (insufficient VRAM) + +**Input:** `replay-oom.log` + +| Field | Value | +|-------|-------| +| outcome | ABORT | +| fault_class | replay_oom | +| failing_call_index | (from `Fatal: T* Event *`) | +| failing_api | hipMalloc | + +**Interpretation:** Replay aborted for lack of device memory — free VRAM or reduce conflicting workloads before attributing to capture fidelity. + +--- + +## Example 4 — multi-run sweep summary + +```bash +python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py \ + --log replay-gpu0.log \ + --sweep-tsv replay-sweep.summary.tsv \ + --format markdown +``` + +Use when several replays of the same archive were run (e.g. across GPUs). Consistent `fault_class` and kernel across runs suggests deterministic replay divergence, not GPU-specific hardware variation. diff --git a/skills/hrr-replay-findings/reference.md b/skills/hrr-replay-findings/reference.md new file mode 100644 index 0000000..f1e6f73 --- /dev/null +++ b/skills/hrr-replay-findings/reference.md @@ -0,0 +1,113 @@ +# 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` / `libamdhip64` build and **MAF** on another. When comparing runs, record the HIP library used at capture time (if logged) and the playback binary used at replay time. Do not assume one outcome generalizes across builds without evidence. + +## 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/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 From 4360b8d761dd36d56103da94c0b1e452020b7d0f Mon Sep 17 00:00:00 2001 From: AMD HRR Team Date: Wed, 15 Jul 2026 15:38:50 +0000 Subject: [PATCH 2/6] Extend hrr-replay-findings to run replay, not only parse logs. Add run_hrr_replay.sh (native or Docker) with optional --analyze step; update SKILL workflow to archive-first replay then structured finding. --- skills/hrr-replay-findings/SKILL.md | 138 +++++++++++------- skills/hrr-replay-findings/examples.md | 26 +++- .../scripts/run_hrr_replay.sh | 122 ++++++++++++++++ 3 files changed, 222 insertions(+), 64 deletions(-) create mode 100755 skills/hrr-replay-findings/scripts/run_hrr_replay.sh diff --git a/skills/hrr-replay-findings/SKILL.md b/skills/hrr-replay-findings/SKILL.md index f84826f..e14d547 100644 --- a/skills/hrr-replay-findings/SKILL.md +++ b/skills/hrr-replay-findings/SKILL.md @@ -1,78 +1,102 @@ --- name: hrr-replay-findings description: >- - Read-only analysis of HIP Runtime Replay (HRR) recordings and replay logs. + Run HIP Runtime Replay (HRR) against a capture archive and analyze the result. Classifies replay divergence or GPU faults (OOB/read-only write, illegal access, NaN/Inf D2H divergence, hang, OOM), extracts fault address, failing event/call index, and implicated kernel/allocation. Summarizes what an HRR capture holds. - Use when given an HRR archive (capture.hrr/pid-*), hrr-playback log, capture log, - multi-replay sweep TSV, Memory access fault output, or when triaging replay PASS - vs MAF across workloads (vLLM, PyTorch finetune, hipBLASLt StreamK). + Use when given an HRR archive (capture.hrr/pid-*), when asked to replay a + recording, or when triaging replay PASS vs MAF across workloads (vLLM, PyTorch, + hipBLASLt StreamK). --- -# HRR Replay Findings (read-only) +# HRR Replay Findings -Turn **replay output + optional archive metadata** into a **structured finding**. -This skill is **read-only** — it does not patch, rebuild, or re-run capture unless -the user explicitly asks in a follow-up. +**Run replay** on an HRR archive, then produce a **structured finding** from the output. +Analysis is read-only (no patches); replay needs a GPU and a built `hrr-playback`. ## When to use -- Replay or capture log shows `[HRR] PASS`, `Memory access fault`, `Fatal:`, or `d2h_fail` -- Comparing replay outcomes across GPUs, runs, or playback builds -- First triage step **before** end-to-end crash analysis (patch/build/validate workflow) +- User provides an HRR archive (`capture.hrr/pid-*`) and wants replay + triage +- Compare replay across GPUs or playback builds +- First step before deeper crash root-cause work ## Quick workflow ``` -1. Identify inputs (archive, replay log, capture log, sweep TSV) -2. Run analyze_replay_finding.py (below) -3. Present the structured finding + capture explainer (reference.md) -4. Only if user asks: reproduce, bisect, or patch +1. Resolve archive (largest events.bin under capture.hrr/pid-*) +2. Run replay → save log (run_hrr_replay.sh) +3. Parse log → structured finding (analyze_replay_finding.py) +4. Interpret + recommend next steps ``` -## Step 1 — Gather inputs +## Step 0 — Prerequisites -| Input | How to find | -|-------|-------------| -| **Archive** | Largest `events.bin` under `capture.hrr/pid-*` | -| **Replay log** | `hrr-playback` stdout/stderr (saved to a file) | -| **Capture log** | Application log from the captured run; optional `[capture] HIP_SO=` line | -| **Sweep TSV** | Tab-separated summary from multiple replay runs (`*.summary.tsv`) | +| Requirement | How to verify | +|-------------|---------------| +| `hrr-playback` built | `HRR_PLAYBACK` points to executable, or on `PATH` | +| GPU free (~full VRAM for large captures) | `rocm-smi --showmeminfo vram` | +| Docker replay (recommended) | `HIP_SO` + `HSA_SO` from CLR/rocr build; `IMAGE` env | -## Step 2 — Run the parser (default) +**This workspace** (if `scripts/maf-hrr-docker-playback.sh` exists): set `HRR_REPO_ROOT` to the repo root and the skill runner delegates to that script. + +## Step 1 — Archive metadata (no GPU) ```bash -SCRIPT=skills/hrr-replay-findings/scripts/analyze_replay_finding.py +SKILL=skills/hrr-replay-findings # or .cursor/skills/hrr-replay-findings in a project -python3 "$SCRIPT" \ - --log /path/to/replay.log \ - --log /path/to/capture.log \ +export HRR_PLAYBACK=/path/to/hrr-playback +./$SKILL/scripts/run_hrr_replay.sh \ --archive /path/to/capture.hrr/pid-NNN \ - --hrr-playback /path/to/hrr-playback \ - --format markdown \ - -o finding.md + --info +``` + +## Step 2 — Run replay + save log + +```bash +export HRR_REPO_ROOT=/path/to/hrr-repo # optional: use repo docker helper +export HRR_PLAYBACK=/path/to/hrr-playback +export HIP_SO=/path/to/libamdhip64.so.7.* +export HSA_SO=/path/to/libhsa-runtime64.so.1 +export GPU=1 + +./$SKILL/scripts/run_hrr_replay.sh \ + --archive /path/to/capture.hrr/pid-NNN \ + --log /path/to/replay.log \ + --analyze ``` -Optional multi-run table: +`--analyze` runs the parser and writes `replay.finding.md` next to the log. + +**Native replay** (no Docker): `HRR_REPLAY_MODE=native` or omit `HIP_SO` with GPU device visible. + +**Extra hrr-playback flags** pass after `--`: ```bash -python3 "$SCRIPT" \ - --log replay-gpu0.log --log replay-gpu1.log \ - --sweep-tsv replay-sweep.summary.tsv \ - --format json +./$SKILL/scripts/run_hrr_replay.sh --archive ... --log replay.log -- --single-thread ``` -## Step 3 — Report template +## Step 3 — Analyze an existing log only + +If replay already ran: + +```bash +python3 $SKILL/scripts/analyze_replay_finding.py \ + --log replay.log \ + --archive capture.hrr/pid-NNN \ + --hrr-playback "$HRR_PLAYBACK" \ + --format markdown \ + -o finding.md +``` -Use the parser output and add **one paragraph** of interpretation: +## Step 4 — Report template ```markdown ## Interpretation - Fault class `read_only_page_fault` on `Cijk_*_SK3_*` → likely hipBLASLt StreamK edge/OOB write. -- `d2h_fail=0` at MAF → replay diverged on-GPU before any host numerical check failed. -- Archive incomplete (no EOF trailer) → common for crash captures with event recovery enabled. +- `d2h_fail=0` at MAF → on-GPU divergence before host numerical mismatch. +- Archive incomplete (no EOF trailer) → common for crash captures with event recovery. ``` ## Fault taxonomy @@ -83,7 +107,7 @@ Use the parser output and add **one paragraph** of interpretation: | `read_only_page_fault` | Write to non-writable page | `Reason: Write access to a read-only page` | | `illegal_memory_access` | Other GPU memory fault | `Memory access fault`, `MEMORY_FAULT` | | `nan_inf_divergence` | Host saw bad numerics | `d2h_fail>0`, `[HRR] FAIL` | -| `hang` | Queue/device hang without clean summary | `HSA_STATUS_ERROR_*`, no PASS/MAF line | +| `hang` | Queue/device hang | `HSA_STATUS_ERROR_*`, no PASS/MAF | | `replay_oom` | Replay ran out of VRAM | `out of memory`, `hipMalloc` + `Fatal:` | | `replay_fatal_api` | HIP API error aborted replay | `[HRR] Fatal: T* Event *` | | `replay_aborted` | Aborted without classified fault | `aborting replay` | @@ -91,30 +115,32 @@ Use the parser output and add **one paragraph** of interpretation: ## Capture explainer (short) -An HRR **capture** is a time-ordered **event stream** plus **sidecar blobs**: +- **events.bin** — HIP API stream (alloc, memcpy, launches, sync), kernargs, optional D2H snapshots +- **blobs/** — code objects, graphs, memcpy payloads +- **Trailer** — present on clean exit; absent on crash (reader recovers complete events) + +Replay remaps recorded device pointers to live allocations. Divergence = replay state ≠ capture at some event index. -- **events.bin** — HIP API calls (alloc, memcpy, kernel launch, sync), kernel names, kernarg payloads, optional D2H snapshots -- **blobs/** — code objects, graph capture data, memcpy payloads -- **Trailer** — clean shutdown marker when capture exits normally; **absent** when the app crashed (reader recovers complete events) +See [reference.md](reference.md). -Replay **re-executes** API sequence on a live GPU, remapping recorded device pointers to new allocations. Divergence means replay state ≠ capture state at some event index — not necessarily that capture missed data. +## Agent instructions -Full layout: [reference.md](reference.md). +When the user gives an **archive only**: -## Cross-workload notes +1. Run `--info` if `hrr-playback` is available +2. Check GPU VRAM; free device or pick another `GPU=` if OOM risk +3. Run `run_hrr_replay.sh --archive ... --log ... --analyze` +4. Present finding + interpretation; do not guess root cause without evidence -- **hipBLASLt / Tensile `Cijk_*`**: decode `MTxx`, `_SK3_` = StreamK; see reference.md -- **PyTorch `_*` kernels**: last `[HRR progress] last=` demangled name is the active kernel before fault -- **Suballoc OOB lines**: `SUBALLOC OOB: kernel arg[N]` — pointer in segment but outside active tensor block -- **D2H checks**: `d2h_pass` / `d2h_fail` — host numerical validation; independent of on-GPU MAF +When the user gives a **log only**: skip replay; run `analyze_replay_finding.py`. ## Do not -- Guess ROCm commit, hipBLASLt tag, or root cause without evidence -- Treat `PASS` on one archive as proof all archives pass on the same build -- Modify archives (use `hrr-playback --repair` only when user requests) +- Guess ROCm commit or root cause without evidence +- Modify archives unless user requests `--repair` +- Assume one PASS archive means all archives pass ## Additional resources -- Capture layout and log patterns: [reference.md](reference.md) -- Worked examples: [examples.md](examples.md) +- [reference.md](reference.md) — log patterns, kernel name decode +- [examples.md](examples.md) — sample commands and findings diff --git a/skills/hrr-replay-findings/examples.md b/skills/hrr-replay-findings/examples.md index c6bed09..e2f6b69 100644 --- a/skills/hrr-replay-findings/examples.md +++ b/skills/hrr-replay-findings/examples.md @@ -1,10 +1,24 @@ # Examples -## Example 1 — GPU memory fault during replay (StreamK GEMM) +## Example 1 — Run replay then analyze -**Input:** `replay.log` + archive `capture.hrr/pid-12345` +```bash +export HRR_PLAYBACK=/path/to/hrr-playback +export HIP_SO=/path/to/libamdhip64.so.7.* +export HSA_SO=/path/to/libhsa-runtime64.so.1 +export GPU=0 + +skills/hrr-replay-findings/scripts/run_hrr_replay.sh \ + --archive capture.hrr/pid-12345 \ + --log replay.log \ + --analyze +``` -**Command:** +Produces `replay.log` and `replay.finding.md`. + +--- + +## Example 2 — Analyze log only (replay already ran) ```bash python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py \ @@ -20,15 +34,11 @@ python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py \ | outcome | MAF | | fault_class | read_only_page_fault | | kernel_name | Cijk_..._MT128x192x128_..._SK3_... | -| fault_address | 0x7b9224c09000 | -| failing_event_seq | (from last `[HRR progress]` line) | | d2h_fail | 0 | -**Interpretation:** On-GPU fault at a hipBLASLt StreamK GEMM; `d2h_fail=0` means host numerical checks had not failed yet — not a NaN/Inf divergence. - --- -## Example 2 — clean replay pass +## Example 3 — clean replay pass **Input:** `replay-pass.log` 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..909a2a0 --- /dev/null +++ b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh @@ -0,0 +1,122 @@ +#!/usr/bin/env bash +# Run HRR replay against an archive and optionally analyze the log. +# Portable wrapper: native hrr-playback or Docker (ROCm/vLLM image + HIP inject). +# +# Usage: +# GPU=1 ./run_hrr_replay.sh --archive capture.hrr/pid-123 --log replay.log +# ./run_hrr_replay.sh --archive capture.hrr/pid-123 --info +# ./run_hrr_replay.sh --archive capture.hrr/pid-123 --analyze +# +# Env: +# HRR_PLAYBACK path to hrr-playback (required unless on PATH) +# HIP_SO, HSA_SO for Docker inject mode (libamdhip64 + libhsa) +# GPU device index (default 0) +# IMAGE Docker image (default rocm/vllm rocm7.13 gfx950) +# HRR_REPLAY_MODE auto|native|docker (default auto) +# HRR_REPO_ROOT if set and contains scripts/maf-hrr-docker-playback.sh, use it +set -euo pipefail + +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +ANALYZER="$SCRIPT_DIR/analyze_replay_finding.py" + +ARCHIVE="" +LOG="" +DO_ANALYZE=0 +DO_INFO=0 +MODE="${HRR_REPLAY_MODE:-auto}" +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 ;; + --native) MODE=native; shift ;; + --docker) MODE=docker; 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() { + if [[ -n "${HRR_PLAYBACK:-}" && -x "$HRR_PLAYBACK" ]]; then + echo "$HRR_PLAYBACK" + return + fi + if command -v hrr-playback >/dev/null 2>&1; then + command -v hrr-playback + return + fi + echo "" +} + +HRR_PLAY="$(resolve_playback)" +GPU="${GPU:-0}" + +if [[ "$DO_INFO" == "1" ]]; then + [[ -n "$HRR_PLAY" ]] || { echo "error: hrr-playback not found; set HRR_PLAYBACK" >&2; exit 1; } + exec "$HRR_PLAY" "$ARCHIVE" --info "${EXTRA_ARGS[@]}" +fi + +if [[ -z "$LOG" ]]; then + LOG="hrr-replay-$(basename "$ARCHIVE")-gpu${GPU}-$(date -u +%Y%m%dT%H%M%SZ).log" +fi +mkdir -p "$(dirname "$LOG")" 2>/dev/null || true + +# Prefer workspace docker helper when available (full inject recipe). +if [[ -n "${HRR_REPO_ROOT:-}" && -f "$HRR_REPO_ROOT/scripts/maf-hrr-docker-playback.sh" ]]; then + echo "[run_hrr_replay] using $HRR_REPO_ROOT/scripts/maf-hrr-docker-playback.sh" + sudo docker rm -f "${HRR_NAME:-maf-hrr-playback}" 2>/dev/null || true + set +e + sudo -E GPU="$GPU" bash "$HRR_REPO_ROOT/scripts/maf-hrr-docker-playback.sh" \ + "$ARCHIVE" "${EXTRA_ARGS[@]}" 2>&1 | tee "$LOG" + RC=${PIPESTATUS[0]} + set -e +elif [[ "$MODE" == "docker" || ( "$MODE" == "auto" && -n "${HIP_SO:-}" ) ]]; then + [[ -n "$HRR_PLAY" ]] || { echo "error: hrr-playback required for docker mode; set HRR_PLAYBACK" >&2; exit 1; } + IMAGE="${IMAGE:-rocm/vllm:rocm7.13.0_gfx950-dcgpu_ubuntu24.04_py3.13_pytorch_2.10.0_vllm_0.19.1}" + HIP_DST=/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib/libamdhip64.so.7 + HSA_DST=/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib/libhsa-runtime64.so.1 + ROC_LIB=/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib + CORES_HOST="${CORES_HOST:-/tmp/hrr-cores}" + mkdir -p "$CORES_HOST" + docker_env=( -e ROCR_VISIBLE_DEVICES="$GPU" -e HRR_INJECT=1 -e LD_LIBRARY_PATH="$ROC_LIB:/opt/rocm/lib" ) + vols=( -v "$ARCHIVE":/data/hrr-archive:ro -v "$HRR_PLAY":/opt/hrr-tools/hrr-playback:ro -v "$CORES_HOST":/data/cores ) + [[ -f "${HIP_SO:-}" ]] && vols+=( -v "$HIP_SO":"$HIP_DST":ro ) + [[ -f "${HSA_SO:-}" ]] && vols+=( -v "$HSA_SO":"$HSA_DST":ro ) + echo "[run_hrr_replay] docker image=$IMAGE GPU=$GPU archive=$ARCHIVE" + sudo docker rm -f "${HRR_NAME:-maf-hrr-playback}" 2>/dev/null || true + set +e + sudo -E docker run --rm --privileged --init \ + --name "${HRR_NAME:-maf-hrr-playback}" \ + --device /dev/kfd -v /dev/dri:/dev/dri --shm-size=4g \ + --security-opt seccomp=unconfined --ulimit core=-1:-1 \ + "${docker_env[@]}" "${vols[@]}" \ + "$IMAGE" /opt/hrr-tools/hrr-playback /data/hrr-archive "${EXTRA_ARGS[@]}" \ + 2>&1 | tee "$LOG" + RC=${PIPESTATUS[0]} + set -e +else + [[ -n "$HRR_PLAY" ]] || { echo "error: hrr-playback not found; set HRR_PLAYBACK or use --docker with HIP_SO" >&2; exit 1; } + echo "[run_hrr_replay] native GPU=$GPU archive=$ARCHIVE" + set +e + ROCR_VISIBLE_DEVICES="$GPU" "$HRR_PLAY" "$ARCHIVE" "${EXTRA_ARGS[@]}" 2>&1 | tee "$LOG" + RC=${PIPESTATUS[0]} + set -e +fi + +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" From 1524754c8cd44550d187af30614093c2e8c89668 Mon Sep 17 00:00:00 2001 From: AMD HRR Team Date: Wed, 15 Jul 2026 15:41:54 +0000 Subject: [PATCH 3/6] Make hrr-replay-findings customer-facing: no source or GPU setup. Remove HRR_REPO_ROOT and internal lab script deps; auto-pick GPU by free VRAM; default docker replay uses stock container ROCm without HIP_SO paths. --- skills/hrr-replay-findings/SKILL.md | 158 ++++++------------ skills/hrr-replay-findings/examples.md | 60 ++----- skills/hrr-replay-findings/reference.md | 6 +- .../scripts/run_hrr_replay.sh | 138 ++++++++++----- 4 files changed, 170 insertions(+), 192 deletions(-) diff --git a/skills/hrr-replay-findings/SKILL.md b/skills/hrr-replay-findings/SKILL.md index e14d547..5263a06 100644 --- a/skills/hrr-replay-findings/SKILL.md +++ b/skills/hrr-replay-findings/SKILL.md @@ -1,146 +1,98 @@ --- name: hrr-replay-findings description: >- - Run HIP Runtime Replay (HRR) against a capture archive and analyze the result. - Classifies replay divergence or GPU faults (OOB/read-only write, illegal access, - NaN/Inf D2H divergence, hang, OOM), extracts fault address, failing event/call - index, and implicated kernel/allocation. Summarizes what an HRR capture holds. - Use when given an HRR archive (capture.hrr/pid-*), when asked to replay a - recording, or when triaging replay PASS vs MAF across workloads (vLLM, PyTorch, - hipBLASLt StreamK). + Run HIP Runtime Replay (HRR) against a customer capture archive and analyze + the result. Classifies replay faults (OOB/read-only write, illegal access, + NaN/Inf divergence, hang, OOM) and extracts fault address, failing event index, + and implicated kernel. Requires only the archive and a shipped hrr-playback + binary — no source checkout or GPU configuration knowledge. Use when given an + HRR archive (capture.hrr/pid-*), a replay request, or Memory access fault output. --- # HRR Replay Findings -**Run replay** on an HRR archive, then produce a **structured finding** from the output. -Analysis is read-only (no patches); replay needs a GPU and a built `hrr-playback`. +Run replay on an HRR **archive**, then emit a **structured finding**. The user +supplies the recording; the agent handles playback discovery, GPU selection, and +analysis. -## When to use +## What the user provides -- User provides an HRR archive (`capture.hrr/pid-*`) and wants replay + triage -- Compare replay across GPUs or playback builds -- First step before deeper crash root-cause work +| Input | Required? | +|-------|-----------| +| HRR archive (`capture.hrr/pid-*` with `events.bin`) | **Yes** | +| `hrr-playback` binary (on `PATH` or path via `HRR_PLAYBACK`) | **Yes** for replay | +| Source code, ROCm build tree, GPU index | **No** | -## Quick workflow +Optional: container image (`IMAGE`) if not using the default ROCm/vLLM image. -``` -1. Resolve archive (largest events.bin under capture.hrr/pid-*) -2. Run replay → save log (run_hrr_replay.sh) -3. Parse log → structured finding (analyze_replay_finding.py) -4. Interpret + recommend next steps -``` - -## Step 0 — Prerequisites - -| Requirement | How to verify | -|-------------|---------------| -| `hrr-playback` built | `HRR_PLAYBACK` points to executable, or on `PATH` | -| GPU free (~full VRAM for large captures) | `rocm-smi --showmeminfo vram` | -| Docker replay (recommended) | `HIP_SO` + `HSA_SO` from CLR/rocr build; `IMAGE` env | +## Quick workflow (agent) -**This workspace** (if `scripts/maf-hrr-docker-playback.sh` exists): set `HRR_REPO_ROOT` to the repo root and the skill runner delegates to that script. - -## Step 1 — Archive metadata (no GPU) - -```bash -SKILL=skills/hrr-replay-findings # or .cursor/skills/hrr-replay-findings in a project - -export HRR_PLAYBACK=/path/to/hrr-playback -./$SKILL/scripts/run_hrr_replay.sh \ - --archive /path/to/capture.hrr/pid-NNN \ - --info +``` +1. Find archive — largest events.bin under capture.hrr/pid-* +2. Find hrr-playback — PATH, HRR_PLAYBACK, or ask user once if missing +3. run_hrr_replay.sh --archive ... --analyze (GPU auto-selected) +4. Present finding + plain-language interpretation ``` -## Step 2 — Run replay + save log +## Run replay + analyze ```bash -export HRR_REPO_ROOT=/path/to/hrr-repo # optional: use repo docker helper -export HRR_PLAYBACK=/path/to/hrr-playback -export HIP_SO=/path/to/libamdhip64.so.7.* -export HSA_SO=/path/to/libhsa-runtime64.so.1 -export GPU=1 +SKILL=skills/hrr-replay-findings +# Only set HRR_PLAYBACK if hrr-playback is not already on PATH ./$SKILL/scripts/run_hrr_replay.sh \ --archive /path/to/capture.hrr/pid-NNN \ - --log /path/to/replay.log \ --analyze ``` -`--analyze` runs the parser and writes `replay.finding.md` next to the log. - -**Native replay** (no Docker): `HRR_REPLAY_MODE=native` or omit `HIP_SO` with GPU device visible. +Writes `hrr-replay-pid-NNN-.log` and `.finding.md` in the current directory. -**Extra hrr-playback flags** pass after `--`: +Archive metadata only (no GPU): ```bash -./$SKILL/scripts/run_hrr_replay.sh --archive ... --log replay.log -- --single-thread +./$SKILL/scripts/run_hrr_replay.sh --archive /path/to/capture.hrr/pid-NNN --info ``` -## Step 3 — Analyze an existing log only +## Analyze an existing log -If replay already ran: +If replay already ran elsewhere: ```bash python3 $SKILL/scripts/analyze_replay_finding.py \ --log replay.log \ - --archive capture.hrr/pid-NNN \ - --hrr-playback "$HRR_PLAYBACK" \ - --format markdown \ - -o finding.md + --archive /path/to/capture.hrr/pid-NNN \ + --format markdown -o finding.md ``` -## Step 4 — Report template +## Agent instructions -```markdown -## Interpretation +**Do not** ask the user for GPU index, source paths, or HIP library paths unless replay fails and support provided a custom library bundle. -- Fault class `read_only_page_fault` on `Cijk_*_SK3_*` → likely hipBLASLt StreamK edge/OOB write. -- `d2h_fail=0` at MAF → on-GPU divergence before host numerical mismatch. -- Archive incomplete (no EOF trailer) → common for crash captures with event recovery. -``` +1. Locate archive (`pid-*` with largest `events.bin`). +2. Resolve `hrr-playback` (`command -v hrr-playback` or `HRR_PLAYBACK`). +3. Run `run_hrr_replay.sh --archive ... --analyze`. +4. If `replay_oom`: auto-retry on another GPU (script picks most free VRAM) or report insufficient VRAM. +5. Present finding; do not invent root cause without log evidence. + +**Do not** require checkout of TheRock, CLR, or internal lab scripts. ## Fault taxonomy -| `fault_class` | Meaning | Typical log signals | -|---------------|---------|---------------------| -| `replay_pass` | Replay finished cleanly | `[HRR] PASS`, `d2h_fail=0` | -| `read_only_page_fault` | Write to non-writable page | `Reason: Write access to a read-only page` | -| `illegal_memory_access` | Other GPU memory fault | `Memory access fault`, `MEMORY_FAULT` | -| `nan_inf_divergence` | Host saw bad numerics | `d2h_fail>0`, `[HRR] FAIL` | -| `hang` | Queue/device hang | `HSA_STATUS_ERROR_*`, no PASS/MAF | -| `replay_oom` | Replay ran out of VRAM | `out of memory`, `hipMalloc` + `Fatal:` | -| `replay_fatal_api` | HIP API error aborted replay | `[HRR] Fatal: T* Event *` | -| `replay_aborted` | Aborted without classified fault | `aborting replay` | -| `unknown` | Insufficient signals | — | +| `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 during replay | +| `replay_fatal_api` | HIP API error stopped replay | +| `unknown` | Insufficient log data | ## Capture explainer (short) -- **events.bin** — HIP API stream (alloc, memcpy, launches, sync), kernargs, optional D2H snapshots -- **blobs/** — code objects, graphs, memcpy payloads -- **Trailer** — present on clean exit; absent on crash (reader recovers complete events) - -Replay remaps recorded device pointers to live allocations. Divergence = replay state ≠ capture at some event index. - -See [reference.md](reference.md). - -## Agent instructions - -When the user gives an **archive only**: - -1. Run `--info` if `hrr-playback` is available -2. Check GPU VRAM; free device or pick another `GPU=` if OOM risk -3. Run `run_hrr_replay.sh --archive ... --log ... --analyze` -4. Present finding + interpretation; do not guess root cause without evidence - -When the user gives a **log only**: skip replay; run `analyze_replay_finding.py`. - -## Do not - -- Guess ROCm commit or root cause without evidence -- Modify archives unless user requests `--repair` -- Assume one PASS archive means all archives pass - -## Additional resources +- **events.bin** — recorded HIP API sequence (alloc, memcpy, kernel launches, sync) +- **blobs/** — code objects and sidecar payloads +- **Trailer** — missing when the original run crashed; reader still recovers complete events -- [reference.md](reference.md) — log patterns, kernel name decode -- [examples.md](examples.md) — sample commands and findings +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 index e2f6b69..36a6f2c 100644 --- a/skills/hrr-replay-findings/examples.md +++ b/skills/hrr-replay-findings/examples.md @@ -1,80 +1,50 @@ # Examples -## Example 1 — Run replay then analyze +## Example 1 — Archive only (typical customer handoff) + +Customer sends `capture.hrr/pid-1842/`. Support ships matching `hrr-playback`. ```bash -export HRR_PLAYBACK=/path/to/hrr-playback -export HIP_SO=/path/to/libamdhip64.so.7.* -export HSA_SO=/path/to/libhsa-runtime64.so.1 -export GPU=0 +export HRR_PLAYBACK=/opt/rocm/bin/hrr-playback # only if not on PATH skills/hrr-replay-findings/scripts/run_hrr_replay.sh \ - --archive capture.hrr/pid-12345 \ - --log replay.log \ + --archive capture.hrr/pid-1842 \ --analyze ``` -Produces `replay.log` and `replay.finding.md`. +No GPU index, no source tree, no HIP paths. --- -## Example 2 — Analyze log only (replay already ran) +## Example 2 — Analyze log only ```bash python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py \ --log replay.log \ - --archive capture.hrr/pid-12345 \ + --archive capture.hrr/pid-1842 \ --format markdown ``` -**Expected finding (abridged):** - -| Field | Value | -|-------|-------| -| outcome | MAF | -| fault_class | read_only_page_fault | -| kernel_name | Cijk_..._MT128x192x128_..._SK3_... | -| d2h_fail | 0 | - --- -## Example 3 — clean replay pass - -**Input:** `replay-pass.log` +## Example 3 — Expected MAF finding (abridged) | Field | Value | |-------|-------| -| outcome | PASS | -| fault_class | replay_pass | -| kernels_launched | (from replay summary) | +| outcome | MAF | +| fault_class | read_only_page_fault | +| kernel_name | Cijk_..._SK3_... | | d2h_fail | 0 | -**Interpretation:** Replay completed without GPU fault or D2H mismatch. +`d2h_fail=0` → fault happened on GPU before host numerical checks failed. --- -## Example 3 — replay OOM (insufficient VRAM) - -**Input:** `replay-oom.log` +## Example 4 — OOM (environment, not capture bug) | Field | Value | |-------|-------| -| outcome | ABORT | | fault_class | replay_oom | -| failing_call_index | (from `Fatal: T* Event *`) | | failing_api | hipMalloc | -**Interpretation:** Replay aborted for lack of device memory — free VRAM or reduce conflicting workloads before attributing to capture fidelity. - ---- - -## Example 4 — multi-run sweep summary - -```bash -python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py \ - --log replay-gpu0.log \ - --sweep-tsv replay-sweep.summary.tsv \ - --format markdown -``` - -Use when several replays of the same archive were run (e.g. across GPUs). Consistent `fault_class` and kernel across runs suggests deterministic replay divergence, not GPU-specific hardware variation. +Retry after freeing GPU memory; script auto-picks the GPU with most free VRAM. diff --git a/skills/hrr-replay-findings/reference.md b/skills/hrr-replay-findings/reference.md index f1e6f73..86eef22 100644 --- a/skills/hrr-replay-findings/reference.md +++ b/skills/hrr-replay-findings/reference.md @@ -100,7 +100,11 @@ Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x192x128_..._SK3_..._WS64_WG16 ## Playback build fidelity -The same archive can **PASS** on one `hrr-playback` / `libamdhip64` build and **MAF** on another. When comparing runs, record the HIP library used at capture time (if logged) and the playback binary used at replay time. Do not assume one outcome generalizes across builds without evidence. +The same archive can **PASS** on one `hrr-playback` build and fault on another. Record which playback binary (and container image, if used) produced each result. Do not assume one outcome generalizes. + +### Advanced (support bundles only) + +If AMD shipped matched `libamdhip64` / `libhsa` libraries with the playback tool, set `HIP_SO` and `HSA_SO` when running docker replay. **Customers are not expected to build or locate these themselves.** ## Parser script diff --git a/skills/hrr-replay-findings/scripts/run_hrr_replay.sh b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh index 909a2a0..d3a5304 100755 --- a/skills/hrr-replay-findings/scripts/run_hrr_replay.sh +++ b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh @@ -1,19 +1,19 @@ #!/usr/bin/env bash # Run HRR replay against an archive and optionally analyze the log. -# Portable wrapper: native hrr-playback or Docker (ROCm/vLLM image + HIP inject). +# +# Customer-facing: needs only an archive directory and hrr-playback (on PATH or +# HRR_PLAYBACK). No source tree, no GPU index, no HIP library paths required. # # Usage: -# GPU=1 ./run_hrr_replay.sh --archive capture.hrr/pid-123 --log replay.log -# ./run_hrr_replay.sh --archive capture.hrr/pid-123 --info # ./run_hrr_replay.sh --archive capture.hrr/pid-123 --analyze +# ./run_hrr_replay.sh --archive capture.hrr/pid-123 --info # -# Env: -# HRR_PLAYBACK path to hrr-playback (required unless on PATH) -# HIP_SO, HSA_SO for Docker inject mode (libamdhip64 + libhsa) -# GPU device index (default 0) -# IMAGE Docker image (default rocm/vllm rocm7.13 gfx950) -# HRR_REPLAY_MODE auto|native|docker (default auto) -# HRR_REPO_ROOT if set and contains scripts/maf-hrr-docker-playback.sh, use it +# Optional env (agent may set; user does not need to know): +# HRR_PLAYBACK path to hrr-playback if not on PATH +# GPU force device index; otherwise auto-pick most free VRAM +# IMAGE container for docker replay (default: ROCm/vLLM image) +# HRR_REPLAY_MODE auto|native|docker +# HIP_SO, HSA_SO advanced: override container HIP/HSA libs (support bundles only) set -euo pipefail SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" @@ -55,66 +55,118 @@ resolve_playback() { echo "" } +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 + 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 (set GPU=N to override)" >&2 + echo "0" +} + HRR_PLAY="$(resolve_playback)" -GPU="${GPU:-0}" +GPU="$(pick_gpu)" if [[ "$DO_INFO" == "1" ]]; then - [[ -n "$HRR_PLAY" ]] || { echo "error: hrr-playback not found; set HRR_PLAYBACK" >&2; exit 1; } + [[ -n "$HRR_PLAY" ]] || { + echo "error: hrr-playback not found. Install it on PATH or set HRR_PLAYBACK=/path/to/hrr-playback" >&2 + exit 1 + } exec "$HRR_PLAY" "$ARCHIVE" --info "${EXTRA_ARGS[@]}" fi if [[ -z "$LOG" ]]; then - LOG="hrr-replay-$(basename "$ARCHIVE")-gpu${GPU}-$(date -u +%Y%m%dT%H%M%SZ).log" + LOG="hrr-replay-$(basename "$ARCHIVE")-$(date -u +%Y%m%dT%H%M%SZ).log" fi mkdir -p "$(dirname "$LOG")" 2>/dev/null || true -# Prefer workspace docker helper when available (full inject recipe). -if [[ -n "${HRR_REPO_ROOT:-}" && -f "$HRR_REPO_ROOT/scripts/maf-hrr-docker-playback.sh" ]]; then - echo "[run_hrr_replay] using $HRR_REPO_ROOT/scripts/maf-hrr-docker-playback.sh" - sudo docker rm -f "${HRR_NAME:-maf-hrr-playback}" 2>/dev/null || true +use_docker=0 +if [[ "$MODE" == "docker" ]]; then + use_docker=1 +elif [[ "$MODE" == "native" ]]; then + use_docker=0 +elif command -v docker >/dev/null 2>&1 && [[ ! -r /dev/kfd ]]; then + use_docker=1 +else + use_docker=0 +fi + +run_native() { + [[ -n "$HRR_PLAY" ]] || { + echo "error: hrr-playback not found. Install on PATH or set HRR_PLAYBACK." >&2 + exit 1 + } + echo "[run_hrr_replay] native playback GPU=$GPU archive=$ARCHIVE" set +e - sudo -E GPU="$GPU" bash "$HRR_REPO_ROOT/scripts/maf-hrr-docker-playback.sh" \ - "$ARCHIVE" "${EXTRA_ARGS[@]}" 2>&1 | tee "$LOG" - RC=${PIPESTATUS[0]} - set -e -elif [[ "$MODE" == "docker" || ( "$MODE" == "auto" && -n "${HIP_SO:-}" ) ]]; then - [[ -n "$HRR_PLAY" ]] || { echo "error: hrr-playback required for docker mode; set HRR_PLAYBACK" >&2; exit 1; } + ROCR_VISIBLE_DEVICES="$GPU" "$HRR_PLAY" "$ARCHIVE" "${EXTRA_ARGS[@]}" 2>&1 | tee "$LOG" + return "${PIPESTATUS[0]}" +} + +run_docker() { + [[ -n "$HRR_PLAY" ]] || { + echo "error: hrr-playback not found. Set HRR_PLAYBACK to the shipped playback binary." >&2 + exit 1 + } + command -v docker >/dev/null 2>&1 || { echo "error: docker not found" >&2; exit 1; } IMAGE="${IMAGE:-rocm/vllm:rocm7.13.0_gfx950-dcgpu_ubuntu24.04_py3.13_pytorch_2.10.0_vllm_0.19.1}" - HIP_DST=/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib/libamdhip64.so.7 - HSA_DST=/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib/libhsa-runtime64.so.1 ROC_LIB=/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib - CORES_HOST="${CORES_HOST:-/tmp/hrr-cores}" - mkdir -p "$CORES_HOST" - docker_env=( -e ROCR_VISIBLE_DEVICES="$GPU" -e HRR_INJECT=1 -e LD_LIBRARY_PATH="$ROC_LIB:/opt/rocm/lib" ) - vols=( -v "$ARCHIVE":/data/hrr-archive:ro -v "$HRR_PLAY":/opt/hrr-tools/hrr-playback:ro -v "$CORES_HOST":/data/cores ) - [[ -f "${HIP_SO:-}" ]] && vols+=( -v "$HIP_SO":"$HIP_DST":ro ) - [[ -f "${HSA_SO:-}" ]] && vols+=( -v "$HSA_SO":"$HSA_DST":ro ) - echo "[run_hrr_replay] docker image=$IMAGE GPU=$GPU archive=$ARCHIVE" - sudo docker rm -f "${HRR_NAME:-maf-hrr-playback}" 2>/dev/null || true + docker_env=( -e ROCR_VISIBLE_DEVICES="$GPU" -e LD_LIBRARY_PATH="$ROC_LIB:/opt/rocm/lib" ) + vols=( -v "$ARCHIVE":/data/hrr-archive:ro -v "$HRR_PLAY":/opt/hrr-tools/hrr-playback:ro ) + # Stock container HIP/HSA by default — no source build or library paths required. + if [[ -f "${HIP_SO:-}" ]]; then + docker_env+=( -e HRR_INJECT=1 ) + vols+=( -v "$HIP_SO":/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib/libamdhip64.so.7:ro ) + [[ -f "${HSA_SO:-}" ]] && vols+=( -v "$HSA_SO":/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib/libhsa-runtime64.so.1:ro ) + echo "[run_hrr_replay] docker with custom HIP/HSA overlay (support bundle)" >&2 + else + echo "[run_hrr_replay] docker with stock container ROCm stack" >&2 + fi + echo "[run_hrr_replay] image=$IMAGE GPU=$GPU archive=$ARCHIVE" + sudo docker rm -f "${HRR_NAME:-hrr-replay}" 2>/dev/null || true set +e sudo -E docker run --rm --privileged --init \ - --name "${HRR_NAME:-maf-hrr-playback}" \ + --name "${HRR_NAME:-hrr-replay}" \ --device /dev/kfd -v /dev/dri:/dev/dri --shm-size=4g \ --security-opt seccomp=unconfined --ulimit core=-1:-1 \ "${docker_env[@]}" "${vols[@]}" \ "$IMAGE" /opt/hrr-tools/hrr-playback /data/hrr-archive "${EXTRA_ARGS[@]}" \ 2>&1 | tee "$LOG" - RC=${PIPESTATUS[0]} - set -e + return "${PIPESTATUS[0]}" +} + +if [[ "$use_docker" == "1" ]]; then + run_docker + RC=$? else - [[ -n "$HRR_PLAY" ]] || { echo "error: hrr-playback not found; set HRR_PLAYBACK or use --docker with HIP_SO" >&2; exit 1; } - echo "[run_hrr_replay] native GPU=$GPU archive=$ARCHIVE" - set +e - ROCR_VISIBLE_DEVICES="$GPU" "$HRR_PLAY" "$ARCHIVE" "${EXTRA_ARGS[@]}" 2>&1 | tee "$LOG" - RC=${PIPESTATUS[0]} - set -e + run_native + RC=$? fi 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" \ + python3 "$ANALYZER" --log "$LOG" --archive "$ARCHIVE" \ + ${HRR_PLAY:+--hrr-playback "$HRR_PLAY"} \ --format markdown -o "$FINDING" echo "[run_hrr_replay] finding=$FINDING" fi From 7228ee219da32b82c64b4ea8b502071a1921140d Mon Sep 17 00:00:00 2001 From: AMD HRR Team Date: Wed, 15 Jul 2026 15:51:24 +0000 Subject: [PATCH 4/6] Remove Docker from hrr-replay-findings; native GPU replay only. Replay runs hrr-playback on the host with auto GPU selection; no container image or docker dependency in the skill or run_hrr_replay.sh. --- skills/hrr-replay-findings/SKILL.md | 26 ++++-- skills/hrr-replay-findings/reference.md | 6 +- .../scripts/run_hrr_replay.sh | 92 ++++--------------- 3 files changed, 37 insertions(+), 87 deletions(-) diff --git a/skills/hrr-replay-findings/SKILL.md b/skills/hrr-replay-findings/SKILL.md index 5263a06..aa6081c 100644 --- a/skills/hrr-replay-findings/SKILL.md +++ b/skills/hrr-replay-findings/SKILL.md @@ -21,9 +21,7 @@ analysis. |-------|-----------| | HRR archive (`capture.hrr/pid-*` with `events.bin`) | **Yes** | | `hrr-playback` binary (on `PATH` or path via `HRR_PLAYBACK`) | **Yes** for replay | -| Source code, ROCm build tree, GPU index | **No** | - -Optional: container image (`IMAGE`) if not using the default ROCm/vLLM image. +| Source code, ROCm build tree, GPU index, Docker | **No** | ## Quick workflow (agent) @@ -66,15 +64,23 @@ python3 $SKILL/scripts/analyze_replay_finding.py \ ## Agent instructions -**Do not** ask the user for GPU index, source paths, or HIP library paths unless replay fails and support provided a custom library bundle. +**Execute immediately.** Do not say you will "locate the skill", "inspect expectations", or "look up documentation" — run the commands below in the same turn. + +When the user gives an **archive** (and optionally `hrr-playback` path): + +```bash +export HRR_PLAYBACK= # skip if hrr-playback is on PATH +SKILL/scripts/run_hrr_replay.sh --archive --analyze +``` + +1. If `hrr-playback` path omitted: `command -v hrr-playback` or ask **once** for the binary path. +2. Run `run_hrr_replay.sh --archive ... --analyze` (GPU is auto-selected). +3. Read the generated `.finding.md` and summarize for the user. +4. On `replay_oom`: report insufficient VRAM; do not blame the capture. -1. Locate archive (`pid-*` with largest `events.bin`). -2. Resolve `hrr-playback` (`command -v hrr-playback` or `HRR_PLAYBACK`). -3. Run `run_hrr_replay.sh --archive ... --analyze`. -4. If `replay_oom`: auto-retry on another GPU (script picks most free VRAM) or report insufficient VRAM. -5. Present finding; do not invent root cause without log evidence. +When the user gives a **log only**: run `analyze_replay_finding.py` on that log. -**Do not** require checkout of TheRock, CLR, or internal lab scripts. +**Never** require source checkout, GPU index, or HIP library paths from the user. ## Fault taxonomy diff --git a/skills/hrr-replay-findings/reference.md b/skills/hrr-replay-findings/reference.md index 86eef22..b9248f9 100644 --- a/skills/hrr-replay-findings/reference.md +++ b/skills/hrr-replay-findings/reference.md @@ -100,11 +100,7 @@ Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x192x128_..._SK3_..._WS64_WG16 ## Playback build fidelity -The same archive can **PASS** on one `hrr-playback` build and fault on another. Record which playback binary (and container image, if used) produced each result. Do not assume one outcome generalizes. - -### Advanced (support bundles only) - -If AMD shipped matched `libamdhip64` / `libhsa` libraries with the playback tool, set `HIP_SO` and `HSA_SO` when running docker replay. **Customers are not expected to build or locate these themselves.** +The same archive can **PASS** on one `hrr-playback` build and fault on another. Record which playback binary produced each result. ## Parser script diff --git a/skills/hrr-replay-findings/scripts/run_hrr_replay.sh b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh index d3a5304..f05526a 100755 --- a/skills/hrr-replay-findings/scripts/run_hrr_replay.sh +++ b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh @@ -1,19 +1,16 @@ #!/usr/bin/env bash -# Run HRR replay against an archive and optionally analyze the log. +# Run HRR replay on the host GPU and optionally analyze the log. # -# Customer-facing: needs only an archive directory and hrr-playback (on PATH or -# HRR_PLAYBACK). No source tree, no GPU index, no HIP library paths required. +# Needs: archive directory + hrr-playback (on PATH or HRR_PLAYBACK). +# No Docker, no source tree, no GPU index (auto-picked by free VRAM). # # Usage: # ./run_hrr_replay.sh --archive capture.hrr/pid-123 --analyze # ./run_hrr_replay.sh --archive capture.hrr/pid-123 --info # -# Optional env (agent may set; user does not need to know): +# Optional env: # HRR_PLAYBACK path to hrr-playback if not on PATH # GPU force device index; otherwise auto-pick most free VRAM -# IMAGE container for docker replay (default: ROCm/vLLM image) -# HRR_REPLAY_MODE auto|native|docker -# HIP_SO, HSA_SO advanced: override container HIP/HSA libs (support bundles only) set -euo pipefail SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" @@ -23,7 +20,6 @@ ARCHIVE="" LOG="" DO_ANALYZE=0 DO_INFO=0 -MODE="${HRR_REPLAY_MODE:-auto}" EXTRA_ARGS=() while [[ $# -gt 0 ]]; do @@ -32,8 +28,6 @@ while [[ $# -gt 0 ]]; do --log) LOG="$2"; shift 2 ;; --analyze) DO_ANALYZE=1; shift ;; --info) DO_INFO=1; shift ;; - --native) MODE=native; shift ;; - --docker) MODE=docker; shift ;; --) shift; EXTRA_ARGS+=("$@"); break ;; *) EXTRA_ARGS+=("$1"); shift ;; esac @@ -88,78 +82,32 @@ GPU="$(pick_gpu)" if [[ "$DO_INFO" == "1" ]]; then [[ -n "$HRR_PLAY" ]] || { - echo "error: hrr-playback not found. Install it on PATH or set HRR_PLAYBACK=/path/to/hrr-playback" >&2 + echo "error: hrr-playback not found. Set HRR_PLAYBACK=/path/to/hrr-playback" >&2 exit 1 } exec "$HRR_PLAY" "$ARCHIVE" --info "${EXTRA_ARGS[@]}" fi -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 - -use_docker=0 -if [[ "$MODE" == "docker" ]]; then - use_docker=1 -elif [[ "$MODE" == "native" ]]; then - use_docker=0 -elif command -v docker >/dev/null 2>&1 && [[ ! -r /dev/kfd ]]; then - use_docker=1 -else - use_docker=0 -fi - -run_native() { - [[ -n "$HRR_PLAY" ]] || { - echo "error: hrr-playback not found. Install on PATH or set HRR_PLAYBACK." >&2 - exit 1 - } - echo "[run_hrr_replay] native playback GPU=$GPU archive=$ARCHIVE" - set +e - ROCR_VISIBLE_DEVICES="$GPU" "$HRR_PLAY" "$ARCHIVE" "${EXTRA_ARGS[@]}" 2>&1 | tee "$LOG" - return "${PIPESTATUS[0]}" +[[ -n "$HRR_PLAY" ]] || { + echo "error: hrr-playback not found. Set HRR_PLAYBACK=/path/to/hrr-playback" >&2 + exit 1 } -run_docker() { - [[ -n "$HRR_PLAY" ]] || { - echo "error: hrr-playback not found. Set HRR_PLAYBACK to the shipped playback binary." >&2 - exit 1 - } - command -v docker >/dev/null 2>&1 || { echo "error: docker not found" >&2; exit 1; } - IMAGE="${IMAGE:-rocm/vllm:rocm7.13.0_gfx950-dcgpu_ubuntu24.04_py3.13_pytorch_2.10.0_vllm_0.19.1}" - ROC_LIB=/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib - docker_env=( -e ROCR_VISIBLE_DEVICES="$GPU" -e LD_LIBRARY_PATH="$ROC_LIB:/opt/rocm/lib" ) - vols=( -v "$ARCHIVE":/data/hrr-archive:ro -v "$HRR_PLAY":/opt/hrr-tools/hrr-playback:ro ) - # Stock container HIP/HSA by default — no source build or library paths required. - if [[ -f "${HIP_SO:-}" ]]; then - docker_env+=( -e HRR_INJECT=1 ) - vols+=( -v "$HIP_SO":/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib/libamdhip64.so.7:ro ) - [[ -f "${HSA_SO:-}" ]] && vols+=( -v "$HSA_SO":/opt/python/lib/python3.13/site-packages/_rocm_sdk_core/lib/libhsa-runtime64.so.1:ro ) - echo "[run_hrr_replay] docker with custom HIP/HSA overlay (support bundle)" >&2 - else - echo "[run_hrr_replay] docker with stock container ROCm stack" >&2 - fi - echo "[run_hrr_replay] image=$IMAGE GPU=$GPU archive=$ARCHIVE" - sudo docker rm -f "${HRR_NAME:-hrr-replay}" 2>/dev/null || true - set +e - sudo -E docker run --rm --privileged --init \ - --name "${HRR_NAME:-hrr-replay}" \ - --device /dev/kfd -v /dev/dri:/dev/dri --shm-size=4g \ - --security-opt seccomp=unconfined --ulimit core=-1:-1 \ - "${docker_env[@]}" "${vols[@]}" \ - "$IMAGE" /opt/hrr-tools/hrr-playback /data/hrr-archive "${EXTRA_ARGS[@]}" \ - 2>&1 | tee "$LOG" - return "${PIPESTATUS[0]}" +[[ -r /dev/kfd ]] || { + echo "error: /dev/kfd not accessible — AMD GPU driver required for native replay" >&2 + exit 1 } -if [[ "$use_docker" == "1" ]]; then - run_docker - RC=$? -else - run_native - RC=$? +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] native GPU=$GPU archive=$ARCHIVE" +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" From 275a5b4b31ba7a186e6e4b23979d132a99425e3d Mon Sep 17 00:00:00 2001 From: AMD HRR Team Date: Wed, 15 Jul 2026 15:54:16 +0000 Subject: [PATCH 5/6] Hrr-replay-findings: discover ROCm paths; user only supplies archive. Auto-find hrr-playback on PATH and /opt/rocm/bin; auto LD_LIBRARY_PATH; ask user once only for archive or playback if discovery fails. No Docker, no hardcoded lab paths, no run-this-script user instructions. --- skills/hrr-replay-findings/SKILL.md | 115 +++++++++--------- skills/hrr-replay-findings/examples.md | 52 +++----- skills/hrr-replay-findings/reference.md | 10 +- .../scripts/run_hrr_replay.sh | 77 +++++++----- 4 files changed, 129 insertions(+), 125 deletions(-) diff --git a/skills/hrr-replay-findings/SKILL.md b/skills/hrr-replay-findings/SKILL.md index aa6081c..b40a71e 100644 --- a/skills/hrr-replay-findings/SKILL.md +++ b/skills/hrr-replay-findings/SKILL.md @@ -1,86 +1,86 @@ --- name: hrr-replay-findings description: >- - Run HIP Runtime Replay (HRR) against a customer capture archive and analyze - the result. Classifies replay faults (OOB/read-only write, illegal access, - NaN/Inf divergence, hang, OOM) and extracts fault address, failing event index, - and implicated kernel. Requires only the archive and a shipped hrr-playback - binary — no source checkout or GPU configuration knowledge. Use when given an - HRR archive (capture.hrr/pid-*), a replay request, or Memory access fault output. + 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 -Run replay on an HRR **archive**, then emit a **structured finding**. The user -supplies the recording; the agent handles playback discovery, GPU selection, and -analysis. +Replay an HRR archive on the host GPU, then produce a structured finding. -## What the user provides +## What the user should say -| Input | Required? | -|-------|-----------| -| HRR archive (`capture.hrr/pid-*` with `events.bin`) | **Yes** | -| `hrr-playback` binary (on `PATH` or path via `HRR_PLAYBACK`) | **Yes** for replay | -| Source code, ROCm build tree, GPU index, Docker | **No** | +The user only needs to point at the recording. Examples: -## Quick workflow (agent) +- *"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) -``` -1. Find archive — largest events.bin under capture.hrr/pid-* -2. Find hrr-playback — PATH, HRR_PLAYBACK, or ask user once if missing -3. run_hrr_replay.sh --archive ... --analyze (GPU auto-selected) -4. Present finding + plain-language interpretation -``` +The user should **not** need to name scripts, set env vars, pick a GPU, or know where ROCm is installed. -## Run replay + analyze +## What to ask the user (only if missing) -```bash -SKILL=skills/hrr-replay-findings +| 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`). -# Only set HRR_PLAYBACK if hrr-playback is not already on PATH -./$SKILL/scripts/run_hrr_replay.sh \ - --archive /path/to/capture.hrr/pid-NNN \ - --analyze +## 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 ``` -Writes `hrr-replay-pid-NNN-.log` and `.finding.md` in the current directory. +**Execute in the same turn** — do not narrate planning steps. -Archive metadata only (no GPU): +### Discover `hrr-playback` (in order) -```bash -./$SKILL/scripts/run_hrr_replay.sh --archive /path/to/capture.hrr/pid-NNN --info -``` +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 -## Analyze an existing log +If the user gives `capture.hrr/` without a `pid-*` child, pick the `pid-*` directory whose `events.bin` is largest. -If replay already ran elsewhere: +### Commands (agent runs these — not the user) ```bash -python3 $SKILL/scripts/analyze_replay_finding.py \ - --log replay.log \ - --archive /path/to/capture.hrr/pid-NNN \ - --format markdown -o finding.md -``` +SKILL=/hrr-replay-findings # from .cursor/skills or installed skills dir -## Agent instructions +# Metadata only (~seconds, no full replay): +"$SKILL/scripts/run_hrr_replay.sh" --archive --info -**Execute immediately.** Do not say you will "locate the skill", "inspect expectations", or "look up documentation" — run the commands below in the same turn. +# Full replay + structured finding: +"$SKILL/scripts/run_hrr_replay.sh" --archive --analyze +``` -When the user gives an **archive** (and optionally `hrr-playback` path): +Log-only (no replay): ```bash -export HRR_PLAYBACK= # skip if hrr-playback is on PATH -SKILL/scripts/run_hrr_replay.sh --archive --analyze +python3 "$SKILL/scripts/analyze_replay_finding.py" \ + --log --archive --format markdown -o finding.md ``` -1. If `hrr-playback` path omitted: `command -v hrr-playback` or ask **once** for the binary path. -2. Run `run_hrr_replay.sh --archive ... --analyze` (GPU is auto-selected). -3. Read the generated `.finding.md` and summarize for the user. -4. On `replay_oom`: report insufficient VRAM; do not blame the capture. - -When the user gives a **log only**: run `analyze_replay_finding.py` on that log. +## System assumptions -**Never** require source checkout, GPU index, or HIP library paths from the user. +| 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 @@ -91,14 +91,13 @@ When the user gives a **log only**: run `analyze_replay_finding.py` on that log. | `illegal_memory_access` | Other GPU memory fault | | `nan_inf_divergence` | D2H numerical mismatch | | `hang` | Device/queue hang | -| `replay_oom` | Out of VRAM during replay | +| `replay_oom` | Out of VRAM — report environment issue | | `replay_fatal_api` | HIP API error stopped replay | -| `unknown` | Insufficient log data | ## Capture explainer (short) -- **events.bin** — recorded HIP API sequence (alloc, memcpy, kernel launches, sync) -- **blobs/** — code objects and sidecar payloads -- **Trailer** — missing when the original run crashed; reader still recovers complete events +- **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 index 36a6f2c..8af9a52 100644 --- a/skills/hrr-replay-findings/examples.md +++ b/skills/hrr-replay-findings/examples.md @@ -1,50 +1,32 @@ # Examples -## Example 1 — Archive only (typical customer handoff) +## What the user says -Customer sends `capture.hrr/pid-1842/`. Support ships matching `hrr-playback`. +**Full replay + analysis** -```bash -export HRR_PLAYBACK=/opt/rocm/bin/hrr-playback # only if not on PATH +> Replay and analyze my HRR archive at `/data/crash/capture.hrr/pid-1842` -skills/hrr-replay-findings/scripts/run_hrr_replay.sh \ - --archive capture.hrr/pid-1842 \ - --analyze -``` +**Archive summary only** -No GPU index, no source tree, no HIP paths. +> Summarize this HRR capture: `capture.hrr/pid-1842` ---- +**Existing log** -## Example 2 — Analyze log only +> Analyze this HRR replay log: `replay.log` (archive is `capture.hrr/pid-1842`) -```bash -python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py \ - --log replay.log \ - --archive capture.hrr/pid-1842 \ - --format markdown -``` +The user does not mention scripts, `HRR_PLAYBACK`, or GPU numbers. ---- +## What the agent does -## Example 3 — Expected MAF finding (abridged) +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 -| Field | Value | -|-------|-------| -| outcome | MAF | -| fault_class | read_only_page_fault | -| kernel_name | Cijk_..._SK3_... | -| d2h_fail | 0 | +## If hrr-playback is not in a standard location -`d2h_fail=0` → fault happened on GPU before host numerical checks failed. +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. -## Example 4 — OOM (environment, not capture bug) - -| Field | Value | -|-------|-------| -| fault_class | replay_oom | -| failing_api | hipMalloc | - -Retry after freeing GPU memory; script auto-picks the GPU with most free VRAM. +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 index b9248f9..4cef900 100644 --- a/skills/hrr-replay-findings/reference.md +++ b/skills/hrr-replay-findings/reference.md @@ -102,7 +102,15 @@ Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_SAV_UserArgs_MT128x192x128_..._SK3_..._WS64_WG16 The same archive can **PASS** on one `hrr-playback` build and fault on another. Record which playback binary produced each result. -## Parser script +## 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. ```bash python3 skills/hrr-replay-findings/scripts/analyze_replay_finding.py --help diff --git a/skills/hrr-replay-findings/scripts/run_hrr_replay.sh b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh index f05526a..8987128 100755 --- a/skills/hrr-replay-findings/scripts/run_hrr_replay.sh +++ b/skills/hrr-replay-findings/scripts/run_hrr_replay.sh @@ -1,20 +1,10 @@ #!/usr/bin/env bash # Run HRR replay on the host GPU and optionally analyze the log. -# -# Needs: archive directory + hrr-playback (on PATH or HRR_PLAYBACK). -# No Docker, no source tree, no GPU index (auto-picked by free VRAM). -# -# Usage: -# ./run_hrr_replay.sh --archive capture.hrr/pid-123 --analyze -# ./run_hrr_replay.sh --archive capture.hrr/pid-123 --info -# -# Optional env: -# HRR_PLAYBACK path to hrr-playback if not on PATH -# GPU force device index; otherwise auto-pick most free VRAM 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="" @@ -38,24 +28,49 @@ ARCHIVE="$(readlink -f "$ARCHIVE" 2>/dev/null || realpath "$ARCHIVE" 2>/dev/null [[ -d "$ARCHIVE" ]] || { echo "error: archive not found: $ARCHIVE" >&2; exit 1; } resolve_playback() { - if [[ -n "${HRR_PLAYBACK:-}" && -x "$HRR_PLAYBACK" ]]; then - echo "$HRR_PLAYBACK" - return - fi + local c candidates=() + [[ -n "${HRR_PLAYBACK:-}" ]] && candidates+=("$HRR_PLAYBACK") if command -v hrr-playback >/dev/null 2>&1; then - command -v hrr-playback - return + 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 + local best="" best_free=-1 idx free while read -r idx free; do [[ -n "$idx" ]] || continue if (( free > best_free )); then @@ -73,28 +88,28 @@ pick_gpu() { return fi fi - echo "[run_hrr_replay] default GPU 0 (set GPU=N to override)" >&2 + 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 - [[ -n "$HRR_PLAY" ]] || { - echo "error: hrr-playback not found. Set HRR_PLAYBACK=/path/to/hrr-playback" >&2 - exit 1 - } exec "$HRR_PLAY" "$ARCHIVE" --info "${EXTRA_ARGS[@]}" fi -[[ -n "$HRR_PLAY" ]] || { - echo "error: hrr-playback not found. Set HRR_PLAYBACK=/path/to/hrr-playback" >&2 - exit 1 -} - [[ -r /dev/kfd ]] || { - echo "error: /dev/kfd not accessible — AMD GPU driver required for native replay" >&2 + echo "error: /dev/kfd not accessible — AMD GPU driver required" >&2 exit 1 } @@ -103,7 +118,7 @@ if [[ -z "$LOG" ]]; then fi mkdir -p "$(dirname "$LOG")" 2>/dev/null || true -echo "[run_hrr_replay] native GPU=$GPU archive=$ARCHIVE" +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]} @@ -114,7 +129,7 @@ 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_PLAY:+--hrr-playback "$HRR_PLAY"} \ + --hrr-playback "$HRR_PLAY" \ --format markdown -o "$FINDING" echo "[run_hrr_replay] finding=$FINDING" fi From 2e57fcaabfb1b419b42c0971d1d64707277922db Mon Sep 17 00:00:00 2001 From: AMD HRR Team Date: Wed, 15 Jul 2026 15:55:59 +0000 Subject: [PATCH 6/6] Restore Parser script section header in hrr-replay-findings reference. --- skills/hrr-replay-findings/reference.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/skills/hrr-replay-findings/reference.md b/skills/hrr-replay-findings/reference.md index 4cef900..5cbc4e0 100644 --- a/skills/hrr-replay-findings/reference.md +++ b/skills/hrr-replay-findings/reference.md @@ -112,6 +112,8 @@ The same archive can **PASS** on one `hrr-playback` build and fault on another. `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 ```