Repository navigation
[AIRUNTIME-2809] Move HRR into standalone project - #11157
Merged
Merged
Conversation
github-actions
Bot
added
documentation
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
mangupta
force-pushed
the
users/mangupta/hrr-migration
branch
from
September 3, 2026 12:06
62d9eb8 to
b2aa0a4
Compare
mangupta
marked this pull request as ready for review
September 3, 2026 12:08
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
There are confirmed build/behavior issues that can abort replay or break optional builds (notably HIP-language gating for integration tests on non-amd HIP_PLATFORM, unsafe replay generation for hipMemGetDefaultMemPool, and robustness fixes in process output capture and archive parsing).
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Pull request overview
This PR restructures HIP Record & Replay (HRR) by extracting the portable components (wire format header, playback tool/library, generator, docs, skills, and tests) into a standalone CMake project under projects/hrr, while keeping the runtime capture layer inside CLR (projects/clr) where it is built into amdhip64.
Changes:
- Introduces a standalone
projects/hrrCMake project (shared wire-format header,hrr_readerlibrary,hrr-playbacktool, and unit/integration tests). - Updates the CLR capture layer to consume the shared wire-format header from
projects/hrr/includeand adjusts generator default output paths accordingly. - Migrates/updates triage skills + scripts, adds a minimal hip-tests capture smoke sentinel, and wires HRR into labeling/CODEOWNERS and TheRock CI matrix.
File summaries
| File | Description |
|---|---|
| projects/hrr/tools/gen_hrr_api_args.py | Updates generator defaults for split output layout and extends NOOP handling for standalone playback builds. |
| projects/hrr/tests/unit/hrr_recovery_test.cc | Switches to local HRR test harness headers and updated include path for shared wire-format header. |
| projects/hrr/tests/unit/hrr_playback_test.cc | Migrates test macros to HRR-local harness (HRR_TEST_CASE). |
| projects/hrr/tests/unit/hrr_format_test.cc | Migrates includes/macros to standalone HRR test harness and shared header path. |
| projects/hrr/tests/unit/CMakeLists.txt | Adds standalone HRR unit test target and wiring (Catch2 + hip::host + optional hrr-playback dependency). |
| projects/hrr/tests/integration/hrr_spt_workload_test.cc | Migrates includes and HIP-check macro to HRR-local harness. |
| projects/hrr/tests/integration/hrr_roundtrip_test.cc | Migrates process helper and macros to HRR-local harness; adjusts skips and Windows gating. |
| projects/hrr/tests/integration/hrr_repro_workload_test.cc | Migrates includes and HIP-check macro to HRR-local harness for GPU workloads. |
| projects/hrr/tests/integration/CMakeLists.txt | Adds standalone HRR integration test target, HIP-language compilation setup, and per-source default stream flags. |
| projects/hrr/tests/hrr_test_process.hh | Adds standalone subprocess helper replacing hip-tests SpawnProc. |
| projects/hrr/tests/hrr_test_common.hh | Adds standalone Catch2/HIP convenience macros for HRR tests. |
| projects/hrr/tests/CMakeLists.txt | Adds HRR tests superproject setup (Catch2 discovery/fetch + hip package + integration-tests option). |
| projects/hrr/skills/decode-and-triage/SKILL.md | Updates skill documentation paths and playback discovery/build behavior for standalone HRR. |
| projects/hrr/skills/decode-and-triage/scripts/triage_archive.sh | Updates LD_LIBRARY_PATH discovery to match installed/standalone layouts and in-repo ROCR fallback path. |
| projects/hrr/skills/decode-and-triage/scripts/triage_archive.ps1 | Updates playback discovery/build guidance and error text for standalone HRR. |
| projects/hrr/skills/decode-and-triage/scripts/test_check_replay_compat.py | Adjusts tests for new playback path layout and adds mocks for tool detection. |
| projects/hrr/skills/decode-and-triage/scripts/test_analyze_replay_finding.py | Adds unit tests for the new replay finding analyzer script. |
| projects/hrr/skills/decode-and-triage/scripts/replay_docker.sh | Updates overlay mounting semantics and playback/lib discovery variables for container replay. |
| projects/hrr/skills/decode-and-triage/scripts/ensure_playback.sh | Adds standalone playback locator/builder script for Linux. |
| projects/hrr/skills/decode-and-triage/scripts/ensure_playback.ps1 | Adds standalone playback locator/builder script for Windows PowerShell. |
| projects/hrr/skills/decode-and-triage/scripts/check_replay_compat.py | Updates installed-layout lib resolution (playback in <prefix>/bin, libs in <prefix>/lib). |
| projects/hrr/skills/decode-and-triage/scripts/analyze_replay_finding.py | Adds standalone analyzer to parse replay/capture logs into JSON/Markdown findings. |
| projects/hrr/README.md | Updates build and usage docs for standalone playback and emphasizes capture/replay commit alignment. |
| projects/hrr/playback/hrr_reader.h | Updates shared header include path (hrr/hrr_api_args.h). |
| projects/hrr/playback/hrr_reader.cpp | Adds standalone archive reader implementation with recovery logic and event decoding helpers. |
| projects/hrr/playback/hip_playback.h | Updates shared header include path for playback context helpers. |
| projects/hrr/playback/hip_playback.cpp | Updates shared header include path for playback implementation. |
| projects/hrr/playback/CMakeLists.txt | Adds standalone playback build (CPU-only reader always; HIP-dependent tool gated by option). |
| projects/hrr/DESIGN.md | Updates architecture/build/docs to reflect split CLR capture vs standalone HRR playback project. |
| projects/hrr/CMakeLists.txt | Adds top-level standalone HRR project (install wire-format header; conditionally add playback/tests). |
| projects/hip/docs/how-to/debugging.rst | Updates HRR documentation link to point at new projects/hrr/README.md. |
| projects/hip-tests/catch/unit/hrr/hrr_capture_smoke_test.cc | Adds minimal hip-tests sentinel validating capture writes an archive when env var is set. |
| projects/hip-tests/catch/unit/hrr/CMakeLists.txt | Replaces full HRR suite with just the capture smoke sentinel in hip-tests. |
| projects/hip-tests/catch/config/configs/unit/hrr.yaml | Updates hip-tests HRR config to only run the capture smoke test entries. |
| projects/clr/hipamd/src/hrr/hip_capture.cpp | Updates include path to shared wire-format header (hrr/hrr_api_args.h). |
| projects/clr/hipamd/src/hrr/hip_capture_writer.h | Updates include path to shared wire-format header (hrr/hrr_api_args.h). |
| projects/clr/hipamd/src/hrr/hip_capture_generated.cpp | Updates include path and includes newly generated capture shims reflecting updated API list. |
| projects/clr/hipamd/src/hrr/CMakeLists.txt | Adds include path for shared wire-format header from projects/hrr/include and stops building playback here. |
| .github/scripts/therock_matrix.py | Adds HRR to TheRock project matrix with build-only coverage settings. |
| .github/labeler.yml | Adds labeler rule for projects/hrr/**. |
| .github/CODEOWNERS | Adds CODEOWNERS entry for projects/hrr/ and adjusts rccl ownership lines. |
| projects/clr/hipamd/src/hrr/skills/decode-and-triage/scripts/ensure_playback.sh | Removed (migrated to projects/hrr/skills/...). |
| projects/clr/hipamd/src/hrr/skills/decode-and-triage/scripts/ensure_playback.ps1 | Removed (migrated to projects/hrr/skills/...). |
| projects/clr/hipamd/src/hrr/playback/CMakeLists.txt | Removed (playback now builds from standalone projects/hrr/playback). |
Review details
Suppressed comments (2)
projects/hrr/skills/decode-and-triage/scripts/triage_archive.ps1:362
- This error message suggests installing hrr-playback only under HIP_PATH\bin, but ensure_playback.ps1 also searches ROCM_PATH\bin. Mention both to avoid misleading users who set ROCM_PATH only.
projects/hrr/tools/gen_hrr_api_args.py:675 - The generator currently emits hipMemGetDefaultMemPool with a null hipMemLocation* at replay time, which is likely to return hipErrorInvalidValue and abort the entire replay (dispatch_event treats any non-hipSuccess as fatal). Since the generator cannot safely serialize hipMemLocation*, this API should be treated as a NOOP playback handler (similar to other query-only APIs with stale/unsupported pointer args).
- Files reviewed: 42/48 changed files
- Comments generated: 2
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
jayhawk-commits
requested changes
Sep 3, 2026
jayhawk-commits
left a comment
jayhawk-commits
left a comment
Collaborator
There was a problem hiding this comment.
Deferring feedback on the clr changes to hip runtime team. The CODEOWNERS changed rccl entries. Please undo those.
mangupta
force-pushed
the
users/mangupta/hrr-migration
branch
from
September 3, 2026 16:15
b2aa0a4 to
4a237fc
Compare
Contributor
Author
|
Addressed the latest actionable feedback in
The runtime-capture relocation proposal remains deferred to a separate effort. Verification:
|
jayhawk-commits
approved these changes
Sep 4, 2026
jayhawk-commits
left a comment
jayhawk-commits
left a comment
Collaborator
There was a problem hiding this comment.
Approving the .github directory changes
chrispaquot
approved these changes
Sep 7, 2026
Move the portable HRR wire format, playback tool, generator, docs, skills, and HRR-owned tests under projects/hrr while keeping runtime capture compiled into CLR.\n\nReplace the hip-tests HRR suite with a smoke-only capture sentinel, add HRR label/CODEOWNERS metadata, and update the decode/triage helper scripts for standalone hrr-playback builds.\n\nCo-Authored-By: Claude <noreply@anthropic.com>
Updated HRR reviewers
Require playback when building HRR integration tests and preserve the old Windows skip behavior for the driver memcpy round-trip cases.\n\nCo-Authored-By: Claude <noreply@anthropic.com>
Integration tests must load capture-enabled libamdhip64 from the same commit as projects/hrr playback. Without it, SDK capture writes compact wire-format payloads that in-tree hrr-playback rejects. Add HrrCaptureLib.cmake to resolve HRR_CLR_LIB, prepend it to LD_LIBRARY_PATH in capture/replay subprocesses, set RUNPATH on test binaries, fail configure when unset, document the pairing requirement, and add CI that builds amdhip64 before running HRR ctest. Co-authored-by: Cursor <cursoragent@cursor.com>
This reverts commit 9f80886.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…hrr-playback-version #11157 has landed, so the HRR standalone-project layout is now in develop. This branch already carried that layout by merging #11157's branch; develop arrives with the same content as a squash commit, so both sides performed the same renames independently and git pairs them up. Two conflicts, both this branch's own additions that develop has no counterpart for, so both keep this side: * projects/hrr/playback/CMakeLists.txt — the HRR_BUILD_TYPE compile definition and the build-identity stamp (HrrVersionStamp.cmake, a new file here) that --version reports. * projects/hrr/tests/unit/hrr_playback_test.cc — the "hrr_test_process.hh" include. Unit_HRR_Playback_VersionOption spawns hrr-playback as a subprocess, which develop's copy of this file does not do. The hip-tests registration is gone with the migration: hrr_playback_test.cc is already in projects/hrr/tests/unit/CMakeLists.txt, and the suite runs under CTest as hrr-unit-tests, so the hrr.yaml entry this branch used to add is no longer needed. Note that its amd_windows/amd_linux/amd_wsl disable goes with it.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…ersion' into users/danielbolsson/hrr-api-ids Carries develop, now containing the merged HRR standalone-project layout (#11157), up the stack from the parent branch. No conflicts. Re-running gen_hrr_api_args.py over the merged generator reproduces hrr_api_args.h, hip_capture_generated.cpp and hip_playback_generated.cpp byte-for-byte, so the merged generated files are consistent with their source.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
… into users/danielbolsson/hrr-regions Carries develop, now containing the merged HRR standalone-project layout (#11157), up the stack from the parent branch. No conflicts. gen_hrr_api_args.py re-run over the merged generator reproduces the generated header and shims unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
… into users/danielbolsson/hrr-regions-memcpy Carries develop, now containing the merged HRR standalone-project layout (#11157), up the stack from the parent branch. No conflicts.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…memcpy' into users/danielbolsson/hrr-test-harness Carries develop, now containing the merged HRR standalone-project layout (#11157), up the stack from the parent branch. No conflicts: the shared roundtrip harness this branch folded into projects/hrr/tests/hrr_test_common.hh merges unchanged against develop's copy.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…ness' into users/danielbolsson/hrr-api-matrix Carries develop, now containing the merged HRR standalone-project layout (#11157), up the stack from the parent branch. No conflicts: the API-matrix tests, their registration in projects/hrr/tests/integration/CMakeLists.txt, and the SpawnProc::runWithTimeout the watchdog is rebuilt on all merge unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…ix' into users/danielbolsson/hrr-api-payloads Carries develop, now containing the merged HRR standalone-project layout (#11157), up the stack from the parent branch. No conflicts. gen_hrr_api_args.py re-run over the merged generator reproduces the generated header and both shims unchanged, including this branch's _AMDHIP_DEF_IN fix (the Windows export list stays at projects/clr/hipamd/src/amdhip.def.in, next to the amdhip64 sources it is linked into, rather than moving with the tool).
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…oads' into users/danielbolsson/hrr-scope-exclusions Carries develop, now containing the merged HRR standalone-project layout (#11157), up the stack from the parent branch. No conflicts: the UNREPLAYABLE replay class and its marker in projects/hrr/tests/hrr_test_common.hh merge unchanged, and re-running gen_hrr_api_args.py reproduces the generated files.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
…clusions' into users/danielbolsson/hrr-kernel-args Carries develop, now containing the merged HRR standalone-project layout (#11157), up the stack from the parent branch. No conflicts: hrr_kernel_args_test.cc and its HRR_GPU_WORKLOAD_SOURCES entry in projects/hrr/tests/integration/CMakeLists.txt merge unchanged, and re-running gen_hrr_api_args.py reproduces the generated files.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
#11157 has landed, so this PR's base auto-retargeted from the migration branch back to develop. Merging develop in replaces the migration content this branch carried from the (now deleted) migration branch with develop's own copy. Two conflicts, both this branch's own additions where develop has the pre-existing form, so both keep this side: * hrr_repro_workload_test.cc — the HSA-vs-Windows split for the never-written source allocation, and the <hsa/hsa_ext_amd.h> include it needs. Develop's ENABLE_ADDRESS_SANITIZER guard around the zeroed-allocation precondition is untouched and still in place. * hrr_workload_test.cc — the two hipDeviceSynchronize() drains after hipDrvMemcpy3D and hipDrvMemcpy2DUnaligned.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
#11157 has landed, so this PR's base auto-retargeted from the migration branch back to develop. Merging develop in replaces the migration content this branch carried from the (now deleted) migration branch with develop's own copy. One conflict, in hrr_workload_test.cc: this branch guards the hipExtStreamWriteValueIncrement call with #ifdef so the suite still builds against SDK headers that predate the extension, where develop calls it unconditionally. The guard is kept. hip_capture.cpp and hip_global.cpp stay in CLR: the capture layer compiles into amdhip64 and did not move with the migration.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
#11157 has landed, so this PR's base auto-retargeted from the migration branch back to develop. Merging develop in replaces the migration content this branch was created on with develop's own copy. Four conflicts, all the same shape: this branch is the migration plus #9758's fixes, develop is the migration alone, and the merge base predates both. Every hunk keeps this side, which contains develop's line and the fix on top — analyze_replay_finding.py (20 hunks), test_analyze_replay_finding.py (2), ensure_playback.sh (1) and triage_archive.sh (1). The resulting diff against develop is unchanged at 15 files, +1061/-27, all under projects/hrr/skills/. All 48 unit tests pass and both shell scripts pass bash -n.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
#11157 has landed, so the HRR standalone-project layout is now in develop and this branch's tests move with it into projects/hrr/tests/integration. Resolution notes: * hrr_workload_test.cc — this branch's removal of the _spt exercise blocks is kept in all seven conflicting regions. They were converted into standalone hidden [.][hrr-direct] cases in hrr_spt_workload_test.cc, which is where develop's copies of them would otherwise be duplicated. No _spt call remains in hrr_workload_test.cc. * hrr_roundtrip_test.cc — this branch's four new _spt roundtrip cases are kept. * The relocated sources are rewritten onto the HRR harness: HIP_TEST_CASE -> HRR_TEST_CASE (5) and HIP_CHECK -> HRR_HIP_CHECK (122 across the workload and SPT translation units). * The 32 hrr.yaml entries and the catch/unit/hrr/CMakeLists.txt edit are dropped in favour of projects/hrr/tests/integration/CMakeLists.txt, which already builds hrr_spt_workload_test.cc with the per-thread-default-stream compile option this conversion depends on. Develop's hip-tests capture sentinel is taken unchanged.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
#11157 has landed. This branch's DESIGN.md edit follows the file to projects/hrr/DESIGN.md automatically. One conflict, in projects/hip/docs/how-to/debugging.rst: both sides rewrote the same introductory paragraph. This branch's version is kept — it explains what the archive makes possible rather than restating what HRR is, and it cross-references :doc:`/how-to/hip_record_replay`, the how-to page this PR adds and registers in the Sphinx TOC. Develop's version pointed at projects/hrr/README.md instead, which the new page supersedes.
danielbolsson
added a commit
that referenced
this pull request
Sep 9, 2026
The merge that brought the HRR migration (#11157) into this branch renamed gen_hrr_api_args.py out of projects/clr and, resolving that against this branch's own edits, kept hipDrvGraphMemcpyNodeSetParams in NOOP_PLAYBACK_APIS as well as adding it to MANUAL_PLAYBACK_APIS. An API is meant to be in exactly one classification set; the handler emitter tests NOOP before MANUAL, so replay got the no-op stub and the hand-written handler in hip_playback.cpp was never reachable. Unit_HRR_ApiMatrix_T2_Roundtrip observes NOOP where it expects REAL. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
danielbolsson
added a commit
that referenced
this pull request
Sep 11, 2026
…inding (#11342) ## Summary The `decode-and-triage` skill was run end to end against real archives on a gfx950 host, and reviewed against the HRR runbooks. It did not survive either exercise: on a crashing archive it produced no finding at all, and on the archives where it did produce one, the verdict was often wrong. This is a skill-only change. No capture or playback source is touched. It is based on [#11157](#11157), which moves the skill to `projects/hrr/skills/`, so the diff here is the fixes alone. The base returns to `develop` once that one lands. > **This replaces [#9758](#9758 That PR was stacked on [#10535](#10535), which moves the same skill files to `projects/clr/.claude/skills/` instead. #10535 and #11157 are a rename/rename conflict over the same sources and cannot both land; #11157's destination is the one that puts the skill under `@ROCm/hrr-reviewers`, so #10535 is closed and the fixes are replayed here. > > Two files needed real merging rather than a path rewrite — `ensure_playback.sh` and `triage_archive.sh`; see the commit message for what was kept from each side. All 48 unit tests pass. ## What was wrong | # | Symptom on a real archive | Cause | |---|---|---| | 1 | Native replay impossible even with `hrr-playback` supplied | `export_playback_env` ended in a trailing `&&`, which under `set -e` became the function's exit status and killed the script before it printed the path | | 2 | A replay that faulted produced a log and no finding | `run_native_replay` re-armed `set -e` inside itself, cancelling the `set +e` the caller wraps it in, so any non-zero replay exit ended the script before the analyzer ran | | 3 | A hung replay produced nothing and never returned | no bound on the replay, so the analyzer was never reached | | 4 | Metadata-only triage reported no events, kernels or completeness | `--no-replay` never resolved a playback binary, so `--info` never ran | | 5 | A packaged playback failed to load its own libamdhip64 | the library path handled only a playback inside a CLR build tree, not `bin/` and `lib/` siblings | | 6 | A queue abort reported as a hang with an unknown kernel | `hang` was reachable from any `HSA_STATUS_ERROR_*`, and the abort line carries a kernel but no faulting address, so the memory-fault regex could not see it | | 7 | A read-only page fault reported as `replay_fatal_api` | the generic abort branch was tested before the memory-fault branch | | 8 | A clean replay reported a hipBLASLt GEMM as its kernel | the name was matched out of the archive's own kernel listing | | 9 | A fault with no attribution reported no kernel at all | the `--info` kernel table was not parsed, so a single-kernel archive named nothing | | 10 | The kernel name dropped on some ROCm builds | `RE_MEM_FAULT_ERR` required a `host:` field that not every build emits | | 11 | Replay picked the busiest GPU and reported its full VRAM as free | the awk matched `Used Memory` and `Total Memory` as separate patterns, but the used line reads `VRAM Total Used Memory` and matches both | | 12 | Replay could land on a device other than the one picked | the native path masked with `ROCR_VISIBLE_DEVICES`, which re-indexes underneath a HIP mask; the Windows and container paths already used `HIP_VISIBLE_DEVICES` | | 13 | `Complete:`, `Events:`, `Kernels:` and both format versions unreported | uppercase-only matcher, unparsed labelled lines, versions recorded only in prose | | 14 | `Last progress kernel` always empty | assigned unconditionally, so the `--info` pass erased what the replay log had recorded | | 15 | A replay killed by a signal reported no class and still named a kernel | the log ends where the process died, so the truncation read as insufficient signal and the kernel was then inferred from a single-kernel archive | Two of these deserve a note. **Nothing beats a wrong verdict for cost.** Defects 6 and 8 are the expensive ones: one reports a hang that never happened and drops the kernel the runtime had just named, the other prints a GEMM next to a `PASS`, where it reads as a culprit. Both send a kernel developer after a defect the run never observed. **A hang is now bounded, not detected.** The replay is stopped after `HRR_REPLAY_TIMEOUT` seconds, default 1800, and the runner records that in the log so the finding says `hang` rather than nothing. The bound goes around the playback process, not around a container that may wrap it: a timeout on the container leaves the replay running inside it. ## What changed in behaviour - Replay runs with `--sync-after-launch` by default; `--no-sync` opts out. Measured cost on a real vLLM capture of 13233 kernels: 41.7 s against 39.6 s. - A clean replay implicates no kernel. A fault with no attribution in the log takes the kernel from an archive holding exactly one, marked in the finding as inferred; more than one stays unknown, and names truncated by the `--info` column width are never used. - New outcome `HANG`. `HSA_STATUS_ERROR_MEMORY_FAULT` classifies as `illegal_memory_access`, the other HSA statuses as `replay_aborted`, and the raw status goes in the notes. - A replay process that dies on a signal is its own finding: outcome `ABORT`, fault class `replay_crashed`, no kernel implicated, because nothing about the workload was established. The runner records the signal in the log the way it already records a timeout, and the classifier ranks that marker last, so a log which explains its own fault is still believed first: a memory fault also dies on a signal. - A fault on a `<<<>>>`-launched ATen kernel keeps its class and gains a caveat. Capture records the pointers these kernels pass inside by-value structs and replay translates them, so such a fault can be genuine and reclassifying it would suppress a real finding; but the translation is a heuristic and older archives carry no offsets, so the skill asks for the original failure signature instead of resolving the ambiguity. ## Testing Fixtures are unedited tool output recorded on a gfx950 host, so a change in what `hrr-playback` prints fails a test rather than silently degrading a finding. They sit under the skill's `evals/fixtures/`, with the recorded evidence rather than beside the scripts that read it. Device selection, playback discovery and the library layout are covered without a GPU by lifting the shell functions out of the script and running them against recorded input. Unit suite: 24 tests on `develop`, 48 here. On this head 46 pass, 1 skips below bash 4.4, and 1 fails: `test_probe_host_reads_comgr`, which fails identically on the unmodified `develop` copies of the same two files and is untouched here. Each fix was also removed in turn to confirm the test guarding it fails without it, since a test that passes either way proves nothing. The exit-code plumbing behind the new class was checked against a real signal death through the same pipe-to-`tee` shape the runner uses: `PIPESTATUS[0]` carries 139 for `SIGSEGV` and 134 for `SIGABRT`. End to end with `triage_archive.sh` on a gfx950 host: | Archive | Result | |---|---| | vLLM inference capture, format v4 | 13233 kernels replayed, 32 device-to-host checks exact, `PASS` / `replay_pass`, no kernel implicated, `Complete: yes`, 185469 events | | Capture whose workload faults | fault reproduced, run survives it, `MAF` / `illegal_memory_access` with fault address, kernel `crash_oob` and the dispatch grid. Previously: no finding at all | | PyTorch capture indexing past its embedding table | replay reproduces the capture-time failure exactly, same `HSA_STATUS_ERROR_EXCEPTION` on the same kernel. `ABORT` / `replay_aborted`, full `indexSelectSmallIndex` signature, ATen caveat. Previously: `UNKNOWN` / `hang`, kernel unknown | | Hanging capture, 45 s bound | `HANG` / `hang` with a finding. Previously: ran until killed, no finding | | Archive the playback build cannot read | `version_mismatch` with archive version 3 and reader version 4, no kernel invented | | Metadata-only, no GPU | 5 events, 1 kernel, `crash_oob`, `Complete: no`. Previously: every field `n/a` | ## Known gaps - Multi-process archives are untested: none of the 13 archives available here has more than one `pid-*` directory, and `load_capture_metadata` reads only `<archive>/manifest.json`, so an archive root loses the preflight and the capture fields. - A hang is inferred from the bound firing, not from any signal the replay emits; a replay legitimately slower than the timeout will be reported as a hang. - No live `hrr-playback` was made to die on a signal, so that path rests on the runner and analyser tests plus the exit-code check above rather than on a recorded run. - Windows was not exercised. `triage_archive.ps1` gains no changes here beyond already using the HIP mask, and the PowerShell runner records no equivalent marker yet. ## Issue Tracking JIRA ID: AIRUNTIME-2695 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Maneesh Gupta <maneesh.gupta@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
danielbolsson
added a commit
that referenced
this pull request
Sep 16, 2026
…#11686) ## Summary `projects/hrr/tests/integration/hrr_repro_workload_test.cc` calls `hsa_amd_memory_fill` directly, but the `hrr-integration-tests` target links no ROCr. `hip::host` resolves ROCr for its own needs and does not carry it in its link interface, so the call has no definition at link time. Name the runtime on the target, guarded by `NOT WIN32` because Windows has no ROCr and the call site compiles out there. ## Why this was not caught `HRR_BUILD_INTEGRATION_TESTS` defaults to `OFF`, so nothing in CI had ever compiled the target. The migration in #11157 landed the `hsa_amd_memory_fill` call on develop on 2026-09-09; this link fix was written on 2026-09-10 and stayed on a topic branch. ## How it surfaced ROCm/TheRock#7669 turns `HRR_BUILD_INTEGRATION_TESTS` on, making it the first build anywhere to compile the target. It fails against the current develop gitlink: ``` ld.lld: error: undefined symbol: hsa_amd_memory_fill >>> referenced by hrr_repro_workload_test.cc >>> tests/integration/CMakeFiles/hrr-integration-tests.dir/hrr_repro_workload_test.cc.o ``` That CI failure cannot be fixed inside TheRock: it builds rocm-systems at a pinned gitlink, so this has to land here first and the pin move past it. ## Validation - The integration target links and the full suite passes with this change on `users/danielbolsson/hrr-development`: 23 unit and 70 integration tests green locally on gfx950, and green across gfx94x, gfx950, gfx103x, gfx110x, gfx1150, gfx1151 and gfx120x in CI. - No functional change. One CMake block, build configuration only. Made with [Cursor](https://cursor.com) Co-authored-by: Cursor <cursoragent@cursor.com>
abchoudh-amd
pushed a commit
that referenced
this pull request
Sep 16, 2026
…inding (#11342) ## Summary The `decode-and-triage` skill was run end to end against real archives on a gfx950 host, and reviewed against the HRR runbooks. It did not survive either exercise: on a crashing archive it produced no finding at all, and on the archives where it did produce one, the verdict was often wrong. This is a skill-only change. No capture or playback source is touched. It is based on [#11157](#11157), which moves the skill to `projects/hrr/skills/`, so the diff here is the fixes alone. The base returns to `develop` once that one lands. > **This replaces [#9758](#9758 That PR was stacked on [#10535](#10535), which moves the same skill files to `projects/clr/.claude/skills/` instead. #10535 and #11157 are a rename/rename conflict over the same sources and cannot both land; #11157's destination is the one that puts the skill under `@ROCm/hrr-reviewers`, so #10535 is closed and the fixes are replayed here. > > Two files needed real merging rather than a path rewrite — `ensure_playback.sh` and `triage_archive.sh`; see the commit message for what was kept from each side. All 48 unit tests pass. ## What was wrong | # | Symptom on a real archive | Cause | |---|---|---| | 1 | Native replay impossible even with `hrr-playback` supplied | `export_playback_env` ended in a trailing `&&`, which under `set -e` became the function's exit status and killed the script before it printed the path | | 2 | A replay that faulted produced a log and no finding | `run_native_replay` re-armed `set -e` inside itself, cancelling the `set +e` the caller wraps it in, so any non-zero replay exit ended the script before the analyzer ran | | 3 | A hung replay produced nothing and never returned | no bound on the replay, so the analyzer was never reached | | 4 | Metadata-only triage reported no events, kernels or completeness | `--no-replay` never resolved a playback binary, so `--info` never ran | | 5 | A packaged playback failed to load its own libamdhip64 | the library path handled only a playback inside a CLR build tree, not `bin/` and `lib/` siblings | | 6 | A queue abort reported as a hang with an unknown kernel | `hang` was reachable from any `HSA_STATUS_ERROR_*`, and the abort line carries a kernel but no faulting address, so the memory-fault regex could not see it | | 7 | A read-only page fault reported as `replay_fatal_api` | the generic abort branch was tested before the memory-fault branch | | 8 | A clean replay reported a hipBLASLt GEMM as its kernel | the name was matched out of the archive's own kernel listing | | 9 | A fault with no attribution reported no kernel at all | the `--info` kernel table was not parsed, so a single-kernel archive named nothing | | 10 | The kernel name dropped on some ROCm builds | `RE_MEM_FAULT_ERR` required a `host:` field that not every build emits | | 11 | Replay picked the busiest GPU and reported its full VRAM as free | the awk matched `Used Memory` and `Total Memory` as separate patterns, but the used line reads `VRAM Total Used Memory` and matches both | | 12 | Replay could land on a device other than the one picked | the native path masked with `ROCR_VISIBLE_DEVICES`, which re-indexes underneath a HIP mask; the Windows and container paths already used `HIP_VISIBLE_DEVICES` | | 13 | `Complete:`, `Events:`, `Kernels:` and both format versions unreported | uppercase-only matcher, unparsed labelled lines, versions recorded only in prose | | 14 | `Last progress kernel` always empty | assigned unconditionally, so the `--info` pass erased what the replay log had recorded | | 15 | A replay killed by a signal reported no class and still named a kernel | the log ends where the process died, so the truncation read as insufficient signal and the kernel was then inferred from a single-kernel archive | Two of these deserve a note. **Nothing beats a wrong verdict for cost.** Defects 6 and 8 are the expensive ones: one reports a hang that never happened and drops the kernel the runtime had just named, the other prints a GEMM next to a `PASS`, where it reads as a culprit. Both send a kernel developer after a defect the run never observed. **A hang is now bounded, not detected.** The replay is stopped after `HRR_REPLAY_TIMEOUT` seconds, default 1800, and the runner records that in the log so the finding says `hang` rather than nothing. The bound goes around the playback process, not around a container that may wrap it: a timeout on the container leaves the replay running inside it. ## What changed in behaviour - Replay runs with `--sync-after-launch` by default; `--no-sync` opts out. Measured cost on a real vLLM capture of 13233 kernels: 41.7 s against 39.6 s. - A clean replay implicates no kernel. A fault with no attribution in the log takes the kernel from an archive holding exactly one, marked in the finding as inferred; more than one stays unknown, and names truncated by the `--info` column width are never used. - New outcome `HANG`. `HSA_STATUS_ERROR_MEMORY_FAULT` classifies as `illegal_memory_access`, the other HSA statuses as `replay_aborted`, and the raw status goes in the notes. - A replay process that dies on a signal is its own finding: outcome `ABORT`, fault class `replay_crashed`, no kernel implicated, because nothing about the workload was established. The runner records the signal in the log the way it already records a timeout, and the classifier ranks that marker last, so a log which explains its own fault is still believed first: a memory fault also dies on a signal. - A fault on a `<<<>>>`-launched ATen kernel keeps its class and gains a caveat. Capture records the pointers these kernels pass inside by-value structs and replay translates them, so such a fault can be genuine and reclassifying it would suppress a real finding; but the translation is a heuristic and older archives carry no offsets, so the skill asks for the original failure signature instead of resolving the ambiguity. ## Testing Fixtures are unedited tool output recorded on a gfx950 host, so a change in what `hrr-playback` prints fails a test rather than silently degrading a finding. They sit under the skill's `evals/fixtures/`, with the recorded evidence rather than beside the scripts that read it. Device selection, playback discovery and the library layout are covered without a GPU by lifting the shell functions out of the script and running them against recorded input. Unit suite: 24 tests on `develop`, 48 here. On this head 46 pass, 1 skips below bash 4.4, and 1 fails: `test_probe_host_reads_comgr`, which fails identically on the unmodified `develop` copies of the same two files and is untouched here. Each fix was also removed in turn to confirm the test guarding it fails without it, since a test that passes either way proves nothing. The exit-code plumbing behind the new class was checked against a real signal death through the same pipe-to-`tee` shape the runner uses: `PIPESTATUS[0]` carries 139 for `SIGSEGV` and 134 for `SIGABRT`. End to end with `triage_archive.sh` on a gfx950 host: | Archive | Result | |---|---| | vLLM inference capture, format v4 | 13233 kernels replayed, 32 device-to-host checks exact, `PASS` / `replay_pass`, no kernel implicated, `Complete: yes`, 185469 events | | Capture whose workload faults | fault reproduced, run survives it, `MAF` / `illegal_memory_access` with fault address, kernel `crash_oob` and the dispatch grid. Previously: no finding at all | | PyTorch capture indexing past its embedding table | replay reproduces the capture-time failure exactly, same `HSA_STATUS_ERROR_EXCEPTION` on the same kernel. `ABORT` / `replay_aborted`, full `indexSelectSmallIndex` signature, ATen caveat. Previously: `UNKNOWN` / `hang`, kernel unknown | | Hanging capture, 45 s bound | `HANG` / `hang` with a finding. Previously: ran until killed, no finding | | Archive the playback build cannot read | `version_mismatch` with archive version 3 and reader version 4, no kernel invented | | Metadata-only, no GPU | 5 events, 1 kernel, `crash_oob`, `Complete: no`. Previously: every field `n/a` | ## Known gaps - Multi-process archives are untested: none of the 13 archives available here has more than one `pid-*` directory, and `load_capture_metadata` reads only `<archive>/manifest.json`, so an archive root loses the preflight and the capture fields. - A hang is inferred from the bound firing, not from any signal the replay emits; a replay legitimately slower than the timeout will be reported as a hang. - No live `hrr-playback` was made to die on a signal, so that path rests on the runner and analyser tests plus the exit-code check above rather than on a recorded run. - Windows was not exercised. `triage_archive.ps1` gains no changes here beyond already using the HIP mask, and the PowerShell runner records no equivalent marker yet. ## Issue Tracking JIRA ID: AIRUNTIME-2695 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Maneesh Gupta <maneesh.gupta@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
dlamd1dai
pushed a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 16, 2026
## Summary - Move HRR portable layer (wire format, playback, generator, docs, skills, tests) to `projects/hrr`. - Keep runtime capture in CLR and add a hip-tests capture smoke sentinel. - Add HRR label/CODEOWNERS metadata and update docs/skills for standalone playback builds. - Add build-only TheRock CI matrix coverage for `projects/hrr` changes and make GPU integration tests explicit opt-in. ## Validation - Static checks: no flat `hrr_api_args.h` includes; no stale old playback paths; no conflict markers or pycache artifacts. - TheRock CI matrix tests: 32 passed. - HRR script tests: 24 passed. - Generator idempotency: checksum-stable. - HRR unit CTest passed. - HRR integration and `hrr-playback` build passed locally with HIP compiler setup. - hip-tests `Unit_HRR_CaptureSmoke` passed locally. ## Pending / follow-up - Full CLR `amdhip64` build needs a matching ROCm/HSA toolchain; the local attempt hit an unrelated HSA header mismatch. - HRR integration runtime parity/datacenter validation remains environment-dependent. - TheRock orchestration/packaging is paired in ROCm/TheRock#7669. - Runtime capture relocation remains intentionally out of scope for this PR. ## Related - ROCm/TheRock#7669 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
pushed a commit
to dlamd1dai/rocm-systems
that referenced
this pull request
Sep 16, 2026
…inding (ROCm#11342) ## Summary The `decode-and-triage` skill was run end to end against real archives on a gfx950 host, and reviewed against the HRR runbooks. It did not survive either exercise: on a crashing archive it produced no finding at all, and on the archives where it did produce one, the verdict was often wrong. This is a skill-only change. No capture or playback source is touched. It is based on [ROCm#11157](ROCm#11157), which moves the skill to `projects/hrr/skills/`, so the diff here is the fixes alone. The base returns to `develop` once that one lands. > **This replaces [ROCm#9758](ROCm#9758 That PR was stacked on [ROCm#10535](ROCm#10535), which moves the same skill files to `projects/clr/.claude/skills/` instead. ROCm#10535 and ROCm#11157 are a rename/rename conflict over the same sources and cannot both land; ROCm#11157's destination is the one that puts the skill under `@ROCm/hrr-reviewers`, so ROCm#10535 is closed and the fixes are replayed here. > > Two files needed real merging rather than a path rewrite — `ensure_playback.sh` and `triage_archive.sh`; see the commit message for what was kept from each side. All 48 unit tests pass. ## What was wrong | # | Symptom on a real archive | Cause | |---|---|---| | 1 | Native replay impossible even with `hrr-playback` supplied | `export_playback_env` ended in a trailing `&&`, which under `set -e` became the function's exit status and killed the script before it printed the path | | 2 | A replay that faulted produced a log and no finding | `run_native_replay` re-armed `set -e` inside itself, cancelling the `set +e` the caller wraps it in, so any non-zero replay exit ended the script before the analyzer ran | | 3 | A hung replay produced nothing and never returned | no bound on the replay, so the analyzer was never reached | | 4 | Metadata-only triage reported no events, kernels or completeness | `--no-replay` never resolved a playback binary, so `--info` never ran | | 5 | A packaged playback failed to load its own libamdhip64 | the library path handled only a playback inside a CLR build tree, not `bin/` and `lib/` siblings | | 6 | A queue abort reported as a hang with an unknown kernel | `hang` was reachable from any `HSA_STATUS_ERROR_*`, and the abort line carries a kernel but no faulting address, so the memory-fault regex could not see it | | 7 | A read-only page fault reported as `replay_fatal_api` | the generic abort branch was tested before the memory-fault branch | | 8 | A clean replay reported a hipBLASLt GEMM as its kernel | the name was matched out of the archive's own kernel listing | | 9 | A fault with no attribution reported no kernel at all | the `--info` kernel table was not parsed, so a single-kernel archive named nothing | | 10 | The kernel name dropped on some ROCm builds | `RE_MEM_FAULT_ERR` required a `host:` field that not every build emits | | 11 | Replay picked the busiest GPU and reported its full VRAM as free | the awk matched `Used Memory` and `Total Memory` as separate patterns, but the used line reads `VRAM Total Used Memory` and matches both | | 12 | Replay could land on a device other than the one picked | the native path masked with `ROCR_VISIBLE_DEVICES`, which re-indexes underneath a HIP mask; the Windows and container paths already used `HIP_VISIBLE_DEVICES` | | 13 | `Complete:`, `Events:`, `Kernels:` and both format versions unreported | uppercase-only matcher, unparsed labelled lines, versions recorded only in prose | | 14 | `Last progress kernel` always empty | assigned unconditionally, so the `--info` pass erased what the replay log had recorded | | 15 | A replay killed by a signal reported no class and still named a kernel | the log ends where the process died, so the truncation read as insufficient signal and the kernel was then inferred from a single-kernel archive | Two of these deserve a note. **Nothing beats a wrong verdict for cost.** Defects 6 and 8 are the expensive ones: one reports a hang that never happened and drops the kernel the runtime had just named, the other prints a GEMM next to a `PASS`, where it reads as a culprit. Both send a kernel developer after a defect the run never observed. **A hang is now bounded, not detected.** The replay is stopped after `HRR_REPLAY_TIMEOUT` seconds, default 1800, and the runner records that in the log so the finding says `hang` rather than nothing. The bound goes around the playback process, not around a container that may wrap it: a timeout on the container leaves the replay running inside it. ## What changed in behaviour - Replay runs with `--sync-after-launch` by default; `--no-sync` opts out. Measured cost on a real vLLM capture of 13233 kernels: 41.7 s against 39.6 s. - A clean replay implicates no kernel. A fault with no attribution in the log takes the kernel from an archive holding exactly one, marked in the finding as inferred; more than one stays unknown, and names truncated by the `--info` column width are never used. - New outcome `HANG`. `HSA_STATUS_ERROR_MEMORY_FAULT` classifies as `illegal_memory_access`, the other HSA statuses as `replay_aborted`, and the raw status goes in the notes. - A replay process that dies on a signal is its own finding: outcome `ABORT`, fault class `replay_crashed`, no kernel implicated, because nothing about the workload was established. The runner records the signal in the log the way it already records a timeout, and the classifier ranks that marker last, so a log which explains its own fault is still believed first: a memory fault also dies on a signal. - A fault on a `<<<>>>`-launched ATen kernel keeps its class and gains a caveat. Capture records the pointers these kernels pass inside by-value structs and replay translates them, so such a fault can be genuine and reclassifying it would suppress a real finding; but the translation is a heuristic and older archives carry no offsets, so the skill asks for the original failure signature instead of resolving the ambiguity. ## Testing Fixtures are unedited tool output recorded on a gfx950 host, so a change in what `hrr-playback` prints fails a test rather than silently degrading a finding. They sit under the skill's `evals/fixtures/`, with the recorded evidence rather than beside the scripts that read it. Device selection, playback discovery and the library layout are covered without a GPU by lifting the shell functions out of the script and running them against recorded input. Unit suite: 24 tests on `develop`, 48 here. On this head 46 pass, 1 skips below bash 4.4, and 1 fails: `test_probe_host_reads_comgr`, which fails identically on the unmodified `develop` copies of the same two files and is untouched here. Each fix was also removed in turn to confirm the test guarding it fails without it, since a test that passes either way proves nothing. The exit-code plumbing behind the new class was checked against a real signal death through the same pipe-to-`tee` shape the runner uses: `PIPESTATUS[0]` carries 139 for `SIGSEGV` and 134 for `SIGABRT`. End to end with `triage_archive.sh` on a gfx950 host: | Archive | Result | |---|---| | vLLM inference capture, format v4 | 13233 kernels replayed, 32 device-to-host checks exact, `PASS` / `replay_pass`, no kernel implicated, `Complete: yes`, 185469 events | | Capture whose workload faults | fault reproduced, run survives it, `MAF` / `illegal_memory_access` with fault address, kernel `crash_oob` and the dispatch grid. Previously: no finding at all | | PyTorch capture indexing past its embedding table | replay reproduces the capture-time failure exactly, same `HSA_STATUS_ERROR_EXCEPTION` on the same kernel. `ABORT` / `replay_aborted`, full `indexSelectSmallIndex` signature, ATen caveat. Previously: `UNKNOWN` / `hang`, kernel unknown | | Hanging capture, 45 s bound | `HANG` / `hang` with a finding. Previously: ran until killed, no finding | | Archive the playback build cannot read | `version_mismatch` with archive version 3 and reader version 4, no kernel invented | | Metadata-only, no GPU | 5 events, 1 kernel, `crash_oob`, `Complete: no`. Previously: every field `n/a` | ## Known gaps - Multi-process archives are untested: none of the 13 archives available here has more than one `pid-*` directory, and `load_capture_metadata` reads only `<archive>/manifest.json`, so an archive root loses the preflight and the capture fields. - A hang is inferred from the bound firing, not from any signal the replay emits; a replay legitimately slower than the timeout will be reported as a hang. - No live `hrr-playback` was made to die on a signal, so that path rests on the runner and analyser tests plus the exit-code check above rather than on a recorded run. - Windows was not exercised. `triage_archive.ps1` gains no changes here beyond already using the HIP mask, and the PowerShell runner records no equivalent marker yet. ## Issue Tracking JIRA ID: AIRUNTIME-2695 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Maneesh Gupta <maneesh.gupta@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
anugodavar
pushed a commit
that referenced
this pull request
Sep 17, 2026
## Summary - Move HRR portable layer (wire format, playback, generator, docs, skills, tests) to `projects/hrr`. - Keep runtime capture in CLR and add a hip-tests capture smoke sentinel. - Add HRR label/CODEOWNERS metadata and update docs/skills for standalone playback builds. - Add build-only TheRock CI matrix coverage for `projects/hrr` changes and make GPU integration tests explicit opt-in. ## Validation - Static checks: no flat `hrr_api_args.h` includes; no stale old playback paths; no conflict markers or pycache artifacts. - TheRock CI matrix tests: 32 passed. - HRR script tests: 24 passed. - Generator idempotency: checksum-stable. - HRR unit CTest passed. - HRR integration and `hrr-playback` build passed locally with HIP compiler setup. - hip-tests `Unit_HRR_CaptureSmoke` passed locally. ## Pending / follow-up - Full CLR `amdhip64` build needs a matching ROCm/HSA toolchain; the local attempt hit an unrelated HSA header mismatch. - HRR integration runtime parity/datacenter validation remains environment-dependent. - TheRock orchestration/packaging is paired in ROCm/TheRock#7669. - Runtime capture relocation remains intentionally out of scope for this PR. ## Related - ROCm/TheRock#7669 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com>
andraghetti
changed the title
1 task done
danielbolsson
added a commit
to ROCm/TheRock
that referenced
this pull request
Oct 28, 2026
…8965) ## Motivation HIP Record and Replay (HRR) captures HIP API calls from inside `amdhip64` and replays them with `hrr-playback`. Since HRR moved out of CLR into `rocm-systems/projects/hrr` (ROCm/rocm-systems#11157), TheRock no longer builds or ships `hrr-playback`. This PR: - builds `hrr-playback` as a new `hrr` subproject and `core-hrr` artifact, and packages it as `amdrocm-hrr`; - maps `projects/hrr` to an artifact, so a rocm-systems PR that touches only HRR builds one stage instead of all of them. With rocm-systems' `.github/scripts/compute_build_stages.py --changed-projects projects/hrr`, current `main` logs `Unrecognized project(s) ['hrr'] -> build all stages`. With this branch it selects `compiler-runtime` only and skips Windows (`windows_amdgpu_families=none`), because `core-hrr` is Linux-only; - adds a testing-only `hrr-capture` build variant. Release `amdhip64` compiles capture out, so HRR's GPU tests need a separate build. The variant is the release build plus a `HIP_HRR_CAPTURE` flag; release builds and their artifacts are unchanged. This supersedes #7669 by @mangupta. His two commits are cherry-picked with his authorship, followed by the corrections below. JIRA ID : AIRUNTIME-2871 ## Technical Details **Build and packaging (from #7669, corrected)** - `core-hrr`: `hip-runtime` group, target-neutral, depends on `core-hip`, components `run` (`bin/hrr-playback`), `dev` (`include/hrr`), `dbg` and `doc`. - `core-hrr` is disabled on Windows (`disable_platforms = ["windows"]`). Nothing uses `hrr-playback` there: `amdrocm-hrr` is a Linux package, and no Windows build compiles capture into `amdhip64`. Windows builds and tarballs therefore have no `hrr-playback.exe`. - The subproject builds playback only (`HRR_BUILD_TESTS=OFF`, no catch2). The HRR test suites are not installable yet, and the integration suite needs a capture-enabled `amdhip64`. - Dropped from #7669: the `rocjitsu`/`mirage` `artifact_deps` change (unrelated), the `core-hiptests -> core-hrr` edge and the `hip-tests` runtime dependency on `hrr` (the hip-tests HRR smoke test does not run `hrr-playback`), and the empty `lib`/`test` components. - New `amdrocm-hrr` package (`run`, `dev`, `doc`; depends on `amdrocm-runtime`; MIT). `amdrocm-developer-tools` depends on it, as it does on `amdrocm-debugger`. Nothing is added to `amdrocm-runtime*`. - `install_rocm_from_artifacts.py --hrr` installs `core-hrr_run`. - Regenerated `test_tools/therock_consumer_graph.json` and added `core-hrr` to `build_tools/artifact_subprojects.json`. **`HIP_HRR_CAPTURE` flag and `hrr-capture` variant** - `FLAGS.cmake`: `HIP_HRR_CAPTURE`, default `OFF`, adds `HIP_HRR_CAPTURE_ENABLED=1` to `hip-clr`. That is the definition CLR's capture code checks today. `flags.md` recommends `ROCMBuildFlags.cmake` plus a private header for new C/C++ consumers. CLR has no such consumer yet, so this uses the scoped `CPP_DEFINES` path, as `HSA_HOTSWAP_DEFAULT_ON` does for ROCR-Runtime. It can switch once CLR resolves the flag itself. - `CMakePresets.json`: `linux-release-hrr-capture` (Release plus the flag). - `amdgpu_family_matrix.py`: an `hrr-capture` Linux variant with suffix `hrr-capture` (artifact group `multi-arch-hrr-capture`). The name avoids `san`, which `select_build_runner` routes to sanitizer runners. It is listed for gfx90a, gfx94x, gfx950, gfx103x, gfx110x, gfx1150, gfx1151, gfx1153, gfx120x and gfx125x. - No workflow selects the variant yet. `configure_multi_arch_ci.py` needs no special case: families are filtered by `build_variants`, and tests follow the normal trigger rules. **CI path handling** TheRock's CI code has no HRR-specific path rules. rocm-systems lists its HRR workflows, and the HRR paths TheRock never reads (`projects/hrr/tests`, `tools`, `skills`), in its own `.github/skip-ci-config.toml` ([ROCm/rocm-systems#13256](ROCm/rocm-systems#13256)). `configure_multi_arch_ci.py` already reads that file, so a rocm-systems PR that changes only those paths skips TheRock CI. An earlier revision of this PR changed other things that are no longer here: - it aliased those workflows to `projects/hrr` in `configure_external_repo_ci.py`; - it passed external-repo changed files into stage reuse and narrowed the rebuilt stages. The stage-reuse part changes reuse for every external repo, and it needs its own design review. It made no difference for HRR-only changes, because rocm-systems already limits those to `compiler-runtime` through `build_stages`. **Dependency added** - License: MIT (SPDX headers in `projects/hrr`). - Owner: `@ROCm/hrr-reviewers`, the code owners of `projects/hrr` in rocm-systems. - Build impact: host-only C++ with no device code. Built locally, playback is 5 C++ objects that compile in about 4 s with `-j4`; `hrr-playback` is 0.9 MB (0.7 MB stripped). **Follow-ups, not in this PR** - A `core-hrrtests` artifact once `projects/hrr` installs its test binaries and configs so their paths resolve from the installed tree. - A `test_hrr` component in `fetch_test_configurations.py`, following `miopen-dbsync`: `linux_cpu_runner: True` and `test_types`, running the suite under rocJITsu. That uses CPU runners, so it needs no GPU-runner capacity plan. - A caller for the `hrr-capture` variant, which nothing selects yet. It would request `build_variant: hrr-capture` with `build_stages` limited to `compiler-runtime`. Release builds stay unchanged, and that request does not rebuild `math-libs`. - Until those land, rocm-systems HRR CI keeps building its own capture-enabled `amdhip64` on top of TheRock release artifacts. ## Test Plan - Unit tests: `build_tools/tests/install_rocm_from_artifacts_test.py`, `build_tools/github_actions/tests/configure_multi_arch_ci_test.py`, `amdgpu_family_matrix_test.py` and `check_consumer_graph_drift_test.py`. New tests: - `--hrr` installs `core-hrr_run`; - `hrr-capture` selects only Linux families that list it, uses its preset and suffix, and keeps presubmit test runners; - every listed variant is defined, and every variant preset exists in `CMakePresets.json`. Removing the new preset makes this test fail. - `build_tools/tests/configure_external_repo_ci_test.py`: a `therock*` workflow change on rocm-systems runs all tests. That test was added in an earlier revision and is kept, because no other test reaches that branch. - All of `build_tools/tests`, `build_tools/github_actions/tests` and `test_tools/tests`. - `fetch_sources.py`, then a full configure (`-DTHEROCK_ENABLE_ALL=ON -DTHEROCK_AMDGPU_FAMILIES=gfx94X-dcgpu`) in the pinned manylinux build image, then `check_consumer_graph_drift.py --check`. - `generate_subproject_manifest.py --verify` and `build_tools/tests/configure_stage_test.py` in the same image, on this branch and on `main`. - `compute_build_stages.py --changed-projects projects/hrr` from rocm-systems against `main` and this branch. - `pre-commit` on the changed files; bandit with `build_tools/scan_tools/bandit.yml`. ## Test Result - Unit tests from the earlier local run (install, family matrix, consumer graph, and `configure_multi_arch_ci` as it was then): 205 passed, 1 skipped (an existing skip). - At the head of this branch, `build_tools/tests`, `build_tools/github_actions/tests` and `test_tools/tests`: 2372 passed, 17 skipped. Two files were left out because they also fail on `main` on the test machine (host ASAN runtime and stub compilation): `exe_stub_gen_test.py` and `configure_asan_env_test.py`. - Configure succeeds and reports `HIP_HRR_CAPTURE = OFF` and `CORE_HRR` enabled. The consumer graph is up to date. - The committed `artifact_subprojects.json` is already stale on `main`: it predates several artifacts, and the generator now writes a different layout. This PR adds only the `core-hrr` entry rather than regenerating the whole file. As a result, `configure_stage_test.py` has the same single failure (`test_artifact_subprojects_matches_cmake`) on this branch as on `main`, with 15 passed. CI skips that test when submodules are not fetched. - `compute_build_stages.py`: all stages on `main`. On this branch, `compiler-runtime` only, with `windows_amdgpu_families=none`. - The configure and consumer-graph results above predate the Linux-only change. The Linux graph does not depend on Windows availability, and this PR's consumer graph drift check covers it. - `pre-commit` and bandit: clean. - Not run locally: a full TheRock build of the `hrr` subproject, or a build with the `hrr-capture` preset. TheRock CI covers the first. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/CONTRIBUTING.md. --------- Signed-off-by: Daniel Olsson <daniel.olsson@amd.com> Co-authored-by: Maneesh Gupta <maneesh.gupta@amd.com> Co-authored-by: Cursor Agent <cursoragent@cursor.com> Co-authored-by: Claude Sonnet 5.5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
projects/hrr.projects/hrrchanges and make GPU integration tests explicit opt-in.Validation
hrr_api_args.hincludes; no stale old playback paths; no conflict markers or pycache artifacts.hrr-playbackbuild passed locally with HIP compiler setup.Unit_HRR_CaptureSmokepassed locally.Pending / follow-up
amdhip64build needs a matching ROCm/HSA toolchain; the local attempt hit an unrelated HSA header mismatch.Related
🤖 Generated with Claude Code