Repository navigation
feat(hrr): external region annotations for memory HRR never saw - #10225
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
I noticed the four new region tests are not in |
andraghetti
left a comment
andraghetti
left a comment
There was a problem hiding this comment.
My mistake, these are plain TEST_CASE with the tags in the source, so ctest finds them without a YAML entry. Only thing we lose is turning them off per platform, since disabled: [amd_windows] and the rest only work through the YAML. Fine for now, just something to remember if one of them ever needs skipping. Approving.
Something I'm not sure about: when an argument doesn't resolve to any allocation, it now goes to the kernel as the recorded address instead of null, and that part isn't behind a flag (--warn-untranslated-args only changes what gets printed). With annotations it never comes up, the segment gets materialised and the pointer translates. Without them, that old address could well be mapped in the replay process, since the VA layout tends to repeat between runs. Then the kernel writes into some unrelated buffer and we see nothing, instead of a null fault we'd spot straight away. Did you mean it for every archive, or would you rather keep null when there are no regions?
…ns' v5 Two independent branches were both claiming HRR_VERSION 5 for different wire formats. hrr-regions (#10225) derives hrr_api_id_t from HipDispatchTable member order, which shifts most API IDs but leaves struct layout alone. This branch adds DEREF_FIELDS, which grows event payloads for ~50 APIs but leaves IDs alone. HRR_MAGIC and HRR_VERSION are the only capture/replay cross-checks, so once both landed an archive from one would be silently accepted by the other's reader and every event past the first shifted ID would dispatch to the wrong handler. #10225 carries v5 to develop. This branch takes v6 and picks up the ID derivation when it rebases onto develop after #10225 lands. Only the version constant and its documentation change; the generated shims are untouched. Note that the checked-in generated files are already stale relative to gen_hrr_api_args.py (a develop merge added APIs without a regenerate), so hrr_api_args.h is edited to exactly what the generator emits for this block rather than by re-running it, which would have mixed ~1500 unrelated lines into this commit. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
|
Review follow-up in HIGH — API IDs. HIGH — MEDIUM — Minors. Framing test checks all three records; OOB assertion matches |
…y device Review follow-up on #10225. Embedded pointers took a different path from whole-pointer arguments: a struct field pointing into a segment that bypassed HIP returned early on translation failure instead of materialising the segment, so a framework handing its kernels a struct of tensor bases got none of the benefit of annotating its allocator. Both paths now materialise on the same evidence, which is still an address the producer declared, so a mis-flagged scalar is left alone as before. The region record's device ordinal was parsed and then dropped. The live sets now carry it and materialisation allocates on that device rather than on whichever one happened to be current when the first pointer missed, which put the buffer on the wrong GPU as soon as the recorded program used more than one. Lookup stays keyed on the VA alone: sidecars come from a single recorded process, so an address identifies its allocation without the ordinal. An untranslated kernel-argument pointer goes back to reaching the GPU as null, as on develop. Forwarding the capture-time address risked hitting a live buffer, since the driver tends to reproduce a VA layout across runs, turning a fault into silent corruption. The count and warning behind --warn-untranslated-args stay. Also fixes a build break from ce9ccbb: Unit_HRR_Recovery_ImplausiblePayloadLength uses HIP_TEST_CASE, which needs an hrr.yaml entry it did not have, so HrrTest failed to link and ctest silently ran a stale v4 binary. Co-authored-by: Cursor <cursoragent@cursor.com>
|
Review follow-up in @andraghetti — untranslated args. You were right, and I have reverted it to Worth being clear that an untranslated pointer is a defect either way, which is why it should be loud. Three things produce one: HRR missed an allocation API (a capture bug), the allocation bypassed HIP entirely (not a bug — the case these annotations exist to cover), or the capture-side value heuristic mis-flagged a scalar (harmless). You cannot tell which from the address, so faulting at first use is the right default. The count and one-time warning behind @LeoZhao-BJ — both comments fixed, replies inline. Embedded pointers now materialise bypassed segments like whole-pointer arguments do, with a regression test that I verified fails without the fix; and the region One unrelated thing this shook out: |
…y device Review follow-up on #10225. Embedded pointers took a different path from whole-pointer arguments: a struct field pointing into a segment that bypassed HIP returned early on translation failure instead of materialising the segment, so a framework handing its kernels a struct of tensor bases got none of the benefit of annotating its allocator. Both paths now materialise on the same evidence, which is still an address the producer declared, so a mis-flagged scalar is left alone as before. The region record's device ordinal was parsed and then dropped. The live sets now carry it and materialisation allocates on that device rather than on whichever one happened to be current when the first pointer missed, which put the buffer on the wrong GPU as soon as the recorded program used more than one. Lookup stays keyed on the VA alone: sidecars come from a single recorded process, so an address identifies its allocation without the ordinal. An untranslated kernel-argument pointer goes back to reaching the GPU as null, as on develop. Forwarding the capture-time address risked hitting a live buffer, since the driver tends to reproduce a VA layout across runs, turning a fault into silent corruption. The count and warning behind --warn-untranslated-args stay. Also fixes a build break from ce9ccbb: Unit_HRR_Recovery_ImplausiblePayloadLength uses HIP_TEST_CASE, which needs an hrr.yaml entry it did not have, so HrrTest failed to link and ctest silently ran a stale v4 binary. Co-authored-by: Cursor <cursoragent@cursor.com>
…ns' v5 Two independent branches were both claiming HRR_VERSION 5 for different wire formats. hrr-regions (#10225) derives hrr_api_id_t from HipDispatchTable member order, which shifts most API IDs but leaves struct layout alone. This branch adds DEREF_FIELDS, which grows event payloads for ~50 APIs but leaves IDs alone. HRR_MAGIC and HRR_VERSION are the only capture/replay cross-checks, so once both landed an archive from one would be silently accepted by the other's reader and every event past the first shifted ID would dispatch to the wrong handler. derivation when it rebases onto develop after #10225 lands. Only the version constant and its documentation change; the generated shims are untouched. Note that the checked-in generated files are already stale relative to gen_hrr_api_args.py (a develop merge added APIs without a regenerate), so hrr_api_args.h is edited to exactly what the generator emits for this block rather than by re-running it, which would have mixed ~1500 unrelated lines into this commit. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
e209404 to
3c4a6be
Compare
|
since 2 issues found were tracked by other 2 PRs, approve it here. |
Reset events_failed (and the new code-object counter) after the kernel-filter warm-up pass so --continue-on-error does not double-count. Skip pid-* dirs that have no events.bin instead of treating them as repair failures. Fold unloadable pre-load modules into the replay result. Print "unknown" for build type when $<CONFIG> is empty, and wrap the loaded-library report so Windows does not compile dl_iterate_phdr. Co-authored-by: Cursor <cursoragent@cursor.com>
…ux/wsl Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com>
…ersion' into users/danielbolsson/hrr-api-ids
… into users/danielbolsson/hrr-regions
There was a problem hiding this comment.
🟡 Changes recommended
There are a few concrete correctness/robustness bugs in the new region producer and guard-block budgeting logic that should be fixed before approval.
Once you've addressed the issues Copilot identified, you can request another Copilot review.
Pull request overview
This PR adds an external “region annotations” channel to HRR replay so that memory ranges the HIP dispatch table never observed (sub-allocations inside a segment and allocations that bypass HIP entirely) can be described by an out-of-process producer and consumed by hrr-playback for pointer classification and optional guard-page diagnostics. It also registers/extends unit coverage for torn/implausible record recovery and the regions feature.
Changes:
- Add region sidecar format (
*.hrrr), producer documentation + reference PyTorch producer, and replay-side ingestion/classification/materialization. - Add replay flags for region reporting/strictness, untranslated kernel-arg reporting, and guard modes (
--guard-segments,--guard-blocks+ tuning flags). - Add and register Catch2 unit tests for region sidecars and implausible payload-length torn-tail recovery.
File summaries
| File | Description |
|---|---|
| projects/hip-tests/catch/unit/hrr/hrr_regions_test.cc | New unit tests covering region stream framing and replay behavior with synthetic sidecars. |
| projects/hip-tests/catch/unit/hrr/hrr_recovery_test.cc | Adds recovery test for implausible record size handling. |
| projects/hip-tests/catch/unit/hrr/CMakeLists.txt | Registers the new regions unit test source in the HRR unit test target. |
| projects/hip-tests/catch/config/configs/unit/hrr.yaml | Registers previously-unregistered HRR tests so they run under ctest. |
| projects/clr/hipamd/src/hrr/README.md | Documents new playback flags and introduces external region annotations at a high level. |
| projects/clr/hipamd/src/hrr/producers/README.md | Defines the producer-side contract and sidecar wire format + durability rules. |
| projects/clr/hipamd/src/hrr/producers/pytorch/hrr_torch_regions.py | Reference producer that emits regions for PyTorch’s HIP caching allocator. |
| projects/clr/hipamd/src/hrr/playback/hrr_region_map.h | Declares replay-side region timeline/map API (load/advance/classify/materialize). |
| projects/clr/hipamd/src/hrr/playback/hrr_region_map.cpp | Implements merging/sorting of sidecars, timeline cursor, and lazy segment materialization. |
| projects/clr/hipamd/src/hrr/playback/hrr_reader.h | Exposes shared record-stream primitives (framing + sidecar discovery). |
| projects/clr/hipamd/src/hrr/playback/hrr_reader.cpp | Implements shared framing/torn-tail recovery and region-stream discovery. |
| projects/clr/hipamd/src/hrr/playback/hrr_playback.cpp | Loads/prints region info and advances region cursor per event dispatch. |
| projects/clr/hipamd/src/hrr/playback/hip_playback.h | Extends PlaybackContext with region/guard state and new helper APIs. |
| projects/clr/hipamd/src/hrr/playback/hip_playback.cpp | Implements region pointer checks, segment materialization, and guard-page mechanics. |
| projects/clr/hipamd/src/hrr/playback/CMakeLists.txt | Adds hrr_region_map.cpp to the hrr-playback build. |
| projects/clr/hipamd/src/hrr/hrr_regions.h | Defines the region record/batch wire format and helpers used by producers/readers. |
| projects/clr/hipamd/src/hrr/DESIGN.md | Adds design/architecture documentation for external region annotations and guard modes. |
Review details
- Files reviewed: 17/17 changed files
- Comments generated: 3
- Review effort level: Lite
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Rebases the --version work onto the standalone projects/hrr layout from PR #11157. * playback/CMakeLists.txt: hand-port the version-stamp block, since #11157 rewrote the CLR file below git's rename threshold (D+A, not R) so the edits could not follow automatically. * tests/unit/hrr_playback_test.cc: port to the standalone harness -- HIP_TEST_CASE -> HRR_TEST_CASE, hip::SpawnProc -> hrr::test::SpawnProc, hip_test_{common,process}.hh -> hrr_test_{common,process}.hh, "hrr_api_args.h" -> "hrr/hrr_api_args.h". * hrr.yaml: entry dropped rather than re-expressed. Catch2 auto-discovers the case from the already-registered source; per team direction only the CaptureSmoke sentinel stays in hip-tests.
Brings in the HRR standalone-project migration (#11157) via #10221, which moves the playback tool, generated sources and API-args header out of projects/clr/hipamd/src/hrr/ into projects/hrr/. Conflicts were confined to the three generated outputs (hip_capture_generated.cpp, hrr_api_args.h, hip_playback_generated.cpp) and were resolved by re-running projects/hrr/tools/gen_hrr_api_args.py rather than merging them by hand; the generator resolves its output paths relative to its own location, so a single run emits all three at the post-migration destinations. The capture shim stays at projects/clr/hipamd/src/hrr/hip_capture_generated.cpp by design: it compiles into amdhip64.
Follow the HRR relocation from projects/clr/hipamd/src/hrr and
projects/hip-tests/catch/unit/hrr into projects/hrr.
- hrr_regions.h -> projects/hrr/include/hrr/, consumers now include
"hrr/hrr_regions.h" via the hrr::api_args interface target
- hrr_region_map.{cpp,h} -> projects/hrr/playback/, registered in the
hrr-playback target
- producers/ (README + hrr_torch_regions.py) -> projects/hrr/producers/
- hrr_regions_test.cc -> projects/hrr/tests/integration/, ported to the
standalone harness (hrr_test_common.hh / hrr_test_process.hh,
hrr::test::SpawnProc, HRR_HIP_CHECK) and registered in
HRR_GPU_WORKLOAD_SOURCES
- hrr::test::SpawnProc gains capture_stderr, needed by the playback
subprocess helper
- generated files regenerated with tools/gen_hrr_api_args.py
hip-tests keeps only the frozen capture smoke sentinel; hrr.yaml and
catch/unit/hrr/CMakeLists.txt are untouched by this branch.
The standalone harness in projects/hrr/tests/hrr_test_common.hh defines HRR_TEST_CASE; HIP_TEST_CASE only exists in the hip-tests harness and does not compile here.
…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.
…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.
… 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.
The comment on the reserved fast-path in generate_build_table claimed the slot holds "the nullptr from the real table". That is only true for a retired void* slot. reserved is also set when a typedef cannot be found or cannot be parsed, and in those cases the dispatch table member is a live function pointer, so skipping the override keeps a working pass-through rather than a nullptr. Behaviour is unchanged and the generated output is byte-identical; only the generator-side comment was wrong. The equivalent comment in generate_args_struct already says "retired/unparsable".
#10221 has merged, so this PR's base auto-retargeted from its branch to develop. Because #10221 was squash-merged, its commits on this branch have no ancestry in common with develop's copy, which rendered its seven files as changes here. Merging develop settles that: the diff is back to the generator and its three generated outputs. #10306 has also merged in the meantime and arrives with the same merge.
… into users/danielbolsson/hrr-regions Carries the parent branch down the stack. It brings develop with it, now containing the merged #10221 and #10306, plus the review fix on #10632 that corrects the reserved-slot comment in gen_hrr_api_args.py. No conflicts. Re-running the generator over the merged tree produces an empty diff, so the checked-in artifacts remain generator output.
Picks up five commits, including #10454, which landed the driver-memcpy D2D ordering fix in the same test files this stack works on. No conflicts here; the generator reproduces the checked-in artifacts unchanged.
… into users/danielbolsson/hrr-regions Carries develop down the stack from the parent branch, picking up #10454. No conflicts; the generator reproduces the checked-in artifacts unchanged.
… table gen_hrr_api_args.py built playback argument casts from p.raw_type, which is the dispatch-table (hip_api_trace.hpp) signature. Capture is right to use it — capture calls through the table. Playback is not: it calls the public API by name, a different declaration that is free to diverge. hip_api_trace.hpp is append-only ABI; hip_runtime_api.h may re-declare an API with a new signature under a HIP_FORCE_API_VERSION guard. That divergence is already present. hipMemAdvise takes `int device` in the table, but at HIP_FORCE_API_VERSION >= 800 the public API takes a hipMemLocation struct, so the generated `(int)a->device` does not match the function it is passed to when built against HIP 8.x. Fix the class, not the instance: * parse_public_api_signatures() reads hip_runtime_api.h and records each non-template declaration with its parameter types and whether it sits inside a version guard. Template overloads are skipped — they coexist with the C declaration rather than replacing it, so they never change which function a C-typed call resolves to. Simple typedef aliases are resolved, so hipDeviceProp_tR0600 compares equal to hipDeviceProp_t. * validate_playback_signatures() requires that every generated playback call has a public declaration that both matches the table signature and is unconditionally visible. Generation fails otherwise, naming the API, both signatures, their line numbers and guard status. Today it clears 193 APIs and flags exactly one. * _PLAYBACK_ARG_BRIDGES is the explicit, reviewed resolution for an API that legitimately diverges. hipMemAdvise rebuilds a hipMemLocation from the recorded int device ordinal under the same guard the public header uses, so the recorded device survives rather than being replaced by a zeroed struct. Verified against the in-tree headers: the pre-fix call fails to compile on HIP 8.x and the new one compiles on both 7.x and 8.x, with no change on 7.x. The generator is idempotent — re-running produces no further diff. Regenerated outputs are committed alongside. hip_capture_generated.cpp is byte-identical; the hrr_api_args.h hunk is a stale comment the previous generator would also have rewritten (hipChooseDevice's table parameter is spelled hipDeviceProp_tR0600), not a struct-layout change. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…/hrr-regions Propagates ef76664 (derive playback casts from the public API, not the dispatch table) down the HRR stack, one branch at a time so no child's diff is ever computed against a base that lacks it. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… into users/danielbolsson/hrr-regions
The pytorch producer opens regions/pytorch.hrrr for append and wrote the 8-byte file header on every open. DESIGN.md explicitly supports a process re-opening its own writer on resume, so that second header lands in the middle of the record stream. A reader parses those bytes as an hrr_event_header, reads event_type 0x5248 (the file magic) instead of 0xFFFD, takes a garbage payload length, and discards every record after that point as a torn tail. The capture looks healthy and silently loses everything written after the reopen. Write the header only when the file is empty.
TheRock Submodule Bump ActivityNewest first
|
Motivation
HRR interposes the HIP dispatch table, so it records the memory that crosses a HIP API and nothing else. Two common cases leave replay with a device VA range it cannot account for:
hipMalloc(PyTorch's HIP caching allocator)Pointers into those ranges resolve nowhere at replay, so the capture is unusable.
A producer outside the runtime writes the ranges down as
pid-<pid>/regions/<name>.hrrr, andhrr-playbackloads them automatically.Technical Details
Fidelity. With annotations present and no guard flag, replay's memory layout is exactly what it would have been without them. The annotations are read, not acted on, except that a segment HIP never saw now gets allocated, so pointers into it resolve instead of reaching the GPU as an address from another process. Fidelity therefore only increases.
The two
--guard-*flags are the deliberate exception: they move memory so an out-of-bounds access faults instead of landing in a live neighbour, and are off by default for that reason.--guard-blocksrestores every guarded block and releases the relocation before the next event, so the divergence is confined to the launch under examination.RegionMapis rewound after kernel-filter warm-up.advance_tois monotonic, so the timed pass would otherwise classify against an empty live-set, giving--regions-strictphantom failures and leaving--guard-blocksguarding nothing.Includes the reference PyTorch producer and the sidecar format documentation.
This PR has been rebased into a stack
It previously carried 6,225 lines of generated API-ID churn that had nothing to do with regions. That is now its own PR and this one is zero generated lines.
AIRUNTIME-2699), which renumbers the API IDs and freezes v5.--continue-on-error.Region tests were never registered
While splitting this out I found that none of the five
Unit_HRR_Regions_*cases had anhrr.yamlentry, so ctest never registered them and this feature was shipping its own tests unrun. This answers @andraghetti's question on this PR, which my earlier reply only partly addressed — that reply coveredUnit_HRR_Recovery_ImplausiblePayloadLength, a different test.They use plain
TEST_CASErather thanHIP_TEST_CASE, so unlike the recovery test nothing failed to compile and the gap stayed silent.check_sources.pydid flag all five, but it is a warning rather than an error unlessTHEROCK_REQUIRE_HIP_YAML_ENTRIESis set.Issue Tracking
AIRUNTIME-2658
Test Plan
Full HRR Catch2 unit suite, GPU-free and GPU passes, on gfx950.
Because this PR is a rebase rather than new work, the stronger check is that the split lost nothing: the tree at this tip must differ from the pre-split branch only by the API-ID regeneration, #10221's content, and the develop delta.
Test Result
Unit_HRR_MemcpyPeerRoundtrip, needs a second device).[hrr]-tagged region tests now actually run. Verified individually:Unit_HRR_Regions_StreamFraming,_Roundtrip,_SegmentMaterializationand_EmbeddedPointerMaterializationall pass.hipInitDevice.Submission Checklist