Skip to content

fix(hrr): record the code object HIP loads, not the whole fat binary - #13698

Open
andraghetti wants to merge 10 commits into
developfrom
users/andraghetti/hrr-fatbin-code-object-only
Open

andraghetti wants to merge 10 commits into
developfrom
users/andraghetti/hrr-fatbin-code-object-only

Conversation

@andraghetti

@andraghetti andraghetti commented Oct 18, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Capture stored every registered fat binary as a blob. It took the blob's size from the bundle header: the furthest offset + size of the entries, or the compressed totalSize. HIP receives only a pointer, so nothing bounds that value. #12260 added header checks, but an entry declaring offset = 0, size = 1 TiB passed all of them. Capture then read a terabyte from a pointer whose real extent is unknown. HIP itself reads only the header and the code object it picks for this GPU. mangupta raised this in the review of #12260.

Replay already resolves kernels and globals in the code objects the archive holds. So this change stops storing bundles instead of bounding their size.

Technical Details

  • Capture no longer reads the registered bundle. compute_bundle_size() and its header checks are removed.
  • The __hipRegisterFatBinary event is still written, with blob size zero. Its blob hash names the code object capture recorded for it, or is zero. The event layout does not change.
  • Capture records the code object HIP extracted from the fat binary for the current device, or device 0 before one is current. It reads the span through FatBinaryInfo::GetProgram() and amd::Program::binary(). That span's size is the one the bundle entry declares. The code object goes through the same hash_program_image() path that kernel launches and module loads use, which now hashes any ELF only as far as its headers reach, and only within the memory that can be read. It probes the ELF header, then the header tables, then the extent they give, page by page: a one-byte write() into a pipe on Linux, which fails with EFAULT on a page that cannot be read, and VirtualQuery on Windows. An ELF whose headers reach further is not recorded.
  • Fat binaries registered before capture starts are swept after the __device__ globals. Resolving a global makes HIP load the code object that holds it, and the fat binary sweep then records that code object.
  • StatCO::ForEachFatBinaryBlob becomes ForEachFatBinary. Its callback also receives the FatBinaryInfo*, which is null until HIP digests the fat binary.
  • Replay keeps one module per code object and device, keyed <hash>@<device>, because HIP builds a loaded code object for one device. Before it replays any event, it loads each code object on every device a registration naming it, a module load or a kernel launch used it on. Kernel and __device__ global lookups search only the modules on the current device. Archives that still carry fat binary blobs replay through the existing handler, which keys its module the same way. A code object that fails to load on one of its devices aborts the replay, as any pre-load failure does, unless --continue-on-error is given. Every pass starts on device 0, so the timed pass after a --kernel-filter warm-up does not start on the device the warm-up ended on. The reader decodes every launch id capture writes (_spt, cooperative, hipDrvLaunchKernelEx, hipLaunchKernelExC), so each launch counts as a use on its device. Context push/pop and graph kernel nodes do not count yet: AIRUNTIME-3028.
  • hrr_run_roundtrip() now requires a blob only when it also requires a device-to-host check. On a GPU without image support, the texture workload copies nothing back. Its fat binary was the only blob in its archive.
  • Unit_HRR_FatBinHeaderChecks is rewritten. It checks that no register event carries a blob. It checks that each bundle HIP loaded has its code object in the archive as a bare ELF, found by kernel name. It checks that no archive file holds a bundle, or bytes from a bundle HIP never loaded. A bundle whose host entry declares 1 TiB is registered both before and after HIP initialises. Both captures must succeed. So must one of a code object whose entry declares 1 TiB and whose unloaded section does too, in front of a page that cannot be read, on Linux and on Windows; that code object is not stored. The archive is then rewritten as one from before this change, the code object as a bundle blob and its launch with no code object hash, and must replay.
  • Unit_HRR_FatBinLoadDevice needs two GPUs of one kind, and skips otherwise. It registers a fat binary with device 1 current and launches its kernel on device 0. It and its workload are listed in tests/config/integration.yaml. It also writes a __device__ global on device 0 and reads it with a kernel on both devices. Replay must pass every device-to-host check, also with --kernel-filter hrr_fatbin, whose warm-up pass ends on device 1.
  • Unit_HRR_FatBinGlobalCopy registers a fat binary with a __device__ global before HIP initialises, copies to the global and back, and launches nothing. The archive must hold the global's code object, and replay must bind the global.
  • Unit_HRR_Format_EveryLaunchIdDecodes checks that the reader decodes each launch id.
  • PRIVACY.md no longer lists fat binary bundles in blobs/.
  • DESIGN.md describes the new behaviour.

The Windows build was not run.

Issue Tracking

JIRA ID: AIRUNTIME-2858

Test Plan

  • HRR build with capture compiled into CLR, unit and integration ctest, then every Catch2 case one at a time, hidden cases included, on MI350X (gfx950), one GPU. The gate ran on head c0cee0c473, the merge of develop 08c6e8f7c9, and every case has the same status as develop's own run at that commit (26 hidden cases fail on both). Unit_HRR_FatBinHeaderChecks, Unit_HRR_FatBinLoadDevice, Unit_HRR_FatBinGlobalCopy and Unit_HRR_CodeObjectScrub also ran on c0cee0c with two MI350X visible, and pass. Unit_HRR_CodeObjectScrub fails with the scrub taken out of write_code_object, for the fat binary's code object as well as the module's. hrr-handler-tests ran on 59ca9e736f; no handler changed since.
  • Negative control: the new tests with the old capture code, hip_code_object.* and DESIGN.md put back, same gate. It ran on an earlier head of this branch, which differs from 07b0f274ee only in hrr_run_roundtrip(), which the new test does not use.
  • Negative controls on 56d2270, two GPUs: without the readable-memory bound, the capture child of Unit_HRR_FatBinHeaderChecks segfaults (exit 139). With global lookup on any device, the two reads of Unit_HRR_FatBinLoadDevice fail their device-to-host checks. With the old-archive handler keying its module without a device, the rewritten archive fails: kernel hrr_fatbin_live_probe not found.
  • Negative controls on the code of 7136c395a9, two GPUs, run before Unit_HRR_FatBinGlobalCopy took its final form: with the reader decoding only the four old launch ids, Unit_HRR_Format_EveryLaunchIdDecodes fails at hipLaunchKernel_spt. With each pass starting where the previous one ended, the --kernel-filter replay of Unit_HRR_FatBinLoadDevice fails. With readable_extent() trusting the declared size, the capture child of Unit_HRR_FatBinHeaderChecks segfaults (exit 139). On 7136c39 itself, with fat binaries swept before globals, Unit_HRR_FatBinGlobalCopy fails: no code object holds the global, and replay reports it is in no module.
  • Baseline: develop at 5b4535b7ba, same runner and same scripts.

Test Result

Check This PR (7136c39) develop 5b4535b
Build, unit ctest pass pass
Integration ctest pass pass
Unit Catch2 cases 76 of 76 pass 75 of 75 pass
Visible integration cases 125 pass, 5 skip 124 pass, 4 skip
Hidden integration cases 122 of 155 pass 122 of 153 pass

Every case has the same status as on develop. The new skips are Unit_HRR_FatBinLoadDevice and its workload, which need two GPUs, and the Unit_HRR_FatBinGlobalCopy_Direct workload, which needs its driver's bundle. With two GPUs visible, Unit_HRR_FatBinLoadDevice passes. The first head of this branch failed Unit_HRR_TextureRoundtrip, because its archive held no blob on MI350X. The hrr_run_roundtrip() change above fixes that. The 33 hidden cases that do not pass are mostly _Direct workloads that need their driver's setup. They include FirstMallocRoundtrip, which also fails on develop.

Submission Checklist

Capture stored every registered fat binary as a blob whose size came
from the bundle header: the furthest offset + size of its entries, or
the compressed totalSize. HIP is handed only a pointer, so nothing
bounds that value, and an entry declaring offset 0, size 1 TiB made
capture read far past the image, much further than HIP itself, which
reads the header and the code object it picks for this GPU.

Stop storing bundles. The register event is still written, without a
blob, and capture records the code object HIP extracted for the
current device, with the size COMGR gave it, through the same writer
the launch path uses. Fat binaries registered before capture started
are swept after the __device__ globals, since resolving a global is
what makes HIP load its code object. Replay pre-loads every code
object, so kernels and globals resolve as before; older archives with
blobs still replay. The header checks in compute_bundle_size go with
it.

The fat binary integration test now checks that no register event
carries a blob, that each loaded bundle's code object is in the
archive as a bare ELF, and that a bundle with a terabyte host entry is
captured. The roundtrip helper now asks for a blob only when it asks
for a D2H check: on a GPU without image support the texture workload
copies nothing back, and its fat binary was the only blob it had.

Jira: AIRUNTIME-2858
@andraghetti
andraghetti requested review from a team as code owners October 18, 2026 13:06
Copilot AI balanced review requested due to automatic review settings October 18, 2026 13:06
@therock-pr-bot

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Untrusted selected-entry lengths can still cause out-of-bounds reads, and device-specific preloading can break heterogeneous-GPU replay.

2 open issues
What changed in this PR

Replaces whole-fat-binary capture with HIP-selected code objects, avoiding reads of unused bundle entries while preserving registration events.

Changes:

  • Records selected code objects and sweeps globals before fat binaries.
  • Expands malformed-bundle tests and permits blob-free roundtrips when appropriate.
  • Updates capture documentation.
File Description
projects/​hrr/​tests/​integration/​hrr_fatbin_header_test.cc Checks standalone ELF capture and oversized unused entries.
projects/​hrr/​tests/​integration/​CMakeLists.txt Updates test-description comments.
projects/​hrr/​tests/​include/​hrr_test_common.hh Makes blob requirements conditional.
projects/​hrr/​DESIGN.md Documents code-object capture.
projects/​clr/​hipamd/​src/​hrr/​hip_capture.cpp Replaces bundle recording and reorders initialization sweeps.
projects/​clr/​hipamd/​src/​hip_code_object.hpp Updates the fat-binary iterator interface.
projects/​clr/​hipamd/​src/​hip_code_object.cpp Passes fat-binary information to callbacks.

Review effort: Balanced

Comment thread projects/clr/hipamd/src/hrr/hip_capture.cpp Outdated
Comment thread projects/clr/hipamd/src/hrr/hip_capture.cpp Outdated
…ers reach

HIP keeps the size the bundle entry declares for the code object it picks,
and nothing bounds it, so a GPU entry declaring a terabyte registered fine
and capture then hashed past the image. Capture now takes the extent from
the ELF's own section and program headers, checked against the declared
size, and records nothing for an image that is not an ELF.

The fat binary test registers such a bundle in front of an unreadable page.
…stered on

Replay pre-loads every code object on the current device before it replays
any event, so a code object a fat binary registration loaded on device 1
went to device 0, and on a host with GPUs of more than one kind replay
aborted. The register event now names its code object (blob hash set, blob
size 0), and pre-load puts it on the device current on the registering
thread at that event, from the hipSetDevice events before it.

A test registers a fat binary with device 1 current and checks where replay
loads its code object.

@mangupta mangupta left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🤖 AI-assisted pre-review (for @mangupta)

Comment thread projects/hrr/playback/hrr_playback.cpp Outdated
Pre-load put a code object a fat binary registration named only on the
registering device, and replay kept one module per code object, so a kernel
from it launched on another device got the module of the first. On GPUs of
one kind a library registered with device 1 current and launched on device 0
uses the same code object on both. Replay now keeps a module per code object
and device, and pre-loads each code object on every device a registration, a
module load or a kernel launch used it on.

The fat binary load-device test now launches the kernel on device 0 after
registering with device 1 current, and replays it.
@mangupta

Copy link
Copy Markdown
Contributor

🤖 AI-assisted pre-review (for @mangupta)

Re-reviewed the updated head (5b08ef8d43b3) — no new findings at or above the posting threshold.

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 Changes recommended

Unresolved image-bound safety, replay lookup regressions, and incompatible-GPU test assumptions need correction.

4 open issues
2 resolved since last review

Review effort: Balanced

Comment thread projects/clr/hipamd/src/hrr/hip_capture.cpp Outdated
Comment thread projects/hrr/playback/hip_playback.cpp
Comment thread projects/hrr/playback/hip_playback.cpp Outdated
Comment thread projects/hrr/tests/integration/hrr_fatbin_header_test.cc

@danielbolsson danielbolsson left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Review — needs work

Full review of the whole change (839470366c…5b08ef8d43, four commits), read as a single context rather than a fan-out.
Description names AIRUNTIME-2858 (not retrievable from here); the description was read after the findings were settled.

Blocking

  • Legacy-blob archives lose their name-scan fallback for kernels — projects/hrr/playback/hip_playback.cpp:653-657, :1600
    load_module now keys co_modules as <hash>@<dev>, and the name-only fallback in resolve_kernel_function skips every key that doesn't end in the current device suffix. playback___hipRegisterFatBinary (still the handler for pre-change archives, which carry a bundle blob) inserts with the bare hex key (:1600), so those modules are never searched. A pre-change archive with a launch that has no recorded co_hash (the case that fallback exists for) now ends in "Kernel not found in any loaded module". The PR description and DESIGN.md both say old archives "replay through the existing handler"; nothing tests it. (module_map-based lookup and the symbol scan at :1623 still see them, so this is specific to the kernel scan.)
    Confidence 85 · logic · Fix: key the blob module with the same device suffix it was loaded on (or accept unsuffixed keys in the scan), and add a replay test over an archive with a bundle blob.

  • The description misstates the change — PR description, Technical Details
    It says "Replay does not change." Replay changes substantially: code_object_devices and per-device pre-load in hrr_playback.cpp:863-957, device-suffixed module/function keys in hip_playback.cpp, a new two-GPU test (Unit_HRR_FatBinLoadDevice), and hash_program_image now truncates every ELF it hashes to its header-bounded size (hip_capture.cpp:955-969), which applies to module-load and launch hashing too, not only fat binaries. Test Plan/Result are also pinned to head 07b0f274ee, the first of four commits; the later two commits (per-device pre-load) are not covered by the stated results. Remedy: rewrite the description to the current four-commit scope and re-run the gate on the current head.
    Confidence 85 · mechanical · Fix: update the description; the code is not what's wrong here.

  • Failing checks on the head — therock-pr-bot, Multi-Arch CI Summary, Windows::release / Test gfx110X-all / Test Sanity Check
    All three report FAILURE. I did not establish whether they are related (the Windows job looks like a sanity/infra failure), and no HRR test job appears in the rollup at all, so the HRR behaviour is only evidenced by the author's local run on the old head. The author should confirm they are unrelated or green them.
    Confidence 70 · architectural

Non-blocking

  • Stale comment on the register handler — projects/hrr/playback/hip_playback.cpp:1528-1533 (and :1558)
    It says the handler loads the blob and stores it "keyed by the full 32-char hex hash … consistent with load_module()". load_module is no longer keyed that way, and for new-format events (blob_size == 0) the handler returns immediately and the pre-load does the work. Say so where the early return is.
    Confidence 85 · mechanical

  • A failed device switch is reported as a load on that device — projects/hrr/playback/hrr_playback.cpp:949-956
    If hipSetDevice(dev) fails, moved is false but load_module still runs on start_dev, the log line has already said "on device N", and the result is cached under @start_dev — so the module the later launch on dev needs is never loaded and nothing says so.
    Confidence 70 · logic · Fix: skip/report when the switch fails.

Decisions for the author

  • Pre-load failure on a secondary device is fatal — tradeoff
    A code object that fails to load on one of its recorded devices aborts the replay (without --continue-on-error), while the comment and DESIGN.md say that on mixed-GPU hosts it "may load on no other". Either that case is expected and should be non-fatal, or it is a real failure and the wording should not suggest otherwise.
  • Two-GPU test is skipped on single-GPU runners — tradeoff
    Unit_HRR_FatBinLoadDevice returns HRR_SKIP below two GPUs, so the per-device pre-load is exercised only on multi-GPU hardware. Confirm that is where this gets run.

Positive signals

  • Using the size HIP/COMGR already holds, bounded by the ELF's own headers, instead of trusting the bundle header is the right fix for the 1 TiB entry, and the guard-page test makes it provable.
  • The negative-control run (new tests against old capture code segfaulting) is exactly the evidence a test change like this needs.

Deployment notes

  • Capture output format changes: __hipRegisterFatBinary events now carry a code-object hash with blob_size == 0 and no bundle blob; older replay builds will not resolve kernels from such archives via that event.

What this covered

Every file in the change read completely, plus the callers and cases around the changed symbols in hip_playback.cpp/hrr_playback.cpp, elf.cpp::getElfSize and hip_platform.cpp::__hipRegisterFatBinary. Nothing was built or run; findings are from reading, and CI status was taken from the check rollup.

Did not run: fan-out (single context — degraded, coverage is lower than a split review), history, prior-changes, code-comments and agent-instruction passes; the ticket AIRUNTIME-2858 was not retrieved. Not reconciled against any discussion here.

…ules to their device

An ELF's own section table can declare as much as the bundle entry does, so a
section of 1 TiB passed getElfSize() against the entry's 1 TiB and capture
hashed past the image. Capture now finds how much of the span can be read,
page by page with process_vm_readv, and refuses an ELF whose headers reach
past that.

Replay keeps a module per code object and device, but resolved a __device__
global in a module of any device, and the handler for archives that still
carry a bundle blob keyed its module without a device, so the kernel name
scan skipped it. Both now use the current device.

The fat binary header test registers a code object whose unloaded section
declares 1 TiB, launches a kernel, and replays the archive rewritten as an
old one: the code object as a bundle blob and a launch without a code object
hash. The two-GPU test writes a __device__ global on device 0 and reads it on
both devices, and both of its halves skip on GPUs of two kinds.
@mangupta

Copy link
Copy Markdown
Contributor

🤖 AI-assisted pre-review (for @mangupta)

Re-reviewed the updated head (56d2270dd3a1) — no new findings at or above the posting threshold.

… the two-GPU cases

A code object is pre-loaded only on the devices capture used it on, and a
failure there is fatal like any pre-load failure; the comments and DESIGN.md
read as if it were expected. Say why the device switch cannot fail. List
Unit_HRR_FatBinLoadDevice and its workload in integration.yaml, which must
name every case the binary registers.
@mangupta

Copy link
Copy Markdown
Contributor

🤖 AI-assisted pre-review (for @mangupta)

Re-reviewed the updated head (4485b419b248) — no new findings at or above the posting threshold.

@danielbolsson danielbolsson left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Re-checked at 4485b41: the old-archive handler now keys its module per device and the rewritten-archive replay covers it, the readable-extent probe bounds the overstated-section case, global lookup is per device, and the description now matches the replay changes.

Nits:

  • projects/clr/hipamd/src/hrr/hip_capture.cpp:959 - a non-EFAULT process_vm_readv failure (e.g. EPERM under a strict seccomp profile) returns 0, so every ELF is dropped with a "does not fit the 0 bytes" warning; fall back to the declared-size bound or log the probe failure distinctly.
  • projects/hrr/tests/integration/hrr_fatbin_header_test.cc:592 - the device-0 pre-load from the launch event is not observable: without it resolve_kernel_function lazily loads by co_hash on device 0 and replay still passes; assert the pre-load log line or soften the comment.

@andraghetti
andraghetti requested a balanced review from Copilot October 20, 2026 15:37

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Comment thread projects/clr/hipamd/src/hrr/hip_capture.cpp Outdated
Comment thread projects/hrr/playback/hrr_playback.cpp
Comment thread projects/hrr/playback/hrr_playback.cpp
Comment thread projects/hrr/playback/hrr_playback.cpp Outdated
…ss on device 0

Capture finds how much of a code object can be read with a pipe write()
on Linux, which no sandbox refuses, and with VirtualQuery on Windows,
probing the ELF header, then its header tables, then the extent they
give. The overstated-section case now has a guard page on Windows too.

Replay starts every pass on device 0, so a warm-up pass that ended on
device 1 no longer moves the timed pass; the reader decodes every launch
id capture writes, so each launch gives its code object's device. A
__device__ global copied with nothing launched replays (startup sweep
order). PRIVACY.md: fat binaries no longer go to blobs/ (AIRUNTIME-2858).

New cases: Unit_HRR_Format_EveryLaunchIdDecodes,
Unit_HRR_FatBinGlobalCopy, Unit_HRR_FatBinGlobalCopy_Direct.
@mangupta

Copy link
Copy Markdown
Contributor

🤖 AI-assisted pre-review (for @mangupta)

Re-reviewed the updated head (4ec383d0bd14) — no new findings at or above the posting threshold.

@danielbolsson danielbolsson left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Re-checked c0cee0c on top of my approval of 59ca9e7: readable_extent() probes with write() to a pipe on Linux and VirtualQuery on Windows, the ELF is probed header, tables, then extent, each replay pass starts on device 0, and the reader decodes the _spt/cooperative/Ex launches; the head since is a develop merge. Nothing blocking.

@mangupta

Copy link
Copy Markdown
Contributor

🤖 AI-assisted pre-review (for @mangupta)

Re-reviewed the updated head (c0cee0c47370) — no new findings at or above the posting threshold.

@andraghetti andraghetti reopened this Oct 21, 2026
@andraghetti andraghetti reopened this Oct 21, 2026
@andraghetti
andraghetti requested a review from a team October 22, 2026 11:46

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants