diff --git a/.claude-plugin/marketplace.json b/.claude-plugin/marketplace.json
index de99b2f..3535f65 100644
--- a/.claude-plugin/marketplace.json
+++ b/.claude-plugin/marketplace.json
@@ -17,9 +17,10 @@
         "./skills/local-ai-use",
         "./skills/local-ai-app-integration",
         "./skills/serving-llms-on-instinct",
-        "./skills/tracelens-analysis-orchestrator"
+        "./skills/tracelens-analysis-orchestrator",
+        "./skills/hrr-replay-analysis"
       ],
-      "description": "AMD's verified Agent Skills in one plugin: route image/audio through local AI on Ryzen AI, serve LLMs on AMD Instinct GPUs with vLLM, and analyze GPU kernel and PyTorch trace performance."
+      "description": "AMD's verified Agent Skills in one plugin: route image/audio through local AI on Ryzen AI, serve LLMs on AMD Instinct GPUs with vLLM, analyze GPU kernel and PyTorch trace performance, and replay HIP Record and Replay captures to triage GPU crashes."
     }
   ]
 }
diff --git a/.codex-plugin/plugin.json b/.codex-plugin/plugin.json
index cd38ee1..c15ef84 100644
--- a/.codex-plugin/plugin.json
+++ b/.codex-plugin/plugin.json
@@ -21,12 +21,13 @@
     "./skills/local-ai-use",
     "./skills/local-ai-app-integration",
     "./skills/serving-llms-on-instinct",
-    "./skills/tracelens-analysis-orchestrator"
+    "./skills/tracelens-analysis-orchestrator",
+    "./skills/hrr-replay-analysis"
   ],
   "interface": {
     "displayName": "AMD Skills",
     "shortDescription": "Agent Skills for AMD-optimized workflows.",
-    "longDescription": "AMD's verified Agent Skills in one plugin: route image/audio through local AI on Ryzen AI, serve LLMs on AMD Instinct GPUs with vLLM, and analyze GPU kernel and PyTorch trace performance.",
+    "longDescription": "AMD's verified Agent Skills in one plugin: route image/audio through local AI on Ryzen AI, serve LLMs on AMD Instinct GPUs with vLLM, analyze GPU kernel and PyTorch trace performance, and replay HIP Record and Replay captures to triage GPU crashes.",
     "developerName": "AMD",
     "category": "Developer Tools",
     "capabilities": [
diff --git a/.cursor-plugin/marketplace.json b/.cursor-plugin/marketplace.json
index 62c7697..542b372 100644
--- a/.cursor-plugin/marketplace.json
+++ b/.cursor-plugin/marketplace.json
@@ -16,9 +16,10 @@
         "./skills/local-ai-use",
         "./skills/local-ai-app-integration",
         "./skills/serving-llms-on-instinct",
-        "./skills/tracelens-analysis-orchestrator"
+        "./skills/tracelens-analysis-orchestrator",
+        "./skills/hrr-replay-analysis"
       ],
-      "description": "AMD's verified Agent Skills in one plugin: route image/audio through local AI on Ryzen AI, serve LLMs on AMD Instinct GPUs with vLLM, and analyze GPU kernel and PyTorch trace performance.",
+      "description": "AMD's verified Agent Skills in one plugin: route image/audio through local AI on Ryzen AI, serve LLMs on AMD Instinct GPUs with vLLM, analyze GPU kernel and PyTorch trace performance, and replay HIP Record and Replay captures to triage GPU crashes.",
       "author": {
         "name": "AMD"
       },
diff --git a/README.md b/README.md
index c9fb0da..e1f28c9 100644
--- a/README.md
+++ b/README.md
@@ -83,7 +83,7 @@ Cross-stack skills, from client to cloud.
 | `rocm-doctor` | Diagnose ROCm / PyTorch / llama.cpp failures on AMD GPUs against a fixed list of known misconfigurations. | _planned_ |
 | `hyperloom-kernel-optimizer` | Autonomously optimizes LLM inference on AMD GPUs. | _planned_ |
 | `vllm-semantic-router` | Setup a vLLM router that semantically maps your request to the best available platform. | _planned_ |
-| `hrr-replay-analysis` | Record, replay, and analyze GPU workload behavior on ROCm across AMD Instinct, Radeon, and Ryzen hardware using HIP Record and Replay archives. | _planned_ |
+| [`hrr-replay-analysis`](skills/hrr-replay-analysis/SKILL.md) | Record, replay, and analyze GPU workload behavior on ROCm across AMD Instinct, Radeon, and Ryzen hardware using HIP Record and Replay archives. | in-repo |
 
 ### Server-Native
 
diff --git a/skills/hrr-replay-analysis/SKILL.md b/skills/hrr-replay-analysis/SKILL.md
new file mode 100644
index 0000000..195c711
--- /dev/null
+++ b/skills/hrr-replay-analysis/SKILL.md
@@ -0,0 +1,214 @@
+---
+name: hrr-replay-analysis
+description: >-
+  Replays a HIP Record and Replay (HRR) capture of a GPU workload and reports
+  what went wrong as a structured finding: fault class (write to a read-only
+  page, illegal memory access, NaN/Inf output divergence, hang, out of memory,
+  clean pass), fault address, failing event index, and the implicated kernel.
+  Also records a new capture when the user has none yet. Use when the user
+  supplies an HRR archive, a `capture.hrr` directory or a `pid-*/events.bin`
+  tree, asks to replay, summarize or triage a GPU recording, mentions
+  hrr-playback or HIP Record and Replay, or brings a "Memory access fault by
+  GPU" or "read-only page" failure from a vLLM, PyTorch or hipBLASLt run. Do not
+  use to patch or rebuild a kernel, to drive rocgdb, or to verify a fix with
+  `--replace-kernel`; this skill stops at the finding and hands off.
+allowed-tools: Bash, Read
+---
+
+# HRR Replay Analysis
+
+Replay an HRR archive on the host GPU, then produce a structured finding. An HRR
+archive is a deterministic record of a real HIP workload, so the same fault can
+be reproduced on a different machine from the recording alone.
+
+## 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`"*
+- *"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` 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`).
+
+Ask for the original failure signature (the user's `Memory access fault` line or
+`serve.log`) only when the replay itself comes back clean. A clean replay of a
+crashing workload is a finding in its own right, and the user's log is what
+tells you which fault you were supposed to see.
+
+## Agent workflow
+
+```
+1. Resolve archive — see below; a bare capture.hrr root needs a pid-* pick
+2. Discover hrr-playback (see below); ask the user only if not found
+3. Read metadata first: run_hrr_replay.sh --archive <dir> --info
+4. Replay and analyze: run_hrr_replay.sh --archive <dir> --analyze
+5. Read the generated .finding.md and explain it in plain language
+```
+
+**Execute in the same turn** — do not narrate planning steps.
+
+Step 3 is cheap and needs no GPU, so always do it before a full replay: it
+catches an unreadable archive, a wrong `pid-*`, or a reader/format mismatch
+before spending GPU time.
+
+Step 4 replays with `--sync-after-launch`, which `run_hrr_replay.sh` adds for
+you. Without it the GPU is serialized only once at the end, so a fault is
+reported but not attributed and the finding has no failing event and no kernel.
+Pass `--no-sync` only when throughput matters more than attribution.
+
+### 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.
+
+### Resolve the archive
+
+A capture is either a single process directory or a multi-process tree:
+
+```
+capture.hrr/
+  pid-138/          # often the init/parent process — small
+  pid-680/
+    events.bin      # the workload stream
+    blobs/
+    code_objects/
+```
+
+Run `--info` on the **archive root** first when the path has `pid-*` children:
+that prints the process table (PID, parent PID, complete, events, blobs) so you
+can name the workload process instead of guessing. Then pick the `pid-*`
+directory with the largest `events.bin` and use it for the replay:
+
+```bash
+find <capture.hrr> -name events.bin -printf '%s %h\n' 2>/dev/null | sort -rn | head -1
+```
+
+If the user already gave a `pid-*` path, use it directly.
+
+### Commands (the agent runs these, not the user)
+
+```bash
+SKILL=<path-to>/hrr-replay-analysis   # installed skill directory
+
+# Metadata only (seconds, no GPU needed):
+"$SKILL/scripts/run_hrr_replay.sh" --archive <archive-or-pid-dir> --info
+
+# Full replay + structured finding:
+"$SKILL/scripts/run_hrr_replay.sh" --archive <pid-dir> --analyze
+```
+
+Log-only (no replay):
+
+```bash
+python3 "$SKILL/scripts/analyze_replay_finding.py" \
+  --log <replay.log> --archive <pid-dir> --format markdown -o finding.md
+```
+
+## System assumptions
+
+| Assumption | Default |
+|------------|---------|
+| ROCm install | `/opt/rocm` (`$ROCM_PATH` override) |
+| GPU | Auto-pick the device with the most free VRAM (`GPU=<index>` override) |
+| Replay mode | Native host (`/dev/kfd`); no Docker |
+
+## Fault taxonomy
+
+| `fault_class` | Meaning |
+|---------------|---------|
+| `replay_pass` | Clean replay, all device-to-host checks passed |
+| `read_only_page_fault` | Write to a read-only page |
+| `illegal_memory_access` | Other GPU memory fault |
+| `nan_inf_divergence` | Device-to-host numerical mismatch against the recording |
+| `hang` | Device or queue hang |
+| `replay_oom` | Out of VRAM on the replay machine — an environment issue, not the recorded bug |
+| `replay_fatal_api` | A HIP API returned an error and stopped replay |
+| `archive_version_mismatch` | The archive format and this `hrr-playback` disagree; nothing was replayed |
+
+## A faulting ATen kernel needs the original failure signature
+
+PyTorch and vLLM reach the GPU through `<<<>>>` (`hipLaunchByPtr`), and those
+kernels pass device pointers inside by-value structs: `vectorized_elementwise_kernel`
+takes a `std::array<char*,N>`, `reduce_kernel` a config struct with pointers at
+arbitrary offsets. Capture records those offsets and replay translates them, with
+a defensive rescan for any the capture-time heuristic missed, so on a current
+build these kernels replay faithfully and a fault on one is a real finding.
+
+Two things still make such a fault ambiguous: the detector is a value-based
+heuristic, and an archive recorded before that support landed carries no offsets
+at all, in which case replay launches with a capture-time address and faults on a
+workload that is fine. Both look identical to a workload defect.
+
+So the fault class stays what the evidence says (`illegal_memory_access`), the
+finding carries a note, and you **ask the user for their original failure
+signature** before calling it their bug. If their run never faulted here, suspect
+the recording. Faults on `hipModuleLaunchKernel`-launched kernels (hipBLASLt
+GEMMs, custom HIP kernels) do not carry this ambiguity.
+
+## Reading `--info`
+
+`Complete: NO` is the **expected** signature of a capture whose workload
+crashed: the trailer is written on clean shutdown only. The reader recovers
+every complete record and the archive is still replayable, so do not report
+`Complete: NO` as corruption. `Recovered: N events` tells you how much survived.
+
+## When the archive will not open
+
+`[HRR] Version mismatch: file=N reader=M` means this `hrr-playback` cannot read
+this archive. Stop and report `archive_version_mismatch`: replaying is
+impossible until the versions line up. `file` newer than `reader` needs a newer
+playback build; `reader` newer than `file` needs the older matching build or a
+fresh capture. Never present a mismatch as a workload result.
+
+## Recording a new capture
+
+When the user has a failing workload but no archive, capture is an
+environment-variable change to their existing run — no code edit, no rebuild:
+
+```bash
+HIP_HRR_CAPTURE_OUTPUT=./capture.hrr <their normal command>
+```
+
+Capture survives a crash, so let the workload fail as it did before. The archive
+is what they then replay or hand over.
+
+## Guardrails
+
+- **Report only what the evidence supports.** The archive does not record the
+  GPU SKU, library commits, or the full software stack. If the faulting kernel
+  is not in the log or the archive, say it is unknown rather than naming a
+  plausible one.
+- **Distinguish the recorded bug from the replay environment.** `replay_oom` and
+  `archive_version_mismatch` are properties of the replay machine.
+- **Do not pick a GPU on a shared host without confirming.** Replay takes a
+  whole device; `--info` does not.
+- **Note the playback build.** The same archive can pass on one `hrr-playback`
+  build and fault on another, so record which binary produced the result.
+- **Privacy.** Archives can contain prompts and payload data. Say so before the
+  user forwards one outside their organization.
+
+## Out of scope — hand off after the finding
+
+Kernel patching, rebuilds, `--replace-kernel` verification, rocgdb and core-dump
+inspection, and bisecting a divergent kernel all belong to deeper crash
+analysis. Produce the finding, name the next action, and stop.
+
+See [reference.md](reference.md) for archive layout, log-line formats and flags,
+and [examples.md](examples.md) for worked prompts.
diff --git a/skills/hrr-replay-analysis/evals/evals.py b/skills/hrr-replay-analysis/evals/evals.py
new file mode 100644
index 0000000..d6359b1
--- /dev/null
+++ b/skills/hrr-replay-analysis/evals/evals.py
@@ -0,0 +1,212 @@
+# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# See LICENSE for license information.
+
+"""Behavioral tests for the `hrr-replay-analysis` skill.
+
+Two tiers, because triage and replay have different hardware needs:
+
+* **Recorded-evidence tier** (default). Runs against replay logs checked into
+  ``evals/fixtures/``, so it grades the agent's *reasoning* -- does it name the
+  right kernel, the right fault address, the right fault class -- on a machine
+  with no GPU. These are the tests CI can run today.
+* **Live-replay tier**. Replays a real archive on a ROCm host. Skipped unless
+  ``/dev/kfd``, an ``hrr-playback`` binary, and ``HRR_EVAL_ARCHIVE`` are all
+  present, following the same skip-when-unavailable pattern as `local-ai-use`.
+
+Run locally (needs the `claude` CLI authenticated):
+
+    cd eval/behavioral
+    python -m pytest -c pytest.ini -p conftest \
+        ../../skills/hrr-replay-analysis/evals/evals.py
+
+Each check on `run` prints a `[PASS]`/`[FAIL]` line and raises on failure.
+`logs_contains` / `workspace_contains` are deterministic; `should` /
+`should_not` are graded by an LLM judge over the captured evidence.
+"""
+
+from __future__ import annotations
+
+import os
+import shutil
+from pathlib import Path
+
+import pytest
+
+from harness import claude
+
+SKILL = "hrr-replay-analysis"
+FIXTURES = Path(__file__).resolve().parent / "fixtures"
+
+# The fault recorded in fixtures/read_only_page_fault: a StreamK GEMM writing
+# past the tile edge into a read-only page. The agent has to recover all three.
+FAULT_ADDRESS = "0x7f2a4c000000"
+FAULT_KERNEL = "Cijk_Ailk_Bljk_HHS_BH_MT128x128x16_MI16x16x16x1_SK3_WGM8"
+FAULT_CLASS = "read_only_page_fault"
+
+
+def _stage(agent, fixture: str, name: str = "replay.log") -> Path:
+    """Copy a recorded fixture into the agent's workspace."""
+    dest = agent.workspace / name
+    shutil.copyfile(FIXTURES / fixture / "replay.log", dest)
+    return dest
+
+
+def _finding_text(agent, filename: str = "finding.md") -> str:
+    path = agent.workspace / filename
+    assert path.is_file(), f"{filename} was not written to {agent.workspace}"
+    return path.read_text(encoding="utf-8", errors="replace")
+
+
+def _live_replay_skip_reason() -> str | None:
+    """Why the live-replay tier cannot run here, or None when it can."""
+    if not Path("/dev/kfd").exists():
+        return "no /dev/kfd: this host has no ROCm GPU"
+    rocm = Path(os.environ.get("ROCM_PATH", "/opt/rocm"))
+    if not (shutil.which("hrr-playback") or (rocm / "bin" / "hrr-playback").is_file()):
+        return "hrr-playback not found on PATH or under $ROCM_PATH/bin"
+    archive = os.environ.get("HRR_EVAL_ARCHIVE", "").strip()
+    if not archive:
+        return "HRR_EVAL_ARCHIVE is not set to an HRR archive directory"
+    if not Path(archive).is_dir():
+        return f"HRR_EVAL_ARCHIVE does not exist: {archive}"
+    return None
+
+
+def test_triage_read_only_page_fault_from_log():
+    """The core case: a recorded fault log becomes a correct structured finding."""
+    with claude("opus", skill=SKILL) as agent:
+        _stage(agent, "read_only_page_fault")
+
+        run = agent.prompt(
+            "Analyze this replay log from an HRR run and write the structured "
+            "finding to finding.md: replay.log"
+        )
+
+        run.logs_contains(SKILL)
+        run.workspace_contains("finding.md")
+
+        finding = _finding_text(agent)
+        assert FAULT_CLASS in finding, f"fault class missing from finding:\n{finding}"
+        assert FAULT_ADDRESS in finding, f"fault address missing from finding:\n{finding}"
+        assert FAULT_KERNEL in finding, f"faulting kernel missing from finding:\n{finding}"
+
+        run.should("Report the failure as a write to a read-only page")
+        run.should(f"Name {FAULT_KERNEL} as the implicated kernel")
+        run.should(f"Report the faulting address {FAULT_ADDRESS}")
+        run.should("Explain the finding in plain language to the user")
+
+        run.should_not("Re-record the workload instead of analyzing the log provided")
+        run.should_not("Attempt to patch, rebuild or replace the kernel")
+        run.should_not("Invent a root cause the log does not support")
+
+
+def test_version_mismatch_is_not_reported_as_a_workload_failure():
+    """False-positive guard: an unreadable archive is a tooling problem.
+
+    The skill is explicit that a format mismatch means nothing was replayed, so
+    reporting it as a crash would send a kernel developer after a bug that this
+    run never observed.
+    """
+    with claude("opus", skill=SKILL) as agent:
+        _stage(agent, "version_mismatch")
+
+        run = agent.prompt("Analyze this HRR replay log and tell me what happened: replay.log")
+
+        run.logs_contains(SKILL)
+
+        run.should("Report that the archive format and the hrr-playback build disagree")
+        run.should("State that nothing was replayed")
+        run.should("Say that a matching hrr-playback build or a fresh capture is needed")
+
+        run.should_not("Report a GPU fault, a crash, or any workload-level failure")
+        run.should_not("Name a faulting kernel")
+
+
+def test_clean_replay_asks_for_the_original_failure_signature():
+    """A clean replay of a workload the user says crashed is itself a finding."""
+    with claude("opus", skill=SKILL) as agent:
+        _stage(agent, "clean_pass")
+
+        run = agent.prompt(
+            "This workload crashed for our customer with a memory access fault, "
+            "but here is the replay log from the capture: replay.log. "
+            "What happened?"
+        )
+
+        run.logs_contains(SKILL)
+
+        run.should("Report the replay as a clean pass with all device-to-host checks passing")
+        run.should(
+            "Treat the clean replay of a crashing workload as a finding, for example by "
+            "asking for the original failure signature or noting the fault did not reproduce"
+        )
+
+        run.should_not("Claim to have reproduced the customer's memory access fault")
+        # The archive listing in this fixture contains a GEMM kernel. A clean
+        # replay implicates none of them, so naming one would read as a culprit.
+        run.should_not("Name any kernel as implicated in a failure")
+
+
+def test_aten_chevron_fault_is_reported_but_not_asserted_as_a_customer_bug():
+    """The ATen fault is a real finding that must not be stated as a verdict.
+
+    ATen kernels launched through `<<<>>>` pass device pointers inside by-value
+    structs. Replay translates those, so such a fault can be genuine, but the
+    translation is heuristic and older archives lack the recorded offsets, so it
+    can equally be an artefact of the recording. The agent must report the fault
+    it observed and ask for the original failure signature, rather than resolving
+    the ambiguity in either direction on its own.
+    """
+    with claude("opus", skill=SKILL) as agent:
+        _stage(agent, "aten_chevron_limitation")
+
+        run = agent.prompt("Analyze this HRR replay log and tell me what happened: replay.log")
+
+        run.logs_contains(SKILL)
+
+        run.should("Report the memory fault and name the faulting ATen kernel")
+        run.should("Flag that a <<<>>>-launched kernel faulting at replay may come from the recording rather than the workload")
+        run.should("Ask for the original failure signature before concluding anything about the workload")
+
+        run.should_not("State as settled fact that the customer's workload has a memory bug")
+        run.should_not("State as settled fact that the fault is only an artefact of the recording")
+
+
+def test_skill_does_not_activate_on_an_unrelated_prompt():
+    """False-activation screen: an unrelated request must not pull in the skill."""
+    with claude("opus", skill=SKILL) as agent:
+        (agent.workspace / "notes.txt").write_text(
+            "Sprint notes: rotate the API token, then update the changelog.\n",
+            encoding="utf-8",
+        )
+
+        run = agent.prompt("Read notes.txt and summarize it in one sentence.")
+
+        run.should_not("Run the HRR replay analysis scripts")
+        run.should_not("Treat notes.txt as a GPU capture, replay log or HRR archive")
+
+
+@pytest.mark.skipif(
+    _live_replay_skip_reason() is not None,
+    reason=f"live replay unavailable: {_live_replay_skip_reason()}",
+)
+def test_live_replay_of_a_real_archive():
+    """Live-replay tier: replay a real archive end to end on a ROCm host.
+
+    Set ``HRR_EVAL_ARCHIVE`` to an archive root or a ``pid-*`` directory. The
+    archive is not checked in, so this stays opt-in and out of CI until a GPU
+    runner that can reach the agent API exists.
+    """
+    archive = os.environ["HRR_EVAL_ARCHIVE"]
+
+    with claude("opus", skill=SKILL) as agent:
+        run = agent.prompt(f"Replay and analyze this HRR archive: {archive}")
+
+        run.logs_contains(SKILL)
+
+        run.should("Read the archive metadata with --info before replaying it")
+        run.should("Produce a structured finding with an outcome and a fault class")
+        run.should("Record which hrr-playback build produced the result")
+
+        run.should_not("Ask the user for a GPU index, a ROCm path or a Docker command")
diff --git a/skills/hrr-replay-analysis/evals/fixtures/aten_chevron_limitation/replay.log b/skills/hrr-replay-analysis/evals/fixtures/aten_chevron_limitation/replay.log
new file mode 100644
index 0000000..5609609
--- /dev/null
+++ b/skills/hrr-replay-analysis/evals/fixtures/aten_chevron_limitation/replay.log
@@ -0,0 +1,14 @@
+[run_hrr_replay] playback=/opt/rocm/bin/hrr-playback
+[run_hrr_replay] adding --sync-after-launch (disable with --no-sync)
+[run_hrr_replay] GPU=3 archive=/data/llama_run.hrr/pid-680
+[HRR] Archive : 177203 events, 8214 kernels, 4733 blobs, 41 code objects
+[HRR] Complete:     yes (clean shutdown)
+[HRR] [replay] Device state OK — beginning replay
+[HRR] Event 41188: hipMallocAsync
+[HRR] Event 41202: hipModuleLaunchKernel   -> Kernel '_ZN2at6native29vectorized_elementwise_kernelILi4EZZZNS0_18add_kernel_cudaERNS_18TensorIteratorBaseEENKUlvE_clEvENKUlvE2_clEvEUlffE_NS_6detail5ArrayIPcLi3EEEEEviT0_T1_' OK
+[HRR] Event 41266: hipMemcpyAsync
+[HRR] Event 41310: hipLaunchKernel
+:0:rocdevice.cpp :3608: Memory Fault Error [GPU index: 0, faulting addr: 0x7f8c11a00000, kernel: _ZN2at6native29vectorized_elementwise_kernelILi4EZZZNS0_18mul_kernel_cudaERNS_18TensorIteratorBaseEENKUlvE_clEvENKUlvE2_clEvEUlffE_NS_6detail5ArrayIPcLi3EEEEEviT0_T1_]
+Memory access fault by GPU node-4 (Agent handle: 0x55d1a2f00c40) on address 0x7f8c11a00000. Reason: Unknown.
+Aborted (core dumped)
+[run_hrr_replay] log=hrr-replay-pid-680-20260805T120400Z.log exit=134
diff --git a/skills/hrr-replay-analysis/evals/fixtures/clean_pass/replay.log b/skills/hrr-replay-analysis/evals/fixtures/clean_pass/replay.log
new file mode 100644
index 0000000..d51bf1a
--- /dev/null
+++ b/skills/hrr-replay-analysis/evals/fixtures/clean_pass/replay.log
@@ -0,0 +1,18 @@
+[HRR] hrr-playback 0.4.0 (format reader v4)
+[HRR] Archive : 412773 events, 1904 kernels, 1904 blobs, 42 code objects
+[HRR] Replaying pid-1842 on GPU node-4
+[HRR progress] seq=412773 elapsed=12.4s kernels=1904 d2h_pass=401 d2h_fail=0 d2h_attempted=401 last="_ZN2at6native18elementwise_kernelILi128ELi2EEEvi"
+D2H checks   : 401 pass, 0 fail, 0 skipped
+[HRR] PASS
+
+HRR Archive: pid-1842
+========================================
+Complete:     yes (clean shutdown)
+Events:       412773
+Kernels:      1904
+
+Kernel Summary (first 20):
+  ID   Kernel                                             Grid            Block           SharedMem
+  --   ------                                             ----            -----           ---------
+  0    Cijk_Alik_Bljk_BBS_BH_Bias_HA_S_MT128x128x12       [2048,1,1]      [256,1,1]       0
+  1    _ZN2at6native18elementwise_kernelILi128ELi2EEEvi    [512,1,1]       [128,1,1]       0
diff --git a/skills/hrr-replay-analysis/evals/fixtures/read_only_page_fault/replay.log b/skills/hrr-replay-analysis/evals/fixtures/read_only_page_fault/replay.log
new file mode 100644
index 0000000..e9c04d0
--- /dev/null
+++ b/skills/hrr-replay-analysis/evals/fixtures/read_only_page_fault/replay.log
@@ -0,0 +1,10 @@
+[HRR] hrr-playback 0.4.0 (format reader v4)
+[HRR] Archive : 9109510 events, 11239 kernels, 11239 blobs, 187 code objects
+[HRR] Replaying pid-680 on GPU node-6
+[HRR progress] seq=4210000 elapsed=31.2s kernels=5120 d2h_pass=180 d2h_fail=0 d2h_attempted=180 last="Cijk_Ailk_Bljk_HHS_BH_MT64x64x32_MI16x16x16x1_WGM4"
+[HRR progress] seq=9109402 elapsed=68.9s kernels=11238 d2h_pass=396 d2h_fail=0 d2h_attempted=396 last="Cijk_Ailk_Bljk_HHS_BH_MT128x128x16_MI16x16x16x1_SK3_WGM8"
+[HRR] Dispatch: kernarg_address=0x7f2a3d201000 grid=[2048, 1, 1], workgroup=[256, 1, 1]
+Memory access fault by GPU node-6 (Agent handle: 0x55f1c02a4b80) on address 0x7f2a4c000000. Reason: Write to a read-only page.
+Memory Fault Error [host: gpu-node, GPU index: 2, faulting addr: 0x7f2a4c000000, kernel: Cijk_Ailk_Bljk_HHS_BH_MT128x128x16_MI16x16x16x1_SK3_WGM8]
+[HRR] Fatal: GPU error after T0 Event 9109510 (hipModuleLaunchKernel): 4 (hipErrorLaunchFailure)
+[HRR] aborting replay
diff --git a/skills/hrr-replay-analysis/evals/fixtures/version_mismatch/replay.log b/skills/hrr-replay-analysis/evals/fixtures/version_mismatch/replay.log
new file mode 100644
index 0000000..3614084
--- /dev/null
+++ b/skills/hrr-replay-analysis/evals/fixtures/version_mismatch/replay.log
@@ -0,0 +1,4 @@
+[HRR] hrr-playback 0.3.1 (format reader v3)
+[HRR] Opening archive pid-680
+[HRR] Version mismatch: file=4 reader=3
+[HRR] archive cannot be read by this build; nothing was replayed
diff --git a/skills/hrr-replay-analysis/examples.md b/skills/hrr-replay-analysis/examples.md
new file mode 100644
index 0000000..a924a3c
--- /dev/null
+++ b/skills/hrr-replay-analysis/examples.md
@@ -0,0 +1,63 @@
+# 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`
+
+**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 at `/opt/rocm/bin/hrr-playback`
+2. If missing, asks: *"Where is hrr-playback installed?"*
+3. Runs `run_hrr_replay.sh --archive ... --info` to check the archive opens
+4. Runs `run_hrr_replay.sh --archive ... --analyze`
+5. Presents the finding
+
+## Multi-process capture
+
+User provides the root, not a process directory:
+
+> Here's the capture: `/data/crash/capture.hrr`
+
+The agent runs `--info` on the root, reads the process table, and reports which
+process it chose before replaying:
+
+> The capture holds two processes: `pid-138` (4,211 events, complete) and
+> `pid-680` (13,118,764 events, `Complete: NO`). `pid-680` is the workload and
+> the crash-truncated one, so I'm replaying that.
+
+## Clean replay of a crashing workload
+
+Replay returns `replay_pass` but the user's original run faulted. That is a
+finding, not a dead end: report that the recording does not reproduce the fault
+on this machine and this playback build, then ask for the original fault line to
+confirm which failure was expected.
+
+## Archive will not open
+
+```
+[HRR] Version mismatch: file=4 reader=3
+```
+
+Report `archive_version_mismatch` and stop. The archive is newer than the
+installed `hrr-playback`, so a newer playback build is needed before any replay
+result means anything.
+
+## If hrr-playback is not in a standard location
+
+User answers: *"It's in `/opt/amd-hrr/bin/hrr-playback`"*
+
+The agent sets `HRR_PLAYBACK=/opt/amd-hrr/bin/hrr-playback` for that run only and
+re-runs. If a `lib/` directory sits beside `bin/`, the runner picks it up
+automatically.
diff --git a/skills/hrr-replay-analysis/reference.md b/skills/hrr-replay-analysis/reference.md
new file mode 100644
index 0000000..af0b58c
--- /dev/null
+++ b/skills/hrr-replay-analysis/reference.md
@@ -0,0 +1,257 @@
+# HRR capture and replay reference
+
+- [Directory layout](#directory-layout)
+- [Archive summary output](#archive-summary-output)
+- [Completeness markers](#completeness-markers)
+- [Version mismatch](#version-mismatch)
+- [events.bin record model](#eventsbin-record-model-conceptual)
+- [Replay log lines](#replay-log-lines)
+- [Playback flags](#playback-flags)
+- [Capture environment variables](#capture-environment-variables)
+- [Tensile / hipBLASLt kernel names](#tensile--hipblaslt-kernel-name-cheat-sheet)
+- [Playback build fidelity](#playback-build-fidelity)
+- [ROCm layout](#rocm-layout-assumed)
+- [Parser script](#parser-script)
+
+## Directory layout
+
+```
+capture.hrr/
+  pid-<capture_pid>/
+    events.bin          # primary event stream (may be GB-scale)
+    blobs/              # code objects, graphs, memcpy payloads
+    code_objects/
+    (optional metadata files)
+```
+
+A capture of a multi-process workload holds one `pid-*` directory per process.
+Run `--info` on the archive **root** for the process table, then pick the
+`pid-*` directory with the **largest `events.bin`** for the faulting workload —
+the smaller ones are usually short-lived parents or init processes.
+
+## Archive summary output
+
+`hrr-playback <archive> --info` on the root of a multi-process archive:
+
+```
+HRR Archive Root: /data/capture.hrr
+========================================
+Capture Mode: in-tree
+Owner PID:    680
+Processes:    2
+
+  PID          Parent PID   Complete   Events       Blobs      Path
+  ---          ----------   --------   ------       -----      ----
+  138          1            yes        4211         12         /data/capture.hrr/pid-138
+  680          138          NO         13118764     8034       /data/capture.hrr/pid-680
+```
+
+A single process directory:
+
+```
+HRR Archive: /data/capture.hrr/pid-680
+========================================
+Complete:     NO (no shutdown trailer; capture likely crashed)
+Recovered:    13118764 events
+Events:       13118764
+Kernels:      797227
+Blobs:        8034
+Code Objects: 41
+```
+
+`--info` needs no GPU. Add `--events` for the full event log (very large output).
+
+## Completeness markers
+
+| Signal | Meaning |
+|--------|---------|
+| `Complete:     yes (clean shutdown)` | Shutdown trailer present |
+| `Complete:     NO (crash-truncated; trailing torn record discarded)` | Crash capture; last record was partial and was dropped |
+| `Complete:     NO (no shutdown trailer; capture likely crashed)` | Crash capture; every complete record kept |
+| `recovered N events` (stderr) | Reader repaired a torn tail and kept N events |
+
+The value is lowercase `yes` or upper-case `NO`, each followed by an
+explanation, so match it case-insensitively.
+
+Crash captures are **expected** to lack a trailer and still replay. `Complete:
+NO` is not corruption.
+
+## Version mismatch
+
+```
+[HRR] Version mismatch: file=4 reader=3
+```
+
+Printed to stderr by the archive reader when the on-disk format version and the
+`hrr-playback` build disagree. Nothing is replayed. `file` > `reader` needs a
+newer playback binary; `reader` > `file` needs the matching older binary or a
+fresh capture.
+
+## 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
+device-to-host snapshot descriptors.
+
+## Replay log lines
+
+### Progress
+
+```
+[HRR progress] elapsed_s=612.4 seq=13118764 kernels=797227 d2h_pass=4303 d2h_fail=0 d2h_attempted=4315 last="Cijk_..."
+```
+
+- `seq` — last replayed event sequence number (use as **failing_call_index**
+  proxy when a fault follows)
+- `kernels` — kernel launch count so far
+- `d2h_*` — device-to-host validation counters
+
+### Final summary
+
+```
+[HRR]   D2H checks     : 4303 pass (4300 exact, 3 within tol), 0 fail, 12 skipped
+[HRR] PASS
+```
+
+`[HRR] PASS` / `[HRR] FAIL` is the replay verdict; the process exit code is 0
+only when no device-to-host check failed.
+
+### GPU memory fault (ROCr)
+
+```
+Memory access fault by GPU node-4 (Agent handle: 0x...) on address 0x7f2c0a800000. Reason: Write access to a read-only page.
+:0:rocdevice.cpp:NNNN: Memory Fault Error [..., faulting addr: 0x7f2c0a800000, kernel: Cijk_...]
+```
+
+The bracket's leading fields vary by ROCm build. Both of these appear in the
+wild and both are parsed:
+
+```
+Memory Fault Error [GPU index: 0, faulting addr: 0x..., kernel: hrr_fault_kernel]
+Memory Fault Error [host: h1, GPU index: 2, faulting addr: 0x..., kernel: Cijk_...]
+```
+
+Extract: **fault_address**, **kernel_name**, **gpu_node**, **fault_reason**.
+
+A GPU fault usually also trips an HRR abort line (below). The memory fault is
+the finding; the abort is its consequence.
+
+### Hang analysis block
+
+```
+Dispatch Header = 0x..., grid=[...], workgroup=[...], kernarg_address=0x..., kernel_obj=0x...
+```
+
+Extract: **kernarg_address**, **grid**, **workgroup** — ties the fault to the
+launch packet.
+
+### Per-event progress (`--sync-after-launch`)
+
+```
+[HRR] Event 1304: hipModuleLaunchKernel   -> Kernel '..._FillFunctor...' OK
+[HRR] Event 1324: hipModuleLaunchKernel
+```
+
+The last of these before a fault is the failing dispatch. A GPU fault aborts the
+process before HRR writes its own `Fatal` line, so on a hard memory fault this
+is often the only record of **failing_call_index** and the implicated kernel.
+
+### Fatal abort
+
+```
+[HRR] Fatal: T146 Event 9268 (hipMalloc) returned 2 (out of memory) — aborting replay
+[HRR] Fatal: GPU error after T146 Event 9268 (hipLaunchKernel): 1 (invalid argument) — aborting
+```
+
+Extract: **failing_thread**, **failing_call_index**, **failing_api**. An `out of
+memory` here is a replay-environment limit, not the recorded defect.
+
+### Sub-allocation fidelity (optional playback feature)
+
+```
+[HRR] SUBALLOC OOB: kernel arg[10] rec 0x... resolves inside a captured segment but in no active tensor block
+```
+
+A high count on one `arg[N]` with `d2h_fail=0` and a later memory fault points
+at a **stale or out-of-bounds device pointer** in the kernel arguments rather
+than host numerics.
+
+## Playback flags
+
+| Flag | Use |
+|------|-----|
+| `--info` | Archive summary; no GPU required |
+| `--info --events` | Full event log (very large) |
+| `--sync-after-launch` | Synchronize after every launch → attribute a fault to a kernel. `run_hrr_replay.sh` adds this by default; `--no-sync` opts out, and `--timing` opts out on its own |
+| `--sync-after-event` | Synchronize after every event → attribute a fault to any event (slowest) |
+| `--single-thread` | Replay on one thread; deterministic ordering |
+| `--repair` | Rewrite a torn archive tail in place |
+
+`--replace-kernel` exists but belongs to fix verification, which is out of scope
+for this skill.
+
+## Capture environment variables
+
+Capture is enabled on the user's unmodified workload:
+
+```bash
+HIP_HRR_CAPTURE_OUTPUT=./capture.hrr <their normal command>
+```
+
+Capture is crash-durable: let the workload fail as it did originally.
+
+## 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 |
+
+A `read_only_page_fault` on a StreamK GEMM points at the edge tile, the
+workspace pointer (`AddressWS`) or the output store (`AddressD`).
+
+## 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 `<prefix>/bin/`, `<prefix>/lib` is added automatically.
+
+## Parser script
+
+```bash
+python3 scripts/analyze_replay_finding.py --help
+```
+
+Outputs a 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`, `archive_complete`, …
diff --git a/skills/hrr-replay-analysis/scripts/analyze_replay_finding.py b/skills/hrr-replay-analysis/scripts/analyze_replay_finding.py
new file mode 100755
index 0000000..38353f2
--- /dev/null
+++ b/skills/hrr-replay-analysis/scripts/analyze_replay_finding.py
@@ -0,0 +1,513 @@
+#!/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]+)"
+)
+# The leading fields of this bracket vary by ROCm build -- some emit `host:`,
+# some start at `GPU index:` -- so anchor on the two fields actually consumed
+# rather than on the whole prefix.
+RE_MEM_FAULT_ERR = re.compile(
+    r"Memory Fault Error \[[^\]]*?faulting addr: (0x[0-9a-fA-F]+), 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"
+)
+# `--info` prints `Complete:     yes (clean shutdown)` or `Complete:     NO (...)`,
+# so the verdict is case-mixed and always followed by an explanation.
+RE_ARCHIVE_COMPLETE = re.compile(r"Complete:\s+(yes|no)\b", re.IGNORECASE)
+# `--info` reports the archive as labelled fields, one per line; a replay run
+# reports the same totals on a single `[HRR] Archive :` line.
+RE_INFO_EVENTS = re.compile(r"^Events:\s+(\d+)\s*$", re.MULTILINE)
+RE_INFO_KERNELS = re.compile(r"^Kernels:\s+(\d+)\s*$", re.MULTILINE)
+# Short archives print no `Kernels:` total and report launches only in the API
+# call-count block. Without this the total stays unknown and the single-kernel
+# inference in finalize() can never fire.
+RE_INFO_LAUNCH_COUNT = re.compile(r"\bhip\w*LaunchKernel\s+(\d+)\b")
+# Rows of the `--info` "Kernel Summary" table: name, grid, block, with an
+# optional leading id (some builds omit the id column). A memory fault can kill
+# the replay before any per-launch attribution reaches the log, in which case
+# the archive's kernel list is the only record of what ran. The name must start
+# with a letter or underscore, which is what keeps a bare id from being read as
+# a symbol now that the id is optional.
+RE_INFO_KERNEL_ROW = re.compile(
+    r"^[ \t]*(?:\d+[ \t]+)?([A-Za-z_][^\s\[]*)[ \t]+\[[\d,\s]+\][ \t]+\[[\d,\s]+\]",
+    re.MULTILINE,
+)
+# `--sync-after-launch` and `--sync-after-event` print one line per replayed
+# event. A GPU fault tears the process down before HRR writes its own Fatal
+# line, so the last of these is often the only record of the failing dispatch.
+RE_EVENT_PROGRESS = re.compile(
+    r"^[ \t]*(?:\[HRR\][ \t]*)?Event (\d+):[ \t]*(\w+)"
+    r"(?:[^\n]*?->[ \t]*Kernel '([^']+)')?",
+    re.MULTILINE,
+)
+# PyTorch/ATen kernels reach the GPU through `<<<>>>` (hipLaunchByPtr) and pass
+# device pointers inside by-value structs. Capture records those (arg encoding
+# `value_kind == 3`) and replay translates them, plus a defensive rescan, so
+# these kernels do replay faithfully on a current build. The detector is still a
+# value-based heuristic, and an archive recorded before it landed carries no
+# such offsets at all, so a fault on one of these symbols earns a caveat in the
+# finding -- not a different verdict.
+RE_ATEN_CHEVRON = re.compile(r"_ZN2at6native|at::native::")
+# Emitted by the archive reader when the on-disk format and this hrr-playback
+# build disagree. Nothing is replayed, so this outranks every other signal.
+RE_VERSION_MISMATCH = re.compile(r"\[HRR\] Version mismatch: file=(\d+) reader=(\d+)")
+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
+    last_event_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_kernel_names: list[str] = field(default_factory=list)
+    archive_complete: str | None = None
+    archive_format_version: int | None = None
+    reader_format_version: int | 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:
+    # Order matters. A GPU memory fault also trips HRR's own abort line and an
+    # HSA_STATUS_ERROR_MEMORY_FAULT, so the fault has to be classified before
+    # the generic abort and hang branches or every fault reads as an API error.
+    if RE_VERSION_MISMATCH.search(text):
+        return "archive_version_mismatch"
+    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_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):
+        return "replay_fatal_api"
+    if RE_HANG.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).lower()
+
+    m = RE_INFO_EVENTS.search(text)
+    if m:
+        finding.archive_events = int(m.group(1))
+    m = RE_INFO_KERNELS.search(text)
+    if m:
+        finding.archive_kernels = int(m.group(1))
+    elif finding.archive_kernels is None:
+        m = RE_INFO_LAUNCH_COUNT.search(text)
+        if m:
+            finding.archive_kernels = int(m.group(1))
+
+    for name in RE_INFO_KERNEL_ROW.findall(text):
+        # The `--info` table truncates long names to its column width, and a
+        # truncated symbol is worse than none: it cannot be looked up or given
+        # to a kernel developer.
+        if name.endswith("...") or name in finding.archive_kernel_names:
+            continue
+        finding.archive_kernel_names.append(name)
+
+    m = RE_VERSION_MISMATCH.search(text)
+    if m:
+        finding.archive_format_version = int(m.group(1))
+        finding.reader_format_version = int(m.group(2))
+        finding.notes.append(
+            f"archive format v{m.group(1)} cannot be read by this hrr-playback "
+            f"(reader v{m.group(2)}); nothing was replayed"
+        )
+
+    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
+
+    last_event = None
+    for m in RE_EVENT_PROGRESS.finditer(text):
+        last_event = m
+    if last_event is not None:
+        # An HRR Fatal line names the failing event exactly; this is only the
+        # last event that started, so it never overrides one.
+        if finding.failing_call_index is None:
+            finding.failing_call_index = int(last_event.group(1))
+            finding.failing_api = last_event.group(2)
+        finding.last_event_kernel = last_event.group(3)
+
+    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))
+
+    finding.kernel_family = _kernel_family(finding.kernel_name)
+    return finding
+
+
+def finalize(finding: Finding, corpus: str) -> Finding:
+    """Set the verdict from every input at once.
+
+    Classification has to run over the whole corpus rather than per input: a
+    `--analyze` run parses the replay log and then the archive `--info` dump,
+    and the `--info` text alone carries no verdict, so classifying per input
+    would let it overwrite the replay's outcome with UNKNOWN.
+    """
+    if RE_VERSION_MISMATCH.search(corpus):
+        finding.outcome = "UNREADABLE"
+    elif RE_MAF.search(corpus) or RE_MEM_FAULT_ERR.search(corpus):
+        finding.outcome = "MAF"
+    elif RE_PASS.search(corpus):
+        finding.outcome = "PASS"
+    elif RE_FAIL.search(corpus):
+        finding.outcome = "FAIL"
+    elif "aborting replay" in corpus or RE_FATAL_EVENT.search(corpus):
+        finding.outcome = "ABORT"
+    else:
+        finding.outcome = "UNKNOWN"
+
+    finding.fault_class = _classify(corpus, finding)
+
+    # A clean replay has no implicated kernel. The archive still lists the
+    # kernels it ran, and a GEMM name matched out of that listing would sit in
+    # the report next to a PASS as if it were a culprit.
+    if finding.fault_class == "replay_pass":
+        finding.kernel_name = None
+        finding.kernel_family = None
+        finding.last_event_kernel = None
+        return finding
+
+    # Under --sync-after-launch the last launch to start is the one that
+    # faulted, so it stands in when the runtime's fault line carried no kernel.
+    if not finding.kernel_name and finding.last_event_kernel:
+        finding.kernel_name = finding.last_event_kernel
+        finding.notes.append(
+            f"kernel name taken from the last launch to start before the fault "
+            f"({finding.last_event_kernel}); the runtime fault line named no "
+            f"kernel. Valid only because the replay ran with "
+            f"--sync-after-launch."
+        )
+
+    # A memory fault can tear the process down before HRR attributes the
+    # failing dispatch, leaving a replay log with no kernel at all. When the
+    # archive holds exactly one kernel, that kernel is the one that faulted.
+    # Anything more than one stays unknown: guessing among several would be
+    # exactly the plausible-but-unevidenced answer the skill forbids.
+    if (
+        not finding.kernel_name
+        and finding.archive_kernels == 1
+        and len(finding.archive_kernel_names) == 1
+    ):
+        finding.kernel_name = finding.archive_kernel_names[0]
+        finding.notes.append(
+            f"kernel name inferred from the archive, which contains exactly one "
+            f"kernel ({finding.kernel_name}); the replay log carried no "
+            f"per-launch attribution. Re-run with --sync-after-launch to confirm "
+            f"the faulting dispatch directly."
+        )
+
+    # A `<<<>>>`-launched ATen kernel passes device pointers inside by-value
+    # structs. Current capture records those offsets and replay translates them,
+    # so this is not automatically a recording artefact, but the detector is a
+    # heuristic and an archive taken before it landed carries no offsets at all.
+    # Both failure modes look exactly like a workload fault, so flag the
+    # ambiguity rather than resolving it either way.
+    if finding.fault_class in ("illegal_memory_access", "read_only_page_fault") and (
+        RE_ATEN_CHEVRON.search(finding.kernel_name or "")
+    ):
+        finding.notes.append(
+            "the faulting kernel is an ATen kernel launched through <<<>>> "
+            "(hipLaunchByPtr), which passes device pointers inside by-value "
+            "structs. Replay translates those via a value-based heuristic, and "
+            "an archive recorded before that support landed has none recorded "
+            "at all, so an untranslated pointer here would fault exactly like a "
+            "workload defect. Confirm against the user's original failure "
+            "signature before reporting this as their bug."
+        )
+
+    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'}`",
+        f"- **Last launch before fault**: `{f.last_event_kernel or 'n/a'}`",
+        "",
+        "## Archive / capture",
+        f"- **Events**: {f.archive_events or 'n/a'}",
+        f"- **Kernels (archive)**: {f.archive_kernels or 'n/a'}",
+        f"- **Kernel names (archive)**: {', '.join(f.archive_kernel_names) 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")
+    chunks: list[str] = []
+    for log_path in args.log:
+        p = Path(log_path)
+        if not p.is_file():
+            finding.notes.append(f"log not found: {p}")
+            continue
+        text = p.read_text(encoding="utf-8", errors="replace")
+        chunks.append(text)
+        parse_text(text, str(p), finding)
+
+    if args.archive:
+        arch = Path(args.archive)
+        info = run_archive_info(arch, args.hrr_playback)
+        if info:
+            chunks.append(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))
+
+    finalize(finding, "\n".join(chunks))
+
+    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-analysis/scripts/run_hrr_replay.sh b/skills/hrr-replay-analysis/scripts/run_hrr_replay.sh
new file mode 100755
index 0000000..6354d5a
--- /dev/null
+++ b/skills/hrr-replay-analysis/scripts/run_hrr_replay.sh
@@ -0,0 +1,175 @@
+#!/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
+NO_SYNC=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 ;;
+    --no-sync) NO_SYNC=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 '
+      # Every value line is prefixed with GPU[N]. Total and Used arrive on
+      # separate lines whose order is not guaranteed, so collect both per
+      # device and subtract at END rather than on whichever line lands last.
+      # "VRAM Total Used Memory" must be tested before "VRAM Total Memory",
+      # since the used line also contains the word Total.
+      /GPU\[/ {
+        id = $1; gsub(/[^0-9]/, "", id)
+        if ($0 ~ /Total Used Memory/) used[id] = $NF
+        else if ($0 ~ /Total Memory/) total[id] = $NF
+      }
+      END {
+        for (id in total)
+          if (id in used) print id, total[id] - used[id]
+      }
+    ')
+    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)"
+
+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
+
+# --info reads the archive only, so it must work on a host with no GPU: resolve
+# a device after this point, never before.
+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
+}
+
+GPU="$(pick_gpu)"
+
+# Default replay serializes the GPU once at the end, so a fault is reported but
+# not attributed to a launch -- the finding then has no failing event and no
+# kernel. Synchronizing after every launch is what makes the last launch line
+# before the fault the culprit. Opt out with --no-sync when throughput matters
+# more than attribution (a long soak). --timing opts out on its own: serializing
+# every launch is exactly what a timing run must not do.
+PLAY_ARGS=(${EXTRA_ARGS[@]+"${EXTRA_ARGS[@]}"})
+if [[ "$NO_SYNC" == "0" ]]; then
+  skip_sync=0
+  for a in ${PLAY_ARGS[@]+"${PLAY_ARGS[@]}"}; do
+    case "$a" in --sync-after-launch|--sync-after-event|--timing) skip_sync=1 ;; esac
+  done
+  if [[ "$skip_sync" == "0" ]]; then
+    PLAY_ARGS+=("--sync-after-launch")
+    echo "[run_hrr_replay] adding --sync-after-launch (disable with --no-sync)" >&2
+  fi
+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
+
+echo "[run_hrr_replay] GPU=$GPU archive=$ARCHIVE" >&2
+set +e
+# hrr-playback is a HIP program, so mask with HIP_VISIBLE_DEVICES. Setting
+# ROCR_VISIBLE_DEVICES as well re-indexes the devices underneath the HIP mask,
+# which can land the replay on a different card than the one selected here.
+HIP_VISIBLE_DEVICES="$GPU" "$HRR_PLAY" "$ARCHIVE" \
+  ${PLAY_ARGS[@]+"${PLAY_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-analysis/skill-card.md b/skills/hrr-replay-analysis/skill-card.md
new file mode 100644
index 0000000..7af7169
--- /dev/null
+++ b/skills/hrr-replay-analysis/skill-card.md
@@ -0,0 +1,17 @@
+# Skill Card
+
+## Description
+
+Records, replays, and analyzes GPU workload behavior on ROCm from HIP Record and
+Replay (HRR) archives, producing a structured finding: fault class (write to a
+read-only page, illegal memory access, NaN/Inf divergence, hang, out of memory,
+clean pass), fault address, failing event index, and the implicated kernel and
+kernel arguments. Covers vLLM, PyTorch and hipBLASLt workloads.
+
+## Owner
+
+AMD
+
+## License
+
+MIT